Skip to content

ggml: add SM70 small-batch Q6_K kernels #23

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

Merged
novkien merged 1 commit into from
Oct 4, 2026
Merged

ggml: add SM70 small-batch Q6_K kernels #23

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
81 changes: 76 additions & 5 deletions ggml/src/ggml-cuda/mmvq.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,14 @@
#include <cstdint>
#include <type_traits>

#ifndef GGML_CUDA_VOLTA_Q6_W4R4
#define GGML_CUDA_VOLTA_Q6_W4R4 1
#endif

#ifndef GGML_CUDA_VOLTA_Q6_W1R4
#define GGML_CUDA_VOLTA_Q6_W1R4 1
#endif

// only enabled on DGX Spark, where it is a gain on every type below. On the higher-bandwidth parts the kernel
// has little exposed latency left to hide and the extra requests cost more than they save.
// For perf data, see https://github.com/ggml-org/llama.cpp/pull/26705#issuecomment-5569335031
Expand Down Expand Up @@ -596,7 +604,7 @@ static constexpr __host__ __device__ int calc_rows_per_block(int ncols_dst, int
return 1;
}

template <ggml_type type, int ncols_dst, bool has_fusion, bool small_k = false, bool halve_iters = false>
template <ggml_type type, int ncols_dst, bool has_fusion, bool small_k = false, bool halve_iters = false, int rows_per_block_override = 0>
__launch_bounds__(calc_nwarps(type, ncols_dst, get_device_table_id(), small_k, halve_iters)*ggml_cuda_get_physical_warp_size(), 1)
static __global__ void mul_mat_vec_q(
const void * vx_ptr, const void * vy_ptr, const int32_t * ids_ptr, const ggml_cuda_mm_fusion_args_device fusion, float * dst_ptr,
Expand All @@ -615,7 +623,8 @@ static __global__ void mul_mat_vec_q(
constexpr int vdr = get_vdr_mmvq(type);
constexpr mmvq_parameter_table_id table_id = get_device_table_id();
constexpr int nwarps = calc_nwarps(type, ncols_dst, table_id, small_k, halve_iters);
constexpr int rows_per_cuda_block = calc_rows_per_block(ncols_dst, table_id, small_k, nwarps);
constexpr int rows_per_cuda_block = rows_per_block_override == 0 ?
calc_rows_per_block(ncols_dst, table_id, small_k, nwarps) : rows_per_block_override;
constexpr int warp_size = ggml_cuda_get_physical_warp_size();

constexpr vec_dot_q_cuda_t vec_dot_q_cuda = get_vec_dot_q_cuda(type);
Expand Down Expand Up @@ -1007,7 +1016,7 @@ static std::pair<dim3, dim3> calc_launch_params(
return {block_nums, block_dims};
}

template<ggml_type type, int c_ncols_dst, bool small_k = false, bool halve_iters = false>
template<ggml_type type, int c_ncols_dst, bool small_k = false, bool halve_iters = false, int rows_per_block_override = 0>
static void mul_mat_vec_q_switch_fusion(
const void * vx, const void * vy, const int32_t * ids, const ggml_cuda_mm_fusion_args_device fusion, float * dst,
const uint32_t ncols_x, const uint3 nchannels_y, const uint32_t stride_row_x, const uint32_t stride_col_y,
Expand All @@ -1022,7 +1031,7 @@ static void mul_mat_vec_q_switch_fusion(
if constexpr (c_ncols_dst == 1) {
if (has_fusion) {
const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(block_nums, block_dims, nbytes_shared, stream);
ggml_cuda_kernel_launch(mul_mat_vec_q<type, c_ncols_dst, true, small_k, halve_iters>, launch_params,
ggml_cuda_kernel_launch(mul_mat_vec_q<type, c_ncols_dst, true, small_k, halve_iters, rows_per_block_override>, launch_params,
vx, vy, ids, fusion, dst, ncols_x, nchannels_y, stride_row_x, stride_col_y, stride_col_dst,
channel_ratio, stride_channel_x, stride_channel_y, stride_channel_dst,
sample_ratio, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride);
Expand All @@ -1033,7 +1042,7 @@ static void mul_mat_vec_q_switch_fusion(
GGML_ASSERT(!has_fusion && "fusion only supported for ncols_dst=1");

const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(block_nums, block_dims, nbytes_shared, stream);
ggml_cuda_kernel_launch(mul_mat_vec_q<type, c_ncols_dst, false, small_k, halve_iters>, launch_params,
ggml_cuda_kernel_launch(mul_mat_vec_q<type, c_ncols_dst, false, small_k, halve_iters, rows_per_block_override>, launch_params,
vx, vy, ids, fusion, dst, ncols_x, nchannels_y, stride_row_x, stride_col_y, stride_col_dst,
channel_ratio, stride_channel_x, stride_channel_y, stride_channel_dst,
sample_ratio, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride);
Expand Down Expand Up @@ -1072,6 +1081,26 @@ static void mul_mat_vec_q_moe_launch(
}
}

#if GGML_CUDA_VOLTA_Q6_W4R4
static bool is_volta_q6_w4r4_shape(const int ncols_x, const int nrows_x) {
switch (ncols_x) {
case 2560:
return nrows_x == 6144 || nrows_x == 10240 || nrows_x == 12288;
case 5120:
return nrows_x == 1024 || nrows_x == 6144 || nrows_x == 10240 ||
nrows_x == 12288 || nrows_x == 17408;
case 6144:
return nrows_x == 2560 || nrows_x == 5120;
case 10240:
return nrows_x == 5120;
case 17408:
return nrows_x == 5120;
default:
return false;
}
}
#endif

template <ggml_type type>
static void mul_mat_vec_q_switch_ncols_dst(
const void * vx, const void * vy, const int32_t * ids, const ggml_cuda_mm_fusion_args_device fusion, float * dst,
Expand Down Expand Up @@ -1177,6 +1206,27 @@ static void mul_mat_vec_q_switch_ncols_dst(
// static, else MSVC lambda capture breaks the constexpr uses below
static constexpr int c_ncols_dst = 1;

#if GGML_CUDA_VOLTA_Q6_W1R4
if constexpr (type == GGML_TYPE_Q6_K) {
// Candidate for the exact Flash MTP Q6_K vocabulary-head shape.
const bool has_fusion = fusion.gate != nullptr || fusion.x_bias != nullptr || fusion.gate_bias != nullptr ||
fusion.x_scale != nullptr || fusion.gate_scale != nullptr;
if (GGML_CUDA_CC_IS_NVIDIA(cc) && cc == GGML_CUDA_CC_VOLTA &&
ncols_x == 2560 && nrows_x == 248320 && !has_ids && !has_fusion) {
constexpr int rows_per_block = 4;
const int nwarps = calc_nwarps(type, c_ncols_dst, table_id, false, false);
const dim3 block_nums(nrows_x / rows_per_block, nchannels_dst, nsamples_dst);
const dim3 block_dims(warp_size, nwarps, 1);
mul_mat_vec_q_switch_fusion<type, c_ncols_dst, false, false, rows_per_block>(
vx, vy, ids, fusion, dst, ncols_x, nchannels_y_fd, stride_row_x, stride_col_y, stride_col_dst,
channel_ratio_fd, stride_channel_x, stride_channel_y, stride_channel_dst,
sample_ratio_fd, stride_sample_x, stride_sample_y, stride_sample_dst,
block_nums, block_dims, 0, ids_stride, stream);
return;
}
}
#endif

// Tag types keep the flags compile-time, so __launch_bounds__ matches what is launched.
const auto launch = [&](auto small_k_tag, auto halve_iters_tag) {
constexpr bool c_small_k = decltype(small_k_tag)::value;
Expand Down Expand Up @@ -1222,6 +1272,27 @@ static void mul_mat_vec_q_switch_ncols_dst(
} break;
case 4: {
constexpr int c_ncols_dst = 4;
#if GGML_CUDA_VOLTA_Q6_W4R4
if constexpr (type == GGML_TYPE_Q6_K) {
const bool has_fusion = fusion.gate != nullptr || fusion.x_bias != nullptr || fusion.gate_bias != nullptr ||
fusion.x_scale != nullptr || fusion.gate_scale != nullptr;
// These exact Q6_K K/M pairs passed the paired operator screen.
if (GGML_CUDA_CC_IS_NVIDIA(cc) && cc == GGML_CUDA_CC_VOLTA &&
!has_ids && !has_fusion && nrows_x % 4 == 0 &&
is_volta_q6_w4r4_shape(ncols_x, nrows_x)) {
constexpr int rows_per_block = 4;
const int nwarps = calc_nwarps(type, c_ncols_dst, table_id);
const dim3 block_nums(nrows_x / rows_per_block, nchannels_dst, nsamples_dst);
const dim3 block_dims(warp_size, nwarps, 1);
mul_mat_vec_q_switch_fusion<type, c_ncols_dst, false, false, rows_per_block>(
vx, vy, ids, fusion, dst, ncols_x, nchannels_y_fd, stride_row_x, stride_col_y, stride_col_dst,
channel_ratio_fd, stride_channel_x, stride_channel_y, stride_channel_dst,
sample_ratio_fd, stride_sample_x, stride_sample_y, stride_sample_dst,
block_nums, block_dims, 0, ids_stride, stream);
return;
}
}
#endif
std::pair<dim3, dim3> dims = calc_launch_params<type>(c_ncols_dst, nrows_x, nchannels_dst, nsamples_dst, warp_size, table_id);
mul_mat_vec_q_switch_fusion<type, c_ncols_dst>(vx, vy, ids, fusion, dst, ncols_x, nchannels_y_fd, stride_row_x, stride_col_y, stride_col_dst,
channel_ratio_fd, stride_channel_x, stride_channel_y, stride_channel_dst,
Expand Down
74 changes: 74 additions & 0 deletions tests/test-backend-ops.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10268,6 +10268,43 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
test_cases.emplace_back(new test_mul_mat(
GGML_TYPE_BF16, GGML_TYPE_F32, 10240, 4, 320, { 2, 1 }, { 2, 1 }));

// Flash/MTP Q6_K W4 shapes, including the measured small-M fallback cases.
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 512, 4, 2560, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 640, 4, 2560, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 6144, 4, 2560, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 10240, 4, 2560, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 12288, 4, 2560, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 2560, 4, 6144, { 1, 1 }, { 1, 1 }));

// Small-M draft shape plus row and width fallback controls.
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 320, 4, 10240, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 6145, 4, 2560, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 6144, 3, 2560, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 6144, 5, 2560, { 1, 1 }, { 1, 1 }));

// Actual Qwen-27B Q6_K W4 tensor shapes.
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 17408, 4, 5120, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 5120, 4, 17408, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 5120, 4, 6144, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 10240, 4, 5120, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 6144, 4, 5120, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 1024, 4, 5120, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 12288, 4, 5120, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 5120, 4, 10240, { 1, 1 }, { 1, 1 }));

// Q6_K W4 layout controls: padded K view and batch broadcast.
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32,
6144, 4, 2560, { 1, 1 }, { 1, 1 }, { 0, 1, 2, 3 }, 2816));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32,
6144, 4, 2560, { 2, 1 }, { 2, 1 }));

// Candidate Flash Next MTP Q6_K vocabulary head plus row/width/K fallback controls.
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 248320, 1, 2560, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 248320, 2, 2560, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 248320, 4, 2560, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 248319, 1, 2560, { 1, 1 }, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 248320, 1, 2816, { 1, 1 }, { 1, 1 }));

// The SYCL backend picks between one and two output rows per subgroup by row count when there
// are two destination columns (Q4_K_MMVQ_ROW_PAIR_MIN_NROWS in ggml-sycl/mmvq.cpp). Cover both
// sides of that boundary, including an odd row count above it for the row-pair tail.
Expand Down Expand Up @@ -11666,6 +11703,43 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
test_cases.emplace_back(new test_mul_mat(
GGML_TYPE_BF16, GGML_TYPE_F32, 10240, 4, 320, {1, 1}, {1, 1}));

// Flash/MTP Q6_K W4 shapes, including the measured small-M fallback cases.
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 512, 4, 2560, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 640, 4, 2560, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 6144, 4, 2560, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 10240, 4, 2560, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 12288, 4, 2560, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 2560, 4, 6144, {1, 1}, {1, 1}));

// Small-M draft shape plus row and width fallback controls.
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 320, 4, 10240, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 6145, 4, 2560, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 6144, 3, 2560, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 6144, 5, 2560, {1, 1}, {1, 1}));

// Actual Qwen-27B Q6_K W4 tensor shapes.
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 17408, 4, 5120, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 5120, 4, 17408, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 5120, 4, 6144, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 10240, 4, 5120, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 6144, 4, 5120, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 1024, 4, 5120, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 12288, 4, 5120, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 5120, 4, 10240, {1, 1}, {1, 1}));

// Q6_K W4 layout controls: padded K view and batch broadcast.
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32,
6144, 4, 2560, {1, 1}, {1, 1}, {0, 1, 2, 3}, 2816));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32,
6144, 4, 2560, {2, 1}, {2, 1}));

// Candidate Flash Next MTP Q6_K vocabulary head plus row/width/K fallback controls.
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 248320, 1, 2560, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 248320, 2, 2560, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 248320, 4, 2560, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 248319, 1, 2560, {1, 1}, {1, 1}));
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 248320, 1, 2816, {1, 1}, {1, 1}));

// qwen3-30b-a3b
for (int bs : {1, 4, 8, 32, 64, 128, 256, 512}) {
for (ggml_type type_a : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_Q4_0, GGML_TYPE_Q8_0, GGML_TYPE_Q4_K, GGML_TYPE_Q6_K, GGML_TYPE_IQ2_XS, GGML_TYPE_IQ4_XS}) {
Expand Down
Loading