From a70d5dd1b74c00c5aca076c9f614ad520c0f2a2e Mon Sep 17 00:00:00 2001 From: lichang Date: Mon, 22 Jun 2026 20:06:38 -0600 Subject: [PATCH 01/22] Add kernel dimension logging for profiling analysis MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Instrumented CUDA kernels to log actual M×N×K dimensions to stderr for performance analysis and TFLOPS/bandwidth calculations. Changes: - quantize.cu: Log M, K, grid for quantize_mmq_q8_1 kernel - mmq.cu: Log M, N, K, ne dimensions for mul_mat_q (GEMM) kernel - mmvq.cu: Log M, K dimensions for mul_mat_vec_q (GEMV) kernel Usage: llama-bench ... 2>&1 | tee profile.log The logged dimensions are matched with rocprofv3 timing data to calculate: - TFLOPS: From actual M×N×K dimensions (2×M×N×K FLOPs) - Bandwidth: From actual data transfer patterns This instrumentation enabled accurate performance analysis showing: - GEMM kernels: 8.83-43.58 TFLOPS (varies by quantization type) - Quantize kernels: 284 GB/s bandwidth (data in L2 cache) - GEMV kernels: 0.05-5.52 TFLOPS (memory-bound decode) Co-Authored-By: Claude Sonnet 4 --- ggml/src/ggml-cuda/mmq.cu | 10 ++++++++++ ggml/src/ggml-cuda/mmvq.cu | 5 +++++ ggml/src/ggml-cuda/quantize.cu | 5 +++++ 3 files changed, 20 insertions(+) diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index e1add5e03316..f71c134e7591 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -157,6 +157,12 @@ void ggml_cuda_mul_mat_q( ne02, ne12, s02, s12, s2, ne03, ne13, s03, s13, s3, use_stream_k, ne1}; + + // Log kernel dimensions for profiling analysis + // Matrix operation: C[M,N] = A[M,K] × B[K,N] where M=ne1, N=ne00, K=ne10 + fprintf(stderr, "[MUL_MAT_Q] M=%ld N=%ld K=%ld ne00=%ld ne01=%ld ne1=%ld ne11=%ld ne02=%ld ne12=%ld ne03=%ld ne13=%ld type=%d\n", + ne1, ne00, ne10, ne00, ne01, ne1, ne11, ne02, ne12, ne03, ne13, src0->type); + ggml_cuda_mul_mat_q_switch_type(ctx, args, stream); return; } @@ -219,6 +225,10 @@ void ggml_cuda_mul_mat_q( ne03, ne13, s03, s13, s3, use_stream_k, ne12}; + // Log kernel dimensions for profiling analysis (batched/MoE path) + fprintf(stderr, "[MUL_MAT_Q_MoE] M=%ld N=%ld K=%ld ne_get_rows=%ld ne00=%ld ne01=%ld ne02=%ld ne12=%ld type=%d\n", + ne_get_rows, ne00, ne10, ne_get_rows, ne00, ne01, ne02, ne12, src0->type); + ggml_cuda_mul_mat_q_switch_type(ctx, args, stream); } diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index fe44a58da918..097e8d49abcc 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -1213,6 +1213,11 @@ void ggml_cuda_mul_mat_vec_q( const int64_t ids_stride = ids ? ids->nb[1] / ggml_type_size(ids->type) : 0; + // Log kernel dimensions for profiling analysis + // Matrix-vector operation: y[M] = A[M,K] × x[K] where M=ne00, K=ne10 + fprintf(stderr, "[MUL_MAT_VEC_Q] M=%ld K=%ld ne00=%ld ne01=%ld ne10=%ld ne11=%ld ne02=%ld ne12=%ld ne03=%ld ne13=%ld type=%d\n", + ne00, ne10, ne00, ne01, ne10, ne11, ne02, ne12, ne03, ne13, src0->type); + mul_mat_vec_q_switch_type( src0->data, src0->type, src1_q8_1.get(), ids_d, fusion_local, dst_d, ne00, ne01, ncols_dst, s01, stride_col_y, stride_col_dst, diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu index 39a500a17041..c3c401bf9042 100644 --- a/ggml/src/ggml-cuda/quantize.cu +++ b/ggml/src/ggml-cuda/quantize.cu @@ -400,6 +400,11 @@ void quantize_mmq_q8_1_cuda( const int64_t block_num_y = (ne0 + 4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ - 1) / (4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ); const dim3 num_blocks(ne1, block_num_y, ne2*ne3); const dim3 block_size(CUDA_QUANTIZE_BLOCK_SIZE_MMQ, 1, 1); + + // Log kernel dimensions for profiling analysis + fprintf(stderr, "[QUANTIZE_MMQ_Q8_1] ne00=%ld ne0=%ld ne1=%ld ne2=%ld ne3=%ld grid=(%u,%u,%u) type=%d\n", + ne00, ne0, ne1, ne2, ne3, num_blocks.x, num_blocks.y, num_blocks.z, type_src0); + switch (mmq_get_q8_1_ds_layout(type_src0)) { case MMQ_Q8_1_DS_LAYOUT_D4: quantize_mmq_q8_1 From 3c63ba2263c6f8a6637e3a082b42971b9869cfe2 Mon Sep 17 00:00:00 2001 From: lichang Date: Tue, 23 Jun 2026 13:04:08 -0600 Subject: [PATCH 02/22] Fix M/N/K dimension logging for MMQ and MMVQ profiling. Log GEMM as M=ne11, N=ne01, K=ne00 and GEMV as M=ne01, K=ne00 so instrumented shapes match actual matmul dimensions. Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cu | 7 ++++--- ggml/src/ggml-cuda/mmvq.cu | 5 +++-- 2 files changed, 7 insertions(+), 5 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index f71c134e7591..3c3272501968 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -159,9 +159,10 @@ void ggml_cuda_mul_mat_q( use_stream_k, ne1}; // Log kernel dimensions for profiling analysis - // Matrix operation: C[M,N] = A[M,K] × B[K,N] where M=ne1, N=ne00, K=ne10 + // dst[ne01, ne11] = src0[ne00, ne01] @ src1[ne10, ne11]; ne00==ne10 (K) + // GEMM C[M,N] = W[N,K] @ X[K,M]: M=ne11, N=ne01, K=ne00 fprintf(stderr, "[MUL_MAT_Q] M=%ld N=%ld K=%ld ne00=%ld ne01=%ld ne1=%ld ne11=%ld ne02=%ld ne12=%ld ne03=%ld ne13=%ld type=%d\n", - ne1, ne00, ne10, ne00, ne01, ne1, ne11, ne02, ne12, ne03, ne13, src0->type); + ne11, ne01, ne00, ne00, ne01, ne1, ne11, ne02, ne12, ne03, ne13, src0->type); ggml_cuda_mul_mat_q_switch_type(ctx, args, stream); return; @@ -227,7 +228,7 @@ void ggml_cuda_mul_mat_q( // Log kernel dimensions for profiling analysis (batched/MoE path) fprintf(stderr, "[MUL_MAT_Q_MoE] M=%ld N=%ld K=%ld ne_get_rows=%ld ne00=%ld ne01=%ld ne02=%ld ne12=%ld type=%d\n", - ne_get_rows, ne00, ne10, ne_get_rows, ne00, ne01, ne02, ne12, src0->type); + ne11_flat, ne01, ne00, ne_get_rows, ne00, ne01, ne02, ne12, src0->type); ggml_cuda_mul_mat_q_switch_type(ctx, args, stream); } diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 097e8d49abcc..d94db2993685 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -1214,9 +1214,10 @@ void ggml_cuda_mul_mat_vec_q( const int64_t ids_stride = ids ? ids->nb[1] / ggml_type_size(ids->type) : 0; // Log kernel dimensions for profiling analysis - // Matrix-vector operation: y[M] = A[M,K] × x[K] where M=ne00, K=ne10 + // dst[ne01] = src0[ne00, ne01] @ src1[ne10]; ne00==ne10 (K), ne11==1 for decode + // GEMV y[M] = W[M,K] @ x[K]: M=ne01, K=ne00 fprintf(stderr, "[MUL_MAT_VEC_Q] M=%ld K=%ld ne00=%ld ne01=%ld ne10=%ld ne11=%ld ne02=%ld ne12=%ld ne03=%ld ne13=%ld type=%d\n", - ne00, ne10, ne00, ne01, ne10, ne11, ne02, ne12, ne03, ne13, src0->type); + ne01, ne00, ne00, ne01, ne10, ne11, ne02, ne12, ne03, ne13, src0->type); mul_mat_vec_q_switch_type( src0->data, src0->type, src1_q8_1.get(), ids_d, fusion_local, dst_d, ne00, From 6051c1f358922cb38301fd42e3ec91ab0705c725 Mon Sep 17 00:00:00 2001 From: lichang Date: Tue, 23 Jun 2026 14:37:36 -0600 Subject: [PATCH 03/22] Improve RDNA3.5 MMQ occupancy by reducing mmq_x when LDS limits WGs/CU. On gfx115x, mmq_x=128 uses more than half of per-CU LDS and caps occupancy at one workgroup per CU. Select a smaller mmq_x that fits two workgroups without doubling M-tiles, improving P0 prefill TFLOPS ~25% on gfx1151. Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 54 ++++++++++++++++++++++++++++++++++++++ 1 file changed, 54 insertions(+) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index b58ac9e7b428..5dea7a788ca1 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -6,6 +6,8 @@ #include #include +#include +#include using namespace ggml_cuda_mma; @@ -3957,6 +3959,7 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a const int warp_size = ggml_cuda_info().devices[id].warp_size; const int nwarps = mmq_get_nwarps_host(cc, warp_size); const int mmq_y = get_mmq_y_host(cc); + const size_t smpbo = ggml_cuda_info().devices[id].smpbo; const dim3 block_dims(warp_size, nwarps, 1); @@ -3970,6 +3973,19 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a const int ntzw = args.nchannels_y * args.nsamples_y; const dim3 block_nums_xy_tiling(nty, ntx, ntzw); + const auto log_launch_config = [&](const dim3 & grid, const bool need_check, const char * path) { + fprintf(stderr, + "[MUL_MAT_Q_LAUNCH] type=%d mmq_x=%d mmq_y=%d nwarps=%d warp_size=%d " + "nbytes_shared=%d smpbo=%zu grid=(%u,%u,%u) block=(%u,%u,%u) " + "ntx=%d nty=%d ntzw=%d ncols_max=%ld nrows_x=%ld use_stream_k=%d need_check=%d path=%s\n", + type, mmq_x, mmq_y, nwarps, warp_size, + nbytes_shared, smpbo, + grid.x, grid.y, grid.z, + block_dims.x, block_dims.y, block_dims.z, + ntx, nty, ntzw, args.ncols_max, args.nrows_x, + args.use_stream_k ? 1 : 0, need_check ? 1 : 0, path); + }; + GGML_ASSERT(args.nchannels_y % args.nchannels_x == 0); GGML_ASSERT(args.nsamples_y % args.nsamples_x == 0); const int channel_ratio = args.nchannels_y / args.nchannels_x; @@ -3985,6 +4001,7 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a if (!args.use_stream_k) { if (args.nrows_x % mmq_y == 0) { constexpr bool need_check = false; + log_launch_config(block_nums_xy_tiling, need_check, "tiling"); mul_mat_q<<>> (args.x, args.y, args.ids_dst, args.expert_bounds, args.dst, nullptr, blocks_per_ne00_fd, args.nrows_x, args.ncols_dst, args.stride_row_x, args.ncols_y, args.nrows_dst, @@ -3993,6 +4010,7 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a ntx_fd); } else { constexpr bool need_check = true; + log_launch_config(block_nums_xy_tiling, need_check, "tiling"); mul_mat_q<<>> (args.x, args.y, args.ids_dst, args.expert_bounds, args.dst, nullptr, blocks_per_ne00_fd, args.nrows_x, args.ncols_dst, args.stride_row_x, args.ncols_y, args.nrows_dst, @@ -4025,6 +4043,7 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a if (args.nrows_x % mmq_y == 0) { constexpr bool need_check = false; + log_launch_config(block_nums_stream_k, need_check, "stream_k"); mul_mat_q<<>> (args.x, args.y, args.ids_dst, args.expert_bounds, args.dst, tmp_fixup.ptr, blocks_per_ne00_fd, args.nrows_x, args.ncols_dst, args.stride_row_x, args.ncols_y, args.nrows_dst, @@ -4043,6 +4062,7 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a ntx_fd); } else { constexpr bool need_check = true; + log_launch_config(block_nums_stream_k, need_check, "stream_k"); mul_mat_q<<>> (args.x, args.y, args.ids_dst, args.expert_bounds, args.dst, tmp_fixup.ptr, blocks_per_ne00_fd, args.nrows_x, args.ncols_dst, args.stride_row_x, args.ncols_y, args.nrows_dst, @@ -4091,6 +4111,40 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda } } +#if defined(GGML_USE_HIP) + // RDNA3.5 (gfx115x): mmq_x=128 uses ~59% of LDS/CU → 1 WG/CU. Pick the smallest + // mmq_x with nbytes_shared <= smpbo/2 (2 WGs/CU) without doubling M-tile count. + if (GGML_CUDA_CC_IS_RDNA3_5(cc) && mmq_x_best > 0) { + const size_t lds_dual_wg = smpbo / 2; + const size_t nbytes_best = mmq_get_nbytes_shared(mmq_x_best, mmq_y, cc, warp_size, nwarps); + if (nbytes_best > lds_dual_wg) { + int mmq_x_lds = 0; + size_t nbytes_lds_min = SIZE_MAX; + for (int mmq_x = 8; mmq_x <= mmq_x_max; mmq_x += 8) { + const int granularity = mmq_get_granularity_host(mmq_x, cc); + if (mmq_x % granularity != 0) { + continue; + } + const size_t nbytes = mmq_get_nbytes_shared(mmq_x, mmq_y, cc, warp_size, nwarps); + if (nbytes > smpbo || nbytes > lds_dual_wg) { + continue; + } + const int ntiles_x = (args.ncols_max + mmq_x - 1) / mmq_x; + if (ntiles_x > ntiles_x_best * 2) { + continue; + } + if (nbytes < nbytes_lds_min) { + mmq_x_lds = mmq_x; + nbytes_lds_min = nbytes; + } + } + if (mmq_x_lds > 0) { + mmq_x_best = mmq_x_lds; + } + } + } +#endif // GGML_USE_HIP + switch (mmq_x_best) { case 8: launch_mul_mat_q(ctx, args, stream); From 36249b0c9c4db15db04278fe93a51b5d52aa956d Mon Sep 17 00:00:00 2001 From: lichang Date: Tue, 23 Jun 2026 15:36:58 -0600 Subject: [PATCH 04/22] Improve RDNA3.5 MMQ prefill via LDS layout and activation load pipelining. Restore WMMA granularity at mmq_x>=64, widen tile_y padding to cut LDS bank conflicts, and software-pipeline Q8_1 activation loads into registers during vec_dot to hide global memory latency on gfx115x. Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 171 ++++++++++++++++++++++++++++++++----- 1 file changed, 149 insertions(+), 22 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 5dea7a788ca1..09cd770583ba 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -8,6 +8,7 @@ #include #include #include +#include using namespace ggml_cuda_mma; @@ -277,8 +278,26 @@ static constexpr __host__ __device__ int mmq_get_mma_tile_x_k(ggml_type type) { #define MMQ_TILE_Y_K (MMQ_TILE_NE_K + MMQ_TILE_NE_K / QI8_1) #define MMQ_TILE_Y_FP4_K MMQ_TILE_Y_K +static int mmq_get_tile_y_k_padded_host(const int mmq_x, const int cc, const int warp_size, const int nwarps) { + const int pad = GGML_CUDA_CC_IS_RDNA3_5(cc) ? 2*nwarps*warp_size : nwarps*warp_size; + return GGML_PAD(mmq_x*MMQ_TILE_Y_K, pad); +} + +#if defined(RDNA3_5) +static constexpr __device__ int mmq_get_tile_y_k_padded_device(const int mmq_x, const int nwarps, const int warp_size) { + return GGML_PAD(mmq_x*MMQ_TILE_Y_K, 2*nwarps*warp_size); +} +#else +static constexpr __device__ int mmq_get_tile_y_k_padded_device(const int mmq_x, const int nwarps, const int warp_size) { + return GGML_PAD(mmq_x*MMQ_TILE_Y_K, nwarps*warp_size); +} +#endif // RDNA3_5 + static int mmq_get_granularity_host(const int mmq_x, const int cc) { if (amd_mfma_available(cc) || amd_wmma_available(cc)) { + if (GGML_CUDA_CC_IS_RDNA3_5(cc) && mmq_x >= 64) { + return 32; + } return mmq_x >= 128 ? 32 : 16; } else if (turing_mma_available(cc) && mmq_x >= 48) { return 16; @@ -289,7 +308,11 @@ static int mmq_get_granularity_host(const int mmq_x, const int cc) { #if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) static constexpr __device__ int mmq_get_granularity_device(const int mmq_x) { +#if defined(RDNA3_5) + return mmq_x >= 64 ? 32 : 16; +#else return mmq_x >= 128 ? 32 : 16; +#endif // RDNA3_5 } #elif defined(TURING_MMA_AVAILABLE) static constexpr __device__ int mmq_get_granularity_device(const int mmq_x) { @@ -3455,6 +3478,40 @@ struct mmq_type_traits { static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_q8_0_q8_1_dp4a; }; +#if defined(RDNA3_5) +// Software-pipeline activation tile loads: issue global loads into registers, run WMMA, +// then store to LDS. Overlaps memory latency with vec_dot on gfx115x (no ISA prefetch). +template +static __device__ __forceinline__ void mmq_tile_y_load_global( + int * __restrict__ tile_y, const int * __restrict__ by) { +#pragma unroll + for (int c = 0; c < nchunks; ++c) { + const int l = c*(nwarps*warp_size) + threadIdx.y*warp_size + threadIdx.x; + tile_y[l] = by[l]; + } +} + +template +static __device__ __forceinline__ void mmq_tile_y_load_global_to_regs( + const int * __restrict__ by, int (&cache)[nchunks]) { +#pragma unroll + for (int c = 0; c < nchunks; ++c) { + const int l = c*(nwarps*warp_size) + threadIdx.y*warp_size + threadIdx.x; + cache[c] = by[l]; + } +} + +template +static __device__ __forceinline__ void mmq_tile_y_store_regs( + int * __restrict__ tile_y, const int (&cache)[nchunks]) { +#pragma unroll + for (int c = 0; c < nchunks; ++c) { + const int l = c*(nwarps*warp_size) + threadIdx.y*warp_size + threadIdx.x; + tile_y[l] = cache[c]; + } +} +#endif // RDNA3_5 + template static __device__ __forceinline__ void mul_mat_q_process_tile( const char * __restrict__ x, const int offset_x, const int * __restrict__ y, @@ -3470,7 +3527,7 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( extern __shared__ int data_mul_mat_q[]; int * tile_y = data_mul_mat_q + mmq_x; - int * tile_x = tile_y + GGML_PAD(mmq_x*MMQ_TILE_Y_K, nwarps*warp_size); + int * tile_x = tile_y + mmq_get_tile_y_k_padded_device(mmq_x, nwarps, warp_size); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) constexpr vec_dot_mmq_t vec_dot = mmq_type_traits::vec_dot_mma; @@ -3494,39 +3551,93 @@ 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); - { - const int * by0 = y + ncols_y * (kb0 * qk / ne_block) * sz; -#pragma unroll - for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { - int l = l0 + threadIdx.y*warp_size + threadIdx.x; +#if defined(RDNA3_5) + constexpr int tile_y_elems = mmq_x*MMQ_TILE_Y_K; + constexpr int tile_y_load_stride = nwarps*warp_size; + if constexpr (mmq_x <= 64 && tile_y_elems % tile_y_load_stride == 0) { + constexpr int tile_y_nchunks = tile_y_elems/tile_y_load_stride; + + int y0_next_cache[tile_y_nchunks]; + bool have_y0_prefetch = false; + + 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); - tile_y[l] = by0[l]; + const int yk = kb0 * qk / ne_block; + const int * by0 = y + ncols_y * yk * sz; + const int * by1 = y + ncols_y * (yk + 1) * sz; + + if (have_y0_prefetch) { + mmq_tile_y_store_regs(tile_y, y0_next_cache); + have_y0_prefetch = false; + } else { + mmq_tile_y_load_global(tile_y, by0); } + + __syncthreads(); + + int y1_cache[tile_y_nchunks]; + mmq_tile_y_load_global_to_regs(by1, y1_cache); + + vec_dot(tile_x, tile_y, sum, 0); + + __syncthreads(); + + mmq_tile_y_store_regs(tile_y, y1_cache); + + __syncthreads(); + + const int kb0_next = kb0 + blocks_per_iter; + if (kb0_next < kb0_stop) { + const int * by0_next = y + ncols_y * (kb0_next * qk / ne_block) * sz; + mmq_tile_y_load_global_to_regs(by0_next, y0_next_cache); + have_y0_prefetch = true; + } + + vec_dot(tile_x, tile_y, sum, MMQ_TILE_NE_K); + + __syncthreads(); } + } else +#endif // RDNA3_5 + { + 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); - __syncthreads(); + const int yk = kb0 * qk / ne_block; + const int * by0 = y + ncols_y * yk * sz; + const int * by1 = y + ncols_y * (yk + 1) * sz; - vec_dot(tile_x, tile_y, sum, 0); + { +#pragma unroll + for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { + int l = l0 + threadIdx.y*warp_size + threadIdx.x; - __syncthreads(); + tile_y[l] = by0[l]; + } + } - { - const int * by0 = y + ncols_y * ((kb0 * qk / ne_block) * sz + sz); + __syncthreads(); + + vec_dot(tile_x, tile_y, sum, 0); + + __syncthreads(); + + { #pragma unroll - for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { - int l = l0 + threadIdx.y*warp_size + threadIdx.x; + for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { + int l = l0 + threadIdx.y*warp_size + threadIdx.x; - tile_y[l] = by0[l]; + tile_y[l] = by1[l]; + } } - } - __syncthreads(); + __syncthreads(); - vec_dot(tile_x, tile_y, sum, MMQ_TILE_NE_K); + vec_dot(tile_x, tile_y, sum, MMQ_TILE_NE_K); - __syncthreads(); + __syncthreads(); + } } if (fixup) { @@ -3948,7 +4059,10 @@ static size_t mmq_get_nbytes_shared(const int mmq_x, const int mmq_y, const int const size_t nbs_ids = mmq_x*sizeof(int); const size_t nbs_x = (turing_mma_available(cc) || amd_mfma_available(cc) || amd_wmma_available(cc)) ? mmq_y*mmq_tile_x_k*sizeof(int) : txs.qs*sizeof(int) + txs.dm*sizeof(half2) + txs.sc*sizeof(int); const size_t nbs_y = mmq_x * (sizeof(block_q8_1_mmq)); - return nbs_ids + nbs_x + GGML_PAD(nbs_y, nwarps*warp_size*sizeof(int)); + const int tile_y_k_padded = mmq_get_tile_y_k_padded_host(mmq_x, cc, warp_size, nwarps); + const size_t nbs_y_padded = std::max(nbs_y, (size_t) tile_y_k_padded*sizeof(int)); + const int pad = GGML_CUDA_CC_IS_RDNA3_5(cc) ? 2*nwarps*warp_size : nwarps*warp_size; + return nbs_ids + nbs_x + GGML_PAD(nbs_y_padded, pad*sizeof(int)); } template @@ -4145,6 +4259,19 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda } #endif // GGML_USE_HIP +#if defined(GGML_USE_HIP) + if (GGML_CUDA_CC_IS_RDNA3_5(cc)) { + if (const char * force = getenv("MMQ_FORCE_MMQ_X")) { + const int mmq_x_force = atoi(force); + if (mmq_x_force >= 8 && mmq_x_force <= mmq_x_max && + mmq_x_force % mmq_get_granularity_host(mmq_x_force, cc) == 0 && + mmq_get_nbytes_shared(mmq_x_force, mmq_y, cc, warp_size, nwarps) <= smpbo) { + mmq_x_best = mmq_x_force; + } + } + } +#endif // GGML_USE_HIP + switch (mmq_x_best) { case 8: launch_mul_mat_q(ctx, args, stream); From 7248ec97ca8a04c8942e299a8f7bade3252c7816 Mon Sep 17 00:00:00 2001 From: lichang Date: Tue, 23 Jun 2026 16:50:41 -0600 Subject: [PATCH 05/22] Add opt-in HIP MMQ phase profiling to split dequant vs MMA time. Enables MMQ_PROFILE_PHASES=1 to report load_tiles, y_act, and vec_dot cycle shares per launch, guiding RDNA3.5 prefill optimization work. Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 232 +++++++++++++++++++++++++------ ggml/src/ggml-cuda/vendors/hip.h | 1 + 2 files changed, 189 insertions(+), 44 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 09cd770583ba..ea014766318e 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -12,6 +12,114 @@ using namespace ggml_cuda_mma; +#if defined(GGML_USE_HIP) +enum mmq_profile_phase { + MMQ_PROFILE_LOAD_TILES = 0, + MMQ_PROFILE_Y_ACT = 1, + MMQ_PROFILE_VEC_DOT = 2, + MMQ_PROFILE_KB_ITERS = 3, + MMQ_PROFILE_N = 4, +}; + +static __device__ __forceinline__ bool mmq_profile_want_sample(uint64_t * const profile) { + return profile && blockIdx.x == 0 && blockIdx.y == 0 && blockIdx.z == 0 + && threadIdx.x == 0 && threadIdx.y == 0; +} + +static __device__ __forceinline__ void mmq_profile_kb_iter(uint64_t * const profile) { + if (!mmq_profile_want_sample(profile)) { + return; + } + ++profile[MMQ_PROFILE_KB_ITERS]; +} + +static __device__ __forceinline__ void mmq_profile_phase_begin(uint64_t * const profile, uint64_t & t0) { + if (!mmq_profile_want_sample(profile)) { + return; + } + t0 = clock64(); +} + +static __device__ __forceinline__ void mmq_profile_phase_end(uint64_t * const profile, const uint64_t t0, const int phase) { + if (!mmq_profile_want_sample(profile)) { + return; + } + profile[phase] += clock64() - t0; +} + +static bool mmq_profile_phases_enabled() { + static int enabled = -1; + if (enabled < 0) { + enabled = getenv("MMQ_PROFILE_PHASES") ? 1 : 0; + } + return enabled; +} + +struct mmq_profile_guard { + ggml_cuda_pool_alloc buf; + cudaStream_t stream = nullptr; + int type = 0; + int mmq_x = 0; + int64_t nrows_x = 0; + int64_t ncols_max = 0; + bool active = false; + + mmq_profile_guard(ggml_cuda_pool & pool, const bool enable, const int type_in, const int mmq_x_in, + const int64_t nrows_x_in, const int64_t ncols_max_in, cudaStream_t stream_in) + : buf(pool), stream(stream_in), type(type_in), mmq_x(mmq_x_in), nrows_x(nrows_x_in), ncols_max(ncols_max_in), active(enable) { + if (!active) { + return; + } + buf.alloc(MMQ_PROFILE_N); + CUDA_CHECK(cudaMemset(buf.get(), 0, MMQ_PROFILE_N*sizeof(uint64_t))); + } + + uint64_t * ptr() { + return active ? buf.get() : nullptr; + } + + ~mmq_profile_guard() { + if (!active) { + return; + } + CUDA_CHECK(cudaStreamSynchronize(stream)); + uint64_t h[MMQ_PROFILE_N] = {}; + CUDA_CHECK(cudaMemcpy(h, buf.get(), sizeof(h), cudaMemcpyDeviceToHost)); + const uint64_t total = h[MMQ_PROFILE_LOAD_TILES] + h[MMQ_PROFILE_Y_ACT] + h[MMQ_PROFILE_VEC_DOT]; + if (total > 0) { + fprintf(stderr, + "[MMQ_PROFILE] type=%d mmq_x=%d nrows_x=%ld ncols_max=%ld " + "load_tiles=%.1f%% y_act=%.1f%% vec_dot=%.1f%% " + "kb_iters=%llu cycles_load=%llu cycles_y=%llu cycles_vec=%llu cycles_total=%llu\n", + type, mmq_x, (long) nrows_x, (long) ncols_max, + 100.0*h[MMQ_PROFILE_LOAD_TILES]/total, + 100.0*h[MMQ_PROFILE_Y_ACT]/total, + 100.0*h[MMQ_PROFILE_VEC_DOT]/total, + (unsigned long long) h[MMQ_PROFILE_KB_ITERS], + (unsigned long long) h[MMQ_PROFILE_LOAD_TILES], + (unsigned long long) h[MMQ_PROFILE_Y_ACT], + (unsigned long long) h[MMQ_PROFILE_VEC_DOT], + (unsigned long long) total); + } + } +}; +#else +static __device__ __forceinline__ void mmq_profile_kb_iter(uint64_t *) {} + +static __device__ __forceinline__ void mmq_profile_phase_begin(uint64_t *, uint64_t &) {} + +static __device__ __forceinline__ void mmq_profile_phase_end(uint64_t *, const uint64_t, const int) {} + +static bool mmq_profile_phases_enabled() { + return false; +} + +struct mmq_profile_guard { + mmq_profile_guard(ggml_cuda_pool &, const bool, const int, const int, const int64_t, const int64_t, cudaStream_t) {} + uint64_t * ptr() { return nullptr; } +}; +#endif // GGML_USE_HIP + #define MMQ_DP4A_MAX_BATCH_SIZE 64 // Max. batch size to use for dp4a MMQ kernels when FP16 tensor cores are available. #define MMQ_ITER_K 256 #define MMQ_ITER_K_FP4 512 @@ -3517,7 +3625,8 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( const char * __restrict__ x, const int offset_x, const int * __restrict__ y, const int * __restrict__ ids_dst, float * __restrict__ dst, float * __restrict__ tmp_fixup, 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) { + const int tile_x_max_i, const int tile_y_max_j, const int kb0_start, const int kb0_stop, + uint64_t * mmq_profile) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = mmq_get_nwarps_device(); @@ -3561,82 +3670,114 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( bool have_y0_prefetch = false; 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); + mmq_profile_kb_iter(mmq_profile); + + { + uint64_t t0 = 0; + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); + load_tiles(x, tile_x, offset_x + kb0, tile_x_max_i, stride_row_x); + __syncthreads(); + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_LOAD_TILES); + } const int yk = kb0 * qk / ne_block; const int * by0 = y + ncols_y * yk * sz; const int * by1 = y + ncols_y * (yk + 1) * sz; - if (have_y0_prefetch) { - mmq_tile_y_store_regs(tile_y, y0_next_cache); - have_y0_prefetch = false; - } else { - mmq_tile_y_load_global(tile_y, by0); - } + { + uint64_t t0 = 0; + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); - __syncthreads(); + if (have_y0_prefetch) { + mmq_tile_y_store_regs(tile_y, y0_next_cache); + have_y0_prefetch = false; + } else { + mmq_tile_y_load_global(tile_y, by0); + } - int y1_cache[tile_y_nchunks]; - mmq_tile_y_load_global_to_regs(by1, y1_cache); + __syncthreads(); - vec_dot(tile_x, tile_y, sum, 0); + int y1_cache[tile_y_nchunks]; + mmq_tile_y_load_global_to_regs(by1, y1_cache); - __syncthreads(); + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_Y_ACT); - mmq_tile_y_store_regs(tile_y, y1_cache); + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); - __syncthreads(); + vec_dot(tile_x, tile_y, sum, 0); - const int kb0_next = kb0 + blocks_per_iter; - if (kb0_next < kb0_stop) { - const int * by0_next = y + ncols_y * (kb0_next * qk / ne_block) * sz; - mmq_tile_y_load_global_to_regs(by0_next, y0_next_cache); - have_y0_prefetch = true; - } + __syncthreads(); - vec_dot(tile_x, tile_y, sum, MMQ_TILE_NE_K); + mmq_tile_y_store_regs(tile_y, y1_cache); - __syncthreads(); + __syncthreads(); + + const int kb0_next = kb0 + blocks_per_iter; + if (kb0_next < kb0_stop) { + const int * by0_next = y + ncols_y * (kb0_next * qk / ne_block) * sz; + mmq_tile_y_load_global_to_regs(by0_next, y0_next_cache); + have_y0_prefetch = true; + } + + vec_dot(tile_x, tile_y, sum, MMQ_TILE_NE_K); + + __syncthreads(); + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_VEC_DOT); + } } } else #endif // RDNA3_5 { 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); + mmq_profile_kb_iter(mmq_profile); + + { + uint64_t t0 = 0; + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); + load_tiles(x, tile_x, offset_x + kb0, tile_x_max_i, stride_row_x); + __syncthreads(); + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_LOAD_TILES); + } const int yk = kb0 * qk / ne_block; const int * by0 = y + ncols_y * yk * sz; const int * by1 = y + ncols_y * (yk + 1) * sz; { + uint64_t t0 = 0; + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); #pragma unroll for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { int l = l0 + threadIdx.y*warp_size + threadIdx.x; tile_y[l] = by0[l]; } + __syncthreads(); + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_Y_ACT); } - __syncthreads(); - - vec_dot(tile_x, tile_y, sum, 0); - - __syncthreads(); - { + uint64_t t0 = 0; + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); + vec_dot(tile_x, tile_y, sum, 0); + __syncthreads(); #pragma unroll for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { int l = l0 + threadIdx.y*warp_size + threadIdx.x; tile_y[l] = by1[l]; } + __syncthreads(); + vec_dot(tile_x, tile_y, sum, MMQ_TILE_NE_K); + __syncthreads(); + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_VEC_DOT); } - - __syncthreads(); - - vec_dot(tile_x, tile_y, sum, MMQ_TILE_NE_K); - - __syncthreads(); } } @@ -3668,7 +3809,7 @@ static __global__ void mul_mat_q( 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) { + const uint3 ntx, uint64_t * mmq_profile) { // Skip unused template specializations for faster compilation: if (mmq_x > get_mmq_x_max_device() || mmq_x % mmq_get_granularity_device(mmq_x) != 0) { @@ -3753,7 +3894,7 @@ static __global__ void mul_mat_q( constexpr bool fixup = false; mul_mat_q_process_tile (x, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, stride_row_x, ncols_y, stride_col_dst, - tile_x_max_i, tile_y_max_j, 0, blocks_per_ne00.z); + tile_x_max_i, tile_y_max_j, 0, blocks_per_ne00.z, mmq_profile); return; } #endif // (defined(GGML_USE_HIP) && !defined(CDNA4) && !defined(CDNA3)) || __CUDA_ARCH__ < GGML_CUDA_CC_VOLTA @@ -3833,7 +3974,7 @@ 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, stride_row_x, ncols_y, stride_col_dst, - tile_x_max_i, tile_y_max_j, kb0_start, kb0_stop); + tile_x_max_i, tile_y_max_j, kb0_start, kb0_stop, mmq_profile); kbc += blocks_per_ne00.z; kbc -= fastmodulo(kbc, blocks_per_ne00); @@ -3902,7 +4043,7 @@ 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, stride_row_x, ncols_y, stride_col_dst, - tile_x_max_i, tile_y_max_j, kb0_start, kb0_stop); + tile_x_max_i, tile_y_max_j, kb0_start, kb0_stop, mmq_profile); } template @@ -4077,6 +4218,9 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a const dim3 block_dims(warp_size, nwarps, 1); + mmq_profile_guard prof_guard(ctx.pool(id), mmq_profile_phases_enabled(), type, mmq_x, + args.nrows_x, args.ncols_max, stream); + const int nbytes_shared = mmq_get_nbytes_shared(mmq_x, mmq_y, cc, warp_size, nwarps); CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); @@ -4121,7 +4265,7 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a 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, - ntx_fd); + ntx_fd, prof_guard.ptr()); } else { constexpr bool need_check = true; log_launch_config(block_nums_xy_tiling, need_check, "tiling"); @@ -4130,7 +4274,7 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a 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, - ntx_fd); + ntx_fd, prof_guard.ptr()); } return; } @@ -4163,7 +4307,7 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a 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, - ntx_fd); + ntx_fd, prof_guard.ptr()); if (!fixup_needed) { return; @@ -4182,7 +4326,7 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a 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, - ntx_fd); + ntx_fd, prof_guard.ptr()); if (!fixup_needed) { return; diff --git a/ggml/src/ggml-cuda/vendors/hip.h b/ggml/src/ggml-cuda/vendors/hip.h index a6115cd80dc4..2ad1207bef98 100644 --- a/ggml/src/ggml-cuda/vendors/hip.h +++ b/ggml/src/ggml-cuda/vendors/hip.h @@ -94,6 +94,7 @@ #define cudaMemcpyDeviceToHost hipMemcpyDeviceToHost #define cudaMemcpyHostToDevice hipMemcpyHostToDevice #define cudaMemcpyKind hipMemcpyKind +#define cudaMemcpyToSymbol hipMemcpyToSymbol #define cudaMemset hipMemset #define cudaMemsetAsync hipMemsetAsync #define cudaMemGetInfo hipMemGetInfo From 7ed8d1ec6a57d7d91688ee1d9c383eba4bdf9527 Mon Sep 17 00:00:00 2001 From: lichang Date: Tue, 23 Jun 2026 17:48:43 -0600 Subject: [PATCH 06/22] Improve RDNA3.5 Q4_K MMQ dequant/vec_dot and fix phase profiling overhead. Add gfx115x load_tiles_q4_K_dm_rdna35 and vec_dot_q4_K_q8_1_mma with per-k dmA hoisting; restore pre-profiling tile loop when MMQ_PROFILE_PHASES is off so extra __syncthreads do not regress P0 TFLOPS. Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 312 ++++++++++++++++++++++++++++++------- 1 file changed, 254 insertions(+), 58 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index ea014766318e..a69e56270d85 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -2233,6 +2233,44 @@ static __device__ __forceinline__ int unpack_scales_q45_K(const int * scales, co ((scales[ksc/2] >> (2 * (ksc % 2))) & 0x30303030); // upper 2 bits } +#if defined(RDNA3_5) +// gfx115x mmq_y=64: 128 threads map 2:1 to rows for scale/dm prep (no warp divergence). +template +static __device__ __forceinline__ void load_tiles_q4_K_dm_rdna35( + const char * __restrict__ x, half2 * __restrict__ x_dm, const int kbx0, const int i_max, const int stride) { + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); + + const int dm_tid = threadIdx.y*warp_size + threadIdx.x; + const int dm_row = dm_tid/2; + const int dm_ksc = dm_tid%2; + + if (dm_row >= mmq_y) { + return; + } + + int i = dm_row; + if (need_check) { + i = min(i, i_max); + } + + const block_q4_K * bxi = (const block_q4_K *) x + kbx0 + i*stride; + const int * scales = (const int *) bxi->scales; + + const int sc32 = unpack_scales_q45_K(scales, dm_ksc + 0); + const int m32 = unpack_scales_q45_K(scales, dm_ksc + 2); + + const uint8_t * sc8 = (const uint8_t *) &sc32; + const uint8_t * m8 = (const uint8_t *) &m32; + + const half2 dm = bxi->dm * make_half2(1.0f, -1.0f); + +#pragma unroll + for (int l = 0; l < 4; ++l) { + x_dm[i*MMQ_MMA_TILE_X_K_Q8_1 + sizeof(int)*dm_ksc + l] = dm*make_half2(sc8[l], m8[l]); + } +} +#endif // RDNA3_5 + template static __device__ __forceinline__ void load_tiles_q4_K( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int nwarps = mmq_get_nwarps_device(); @@ -2248,6 +2286,11 @@ template static __device__ __forceinline__ void loa int * x_sc = (int *) (x_dm + txs.dm); #endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) +#if defined(RDNA3_5) && defined(AMD_WMMA_AVAILABLE) && !defined(AMD_MFMA_AVAILABLE) + // dm first: warms cache lines before qs reads from same block_q4_K. + load_tiles_q4_K_dm_rdna35(x, x_dm, kbx0, i_max, stride); +#endif + constexpr int threads_per_row = MMQ_ITER_K / (4 * QR4_K); constexpr int nrows = warp_size / threads_per_row; const int txi = warp_size > threads_per_row ? threadIdx.x % threads_per_row : threadIdx.x; @@ -2261,7 +2304,11 @@ template static __device__ __forceinline__ void loa } const block_q4_K * bxi = (const block_q4_K *) x + kbx0 + i*stride; +#if defined(RDNA3_5) + const int qs0 = ((const int *) bxi->qs)[txi]; +#else const int qs0 = get_int_b4(bxi->qs, txi); +#endif #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) x_qs[i*MMQ_MMA_TILE_X_K_Q8_1 + 16*(txi/8) + txi % 8 + 0] = (qs0 >> 0) & 0x0F0F0F0F; @@ -2272,19 +2319,20 @@ template static __device__ __forceinline__ void loa } #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) +#if defined(RDNA3_5) && defined(AMD_WMMA_AVAILABLE) && !defined(AMD_MFMA_AVAILABLE) + // dm handled above. +#else constexpr int rows_per_warp = warp_size / 2; #pragma unroll for (int i0 = 0; i0 < mmq_y; i0 += nwarps*rows_per_warp) { -#if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) - // Need if on AMD instead of % because warp_size == 64 - // This causes double work and throughput loss (MI300X) - // H100 loses about 100 t/s with 'if' condition over '%' +#if defined(AMD_MFMA_AVAILABLE) + // Need if on CDNA (warp_size == 64) instead of %. int i = i0 + threadIdx.y*rows_per_warp + threadIdx.x/2; if (i < mmq_y) { #else int i = (i0 + threadIdx.y*rows_per_warp + threadIdx.x/2) % mmq_y; { -#endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) +#endif // defined(AMD_MFMA_AVAILABLE) if (need_check) { i = min(i, i_max); } @@ -2308,6 +2356,7 @@ template static __device__ __forceinline__ void loa } } } +#endif // RDNA3_5 RDNA WMMA dm path #else #pragma unroll for (int i0 = 0; i0 < mmq_y; i0 += nwarps*warp_size) { @@ -2377,6 +2426,72 @@ static __device__ __forceinline__ void vec_dot_q4_K_q8_1_dp4a( } } +template +static __device__ __forceinline__ void vec_dot_q4_K_q8_1_mma( + const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) { +#if defined(RDNA3_5) && defined(AMD_WMMA_AVAILABLE) && !defined(AMD_MFMA_AVAILABLE) + // Hoist dmA per k-slice only — full A+dmA hoist spills registers (~2x slower). + constexpr data_layout input_layout = get_input_data_layout(); + typedef tile<16, 8, int, input_layout> tile_A; + typedef tile<16, 8, int, input_layout> tile_B; + typedef tile<16, 16, int, DATA_LAYOUT_J_MAJOR> tile_C; + + constexpr int granularity = mmq_get_granularity_device(mmq_x); + constexpr int rows_per_warp = granularity; + constexpr int ntx = rows_per_warp/tile_C::I; + + y += (threadIdx.y % ntx) * (tile_C::J*MMQ_TILE_Y_K); + + const int * x_qs = (const int *) x; + const half2 * x_dm = (const half2 *) x_qs + 2*MMQ_TILE_NE_K; + const int * y_qs = (const int *) y + 4; + const half2 * y_dm = (const half2 *) y; + + const int i0 = (threadIdx.y / ntx) * rows_per_warp; + + for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += QI8_1) { + const int k0 = k00 + k01; + + tile_A A[ntx]; + float2 dmA[ntx][tile_C::ne]; + +#pragma unroll + for (int n = 0; n < ntx; ++n) { + load_ldmatrix(A[n], x_qs + (i0 + n*tile_A::I)*MMQ_MMA_TILE_X_K_Q8_1 + k0, MMQ_MMA_TILE_X_K_Q8_1); + +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + const int i = i0 + n*tile_A::I + tile_C::get_i(l); + dmA[n][l] = __half22float2(x_dm[i*MMQ_MMA_TILE_X_K_Q8_1 + k0/QI8_1]); + } + } + +#pragma unroll + for (int j0 = 0; j0 < mmq_x; j0 += ntx*tile_C::J) { + tile_B B; + load_ldmatrix(B, y_qs + j0*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); + + const int j = j0 + tile_C::get_j(0); + const float2 dsB = __half22float2(y_dm[j*MMQ_TILE_Y_K + k01/QI8_1]); + +#pragma unroll + for (int n = 0; n < ntx; ++n) { + tile_C C; + mma(C, A[n], B); + +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + sum[(j0/tile_C::J + n)*tile_C::ne + l] += dmA[n][l].x*dsB.x*C.x[l]; + sum[(j0/tile_C::J + n)*tile_C::ne + l] += dmA[n][l].y*dsB.y; + } + } + } + } +#else + vec_dot_q8_1_q8_1_mma(x, y, sum, k00); +#endif +} + template static __device__ __forceinline__ void load_tiles_q5_K( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int nwarps = mmq_get_nwarps_device(); @@ -3502,7 +3617,7 @@ template struct mmq_type_traits { static constexpr int vdr = VDR_Q4_K_Q8_1_MMQ; static constexpr load_tiles_mmq_t load_tiles = load_tiles_q4_K; - static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_q8_1_q8_1_mma; + static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_q4_K_q8_1_mma; static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_q4_K_q8_1_dp4a; }; @@ -3669,26 +3784,73 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( int y0_next_cache[tile_y_nchunks]; bool have_y0_prefetch = false; - for (int kb0 = kb0_start; kb0 < kb0_stop; kb0 += blocks_per_iter) { - mmq_profile_kb_iter(mmq_profile); + if (mmq_profile) { + for (int kb0 = kb0_start; kb0 < kb0_stop; kb0 += blocks_per_iter) { + mmq_profile_kb_iter(mmq_profile); + + { + uint64_t t0 = 0; + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); + load_tiles(x, tile_x, offset_x + kb0, tile_x_max_i, stride_row_x); + __syncthreads(); + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_LOAD_TILES); + } - { - uint64_t t0 = 0; - __syncthreads(); - mmq_profile_phase_begin(mmq_profile, t0); - load_tiles(x, tile_x, offset_x + kb0, tile_x_max_i, stride_row_x); - __syncthreads(); - mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_LOAD_TILES); - } + const int yk = kb0 * qk / ne_block; + const int * by0 = y + ncols_y * yk * sz; + const int * by1 = y + ncols_y * (yk + 1) * sz; - const int yk = kb0 * qk / ne_block; - const int * by0 = y + ncols_y * yk * sz; - const int * by1 = y + ncols_y * (yk + 1) * sz; + { + uint64_t t0 = 0; + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); - { - uint64_t t0 = 0; - __syncthreads(); - mmq_profile_phase_begin(mmq_profile, t0); + if (have_y0_prefetch) { + mmq_tile_y_store_regs(tile_y, y0_next_cache); + have_y0_prefetch = false; + } else { + mmq_tile_y_load_global(tile_y, by0); + } + + __syncthreads(); + + int y1_cache[tile_y_nchunks]; + mmq_tile_y_load_global_to_regs(by1, y1_cache); + + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_Y_ACT); + + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); + + vec_dot(tile_x, tile_y, sum, 0); + + __syncthreads(); + + mmq_tile_y_store_regs(tile_y, y1_cache); + + __syncthreads(); + + const int kb0_next = kb0 + blocks_per_iter; + if (kb0_next < kb0_stop) { + const int * by0_next = y + ncols_y * (kb0_next * qk / ne_block) * sz; + mmq_tile_y_load_global_to_regs(by0_next, y0_next_cache); + have_y0_prefetch = true; + } + + vec_dot(tile_x, tile_y, sum, MMQ_TILE_NE_K); + + __syncthreads(); + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_VEC_DOT); + } + } + } else { + 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); + + const int yk = kb0 * qk / ne_block; + const int * by0 = y + ncols_y * yk * sz; + const int * by1 = y + ncols_y * (yk + 1) * sz; if (have_y0_prefetch) { mmq_tile_y_store_regs(tile_y, y0_next_cache); @@ -3702,11 +3864,6 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( int y1_cache[tile_y_nchunks]; mmq_tile_y_load_global_to_regs(by1, y1_cache); - mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_Y_ACT); - - __syncthreads(); - mmq_profile_phase_begin(mmq_profile, t0); - vec_dot(tile_x, tile_y, sum, 0); __syncthreads(); @@ -3725,58 +3882,97 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( vec_dot(tile_x, tile_y, sum, MMQ_TILE_NE_K); __syncthreads(); - mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_VEC_DOT); } } } else #endif // RDNA3_5 { - for (int kb0 = kb0_start; kb0 < kb0_stop; kb0 += blocks_per_iter) { - mmq_profile_kb_iter(mmq_profile); + if (mmq_profile) { + for (int kb0 = kb0_start; kb0 < kb0_stop; kb0 += blocks_per_iter) { + mmq_profile_kb_iter(mmq_profile); + + { + uint64_t t0 = 0; + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); + load_tiles(x, tile_x, offset_x + kb0, tile_x_max_i, stride_row_x); + __syncthreads(); + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_LOAD_TILES); + } - { - uint64_t t0 = 0; - __syncthreads(); - mmq_profile_phase_begin(mmq_profile, t0); - load_tiles(x, tile_x, offset_x + kb0, tile_x_max_i, stride_row_x); - __syncthreads(); - mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_LOAD_TILES); - } + const int yk = kb0 * qk / ne_block; + const int * by0 = y + ncols_y * yk * sz; + const int * by1 = y + ncols_y * (yk + 1) * sz; - const int yk = kb0 * qk / ne_block; - const int * by0 = y + ncols_y * yk * sz; - const int * by1 = y + ncols_y * (yk + 1) * sz; + { + uint64_t t0 = 0; + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); +#pragma unroll + for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { + int l = l0 + threadIdx.y*warp_size + threadIdx.x; - { - uint64_t t0 = 0; - __syncthreads(); - mmq_profile_phase_begin(mmq_profile, t0); + tile_y[l] = by0[l]; + } + __syncthreads(); + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_Y_ACT); + } + + { + uint64_t t0 = 0; + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); + vec_dot(tile_x, tile_y, sum, 0); + __syncthreads(); #pragma unroll - for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { - int l = l0 + threadIdx.y*warp_size + threadIdx.x; + for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { + int l = l0 + threadIdx.y*warp_size + threadIdx.x; - tile_y[l] = by0[l]; + tile_y[l] = by1[l]; + } + __syncthreads(); + vec_dot(tile_x, tile_y, sum, MMQ_TILE_NE_K); + __syncthreads(); + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_VEC_DOT); } - __syncthreads(); - mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_Y_ACT); } + } else { + 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); + + const int yk = kb0 * qk / ne_block; + const int * by0 = y + ncols_y * yk * sz; + const int * by1 = y + ncols_y * (yk + 1) * sz; + + { +#pragma unroll + for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { + int l = l0 + threadIdx.y*warp_size + threadIdx.x; + + tile_y[l] = by0[l]; + } + } - { - uint64_t t0 = 0; __syncthreads(); - mmq_profile_phase_begin(mmq_profile, t0); + vec_dot(tile_x, tile_y, sum, 0); + __syncthreads(); + + { #pragma unroll - for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { - int l = l0 + threadIdx.y*warp_size + threadIdx.x; + for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { + int l = l0 + threadIdx.y*warp_size + threadIdx.x; - tile_y[l] = by1[l]; + tile_y[l] = by1[l]; + } } + __syncthreads(); + vec_dot(tile_x, tile_y, sum, MMQ_TILE_NE_K); + __syncthreads(); - mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_VEC_DOT); } } } From fb63447180576d40356b696535fbe4a1c942af55 Mon Sep 17 00:00:00 2001 From: lichang Date: Tue, 23 Jun 2026 18:03:57 -0600 Subject: [PATCH 07/22] Refactor RDNA3.5 Q4_K qs load for WMMA ldmatrix layout. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Split load_tiles_q4_K qs staging into load_tiles_q4_K_qs_wmma_rdna35 (one warp/row, coalesced qs global load) and mmq_q4_K_qs_lds_k helper that documents the pseudo-q8_1 LDS offsets used by load_ldmatrix. Perf (gfx1151, P0 128×12288×4096): neutral vs prior commit — 26.56 vs 26.51 TFLOPS, 485 vs 486 µs, pp128 ~1135 t/s (within run-to-run noise). Layout math is unchanged; this is groundwork for a Q4_K-native MMA tile. Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 43 ++++++++++++++++++++++++++++++++++---- 1 file changed, 39 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index a69e56270d85..67f8f05e5b78 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -2233,6 +2233,11 @@ static __device__ __forceinline__ int unpack_scales_q45_K(const int * scales, co ((scales[ksc/2] >> (2 * (ksc % 2))) & 0x30303030); // upper 2 bits } +// LDS int offset for pseudo-q8_1 qs within one weight row (matches Q4_1 / WMMA ldmatrix layout). +static __device__ __forceinline__ int mmq_q4_K_qs_lds_k(const int txi, const int nibble) { + return ((txi >> 3) << 4) + (txi & 7) + (nibble << 3); +} + #if defined(RDNA3_5) // gfx115x mmq_y=64: 128 threads map 2:1 to rows for scale/dm prep (no warp divergence). template @@ -2269,6 +2274,35 @@ static __device__ __forceinline__ void load_tiles_q4_K_dm_rdna35( x_dm[i*MMQ_MMA_TILE_X_K_Q8_1 + sizeof(int)*dm_ksc + l] = dm*make_half2(sc8[l], m8[l]); } } + +// gfx115x WMMA: one warp per row, coalesced qs global load, nibble expand to ldmatrix-ready LDS. +template +static __device__ __forceinline__ void load_tiles_q4_K_qs_wmma_rdna35( + const char * __restrict__ x, int * __restrict__ x_qs, const int kbx0, const int i_max, const int stride) { + constexpr int nwarps = mmq_get_nwarps_device(); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); + constexpr int qs_per_row = MMQ_ITER_K / (4 * QR4_K); // 32 ints per block_q4_K::qs + static_assert(qs_per_row == 32, "bad Q4_K qs_per_row"); + constexpr int nrows = warp_size / qs_per_row; + static_assert(nrows == 1, "Q4_K RDNA3.5 WMMA qs path expects one row per warp"); + +#pragma unroll + for (int i0 = 0; i0 < mmq_y; i0 += nrows*nwarps) { + int i = i0 + threadIdx.y; + + if (need_check) { + i = min(i, i_max); + } + + const block_q4_K * bxi = (const block_q4_K *) x + kbx0 + i*stride; + const int qs0 = ((const int *) bxi->qs)[threadIdx.x]; + + int * row_qs = x_qs + i*MMQ_MMA_TILE_X_K_Q8_1; + const int kqs = mmq_q4_K_qs_lds_k(threadIdx.x, 0); + row_qs[kqs] = qs0 & 0x0F0F0F0F; + row_qs[kqs+8] = (qs0 >> 4) & 0x0F0F0F0F; + } +} #endif // RDNA3_5 template static __device__ __forceinline__ void load_tiles_q4_K( @@ -2289,8 +2323,8 @@ template static __device__ __forceinline__ void loa #if defined(RDNA3_5) && defined(AMD_WMMA_AVAILABLE) && !defined(AMD_MFMA_AVAILABLE) // dm first: warms cache lines before qs reads from same block_q4_K. load_tiles_q4_K_dm_rdna35(x, x_dm, kbx0, i_max, stride); -#endif - + load_tiles_q4_K_qs_wmma_rdna35(x, x_qs, kbx0, i_max, stride); +#else constexpr int threads_per_row = MMQ_ITER_K / (4 * QR4_K); constexpr int nrows = warp_size / threads_per_row; const int txi = warp_size > threads_per_row ? threadIdx.x % threads_per_row : threadIdx.x; @@ -2311,12 +2345,13 @@ template static __device__ __forceinline__ void loa #endif #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) - x_qs[i*MMQ_MMA_TILE_X_K_Q8_1 + 16*(txi/8) + txi % 8 + 0] = (qs0 >> 0) & 0x0F0F0F0F; - x_qs[i*MMQ_MMA_TILE_X_K_Q8_1 + 16*(txi/8) + txi % 8 + 8] = (qs0 >> 4) & 0x0F0F0F0F; + x_qs[i*MMQ_MMA_TILE_X_K_Q8_1 + mmq_q4_K_qs_lds_k(txi, 0)] = (qs0 >> 0) & 0x0F0F0F0F; + x_qs[i*MMQ_MMA_TILE_X_K_Q8_1 + mmq_q4_K_qs_lds_k(txi, 1)] = (qs0 >> 4) & 0x0F0F0F0F; #else x_qs[i*(MMQ_TILE_NE_K + 1) + txi] = qs0; #endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) } +#endif // RDNA3_5 WMMA qs path #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) #if defined(RDNA3_5) && defined(AMD_WMMA_AVAILABLE) && !defined(AMD_MFMA_AVAILABLE) From 5924b0a85b5aa6540223e5e900160a81d7ea6bb4 Mon Sep 17 00:00:00 2001 From: lichang Date: Wed, 24 Jun 2026 12:12:18 -0600 Subject: [PATCH 08/22] Gate MMQ kernel dimension logs behind GGML_CUDA_MM_LOG. fprintf for MUL_MAT_Q, MUL_MAT_VEC_Q, QUANTIZE_MMQ_Q8_1, and MUL_MAT_Q_LAUNCH are off by default to avoid CPU overhead in llama-bench TTFT runs; set GGML_CUDA_MM_LOG=1 for profiling scripts. Co-authored-by: Cursor --- ggml/src/ggml-cuda/common.cuh | 9 +++++++++ ggml/src/ggml-cuda/mmq.cu | 16 ++++++++++------ ggml/src/ggml-cuda/mmq.cuh | 3 +++ ggml/src/ggml-cuda/mmvq.cu | 8 +++++--- ggml/src/ggml-cuda/quantize.cu | 8 +++++--- 5 files changed, 32 insertions(+), 12 deletions(-) diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index e6e50e041195..33916b74d982 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -1639,3 +1639,12 @@ static __inline__ void ggml_cuda_kernel_launch(Kernel kernel, const ggml_cuda_ke CUDA_CHECK(cudaGetLastError()); } +// Opt-in stderr logging for MMQ/MMVQ kernel dimensions (set GGML_CUDA_MM_LOG=1). +static inline bool ggml_cuda_mm_log_enabled() { + static int enabled = -1; + if (enabled < 0) { + enabled = getenv("GGML_CUDA_MM_LOG") ? 1 : 0; + } + return enabled; +} + diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 3c3272501968..e400fcefa987 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -158,11 +158,13 @@ void ggml_cuda_mul_mat_q( ne03, ne13, s03, s13, s3, use_stream_k, ne1}; - // Log kernel dimensions for profiling analysis + // Log kernel dimensions for profiling analysis (GGML_CUDA_MM_LOG=1). // dst[ne01, ne11] = src0[ne00, ne01] @ src1[ne10, ne11]; ne00==ne10 (K) // GEMM C[M,N] = W[N,K] @ X[K,M]: M=ne11, N=ne01, K=ne00 - fprintf(stderr, "[MUL_MAT_Q] M=%ld N=%ld K=%ld ne00=%ld ne01=%ld ne1=%ld ne11=%ld ne02=%ld ne12=%ld ne03=%ld ne13=%ld type=%d\n", - ne11, ne01, ne00, ne00, ne01, ne1, ne11, ne02, ne12, ne03, ne13, src0->type); + if (ggml_cuda_mm_log_enabled()) { + fprintf(stderr, "[MUL_MAT_Q] M=%ld N=%ld K=%ld ne00=%ld ne01=%ld ne1=%ld ne11=%ld ne02=%ld ne12=%ld ne03=%ld ne13=%ld type=%d\n", + ne11, ne01, ne00, ne00, ne01, ne1, ne11, ne02, ne12, ne03, ne13, src0->type); + } ggml_cuda_mul_mat_q_switch_type(ctx, args, stream); return; @@ -226,9 +228,11 @@ void ggml_cuda_mul_mat_q( ne03, ne13, s03, s13, s3, use_stream_k, ne12}; - // Log kernel dimensions for profiling analysis (batched/MoE path) - fprintf(stderr, "[MUL_MAT_Q_MoE] M=%ld N=%ld K=%ld ne_get_rows=%ld ne00=%ld ne01=%ld ne02=%ld ne12=%ld type=%d\n", - ne11_flat, ne01, ne00, ne_get_rows, ne00, ne01, ne02, ne12, src0->type); + // Log kernel dimensions for profiling analysis (batched/MoE path, GGML_CUDA_MM_LOG=1). + if (ggml_cuda_mm_log_enabled()) { + fprintf(stderr, "[MUL_MAT_Q_MoE] M=%ld N=%ld K=%ld ne_get_rows=%ld ne00=%ld ne01=%ld ne02=%ld ne12=%ld type=%d\n", + ne11_flat, ne01, ne00, ne_get_rows, ne00, ne01, ne02, ne12, src0->type); + } ggml_cuda_mul_mat_q_switch_type(ctx, args, stream); } diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 67f8f05e5b78..5ef8be3d2d78 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -4463,6 +4463,9 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a const dim3 block_nums_xy_tiling(nty, ntx, ntzw); const auto log_launch_config = [&](const dim3 & grid, const bool need_check, const char * path) { + if (!ggml_cuda_mm_log_enabled()) { + return; + } fprintf(stderr, "[MUL_MAT_Q_LAUNCH] type=%d mmq_x=%d mmq_y=%d nwarps=%d warp_size=%d " "nbytes_shared=%d smpbo=%zu grid=(%u,%u,%u) block=(%u,%u,%u) " diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index d94db2993685..cee73315d5ef 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -1213,11 +1213,13 @@ void ggml_cuda_mul_mat_vec_q( const int64_t ids_stride = ids ? ids->nb[1] / ggml_type_size(ids->type) : 0; - // Log kernel dimensions for profiling analysis + // Log kernel dimensions for profiling analysis (GGML_CUDA_MM_LOG=1). // dst[ne01] = src0[ne00, ne01] @ src1[ne10]; ne00==ne10 (K), ne11==1 for decode // GEMV y[M] = W[M,K] @ x[K]: M=ne01, K=ne00 - fprintf(stderr, "[MUL_MAT_VEC_Q] M=%ld K=%ld ne00=%ld ne01=%ld ne10=%ld ne11=%ld ne02=%ld ne12=%ld ne03=%ld ne13=%ld type=%d\n", - ne01, ne00, ne00, ne01, ne10, ne11, ne02, ne12, ne03, ne13, src0->type); + if (ggml_cuda_mm_log_enabled()) { + fprintf(stderr, "[MUL_MAT_VEC_Q] M=%ld K=%ld ne00=%ld ne01=%ld ne10=%ld ne11=%ld ne02=%ld ne12=%ld ne03=%ld ne13=%ld type=%d\n", + ne01, ne00, ne00, ne01, ne10, ne11, ne02, ne12, ne03, ne13, src0->type); + } mul_mat_vec_q_switch_type( src0->data, src0->type, src1_q8_1.get(), ids_d, fusion_local, dst_d, ne00, diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu index c3c401bf9042..7d23fd8886dd 100644 --- a/ggml/src/ggml-cuda/quantize.cu +++ b/ggml/src/ggml-cuda/quantize.cu @@ -401,9 +401,11 @@ void quantize_mmq_q8_1_cuda( const dim3 num_blocks(ne1, block_num_y, ne2*ne3); const dim3 block_size(CUDA_QUANTIZE_BLOCK_SIZE_MMQ, 1, 1); - // Log kernel dimensions for profiling analysis - fprintf(stderr, "[QUANTIZE_MMQ_Q8_1] ne00=%ld ne0=%ld ne1=%ld ne2=%ld ne3=%ld grid=(%u,%u,%u) type=%d\n", - ne00, ne0, ne1, ne2, ne3, num_blocks.x, num_blocks.y, num_blocks.z, type_src0); + // Log kernel dimensions for profiling analysis (GGML_CUDA_MM_LOG=1). + if (ggml_cuda_mm_log_enabled()) { + fprintf(stderr, "[QUANTIZE_MMQ_Q8_1] ne00=%ld ne0=%ld ne1=%ld ne2=%ld ne3=%ld grid=(%u,%u,%u) type=%d\n", + ne00, ne0, ne1, ne2, ne3, num_blocks.x, num_blocks.y, num_blocks.z, type_src0); + } switch (mmq_get_q8_1_ds_layout(type_src0)) { case MMQ_Q8_1_DS_LAYOUT_D4: From afa1455f8ff800401d2ef1fabb4c06af87b563f0 Mon Sep 17 00:00:00 2001 From: lichang Date: Wed, 24 Jun 2026 13:41:08 -0600 Subject: [PATCH 09/22] cuda: precompute Q4_K vec_dot scale factors before WMMA accumulate Hoist dmA*dsB into per-lane scl[] before mma on RDNA3.5 vec_dot_q4_K_q8_1_mma to separate scale multiply from C accumulation (ATT showed v_fma_mix hotspots). Prefill P0 Q4_K_M 128x12288x4096 (3x profile_mmq_prefill): median 26.59 TFLOPS / 484.7 us vs 26.56 / 485 us baseline; VGPR 224 (+8), LDS conflict ratio 0.348. Assisted-by: Auto Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 13 +++++++++++-- 1 file changed, 11 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 5ef8be3d2d78..6e33e77d2c0e 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -2511,13 +2511,22 @@ static __device__ __forceinline__ void vec_dot_q4_K_q8_1_mma( #pragma unroll for (int n = 0; n < ntx; ++n) { + float2 scl[tile_C::ne]; + +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + scl[l].x = dmA[n][l].x*dsB.x; + scl[l].y = dmA[n][l].y*dsB.y; + } + tile_C C; mma(C, A[n], B); #pragma unroll for (int l = 0; l < tile_C::ne; ++l) { - sum[(j0/tile_C::J + n)*tile_C::ne + l] += dmA[n][l].x*dsB.x*C.x[l]; - sum[(j0/tile_C::J + n)*tile_C::ne + l] += dmA[n][l].y*dsB.y; + const int si = (j0/tile_C::J + n)*tile_C::ne + l; + sum[si] += scl[l].x*C.x[l]; + sum[si] += scl[l].y; } } } From 46b38f254fbff1db97bcedfa0ba94972a5ee2d5e Mon Sep 17 00:00:00 2001 From: lichang Date: Wed, 24 Jun 2026 17:56:02 -0600 Subject: [PATCH 10/22] cuda: fix MMQ phase profiler clock64 deltas on HIP clock64() can move backward across __syncthreads, which inflated phase cycle counts and produced garbage load_tiles percentages; discard small backward glitches and skip corrupt host-side totals. Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 19 +++++++++++++++++-- 1 file changed, 17 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 6e33e77d2c0e..202333449a2c 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -40,11 +40,25 @@ static __device__ __forceinline__ void mmq_profile_phase_begin(uint64_t * const t0 = clock64(); } +// Safe delta for clock64(): unsigned (t1 - t0) wraps to ~2^64 when t1 < t0 (non-monotonic +// timestamp across __syncthreads / wave scheduling). Discard small backward glitches; keep +// genuine 64-bit counter wrap (t0 near UINT64_MAX, t1 small). +static __device__ __forceinline__ uint64_t mmq_profile_phase_cycles(const uint64_t t0, const uint64_t t1) { + if (t1 >= t0) { + return t1 - t0; + } + const uint64_t backward = t0 - t1; + if (backward > (UINT64_MAX / 2)) { + return t1 - t0; // unsigned wrap + } + return 0; +} + static __device__ __forceinline__ void mmq_profile_phase_end(uint64_t * const profile, const uint64_t t0, const int phase) { if (!mmq_profile_want_sample(profile)) { return; } - profile[phase] += clock64() - t0; + profile[phase] += mmq_profile_phase_cycles(t0, clock64()); } static bool mmq_profile_phases_enabled() { @@ -86,7 +100,8 @@ struct mmq_profile_guard { uint64_t h[MMQ_PROFILE_N] = {}; CUDA_CHECK(cudaMemcpy(h, buf.get(), sizeof(h), cudaMemcpyDeviceToHost)); const uint64_t total = h[MMQ_PROFILE_LOAD_TILES] + h[MMQ_PROFILE_Y_ACT] + h[MMQ_PROFILE_VEC_DOT]; - if (total > 0) { + // Skip corrupt rows (overflow / legacy bad deltas) — see mmq_profile_phase_cycles. + if (total > 0 && total < 5000000000ULL) { fprintf(stderr, "[MMQ_PROFILE] type=%d mmq_x=%d nrows_x=%ld ncols_max=%ld " "load_tiles=%.1f%% y_act=%.1f%% vec_dot=%.1f%% " From bee2b6dd185b182be7976c8044725df7290ca1f3 Mon Sep 17 00:00:00 2001 From: lichang Date: Thu, 25 Jun 2026 13:23:58 -0600 Subject: [PATCH 11/22] cuda: optimize RDNA3.5 tiny-M MMQ for gate Q8_0 prefill MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Use mmq_x=128 for small nrows_x grids instead of dual-WG LDS downsizing, and hoist dA in Q8_0/Q4_0 vec_dot on RDNA3.5 WMMA. Gate 128×32×4096 drops ~209µs to ~105µs with no P0 Q4_K regression. Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 165 ++++++++++++++++++++++++++++++++++++- 1 file changed, 163 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 202333449a2c..bb1564070bb7 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -1451,6 +1451,133 @@ static __device__ __forceinline__ void vec_dot_q8_0_q8_1_mma( #endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) } +template +static __device__ __forceinline__ void vec_dot_q4_0_q8_1_mma( + const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) { +#if defined(RDNA3_5) && defined(AMD_WMMA_AVAILABLE) && !defined(AMD_MFMA_AVAILABLE) + constexpr data_layout input_layout = get_input_data_layout(); + typedef tile<16, 8, int, input_layout> tile_A; + typedef tile<16, 8, int, input_layout> tile_B; + typedef tile<16, 16, int, DATA_LAYOUT_J_MAJOR> tile_C; + + constexpr int granularity = mmq_get_granularity_device(mmq_x); + constexpr int rows_per_warp = granularity; + constexpr int ntx = rows_per_warp/tile_C::I; + + y += (threadIdx.y % ntx) * (tile_C::J*MMQ_TILE_Y_K); + + const int * x_qs = (const int *) x; + const float * x_df = (const float *) x_qs + 2*MMQ_TILE_NE_K; + const int * y_qs = (const int *) y + 4; + const half2 * y_ds = (const half2 *) y; + + const int i0 = (threadIdx.y / ntx) * rows_per_warp; + + for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += QI8_0) { + const int k0 = k00 + k01; + + tile_A A[ntx]; + float dA[ntx][tile_C::ne]; + +#pragma unroll + for (int n = 0; n < ntx; ++n) { + load_ldmatrix(A[n], x_qs + (i0 + n*tile_A::I)*MMQ_MMA_TILE_X_K_Q8_0 + k0, MMQ_MMA_TILE_X_K_Q8_0); + +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + const int i = i0 + n*tile_A::I + tile_C::get_i(l); + dA[n][l] = x_df[i*MMQ_MMA_TILE_X_K_Q8_0 + k0/QI8_0]; + } + } + +#pragma unroll + for (int j0 = 0; j0 < mmq_x; j0 += ntx*tile_C::J) { + tile_B B; + load_ldmatrix(B, y_qs + j0*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); + + const int j = j0 + tile_C::get_j(0); + const float dB = __low2float(y_ds[j*MMQ_TILE_Y_K + k01/QI8_1]); + +#pragma unroll + for (int n = 0; n < ntx; ++n) { + tile_C C; + mma(C, A[n], B); + +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + sum[(j0/tile_C::J + n)*tile_C::ne + l] += C.x[l]*dA[n][l]*dB; + } + } + } + } +#else + vec_dot_q8_0_q8_1_mma(x, y, sum, k00); +#endif +} + +template +static __device__ __forceinline__ void vec_dot_q8_0_q8_1_mma_rdna35( + const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) { +#if defined(RDNA3_5) && defined(AMD_WMMA_AVAILABLE) && !defined(AMD_MFMA_AVAILABLE) + constexpr data_layout input_layout = get_input_data_layout(); + typedef tile<16, 8, int, input_layout> tile_A; + typedef tile<16, 8, int, input_layout> tile_B; + typedef tile<16, 16, int, DATA_LAYOUT_J_MAJOR> tile_C; + + constexpr int granularity = mmq_get_granularity_device(mmq_x); + constexpr int rows_per_warp = granularity; + constexpr int ntx = rows_per_warp/tile_C::I; + + y += (threadIdx.y % ntx) * (tile_C::J*MMQ_TILE_Y_K); + + const int * x_qs = (const int *) x; + const float * x_df = (const float *) x_qs + 2*MMQ_TILE_NE_K; + const int * y_qs = (const int *) y + 4; + const float * y_df = (const float *) y; + + const int i0 = (threadIdx.y / ntx) * rows_per_warp; + + for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += QI8_0) { + const int k0 = k00 + k01; + + tile_A A[ntx]; + float dA[ntx][tile_C::ne]; + +#pragma unroll + for (int n = 0; n < ntx; ++n) { + load_ldmatrix(A[n], x_qs + (i0 + n*tile_A::I)*MMQ_MMA_TILE_X_K_Q8_0 + k0, MMQ_MMA_TILE_X_K_Q8_0); + +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + const int i = i0 + n*tile_A::I + tile_C::get_i(l); + dA[n][l] = x_df[i*MMQ_MMA_TILE_X_K_Q8_0 + k0/QI8_0]; + } + } + +#pragma unroll + for (int j0 = 0; j0 < mmq_x; j0 += ntx*tile_C::J) { + tile_B B; + load_ldmatrix(B, y_qs + j0*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); + + const int j = j0 + tile_C::get_j(0); + const float dB = y_df[j*MMQ_TILE_Y_K + k01/QI8_1]; + +#pragma unroll + for (int n = 0; n < ntx; ++n) { + tile_C C; + mma(C, A[n], B); + +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + sum[(j0/tile_C::J + n)*tile_C::ne + l] += C.x[l]*dA[n][l]*dB; + } + } + } + } +#else + vec_dot_q8_0_q8_1_mma(x, y, sum, k00); +#endif +} template static __device__ __forceinline__ void vec_dot_q8_1_q8_1_dp4a( @@ -3594,7 +3721,7 @@ template struct mmq_type_traits { static constexpr int vdr = VDR_Q4_0_Q8_1_MMQ; static constexpr load_tiles_mmq_t load_tiles = load_tiles_q4_0; - static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_q8_0_q8_1_mma; + static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_q4_0_q8_1_mma; static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_q4_0_q8_1_dp4a; }; @@ -3626,7 +3753,7 @@ template struct mmq_type_traits { static constexpr int vdr = VDR_Q8_0_Q8_1_MMQ; static constexpr load_tiles_mmq_t load_tiles = load_tiles_q8_0; - static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_q8_0_q8_1_mma; + static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_q8_0_q8_1_mma_rdna35; static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_q8_0_q8_1_dp4a; }; @@ -4661,6 +4788,40 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda } #endif // GGML_USE_HIP +#if defined(GGML_USE_HIP) + // Tiny M (e.g. gate 128×32×4096 Q8_0): prefer largest mmq_x within LDS so N is not + // split; occupancy is irrelevant when only a handful of blocks launch. + if (GGML_CUDA_CC_IS_RDNA3_5(cc) && args.nrows_x <= 32) { + const int nty = (args.nrows_x + mmq_y - 1) / mmq_y; + const int ntzw = static_cast(args.nchannels_y * args.nsamples_y); + + int mmq_x_tiny = 0; + int ntiles_tiny = INT_MAX; + for (int mmq_x = mmq_x_max; mmq_x >= 8; mmq_x -= 8) { + const int granularity = mmq_get_granularity_host(mmq_x, cc); + if (mmq_x % granularity != 0) { + continue; + } + const size_t nbytes = mmq_get_nbytes_shared(mmq_x, mmq_y, cc, warp_size, nwarps); + if (nbytes > smpbo) { + continue; + } + const int ntx = (args.ncols_max + mmq_x - 1) / mmq_x; + const int grid_blocks = nty * ntx * ntzw; + if (grid_blocks <= 4 && ntx <= ntiles_tiny) { + mmq_x_tiny = mmq_x; + ntiles_tiny = ntx; + if (ntiles_tiny == 1) { + break; + } + } + } + if (mmq_x_tiny > 0) { + mmq_x_best = mmq_x_tiny; + } + } +#endif // GGML_USE_HIP + #if defined(GGML_USE_HIP) if (GGML_CUDA_CC_IS_RDNA3_5(cc)) { if (const char * force = getenv("MMQ_FORCE_MMQ_X")) { From 6e777cefdaca3612e50a0594c5e98d8781a6bd95 Mon Sep 17 00:00:00 2001 From: lichang Date: Thu, 25 Jun 2026 16:01:56 -0600 Subject: [PATCH 12/22] cuda: hoist Q6_K vec_dot scales on RDNA3.5 WMMA MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Precompute sclA = sc[k01/4]*d per k-slice in vec_dot_q6_K_q8_1_mma instead of reloading scales inside the j-loop. Q6_K FFN down 128×4096×12288 improves ~13.4 to ~15.2 TFLOPS with no P0 regression. Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 61 +++++++++++++++++++++++++++++++++++++- 1 file changed, 60 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index bb1564070bb7..f1395d0be304 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -2959,7 +2959,66 @@ static __device__ __forceinline__ void vec_dot_q6_K_q8_1_dp4a( template static __device__ __forceinline__ void vec_dot_q6_K_q8_1_mma( const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) { -#if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) +#if defined(RDNA3_5) && defined(AMD_WMMA_AVAILABLE) && !defined(AMD_MFMA_AVAILABLE) + // Hoist sclA = sc[k01/4]*d per k-slice — scales were loaded per (n,l) in j-loop. + constexpr data_layout input_layout = get_input_data_layout(); + typedef tile<16, 4, int, input_layout> tile_A; + typedef tile<16, 4, int, input_layout> tile_B; + typedef tile<16, 16, int, DATA_LAYOUT_J_MAJOR> tile_C; + + constexpr int granularity = mmq_get_granularity_device(mmq_x); + constexpr int rows_per_warp = granularity; + constexpr int ntx = rows_per_warp/tile_C::I; + + y += (threadIdx.y % ntx) * (tile_C::J*MMQ_TILE_Y_K); + + const int * x_qs = (const int *) x; + const float * x_df = (const float *) x_qs + MMQ_TILE_NE_K*2; + const int * x_sc = (const int *) x_df + MMQ_TILE_NE_K/QI6_K; + const int * y_qs = (const int *) y + 4; + const float * y_df = (const float *) y; + + const int i0 = (threadIdx.y / ntx) * rows_per_warp; + + for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += 4) { + const int k0 = k00 + k01; + + tile_A A[ntx]; + float sclA[ntx][tile_C::ne]; + +#pragma unroll + for (int n = 0; n < ntx; ++n) { + load_ldmatrix(A[n], x_qs + (i0 + n*tile_A::I)*MMQ_MMA_TILE_X_K_Q6_K + k0, MMQ_MMA_TILE_X_K_Q6_K); + +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + const int i = i0 + n*tile_A::I + tile_C::get_i(l); + const int8_t * sc = (const int8_t *) (x_sc + i*MMQ_MMA_TILE_X_K_Q6_K + k00/16); + sclA[n][l] = (float) sc[k01/4] * x_df[i*MMQ_MMA_TILE_X_K_Q6_K]; + } + } + +#pragma unroll + for (int j0 = 0; j0 < mmq_x; j0 += ntx*tile_C::J) { + tile_B B; + load_ldmatrix(B, y_qs + j0*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); + + const int j = j0 + tile_C::get_j(0); + const float dB = y_df[j*MMQ_TILE_Y_K + k01/QI8_1]; + +#pragma unroll + for (int n = 0; n < ntx; ++n) { + tile_C C; + mma(C, A[n], B); + +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + sum[(j0/tile_C::J + n)*tile_C::ne + l] += C.x[l]*sclA[n][l]*dB; + } + } + } + } +#elif defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) constexpr data_layout input_layout = get_input_data_layout(); typedef tile<16, 4, int, input_layout> tile_A; typedef tile<16, 4, int, input_layout> tile_B; From 4d9b940382762b58b8a4f4ca634b8a3f68ccad32 Mon Sep 17 00:00:00 2001 From: lichang Date: Thu, 25 Jun 2026 16:44:16 -0600 Subject: [PATCH 13/22] cuda: use mmq_x=128 for Q6_K narrow-N prefill on RDNA3.5 MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Prefer ntx=1 tile width when batch ≤128 and LDS allows, improving Q6_K FFN down (~40% faster vs mmq_x=64 on gfx1151). Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 29 +++++++++++++++++++++++++++++ 1 file changed, 29 insertions(+) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index f1395d0be304..5cf6d1b63c69 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -4879,6 +4879,35 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda mmq_x_best = mmq_x_tiny; } } + + // Narrow N (batch ≤128): Q6_K only — prefer mmq_x=128 for ntx=1 when LDS allows. + if (GGML_CUDA_CC_IS_RDNA3_5(cc) && type == GGML_TYPE_Q6_K && args.ncols_max <= mmq_x_max) { + const int ntx_cur = (args.ncols_max + mmq_x_best - 1) / mmq_x_best; + if (ntx_cur > 1) { + const int nty = (args.nrows_x + mmq_y - 1) / mmq_y; + const int ntzw = static_cast(args.nchannels_y * args.nsamples_y); + + for (int mmq_x = mmq_x_max; mmq_x >= 8; mmq_x -= 8) { + const int granularity = mmq_get_granularity_host(mmq_x, cc); + if (mmq_x % granularity != 0) { + continue; + } + const size_t nbytes = mmq_get_nbytes_shared(mmq_x, mmq_y, cc, warp_size, nwarps); + if (nbytes > smpbo) { + continue; + } + const int ntx = (args.ncols_max + mmq_x - 1) / mmq_x; + if (ntx != 1) { + continue; + } + const int grid_blocks = nty * ntx * ntzw; + if (grid_blocks >= 8) { + mmq_x_best = mmq_x; + break; + } + } + } + } #endif // GGML_USE_HIP #if defined(GGML_USE_HIP) From 03927f106c0cbe4c2c67da3446835cb0da24ab97 Mon Sep 17 00:00:00 2001 From: lichang Date: Mon, 29 Jun 2026 19:14:34 -0600 Subject: [PATCH 14/22] cuda: Q6_K MMQ B-prefetch, y-tile pipeline, and vec_dot profiling on RDNA3.5 Prefetch B across j0 with batched WMMA/epilogue, pad tile_y loads for mmq_x<=64 software pipelining, add MMQ_PROFILE_VEC_SUB (off by default), and pick Q6_K mmq_x=128 only when LDS fits dual-WG occupancy (smpbo/2). Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 234 +++++++++++++++++++++++++++++++++---- 1 file changed, 209 insertions(+), 25 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 5cf6d1b63c69..14e785649687 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -14,13 +14,23 @@ using namespace ggml_cuda_mma; #if defined(GGML_USE_HIP) enum mmq_profile_phase { - MMQ_PROFILE_LOAD_TILES = 0, - MMQ_PROFILE_Y_ACT = 1, - MMQ_PROFILE_VEC_DOT = 2, - MMQ_PROFILE_KB_ITERS = 3, - MMQ_PROFILE_N = 4, + MMQ_PROFILE_LOAD_TILES = 0, + MMQ_PROFILE_Y_ACT = 1, + MMQ_PROFILE_VEC_DOT = 2, + MMQ_PROFILE_KB_ITERS = 3, + MMQ_PROFILE_VEC_A = 4, // load_ldmatrix A + sclA hoist + MMQ_PROFILE_VEC_B = 5, // load_ldmatrix B + dB + MMQ_PROFILE_VEC_MMA = 6, // mma() + MMQ_PROFILE_VEC_EPILOGUE = 7, // sum[] FP32 epilogue + MMQ_PROFILE_N = 8, }; +__device__ uint64_t * g_mmq_profile = nullptr; + +static __device__ __forceinline__ void mmq_profile_bind(uint64_t * const profile) { + g_mmq_profile = profile; +} + static __device__ __forceinline__ bool mmq_profile_want_sample(uint64_t * const profile) { return profile && blockIdx.x == 0 && blockIdx.y == 0 && blockIdx.z == 0 && threadIdx.x == 0 && threadIdx.y == 0; @@ -61,6 +71,23 @@ static __device__ __forceinline__ void mmq_profile_phase_end(uint64_t * const pr profile[phase] += mmq_profile_phase_cycles(t0, clock64()); } +// vec_dot sub-phase timers: barrier-wrap so thread (0,0) measures WG wall time, not one lane only. +static __device__ __forceinline__ void mmq_profile_vec_sub_begin(uint64_t & t0) { + if (!g_mmq_profile) { + return; + } + __syncthreads(); + mmq_profile_phase_begin(g_mmq_profile, t0); +} + +static __device__ __forceinline__ void mmq_profile_vec_sub_end(uint64_t & t0, const int phase) { + if (!g_mmq_profile) { + return; + } + __syncthreads(); + mmq_profile_phase_end(g_mmq_profile, t0, phase); +} + static bool mmq_profile_phases_enabled() { static int enabled = -1; if (enabled < 0) { @@ -69,6 +96,18 @@ static bool mmq_profile_phases_enabled() { return enabled; } +static bool mmq_profile_vec_sub_enabled() { + static int enabled = -1; + if (enabled < 0) { + enabled = getenv("MMQ_PROFILE_VEC_SUB") ? 1 : 0; + } + return enabled; +} + +static bool mmq_profile_any_enabled() { + return mmq_profile_phases_enabled() || mmq_profile_vec_sub_enabled(); +} + struct mmq_profile_guard { ggml_cuda_pool_alloc buf; cudaStream_t stream = nullptr; @@ -115,10 +154,41 @@ struct mmq_profile_guard { (unsigned long long) h[MMQ_PROFILE_Y_ACT], (unsigned long long) h[MMQ_PROFILE_VEC_DOT], (unsigned long long) total); + if (mmq_profile_vec_sub_enabled()) { + const uint64_t vec_sub = h[MMQ_PROFILE_VEC_A] + h[MMQ_PROFILE_VEC_B] + + h[MMQ_PROFILE_VEC_MMA] + h[MMQ_PROFILE_VEC_EPILOGUE]; + if (vec_sub > 0 && vec_sub < 5000000000ULL) { + fprintf(stderr, + "[MMQ_PROFILE_VEC] type=%d mmq_x=%d nrows_x=%ld " + "vec_a=%.1f%% vec_b=%.1f%% vec_mma=%.1f%% vec_epi=%.1f%% " + "cycles_a=%llu cycles_b=%llu cycles_mma=%llu cycles_epi=%llu " + "cycles_vec_sub=%llu cycles_vec=%llu vec_sub_pct_of_vec=%.1f%% " + "mma_epi_ratio=%.3f\n", + type, mmq_x, (long) nrows_x, + 100.0*h[MMQ_PROFILE_VEC_A]/vec_sub, + 100.0*h[MMQ_PROFILE_VEC_B]/vec_sub, + 100.0*h[MMQ_PROFILE_VEC_MMA]/vec_sub, + 100.0*h[MMQ_PROFILE_VEC_EPILOGUE]/vec_sub, + (unsigned long long) h[MMQ_PROFILE_VEC_A], + (unsigned long long) h[MMQ_PROFILE_VEC_B], + (unsigned long long) h[MMQ_PROFILE_VEC_MMA], + (unsigned long long) h[MMQ_PROFILE_VEC_EPILOGUE], + (unsigned long long) vec_sub, + (unsigned long long) h[MMQ_PROFILE_VEC_DOT], + 100.0*vec_sub/h[MMQ_PROFILE_VEC_DOT], + (double) h[MMQ_PROFILE_VEC_MMA] / (double) h[MMQ_PROFILE_VEC_EPILOGUE]); + } + } } } }; #else +static __device__ __forceinline__ void mmq_profile_bind(uint64_t *) {} + +static __device__ __forceinline__ void mmq_profile_vec_sub_begin(uint64_t &) {} + +static __device__ __forceinline__ void mmq_profile_vec_sub_end(uint64_t &, const int) {} + static __device__ __forceinline__ void mmq_profile_kb_iter(uint64_t *) {} static __device__ __forceinline__ void mmq_profile_phase_begin(uint64_t *, uint64_t &) {} @@ -129,6 +199,14 @@ static bool mmq_profile_phases_enabled() { return false; } +static bool mmq_profile_vec_sub_enabled() { + return false; +} + +static bool mmq_profile_any_enabled() { + return false; +} + struct mmq_profile_guard { mmq_profile_guard(ggml_cuda_pool &, const bool, const int, const int, const int64_t, const int64_t, cudaStream_t) {} uint64_t * ptr() { return nullptr; } @@ -2980,6 +3058,80 @@ static __device__ __forceinline__ void vec_dot_q6_K_q8_1_mma( const int i0 = (threadIdx.y / ntx) * rows_per_warp; +#if defined(GGML_USE_HIP) + if (g_mmq_profile) { + for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += 4) { + const int k0 = k00 + k01; + + tile_A A[ntx]; + float sclA[ntx][tile_C::ne]; + + uint64_t t_vec_sub = 0; + mmq_profile_vec_sub_begin(t_vec_sub); +#pragma unroll + for (int n = 0; n < ntx; ++n) { + load_ldmatrix(A[n], x_qs + (i0 + n*tile_A::I)*MMQ_MMA_TILE_X_K_Q6_K + k0, MMQ_MMA_TILE_X_K_Q6_K); + +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + const int i = i0 + n*tile_A::I + tile_C::get_i(l); + const int8_t * sc = (const int8_t *) (x_sc + i*MMQ_MMA_TILE_X_K_Q6_K + k00/16); + sclA[n][l] = (float) sc[k01/4] * x_df[i*MMQ_MMA_TILE_X_K_Q6_K]; + } + } + mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_A); + + constexpr int j_step = ntx*tile_C::J; + + tile_B B0; + mmq_profile_vec_sub_begin(t_vec_sub); + load_ldmatrix(B0, y_qs + 0, MMQ_TILE_Y_K); + float dB0 = y_df[tile_C::get_j(0)*MMQ_TILE_Y_K + k01/QI8_1]; + mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_B); + +#pragma unroll + for (int j0 = 0; j0 < mmq_x; j0 += j_step) { + const int j0_next = j0 + j_step; + + tile_B B1; + float dB1 = 0.0f; + mmq_profile_vec_sub_begin(t_vec_sub); + if (j0_next < mmq_x) { + load_ldmatrix(B1, y_qs + j0_next*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); + const int jn = j0_next + tile_C::get_j(0); + dB1 = y_df[jn*MMQ_TILE_Y_K + k01/QI8_1]; + } + mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_B); + + tile_C C[ntx]; + + mmq_profile_vec_sub_begin(t_vec_sub); +#pragma unroll + for (int n = 0; n < ntx; ++n) { + mma(C[n], A[n], B0); + } + mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_MMA); + + mmq_profile_vec_sub_begin(t_vec_sub); +#pragma unroll + for (int n = 0; n < ntx; ++n) { +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + sum[(j0/tile_C::J + n)*tile_C::ne + l] += C[n].x[l]*sclA[n][l]*dB0; + } + } + mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_EPILOGUE); + + if (j0_next < mmq_x) { + B0 = B1; + dB0 = dB1; + } + } + } + return; + } +#endif // GGML_USE_HIP + for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += 4) { const int k0 = k00 + k01; @@ -2998,24 +3150,43 @@ static __device__ __forceinline__ void vec_dot_q6_K_q8_1_mma( } } + constexpr int j_step = ntx*tile_C::J; + + tile_B B0; + load_ldmatrix(B0, y_qs + 0, MMQ_TILE_Y_K); + float dB0 = y_df[tile_C::get_j(0)*MMQ_TILE_Y_K + k01/QI8_1]; + #pragma unroll - for (int j0 = 0; j0 < mmq_x; j0 += ntx*tile_C::J) { - tile_B B; - load_ldmatrix(B, y_qs + j0*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); + for (int j0 = 0; j0 < mmq_x; j0 += j_step) { + const int j0_next = j0 + j_step; - const int j = j0 + tile_C::get_j(0); - const float dB = y_df[j*MMQ_TILE_Y_K + k01/QI8_1]; + tile_B B1; + float dB1 = 0.0f; + if (j0_next < mmq_x) { + load_ldmatrix(B1, y_qs + j0_next*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); + const int jn = j0_next + tile_C::get_j(0); + dB1 = y_df[jn*MMQ_TILE_Y_K + k01/QI8_1]; + } + + tile_C C[ntx]; #pragma unroll for (int n = 0; n < ntx; ++n) { - tile_C C; - mma(C, A[n], B); + mma(C[n], A[n], B0); + } +#pragma unroll + for (int n = 0; n < ntx; ++n) { #pragma unroll for (int l = 0; l < tile_C::ne; ++l) { - sum[(j0/tile_C::J + n)*tile_C::ne + l] += C.x[l]*sclA[n][l]*dB; + sum[(j0/tile_C::J + n)*tile_C::ne + l] += C[n].x[l]*sclA[n][l]*dB0; } } + + if (j0_next < mmq_x) { + B0 = B1; + dB0 = dB1; + } } } #elif defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) @@ -3952,20 +4123,22 @@ struct mmq_type_traits { template static __device__ __forceinline__ void mmq_tile_y_load_global( int * __restrict__ tile_y, const int * __restrict__ by) { + constexpr int valid_elems = mmq_x*MMQ_TILE_Y_K; #pragma unroll for (int c = 0; c < nchunks; ++c) { const int l = c*(nwarps*warp_size) + threadIdx.y*warp_size + threadIdx.x; - tile_y[l] = by[l]; + tile_y[l] = l < valid_elems ? by[l] : 0; } } -template +template static __device__ __forceinline__ void mmq_tile_y_load_global_to_regs( const int * __restrict__ by, int (&cache)[nchunks]) { + constexpr int valid_elems = mmq_x*MMQ_TILE_Y_K; #pragma unroll for (int c = 0; c < nchunks; ++c) { const int l = c*(nwarps*warp_size) + threadIdx.y*warp_size + threadIdx.x; - cache[c] = by[l]; + cache[c] = l < valid_elems ? by[l] : 0; } } @@ -4018,10 +4191,14 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( float sum[mmq_x*mmq_y / (nwarps*warp_size)] = {0.0f}; +#if defined(GGML_USE_HIP) + mmq_profile_bind(mmq_profile); +#endif + constexpr int sz = sizeof(block_q8_1_mmq) / sizeof(int); #if defined(RDNA3_5) - constexpr int tile_y_elems = mmq_x*MMQ_TILE_Y_K; + constexpr int tile_y_elems = mmq_get_tile_y_k_padded_device(mmq_x, nwarps, warp_size); constexpr int tile_y_load_stride = nwarps*warp_size; if constexpr (mmq_x <= 64 && tile_y_elems % tile_y_load_stride == 0) { constexpr int tile_y_nchunks = tile_y_elems/tile_y_load_stride; @@ -4061,7 +4238,7 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( __syncthreads(); int y1_cache[tile_y_nchunks]; - mmq_tile_y_load_global_to_regs(by1, y1_cache); + mmq_tile_y_load_global_to_regs(by1, y1_cache); mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_Y_ACT); @@ -4079,7 +4256,7 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( const int kb0_next = kb0 + blocks_per_iter; if (kb0_next < kb0_stop) { const int * by0_next = y + ncols_y * (kb0_next * qk / ne_block) * sz; - mmq_tile_y_load_global_to_regs(by0_next, y0_next_cache); + mmq_tile_y_load_global_to_regs(by0_next, y0_next_cache); have_y0_prefetch = true; } @@ -4107,7 +4284,7 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( __syncthreads(); int y1_cache[tile_y_nchunks]; - mmq_tile_y_load_global_to_regs(by1, y1_cache); + mmq_tile_y_load_global_to_regs(by1, y1_cache); vec_dot(tile_x, tile_y, sum, 0); @@ -4120,7 +4297,7 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( const int kb0_next = kb0 + blocks_per_iter; if (kb0_next < kb0_stop) { const int * by0_next = y + ncols_y * (kb0_next * qk / ne_block) * sz; - mmq_tile_y_load_global_to_regs(by0_next, y0_next_cache); + mmq_tile_y_load_global_to_regs(by0_next, y0_next_cache); have_y0_prefetch = true; } @@ -4227,6 +4404,10 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( } else { write_back(sum, ids_dst, dst, stride_col_dst, tile_x_max_i, tile_y_max_j); } + +#if defined(GGML_USE_HIP) + mmq_profile_bind(nullptr); +#endif } @@ -4644,7 +4825,8 @@ static size_t mmq_get_nbytes_shared(const int mmq_x, const int mmq_y, const int const int tile_y_k_padded = mmq_get_tile_y_k_padded_host(mmq_x, cc, warp_size, nwarps); const size_t nbs_y_padded = std::max(nbs_y, (size_t) tile_y_k_padded*sizeof(int)); const int pad = GGML_CUDA_CC_IS_RDNA3_5(cc) ? 2*nwarps*warp_size : nwarps*warp_size; - return nbs_ids + nbs_x + GGML_PAD(nbs_y_padded, pad*sizeof(int)); + size_t nbytes = nbs_ids + nbs_x + GGML_PAD(nbs_y_padded, pad*sizeof(int)); + return nbytes; } template @@ -4659,7 +4841,7 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a const dim3 block_dims(warp_size, nwarps, 1); - mmq_profile_guard prof_guard(ctx.pool(id), mmq_profile_phases_enabled(), type, mmq_x, + mmq_profile_guard prof_guard(ctx.pool(id), mmq_profile_any_enabled(), type, mmq_x, args.nrows_x, args.ncols_max, stream); const int nbytes_shared = mmq_get_nbytes_shared(mmq_x, mmq_y, cc, warp_size, nwarps); @@ -4880,12 +5062,14 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda } } - // Narrow N (batch ≤128): Q6_K only — prefer mmq_x=128 for ntx=1 when LDS allows. + // Narrow N (batch ≤128): Q6_K only — prefer mmq_x=128 for ntx=1 when LDS allows + // dual-WG residency (nbytes <= smpbo/2). Forcing 128 at 38 KiB caps occupancy at 1 WG/CU. if (GGML_CUDA_CC_IS_RDNA3_5(cc) && type == GGML_TYPE_Q6_K && args.ncols_max <= mmq_x_max) { const int ntx_cur = (args.ncols_max + mmq_x_best - 1) / mmq_x_best; if (ntx_cur > 1) { const int nty = (args.nrows_x + mmq_y - 1) / mmq_y; const int ntzw = static_cast(args.nchannels_y * args.nsamples_y); + const size_t lds_dual_wg = smpbo / 2; for (int mmq_x = mmq_x_max; mmq_x >= 8; mmq_x -= 8) { const int granularity = mmq_get_granularity_host(mmq_x, cc); @@ -4893,7 +5077,7 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda continue; } const size_t nbytes = mmq_get_nbytes_shared(mmq_x, mmq_y, cc, warp_size, nwarps); - if (nbytes > smpbo) { + if (nbytes > smpbo || nbytes > lds_dual_wg) { continue; } const int ntx = (args.ncols_max + mmq_x - 1) / mmq_x; From cf83d4612eea816d97c148f767f2231e9a853108 Mon Sep 17 00:00:00 2001 From: lichang Date: Mon, 29 Jun 2026 19:21:44 -0600 Subject: [PATCH 15/22] cuda: drop RDNA3.5 activation y-pipeline from MMQ Software register pipelining of tile_y loads showed no measurable gain (~4% y_act slice) and slightly regressed pp128 on gfx1151; keep padded tile_y LDS layout and the simple sequential load path. Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 151 +------------------------------------ 1 file changed, 1 insertion(+), 150 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 14e785649687..f0ceb736b05f 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -4117,42 +4117,6 @@ struct mmq_type_traits { static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_q8_0_q8_1_dp4a; }; -#if defined(RDNA3_5) -// Software-pipeline activation tile loads: issue global loads into registers, run WMMA, -// then store to LDS. Overlaps memory latency with vec_dot on gfx115x (no ISA prefetch). -template -static __device__ __forceinline__ void mmq_tile_y_load_global( - int * __restrict__ tile_y, const int * __restrict__ by) { - constexpr int valid_elems = mmq_x*MMQ_TILE_Y_K; -#pragma unroll - for (int c = 0; c < nchunks; ++c) { - const int l = c*(nwarps*warp_size) + threadIdx.y*warp_size + threadIdx.x; - tile_y[l] = l < valid_elems ? by[l] : 0; - } -} - -template -static __device__ __forceinline__ void mmq_tile_y_load_global_to_regs( - const int * __restrict__ by, int (&cache)[nchunks]) { - constexpr int valid_elems = mmq_x*MMQ_TILE_Y_K; -#pragma unroll - for (int c = 0; c < nchunks; ++c) { - const int l = c*(nwarps*warp_size) + threadIdx.y*warp_size + threadIdx.x; - cache[c] = l < valid_elems ? by[l] : 0; - } -} - -template -static __device__ __forceinline__ void mmq_tile_y_store_regs( - int * __restrict__ tile_y, const int (&cache)[nchunks]) { -#pragma unroll - for (int c = 0; c < nchunks; ++c) { - const int l = c*(nwarps*warp_size) + threadIdx.y*warp_size + threadIdx.x; - tile_y[l] = cache[c]; - } -} -#endif // RDNA3_5 - template static __device__ __forceinline__ void mul_mat_q_process_tile( const char * __restrict__ x, const int offset_x, const int * __restrict__ y, @@ -4197,119 +4161,7 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( constexpr int sz = sizeof(block_q8_1_mmq) / sizeof(int); -#if defined(RDNA3_5) - constexpr int tile_y_elems = mmq_get_tile_y_k_padded_device(mmq_x, nwarps, warp_size); - constexpr int tile_y_load_stride = nwarps*warp_size; - if constexpr (mmq_x <= 64 && tile_y_elems % tile_y_load_stride == 0) { - constexpr int tile_y_nchunks = tile_y_elems/tile_y_load_stride; - - int y0_next_cache[tile_y_nchunks]; - bool have_y0_prefetch = false; - - if (mmq_profile) { - for (int kb0 = kb0_start; kb0 < kb0_stop; kb0 += blocks_per_iter) { - mmq_profile_kb_iter(mmq_profile); - - { - uint64_t t0 = 0; - __syncthreads(); - mmq_profile_phase_begin(mmq_profile, t0); - load_tiles(x, tile_x, offset_x + kb0, tile_x_max_i, stride_row_x); - __syncthreads(); - mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_LOAD_TILES); - } - - const int yk = kb0 * qk / ne_block; - const int * by0 = y + ncols_y * yk * sz; - const int * by1 = y + ncols_y * (yk + 1) * sz; - - { - uint64_t t0 = 0; - __syncthreads(); - mmq_profile_phase_begin(mmq_profile, t0); - - if (have_y0_prefetch) { - mmq_tile_y_store_regs(tile_y, y0_next_cache); - have_y0_prefetch = false; - } else { - mmq_tile_y_load_global(tile_y, by0); - } - - __syncthreads(); - - int y1_cache[tile_y_nchunks]; - mmq_tile_y_load_global_to_regs(by1, y1_cache); - - mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_Y_ACT); - - __syncthreads(); - mmq_profile_phase_begin(mmq_profile, t0); - - vec_dot(tile_x, tile_y, sum, 0); - - __syncthreads(); - - mmq_tile_y_store_regs(tile_y, y1_cache); - - __syncthreads(); - - const int kb0_next = kb0 + blocks_per_iter; - if (kb0_next < kb0_stop) { - const int * by0_next = y + ncols_y * (kb0_next * qk / ne_block) * sz; - mmq_tile_y_load_global_to_regs(by0_next, y0_next_cache); - have_y0_prefetch = true; - } - - vec_dot(tile_x, tile_y, sum, MMQ_TILE_NE_K); - - __syncthreads(); - mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_VEC_DOT); - } - } - } else { - 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); - - const int yk = kb0 * qk / ne_block; - const int * by0 = y + ncols_y * yk * sz; - const int * by1 = y + ncols_y * (yk + 1) * sz; - - if (have_y0_prefetch) { - mmq_tile_y_store_regs(tile_y, y0_next_cache); - have_y0_prefetch = false; - } else { - mmq_tile_y_load_global(tile_y, by0); - } - - __syncthreads(); - - int y1_cache[tile_y_nchunks]; - mmq_tile_y_load_global_to_regs(by1, y1_cache); - - vec_dot(tile_x, tile_y, sum, 0); - - __syncthreads(); - - mmq_tile_y_store_regs(tile_y, y1_cache); - - __syncthreads(); - - const int kb0_next = kb0 + blocks_per_iter; - if (kb0_next < kb0_stop) { - const int * by0_next = y + ncols_y * (kb0_next * qk / ne_block) * sz; - mmq_tile_y_load_global_to_regs(by0_next, y0_next_cache); - have_y0_prefetch = true; - } - - vec_dot(tile_x, tile_y, sum, MMQ_TILE_NE_K); - - __syncthreads(); - } - } - } else -#endif // RDNA3_5 - { - if (mmq_profile) { + if (mmq_profile) { for (int kb0 = kb0_start; kb0 < kb0_stop; kb0 += blocks_per_iter) { mmq_profile_kb_iter(mmq_profile); @@ -4397,7 +4249,6 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( __syncthreads(); } } - } if (fixup) { write_back(sum, ids_dst, tmp_fixup + blockIdx.x*(mmq_x*mmq_y), mmq_y, mmq_y, mmq_x); From d4511bd7484fee0ec922910731d9318e4fad29d6 Mon Sep 17 00:00:00 2001 From: lichang Date: Mon, 29 Jun 2026 19:39:36 -0600 Subject: [PATCH 16/22] cuda: batch Q6_K B loads with load_generic on RDNA3.5 WMMA Load all B tiles per k-slice via load_generic (faster than load_ldmatrix from tile_y), interleave with A on the hot path, and merge the duplicate vec_dot profile loop into one path. Use fast process_tile when only MMQ_PROFILE_VEC_SUB is set. Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 224 +++++++++++++++---------------------- 1 file changed, 90 insertions(+), 134 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index f0ceb736b05f..f66c514db737 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -19,7 +19,7 @@ enum mmq_profile_phase { MMQ_PROFILE_VEC_DOT = 2, MMQ_PROFILE_KB_ITERS = 3, MMQ_PROFILE_VEC_A = 4, // load_ldmatrix A + sclA hoist - MMQ_PROFILE_VEC_B = 5, // load_ldmatrix B + dB + MMQ_PROFILE_VEC_B = 5, // load_generic B + dB from tile_y MMQ_PROFILE_VEC_MMA = 6, // mma() MMQ_PROFILE_VEC_EPILOGUE = 7, // sum[] FP32 epilogue MMQ_PROFILE_N = 8, @@ -27,6 +27,8 @@ enum mmq_profile_phase { __device__ uint64_t * g_mmq_profile = nullptr; +__device__ int g_mmq_profile_outer = 0; + static __device__ __forceinline__ void mmq_profile_bind(uint64_t * const profile) { g_mmq_profile = profile; } @@ -195,6 +197,8 @@ static __device__ __forceinline__ void mmq_profile_phase_begin(uint64_t *, uint6 static __device__ __forceinline__ void mmq_profile_phase_end(uint64_t *, const uint64_t, const int) {} +static __device__ int g_mmq_profile_outer = 0; + static bool mmq_profile_phases_enabled() { return false; } @@ -3058,86 +3062,19 @@ static __device__ __forceinline__ void vec_dot_q6_K_q8_1_mma( const int i0 = (threadIdx.y / ntx) * rows_per_warp; -#if defined(GGML_USE_HIP) - if (g_mmq_profile) { - for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += 4) { - const int k0 = k00 + k01; - - tile_A A[ntx]; - float sclA[ntx][tile_C::ne]; - - uint64_t t_vec_sub = 0; - mmq_profile_vec_sub_begin(t_vec_sub); -#pragma unroll - for (int n = 0; n < ntx; ++n) { - load_ldmatrix(A[n], x_qs + (i0 + n*tile_A::I)*MMQ_MMA_TILE_X_K_Q6_K + k0, MMQ_MMA_TILE_X_K_Q6_K); - -#pragma unroll - for (int l = 0; l < tile_C::ne; ++l) { - const int i = i0 + n*tile_A::I + tile_C::get_i(l); - const int8_t * sc = (const int8_t *) (x_sc + i*MMQ_MMA_TILE_X_K_Q6_K + k00/16); - sclA[n][l] = (float) sc[k01/4] * x_df[i*MMQ_MMA_TILE_X_K_Q6_K]; - } - } - mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_A); - - constexpr int j_step = ntx*tile_C::J; - - tile_B B0; - mmq_profile_vec_sub_begin(t_vec_sub); - load_ldmatrix(B0, y_qs + 0, MMQ_TILE_Y_K); - float dB0 = y_df[tile_C::get_j(0)*MMQ_TILE_Y_K + k01/QI8_1]; - mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_B); - -#pragma unroll - for (int j0 = 0; j0 < mmq_x; j0 += j_step) { - const int j0_next = j0 + j_step; - - tile_B B1; - float dB1 = 0.0f; - mmq_profile_vec_sub_begin(t_vec_sub); - if (j0_next < mmq_x) { - load_ldmatrix(B1, y_qs + j0_next*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); - const int jn = j0_next + tile_C::get_j(0); - dB1 = y_df[jn*MMQ_TILE_Y_K + k01/QI8_1]; - } - mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_B); - - tile_C C[ntx]; - - mmq_profile_vec_sub_begin(t_vec_sub); -#pragma unroll - for (int n = 0; n < ntx; ++n) { - mma(C[n], A[n], B0); - } - mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_MMA); - - mmq_profile_vec_sub_begin(t_vec_sub); -#pragma unroll - for (int n = 0; n < ntx; ++n) { -#pragma unroll - for (int l = 0; l < tile_C::ne; ++l) { - sum[(j0/tile_C::J + n)*tile_C::ne + l] += C[n].x[l]*sclA[n][l]*dB0; - } - } - mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_EPILOGUE); - - if (j0_next < mmq_x) { - B0 = B1; - dB0 = dB1; - } - } - } - return; - } -#endif // GGML_USE_HIP + constexpr int j_step = ntx*tile_C::J; + constexpr int nj = (mmq_x + j_step - 1) / j_step; for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += 4) { const int k0 = k00 + k01; tile_A A[ntx]; float sclA[ntx][tile_C::ne]; + tile_B Btile[nj]; + float dBt[nj]; + uint64_t t_vec_sub = 0; + mmq_profile_vec_sub_begin(t_vec_sub); #pragma unroll for (int n = 0; n < ntx; ++n) { load_ldmatrix(A[n], x_qs + (i0 + n*tile_A::I)*MMQ_MMA_TILE_X_K_Q6_K + k0, MMQ_MMA_TILE_X_K_Q6_K); @@ -3148,45 +3085,55 @@ static __device__ __forceinline__ void vec_dot_q6_K_q8_1_mma( const int8_t * sc = (const int8_t *) (x_sc + i*MMQ_MMA_TILE_X_K_Q6_K + k00/16); sclA[n][l] = (float) sc[k01/4] * x_df[i*MMQ_MMA_TILE_X_K_Q6_K]; } - } - - constexpr int j_step = ntx*tile_C::J; - tile_B B0; - load_ldmatrix(B0, y_qs + 0, MMQ_TILE_Y_K); - float dB0 = y_df[tile_C::get_j(0)*MMQ_TILE_Y_K + k01/QI8_1]; + if (!g_mmq_profile && n < nj) { + const int j0 = n*j_step; + load_generic(Btile[n], y_qs + j0*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); + dBt[n] = y_df[(j0 + tile_C::get_j(0))*MMQ_TILE_Y_K + k01/QI8_1]; + } + } + mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_A); + mmq_profile_vec_sub_begin(t_vec_sub); + if (g_mmq_profile) { #pragma unroll - for (int j0 = 0; j0 < mmq_x; j0 += j_step) { - const int j0_next = j0 + j_step; - - tile_B B1; - float dB1 = 0.0f; - if (j0_next < mmq_x) { - load_ldmatrix(B1, y_qs + j0_next*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); - const int jn = j0_next + tile_C::get_j(0); - dB1 = y_df[jn*MMQ_TILE_Y_K + k01/QI8_1]; + for (int t = 0; t < nj; ++t) { + const int j0 = t*j_step; + load_generic(Btile[t], y_qs + j0*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); + dBt[t] = y_df[(j0 + tile_C::get_j(0))*MMQ_TILE_Y_K + k01/QI8_1]; + } + } else { +#pragma unroll + for (int t = ntx; t < nj; ++t) { + const int j0 = t*j_step; + load_generic(Btile[t], y_qs + j0*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); + dBt[t] = y_df[(j0 + tile_C::get_j(0))*MMQ_TILE_Y_K + k01/QI8_1]; } + } + mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_B); + +#pragma unroll + for (int t = 0; t < nj; ++t) { + const int j0 = t*j_step; tile_C C[ntx]; + mmq_profile_vec_sub_begin(t_vec_sub); #pragma unroll for (int n = 0; n < ntx; ++n) { - mma(C[n], A[n], B0); + mma(C[n], A[n], Btile[t]); } + mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_MMA); + mmq_profile_vec_sub_begin(t_vec_sub); #pragma unroll for (int n = 0; n < ntx; ++n) { #pragma unroll for (int l = 0; l < tile_C::ne; ++l) { - sum[(j0/tile_C::J + n)*tile_C::ne + l] += C[n].x[l]*sclA[n][l]*dB0; + sum[(j0/tile_C::J + n)*tile_C::ne + l] += C[n].x[l]*sclA[n][l]*dBt[t]; } } - - if (j0_next < mmq_x) { - B0 = B1; - dB0 = dB1; - } + mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_EPILOGUE); } } #elif defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) @@ -4161,56 +4108,60 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( constexpr int sz = sizeof(block_q8_1_mmq) / sizeof(int); +#if defined(GGML_USE_HIP) + if (g_mmq_profile_outer) { +#else if (mmq_profile) { - for (int kb0 = kb0_start; kb0 < kb0_stop; kb0 += blocks_per_iter) { - mmq_profile_kb_iter(mmq_profile); +#endif + for (int kb0 = kb0_start; kb0 < kb0_stop; kb0 += blocks_per_iter) { + mmq_profile_kb_iter(mmq_profile); - { - uint64_t t0 = 0; - __syncthreads(); - mmq_profile_phase_begin(mmq_profile, t0); - load_tiles(x, tile_x, offset_x + kb0, tile_x_max_i, stride_row_x); - __syncthreads(); - mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_LOAD_TILES); - } + { + uint64_t t0 = 0; + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); + load_tiles(x, tile_x, offset_x + kb0, tile_x_max_i, stride_row_x); + __syncthreads(); + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_LOAD_TILES); + } - const int yk = kb0 * qk / ne_block; - const int * by0 = y + ncols_y * yk * sz; - const int * by1 = y + ncols_y * (yk + 1) * sz; + const int yk = kb0 * qk / ne_block; + const int * by0 = y + ncols_y * yk * sz; + const int * by1 = y + ncols_y * (yk + 1) * sz; - { - uint64_t t0 = 0; - __syncthreads(); - mmq_profile_phase_begin(mmq_profile, t0); + { + uint64_t t0 = 0; + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); #pragma unroll - for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { - int l = l0 + threadIdx.y*warp_size + threadIdx.x; + for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { + int l = l0 + threadIdx.y*warp_size + threadIdx.x; - tile_y[l] = by0[l]; - } - __syncthreads(); - mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_Y_ACT); + tile_y[l] = by0[l]; } + __syncthreads(); + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_Y_ACT); + } - { - uint64_t t0 = 0; - __syncthreads(); - mmq_profile_phase_begin(mmq_profile, t0); - vec_dot(tile_x, tile_y, sum, 0); - __syncthreads(); + { + uint64_t t0 = 0; + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); + vec_dot(tile_x, tile_y, sum, 0); + __syncthreads(); #pragma unroll - for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { - int l = l0 + threadIdx.y*warp_size + threadIdx.x; + for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { + int l = l0 + threadIdx.y*warp_size + threadIdx.x; - tile_y[l] = by1[l]; - } - __syncthreads(); - vec_dot(tile_x, tile_y, sum, MMQ_TILE_NE_K); - __syncthreads(); - mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_VEC_DOT); + tile_y[l] = by1[l]; } + __syncthreads(); + vec_dot(tile_x, tile_y, sum, MMQ_TILE_NE_K); + __syncthreads(); + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_VEC_DOT); } - } else { + } + } else { 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); @@ -4697,6 +4648,11 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a const int nbytes_shared = mmq_get_nbytes_shared(mmq_x, mmq_y, cc, warp_size, nwarps); +#if defined(GGML_USE_HIP) + const int outer_phases = mmq_profile_phases_enabled() ? 1 : 0; + CUDA_CHECK(cudaMemcpyToSymbol(g_mmq_profile_outer, &outer_phases, sizeof(int))); +#endif + CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); From eb2a7de02d7003a4a94d0f685283ee3e9e16e34c Mon Sep 17 00:00:00 2001 From: lichang Date: Mon, 29 Jun 2026 20:00:56 -0600 Subject: [PATCH 17/22] cuda: restore Q6_K mmq_x=128 for narrow-N FFN down on RDNA3.5 Revert the smpbo/2 gate added in 03927f106 for the Q6_K narrow-N launch heuristic. Q6 at mmq_x=128 (38 KiB, ntx=1) is ~40% faster than mmq_x=64 with ntx=2 even at 1 WG/CU; the dual-WG policy belongs on Q4_K only. Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index f66c514db737..4e8e2444e6d6 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -4869,14 +4869,14 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda } } - // Narrow N (batch ≤128): Q6_K only — prefer mmq_x=128 for ntx=1 when LDS allows - // dual-WG residency (nbytes <= smpbo/2). Forcing 128 at 38 KiB caps occupancy at 1 WG/CU. + // Narrow N (batch ≤128): Q6_K only — prefer mmq_x=128 for ntx=1 when LDS allows. + // Q6_K at mmq_x=128 uses ~38 KiB (1 WG/CU) but beats mmq_x=64 with ntx=2; do not + // require smpbo/2 here — that policy is for Q4_K dual-WG, not Q6 narrow-N. if (GGML_CUDA_CC_IS_RDNA3_5(cc) && type == GGML_TYPE_Q6_K && args.ncols_max <= mmq_x_max) { const int ntx_cur = (args.ncols_max + mmq_x_best - 1) / mmq_x_best; if (ntx_cur > 1) { const int nty = (args.nrows_x + mmq_y - 1) / mmq_y; const int ntzw = static_cast(args.nchannels_y * args.nsamples_y); - const size_t lds_dual_wg = smpbo / 2; for (int mmq_x = mmq_x_max; mmq_x >= 8; mmq_x -= 8) { const int granularity = mmq_get_granularity_host(mmq_x, cc); @@ -4884,7 +4884,7 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda continue; } const size_t nbytes = mmq_get_nbytes_shared(mmq_x, mmq_y, cc, warp_size, nwarps); - if (nbytes > smpbo || nbytes > lds_dual_wg) { + if (nbytes > smpbo) { continue; } const int ntx = (args.ncols_max + mmq_x - 1) / mmq_x; From 49b036ded2167eec7a1055ce3c8cccc817f30696 Mon Sep 17 00:00:00 2001 From: lichang Date: Mon, 29 Jun 2026 23:00:43 -0600 Subject: [PATCH 18/22] cuda: restore Q6_K vec_dot load_ldmatrix path on RDNA3.5 Revert the d4511bd74 B-batch/load_generic vec_dot rewrite that regressed Q6 FFN down from ~13 to ~8 TFLOPS under rocprof trace (~1.5ms to ~1.1ms per launch). Keep the 4d9b94038 j0 loop with load_ldmatrix for B; profile instrumentation remains on a separate g_mmq_profile early-return path only. Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 107 ++++++++++++++++++++++--------------- 1 file changed, 64 insertions(+), 43 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 4e8e2444e6d6..4d43d1be0542 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -3062,19 +3062,69 @@ static __device__ __forceinline__ void vec_dot_q6_K_q8_1_mma( const int i0 = (threadIdx.y / ntx) * rows_per_warp; - constexpr int j_step = ntx*tile_C::J; - constexpr int nj = (mmq_x + j_step - 1) / j_step; +#if defined(GGML_USE_HIP) + if (g_mmq_profile) { + for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += 4) { + const int k0 = k00 + k01; + + tile_A A[ntx]; + float sclA[ntx][tile_C::ne]; + + uint64_t t_vec_sub = 0; + mmq_profile_vec_sub_begin(t_vec_sub); +#pragma unroll + for (int n = 0; n < ntx; ++n) { + load_ldmatrix(A[n], x_qs + (i0 + n*tile_A::I)*MMQ_MMA_TILE_X_K_Q6_K + k0, MMQ_MMA_TILE_X_K_Q6_K); + +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + const int i = i0 + n*tile_A::I + tile_C::get_i(l); + const int8_t * sc = (const int8_t *) (x_sc + i*MMQ_MMA_TILE_X_K_Q6_K + k00/16); + sclA[n][l] = (float) sc[k01/4] * x_df[i*MMQ_MMA_TILE_X_K_Q6_K]; + } + } + mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_A); + +#pragma unroll + for (int j0 = 0; j0 < mmq_x; j0 += ntx*tile_C::J) { + tile_B B; + + mmq_profile_vec_sub_begin(t_vec_sub); + load_ldmatrix(B, y_qs + j0*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); + const int j = j0 + tile_C::get_j(0); + const float dB = y_df[j*MMQ_TILE_Y_K + k01/QI8_1]; + mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_B); + + tile_C C[ntx]; + + mmq_profile_vec_sub_begin(t_vec_sub); +#pragma unroll + for (int n = 0; n < ntx; ++n) { + mma(C[n], A[n], B); + } + mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_MMA); + + mmq_profile_vec_sub_begin(t_vec_sub); +#pragma unroll + for (int n = 0; n < ntx; ++n) { +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + sum[(j0/tile_C::J + n)*tile_C::ne + l] += C[n].x[l]*sclA[n][l]*dB; + } + } + mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_EPILOGUE); + } + } + return; + } +#endif // GGML_USE_HIP for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += 4) { const int k0 = k00 + k01; tile_A A[ntx]; float sclA[ntx][tile_C::ne]; - tile_B Btile[nj]; - float dBt[nj]; - uint64_t t_vec_sub = 0; - mmq_profile_vec_sub_begin(t_vec_sub); #pragma unroll for (int n = 0; n < ntx; ++n) { load_ldmatrix(A[n], x_qs + (i0 + n*tile_A::I)*MMQ_MMA_TILE_X_K_Q6_K + k0, MMQ_MMA_TILE_X_K_Q6_K); @@ -3085,55 +3135,26 @@ static __device__ __forceinline__ void vec_dot_q6_K_q8_1_mma( const int8_t * sc = (const int8_t *) (x_sc + i*MMQ_MMA_TILE_X_K_Q6_K + k00/16); sclA[n][l] = (float) sc[k01/4] * x_df[i*MMQ_MMA_TILE_X_K_Q6_K]; } - - if (!g_mmq_profile && n < nj) { - const int j0 = n*j_step; - load_generic(Btile[n], y_qs + j0*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); - dBt[n] = y_df[(j0 + tile_C::get_j(0))*MMQ_TILE_Y_K + k01/QI8_1]; - } - } - mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_A); - - mmq_profile_vec_sub_begin(t_vec_sub); - if (g_mmq_profile) { -#pragma unroll - for (int t = 0; t < nj; ++t) { - const int j0 = t*j_step; - load_generic(Btile[t], y_qs + j0*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); - dBt[t] = y_df[(j0 + tile_C::get_j(0))*MMQ_TILE_Y_K + k01/QI8_1]; - } - } else { -#pragma unroll - for (int t = ntx; t < nj; ++t) { - const int j0 = t*j_step; - load_generic(Btile[t], y_qs + j0*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); - dBt[t] = y_df[(j0 + tile_C::get_j(0))*MMQ_TILE_Y_K + k01/QI8_1]; - } } - mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_B); #pragma unroll - for (int t = 0; t < nj; ++t) { - const int j0 = t*j_step; + for (int j0 = 0; j0 < mmq_x; j0 += ntx*tile_C::J) { + tile_B B; + load_ldmatrix(B, y_qs + j0*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); - tile_C C[ntx]; + const int j = j0 + tile_C::get_j(0); + const float dB = y_df[j*MMQ_TILE_Y_K + k01/QI8_1]; - mmq_profile_vec_sub_begin(t_vec_sub); #pragma unroll for (int n = 0; n < ntx; ++n) { - mma(C[n], A[n], Btile[t]); - } - mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_MMA); + tile_C C; + mma(C, A[n], B); - mmq_profile_vec_sub_begin(t_vec_sub); -#pragma unroll - for (int n = 0; n < ntx; ++n) { #pragma unroll for (int l = 0; l < tile_C::ne; ++l) { - sum[(j0/tile_C::J + n)*tile_C::ne + l] += C[n].x[l]*sclA[n][l]*dBt[t]; + sum[(j0/tile_C::J + n)*tile_C::ne + l] += C.x[l]*sclA[n][l]*dB; } } - mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_EPILOGUE); } } #elif defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) From 94e9d1dc3f50b97a441d3cf3b13a201b95332385 Mon Sep 17 00:00:00 2001 From: lichang Date: Mon, 29 Jun 2026 23:35:09 -0600 Subject: [PATCH 19/22] cuda: remove always-on MMQ profiling overhead on RDNA3.5 hot path Revert g_mmq_profile_outer/cudaMemcpyToSymbol and Q6 vec_dot global profile branches that regressed Q6_K FFN down ~15% under rocprof trace. Restore RDNA3.5 y-tile software pipeline and MMQ_PROFILE_PHASES-only instrumentation to match 4d9b94038 prefill performance. Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 385 ++++++++++++++++++------------------- 1 file changed, 189 insertions(+), 196 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 4d43d1be0542..7d59cf940fb0 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -14,25 +14,13 @@ using namespace ggml_cuda_mma; #if defined(GGML_USE_HIP) enum mmq_profile_phase { - MMQ_PROFILE_LOAD_TILES = 0, - MMQ_PROFILE_Y_ACT = 1, - MMQ_PROFILE_VEC_DOT = 2, - MMQ_PROFILE_KB_ITERS = 3, - MMQ_PROFILE_VEC_A = 4, // load_ldmatrix A + sclA hoist - MMQ_PROFILE_VEC_B = 5, // load_generic B + dB from tile_y - MMQ_PROFILE_VEC_MMA = 6, // mma() - MMQ_PROFILE_VEC_EPILOGUE = 7, // sum[] FP32 epilogue - MMQ_PROFILE_N = 8, + MMQ_PROFILE_LOAD_TILES = 0, + MMQ_PROFILE_Y_ACT = 1, + MMQ_PROFILE_VEC_DOT = 2, + MMQ_PROFILE_KB_ITERS = 3, + MMQ_PROFILE_N = 4, }; -__device__ uint64_t * g_mmq_profile = nullptr; - -__device__ int g_mmq_profile_outer = 0; - -static __device__ __forceinline__ void mmq_profile_bind(uint64_t * const profile) { - g_mmq_profile = profile; -} - static __device__ __forceinline__ bool mmq_profile_want_sample(uint64_t * const profile) { return profile && blockIdx.x == 0 && blockIdx.y == 0 && blockIdx.z == 0 && threadIdx.x == 0 && threadIdx.y == 0; @@ -73,23 +61,6 @@ static __device__ __forceinline__ void mmq_profile_phase_end(uint64_t * const pr profile[phase] += mmq_profile_phase_cycles(t0, clock64()); } -// vec_dot sub-phase timers: barrier-wrap so thread (0,0) measures WG wall time, not one lane only. -static __device__ __forceinline__ void mmq_profile_vec_sub_begin(uint64_t & t0) { - if (!g_mmq_profile) { - return; - } - __syncthreads(); - mmq_profile_phase_begin(g_mmq_profile, t0); -} - -static __device__ __forceinline__ void mmq_profile_vec_sub_end(uint64_t & t0, const int phase) { - if (!g_mmq_profile) { - return; - } - __syncthreads(); - mmq_profile_phase_end(g_mmq_profile, t0, phase); -} - static bool mmq_profile_phases_enabled() { static int enabled = -1; if (enabled < 0) { @@ -98,18 +69,6 @@ static bool mmq_profile_phases_enabled() { return enabled; } -static bool mmq_profile_vec_sub_enabled() { - static int enabled = -1; - if (enabled < 0) { - enabled = getenv("MMQ_PROFILE_VEC_SUB") ? 1 : 0; - } - return enabled; -} - -static bool mmq_profile_any_enabled() { - return mmq_profile_phases_enabled() || mmq_profile_vec_sub_enabled(); -} - struct mmq_profile_guard { ggml_cuda_pool_alloc buf; cudaStream_t stream = nullptr; @@ -156,61 +115,20 @@ struct mmq_profile_guard { (unsigned long long) h[MMQ_PROFILE_Y_ACT], (unsigned long long) h[MMQ_PROFILE_VEC_DOT], (unsigned long long) total); - if (mmq_profile_vec_sub_enabled()) { - const uint64_t vec_sub = h[MMQ_PROFILE_VEC_A] + h[MMQ_PROFILE_VEC_B] - + h[MMQ_PROFILE_VEC_MMA] + h[MMQ_PROFILE_VEC_EPILOGUE]; - if (vec_sub > 0 && vec_sub < 5000000000ULL) { - fprintf(stderr, - "[MMQ_PROFILE_VEC] type=%d mmq_x=%d nrows_x=%ld " - "vec_a=%.1f%% vec_b=%.1f%% vec_mma=%.1f%% vec_epi=%.1f%% " - "cycles_a=%llu cycles_b=%llu cycles_mma=%llu cycles_epi=%llu " - "cycles_vec_sub=%llu cycles_vec=%llu vec_sub_pct_of_vec=%.1f%% " - "mma_epi_ratio=%.3f\n", - type, mmq_x, (long) nrows_x, - 100.0*h[MMQ_PROFILE_VEC_A]/vec_sub, - 100.0*h[MMQ_PROFILE_VEC_B]/vec_sub, - 100.0*h[MMQ_PROFILE_VEC_MMA]/vec_sub, - 100.0*h[MMQ_PROFILE_VEC_EPILOGUE]/vec_sub, - (unsigned long long) h[MMQ_PROFILE_VEC_A], - (unsigned long long) h[MMQ_PROFILE_VEC_B], - (unsigned long long) h[MMQ_PROFILE_VEC_MMA], - (unsigned long long) h[MMQ_PROFILE_VEC_EPILOGUE], - (unsigned long long) vec_sub, - (unsigned long long) h[MMQ_PROFILE_VEC_DOT], - 100.0*vec_sub/h[MMQ_PROFILE_VEC_DOT], - (double) h[MMQ_PROFILE_VEC_MMA] / (double) h[MMQ_PROFILE_VEC_EPILOGUE]); - } - } } } }; #else -static __device__ __forceinline__ void mmq_profile_bind(uint64_t *) {} - -static __device__ __forceinline__ void mmq_profile_vec_sub_begin(uint64_t &) {} - -static __device__ __forceinline__ void mmq_profile_vec_sub_end(uint64_t &, const int) {} - static __device__ __forceinline__ void mmq_profile_kb_iter(uint64_t *) {} static __device__ __forceinline__ void mmq_profile_phase_begin(uint64_t *, uint64_t &) {} static __device__ __forceinline__ void mmq_profile_phase_end(uint64_t *, const uint64_t, const int) {} -static __device__ int g_mmq_profile_outer = 0; - static bool mmq_profile_phases_enabled() { return false; } -static bool mmq_profile_vec_sub_enabled() { - return false; -} - -static bool mmq_profile_any_enabled() { - return false; -} - struct mmq_profile_guard { mmq_profile_guard(ggml_cuda_pool &, const bool, const int, const int, const int64_t, const int64_t, cudaStream_t) {} uint64_t * ptr() { return nullptr; } @@ -3062,63 +2980,6 @@ static __device__ __forceinline__ void vec_dot_q6_K_q8_1_mma( const int i0 = (threadIdx.y / ntx) * rows_per_warp; -#if defined(GGML_USE_HIP) - if (g_mmq_profile) { - for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += 4) { - const int k0 = k00 + k01; - - tile_A A[ntx]; - float sclA[ntx][tile_C::ne]; - - uint64_t t_vec_sub = 0; - mmq_profile_vec_sub_begin(t_vec_sub); -#pragma unroll - for (int n = 0; n < ntx; ++n) { - load_ldmatrix(A[n], x_qs + (i0 + n*tile_A::I)*MMQ_MMA_TILE_X_K_Q6_K + k0, MMQ_MMA_TILE_X_K_Q6_K); - -#pragma unroll - for (int l = 0; l < tile_C::ne; ++l) { - const int i = i0 + n*tile_A::I + tile_C::get_i(l); - const int8_t * sc = (const int8_t *) (x_sc + i*MMQ_MMA_TILE_X_K_Q6_K + k00/16); - sclA[n][l] = (float) sc[k01/4] * x_df[i*MMQ_MMA_TILE_X_K_Q6_K]; - } - } - mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_A); - -#pragma unroll - for (int j0 = 0; j0 < mmq_x; j0 += ntx*tile_C::J) { - tile_B B; - - mmq_profile_vec_sub_begin(t_vec_sub); - load_ldmatrix(B, y_qs + j0*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); - const int j = j0 + tile_C::get_j(0); - const float dB = y_df[j*MMQ_TILE_Y_K + k01/QI8_1]; - mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_B); - - tile_C C[ntx]; - - mmq_profile_vec_sub_begin(t_vec_sub); -#pragma unroll - for (int n = 0; n < ntx; ++n) { - mma(C[n], A[n], B); - } - mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_MMA); - - mmq_profile_vec_sub_begin(t_vec_sub); -#pragma unroll - for (int n = 0; n < ntx; ++n) { -#pragma unroll - for (int l = 0; l < tile_C::ne; ++l) { - sum[(j0/tile_C::J + n)*tile_C::ne + l] += C[n].x[l]*sclA[n][l]*dB; - } - } - mmq_profile_vec_sub_end(t_vec_sub, MMQ_PROFILE_VEC_EPILOGUE); - } - } - return; - } -#endif // GGML_USE_HIP - for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += 4) { const int k0 = k00 + k01; @@ -4085,6 +3946,41 @@ struct mmq_type_traits { static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_q8_0_q8_1_dp4a; }; + +#if defined(RDNA3_5) +// Software-pipeline activation tile loads: issue global loads into registers, run WMMA, +// then store to LDS. Overlaps memory latency with vec_dot on gfx115x (no ISA prefetch). +template +static __device__ __forceinline__ void mmq_tile_y_load_global( + int * __restrict__ tile_y, const int * __restrict__ by) { +#pragma unroll + for (int c = 0; c < nchunks; ++c) { + const int l = c*(nwarps*warp_size) + threadIdx.y*warp_size + threadIdx.x; + tile_y[l] = by[l]; + } +} + +template +static __device__ __forceinline__ void mmq_tile_y_load_global_to_regs( + const int * __restrict__ by, int (&cache)[nchunks]) { +#pragma unroll + for (int c = 0; c < nchunks; ++c) { + const int l = c*(nwarps*warp_size) + threadIdx.y*warp_size + threadIdx.x; + cache[c] = by[l]; + } +} + +template +static __device__ __forceinline__ void mmq_tile_y_store_regs( + int * __restrict__ tile_y, const int (&cache)[nchunks]) { +#pragma unroll + for (int c = 0; c < nchunks; ++c) { + const int l = c*(nwarps*warp_size) + threadIdx.y*warp_size + threadIdx.x; + tile_y[l] = cache[c]; + } +} +#endif // RDNA3_5 + template static __device__ __forceinline__ void mul_mat_q_process_tile( const char * __restrict__ x, const int offset_x, const int * __restrict__ y, @@ -4123,66 +4019,170 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( float sum[mmq_x*mmq_y / (nwarps*warp_size)] = {0.0f}; -#if defined(GGML_USE_HIP) - mmq_profile_bind(mmq_profile); -#endif - constexpr int sz = sizeof(block_q8_1_mmq) / sizeof(int); -#if defined(GGML_USE_HIP) - if (g_mmq_profile_outer) { -#else - if (mmq_profile) { -#endif - for (int kb0 = kb0_start; kb0 < kb0_stop; kb0 += blocks_per_iter) { - mmq_profile_kb_iter(mmq_profile); +#if defined(RDNA3_5) + constexpr int tile_y_elems = mmq_x*MMQ_TILE_Y_K; + constexpr int tile_y_load_stride = nwarps*warp_size; + if constexpr (mmq_x <= 64 && tile_y_elems % tile_y_load_stride == 0) { + constexpr int tile_y_nchunks = tile_y_elems/tile_y_load_stride; - { - uint64_t t0 = 0; - __syncthreads(); - mmq_profile_phase_begin(mmq_profile, t0); - load_tiles(x, tile_x, offset_x + kb0, tile_x_max_i, stride_row_x); - __syncthreads(); - mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_LOAD_TILES); - } + int y0_next_cache[tile_y_nchunks]; + bool have_y0_prefetch = false; - const int yk = kb0 * qk / ne_block; - const int * by0 = y + ncols_y * yk * sz; - const int * by1 = y + ncols_y * (yk + 1) * sz; + if (mmq_profile) { + for (int kb0 = kb0_start; kb0 < kb0_stop; kb0 += blocks_per_iter) { + mmq_profile_kb_iter(mmq_profile); - { - uint64_t t0 = 0; - __syncthreads(); - mmq_profile_phase_begin(mmq_profile, t0); -#pragma unroll - for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { - int l = l0 + threadIdx.y*warp_size + threadIdx.x; + { + uint64_t t0 = 0; + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); + load_tiles(x, tile_x, offset_x + kb0, tile_x_max_i, stride_row_x); + __syncthreads(); + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_LOAD_TILES); + } + + const int yk = kb0 * qk / ne_block; + const int * by0 = y + ncols_y * yk * sz; + const int * by1 = y + ncols_y * (yk + 1) * sz; + + { + uint64_t t0 = 0; + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); + + if (have_y0_prefetch) { + mmq_tile_y_store_regs(tile_y, y0_next_cache); + have_y0_prefetch = false; + } else { + mmq_tile_y_load_global(tile_y, by0); + } + + __syncthreads(); + + int y1_cache[tile_y_nchunks]; + mmq_tile_y_load_global_to_regs(by1, y1_cache); + + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_Y_ACT); + + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); + + vec_dot(tile_x, tile_y, sum, 0); + + __syncthreads(); + + mmq_tile_y_store_regs(tile_y, y1_cache); - tile_y[l] = by0[l]; + __syncthreads(); + + const int kb0_next = kb0 + blocks_per_iter; + if (kb0_next < kb0_stop) { + const int * by0_next = y + ncols_y * (kb0_next * qk / ne_block) * sz; + mmq_tile_y_load_global_to_regs(by0_next, y0_next_cache); + have_y0_prefetch = true; + } + + vec_dot(tile_x, tile_y, sum, MMQ_TILE_NE_K); + + __syncthreads(); + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_VEC_DOT); } - __syncthreads(); - mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_Y_ACT); } + } else { + 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); + + const int yk = kb0 * qk / ne_block; + const int * by0 = y + ncols_y * yk * sz; + const int * by1 = y + ncols_y * (yk + 1) * sz; + + if (have_y0_prefetch) { + mmq_tile_y_store_regs(tile_y, y0_next_cache); + have_y0_prefetch = false; + } else { + mmq_tile_y_load_global(tile_y, by0); + } - { - uint64_t t0 = 0; __syncthreads(); - mmq_profile_phase_begin(mmq_profile, t0); + + int y1_cache[tile_y_nchunks]; + mmq_tile_y_load_global_to_regs(by1, y1_cache); + vec_dot(tile_x, tile_y, sum, 0); + __syncthreads(); -#pragma unroll - for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { - int l = l0 + threadIdx.y*warp_size + threadIdx.x; - tile_y[l] = by1[l]; - } + mmq_tile_y_store_regs(tile_y, y1_cache); + __syncthreads(); + + const int kb0_next = kb0 + blocks_per_iter; + if (kb0_next < kb0_stop) { + const int * by0_next = y + ncols_y * (kb0_next * qk / ne_block) * sz; + mmq_tile_y_load_global_to_regs(by0_next, y0_next_cache); + have_y0_prefetch = true; + } + vec_dot(tile_x, tile_y, sum, MMQ_TILE_NE_K); + __syncthreads(); - mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_VEC_DOT); } } - } else { + } else +#endif // RDNA3_5 + { + if (mmq_profile) { + for (int kb0 = kb0_start; kb0 < kb0_stop; kb0 += blocks_per_iter) { + mmq_profile_kb_iter(mmq_profile); + + { + uint64_t t0 = 0; + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); + load_tiles(x, tile_x, offset_x + kb0, tile_x_max_i, stride_row_x); + __syncthreads(); + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_LOAD_TILES); + } + + const int yk = kb0 * qk / ne_block; + const int * by0 = y + ncols_y * yk * sz; + const int * by1 = y + ncols_y * (yk + 1) * sz; + + { + uint64_t t0 = 0; + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); +#pragma unroll + for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { + int l = l0 + threadIdx.y*warp_size + threadIdx.x; + + tile_y[l] = by0[l]; + } + __syncthreads(); + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_Y_ACT); + } + + { + uint64_t t0 = 0; + __syncthreads(); + mmq_profile_phase_begin(mmq_profile, t0); + vec_dot(tile_x, tile_y, sum, 0); + __syncthreads(); +#pragma unroll + for (int l0 = 0; l0 < mmq_x * MMQ_TILE_Y_K; l0 += nwarps * warp_size) { + int l = l0 + threadIdx.y*warp_size + threadIdx.x; + + tile_y[l] = by1[l]; + } + __syncthreads(); + vec_dot(tile_x, tile_y, sum, MMQ_TILE_NE_K); + __syncthreads(); + mmq_profile_phase_end(mmq_profile, t0, MMQ_PROFILE_VEC_DOT); + } + } + } else { 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); @@ -4221,19 +4221,17 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( __syncthreads(); } } + } if (fixup) { write_back(sum, ids_dst, tmp_fixup + blockIdx.x*(mmq_x*mmq_y), mmq_y, mmq_y, mmq_x); } else { write_back(sum, ids_dst, dst, stride_col_dst, tile_x_max_i, tile_y_max_j); } - -#if defined(GGML_USE_HIP) - mmq_profile_bind(nullptr); -#endif } + // The mul_mat_q kernel implements "stream-k" work partitioning as described in https://arxiv.org/abs/2301.03598 template @@ -4664,16 +4662,11 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a const dim3 block_dims(warp_size, nwarps, 1); - mmq_profile_guard prof_guard(ctx.pool(id), mmq_profile_any_enabled(), type, mmq_x, + mmq_profile_guard prof_guard(ctx.pool(id), mmq_profile_phases_enabled(), type, mmq_x, args.nrows_x, args.ncols_max, stream); const int nbytes_shared = mmq_get_nbytes_shared(mmq_x, mmq_y, cc, warp_size, nwarps); -#if defined(GGML_USE_HIP) - const int outer_phases = mmq_profile_phases_enabled() ? 1 : 0; - CUDA_CHECK(cudaMemcpyToSymbol(g_mmq_profile_outer, &outer_phases, sizeof(int))); -#endif - CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); CUDA_SET_SHARED_MEMORY_LIMIT((mul_mat_q), nbytes_shared); From 272c12ec04324bb9b94be097e3c05840c551ca8a Mon Sep 17 00:00:00 2001 From: lichang Date: Tue, 30 Jun 2026 00:34:13 -0600 Subject: [PATCH 20/22] cuda: pipeline Q6_K B-tile loads in RDNA3.5 WMMA vec_dot Prefetch the next activation B tile while MMA runs on the current tile in the Q6_K narrow-N path, improving FFN-down kernel throughput without changing launch geometry. Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 35 +++++++++++++++++++++++++++-------- 1 file changed, 27 insertions(+), 8 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 7d59cf940fb0..9fa918b7b4a5 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -2998,24 +2998,43 @@ static __device__ __forceinline__ void vec_dot_q6_K_q8_1_mma( } } + constexpr int j_step = ntx*tile_C::J; + + tile_B B0; + load_ldmatrix(B0, y_qs + 0, MMQ_TILE_Y_K); + float dB0 = y_df[tile_C::get_j(0)*MMQ_TILE_Y_K + k01/QI8_1]; + #pragma unroll - for (int j0 = 0; j0 < mmq_x; j0 += ntx*tile_C::J) { - tile_B B; - load_ldmatrix(B, y_qs + j0*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); + for (int j0 = 0; j0 < mmq_x; j0 += j_step) { + const int j0_next = j0 + j_step; - const int j = j0 + tile_C::get_j(0); - const float dB = y_df[j*MMQ_TILE_Y_K + k01/QI8_1]; + tile_B B1; + float dB1 = 0.0f; + if (j0_next < mmq_x) { + load_ldmatrix(B1, y_qs + j0_next*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); + const int jn = j0_next + tile_C::get_j(0); + dB1 = y_df[jn*MMQ_TILE_Y_K + k01/QI8_1]; + } + + tile_C C[ntx]; #pragma unroll for (int n = 0; n < ntx; ++n) { - tile_C C; - mma(C, A[n], B); + mma(C[n], A[n], B0); + } +#pragma unroll + for (int n = 0; n < ntx; ++n) { #pragma unroll for (int l = 0; l < tile_C::ne; ++l) { - sum[(j0/tile_C::J + n)*tile_C::ne + l] += C.x[l]*sclA[n][l]*dB; + sum[(j0/tile_C::J + n)*tile_C::ne + l] += C[n].x[l]*sclA[n][l]*dB0; } } + + if (j0_next < mmq_x) { + B0 = B1; + dB0 = dB1; + } } } #elif defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) From 46484861e8510f1106431a9954ad20e7c05aa62d Mon Sep 17 00:00:00 2001 From: lichang Date: Tue, 30 Jun 2026 00:48:38 -0600 Subject: [PATCH 21/22] cuda: use dual-WG mmq_x=64 for Q6_K narrow-N on RDNA3.5 Drop the Q6-only mmq_x=128 upgrade so Q6_K FFN uses the same smpbo/2 mmq_x=64 ntx=2 launch as Q4_K, which measures faster on gfx115x prefill. Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 32 ++------------------------------ 1 file changed, 2 insertions(+), 30 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 9fa918b7b4a5..56e230fd7f21 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -4902,36 +4902,8 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda } } - // Narrow N (batch ≤128): Q6_K only — prefer mmq_x=128 for ntx=1 when LDS allows. - // Q6_K at mmq_x=128 uses ~38 KiB (1 WG/CU) but beats mmq_x=64 with ntx=2; do not - // require smpbo/2 here — that policy is for Q4_K dual-WG, not Q6 narrow-N. - if (GGML_CUDA_CC_IS_RDNA3_5(cc) && type == GGML_TYPE_Q6_K && args.ncols_max <= mmq_x_max) { - const int ntx_cur = (args.ncols_max + mmq_x_best - 1) / mmq_x_best; - if (ntx_cur > 1) { - const int nty = (args.nrows_x + mmq_y - 1) / mmq_y; - const int ntzw = static_cast(args.nchannels_y * args.nsamples_y); - - for (int mmq_x = mmq_x_max; mmq_x >= 8; mmq_x -= 8) { - const int granularity = mmq_get_granularity_host(mmq_x, cc); - if (mmq_x % granularity != 0) { - continue; - } - const size_t nbytes = mmq_get_nbytes_shared(mmq_x, mmq_y, cc, warp_size, nwarps); - if (nbytes > smpbo) { - continue; - } - const int ntx = (args.ncols_max + mmq_x - 1) / mmq_x; - if (ntx != 1) { - continue; - } - const int grid_blocks = nty * ntx * ntzw; - if (grid_blocks >= 8) { - mmq_x_best = mmq_x; - break; - } - } - } - } + // Q6_K narrow-N (batch ≤128): use the RDNA3.5 dual-WG path (mmq_x=64, ntx=2) from + // smpbo/2 selection above — faster end-to-end than mmq_x=128 ntx=1 on gfx115x prefill. #endif // GGML_USE_HIP #if defined(GGML_USE_HIP) From 35bd178185405df34602f5cce5c67c466ad04bed Mon Sep 17 00:00:00 2001 From: lichang Date: Tue, 30 Jun 2026 15:03:39 -0600 Subject: [PATCH 22/22] cuda: limit RDNA3.5 dual-WG mmq_x downgrade to K-quants Block quants (Q5_0, Q8_0) regressed when forced to mmq_x=64 for LDS dual-WG occupancy; keep mmq_x=128 for them while retaining mmq_x=64 for K-quants tuned on gfx115x (Q6_K FFN down). Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 26 +++++++++++++++++++++----- 1 file changed, 21 insertions(+), 5 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 56e230fd7f21..0b42888daefe 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -4655,6 +4655,23 @@ struct mmq_args { bool use_stream_k; int64_t ncols_max; }; +#if defined(GGML_USE_HIP) +// RDNA3.5 dual-WG (mmq_x=64, nbytes <= smpbo/2) helps K-quants with large per-tile LDS +// (Q6_K WMMA tuning). Block quants (Q5_0, Q8_0, Q4_0) are faster at mmq_x=128 ntx=1. +static bool mmq_rdna35_dual_wg_eligible(const ggml_type type) { + switch (type) { + case GGML_TYPE_Q2_K: + case GGML_TYPE_Q3_K: + case GGML_TYPE_Q4_K: + case GGML_TYPE_Q5_K: + case GGML_TYPE_Q6_K: + return true; + default: + return false; + } +} +#endif // GGML_USE_HIP + template static size_t mmq_get_nbytes_shared(const int mmq_x, const int mmq_y, const int cc, const int warp_size, const int nwarps) { const tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(type, mmq_y); @@ -4836,9 +4853,9 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda } #if defined(GGML_USE_HIP) - // RDNA3.5 (gfx115x): mmq_x=128 uses ~59% of LDS/CU → 1 WG/CU. Pick the smallest - // mmq_x with nbytes_shared <= smpbo/2 (2 WGs/CU) without doubling M-tile count. - if (GGML_CUDA_CC_IS_RDNA3_5(cc) && mmq_x_best > 0) { + // RDNA3.5 (gfx115x): mmq_x=128 uses ~59% of LDS/CU → 1 WG/CU. For K-quants only, + // pick the smallest mmq_x with nbytes_shared <= smpbo/2 (2 WGs/CU). + if (GGML_CUDA_CC_IS_RDNA3_5(cc) && mmq_x_best > 0 && mmq_rdna35_dual_wg_eligible(type)) { const size_t lds_dual_wg = smpbo / 2; const size_t nbytes_best = mmq_get_nbytes_shared(mmq_x_best, mmq_y, cc, warp_size, nwarps); if (nbytes_best > lds_dual_wg) { @@ -4902,8 +4919,7 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda } } - // Q6_K narrow-N (batch ≤128): use the RDNA3.5 dual-WG path (mmq_x=64, ntx=2) from - // smpbo/2 selection above — faster end-to-end than mmq_x=128 ntx=1 on gfx115x prefill. + // Q6_K / K-quants narrow-N: dual-WG path (mmq_x=64, ntx=2) from smpbo/2 selection. #endif // GGML_USE_HIP #if defined(GGML_USE_HIP)