Skip to content

cuda: PTQ1_0 decode on consumer Ampere/Ada: SoA q8 activations, warp-per-row small-K GEMV (+10.5% TG, +14.7% with #220) #215

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
professorpalmer wants to merge 3 commits into from
+257 −26
Closed

cuda: PTQ1_0 decode on consumer Ampere/Ada: SoA q8 activations, warp-per-row small-K GEMV (+10.5% TG, +14.7% with #220) #215

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
9150e75
cuda: PTQ1_0 decode on consumer Ampere/Ada: SoA q8 activations, warp-…
professorpalmer Sep 19, 2026
68c1e5f
cuda: keep the PTQ1_0 SoA/exact-isum path off on MUSA, matching HIP.
professorpalmer Sep 21, 2026
68ea9b4
cuda: restore the scalar PTQ1_0 vec-dot on MUSA
professorpalmer Sep 21, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
26 changes: 26 additions & 0 deletions ggml/src/ggml-cuda/common.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1026,6 +1026,32 @@ struct ggml_cuda_type_traits<GGML_TYPE_PTQ1_0> {
static constexpr int qi = QI_PTQ1_0;
};

// 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) || defined(GGML_USE_MUSA)
GGML_UNUSED(type_src0);
return false;
#else
return type_src0 == GGML_TYPE_PTQ1_0;
Comment on lines +1030 to +1035

@professorpalmer professorpalmer Sep 21, 2026

Copy link
Copy Markdown
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

You're right: MUSA took the non-HIP branch, so exact-isum and the SoA vec-dot would have gone live. 1b14d26 gates ggml_cuda_q8_1_exact_isum and those vec-dot paths on !HIP && !MUSA. We do not have a MUSA box to re-validate; the old quantizer/vec-dot is what that backend compiled before this PR.

#endif
}

// 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)
#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<GGML_TYPE_Q4_0> {
static constexpr int qk = QK4_0;
Expand Down
125 changes: 118 additions & 7 deletions ggml/src/ggml-cuda/mmvq.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -667,6 +667,81 @@ 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 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;

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<ncols_dst>(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<ncols_dst>(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<warp_size>(acc[j]);
if constexpr (has_fusion && has_gate) {
acc_gate[j] = warp_reduce_sum<warp_size>(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<type == GGML_TYPE_Q1_0, block_q1_0,
Expand Down Expand Up @@ -715,12 +790,29 @@ static __global__ void mul_mat_vec_q(
// x block quant index when casting the quants to int
const int kqs = vdr * (tid % (qi/vdr));

#if !defined(GGML_USE_HIP)
if constexpr (type == GGML_TYPE_PTQ1_0 && ncols_dst > 1 && ncols_dst <= 3) {
#if defined(__CUDA_ARCH__) && !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
if constexpr (type == GGML_TYPE_PTQ1_0) {
const int kbx_prefetch = kbx + blocks_per_iter;
if (kbx_prefetch < blocks_per_row_x && tid % (qi / vdr) == 0) {
# pragma unroll
for (int i = 0; i < rows_per_cuda_block; ++i) {
const block_ptq1_0 * prefetch_ptr =
(const block_ptq1_0 *) vx + kbx_offset + i * stride_row_x + kbx_prefetch;
asm volatile("prefetch.global.L2 [%0];" ::"l"(__cvta_generic_to_global(prefetch_ptr)));
}
}
}
#endif

#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);
# pragma unroll
for (int i = 0; i < rows_per_cuda_block; ++i) {
float dots[ncols_dst];
vec_dot_ptq1_0_q8_1_multi<ncols_dst>(vx, &y[kby], kbx_offset + i * stride_row_x + kbx, kqs,
vec_dot_ptq1_0_q8_1_multi<ncols_dst>(vx, y, kbx_offset + i * stride_row_x + kbx, kbx,
stride_col_y, dots);
# pragma unroll
for (int j = 0; j < ncols_dst; ++j) {
Expand All @@ -729,7 +821,7 @@ static __global__ void mul_mat_vec_q(

if constexpr (has_fusion) {
if constexpr (has_gate) {
vec_dot_ptq1_0_q8_1_multi<ncols_dst>(vgate, &y[kby], kbx_offset + i * stride_row_x + kbx, kqs,
vec_dot_ptq1_0_q8_1_multi<ncols_dst>(vgate, y, kbx_offset + i * stride_row_x + kbx, kbx,
stride_col_y, dots);
# pragma unroll
for (int j = 0; j < ncols_dst; ++j) {
Expand Down Expand Up @@ -897,9 +989,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) && !defined(GGML_USE_MUSA)
if constexpr (type == GGML_TYPE_PTQ1_0) {
// NVIDIA SoA q8: 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);
}
}
}

Expand Down Expand Up @@ -1433,7 +1540,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<char> 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;
Expand Down
32 changes: 28 additions & 4 deletions ggml/src/ggml-cuda/quantize.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -51,6 +51,8 @@ static __device__ __forceinline__ float nvfp4_native_scale_error(
#endif // CUDART_VERSION >= 12080
#endif // defined(BLACKWELL_MMA_AVAILABLE)

// 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 <bool exact_isum>
__launch_bounds__(CUDA_QUANTIZE_BLOCK_SIZE, 1)
static __global__ void quantize_q8_1(
const float * x_ptr, void * vy_ptr,
Expand Down Expand Up @@ -84,14 +86,33 @@ 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<QK8_1>(amax);
sum = warp_reduce_sum<QK8_1>(sum);

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<QK8_1>(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<const int32_t *>(&ds);
return;
}

float sum = xi;
sum = warp_reduce_sum<QK8_1>(sum);

y[ib].qs[iqs] = q;

if (iqs > 0) {
Expand Down Expand Up @@ -648,8 +669,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<true>, launch_params, x, vy, ne00, s01, s02, s03, ne0, ne1, ne2_fastdiv);
} else {
ggml_cuda_kernel_launch(quantize_q8_1<false>, launch_params, x, vy, ne00, s01, s02, s03, ne0, ne1, ne2_fastdiv);
}
}

void quantize_mmq_q8_1_cuda(
Expand Down
Loading