From 9150e75939164574b34f97493a253b0f0f4c668d Mon Sep 17 00:00:00 2001 From: Cary Palmer <24235924+professorpalmer@users.noreply.github.com> Date: Sat, 19 Sep 2026 15:05:56 -0500 Subject: [PATCH 1/3] cuda: PTQ1_0 decode on consumer Ampere/Ada: SoA q8 activations, warp-per-row small-K GEMV (+14% TG) Two changes to the batch-1 PTQ1_0 mat-vec path, found by tracing a Bonsai 2 27B decode graph on an RTX 4070 (sm_89) with CUPTI. Same weights, same bits, output byte-identical. 1. Exact integer sums + warp-transposed (SoA) q8_1 activation layout. quantize_q8_1 stores the int16 sum of each 32-value block in ds.y (instead of the float input sum) and writes the activations grouped by 32 K-blocks with word w of the group contiguous. The ternary vec-dot then accumulates raw trits {0,1,2} with dp4a and corrects once per block, removing a __vsub4 per 4 weights (no native SIMD byte subtract on sm_8x), and a warp-wide load of "word w" in the small-K geometry hits 1 L1 wavefront instead of ~36 (lanes were reading block_q8_1 structs 144 B apart). That LSU pressure, not GDDR, capped every PTQ1_0 GEMV at ~370 GB/s. Bytes per column are unchanged when K is padded to 32*128. Only PTQ1_0 opts in (ggml_cuda_q8_1_exact_isum); other types keep the AoS layout. 2. Warp-per-row geometry for small-K PTQ1_0 GEMVs. The generic small_k loop strides K-blocks across the whole 128-thread block, so for K=5120 (40 blocks) only threads 0..39 ever load anything and two of four warps just hold SM slots at __syncthreads. Each warp now owns one row and its 32 lanes stride that row's K-blocks; shuffle reduction, no shared memory, no barrier. qkv 391 -> 445 GB/s, lm_head 414 -> 460. The recurrent-state gather fusion that was part of the first revision of this PR is now its own PR so that this activation layout can be compared with #218 on its own. RTX 4070 12 GB, Ternary-Bonsai-2-27B-PTQ1_0, llama-bench -fa 1 -ctk/-ctv q4_0, stock clocks, three alternating rounds of r=3: prism @ 9a9394a 51.0 t/s tg128 + 1 + 2 (this PR) 56.4 (+10.5%) + GDN gather fusion (follow-up) 58.5 (+14.7%) test-backend-ops CUDA0 vs CPU: MUL_MAT ptq1_0 45/45 supported cases, MUL_MAT_ID ptq1_0 75/75. 400 greedy tokens identical with the change on and off. --- ggml/src/ggml-cuda/common.cuh | 37 +++++++++ ggml/src/ggml-cuda/mmvq.cu | 132 +++++++++++++++++++++++++++++++-- ggml/src/ggml-cuda/quantize.cu | 30 +++++++- ggml/src/ggml-cuda/vecdotq.cuh | 49 ++++++++---- 4 files changed, 226 insertions(+), 22 deletions(-) diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index ef929d3d7842..b438ad75547d 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -1026,6 +1026,43 @@ struct ggml_cuda_type_traits { static constexpr int qi = QI_PTQ1_0; }; +// For these src0 types the row-quantizer stores the exact integer sum of the q8 values (int16 +// bits) in block_q8_1::ds.y instead of the float input sum, so the ternary vec-dot can accumulate +// raw digits {0,1,2} with dp4a and subtract the bias once per 32-block instead of per 4 weights. +// The quantizer also writes the warp-transposed layout below for these types. +static constexpr __host__ __device__ bool ggml_cuda_q8_1_exact_isum(ggml_type type_src0) { +#if defined(GGML_USE_HIP) + GGML_UNUSED(type_src0); + return false; +#else + return type_src0 == GGML_TYPE_PTQ1_0; +#endif +} + +// Warp-transposed (SoA) q8_1 activation layout for the ternary MMVQ. +// +// One PTQ1_0 K-block (128 weights) consumes 4 block_q8_1 = 36 words (32 qs + 4 ds). In the +// small-K MMVQ geometry each lane owns one K-block, so with the plain AoS layout a warp-wide +// load of "word w" touches 32 lines 144 B apart: ~36 L1 wavefronts per instruction, ~1300 per +// K-iteration against ~250 for the weights themselves. That LSU traffic, not GDDR, capped the +// PTQ1 GEMV near 370 GB/s on Ada. Here K-blocks are grouped by 32 and word w of the group is +// stored contiguously, so the same load is 32 consecutive words = 1 wavefront. +// Bytes per column are unchanged when K is padded to a multiple of 32*128 = 4096. +#define GGML_CUDA_PTQ1_Q8_GROUP_KB 32 +#define GGML_CUDA_PTQ1_Q8_WORDS_PER_KB 36 +#define GGML_CUDA_PTQ1_Q8_GROUP_WORDS (GGML_CUDA_PTQ1_Q8_GROUP_KB * GGML_CUDA_PTQ1_Q8_WORDS_PER_KB) +#define GGML_CUDA_PTQ1_K_PAD (GGML_CUDA_PTQ1_Q8_GROUP_KB * QK_PTQ1_0) + +// Word offset (within one activation column) of word w (0..7 = qs words, 8 = ds) of block_q8_1 ib. +static constexpr __host__ __device__ int ggml_cuda_ptq1_q8_word(int ib, int w) { + const int kb = ib >> 2; + const int sub = ib & 3; + const int g = kb >> 5; + const int lane = kb & 31; + const int ww = w < 8 ? sub * 8 + w : 32 + sub; + return g * GGML_CUDA_PTQ1_Q8_GROUP_WORDS + ww * GGML_CUDA_PTQ1_Q8_GROUP_KB + lane; +} + template<> struct ggml_cuda_type_traits { static constexpr int qk = QK4_0; diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index bf51b61e17b1..5ec5e11a5feb 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -667,6 +667,89 @@ static __global__ void mul_mat_vec_q( const block_q8_1 * y = ((const block_q8_1 *) vy) + sample_y*stride_sample_y + channel_y*stride_channel_y; const int kbx_offset = sample_x*stride_sample_x + channel_x*stride_channel_x + row0*stride_row_x; +#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + if constexpr (small_k && type == GGML_TYPE_PTQ1_0 && nwarps > 1 && rows_per_cuda_block == nwarps) { + // Warp-per-row geometry for small K. + // The generic small_k loop strides K-blocks across the whole block (kbx = tid; kbx += 128), so + // for K=5120 (40 blocks) only threads 0..39 ever load anything: warp 0 is full, warp 1 has + // 8 lanes, warps 2..3 only wait at __syncthreads while holding SM warp slots. Measured + // in-graph with CUPTI on an RTX 4070: 350-390 GB/s for every K=5120 GEMV vs 460 GB/s for the + // K=17408 geometry where all four warps stream. Here warp w owns row row0+w and its 32 lanes + // stride that row's K-blocks: every warp issues loads, the reduction is a shuffle, and + // there is no shared memory or block barrier. Activations are re-read per warp but they + // are 20 KB and L1/L2 resident. + const int warp = threadIdx.y; + const int lane = threadIdx.x; + const bool row_ok = uint32_t(row0 + warp) < stride_col_dst; + + float acc[ncols_dst] = { 0.0f }; + float acc_gate[ncols_dst] = { 0.0f }; + if (row_ok) { + const int kbx_row = kbx_offset + warp * stride_row_x; + for (int kbx = lane; kbx < blocks_per_row_x; kbx += warp_size) { + float dots[ncols_dst]; + vec_dot_ptq1_0_q8_1_multi(vx, y, kbx_row + kbx, kbx, stride_col_y, dots); +#pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + acc[j] += dots[j]; + } + if constexpr (has_fusion && has_gate) { + vec_dot_ptq1_0_q8_1_multi(vgate, y, kbx_row + kbx, kbx, stride_col_y, dots); +#pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + acc_gate[j] += dots[j]; + } + } + } + } +#pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + acc[j] = warp_reduce_sum(acc[j]); + if constexpr (has_fusion && has_gate) { + acc_gate[j] = warp_reduce_sum(acc_gate[j]); + } + } + if (lane == 0 && row_ok) { + float * dst_row = dst + sample_dst*stride_sample_dst + channel_dst*stride_channel_dst + row0 + warp; + [[maybe_unused]] const uint32_t channel_bias = ids ? channel_x : channel_dst; + [[maybe_unused]] const int64_t bias_off = sample_dst*stride_sample_dst + channel_bias*stride_channel_dst + row0 + warp; +#pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + float result = acc[j]; + if constexpr (has_fusion) { + if (use_bias) { + result += ((const float *) fusion.x_bias)[bias_off + j*stride_col_dst]; + } + if constexpr (has_gate) { + float gate_value = acc_gate[j]; + if (use_gate_bias) { + gate_value += ((const float *) fusion.gate_bias)[bias_off + j*stride_col_dst]; + } + switch (active_glu) { + case GGML_GLU_OP_SWIGLU: + result *= ggml_cuda_op_silu_single(gate_value); + break; + case GGML_GLU_OP_GEGLU: + result *= ggml_cuda_op_gelu_single(gate_value); + break; + case GGML_GLU_OP_SWIGLU_OAI: + result = ggml_cuda_op_swiglu_oai_single(gate_value, result); + break; + default: + result = result * gate_value; + break; + } + } + } + dst_row[j*stride_col_dst] = result; + } + } + GGML_UNUSED_VARS(use_gate, use_scale, use_gate_scale, gate_bias, x_bias, x_scale, gate_scale, x_scales, + gate_scales, x_biases, gate_biases, tmp, tmp_gate, vec_dot_q_cuda, blocks_per_iter); + return; + } +#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + if constexpr ((type == GGML_TYPE_Q1_0 || type == GGML_TYPE_Q2_0 || type == GGML_TYPE_PQ2_0) && table_id == MMVQ_PARAMETERS_GB10) { using block_t = std::conditional_t 1 && ncols_dst <= 3) { + // PTQ1_0 activations use the warp-transposed q8 layout (ggml_cuda_ptq1_q8_word), so the + // vec-dot is handed the column base and the K-block index instead of &y[kby]. + if constexpr (type == GGML_TYPE_PTQ1_0) { + GGML_UNUSED(kby); + GGML_UNUSED(kqs); # pragma unroll for (int i = 0; i < rows_per_cuda_block; ++i) { float dots[ncols_dst]; - vec_dot_ptq1_0_q8_1_multi(vx, &y[kby], kbx_offset + i * stride_row_x + kbx, kqs, + vec_dot_ptq1_0_q8_1_multi(vx, y, kbx_offset + i * stride_row_x + kbx, kbx, stride_col_y, dots); # pragma unroll for (int j = 0; j < ncols_dst; ++j) { @@ -729,7 +830,7 @@ static __global__ void mul_mat_vec_q( if constexpr (has_fusion) { if constexpr (has_gate) { - vec_dot_ptq1_0_q8_1_multi(vgate, &y[kby], kbx_offset + i * stride_row_x + kbx, kqs, + vec_dot_ptq1_0_q8_1_multi(vgate, y, kbx_offset + i * stride_row_x + kbx, kbx, stride_col_y, dots); # pragma unroll for (int j = 0; j < ncols_dst; ++j) { @@ -897,9 +998,24 @@ static __global__ void mul_mat_vec_q_moe( const int kby = kbx * (qk/QK8_1); const int kqs = vdr * (threadIdx.x % (qi/vdr)); +#if !defined(GGML_USE_HIP) + if constexpr (type == GGML_TYPE_PTQ1_0) { + // Warp-transposed q8 layout: hand the vec-dot the column base and K-block index. + GGML_UNUSED(kby); + GGML_UNUSED(kqs); #pragma unroll - for (int i = 0; i < c_rows_per_block; ++i) { - tmp[i] += vec_dot_q_cuda(vx, &y[kby], kbx_offset + i*stride_row_x + kbx, kqs); + for (int i = 0; i < c_rows_per_block; ++i) { + float dot; + vec_dot_ptq1_0_q8_1_multi<1>(vx, y, kbx_offset + i*stride_row_x + kbx, kbx, 0, &dot); + tmp[i] += dot; + } + } else +#endif + { +#pragma unroll + for (int i = 0; i < c_rows_per_block; ++i) { + tmp[i] += vec_dot_q_cuda(vx, &y[kby], kbx_offset + i*stride_row_x + kbx, kqs); + } } } @@ -1433,7 +1549,11 @@ void ggml_cuda_mul_mat_vec_q( } } - const int64_t ne10_padded = GGML_PAD(ne10, MATRIX_ROW_PADDING); + int64_t ne10_padded = GGML_PAD(ne10, MATRIX_ROW_PADDING); + if (ggml_cuda_q8_1_exact_isum(src0->type)) { + // Warp-transposed q8 layout needs whole 32-K-block groups per column. + ne10_padded = GGML_PAD(ne10_padded, GGML_CUDA_PTQ1_K_PAD); + } ggml_cuda_pool_alloc src1_q8_1(ctx.pool(), ne13*ne12 * ne11*ne10_padded * sizeof(block_q8_1)/QK8_1); { const int64_t s11 = src1->nb[1] / ts_src1; diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu index fbbc6314ab33..57db61122bd3 100644 --- a/ggml/src/ggml-cuda/quantize.cu +++ b/ggml/src/ggml-cuda/quantize.cu @@ -51,6 +51,11 @@ static __device__ __forceinline__ float nvfp4_native_scale_error( #endif // CUDART_VERSION >= 12080 #endif // defined(BLACKWELL_MMA_AVAILABLE) +// exact_isum: store the integer sum of the quantized values (bit-cast into the ds.y half slot) +// instead of the float sum of the inputs. Ternary vec-dots (PTQ1_0) use it to fold the +// digit bias {0,1,2} -> {-1,0,+1} into one subtraction per 32-block instead of a SIMD byte +// subtract per 4 weights, with results bit-identical to the biased path. +template __launch_bounds__(CUDA_QUANTIZE_BLOCK_SIZE, 1) static __global__ void quantize_q8_1( const float * x_ptr, void * vy_ptr, @@ -92,6 +97,24 @@ static __global__ void quantize_q8_1( const float d = amax / 127.0f; const int8_t q = amax == 0.0f ? 0 : roundf(xi / d); + if constexpr (exact_isum) { + // Warp-transposed layout (see ggml_cuda_ptq1_q8_word); ne0 is padded to GGML_CUDA_PTQ1_K_PAD. + int isum = q; + isum = warp_reduce_sum(isum); + int32_t * yw = (int32_t *) vy; + const int64_t col = i_cont / ne0; + const int64_t col_base = col * (ne0 / QK8_1) * (int64_t) (sizeof(block_q8_1) / 4); + const int ib_col = (int) (i0 / QK8_1); + ((int8_t *) (yw + col_base + ggml_cuda_ptq1_q8_word(ib_col, iqs / 4)))[iqs % 4] = q; + if (iqs > 0) { + return; + } + // |isum| <= 32*127 fits int16; keep the raw bits in the half slot. + const half2 ds = make_half2(__float2half(d), __short_as_half((short) isum)); + yw[col_base + ggml_cuda_ptq1_q8_word(ib_col, 8)] = *reinterpret_cast(&ds); + return; + } + y[ib].qs[iqs] = q; if (iqs > 0) { @@ -648,8 +671,11 @@ void quantize_row_q8_1_cuda( const dim3 num_blocks(block_num_x, ne1, ne2*ne3); const dim3 block_size(CUDA_QUANTIZE_BLOCK_SIZE, 1, 1); const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(num_blocks, block_size, 0, stream); - ggml_cuda_kernel_launch(quantize_q8_1, launch_params, x, vy, ne00, s01, s02, s03, ne0, ne1, ne2_fastdiv); - GGML_UNUSED(type_src0); + if (ggml_cuda_q8_1_exact_isum(type_src0)) { + ggml_cuda_kernel_launch(quantize_q8_1, launch_params, x, vy, ne00, s01, s02, s03, ne0, ne1, ne2_fastdiv); + } else { + ggml_cuda_kernel_launch(quantize_q8_1, launch_params, x, vy, ne00, s01, s02, s03, ne0, ne1, ne2_fastdiv); + } } void quantize_mmq_q8_1_cuda( diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index 65dbbb9573be..23b0c61c34c2 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -805,16 +805,26 @@ static __device__ __forceinline__ float vec_dot_q2_0_q8_1( } #if !defined(GGML_USE_HIP) +// y_col0: activation column 0 base (warp-transposed layout, see ggml_cuda_ptq1_q8_word); +// kbx_x: absolute PTQ1 block index into vbq; kb: K-block index within the row; +// stride_col_y: column stride in block_q8_1 units. template static __device__ __forceinline__ void vec_dot_ptq1_0_q8_1_multi(const void * __restrict__ vbq, - const block_q8_1 * __restrict__ bq8_1, - const int & kbx, - const int & iqs, + const block_q8_1 * __restrict__ y_col0, + const int & kbx_x, + const int & kb, const uint32_t stride_col_y, float * result) { - const block_ptq1_0 * bq = (const block_ptq1_0 *) vbq + kbx; + const block_ptq1_0 * bq = (const block_ptq1_0 *) vbq + kbx_x; int sumi[ncols_dst][4] = {}; + // Lane-coalesced activation words: word ww of this K-block is at yw[ww*32]; 32 lanes of a + // warp own 32 consecutive K-blocks, so every load below is 32 consecutive words. + const int * __restrict__ yw = (const int *) y_col0 + (kb >> 5) * GGML_CUDA_PTQ1_Q8_GROUP_WORDS + (kb & 31); + const uint32_t sy = stride_col_y * (sizeof(block_q8_1) / 4); +# define PTQ1_U(j, sub, m) yw[(j) * sy + ((sub) * 8 + (m)) * GGML_CUDA_PTQ1_Q8_GROUP_KB] +# define PTQ1_DS(j, sub) yw[(j) * sy + (32 + (sub)) * GGML_CUDA_PTQ1_Q8_GROUP_KB] + // Widen four bytes to 16-bit lanes so multiply-by-three cannot carry between bytes. # pragma unroll for (int g = 0; g < 4; ++g) { @@ -829,11 +839,12 @@ static __device__ __forceinline__ void vec_dot_ptq1_0_q8_1_multi(const void * __ v_lo = w_lo & 0x00FF00FF; v_hi = w_hi & 0x00FF00FF; - const int q = __vsub4(__byte_perm(w_lo, w_hi, 0x7531), 0x01010101); + // Raw digits {0,1,2}; the -1 bias is folded into the per-32-block exact q8 sum below. + const int q = __byte_perm(w_lo, w_hi, 0x7531); const int e = t * 16 + 4 * g; # pragma unroll for (int j = 0; j < ncols_dst; ++j) { - const int u = get_int_b4(bq8_1[j * stride_col_y + iqs + (e >> 5)].qs, (e & 31) >> 2); + const int u = PTQ1_U(j, e >> 5, (e & 31) >> 2); sumi[j][e >> 5] = ggml_cuda_dp4a(q, u, sumi[j][e >> 5]); } } @@ -852,11 +863,11 @@ static __device__ __forceinline__ void vec_dot_ptq1_0_q8_1_multi(const void * __ v_lo = w_lo & 0x00FF00FF; v_hi = w_hi & 0x00FF00FF; - const int q = __vsub4(__byte_perm(w_lo, w_hi, 0x7531), 0x01010101); + const int q = __byte_perm(w_lo, w_hi, 0x7531); const int e = 80 + t * 8 + 4 * g; # pragma unroll for (int j = 0; j < ncols_dst; ++j) { - const int u = get_int_b4(bq8_1[j * stride_col_y + iqs + (e >> 5)].qs, (e & 31) >> 2); + const int u = PTQ1_U(j, e >> 5, (e & 31) >> 2); sumi[j][e >> 5] = ggml_cuda_dp4a(q, u, sumi[j][e >> 5]); } } @@ -870,10 +881,10 @@ static __device__ __forceinline__ void vec_dot_ptq1_0_q8_1_multi(const void * __ const uint32_t w1 = v * 3; v = w1 & 0x00FF00FF; - const int q = __vsub4(__byte_perm(w0, w1, 0x7531), 0x01010101); + const int q = __byte_perm(w0, w1, 0x7531); # pragma unroll for (int j = 0; j < ncols_dst; ++j) { - const int u = get_int_b4(bq8_1[j * stride_col_y + iqs + 3].qs, 6 + t / 2); + const int u = PTQ1_U(j, 3, 6 + t / 2); sumi[j][3] = ggml_cuda_dp4a(q, u, sumi[j][3]); } } @@ -884,10 +895,16 @@ static __device__ __forceinline__ void vec_dot_ptq1_0_q8_1_multi(const void * __ float acc = 0.0f; # pragma unroll for (int k = 0; k < 4; ++k) { - acc += __low2float(bq8_1[j * stride_col_y + iqs + k].ds) * (float) sumi[j][k]; + const int ds_bits = PTQ1_DS(j, k); + const half2 ds = *reinterpret_cast(&ds_bits); + // ds.y carries the exact int16 sum of the 32 q8 values (see quantize_q8_1). + const int isum = (int) __half_as_short(__high2half(ds)); + acc += __low2float(ds) * (float) (sumi[j][k] - isum); } result[j] = d * acc; } +# undef PTQ1_U +# undef PTQ1_DS } #endif @@ -969,9 +986,13 @@ static __device__ __forceinline__ float vec_dot_ptq1_0_q8_1(const void * __restr } return (float) bq->d * acc; #else - float result; - vec_dot_ptq1_0_q8_1_multi<1>(vbq, bq8_1, kbx, iqs, 0, &result); - return result; + // Unreachable on CUDA: mul_mat_vec_q routes every PTQ1_0 ncols_dst through + // vec_dot_ptq1_0_q8_1_multi, which needs the column base for the warp-transposed q8 layout. + GGML_UNUSED(vbq); + GGML_UNUSED(bq8_1); + GGML_UNUSED(kbx); + GGML_UNUSED(iqs); + return 0.0f; #endif } From 68c1e5f88b6a278401733a5da664db301ee6e763 Mon Sep 17 00:00:00 2001 From: Cary Palmer <24235924+professorpalmer@users.noreply.github.com> Date: Sun, 20 Sep 2026 20:09:14 -0500 Subject: [PATCH 2/3] cuda: keep the PTQ1_0 SoA/exact-isum path off on MUSA, matching HIP. MUSA took the non-HIP branch and would have used the new quantizer and vec-dot. Also trim the hard-wrapped benchmark comments. --- ggml/src/ggml-cuda/common.cuh | 17 +++-------------- ggml/src/ggml-cuda/mmvq.cu | 19 +++++-------------- ggml/src/ggml-cuda/quantize.cu | 5 +---- ggml/src/ggml-cuda/vecdotq.cuh | 8 +++----- 4 files changed, 12 insertions(+), 37 deletions(-) diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index b438ad75547d..c8e027bb57c6 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -1026,12 +1026,9 @@ struct ggml_cuda_type_traits { static constexpr int qi = QI_PTQ1_0; }; -// For these src0 types the row-quantizer stores the exact integer sum of the q8 values (int16 -// bits) in block_q8_1::ds.y instead of the float input sum, so the ternary vec-dot can accumulate -// raw digits {0,1,2} with dp4a and subtract the bias once per 32-block instead of per 4 weights. -// The quantizer also writes the warp-transposed layout below for these types. +// PTQ1_0 NVIDIA path: row-quantizer stores the exact integer q8 sum in ds.y and the warp-transposed q8 layout. HIP and MUSA keep the old quantizer and vec-dot. static constexpr __host__ __device__ bool ggml_cuda_q8_1_exact_isum(ggml_type type_src0) { -#if defined(GGML_USE_HIP) +#if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) GGML_UNUSED(type_src0); return false; #else @@ -1039,15 +1036,7 @@ static constexpr __host__ __device__ bool ggml_cuda_q8_1_exact_isum(ggml_type ty #endif } -// Warp-transposed (SoA) q8_1 activation layout for the ternary MMVQ. -// -// One PTQ1_0 K-block (128 weights) consumes 4 block_q8_1 = 36 words (32 qs + 4 ds). In the -// small-K MMVQ geometry each lane owns one K-block, so with the plain AoS layout a warp-wide -// load of "word w" touches 32 lines 144 B apart: ~36 L1 wavefronts per instruction, ~1300 per -// K-iteration against ~250 for the weights themselves. That LSU traffic, not GDDR, capped the -// PTQ1 GEMV near 370 GB/s on Ada. Here K-blocks are grouped by 32 and word w of the group is -// stored contiguously, so the same load is 32 consecutive words = 1 wavefront. -// Bytes per column are unchanged when K is padded to a multiple of 32*128 = 4096. +// Warp-transposed (SoA) q8_1 layout for NVIDIA PTQ1_0 MMVQ: word w of each 32-block group is contiguous so a warp load is one L1 wavefront. K padded to a multiple of 4096. #define GGML_CUDA_PTQ1_Q8_GROUP_KB 32 #define GGML_CUDA_PTQ1_Q8_WORDS_PER_KB 36 #define GGML_CUDA_PTQ1_Q8_GROUP_WORDS (GGML_CUDA_PTQ1_Q8_GROUP_KB * GGML_CUDA_PTQ1_Q8_WORDS_PER_KB) diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 5ec5e11a5feb..294b364d3724 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -669,15 +669,7 @@ static __global__ void mul_mat_vec_q( #if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) if constexpr (small_k && type == GGML_TYPE_PTQ1_0 && nwarps > 1 && rows_per_cuda_block == nwarps) { - // Warp-per-row geometry for small K. - // The generic small_k loop strides K-blocks across the whole block (kbx = tid; kbx += 128), so - // for K=5120 (40 blocks) only threads 0..39 ever load anything: warp 0 is full, warp 1 has - // 8 lanes, warps 2..3 only wait at __syncthreads while holding SM warp slots. Measured - // in-graph with CUPTI on an RTX 4070: 350-390 GB/s for every K=5120 GEMV vs 460 GB/s for the - // K=17408 geometry where all four warps stream. Here warp w owns row row0+w and its 32 lanes - // stride that row's K-blocks: every warp issues loads, the reduction is a shuffle, and - // there is no shared memory or block barrier. Activations are re-read per warp but they - // are 20 KB and L1/L2 resident. + // Warp-per-row small-K: warp w owns row row0+w and its lanes stride that row's K-blocks. No smem, no block barrier. const int warp = threadIdx.y; const int lane = threadIdx.x; const bool row_ok = uint32_t(row0 + warp) < stride_col_dst; @@ -812,9 +804,8 @@ static __global__ void mul_mat_vec_q( } #endif -#if !defined(GGML_USE_HIP) - // PTQ1_0 activations use the warp-transposed q8 layout (ggml_cuda_ptq1_q8_word), so the - // vec-dot is handed the column base and the K-block index instead of &y[kby]. +#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + // PTQ1_0 NVIDIA SoA q8 layout: pass the column base and K-block index, not &y[kby]. if constexpr (type == GGML_TYPE_PTQ1_0) { GGML_UNUSED(kby); GGML_UNUSED(kqs); @@ -998,9 +989,9 @@ static __global__ void mul_mat_vec_q_moe( const int kby = kbx * (qk/QK8_1); const int kqs = vdr * (threadIdx.x % (qi/vdr)); -#if !defined(GGML_USE_HIP) +#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) if constexpr (type == GGML_TYPE_PTQ1_0) { - // Warp-transposed q8 layout: hand the vec-dot the column base and K-block index. + // NVIDIA SoA q8: hand the vec-dot the column base and K-block index. GGML_UNUSED(kby); GGML_UNUSED(kqs); #pragma unroll diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu index 57db61122bd3..1aba903d46bd 100644 --- a/ggml/src/ggml-cuda/quantize.cu +++ b/ggml/src/ggml-cuda/quantize.cu @@ -51,10 +51,7 @@ static __device__ __forceinline__ float nvfp4_native_scale_error( #endif // CUDART_VERSION >= 12080 #endif // defined(BLACKWELL_MMA_AVAILABLE) -// exact_isum: store the integer sum of the quantized values (bit-cast into the ds.y half slot) -// instead of the float sum of the inputs. Ternary vec-dots (PTQ1_0) use it to fold the -// digit bias {0,1,2} -> {-1,0,+1} into one subtraction per 32-block instead of a SIMD byte -// subtract per 4 weights, with results bit-identical to the biased path. +// exact_isum: store the integer q8 sum in ds.y so PTQ1_0 can fold digit bias once per 32-block. Bit-identical to the per-nibble subtract. template __launch_bounds__(CUDA_QUANTIZE_BLOCK_SIZE, 1) static __global__ void quantize_q8_1( diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index 23b0c61c34c2..0c674e5c649d 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -804,7 +804,7 @@ static __device__ __forceinline__ float vec_dot_q2_0_q8_1( return d2 * d8 * sumi; } -#if !defined(GGML_USE_HIP) +#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) // y_col0: activation column 0 base (warp-transposed layout, see ggml_cuda_ptq1_q8_word); // kbx_x: absolute PTQ1 block index into vbq; kb: K-block index within the row; // stride_col_y: column stride in block_q8_1 units. @@ -818,8 +818,7 @@ static __device__ __forceinline__ void vec_dot_ptq1_0_q8_1_multi(const void * __ const block_ptq1_0 * bq = (const block_ptq1_0 *) vbq + kbx_x; int sumi[ncols_dst][4] = {}; - // Lane-coalesced activation words: word ww of this K-block is at yw[ww*32]; 32 lanes of a - // warp own 32 consecutive K-blocks, so every load below is 32 consecutive words. + // SoA q8: word ww of this K-block is at yw[ww*32], so a warp load is 32 consecutive words. const int * __restrict__ yw = (const int *) y_col0 + (kb >> 5) * GGML_CUDA_PTQ1_Q8_GROUP_WORDS + (kb & 31); const uint32_t sy = stride_col_y * (sizeof(block_q8_1) / 4); # define PTQ1_U(j, sub, m) yw[(j) * sy + ((sub) * 8 + (m)) * GGML_CUDA_PTQ1_Q8_GROUP_KB] @@ -986,8 +985,7 @@ static __device__ __forceinline__ float vec_dot_ptq1_0_q8_1(const void * __restr } return (float) bq->d * acc; #else - // Unreachable on CUDA: mul_mat_vec_q routes every PTQ1_0 ncols_dst through - // vec_dot_ptq1_0_q8_1_multi, which needs the column base for the warp-transposed q8 layout. + // HIP/MUSA stub: NVIDIA PTQ1_0 uses vec_dot_ptq1_0_q8_1_multi. GGML_UNUSED(vbq); GGML_UNUSED(bq8_1); GGML_UNUSED(kbx); From 68ea9b4d662bbdbb9455419db13826faa7dd15dd Mon Sep 17 00:00:00 2001 From: Cary Palmer <24235924+professorpalmer@users.noreply.github.com> Date: Mon, 21 Sep 2026 12:53:35 -0500 Subject: [PATCH 3/3] cuda: restore the scalar PTQ1_0 vec-dot on MUSA The MUSA gate compiled out the SoA fast path and left the generic entry returning zero. MUSA now takes the same scalar implementation as HIP. The exact-isum quantizer no longer reduces an unused float sum. --- ggml/src/ggml-cuda/quantize.cu | 5 ++-- ggml/src/ggml-cuda/vecdotq.cuh | 53 +++++++++++++++++++++++++++++++++- 2 files changed, 55 insertions(+), 3 deletions(-) diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu index 1aba903d46bd..368546c8ba69 100644 --- a/ggml/src/ggml-cuda/quantize.cu +++ b/ggml/src/ggml-cuda/quantize.cu @@ -86,10 +86,8 @@ static __global__ void quantize_q8_1( ggml_cuda_pdl_sync(); const float xi = i0 < ne00 ? x[i03*s03 + i02*s02 + i01*s01 + i00] : 0.0f; float amax = fabsf(xi); - float sum = xi; amax = warp_reduce_max(amax); - sum = warp_reduce_sum(sum); const float d = amax / 127.0f; const int8_t q = amax == 0.0f ? 0 : roundf(xi / d); @@ -112,6 +110,9 @@ static __global__ void quantize_q8_1( return; } + float sum = xi; + sum = warp_reduce_sum(sum); + y[ib].qs[iqs] = q; if (iqs > 0) { diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index 0c674e5c649d..02c4f062ff43 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -984,8 +984,59 @@ static __device__ __forceinline__ float vec_dot_ptq1_0_q8_1(const void * __restr acc += __low2float(bq8_1[iqs + k].ds) * (float) (sumi[k] - sumu[k]); } return (float) bq->d * acc; +#elif defined(GGML_USE_MUSA) + // MUSA cannot compile the HIP amdgcn path. Scalar decode, same as the pre-#211 generic entry. + const block_ptq1_0 * bq = (const block_ptq1_0 *) vbq + kbx; + int sumi[4] = { 0, 0, 0, 0 }; + +# pragma unroll + for (int m = 0; m < 16; ++m) { + uint32_t v = bq->qs[m]; +# pragma unroll + for (int t = 0; t < 5; ++t) { + const uint32_t w = v * 3; + const int q = (int) (w >> 8) - 1; + v = w & 0xFF; + const int e = t * 16 + m; + sumi[e >> 5] += q * (int) bq8_1[iqs + (e >> 5)].qs[e & 31]; + } + } + +# pragma unroll + for (int m = 0; m < 8; ++m) { + uint32_t v = bq->qs[16 + m]; +# pragma unroll + for (int t = 0; t < 5; ++t) { + const uint32_t w = v * 3; + const int q = (int) (w >> 8) - 1; + v = w & 0xFF; + const int e = 80 + t * 8 + m; + sumi[e >> 5] += q * (int) bq8_1[iqs + (e >> 5)].qs[e & 31]; + } + } + +# pragma unroll + for (int h = 0; h < 2; ++h) { + uint32_t v = bq->qh[h]; +# pragma unroll + for (int t = 0; t < 4; ++t) { + const uint32_t w = v * 3; + const int q = (int) (w >> 8) - 1; + v = w & 0xFF; + const int e = 120 + t * 2 + h; + sumi[e >> 5] += q * (int) bq8_1[iqs + (e >> 5)].qs[e & 31]; + } + } + + float acc = 0.0f; +# pragma unroll + for (int k = 0; k < 4; ++k) { + acc += __low2float(bq8_1[iqs + k].ds) * (float) sumi[k]; + } + return (float) bq->d * acc; #else - // HIP/MUSA stub: NVIDIA PTQ1_0 uses vec_dot_ptq1_0_q8_1_multi. + // NVIDIA PTQ1_0 goes through vec_dot_ptq1_0_q8_1_multi; this generic entry is unused there. + // The SoA _multi signature does not match a plain AoS call, so do not forward. GGML_UNUSED(vbq); GGML_UNUSED(bq8_1); GGML_UNUSED(kbx);