From b1d280ef3be009ecc79610d5ae57da11c1a9a28a Mon Sep 17 00:00:00 2001 From: Anuj Jalota Date: Wed, 30 Sep 2026 23:09:45 +0530 Subject: [PATCH 01/11] cuda: fuse FFN gate/up and GLU in MMQ with interleaved gate/up rows For prompt processing, run the gate and up matmuls of a dense FFN as one mul_mat_q launch and apply the GLU in the write-back. - The loaders take the gate weights as a second pointer and fill the weight tile with up and gate rows interleaved in chunks of 8, so each thread holds an up value and the matching gate value (rows r and r + 8 of an MMA tile). - The MMA loop is unchanged: same tile, warp layout, vec_dot and one sum[] per thread. The fused write-back computes the GLU from registers and writes half the rows. - The fused path uses the regular tiling (no stream-k) with the column tiles innermost so that neighbouring blocks share weight tiles. - Enabled for Q4_K, Q5_K, Q6_K, Q8_0 and IQ4_XS on NVIDIA Turing and newer. --- ggml/src/ggml-cuda/ggml-cuda.cu | 63 ++++++++ ggml/src/ggml-cuda/mmq-load-tiles.cuh | 156 +++++++++++------- ggml/src/ggml-cuda/mmq.cu | 40 +++-- ggml/src/ggml-cuda/mmq.cuh | 225 ++++++++++++++++++++------ ggml/src/ggml-cuda/unary.cuh | 15 ++ tests/test-backend-ops.cpp | 12 ++ 6 files changed, 388 insertions(+), 123 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index f6ff9f893aee..698faad3e65b 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1972,6 +1972,44 @@ static bool ggml_cuda_match_shared_expert(const ggml_cgraph * graph, int routed_ return true; } +// Buffer-independent, so graph_optimize can call it before allocation. +static bool ggml_cuda_mul_mat_q_fusion_matches(const ggml_tensor * up, const ggml_tensor * gate, const ggml_tensor * glu, const int device) { + const ggml_tensor * x_up = up->src[0]; + const ggml_tensor * y = up->src[1]; + + if (up->op != GGML_OP_MUL_MAT || !ggml_cuda_mmq_fusion_supported(x_up->type) || y->type != GGML_TYPE_F32 || + ggml_get_op_params_i32(up, 1) != GGML_HINT_NONE || ggml_get_op_params_i32(gate, 1) != GGML_HINT_NONE) { + return false; + } + + if (!ggml_cuda_should_fuse_mul_mat(up, gate, glu)) { + return false; + } + + // The fused write-back applies SWIGLU_OAI with the default alpha and limit of ggml_cuda_op_swiglu_oai_single. + if (ggml_get_glu_op(glu) == GGML_GLU_OP_SWIGLU_OAI && (ggml_get_op_params_f32(glu, 2) != 1.702f || ggml_get_op_params_f32(glu, 3) != 7.0f)) { + return false; + } + + // The fused launch has the row tiles (64 output rows each) in gridDim.y and channels * samples in gridDim.z. + if ((x_up->ne[1] + 63)/64 > 65535 || y->ne[2]*y->ne[3] > 65535) { + return false; + } + + // NVIDIA only: the fused write-back relies on the MMA accumulator layout. + const int cc = ggml_cuda_info().devices[device].cc; + return turing_mma_available(cc) && !ggml_cuda_should_use_mmvq(x_up->type, cc, y->ne[1]) && ggml_cuda_should_use_mmq(x_up->type, cc, y->ne[1], 0); +} + +static bool ggml_cuda_should_fuse_mul_mat_q(const ggml_tensor * up, const ggml_tensor * gate, const ggml_tensor * glu) { + // The padding of a compute-buffer weight that is a view cannot be cleared, see ggml_cuda_mul_mat. + const auto bad_padding_clear = [](const ggml_tensor * x) { + return ggml_backend_buffer_get_usage(x->buffer) == GGML_BACKEND_BUFFER_USAGE_COMPUTE && + ggml_nbytes(x) != ggml_backend_buffer_get_alloc_size(x->buffer, x) && x->view_src; + }; + return ggml_cuda_mul_mat_q_fusion_matches(up, gate, glu, ggml_cuda_get_device()) && !bad_padding_clear(up->src[0]) && !bad_padding_clear(gate->src[0]); +} + static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { GGML_TENSOR_BINARY_OP_LOCALS @@ -4349,6 +4387,18 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph fused_node_count = 3; break; } + + if (ggml_cuda_should_fuse_mul_mat_q(up, gate, glu)) { + ggml_cuda_mm_fusion_args_host fusion_data{}; + fusion_data.gate = gate->src[0]; + fusion_data.glu_op = ggml_get_glu_op(glu); + fusion_data.glu_limit = ggml_get_op_params_f32(glu, 3); + + ggml_cuda_mul_mat_q(*cuda_ctx, src0, src1, ids, glu, &fusion_data); + fused_mul_mat_vec = true; + fused_node_count = 3; + break; + } } } @@ -4934,6 +4984,19 @@ static void ggml_backend_cuda_graph_optimize(ggml_backend_t backend, ggml_cgraph } } for (int i = 0; i < cgraph->n_nodes; ++i) { + // Keep the inputs of the fused gate/up/GLU MMQ kernel alive until the GLU so that they cannot alias its output (under op offload the weights are buffer copies). + if (i + 2 < cgraph->n_nodes && cgraph->nodes[i + 2]->op == GGML_OP_GLU) { + ggml_tensor * glu = cgraph->nodes[i + 2]; + ggml_tensor * gate = glu->src[0]; + ggml_tensor * up = glu->src[1]; + const bool ok = (gate == cgraph->nodes[i] && up == cgraph->nodes[i + 1]) || (gate == cgraph->nodes[i + 1] && up == cgraph->nodes[i]); + if (ok && ggml_cuda_mul_mat_q_fusion_matches(up, gate, glu, cuda_ctx->device)) { + add_alloc_deps(i, i + 2); + i += 2; + continue; + } + } + ggml_cuda_moe_weighted_reduction_match match; if (ggml_cuda_match_moe_weighted_reduction(cgraph, i, match)) { params->add_alloc_dep(params->user_data, const_cast(match.experts), match.dst); diff --git a/ggml/src/ggml-cuda/mmq-load-tiles.cuh b/ggml/src/ggml-cuda/mmq-load-tiles.cuh index b37d88c16b51..66e3e58df5b7 100644 --- a/ggml/src/ggml-cuda/mmq-load-tiles.cuh +++ b/ggml/src/ggml-cuda/mmq-load-tiles.cuh @@ -5,7 +5,9 @@ #include "mmq.cuh" template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q1_0( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + GGML_UNUSED(x_gate); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -116,7 +118,9 @@ template static __device__ __forceinline_ } template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q2_0( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + GGML_UNUSED(x_gate); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -205,7 +209,9 @@ template static __device__ __forceinline_ } template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q4_0( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + GGML_UNUSED(x_gate); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -268,7 +274,9 @@ template static __device__ __forceinline_ } template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q4_1( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + GGML_UNUSED(x_gate); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -331,7 +339,9 @@ template static __device__ __forceinline_ } template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q5_0( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + GGML_UNUSED(x_gate); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -411,7 +421,9 @@ template static __device__ __forceinline_ } template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q5_1( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + GGML_UNUSED(x_gate); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -488,8 +500,8 @@ template static __device__ __forceinline_ } } -template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q8_0( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { +template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q8_0( + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -515,11 +527,11 @@ template static __device__ __forceinline_ for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); - if (fallback) { + if (fallback && !has_fusion) { i = min(i, i_max); } - const block_q8_0 * bxi = (const block_q8_0 *) x + kbx0 + i*stride + kbx; + const block_q8_0 * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride) + kbx; #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) x_qs[i*sram_stride + 0 + txi] = get_int_b2(bxi[0].qs, kqsx); @@ -538,11 +550,11 @@ template static __device__ __forceinline_ for (int i0 = 0; i0 < I; i0 += nwarps * rows_per_warp) { int i = i0 + threadIdx.y * rows_per_warp + threadIdx.x / blocks_per_tile_x_row; - if (fallback) { + if (fallback && !has_fusion) { i = min(i, i_max); } - const block_q8_0 * bxi = (const block_q8_0 *) x + kbx0 + i*stride + kbxd; + const block_q8_0 * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride) + kbxd; #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) x_df[i*sram_stride + kbxd] = bxi->d; @@ -555,7 +567,9 @@ template static __device__ __forceinline_ // --------------------------------------------------------------------------------------------- template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q2_K( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + GGML_UNUSED(x_gate); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -616,7 +630,9 @@ template static __device__ __forceinline_ } template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q3_K( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + GGML_UNUSED(x_gate); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -728,8 +744,8 @@ static __device__ __forceinline__ int unpack_scales_q45_K(const int * scales, co ((scales[ksc/2] >> (2 * (ksc % 2))) & 0x30303030); // upper 2 bits } -template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q4_K( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { +template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q4_K( + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -753,11 +769,11 @@ template static __device__ __forceinline_ for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); - if (fallback) { + if (fallback && !has_fusion) { i = min(i, i_max); } - const block_q4_K * bxi = (const block_q4_K *) x + kbx0 + i*stride; + const block_q4_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); const int qs0 = get_int_b4(bxi->qs, txi); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) @@ -782,11 +798,11 @@ template static __device__ __forceinline_ int i = (i0 + threadIdx.y*rows_per_warp + threadIdx.x/2) % I; { #endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) - if (fallback) { + if (fallback && !has_fusion) { i = min(i, i_max); } - const block_q4_K * bxi = (const block_q4_K *) x + kbx0 + i*stride; + const block_q4_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); const int * scales = (const int *) bxi->scales; const int ksc = threadIdx.x % 2; @@ -810,11 +826,11 @@ template static __device__ __forceinline_ for (int i0 = 0; i0 < I; i0 += nwarps*warp_size) { int i = (i0 + threadIdx.y*warp_size + threadIdx.x) % I; - if (fallback) { + if (fallback && !has_fusion) { i = min(i, i_max); } - const block_q4_K * bxi = (const block_q4_K *) x + kbx0 + i*stride; + const block_q4_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); x_dm[i] = bxi->dm; } @@ -823,11 +839,11 @@ template static __device__ __forceinline_ for (int i0 = 0; i0 < I; i0 += nwarps*rows_per_warp) { int i = (i0 + threadIdx.y*rows_per_warp + threadIdx.x/(MMQ_TILE_NE_K/8)) % I; - if (fallback) { + if (fallback && !has_fusion) { i = min(i, i_max); } - const block_q4_K * bxi = (const block_q4_K *) x + kbx0 + i*stride + (threadIdx.x % (MMQ_TILE_NE_K/8)) / (QI4_K/8); + const block_q4_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride) + (threadIdx.x % (MMQ_TILE_NE_K/8)) / (QI4_K/8); const int * scales = (const int *) bxi->scales; @@ -839,8 +855,8 @@ template static __device__ __forceinline_ #endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) } -template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q5_K( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { +template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q5_K( + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -864,11 +880,11 @@ template static __device__ __forceinline_ for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); - if (fallback) { + if (fallback && !has_fusion) { i = min(i, i_max); } - const block_q5_K * bxi = (const block_q5_K *) x + kbx0 + i*stride; + const block_q5_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); const int ky = QR5_K*txi; const int ql = get_int_b4(bxi->qs, txi); @@ -905,11 +921,11 @@ template static __device__ __forceinline_ int i = (i0 + threadIdx.y*rows_per_warp + threadIdx.x/2) % I; { #endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) - if (fallback) { + if (fallback && !has_fusion) { i = min(i, i_max); } - const block_q5_K * bxi = (const block_q5_K *) x + kbx0 + i*stride; + const block_q5_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); const int * scales = (const int *) bxi->scales; const int ksc = threadIdx.x % 2; @@ -933,11 +949,11 @@ template static __device__ __forceinline_ for (int i0 = 0; i0 < I; i0 += nwarps*warp_size) { int i = (i0 + threadIdx.y*warp_size + threadIdx.x) % I; - if (fallback) { + if (fallback && !has_fusion) { i = min(i, i_max); } - const block_q5_K * bxi = (const block_q5_K *) x + kbx0 + i*stride; + const block_q5_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); x_dm[i] = bxi->dm; } @@ -947,11 +963,11 @@ template static __device__ __forceinline_ for (int i0 = 0; i0 < I; i0 += nwarps*rows_per_warp) { int i = (i0 + threadIdx.y*rows_per_warp + threadIdx.x/(MMQ_TILE_NE_K/8)) % I; - if (fallback) { + if (fallback && !has_fusion) { i = min(i, i_max); } - const block_q5_K * bxi = (const block_q5_K *) x + kbx0 + i*stride; + const block_q5_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); const int * scales = (const int *) bxi->scales; @@ -963,8 +979,8 @@ template static __device__ __forceinline_ #endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) } -template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q6_K( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { +template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q6_K( + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -989,11 +1005,11 @@ template static __device__ __forceinline_ for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); - if (fallback) { + if (fallback && !has_fusion) { i = min(i, i_max); } - const block_q6_K * bxi = (const block_q6_K *) x + kbx0 + i*stride; + const block_q6_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); const int ql = get_int_b2(bxi->ql, txi); const int ql0 = (ql >> 0) & 0x0F0F0F0F; @@ -1019,11 +1035,11 @@ template static __device__ __forceinline_ for (int i0 = 0; i0 < I; i0 += nwarps*warp_size) { int i = (i0 + threadIdx.y*warp_size + threadIdx.x) % I; - if (fallback) { + if (fallback && !has_fusion) { i = min(i, i_max); } - const block_q6_K * bxi = (const block_q6_K *) x + kbx0 + i*stride; + const block_q6_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) x_df[i*sram_stride] = bxi->d; @@ -1037,11 +1053,11 @@ template static __device__ __forceinline_ for (int i0 = 0; i0 < I; i0 += nwarps*rows_per_warp) { int i = (i0 + threadIdx.y*rows_per_warp + threadIdx.x/(MMQ_TILE_NE_K/8)) % I; - if (fallback) { + if (fallback && !has_fusion) { i = min(i, i_max); } - const block_q6_K * bxi = (const block_q6_K *) x + kbx0 + i*stride + (threadIdx.x % (MMQ_TILE_NE_K/8)) / 4; + const block_q6_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride) + (threadIdx.x % (MMQ_TILE_NE_K/8)) / 4; #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) x_sc[i*sram_stride + threadIdx.x%4] = get_int_b2(bxi->scales, threadIdx.x % (MMQ_TILE_NE_K/8)); @@ -1054,7 +1070,9 @@ template static __device__ __forceinline_ // --------------------------------------------------------------------------------------------- template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq1_s( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + GGML_UNUSED(x_gate); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -1116,7 +1134,9 @@ template static __device__ __forceinline_ } template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq2_xxs( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + GGML_UNUSED(x_gate); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -1180,7 +1200,9 @@ template static __device__ __forceinline_ } template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq2_xs( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + GGML_UNUSED(x_gate); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -1245,7 +1267,9 @@ template static __device__ __forceinline_ } template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq2_s( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + GGML_UNUSED(x_gate); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -1313,7 +1337,9 @@ template static __device__ __forceinline_ } template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq3_xxs( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + GGML_UNUSED(x_gate); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -1377,7 +1403,9 @@ template static __device__ __forceinline_ } template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq3_s( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + GGML_UNUSED(x_gate); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -1445,8 +1473,8 @@ template static __device__ __forceinline_ } } -template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq4_xs( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { +template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq4_xs( + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -1469,11 +1497,11 @@ template static __device__ __forceinline_ for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); - if (fallback) { + if (fallback && !has_fusion) { i = min(i, i_max); } - const block_iq4_xs * bxi = (const block_iq4_xs *) x + kbx0 + i*stride; + const block_iq4_xs * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); const int aux_q4 = get_int_b4(bxi->qs, kqsx); const int2 v = get_int_from_table_16(aux_q4, kvalues_iq4nl); @@ -1493,11 +1521,11 @@ template static __device__ __forceinline_ for (int i0 = 0; i0 < I; i0 += nwarps * rows_per_warp) { int i = i0 + threadIdx.y * rows_per_warp + threadIdx.x / (MMQ_TILE_NE_K/4); - if (fallback) { + if (fallback && !has_fusion) { i = min(i, i_max); } - const block_iq4_xs * bxi = (const block_iq4_xs *) x + kbx0 + i*stride; + const block_iq4_xs * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); const float d = __half2float(bxi->d); @@ -1513,7 +1541,9 @@ template static __device__ __forceinline_ } template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq4_nl( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + GGML_UNUSED(x_gate); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -1582,7 +1612,9 @@ template static __device__ __forceinline_ // --------------------------------------------------------------------------------------------- template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_mxfp4( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + GGML_UNUSED(x_gate); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); @@ -1649,7 +1681,9 @@ template static __device__ __forceinline_ } template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_mxfp4_fp4( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + GGML_UNUSED(x_gate); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q4) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q4); @@ -1691,7 +1725,9 @@ template static __device__ __forceinline_ } template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_nvfp4( - const char * __restrict__ x, int * __restrict__ x_tile, const int kb0, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kb0, const int i_max, const int stride) { + GGML_UNUSED(x_gate); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, prec_src1) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, prec_src1); @@ -1747,7 +1783,9 @@ template static __de } template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_nvfp4_nvfp4( - const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + GGML_UNUSED(x_gate); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q4) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q4); diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 46efdfe0f73e..56c7777945f6 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -132,11 +132,19 @@ static ggml_prec ggml_cuda_mmq_get_prec_src1(const ggml_tensor * src0, const ggm } void ggml_cuda_mul_mat_q( - ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * ids, ggml_tensor * dst) { + ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * ids, ggml_tensor * dst, + const ggml_cuda_mm_fusion_args_host * fusion) { GGML_ASSERT( src1->type == GGML_TYPE_F32); GGML_ASSERT( dst->type == GGML_TYPE_F32); GGML_ASSERT(!ids || ids->type == GGML_TYPE_I32); // Optional, used for batched GGML_MUL_MAT_ID. + // Fused gate/up/GLU: src0 is the up weight, dst the GLU output. + const ggml_tensor * gate = fusion ? fusion->gate : nullptr; + if (gate) { + GGML_ASSERT(!ids && !fusion->x_bias && !fusion->gate_bias && !fusion->x_scale && !fusion->gate_scale); + GGML_ASSERT(gate->type == src0->type && ggml_are_same_shape(gate, src0) && ggml_are_same_stride(gate, src0)); + } + GGML_TENSOR_BINARY_OP_LOCALS; cudaStream_t stream = ctx.stream(); @@ -158,14 +166,16 @@ void ggml_cuda_mul_mat_q( const float * src1_d = (const float *) src1->data; float * dst_d = (float *) dst->data; - // If src0 is a temporary compute buffer, clear any potential padding. - if (ggml_backend_buffer_get_usage(src0->buffer) == GGML_BACKEND_BUFFER_USAGE_COMPUTE) { - const size_t size_data = ggml_nbytes(src0); - const size_t size_alloc = ggml_backend_buffer_get_alloc_size(src0->buffer, src0); - if (size_alloc > size_data) { - GGML_ASSERT(ggml_is_contiguously_allocated(src0)); - GGML_ASSERT(!src0->view_src); - CUDA_CHECK(cudaMemsetAsync((char *) src0->data + size_data, 0, size_alloc - size_data, stream)); + // If src0 (or the fused gate weight) is a temporary compute buffer, clear any potential padding. + for (const ggml_tensor * t : { src0, gate }) { + if (t && ggml_backend_buffer_get_usage(t->buffer) == GGML_BACKEND_BUFFER_USAGE_COMPUTE) { + const size_t size_data = ggml_nbytes(t); + const size_t size_alloc = ggml_backend_buffer_get_alloc_size(t->buffer, t); + if (size_alloc > size_data) { + GGML_ASSERT(ggml_is_contiguously_allocated(t)); + GGML_ASSERT(!t->view_src); + CUDA_CHECK(cudaMemsetAsync((char *) t->data + size_data, 0, size_alloc - size_data, stream)); + } } } @@ -180,7 +190,8 @@ void ggml_cuda_mul_mat_q( const bool fallback = ggml_cuda_mmq_needs_fallback(ne01); - const ggml_prec prec_src1 = ggml_cuda_mmq_get_prec_src1(src0, dst, cc); + // With fusion dst is the GLU output, not a MUL_MAT; the fused types always quantize src1 to Q8_1. + const ggml_prec prec_src1 = gate ? GGML_PREC_Q8 : ggml_cuda_mmq_get_prec_src1(src0, dst, cc); const bool use_native_fp4 = prec_src1 == GGML_PREC_Q4; const size_t y_block_size = use_native_fp4 ? sizeof(block_fp4_mmq) : sizeof(block_q8_1_mmq); @@ -262,13 +273,20 @@ void ggml_cuda_mul_mat_q( ne11 * ne10_padded * sizeof(block_q8_1) / (QK8_1 * sizeof(int)); const int64_t s13 = ne12*s12; + ggml_cuda_mm_fusion_args_device fusion_device = {}; + if (gate) { + fusion_device.gate = gate->data; + fusion_device.glu_op = fusion->glu_op; + fusion_device.glu_limit = fusion->glu_limit; + } + const mmq_args args = { src0_d, src0->type, (const int *) src1_q8_1.ptr, nullptr, nullptr, dst_d, src0->type == GGML_TYPE_NVFP4 && use_native_fp4 ? src1_scale.ptr : nullptr, ne00, ne01, ne1, s01, ne11, s1, ne02, ne12, s02, s12, s2, ne03, ne13, s03, s13, s3, - ne1, J_best}; + ne1, J_best, fusion_device}; ggml_cuda_mul_mat_q_switch_type(ctx, args, stream, prec_src1); return; } diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index e008a34c08a0..c2b794798ab8 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -1,6 +1,7 @@ #pragma once #include "common.cuh" +#include "unary.cuh" #include #include @@ -10,7 +11,8 @@ #define MMQ_ITER_K_FP4 512 #define MMQ_NWARPS 8 -typedef void (*ggml_cuda_mmq_load_tiles_t)(const char * __restrict__ x, int * x_tile, const int kbx0, const int i_max, const int stride); +typedef void (*ggml_cuda_mmq_load_tiles_t)(const char * __restrict__ x, const char * __restrict__ x_gate, int * x_tile, + const int kbx0, const int i_max, const int stride); typedef void (*ggml_cuda_mmq_vec_dot_t)(const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00); typedef void (*ggml_cuda_mmq_write_back_t)(const float * __restrict__ sum, const int32_t * __restrict__ get_rows_to_sorted, float * __restrict__ dst, const float * __restrict__ y_scale, const int stride, const int i_max, const int j_max); @@ -426,6 +428,32 @@ static __host__ int ggml_cuda_mmq_get_nbytes_shared_x(const ggml_cuda_mmq_config // ------------------------------------------------------------ +// Weight types with a fused gate/up/GLU path (see ggml_cuda_mmq_get_load_tiles); the other types run unfused. +static constexpr __host__ __device__ bool ggml_cuda_mmq_fusion_supported(const ggml_type type) { + return type == GGML_TYPE_Q4_K || type == GGML_TYPE_Q5_K || type == GGML_TYPE_Q6_K || type == GGML_TYPE_Q8_0 || type == GGML_TYPE_IQ4_XS; +} + +// Fused gate/up/GLU: the weight tile alternates chunks of this many up and gate rows, so rows r and r + 8 of an MMA tile (same thread) pair up. +static constexpr int GGML_CUDA_MMQ_FUSION_ROWS = 8; + +// Source of weight tile row i; with has_fusion the interleaved x/x_gate row, clamped here because i_max is in output rows. +template +static __device__ __forceinline__ const block_t * ggml_cuda_mmq_x_row( + const char * __restrict__ x, const char * __restrict__ x_gate, const int kbx0, int i, const int i_max, const int stride) { + if constexpr (has_fusion) { + constexpr int R = GGML_CUDA_MMQ_FUSION_ROWS; + const char * x_src = (i / R) % 2 == 0 ? x : x_gate; + i = (i / (2*R))*R + i % R; + if (fallback) { + i = min(i, i_max); + } + return (const block_t *) x_src + kbx0 + i*stride; + } else { + GGML_UNUSED_VARS(x_gate, i_max); + return (const block_t *) x + kbx0 + i*stride; + } +} + #include "mmq-load-tiles.cuh" #include "mmq-vec-dot.cuh" @@ -519,6 +547,49 @@ static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma( } } +// Fused gate/up/GLU write-back: accumulators l and l + 2 hold the up and gate values of the same output row and column. +template +static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma_glu( + const float * __restrict__ sum, const ggml_glu_op glu_op, const float glu_limit, + const int * __restrict__ ids_dst, float * __restrict__ dst, const int stride, const int i_max, const int j_max) { +#ifdef TURING_MMA_AVAILABLE + typedef tile<16, 8, int> tile_C; + static_assert(tile_C::I == 2*GGML_CUDA_MMQ_FUSION_ROWS, "each C tile must hold one chunk of up rows and one of gate rows"); + + constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback); + constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp. + + const int i0 = (threadIdx.y / ntx) * (ntx*tile_C::I); + +#pragma unroll + for (int j0 = 0; j0 < J; j0 += ntx*tile_C::J) { +#pragma unroll + for (int n = 0; n < ntx; ++n) { +#pragma unroll + for (int l = 0; l < tile_C::ne/2; ++l) { + const int j = j0 + (threadIdx.y % ntx) * tile_C::J + tile_C::get_j(l); + + if (j > j_max) { + continue; + } + + const int i = (i0 + n*tile_C::I)/2 + tile_C::get_i(l); + + if (fallback && i > i_max) { + continue; + } + + const int sum_index = (j0/tile_C::J + n)*tile_C::ne + l; + dst[ids_dst[j]*stride + i] = ggml_cuda_op_glu_single(glu_op, sum[sum_index + 2], sum[sum_index], glu_limit); + } + } + } +#else + GGML_UNUSED_VARS(sum, glu_op, glu_limit, ids_dst, dst, stride, i_max, j_max); + NO_DEVICE_CODE; +#endif // TURING_MMA_AVAILABLE +} + // ------------------------------------------------------------------------------------------------------------------------------------- // TODO remove this struct and use ggml_cuda_mmq_sram_layout instead. @@ -852,9 +923,22 @@ static constexpr __device__ int ggml_cuda_mmq_get_vdr() { return ggml_cuda_mmq_get_util_funcs().vdr; } -template +template static constexpr __device__ ggml_cuda_mmq_load_tiles_t ggml_cuda_mmq_get_load_tiles() { - return ggml_cuda_mmq_get_util_funcs().load_tiles; + if constexpr (!has_fusion) { + return ggml_cuda_mmq_get_util_funcs().load_tiles; + } else if constexpr (type == GGML_TYPE_Q8_0) { + return ggml_cuda_mmq_load_tiles_q8_0; + } else if constexpr (type == GGML_TYPE_Q4_K) { + return ggml_cuda_mmq_load_tiles_q4_K; + } else if constexpr (type == GGML_TYPE_Q5_K) { + return ggml_cuda_mmq_load_tiles_q5_K; + } else if constexpr (type == GGML_TYPE_Q6_K) { + return ggml_cuda_mmq_load_tiles_q6_K; + } else { + static_assert(type == GGML_TYPE_IQ4_XS, "no fused loader for this type"); + return ggml_cuda_mmq_load_tiles_iq4_xs; + } } template @@ -869,19 +953,21 @@ static constexpr __device__ ggml_cuda_mmq_write_back_t ggml_cuda_mmq_get_write_b // --------------------------------------------------------------------------------------------- -template +// has_fusion: x and x_gate are the up and gate weights; tile_x_max_i is then in output rows. +template static __device__ __forceinline__ void mul_mat_q_process_tile( - const char * __restrict__ x, const int offset_x, const int * __restrict__ y, + const char * __restrict__ x, const char * __restrict__ x_gate, const int offset_x, const int * __restrict__ y, const int * __restrict__ ids_dst, float * __restrict__ dst, float * __restrict__ tmp_fixup, - const float * __restrict__ y_scale, + const float * __restrict__ y_scale, const ggml_glu_op glu_op, const float glu_limit, const int stride_row_x, const int ncols_y, const int stride_col_dst, const int tile_x_max_i, const int tile_y_max_j, const int kb0_start, const int kb0_stop) { + static_assert(!(fixup && has_fusion), "the fused gate/up path does not split K"); constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, prec_src1) / warp_size; constexpr int qk = ggml_cuda_type_traits::qk; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, prec_src1); - constexpr ggml_cuda_mmq_load_tiles_t load_tiles = ggml_cuda_mmq_get_load_tiles(); + constexpr ggml_cuda_mmq_load_tiles_t load_tiles = ggml_cuda_mmq_get_load_tiles(); constexpr ggml_cuda_mmq_vec_dot_t vec_dot = ggml_cuda_mmq_get_vec_dot(); constexpr ggml_cuda_mmq_write_back_t write_back = ggml_cuda_mmq_get_write_back(); @@ -906,7 +992,7 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( constexpr int sz = sizeof(block_q8_1_mmq) / sizeof(int); for (int kb0 = kb0_start; kb0 < kb0_stop; kb0 += blocks_per_iter) { - load_tiles(x, tile_x, offset_x + kb0, tile_x_max_i, stride_row_x); + load_tiles(x, x_gate, tile_x, offset_x + kb0, tile_x_max_i, stride_row_x); { const int * by0 = y + ncols_y * (kb0 * qk / ne_block) * sz; #pragma unroll @@ -940,6 +1026,12 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( __syncthreads(); } + if constexpr (has_fusion) { + ggml_cuda_mmq_write_back_mma_glu(sum, glu_op, glu_limit, ids_dst, dst, stride_col_dst, tile_x_max_i, tile_y_max_j); + return; + } + GGML_UNUSED_VARS(glu_op, glu_limit); + if (fixup) { write_back(sum, ids_dst, tmp_fixup + blockIdx.x*(J*I), y_scale, I, I, J); } else { @@ -950,19 +1042,26 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( // The mul_mat_q kernel implements "stream-k" work partitioning as described in https://arxiv.org/abs/2301.03598 -template +template __launch_bounds__(ggml_cuda_mmq_get_nthreads(type, J, fallback, prec_src1), ggml_cuda_mmq_get_occupancy(type, J, fallback, prec_src1)) static __global__ void mul_mat_q( - const char * __restrict__ x, const int * __restrict__ y, const int32_t * __restrict__ ids_dst, + const char * __restrict__ x, const char * __restrict__ x_gate, const int * __restrict__ y, const int32_t * __restrict__ ids_dst, const int32_t * __restrict__ expert_bounds, float * __restrict__ dst, float * __restrict__ tmp_fixup, - const float * __restrict__ y_scale, + const float * __restrict__ y_scale, const ggml_glu_op glu_op, const float glu_limit, const uint3 blocks_per_ne00, const int nrows_x, const int ncols_dst, const int stride_row_x, const int ncols_y, const int stride_col_dst, const uint3 channel_ratio, const uint3 nchannels_y, const int stride_channel_x, const int stride_channel_y, const int stride_channel_dst, const uint3 sample_ratio, const uint3 nsamples_y, const int stride_sample_x, const int stride_sample_y, const int stride_sample_dst, const uint3 ntx) { + // The fused gate/up/GLU write-back is implemented for the NVIDIA MMA accumulator layout. +#ifdef TURING_MMA_AVAILABLE + constexpr bool fusion_available = true; +#else + constexpr bool fusion_available = false; +#endif // TURING_MMA_AVAILABLE + // Skip unused template specializations for faster compilation: - if (ggml_cuda_mmq_get_config(type, J, fallback, prec_src1).type == GGML_TYPE_COUNT) { + if (ggml_cuda_mmq_get_config(type, J, fallback, prec_src1).type == GGML_TYPE_COUNT || (has_fusion && !fusion_available)) { NO_DEVICE_CODE; return; } @@ -990,12 +1089,15 @@ static __global__ void mul_mat_q( } __syncthreads(); - if constexpr (!ggml_cuda_mmq_get_stream_k(type, J, fallback, prec_src1)) { + // The fused path reduces all of K per block (the GLU is nonlinear), writes I/2 output rows per tile, and has the column tiles innermost for weight reuse. + if constexpr (has_fusion || !ggml_cuda_mmq_get_stream_k(type, J, fallback, prec_src1)) { + constexpr int I_dst = has_fusion ? I/2 : I; + const uint2 tmp2 = fast_div_modulo(blockIdx.z, nchannels_y); const int wt = tmp2.x; const int zt = tmp2.y; - const int jt = blockIdx.y; - const int it = blockIdx.x; + const int jt = has_fusion ? blockIdx.x : blockIdx.y; + const int it = has_fusion ? blockIdx.y : blockIdx.x; // Defaults for regular matrix multiplication: int col_low = 0; @@ -1040,22 +1142,22 @@ static __global__ void mul_mat_q( } offset_y += (col_low + jt*J)*(sizeof(block_q8_1_mmq)/sizeof(int)); - offset_dst += it*I; + offset_dst += it*I_dst; const float * y_scale_tile = nullptr; if constexpr (type == GGML_TYPE_NVFP4) { offset_y_scale += col_low + jt*J; y_scale_tile = y_scale ? y_scale + offset_y_scale : nullptr; } - const int tile_x_max_i = nrows_x - it*I - 1; + const int tile_x_max_i = nrows_x - it*I_dst - 1; const int tile_y_max_j = col_diff - jt*J - 1; - const int offset_x = fastdiv(wt, sample_ratio)*stride_sample_x + fastdiv(zt, channel_ratio)*stride_channel_x + it*I*stride_row_x; + const int offset_x = fastdiv(wt, sample_ratio)*stride_sample_x + fastdiv(zt, channel_ratio)*stride_channel_x + it*I_dst*stride_row_x; constexpr bool fixup = false; - mul_mat_q_process_tile - (x, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, y_scale_tile, - stride_row_x, ncols_y, stride_col_dst, + mul_mat_q_process_tile + (x, x_gate, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, y_scale_tile, + glu_op, glu_limit, stride_row_x, ncols_y, stride_col_dst, tile_x_max_i, tile_y_max_j, 0, blocks_per_ne00.z); return; } @@ -1148,8 +1250,8 @@ static __global__ void mul_mat_q( constexpr bool fixup = false; // All but (potentially) the last iterations write their data to dst rather than the fixup buffer. mul_mat_q_process_tile - (x, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, y_scale_tile, - stride_row_x, ncols_y, stride_col_dst, + (x, x_gate, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, y_scale_tile, + glu_op, glu_limit, stride_row_x, ncols_y, stride_col_dst, tile_x_max_i, tile_y_max_j, kb0_start, kb0_stop); kbc += blocks_per_ne00.z; @@ -1232,8 +1334,8 @@ static __global__ void mul_mat_q( constexpr bool fixup = true; // Last index writes its data to fixup buffer to avoid data races with other blocks. mul_mat_q_process_tile - (x, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, y_scale_tile, - stride_row_x, ncols_y, stride_col_dst, + (x, x_gate, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, y_scale_tile, + glu_op, glu_limit, stride_row_x, ncols_y, stride_col_dst, tile_x_max_i, tile_y_max_j, kb0_start, kb0_stop); } @@ -1383,6 +1485,7 @@ struct mmq_args { int64_t nsamples_x; int64_t nsamples_y; int64_t stride_sample_x; int64_t stride_sample_y; int64_t stride_sample_dst; int64_t ncols_max; int J_best; // Tile width in ne11(dense)/ne12(MoE) direction to use for optimal performance. + ggml_cuda_mm_fusion_args_device fusion = {}; // gate != nullptr selects the fused gate/up/GLU path }; static size_t mmq_get_nbytes_shared(const ggml_cuda_mmq_config & config, const int cc) { @@ -1392,7 +1495,7 @@ static size_t mmq_get_nbytes_shared(const ggml_cuda_mmq_config & config, const i return nbs_ids + nbs_x + GGML_PAD(nbs_y, config.nthreads*sizeof(int)); } -template +template static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) { const int id = ggml_cuda_get_device(); const int cc = ggml_cuda_info().devices[id].cc; @@ -1406,8 +1509,8 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a const dim3 block_dims(warp_size, nwarps, 1); - CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); - CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); + CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); + CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); const int nty = (args.nrows_x + config.I - 1) / config.I; const int ntx = (args.ncols_max + config.J - 1) / config.J; @@ -1426,9 +1529,12 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a const uint3 channel_ratio_fd = init_fastdiv_values(channel_ratio); const uint3 sample_ratio_fd = init_fastdiv_values(sample_ratio); - if (!config.stream_k) { - mul_mat_q<<>> - (args.x, args.y, args.ids_dst, args.expert_bounds, args.dst, nullptr, args.y_scale, + if (has_fusion || !config.stream_k) { + const int nty_fusion = (args.nrows_x + config.I/2 - 1) / (config.I/2); + GGML_ASSERT(!has_fusion || nty_fusion <= 65535); + const dim3 block_nums = has_fusion ? dim3(ntx, nty_fusion, ntzw) : block_nums_xy_tiling; + mul_mat_q<<>> + (args.x, (const char *) args.fusion.gate, args.y, args.ids_dst, args.expert_bounds, args.dst, nullptr, args.y_scale, args.fusion.glu_op, args.fusion.glu_limit, blocks_per_ne00_fd, args.nrows_x, args.ncols_dst, args.stride_row_x, args.ncols_y, args.nrows_dst, channel_ratio_fd, nchannels_y_fd, args.stride_channel_x, args.stride_channel_y, args.stride_channel_dst, sample_ratio_fd, nsamples_y_fd, args.stride_sample_x, args.stride_sample_y, args.stride_sample_dst, @@ -1458,8 +1564,8 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a const dim3 block_nums_fixup(block_nums_stream_k.x, config.I/warp_size, 1); const dim3 block_dims_fixup(block_dims.x, block_dims.y/2, block_dims.z); - mul_mat_q<<>> - (args.x, args.y, args.ids_dst, args.expert_bounds, args.dst, tmp_fixup.ptr, args.y_scale, + mul_mat_q<<>> + (args.x, (const char *) args.fusion.gate, args.y, args.ids_dst, args.expert_bounds, args.dst, tmp_fixup.ptr, args.y_scale, args.fusion.glu_op, args.fusion.glu_limit, blocks_per_ne00_fd, args.nrows_x, args.ncols_dst, args.stride_row_x, args.ncols_y, args.nrows_dst, channel_ratio_fd, nchannels_y_fd, args.stride_channel_x, args.stride_channel_y, args.stride_channel_dst, sample_ratio_fd, nsamples_y_fd, args.stride_sample_x, args.stride_sample_y, args.stride_sample_dst, @@ -1476,56 +1582,56 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a ntx_fd); } -template +template void mul_mat_q_switch_J(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) { switch (args.J_best) { case 8: - launch_mul_mat_q(ctx, args, stream); + launch_mul_mat_q(ctx, args, stream); break; case 16: - launch_mul_mat_q(ctx, args, stream); + launch_mul_mat_q(ctx, args, stream); break; case 24: - launch_mul_mat_q(ctx, args, stream); + launch_mul_mat_q(ctx, args, stream); break; case 32: - launch_mul_mat_q(ctx, args, stream); + launch_mul_mat_q(ctx, args, stream); break; case 40: - launch_mul_mat_q(ctx, args, stream); + launch_mul_mat_q(ctx, args, stream); break; case 48: - launch_mul_mat_q(ctx, args, stream); + launch_mul_mat_q(ctx, args, stream); break; case 56: - launch_mul_mat_q(ctx, args, stream); + launch_mul_mat_q(ctx, args, stream); break; case 64: - launch_mul_mat_q(ctx, args, stream); + launch_mul_mat_q(ctx, args, stream); break; case 72: - launch_mul_mat_q(ctx, args, stream); + launch_mul_mat_q(ctx, args, stream); break; case 80: - launch_mul_mat_q(ctx, args, stream); + launch_mul_mat_q(ctx, args, stream); break; case 88: - launch_mul_mat_q(ctx, args, stream); + launch_mul_mat_q(ctx, args, stream); break; case 96: - launch_mul_mat_q(ctx, args, stream); + launch_mul_mat_q(ctx, args, stream); break; case 104: - launch_mul_mat_q(ctx, args, stream); + launch_mul_mat_q(ctx, args, stream); break; case 112: - launch_mul_mat_q(ctx, args, stream); + launch_mul_mat_q(ctx, args, stream); break; case 120: - launch_mul_mat_q(ctx, args, stream); + launch_mul_mat_q(ctx, args, stream); break; case 128: - launch_mul_mat_q(ctx, args, stream); + launch_mul_mat_q(ctx, args, stream); break; default: fprintf(stderr, "J_best=%d\n", args.J_best); @@ -1533,14 +1639,26 @@ void mul_mat_q_switch_J(ggml_backend_cuda_context & ctx, const mmq_args & args, } } +template +static void mul_mat_q_switch_fusion(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) { + if constexpr (ggml_cuda_mmq_fusion_supported(type)) { + if (args.fusion.gate) { + mul_mat_q_switch_J(ctx, args, stream); + return; + } + } + GGML_ASSERT(!args.fusion.gate); + mul_mat_q_switch_J(ctx, args, stream); +} + template void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) { if (ggml_cuda_mmq_needs_fallback(args.nrows_x)) { constexpr bool fallback = true; - mul_mat_q_switch_J(ctx, args, stream); + mul_mat_q_switch_fusion(ctx, args, stream); } else { constexpr bool fallback = false; - mul_mat_q_switch_J(ctx, args, stream); + mul_mat_q_switch_fusion(ctx, args, stream); } } @@ -1582,6 +1700,7 @@ extern DECL_MMQ_CASE_W4A4(GGML_TYPE_NVFP4); // ------------------------------------------------------------------------------------------------------------------------- void ggml_cuda_mul_mat_q( - ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * ids, ggml_tensor * dst); + ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * ids, ggml_tensor * dst, + const ggml_cuda_mm_fusion_args_host * fusion = nullptr); bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t n_experts); diff --git a/ggml/src/ggml-cuda/unary.cuh b/ggml/src/ggml-cuda/unary.cuh index 04f3af6443a7..c6ffdb9d63e7 100644 --- a/ggml/src/ggml-cuda/unary.cuh +++ b/ggml/src/ggml-cuda/unary.cuh @@ -121,3 +121,18 @@ __device__ __forceinline__ float ggml_cuda_op_swiglu_clamp_single(float gate, fl return ggml_cuda_op_silu_single(gate) * up; } + +__device__ __forceinline__ float ggml_cuda_op_glu_single(ggml_glu_op op, float gate, float up, float limit) { + switch (op) { + case GGML_GLU_OP_SWIGLU: + return up * ggml_cuda_op_silu_single(gate); + case GGML_GLU_OP_GEGLU: + return up * ggml_cuda_op_gelu_single(gate); + case GGML_GLU_OP_SWIGLU_OAI: + return ggml_cuda_op_swiglu_oai_single(gate, up); + case GGML_GLU_OP_SWIGLU_CLAMP: + return ggml_cuda_op_swiglu_clamp_single(gate, up, limit); + default: + return up * gate; + } +} diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index d3194bf47f52..9e1eb2113eb4 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -12034,6 +12034,18 @@ static std::vector> make_test_cases_eval() { false, 16, 8, false, false, true, false, { 1, 1 })); } + // Fused gate/up + GLU on the MMQ path: rows with bounds checks (129) at one and two MMA tiles per warp, and without (256). + for (ggml_type type : { GGML_TYPE_Q4_K, GGML_TYPE_Q5_K, GGML_TYPE_Q6_K, GGML_TYPE_Q8_0, GGML_TYPE_IQ4_XS }) { + for (int64_t m : { 17, 65 }) { + test_cases.emplace_back(new test_mul_mat_vec_fusion(type, GGML_GLU_OP_SWIGLU, m, 129, 512, false, 1, 1, false, false, true, false, { 1, 1 })); + } + test_cases.emplace_back(new test_mul_mat_vec_fusion(type, GGML_GLU_OP_SWIGLU, 128, 256, 512, false, 1, 1, false, false, true, false, { 1, 1 })); + } + for (ggml_glu_op glu_op : { GGML_GLU_OP_GEGLU, GGML_GLU_OP_SWIGLU_OAI, GGML_GLU_OP_SWIGLU_CLAMP }) { + test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_K, glu_op, 17, 129, 512, false, 1, 1, false, false, true, false, { 1, 1 })); + } + test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_K, GGML_GLU_OP_SWIGLU, 33, 257, 512, false, 1, 1, false, false, true, false, { 2, 3 })); + for (auto gate : {GATING_FUNC_SOFTMAX, GATING_FUNC_SIGMOID, GATING_FUNC_SOFTMAX_WEIGHT, GATING_FUNC_SQRT_SOFTPLUS}) { for (bool with_norm : {false, true}) { for (bool bias_probs : {false, true}) { From 6e59326f7d8a9a41b727c33c8ebfa11b10720d72 Mon Sep 17 00:00:00 2001 From: Anuj Jalota Date: Sun, 4 Oct 2026 16:04:41 +0530 Subject: [PATCH 02/11] cuda: assert the fused MMQ gate argument and trim redundant comments --- ggml/src/ggml-cuda/mmq.cu | 4 ++-- ggml/src/ggml-cuda/mmq.cuh | 3 +-- 2 files changed, 3 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 56c7777945f6..0b6ba47b5e4d 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -140,8 +140,8 @@ void ggml_cuda_mul_mat_q( // Fused gate/up/GLU: src0 is the up weight, dst the GLU output. const ggml_tensor * gate = fusion ? fusion->gate : nullptr; - if (gate) { - GGML_ASSERT(!ids && !fusion->x_bias && !fusion->gate_bias && !fusion->x_scale && !fusion->gate_scale); + if (fusion) { + GGML_ASSERT(gate && !ids && !fusion->x_bias && !fusion->gate_bias && !fusion->x_scale && !fusion->gate_scale); GGML_ASSERT(gate->type == src0->type && ggml_are_same_shape(gate, src0) && ggml_are_same_stride(gate, src0)); } diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index c2b794798ab8..5f404db35f03 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -428,7 +428,7 @@ static __host__ int ggml_cuda_mmq_get_nbytes_shared_x(const ggml_cuda_mmq_config // ------------------------------------------------------------ -// Weight types with a fused gate/up/GLU path (see ggml_cuda_mmq_get_load_tiles); the other types run unfused. +// Weight types with a fused gate/up/GLU path (see ggml_cuda_mmq_get_load_tiles). static constexpr __host__ __device__ bool ggml_cuda_mmq_fusion_supported(const ggml_type type) { return type == GGML_TYPE_Q4_K || type == GGML_TYPE_Q5_K || type == GGML_TYPE_Q6_K || type == GGML_TYPE_Q8_0 || type == GGML_TYPE_IQ4_XS; } @@ -1053,7 +1053,6 @@ static __global__ void mul_mat_q( const uint3 sample_ratio, const uint3 nsamples_y, const int stride_sample_x, const int stride_sample_y, const int stride_sample_dst, const uint3 ntx) { - // The fused gate/up/GLU write-back is implemented for the NVIDIA MMA accumulator layout. #ifdef TURING_MMA_AVAILABLE constexpr bool fusion_available = true; #else From 525f87c18d011c4bf175cb8d3d683d3b4495e407 Mon Sep 17 00:00:00 2001 From: Anuj Jalota Date: Fri, 9 Oct 2026 12:14:36 +0530 Subject: [PATCH 03/11] cuda: move the fused MMQ hint check to should_fuse and add TODOs --- ggml/src/ggml-cuda/ggml-cuda.cu | 8 ++++++-- ggml/src/ggml-cuda/mmq.cu | 2 ++ 2 files changed, 8 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 698faad3e65b..ee9fc0ee39eb 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1977,8 +1977,7 @@ static bool ggml_cuda_mul_mat_q_fusion_matches(const ggml_tensor * up, const ggm const ggml_tensor * x_up = up->src[0]; const ggml_tensor * y = up->src[1]; - if (up->op != GGML_OP_MUL_MAT || !ggml_cuda_mmq_fusion_supported(x_up->type) || y->type != GGML_TYPE_F32 || - ggml_get_op_params_i32(up, 1) != GGML_HINT_NONE || ggml_get_op_params_i32(gate, 1) != GGML_HINT_NONE) { + if (up->op != GGML_OP_MUL_MAT || !ggml_cuda_mmq_fusion_supported(x_up->type) || y->type != GGML_TYPE_F32) { return false; } @@ -2002,6 +2001,11 @@ static bool ggml_cuda_mul_mat_q_fusion_matches(const ggml_tensor * up, const ggm } static bool ggml_cuda_should_fuse_mul_mat_q(const ggml_tensor * up, const ggml_tensor * gate, const ggml_tensor * glu) { + // A hinted matmul can be fused, but its unfused path may be faster (e.g. FWHT for GGML_HINT_SRC0_IS_HADAMARD). + if (ggml_get_op_params_i32(up, 1) != GGML_HINT_NONE || ggml_get_op_params_i32(gate, 1) != GGML_HINT_NONE) { + return false; + } + // The padding of a compute-buffer weight that is a view cannot be cleared, see ggml_cuda_mul_mat. const auto bad_padding_clear = [](const ggml_tensor * x) { return ggml_backend_buffer_get_usage(x->buffer) == GGML_BACKEND_BUFFER_USAGE_COMPUTE && diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 0b6ba47b5e4d..cd271fbc8b4a 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -141,6 +141,7 @@ void ggml_cuda_mul_mat_q( // Fused gate/up/GLU: src0 is the up weight, dst the GLU output. const ggml_tensor * gate = fusion ? fusion->gate : nullptr; if (fusion) { + // TODO MoE support GGML_ASSERT(gate && !ids && !fusion->x_bias && !fusion->gate_bias && !fusion->x_scale && !fusion->gate_scale); GGML_ASSERT(gate->type == src0->type && ggml_are_same_shape(gate, src0) && ggml_are_same_stride(gate, src0)); } @@ -191,6 +192,7 @@ void ggml_cuda_mul_mat_q( const bool fallback = ggml_cuda_mmq_needs_fallback(ne01); // With fusion dst is the GLU output, not a MUL_MAT; the fused types always quantize src1 to Q8_1. + // TODO support for Q4 precision if both tensors allow it const ggml_prec prec_src1 = gate ? GGML_PREC_Q8 : ggml_cuda_mmq_get_prec_src1(src0, dst, cc); const bool use_native_fp4 = prec_src1 == GGML_PREC_Q4; From 744d40350716818deb91638a12da6ca69c96bece Mon Sep 17 00:00:00 2001 From: Anuj Jalota Date: Fri, 9 Oct 2026 12:29:16 +0530 Subject: [PATCH 04/11] cuda: pass prec_src1 to the fused MMQ write-back --- ggml/src/ggml-cuda/mmq.cuh | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 5f404db35f03..3e9be8057a6b 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -548,7 +548,7 @@ static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma( } // Fused gate/up/GLU write-back: accumulators l and l + 2 hold the up and gate values of the same output row and column. -template +template static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma_glu( const float * __restrict__ sum, const ggml_glu_op glu_op, const float glu_limit, const int * __restrict__ ids_dst, float * __restrict__ dst, const int stride, const int i_max, const int j_max) { @@ -556,7 +556,7 @@ static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma_glu( typedef tile<16, 8, int> tile_C; static_assert(tile_C::I == 2*GGML_CUDA_MMQ_FUSION_ROWS, "each C tile must hold one chunk of up rows and one of gate rows"); - constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback); + constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback, prec_src1); constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp. const int i0 = (threadIdx.y / ntx) * (ntx*tile_C::I); @@ -1027,7 +1027,7 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( } if constexpr (has_fusion) { - ggml_cuda_mmq_write_back_mma_glu(sum, glu_op, glu_limit, ids_dst, dst, stride_col_dst, tile_x_max_i, tile_y_max_j); + ggml_cuda_mmq_write_back_mma_glu(sum, glu_op, glu_limit, ids_dst, dst, stride_col_dst, tile_x_max_i, tile_y_max_j); return; } GGML_UNUSED_VARS(glu_op, glu_limit); From 4d20b62de587b6dc9a942ae92d692faf2e044b79 Mon Sep 17 00:00:00 2001 From: Anuj Jalota Date: Fri, 9 Oct 2026 17:46:49 +0530 Subject: [PATCH 05/11] cuda: fuse only with efficient tiling and no fallback, merge the GLU write-back --- ggml/src/ggml-cuda/ggml-cuda.cu | 34 ++++++- ggml/src/ggml-cuda/mmq-load-tiles.cuh | 60 ++++++------- ggml/src/ggml-cuda/mmq.cu | 67 +++++++++----- ggml/src/ggml-cuda/mmq.cuh | 122 ++++++++++---------------- tests/test-backend-ops.cpp | 14 +-- 5 files changed, 161 insertions(+), 136 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index ee9fc0ee39eb..f8d3c90dac12 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1972,7 +1972,7 @@ static bool ggml_cuda_match_shared_expert(const ggml_cgraph * graph, int routed_ return true; } -// Buffer-independent, so graph_optimize can call it before allocation. +// Buffer- and batch-size independent, so graph_optimize produces the same graph topology for every ubatch size (see #29986). static bool ggml_cuda_mul_mat_q_fusion_matches(const ggml_tensor * up, const ggml_tensor * gate, const ggml_tensor * glu, const int device) { const ggml_tensor * x_up = up->src[0]; const ggml_tensor * y = up->src[1]; @@ -1981,10 +1981,23 @@ static bool ggml_cuda_mul_mat_q_fusion_matches(const ggml_tensor * up, const ggm return false; } + // The fused kernel is only compiled for weights whose rows fill whole tiles (no fallback). + if (ggml_cuda_mmq_needs_fallback(x_up->ne[1])) { + return false; + } + if (!ggml_cuda_should_fuse_mul_mat(up, gate, glu)) { return false; } + // The fused kernel quantizes src1 to Q8_1, only fuse if the src1 precision of up and gate allows it. + for (const ggml_tensor * mm : { up, gate }) { + const ggml_prec prec_src1 = (ggml_prec) ggml_get_op_params_i32(mm, 3); + if (prec_src1 != GGML_PREC_UNDEFINED && prec_src1 != GGML_PREC_Q8 && prec_src1 != GGML_PREC_Q4) { + return false; + } + } + // The fused write-back applies SWIGLU_OAI with the default alpha and limit of ggml_cuda_op_swiglu_oai_single. if (ggml_get_glu_op(glu) == GGML_GLU_OP_SWIGLU_OAI && (ggml_get_op_params_f32(glu, 2) != 1.702f || ggml_get_op_params_f32(glu, 3) != 7.0f)) { return false; @@ -1996,8 +2009,7 @@ static bool ggml_cuda_mul_mat_q_fusion_matches(const ggml_tensor * up, const ggm } // NVIDIA only: the fused write-back relies on the MMA accumulator layout. - const int cc = ggml_cuda_info().devices[device].cc; - return turing_mma_available(cc) && !ggml_cuda_should_use_mmvq(x_up->type, cc, y->ne[1]) && ggml_cuda_should_use_mmq(x_up->type, cc, y->ne[1], 0); + return turing_mma_available(ggml_cuda_info().devices[device].cc); } static bool ggml_cuda_should_fuse_mul_mat_q(const ggml_tensor * up, const ggml_tensor * gate, const ggml_tensor * glu) { @@ -2011,7 +2023,21 @@ static bool ggml_cuda_should_fuse_mul_mat_q(const ggml_tensor * up, const ggml_t return ggml_backend_buffer_get_usage(x->buffer) == GGML_BACKEND_BUFFER_USAGE_COMPUTE && ggml_nbytes(x) != ggml_backend_buffer_get_alloc_size(x->buffer, x) && x->view_src; }; - return ggml_cuda_mul_mat_q_fusion_matches(up, gate, glu, ggml_cuda_get_device()) && !bad_padding_clear(up->src[0]) && !bad_padding_clear(gate->src[0]); + const int device = ggml_cuda_get_device(); + if (!ggml_cuda_mul_mat_q_fusion_matches(up, gate, glu, device) || bad_padding_clear(up->src[0]) || bad_padding_clear(gate->src[0])) { + return false; + } + + // The batch size checks are done here, not in ggml_cuda_mul_mat_q_fusion_matches. + const ggml_tensor * x_up = up->src[0]; + const ggml_tensor * y = up->src[1]; + const int cc = ggml_cuda_info().devices[device].cc; + if (ggml_cuda_should_use_mmvq(x_up->type, cc, y->ne[1]) || !ggml_cuda_should_use_mmq(x_up->type, cc, y->ne[1], 0)) { + return false; + } + + // The fused kernel cannot use stream-k, so only fuse if its tiling is efficient. + return ggml_cuda_mmq_fusion_is_efficient(x_up->type, x_up->ne[1], y->ne[1], y->ne[2]*y->ne[3], device); } static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { diff --git a/ggml/src/ggml-cuda/mmq-load-tiles.cuh b/ggml/src/ggml-cuda/mmq-load-tiles.cuh index 66e3e58df5b7..f4776173be2a 100644 --- a/ggml/src/ggml-cuda/mmq-load-tiles.cuh +++ b/ggml/src/ggml-cuda/mmq-load-tiles.cuh @@ -527,11 +527,11 @@ template static for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); - if (fallback && !has_fusion) { + if (fallback) { i = min(i, i_max); } - const block_q8_0 * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride) + kbx; + const block_q8_0 * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, stride) + kbx; #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) x_qs[i*sram_stride + 0 + txi] = get_int_b2(bxi[0].qs, kqsx); @@ -550,11 +550,11 @@ template static for (int i0 = 0; i0 < I; i0 += nwarps * rows_per_warp) { int i = i0 + threadIdx.y * rows_per_warp + threadIdx.x / blocks_per_tile_x_row; - if (fallback && !has_fusion) { + if (fallback) { i = min(i, i_max); } - const block_q8_0 * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride) + kbxd; + const block_q8_0 * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, stride) + kbxd; #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) x_df[i*sram_stride + kbxd] = bxi->d; @@ -769,11 +769,11 @@ template static for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); - if (fallback && !has_fusion) { + if (fallback) { i = min(i, i_max); } - const block_q4_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); + const block_q4_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, stride); const int qs0 = get_int_b4(bxi->qs, txi); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) @@ -798,11 +798,11 @@ template static int i = (i0 + threadIdx.y*rows_per_warp + threadIdx.x/2) % I; { #endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) - if (fallback && !has_fusion) { + if (fallback) { i = min(i, i_max); } - const block_q4_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); + const block_q4_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, stride); const int * scales = (const int *) bxi->scales; const int ksc = threadIdx.x % 2; @@ -826,11 +826,11 @@ template static for (int i0 = 0; i0 < I; i0 += nwarps*warp_size) { int i = (i0 + threadIdx.y*warp_size + threadIdx.x) % I; - if (fallback && !has_fusion) { + if (fallback) { i = min(i, i_max); } - const block_q4_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); + const block_q4_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, stride); x_dm[i] = bxi->dm; } @@ -839,11 +839,11 @@ template static for (int i0 = 0; i0 < I; i0 += nwarps*rows_per_warp) { int i = (i0 + threadIdx.y*rows_per_warp + threadIdx.x/(MMQ_TILE_NE_K/8)) % I; - if (fallback && !has_fusion) { + if (fallback) { i = min(i, i_max); } - const block_q4_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride) + (threadIdx.x % (MMQ_TILE_NE_K/8)) / (QI4_K/8); + const block_q4_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, stride) + (threadIdx.x % (MMQ_TILE_NE_K/8)) / (QI4_K/8); const int * scales = (const int *) bxi->scales; @@ -880,11 +880,11 @@ template static for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); - if (fallback && !has_fusion) { + if (fallback) { i = min(i, i_max); } - const block_q5_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); + const block_q5_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, stride); const int ky = QR5_K*txi; const int ql = get_int_b4(bxi->qs, txi); @@ -921,11 +921,11 @@ template static int i = (i0 + threadIdx.y*rows_per_warp + threadIdx.x/2) % I; { #endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) - if (fallback && !has_fusion) { + if (fallback) { i = min(i, i_max); } - const block_q5_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); + const block_q5_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, stride); const int * scales = (const int *) bxi->scales; const int ksc = threadIdx.x % 2; @@ -949,11 +949,11 @@ template static for (int i0 = 0; i0 < I; i0 += nwarps*warp_size) { int i = (i0 + threadIdx.y*warp_size + threadIdx.x) % I; - if (fallback && !has_fusion) { + if (fallback) { i = min(i, i_max); } - const block_q5_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); + const block_q5_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, stride); x_dm[i] = bxi->dm; } @@ -963,11 +963,11 @@ template static for (int i0 = 0; i0 < I; i0 += nwarps*rows_per_warp) { int i = (i0 + threadIdx.y*rows_per_warp + threadIdx.x/(MMQ_TILE_NE_K/8)) % I; - if (fallback && !has_fusion) { + if (fallback) { i = min(i, i_max); } - const block_q5_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); + const block_q5_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, stride); const int * scales = (const int *) bxi->scales; @@ -1005,11 +1005,11 @@ template static for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); - if (fallback && !has_fusion) { + if (fallback) { i = min(i, i_max); } - const block_q6_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); + const block_q6_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, stride); const int ql = get_int_b2(bxi->ql, txi); const int ql0 = (ql >> 0) & 0x0F0F0F0F; @@ -1035,11 +1035,11 @@ template static for (int i0 = 0; i0 < I; i0 += nwarps*warp_size) { int i = (i0 + threadIdx.y*warp_size + threadIdx.x) % I; - if (fallback && !has_fusion) { + if (fallback) { i = min(i, i_max); } - const block_q6_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); + const block_q6_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, stride); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) x_df[i*sram_stride] = bxi->d; @@ -1053,11 +1053,11 @@ template static for (int i0 = 0; i0 < I; i0 += nwarps*rows_per_warp) { int i = (i0 + threadIdx.y*rows_per_warp + threadIdx.x/(MMQ_TILE_NE_K/8)) % I; - if (fallback && !has_fusion) { + if (fallback) { i = min(i, i_max); } - const block_q6_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride) + (threadIdx.x % (MMQ_TILE_NE_K/8)) / 4; + const block_q6_K * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, stride) + (threadIdx.x % (MMQ_TILE_NE_K/8)) / 4; #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) x_sc[i*sram_stride + threadIdx.x%4] = get_int_b2(bxi->scales, threadIdx.x % (MMQ_TILE_NE_K/8)); @@ -1497,11 +1497,11 @@ template static for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); - if (fallback && !has_fusion) { + if (fallback) { i = min(i, i_max); } - const block_iq4_xs * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); + const block_iq4_xs * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, stride); const int aux_q4 = get_int_b4(bxi->qs, kqsx); const int2 v = get_int_from_table_16(aux_q4, kvalues_iq4nl); @@ -1521,11 +1521,11 @@ template static for (int i0 = 0; i0 < I; i0 += nwarps * rows_per_warp) { int i = i0 + threadIdx.y * rows_per_warp + threadIdx.x / (MMQ_TILE_NE_K/4); - if (fallback && !has_fusion) { + if (fallback) { i = min(i, i_max); } - const block_iq4_xs * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, i_max, stride); + const block_iq4_xs * bxi = ggml_cuda_mmq_x_row(x, x_gate, kbx0, i, stride); const float d = __half2float(bxi->d); diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index cd271fbc8b4a..21f6ac6c82bf 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -131,6 +131,36 @@ static ggml_prec ggml_cuda_mmq_get_prec_src1(const ggml_tensor * src0, const ggm return GGML_PREC_Q4; } +// Tile width in ne11 (dense) / ne12 (MoE) direction with the fewest tiles that fits into shared memory, 0 if there is none. +static int ggml_cuda_mmq_get_J_best( + const ggml_type type, const int64_t ncols, const bool fallback, const int cc, const size_t smpbo, const ggml_prec prec_src1, + int * nthreads_best = nullptr) { + int J_best = 0; + int ntiles_J_best = INT_MAX; + + for (int J = 8; J <= 128 && ntiles_J_best > 1; J += 8) { + const ggml_cuda_mmq_config config = ggml_cuda_mmq_get_config(type, J, fallback, cc, prec_src1); + if (config.type == GGML_TYPE_COUNT) { + continue; + } + + if (mmq_get_nbytes_shared(config, cc) > smpbo) { + continue; + } + + const int ntiles_x = (ncols + config.J - 1) / config.J; + + if (ntiles_x < ntiles_J_best) { + J_best = J; + if (nthreads_best) { + *nthreads_best = config.nthreads; + } + ntiles_J_best = ntiles_x; + } + } + return J_best; +} + void ggml_cuda_mul_mat_q( ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * ids, ggml_tensor * dst, const ggml_cuda_mm_fusion_args_host * fusion) { @@ -214,26 +244,7 @@ void ggml_cuda_mul_mat_q( } } - int ntiles_J_best = INT_MAX; - - for (int J = 8; J <= 128 && ntiles_J_best > 1; J += 8) { - const ggml_cuda_mmq_config config = ggml_cuda_mmq_get_config(src0->type, J, fallback, cc, prec_src1); - if (config.type == GGML_TYPE_COUNT) { - continue; - } - - if (mmq_get_nbytes_shared(config, cc) > smpbo) { - continue; - } - - const int ntiles_x = (ncols_opt + config.J - 1) / config.J; - - if (ntiles_x < ntiles_J_best) { - J_best = J; - nthreads_best = config.nthreads; - ntiles_J_best = ntiles_x; - } - } + J_best = ggml_cuda_mmq_get_J_best(src0->type, ncols_opt, fallback, cc, smpbo, prec_src1, &nthreads_best); } GGML_ASSERT(J_best > 0); @@ -505,3 +516,19 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t return (!GGML_CUDA_CC_IS_CDNA(cc)) || ne11 < MMQ_DP4A_MAX_BATCH_SIZE; } + +bool ggml_cuda_mmq_fusion_is_efficient(const ggml_type type, const int64_t nrows_x, const int64_t ncols, const int64_t nchannels, const int device) { + const int cc = ggml_cuda_info().devices[device].cc; + const int nsm = ggml_cuda_info().devices[device].nsm; + const size_t smpbo = ggml_cuda_info().devices[device].smpbo; + + // Same tile width as ggml_cuda_mul_mat_q; the fused kernel has no fallback, quantizes src1 to Q8_1 and writes I/2 output rows per tile. + const bool fallback = false; + const int J_best = ggml_cuda_mmq_get_J_best(type, ncols, fallback, cc, smpbo, GGML_PREC_Q8); + if (J_best == 0) { + return false; + } + const int64_t nrows_tile = ggml_cuda_mmq_get_config(type, J_best, fallback, cc, GGML_PREC_Q8).I / 2; + const int64_t ntiles = (ncols + J_best - 1)/J_best * ((nrows_x + nrows_tile - 1)/nrows_tile) * nchannels; + return ggml_cuda_mmq_tiling_is_efficient(ntiles, nsm); +} diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 3e9be8057a6b..9b7fca3f8fb9 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -15,7 +15,8 @@ typedef void (*ggml_cuda_mmq_load_tiles_t)(const char * __restrict__ x, const ch const int kbx0, const int i_max, const int stride); typedef void (*ggml_cuda_mmq_vec_dot_t)(const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00); typedef void (*ggml_cuda_mmq_write_back_t)(const float * __restrict__ sum, const int32_t * __restrict__ get_rows_to_sorted, - float * __restrict__ dst, const float * __restrict__ y_scale, const int stride, const int i_max, const int j_max); + float * __restrict__ dst, const float * __restrict__ y_scale, const int stride, const int i_max, const int j_max, + const ggml_glu_op glu_op, const float glu_limit); enum mmq_q8_1_ds_layout { MMQ_Q8_1_DS_LAYOUT_D4, @@ -374,6 +375,12 @@ static __host__ bool ggml_cuda_mmq_needs_fallback(const int64_t nrows_x) { return nrows_x % 128 != 0; } +// Whether one CUDA block per tile (tiling) fills the SMs well enough that stream-k is not needed. +static __host__ bool ggml_cuda_mmq_tiling_is_efficient(const int64_t ntiles, const int nsm) { + const int64_t nwaves = (ntiles + nsm - 1) / nsm; + return 100*ntiles / (nsm*nwaves) >= 90; +} + static constexpr __device__ int ggml_cuda_mmq_get_rows_per_warp(ggml_type type, int J, bool fallback, ggml_prec prec_src1) { return ggml_cuda_mmq_get_config(type, J, fallback, prec_src1).rows_per_warp(); } @@ -436,20 +443,16 @@ static constexpr __host__ __device__ bool ggml_cuda_mmq_fusion_supported(const g // Fused gate/up/GLU: the weight tile alternates chunks of this many up and gate rows, so rows r and r + 8 of an MMA tile (same thread) pair up. static constexpr int GGML_CUDA_MMQ_FUSION_ROWS = 8; -// Source of weight tile row i; with has_fusion the interleaved x/x_gate row, clamped here because i_max is in output rows. -template +// Source of weight tile row i; with has_fusion the up or gate row that the interleaved tile holds there. +template static __device__ __forceinline__ const block_t * ggml_cuda_mmq_x_row( - const char * __restrict__ x, const char * __restrict__ x_gate, const int kbx0, int i, const int i_max, const int stride) { + const char * __restrict__ x, const char * __restrict__ x_gate, const int kbx0, const int i, const int stride) { if constexpr (has_fusion) { constexpr int R = GGML_CUDA_MMQ_FUSION_ROWS; const char * x_src = (i / R) % 2 == 0 ? x : x_gate; - i = (i / (2*R))*R + i % R; - if (fallback) { - i = min(i, i_max); - } - return (const block_t *) x_src + kbx0 + i*stride; + return (const block_t *) x_src + kbx0 + ((i / (2*R))*R + i % R)*stride; } else { - GGML_UNUSED_VARS(x_gate, i_max); + GGML_UNUSED(x_gate); return (const block_t *) x + kbx0 + i*stride; } } @@ -459,7 +462,10 @@ static __device__ __forceinline__ const block_t * ggml_cuda_mmq_x_row( template static __device__ __forceinline__ void ggml_cuda_mmq_write_back_dp4a( const float * __restrict__ sum, const int32_t * __restrict__ ids_dst, float * __restrict__ dst, - const float * __restrict__ y_scale, const int stride, const int i_max, const int j_max) { + const float * __restrict__ y_scale, const int stride, const int i_max, const int j_max, + const ggml_glu_op glu_op, const float glu_limit) { + GGML_UNUSED_VARS(glu_op, glu_limit); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, prec_src1) / warp_size; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, prec_src1); @@ -496,16 +502,19 @@ template static __de } } -template +// has_fusion: accumulators l and l + ne/2 hold the up and gate values of one output element, the GLU of the two is written. +template static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma( const float * __restrict__ sum, const int * __restrict__ ids_dst, float * __restrict__ dst, - const float * __restrict__ y_scale, const int stride, const int i_max, const int j_max) { + const float * __restrict__ y_scale, const int stride, const int i_max, const int j_max, + const ggml_glu_op glu_op, const float glu_limit) { #if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) typedef tile<16, 16, int, DATA_LAYOUT_J_MAJOR> tile_C; #else typedef tile<16, 8, int> tile_C; #endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + static_assert(!has_fusion || tile_C::I == 2*GGML_CUDA_MMQ_FUSION_ROWS, "each C tile must hold one chunk of up rows and one of gate rows"); constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback, prec_src1); constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp. @@ -519,20 +528,23 @@ static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma( #pragma unroll for (int n = 0; n < ntx; ++n) { #pragma unroll - for (int l = 0; l < tile_C::ne; ++l) { + for (int l = 0; l < (has_fusion ? tile_C::ne/2 : tile_C::ne); ++l) { const int j = j0 + (threadIdx.y % ntx) * tile_C::J + tile_C::get_j(l); if (j > j_max) { continue; } - const int i = i0 + n*tile_C::I + tile_C::get_i(l); + const int i = (i0 + n*tile_C::I)/(has_fusion ? 2 : 1) + tile_C::get_i(l); if (fallback && i > i_max) { continue; } - if constexpr (type == GGML_TYPE_NVFP4) { + if constexpr (has_fusion) { + const int sum_index = (j0/tile_C::J + n)*tile_C::ne + l; + dst[ids_dst[j]*stride + i] = ggml_cuda_op_glu_single(glu_op, sum[sum_index + tile_C::ne/2], sum[sum_index], glu_limit); + } else if constexpr (type == GGML_TYPE_NVFP4) { if (y_scale_used) { dst[ids_dst[j]*stride + i] = y_scale[j] * sum[(j0/tile_C::J + n)*tile_C::ne + l]; } else { @@ -545,49 +557,8 @@ static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma( } } } -} - -// Fused gate/up/GLU write-back: accumulators l and l + 2 hold the up and gate values of the same output row and column. -template -static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma_glu( - const float * __restrict__ sum, const ggml_glu_op glu_op, const float glu_limit, - const int * __restrict__ ids_dst, float * __restrict__ dst, const int stride, const int i_max, const int j_max) { -#ifdef TURING_MMA_AVAILABLE - typedef tile<16, 8, int> tile_C; - static_assert(tile_C::I == 2*GGML_CUDA_MMQ_FUSION_ROWS, "each C tile must hold one chunk of up rows and one of gate rows"); - - constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback, prec_src1); - constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp. - - const int i0 = (threadIdx.y / ntx) * (ntx*tile_C::I); - -#pragma unroll - for (int j0 = 0; j0 < J; j0 += ntx*tile_C::J) { -#pragma unroll - for (int n = 0; n < ntx; ++n) { -#pragma unroll - for (int l = 0; l < tile_C::ne/2; ++l) { - const int j = j0 + (threadIdx.y % ntx) * tile_C::J + tile_C::get_j(l); - - if (j > j_max) { - continue; - } - const int i = (i0 + n*tile_C::I)/2 + tile_C::get_i(l); - - if (fallback && i > i_max) { - continue; - } - - const int sum_index = (j0/tile_C::J + n)*tile_C::ne + l; - dst[ids_dst[j]*stride + i] = ggml_cuda_op_glu_single(glu_op, sum[sum_index + 2], sum[sum_index], glu_limit); - } - } - } -#else - GGML_UNUSED_VARS(sum, glu_op, glu_limit, ids_dst, dst, stride, i_max, j_max); - NO_DEVICE_CODE; -#endif // TURING_MMA_AVAILABLE + GGML_UNUSED_VARS(y_scale_used, glu_op, glu_limit); } // ------------------------------------------------------------------------------------------------------------------------------------- @@ -946,14 +917,18 @@ static constexpr __device__ ggml_cuda_mmq_vec_dot_t ggml_cuda_mmq_get_vec_dot() return ggml_cuda_mmq_get_util_funcs().vec_dot; } -template +template static constexpr __device__ ggml_cuda_mmq_write_back_t ggml_cuda_mmq_get_write_back() { - return ggml_cuda_mmq_get_util_funcs().write_back; + if constexpr (has_fusion) { + return ggml_cuda_mmq_write_back_mma; + } else { + return ggml_cuda_mmq_get_util_funcs().write_back; + } } // --------------------------------------------------------------------------------------------- -// has_fusion: x and x_gate are the up and gate weights; tile_x_max_i is then in output rows. +// has_fusion: x and x_gate are the up and gate weights. template static __device__ __forceinline__ void mul_mat_q_process_tile( const char * __restrict__ x, const char * __restrict__ x_gate, const int offset_x, const int * __restrict__ y, @@ -962,6 +937,7 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( const int stride_row_x, const int ncols_y, const int stride_col_dst, const int tile_x_max_i, const int tile_y_max_j, const int kb0_start, const int kb0_stop) { static_assert(!(fixup && has_fusion), "the fused gate/up path does not split K"); + static_assert(!(fallback && has_fusion), "the fused gate/up path needs nrows_x % 128 == 0"); constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, prec_src1) / warp_size; @@ -969,7 +945,7 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, prec_src1); constexpr ggml_cuda_mmq_load_tiles_t load_tiles = ggml_cuda_mmq_get_load_tiles(); constexpr ggml_cuda_mmq_vec_dot_t vec_dot = ggml_cuda_mmq_get_vec_dot(); - constexpr ggml_cuda_mmq_write_back_t write_back = ggml_cuda_mmq_get_write_back(); + constexpr ggml_cuda_mmq_write_back_t write_back = ggml_cuda_mmq_get_write_back(); extern __shared__ int data_mul_mat_q[]; int * tile_y = data_mul_mat_q + J; @@ -1026,16 +1002,10 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( __syncthreads(); } - if constexpr (has_fusion) { - ggml_cuda_mmq_write_back_mma_glu(sum, glu_op, glu_limit, ids_dst, dst, stride_col_dst, tile_x_max_i, tile_y_max_j); - return; - } - GGML_UNUSED_VARS(glu_op, glu_limit); - if (fixup) { - write_back(sum, ids_dst, tmp_fixup + blockIdx.x*(J*I), y_scale, I, I, J); + write_back(sum, ids_dst, tmp_fixup + blockIdx.x*(J*I), y_scale, I, I, J, glu_op, glu_limit); } else { - write_back(sum, ids_dst, dst, y_scale, stride_col_dst, tile_x_max_i, tile_y_max_j); + write_back(sum, ids_dst, dst, y_scale, stride_col_dst, tile_x_max_i, tile_y_max_j, glu_op, glu_limit); } } @@ -1509,7 +1479,9 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a const dim3 block_dims(warp_size, nwarps, 1); CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); - CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); + if constexpr (!has_fusion) { + CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); + } const int nty = (args.nrows_x + config.I - 1) / config.I; const int ntx = (args.ncols_max + config.J - 1) / config.J; @@ -1544,13 +1516,11 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a // For the stream-k kernel it is possible to run it with tiling by setting the number of CUDA blocks equal to the number of tiles. // This is worthwhile if the efficiency of tiling is high and skipping the fixup kernel is more important. const int ntiles_dst = ntx * nty * ntzw; - const int tiles_nwaves = (ntiles_dst + nsm - 1) / nsm; - const int tiles_efficiency_percent = 100 * ntiles_dst / (nsm*tiles_nwaves); GGML_ASSERT(ntiles_dst * blocks_per_ne00_fd.z < (1 << 30)); // Assert that variable kbc will not overflow. const int nwork = ntiles_dst * ((args.ncols_x + config.K_vram - 1) / config.K_vram); // Maximum number of CUDA blocks that can do useful work. - const dim3 block_nums_stream_k(GGML_CUDA_CC_IS_NVIDIA(cc) && tiles_efficiency_percent >= 90 ? ntiles_dst : std::min(nsm, nwork), 1, 1); + const dim3 block_nums_stream_k(GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_mmq_tiling_is_efficient(ntiles_dst, nsm) ? ntiles_dst : std::min(nsm, nwork), 1, 1); const bool fixup_needed = ntiles_dst % block_nums_stream_k.x != 0; @@ -1640,7 +1610,7 @@ void mul_mat_q_switch_J(ggml_backend_cuda_context & ctx, const mmq_args & args, template static void mul_mat_q_switch_fusion(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) { - if constexpr (ggml_cuda_mmq_fusion_supported(type)) { + if constexpr (!fallback && ggml_cuda_mmq_fusion_supported(type)) { if (args.fusion.gate) { mul_mat_q_switch_J(ctx, args, stream); return; @@ -1703,3 +1673,5 @@ void ggml_cuda_mul_mat_q( const ggml_cuda_mm_fusion_args_host * fusion = nullptr); bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t n_experts); + +bool ggml_cuda_mmq_fusion_is_efficient(ggml_type type, int64_t nrows_x, int64_t ncols, int64_t nchannels, int device); diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 9e1eb2113eb4..d599454bc4f8 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -12034,17 +12034,17 @@ static std::vector> make_test_cases_eval() { false, 16, 8, false, false, true, false, { 1, 1 })); } - // Fused gate/up + GLU on the MMQ path: rows with bounds checks (129) at one and two MMA tiles per warp, and without (256). + // Fused gate/up + GLU on the MMQ path, which is only used if its tiling fills the GPU: >= 2048 tiles, enough for up to 252 SMs. + // One MMA tile per warp (m = 17, 1 column tile) and two (m = 1024, 8 column tiles); k = 512 for more than one iteration over k. for (ggml_type type : { GGML_TYPE_Q4_K, GGML_TYPE_Q5_K, GGML_TYPE_Q6_K, GGML_TYPE_Q8_0, GGML_TYPE_IQ4_XS }) { - for (int64_t m : { 17, 65 }) { - test_cases.emplace_back(new test_mul_mat_vec_fusion(type, GGML_GLU_OP_SWIGLU, m, 129, 512, false, 1, 1, false, false, true, false, { 1, 1 })); - } - test_cases.emplace_back(new test_mul_mat_vec_fusion(type, GGML_GLU_OP_SWIGLU, 128, 256, 512, false, 1, 1, false, false, true, false, { 1, 1 })); + test_cases.emplace_back(new test_mul_mat_vec_fusion(type, GGML_GLU_OP_SWIGLU, 17, 131072, 256, false, 1, 1, false, false, true, false, { 1, 1 })); + test_cases.emplace_back(new test_mul_mat_vec_fusion(type, GGML_GLU_OP_SWIGLU, 1024, 16384, 256, false, 1, 1, false, false, true, false, { 1, 1 })); } + test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_K, GGML_GLU_OP_SWIGLU, 1024, 16384, 512, false, 1, 1, false, false, true, false, { 1, 1 })); for (ggml_glu_op glu_op : { GGML_GLU_OP_GEGLU, GGML_GLU_OP_SWIGLU_OAI, GGML_GLU_OP_SWIGLU_CLAMP }) { - test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_K, glu_op, 17, 129, 512, false, 1, 1, false, false, true, false, { 1, 1 })); + test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_K, glu_op, 1024, 16384, 256, false, 1, 1, false, false, true, false, { 1, 1 })); } - test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_K, GGML_GLU_OP_SWIGLU, 33, 257, 512, false, 1, 1, false, false, true, false, { 2, 3 })); + test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_K, GGML_GLU_OP_SWIGLU, 1024, 2816, 256, false, 1, 1, false, false, true, false, { 2, 3 })); for (auto gate : {GATING_FUNC_SOFTMAX, GATING_FUNC_SIGMOID, GATING_FUNC_SOFTMAX_WEIGHT, GATING_FUNC_SQRT_SOFTPLUS}) { for (bool with_norm : {false, true}) { From ebadd88c6454bfe99cf21fac204355981a911082 Mon Sep 17 00:00:00 2001 From: Anuj Jalota Date: Tue, 20 Oct 2026 12:29:58 +0530 Subject: [PATCH 06/11] cuda: require has_fusion explicitly, drop the fusion-only src1 precision check --- ggml/src/ggml-cuda/ggml-cuda.cu | 8 --- ggml/src/ggml-cuda/mmq-load-tiles.cuh | 10 +-- ggml/src/ggml-cuda/mmq.cuh | 92 +++++++++++++-------------- 3 files changed, 51 insertions(+), 59 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index f8d3c90dac12..8b8dc6768bcd 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1990,14 +1990,6 @@ static bool ggml_cuda_mul_mat_q_fusion_matches(const ggml_tensor * up, const ggm return false; } - // The fused kernel quantizes src1 to Q8_1, only fuse if the src1 precision of up and gate allows it. - for (const ggml_tensor * mm : { up, gate }) { - const ggml_prec prec_src1 = (ggml_prec) ggml_get_op_params_i32(mm, 3); - if (prec_src1 != GGML_PREC_UNDEFINED && prec_src1 != GGML_PREC_Q8 && prec_src1 != GGML_PREC_Q4) { - return false; - } - } - // The fused write-back applies SWIGLU_OAI with the default alpha and limit of ggml_cuda_op_swiglu_oai_single. if (ggml_get_glu_op(glu) == GGML_GLU_OP_SWIGLU_OAI && (ggml_get_op_params_f32(glu, 2) != 1.702f || ggml_get_op_params_f32(glu, 3) != 7.0f)) { return false; diff --git a/ggml/src/ggml-cuda/mmq-load-tiles.cuh b/ggml/src/ggml-cuda/mmq-load-tiles.cuh index f4776173be2a..8b87e50a91a3 100644 --- a/ggml/src/ggml-cuda/mmq-load-tiles.cuh +++ b/ggml/src/ggml-cuda/mmq-load-tiles.cuh @@ -500,7 +500,7 @@ template static __device__ __forceinline_ } } -template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q8_0( +template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q8_0( const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; @@ -744,7 +744,7 @@ static __device__ __forceinline__ int unpack_scales_q45_K(const int * scales, co ((scales[ksc/2] >> (2 * (ksc % 2))) & 0x30303030); // upper 2 bits } -template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q4_K( +template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q4_K( const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; @@ -855,7 +855,7 @@ template static #endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) } -template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q5_K( +template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q5_K( const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; @@ -979,7 +979,7 @@ template static #endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) } -template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q6_K( +template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q6_K( const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; @@ -1473,7 +1473,7 @@ template static __device__ __forceinline_ } } -template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq4_xs( +template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq4_xs( const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 9b7fca3f8fb9..7ec283bc2598 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -503,7 +503,7 @@ template static __de } // has_fusion: accumulators l and l + ne/2 hold the up and gate values of one output element, the GLU of the two is written. -template +template static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma( const float * __restrict__ sum, const int * __restrict__ ids_dst, float * __restrict__ dst, const float * __restrict__ y_scale, const int stride, const int i_max, const int j_max, @@ -618,7 +618,7 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func case GGML_TYPE_Q8_0: return ggml_cuda_mmq_util_funcs( VDR_Q8_0_Q8_1_MMQ, - ggml_cuda_mmq_load_tiles_q8_0, + ggml_cuda_mmq_load_tiles_q8_0, ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a, ggml_cuda_mmq_write_back_dp4a); // --------------------------------------------------------------------------------------------- @@ -637,19 +637,19 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func case GGML_TYPE_Q4_K: return ggml_cuda_mmq_util_funcs( VDR_Q4_K_Q8_1_MMQ, - ggml_cuda_mmq_load_tiles_q4_K, + ggml_cuda_mmq_load_tiles_q4_K, ggml_cuda_mmq_vec_dot_q4_K_q8_1_dp4a, ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_Q5_K: return ggml_cuda_mmq_util_funcs( VDR_Q5_K_Q8_1_MMQ, - ggml_cuda_mmq_load_tiles_q5_K, + ggml_cuda_mmq_load_tiles_q5_K, ggml_cuda_mmq_vec_dot_q5_K_q8_1_dp4a, ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_Q6_K: return ggml_cuda_mmq_util_funcs( VDR_Q6_K_Q8_1_MMQ, - ggml_cuda_mmq_load_tiles_q6_K, + ggml_cuda_mmq_load_tiles_q6_K, ggml_cuda_mmq_vec_dot_q6_K_q8_1_dp4a, ggml_cuda_mmq_write_back_dp4a); // --------------------------------------------------------------------------------------------- @@ -692,7 +692,7 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func case GGML_TYPE_IQ4_XS: return ggml_cuda_mmq_util_funcs( VDR_IQ4_XS_Q8_1_MMQ, - ggml_cuda_mmq_load_tiles_iq4_xs, + ggml_cuda_mmq_load_tiles_iq4_xs, ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a, ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_IQ4_NL: @@ -729,7 +729,7 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func -1, ggml_cuda_mmq_load_tiles_mxfp4_fp4, ggml_cuda_mmq_vec_dot_fp4_fp4_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); } break; case GGML_TYPE_NVFP4: @@ -738,7 +738,7 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func -1, ggml_cuda_mmq_load_tiles_nvfp4_nvfp4, ggml_cuda_mmq_vec_dot_fp4_fp4_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); } break; default: @@ -754,136 +754,136 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func -1, ggml_cuda_mmq_load_tiles_q1_0, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q2_0: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_q2_0, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q4_0: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_q4_0, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q4_1: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_q4_1, ggml_cuda_mmq_vec_dot_q8_1_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q5_0: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_q5_0, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q5_1: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_q5_1, ggml_cuda_mmq_vec_dot_q8_1_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q8_0: return ggml_cuda_mmq_util_funcs( -1, - ggml_cuda_mmq_load_tiles_q8_0, + ggml_cuda_mmq_load_tiles_q8_0, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); // --------------------------------------------------------------------------------------------- case GGML_TYPE_Q2_K: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_q2_K, ggml_cuda_mmq_vec_dot_q2_K_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q3_K: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_q3_K, ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q4_K: return ggml_cuda_mmq_util_funcs( -1, - ggml_cuda_mmq_load_tiles_q4_K, + ggml_cuda_mmq_load_tiles_q4_K, ggml_cuda_mmq_vec_dot_q8_1_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q5_K: return ggml_cuda_mmq_util_funcs( -1, - ggml_cuda_mmq_load_tiles_q5_K, + ggml_cuda_mmq_load_tiles_q5_K, ggml_cuda_mmq_vec_dot_q8_1_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q6_K: return ggml_cuda_mmq_util_funcs( -1, - ggml_cuda_mmq_load_tiles_q6_K, + ggml_cuda_mmq_load_tiles_q6_K, ggml_cuda_mmq_vec_dot_q6_K_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); // --------------------------------------------------------------------------------------------- case GGML_TYPE_IQ1_S: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_iq1_s, ggml_cuda_mmq_vec_dot_q8_1_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_IQ2_XXS: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_iq2_xxs, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_IQ2_XS: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_iq2_xs, ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_IQ2_S: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_iq2_s, ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_IQ3_XXS: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_iq3_xxs, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_IQ3_S: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_iq3_s, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_IQ4_XS: return ggml_cuda_mmq_util_funcs( -1, - ggml_cuda_mmq_load_tiles_iq4_xs, + ggml_cuda_mmq_load_tiles_iq4_xs, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_IQ4_NL: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_iq4_nl, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); // --------------------------------------------------------------------------------------------- case GGML_TYPE_MXFP4: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_mxfp4, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_NVFP4: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_nvfp4, ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); default: return ggml_cuda_mmq_util_funcs(1, nullptr, nullptr, nullptr); } @@ -894,7 +894,7 @@ static constexpr __device__ int ggml_cuda_mmq_get_vdr() { return ggml_cuda_mmq_get_util_funcs().vdr; } -template +template static constexpr __device__ ggml_cuda_mmq_load_tiles_t ggml_cuda_mmq_get_load_tiles() { if constexpr (!has_fusion) { return ggml_cuda_mmq_get_util_funcs().load_tiles; @@ -917,7 +917,7 @@ static constexpr __device__ ggml_cuda_mmq_vec_dot_t ggml_cuda_mmq_get_vec_dot() return ggml_cuda_mmq_get_util_funcs().vec_dot; } -template +template static constexpr __device__ ggml_cuda_mmq_write_back_t ggml_cuda_mmq_get_write_back() { if constexpr (has_fusion) { return ggml_cuda_mmq_write_back_mma; @@ -929,7 +929,7 @@ static constexpr __device__ ggml_cuda_mmq_write_back_t ggml_cuda_mmq_get_write_b // --------------------------------------------------------------------------------------------- // has_fusion: x and x_gate are the up and gate weights. -template +template static __device__ __forceinline__ void mul_mat_q_process_tile( const char * __restrict__ x, const char * __restrict__ x_gate, const int offset_x, const int * __restrict__ y, const int * __restrict__ ids_dst, float * __restrict__ dst, float * __restrict__ tmp_fixup, @@ -1012,7 +1012,7 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( // The mul_mat_q kernel implements "stream-k" work partitioning as described in https://arxiv.org/abs/2301.03598 -template +template __launch_bounds__(ggml_cuda_mmq_get_nthreads(type, J, fallback, prec_src1), ggml_cuda_mmq_get_occupancy(type, J, fallback, prec_src1)) static __global__ void mul_mat_q( const char * __restrict__ x, const char * __restrict__ x_gate, const int * __restrict__ y, const int32_t * __restrict__ ids_dst, @@ -1218,7 +1218,7 @@ static __global__ void mul_mat_q( const int offset_x = fastdiv(wt, sample_ratio)*stride_sample_x + fastdiv(zt, channel_ratio)*stride_channel_x + it*I*stride_row_x; constexpr bool fixup = false; // All but (potentially) the last iterations write their data to dst rather than the fixup buffer. - mul_mat_q_process_tile + mul_mat_q_process_tile (x, x_gate, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, y_scale_tile, glu_op, glu_limit, stride_row_x, ncols_y, stride_col_dst, tile_x_max_i, tile_y_max_j, kb0_start, kb0_stop); @@ -1302,7 +1302,7 @@ static __global__ void mul_mat_q( const int offset_x = fastdiv(wt, sample_ratio)*stride_sample_x + fastdiv(zt, channel_ratio)*stride_channel_x + it*I*stride_row_x; constexpr bool fixup = true; // Last index writes its data to fixup buffer to avoid data races with other blocks. - mul_mat_q_process_tile + mul_mat_q_process_tile (x, x_gate, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, y_scale_tile, glu_op, glu_limit, stride_row_x, ncols_y, stride_col_dst, tile_x_max_i, tile_y_max_j, kb0_start, kb0_stop); @@ -1464,7 +1464,7 @@ static size_t mmq_get_nbytes_shared(const ggml_cuda_mmq_config & config, const i return nbs_ids + nbs_x + GGML_PAD(nbs_y, config.nthreads*sizeof(int)); } -template +template static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) { const int id = ggml_cuda_get_device(); const int cc = ggml_cuda_info().devices[id].cc; @@ -1478,9 +1478,9 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a const dim3 block_dims(warp_size, nwarps, 1); - CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); + CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); if constexpr (!has_fusion) { - CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); + CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); } const int nty = (args.nrows_x + config.I - 1) / config.I; @@ -1551,7 +1551,7 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a ntx_fd); } -template +template void mul_mat_q_switch_J(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) { switch (args.J_best) { case 8: @@ -1617,7 +1617,7 @@ static void mul_mat_q_switch_fusion(ggml_backend_cuda_context & ctx, const mmq_a } } GGML_ASSERT(!args.fusion.gate); - mul_mat_q_switch_J(ctx, args, stream); + mul_mat_q_switch_J(ctx, args, stream); } template From a11b8801ad97fc527ff994e70293627fb405d4aa Mon Sep 17 00:00:00 2001 From: Anuj Jalota Date: Tue, 20 Oct 2026 17:07:42 +0530 Subject: [PATCH 07/11] cuda: launch the fused MMQ kernel through the stream-k kernel with one block per tile --- ggml/src/ggml-cuda/ggml-cuda.cu | 5 ---- ggml/src/ggml-cuda/mmq.cuh | 48 ++++++++++++++++++--------------- 2 files changed, 26 insertions(+), 27 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 8b8dc6768bcd..1fdf1cc6e9ce 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1995,11 +1995,6 @@ static bool ggml_cuda_mul_mat_q_fusion_matches(const ggml_tensor * up, const ggm return false; } - // The fused launch has the row tiles (64 output rows each) in gridDim.y and channels * samples in gridDim.z. - if ((x_up->ne[1] + 63)/64 > 65535 || y->ne[2]*y->ne[3] > 65535) { - return false; - } - // NVIDIA only: the fused write-back relies on the MMA accumulator layout. return turing_mma_available(ggml_cuda_info().devices[device].cc); } diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 7ec283bc2598..8ef4a7e56b05 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -1040,7 +1040,9 @@ static __global__ void mul_mat_q( constexpr int qk = ggml_cuda_type_traits::qk; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, prec_src1); - const uint32_t nty = (nrows_x + I - 1) / I; // Number of tiles y + // The fused kernel writes the GLU of I/2 up rows and I/2 gate rows per tile. + constexpr int I_dst = has_fusion ? I/2 : I; + const uint32_t nty = (nrows_x + I_dst - 1) / I_dst; // Number of tiles y // Initialize the ids for writing back data with just the index. // For regular matrix multiplications this is never changed. @@ -1058,15 +1060,12 @@ static __global__ void mul_mat_q( } __syncthreads(); - // The fused path reduces all of K per block (the GLU is nonlinear), writes I/2 output rows per tile, and has the column tiles innermost for weight reuse. - if constexpr (has_fusion || !ggml_cuda_mmq_get_stream_k(type, J, fallback, prec_src1)) { - constexpr int I_dst = has_fusion ? I/2 : I; - + if constexpr (!ggml_cuda_mmq_get_stream_k(type, J, fallback, prec_src1)) { const uint2 tmp2 = fast_div_modulo(blockIdx.z, nchannels_y); const int wt = tmp2.x; const int zt = tmp2.y; - const int jt = has_fusion ? blockIdx.x : blockIdx.y; - const int it = has_fusion ? blockIdx.y : blockIdx.x; + const int jt = blockIdx.y; + const int it = blockIdx.x; // Defaults for regular matrix multiplication: int col_low = 0; @@ -1111,20 +1110,20 @@ static __global__ void mul_mat_q( } offset_y += (col_low + jt*J)*(sizeof(block_q8_1_mmq)/sizeof(int)); - offset_dst += it*I_dst; + offset_dst += it*I; const float * y_scale_tile = nullptr; if constexpr (type == GGML_TYPE_NVFP4) { offset_y_scale += col_low + jt*J; y_scale_tile = y_scale ? y_scale + offset_y_scale : nullptr; } - const int tile_x_max_i = nrows_x - it*I_dst - 1; + const int tile_x_max_i = nrows_x - it*I - 1; const int tile_y_max_j = col_diff - jt*J - 1; - const int offset_x = fastdiv(wt, sample_ratio)*stride_sample_x + fastdiv(zt, channel_ratio)*stride_channel_x + it*I_dst*stride_row_x; + const int offset_x = fastdiv(wt, sample_ratio)*stride_sample_x + fastdiv(zt, channel_ratio)*stride_channel_x + it*I*stride_row_x; constexpr bool fixup = false; - mul_mat_q_process_tile + mul_mat_q_process_tile (x, x_gate, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, y_scale_tile, glu_op, glu_limit, stride_row_x, ncols_y, stride_col_dst, tile_x_max_i, tile_y_max_j, 0, blocks_per_ne00.z); @@ -1205,24 +1204,28 @@ static __global__ void mul_mat_q( } offset_y += (col_low + jt * J) * (sizeof(block_q8_1_mmq) / sizeof(int)); - offset_dst += it*I; + offset_dst += it*I_dst; const float * y_scale_tile = nullptr; if constexpr (type == GGML_TYPE_NVFP4) { offset_y_scale += col_low + jt * J; y_scale_tile = y_scale ? y_scale + offset_y_scale : nullptr; } - const int tile_x_max_i = nrows_x - it*I - 1; + const int tile_x_max_i = nrows_x - it*I_dst - 1; const int tile_y_max_j = col_diff - jt*J - 1; - const int offset_x = fastdiv(wt, sample_ratio)*stride_sample_x + fastdiv(zt, channel_ratio)*stride_channel_x + it*I*stride_row_x; + const int offset_x = fastdiv(wt, sample_ratio)*stride_sample_x + fastdiv(zt, channel_ratio)*stride_channel_x + it*I_dst*stride_row_x; constexpr bool fixup = false; // All but (potentially) the last iterations write their data to dst rather than the fixup buffer. - mul_mat_q_process_tile + mul_mat_q_process_tile (x, x_gate, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, y_scale_tile, glu_op, glu_limit, stride_row_x, ncols_y, stride_col_dst, tile_x_max_i, tile_y_max_j, kb0_start, kb0_stop); + if constexpr (has_fusion) { + return; // The fused kernel runs with one block per tile. + } + kbc += blocks_per_ne00.z; kbc -= fastmodulo(kbc, blocks_per_ne00); @@ -1230,7 +1233,7 @@ static __global__ void mul_mat_q( kb0_stop = min(blocks_per_ne00.z, uint32_t(kbc_stop - kbc)); } - if (kbc >= kbc_stop) { + if (has_fusion || kbc >= kbc_stop) { return; } @@ -1483,7 +1486,8 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); } - const int nty = (args.nrows_x + config.I - 1) / config.I; + const int I_dst = has_fusion ? config.I/2 : config.I; // Rows of dst per tile, see mul_mat_q. + const int nty = (args.nrows_x + I_dst - 1) / I_dst; const int ntx = (args.ncols_max + config.J - 1) / config.J; const int ntzw = args.nchannels_y * args.nsamples_y; const dim3 block_nums_xy_tiling(nty, ntx, ntzw); @@ -1500,11 +1504,9 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a const uint3 channel_ratio_fd = init_fastdiv_values(channel_ratio); const uint3 sample_ratio_fd = init_fastdiv_values(sample_ratio); - if (has_fusion || !config.stream_k) { - const int nty_fusion = (args.nrows_x + config.I/2 - 1) / (config.I/2); - GGML_ASSERT(!has_fusion || nty_fusion <= 65535); - const dim3 block_nums = has_fusion ? dim3(ntx, nty_fusion, ntzw) : block_nums_xy_tiling; - mul_mat_q<<>> + if (!config.stream_k) { + GGML_ASSERT(!has_fusion); + mul_mat_q<<>> (args.x, (const char *) args.fusion.gate, args.y, args.ids_dst, args.expert_bounds, args.dst, nullptr, args.y_scale, args.fusion.glu_op, args.fusion.glu_limit, blocks_per_ne00_fd, args.nrows_x, args.ncols_dst, args.stride_row_x, args.ncols_y, args.nrows_dst, channel_ratio_fd, nchannels_y_fd, args.stride_channel_x, args.stride_channel_y, args.stride_channel_dst, @@ -1521,6 +1523,8 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a const int nwork = ntiles_dst * ((args.ncols_x + config.K_vram - 1) / config.K_vram); // Maximum number of CUDA blocks that can do useful work. const dim3 block_nums_stream_k(GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_mmq_tiling_is_efficient(ntiles_dst, nsm) ? ntiles_dst : std::min(nsm, nwork), 1, 1); + // The GLU is nonlinear, so the fused kernel needs the whole K range of a tile in one block. + GGML_ASSERT(!has_fusion || int(block_nums_stream_k.x) == ntiles_dst); const bool fixup_needed = ntiles_dst % block_nums_stream_k.x != 0; From 308407696e2a32d4817bd2033135fc69b1a8f2a8 Mon Sep 17 00:00:00 2001 From: Anuj Jalota Date: Tue, 20 Oct 2026 18:25:24 +0530 Subject: [PATCH 08/11] cuda: inline the MMQ tiling-efficiency check --- ggml/src/ggml-cuda/ggml-cuda.cu | 2 +- ggml/src/ggml-cuda/mmq.cu | 5 ++++- ggml/src/ggml-cuda/mmq.cuh | 10 +++------- 3 files changed, 8 insertions(+), 9 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 1fdf1cc6e9ce..9d40c0af4a6e 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -2023,7 +2023,7 @@ static bool ggml_cuda_should_fuse_mul_mat_q(const ggml_tensor * up, const ggml_t return false; } - // The fused kernel cannot use stream-k, so only fuse if its tiling is efficient. + // The fused kernel needs one block per tile (no stream-k split), so only fuse if the tiling is efficient. return ggml_cuda_mmq_fusion_is_efficient(x_up->type, x_up->ne[1], y->ne[1], y->ne[2]*y->ne[3], device); } diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 21f6ac6c82bf..27558ec77289 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -530,5 +530,8 @@ bool ggml_cuda_mmq_fusion_is_efficient(const ggml_type type, const int64_t nrows } const int64_t nrows_tile = ggml_cuda_mmq_get_config(type, J_best, fallback, cc, GGML_PREC_Q8).I / 2; const int64_t ntiles = (ncols + J_best - 1)/J_best * ((nrows_x + nrows_tile - 1)/nrows_tile) * nchannels; - return ggml_cuda_mmq_tiling_is_efficient(ntiles, nsm); + // Same condition as launch_mul_mat_q uses to run the stream-k kernel with one block per tile. + const int64_t tiles_nwaves = (ntiles + nsm - 1) / nsm; + const int64_t tiles_efficiency_percent = 100 * ntiles / (nsm*tiles_nwaves); + return tiles_efficiency_percent >= 90; } diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 8ef4a7e56b05..007ada77071b 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -375,12 +375,6 @@ static __host__ bool ggml_cuda_mmq_needs_fallback(const int64_t nrows_x) { return nrows_x % 128 != 0; } -// Whether one CUDA block per tile (tiling) fills the SMs well enough that stream-k is not needed. -static __host__ bool ggml_cuda_mmq_tiling_is_efficient(const int64_t ntiles, const int nsm) { - const int64_t nwaves = (ntiles + nsm - 1) / nsm; - return 100*ntiles / (nsm*nwaves) >= 90; -} - static constexpr __device__ int ggml_cuda_mmq_get_rows_per_warp(ggml_type type, int J, bool fallback, ggml_prec prec_src1) { return ggml_cuda_mmq_get_config(type, J, fallback, prec_src1).rows_per_warp(); } @@ -1518,11 +1512,13 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a // For the stream-k kernel it is possible to run it with tiling by setting the number of CUDA blocks equal to the number of tiles. // This is worthwhile if the efficiency of tiling is high and skipping the fixup kernel is more important. const int ntiles_dst = ntx * nty * ntzw; + const int tiles_nwaves = (ntiles_dst + nsm - 1) / nsm; + const int tiles_efficiency_percent = 100 * ntiles_dst / (nsm*tiles_nwaves); GGML_ASSERT(ntiles_dst * blocks_per_ne00_fd.z < (1 << 30)); // Assert that variable kbc will not overflow. const int nwork = ntiles_dst * ((args.ncols_x + config.K_vram - 1) / config.K_vram); // Maximum number of CUDA blocks that can do useful work. - const dim3 block_nums_stream_k(GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_mmq_tiling_is_efficient(ntiles_dst, nsm) ? ntiles_dst : std::min(nsm, nwork), 1, 1); + const dim3 block_nums_stream_k(GGML_CUDA_CC_IS_NVIDIA(cc) && tiles_efficiency_percent >= 90 ? ntiles_dst : std::min(nsm, nwork), 1, 1); // The GLU is nonlinear, so the fused kernel needs the whole K range of a tile in one block. GGML_ASSERT(!has_fusion || int(block_nums_stream_k.x) == ntiles_dst); From 352d51066600266497b5962ca89d110aa9c64a4b Mon Sep 17 00:00:00 2001 From: Anuj Jalota Date: Tue, 20 Oct 2026 18:25:24 +0530 Subject: [PATCH 09/11] tests: one fused gate/up MMQ case per type and GLU op --- tests/test-backend-ops.cpp | 18 +++++++----------- 1 file changed, 7 insertions(+), 11 deletions(-) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index d599454bc4f8..ccc33de45afc 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -12034,17 +12034,13 @@ static std::vector> make_test_cases_eval() { false, 16, 8, false, false, true, false, { 1, 1 })); } - // Fused gate/up + GLU on the MMQ path, which is only used if its tiling fills the GPU: >= 2048 tiles, enough for up to 252 SMs. - // One MMA tile per warp (m = 17, 1 column tile) and two (m = 1024, 8 column tiles); k = 512 for more than one iteration over k. - for (ggml_type type : { GGML_TYPE_Q4_K, GGML_TYPE_Q5_K, GGML_TYPE_Q6_K, GGML_TYPE_Q8_0, GGML_TYPE_IQ4_XS }) { - test_cases.emplace_back(new test_mul_mat_vec_fusion(type, GGML_GLU_OP_SWIGLU, 17, 131072, 256, false, 1, 1, false, false, true, false, { 1, 1 })); - test_cases.emplace_back(new test_mul_mat_vec_fusion(type, GGML_GLU_OP_SWIGLU, 1024, 16384, 256, false, 1, 1, false, false, true, false, { 1, 1 })); - } - test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_K, GGML_GLU_OP_SWIGLU, 1024, 16384, 512, false, 1, 1, false, false, true, false, { 1, 1 })); - for (ggml_glu_op glu_op : { GGML_GLU_OP_GEGLU, GGML_GLU_OP_SWIGLU_OAI, GGML_GLU_OP_SWIGLU_CLAMP }) { - test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_K, glu_op, 1024, 16384, 256, false, 1, 1, false, false, true, false, { 1, 1 })); - } - test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_K, GGML_GLU_OP_SWIGLU, 1024, 2816, 256, false, 1, 1, false, false, true, false, { 2, 3 })); + // Fused gate/up + GLU on the MMQ path, only used if its tiling is efficient: ~1280 tiles per case, enough for up to 158 SMs and for 170 or 188 SMs. + // One case per fused type and GLU op. m = 17: one MMA tile per warp; m = 150: two, and two column tiles with a partial last one; k = 512: two k iterations. + test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_K, GGML_GLU_OP_SWIGLU, 17, 81920, 256, false, 1, 1, false, false, true, false, { 1, 1 })); + test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q5_K, GGML_GLU_OP_GEGLU, 150, 40960, 256, false, 1, 1, false, false, true, false, { 1, 1 })); + test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q6_K, GGML_GLU_OP_SWIGLU_OAI, 150, 40960, 256, false, 1, 1, false, false, true, false, { 1, 1 })); + test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q8_0, GGML_GLU_OP_SWIGLU_CLAMP, 150, 6912, 256, false, 1, 1, false, false, true, false, { 2, 3 })); + test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_IQ4_XS, GGML_GLU_OP_SWIGLU, 150, 40960, 512, false, 1, 1, false, false, true, false, { 1, 1 })); for (auto gate : {GATING_FUNC_SOFTMAX, GATING_FUNC_SIGMOID, GATING_FUNC_SOFTMAX_WEIGHT, GATING_FUNC_SQRT_SOFTPLUS}) { for (bool with_norm : {false, true}) { From 7d9a32d48284db65d4743291f3f03bda55b7d579 Mon Sep 17 00:00:00 2001 From: Anuj Jalota Date: Tue, 20 Oct 2026 23:55:29 +0530 Subject: [PATCH 10/11] cuda: always launch the fused MMQ kernel with one block per tile --- ggml/src/ggml-cuda/mmq.cuh | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 007ada77071b..55375f92efee 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -1477,7 +1477,7 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); if constexpr (!has_fusion) { - CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); + CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); } const int I_dst = has_fusion ? config.I/2 : config.I; // Rows of dst per tile, see mul_mat_q. @@ -1518,9 +1518,8 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a GGML_ASSERT(ntiles_dst * blocks_per_ne00_fd.z < (1 << 30)); // Assert that variable kbc will not overflow. const int nwork = ntiles_dst * ((args.ncols_x + config.K_vram - 1) / config.K_vram); // Maximum number of CUDA blocks that can do useful work. - const dim3 block_nums_stream_k(GGML_CUDA_CC_IS_NVIDIA(cc) && tiles_efficiency_percent >= 90 ? ntiles_dst : std::min(nsm, nwork), 1, 1); - // The GLU is nonlinear, so the fused kernel needs the whole K range of a tile in one block. - GGML_ASSERT(!has_fusion || int(block_nums_stream_k.x) == ntiles_dst); + // The GLU is nonlinear, so the fused kernel always needs the whole K range of a tile in one block. + const dim3 block_nums_stream_k(has_fusion || (GGML_CUDA_CC_IS_NVIDIA(cc) && tiles_efficiency_percent >= 90) ? ntiles_dst : std::min(nsm, nwork), 1, 1); const bool fixup_needed = ntiles_dst % block_nums_stream_k.x != 0; From 6e34e56db9384dfd9f247bd6af5bf5c04bf93f7e Mon Sep 17 00:00:00 2001 From: Anuj Jalota Date: Tue, 20 Oct 2026 23:55:30 +0530 Subject: [PATCH 11/11] tests: shorten the fused gate/up MMQ test comment --- tests/test-backend-ops.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index ccc33de45afc..42dabc1bc9e1 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -12034,8 +12034,7 @@ static std::vector> make_test_cases_eval() { false, 16, 8, false, false, true, false, { 1, 1 })); } - // Fused gate/up + GLU on the MMQ path, only used if its tiling is efficient: ~1280 tiles per case, enough for up to 158 SMs and for 170 or 188 SMs. - // One case per fused type and GLU op. m = 17: one MMA tile per warp; m = 150: two, and two column tiles with a partial last one; k = 512: two k iterations. + // Fused gate/up + GLU on the MMQ path, only used if its tiling is efficient: ~1280 tiles per case, one case per type and GLU op. test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_K, GGML_GLU_OP_SWIGLU, 17, 81920, 256, false, 1, 1, false, false, true, false, { 1, 1 })); test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q5_K, GGML_GLU_OP_GEGLU, 150, 40960, 256, false, 1, 1, false, false, true, false, { 1, 1 })); test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q6_K, GGML_GLU_OP_SWIGLU_OAI, 150, 40960, 256, false, 1, 1, false, false, true, false, { 1, 1 }));