Skip to content

cuda: bound mmvq rows by the per-slot stride for MUL_MAT_ID #258

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
bri-prism merged 1 commit into from
Sep 29, 2026
Merged

cuda: bound mmvq rows by the per-slot stride for MUL_MAT_ID #258

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
7 changes: 5 additions & 2 deletions ggml/src/ggml-cuda/mmvq.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -597,6 +597,9 @@ static __global__ void mul_mat_vec_q(
const uint32_t sample_x = fastdiv(sample_dst, sample_ratio);
const uint32_t sample_y = sample_dst;

// dst rows are contiguous; with ids the per-slot row stride is stride_channel_dst, not stride_col_dst
const uint32_t nrows_dst = ids ? stride_channel_dst : stride_col_dst;

constexpr bool use_gate = has_gate;
bool use_bias = false;
bool use_gate_bias = false;
Expand Down Expand Up @@ -634,7 +637,7 @@ static __global__ void mul_mat_vec_q(
// 2. load only on threads that won't die after partial sum calculation
const uint32_t channel_bias = ids ? channel_x : channel_dst;
if (threadIdx.x < rows_per_cuda_block && threadIdx.y == 0 &&
(rows_per_cuda_block == 1 || uint32_t(row0 + threadIdx.x) < stride_col_dst)) {
(rows_per_cuda_block == 1 || uint32_t(row0 + threadIdx.x) < nrows_dst)) {
if (use_bias) {
x_bias = x_bias + sample_dst * stride_sample_dst + channel_bias * stride_channel_dst + row0;
#pragma unroll
Expand Down Expand Up @@ -805,7 +808,7 @@ static __global__ void mul_mat_vec_q(
}
}

if (threadIdx.x == i && (rows_per_cuda_block == 1 || uint32_t(row0 + i) < stride_col_dst)) {
if (threadIdx.x == i && (rows_per_cuda_block == 1 || uint32_t(row0 + i) < nrows_dst)) {
float result = tmp[j][i];
if constexpr (has_fusion) {
if constexpr (type == GGML_TYPE_NVFP4) {
Expand Down
Loading