Skip to content

ggml-cuda: pad GEMM weight rows to avoid cache-set aliasing #34

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

Draft
roberteg16 wants to merge 1 commit into
base: gfx11
Choose a base branch
from
Draft

ggml-cuda: pad GEMM weight rows to avoid cache-set aliasing #34

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
1 change: 1 addition & 0 deletions ggml/include/ggml.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -647,6 +647,7 @@ extern "C" {
GGML_TENSOR_FLAG_PARAM = 4, // ...contains trainable parameters
GGML_TENSOR_FLAG_LOSS = 8, // ...defines loss for numerical optimization (multiple loss tensors add up)
GGML_TENSOR_FLAG_COMPUTE = 16, // ...must be computed
GGML_TENSOR_FLAG_PAD_ROWS = 32, // ...is a GEMM weight whose row stride a backend may pad to avoid cache-set aliasing
};

enum ggml_tri_type {
Expand Down
9 changes: 9 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 @@ -177,6 +177,15 @@ static int ggml_cuda_highest_compiled_arch(const int arch) {

#define GGML_CUDA_MAX_STREAMS 8

// True when each row is internally contiguous and the higher dimensions are tightly nested,
// while the row stride nb[1] may exceed the packed row size. Such a tensor can be fed to a
// strided GEMM (leading dimension = nb[1] / type_size) without repacking to contiguous.
static inline bool ggml_cuda_is_contiguous_rows(const ggml_tensor * t) {
return t->nb[0] == ggml_type_size(t->type)
&& t->nb[2] == t->nb[1]*t->ne[1]
&& t->nb[3] == t->nb[2]*t->ne[2];
}

[[noreturn]]
void ggml_cuda_error(const char * stmt, const char * func, const char * file, int line, const char * msg);

Expand Down
86 changes: 78 additions & 8 deletions ggml/src/ggml-cuda/ggml-cuda.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -652,6 +652,39 @@ static void * ggml_backend_cuda_buffer_get_base(ggml_backend_buffer_t buffer) {
return ctx->dev_ptr;
}

// One cache line, added to a weight's row stride when the packed row size is a multiple of

@mgehre-amd mgehre-amd Jul 6, 2026

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

I guess this needs some gfx11 qualification as this is not generally applicable to all CUDA GPUs

// 2048 bytes to avoid cache-set aliasing in matmul. Enabled by default; disable with
// GGML_CUDA_NO_PAD_WEIGHTS.
#define GGML_CUDA_CACHE_LINE 128
#define GGML_CUDA_ROW_ALIAS_STRIDE 2048

static bool ggml_cuda_pad_weights_enabled() {
static const bool enabled = getenv("GGML_CUDA_NO_PAD_WEIGHTS") == nullptr;
return enabled;
}

// A float GEMM weight (tagged GGML_TENSOR_FLAG_PAD_ROWS by the loader) whose packed row size
// aliases at GGML_CUDA_ROW_ALIAS_STRIDE and therefore benefits from a one-cache-line row-stride
// pad. Quantized weights are excluded: a 128-byte pad is not a multiple of their block size and
// would misalign the block-indexed matmul kernels.
static bool ggml_cuda_should_pad_weight(const ggml_tensor * tensor) {
if (!ggml_cuda_pad_weights_enabled()) {
return false;
}
if (tensor->view_src != nullptr || !(tensor->flags & GGML_TENSOR_FLAG_PAD_ROWS)
|| ggml_is_quantized(tensor->type) || tensor->ne[1] <= 1) {
return false;
}
return ggml_row_size(tensor->type, tensor->ne[0]) % GGML_CUDA_ROW_ALIAS_STRIDE == 0;
}

// Row stride (nb[1]) in bytes for a weight tensor: the packed row size, plus one cache line
// when padding applies.
static size_t ggml_cuda_padded_nb1(const ggml_tensor * tensor) {
const size_t packed = ggml_row_size(tensor->type, tensor->ne[0]);
return ggml_cuda_should_pad_weight(tensor) ? packed + GGML_CUDA_CACHE_LINE : packed;
}

static enum ggml_status ggml_backend_cuda_buffer_init_tensor(ggml_backend_buffer_t buffer, ggml_tensor * tensor) {
ggml_backend_cuda_buffer_context * ctx = (ggml_backend_cuda_buffer_context *)buffer->context;

Expand All @@ -660,6 +693,19 @@ static enum ggml_status ggml_backend_cuda_buffer_init_tensor(ggml_backend_buffer
return GGML_STATUS_SUCCESS;
}

if (ggml_backend_buffer_get_usage(buffer) != GGML_BACKEND_BUFFER_USAGE_COMPUTE
&& ggml_cuda_should_pad_weight(tensor)) {
tensor->nb[1] = ggml_cuda_padded_nb1(tensor);
tensor->nb[2] = tensor->nb[1] * tensor->ne[1];
tensor->nb[3] = tensor->nb[2] * tensor->ne[2];

// zero the allocation so the inter-row padding gaps hold no NaNs
ggml_cuda_set_device(ctx->device);
CUDA_CHECK(cudaMemset((char *) tensor->data, 0,
ggml_backend_buft_get_alloc_size(buffer->buft, tensor)));
return GGML_STATUS_SUCCESS;
}

if (ggml_is_quantized(tensor->type) && tensor->view_src == nullptr && ggml_backend_buffer_get_usage(buffer) != GGML_BACKEND_BUFFER_USAGE_COMPUTE) {
// initialize padding to 0 to avoid possible NaN values
const size_t original_size = ggml_nbytes(tensor);
Expand Down Expand Up @@ -816,6 +862,16 @@ static size_t ggml_backend_cuda_buffer_type_get_alloc_size(ggml_backend_buffer_t
}
}

// K-dimension cache-line padding: reserve the padded per-row stride for weight matrices.
if (ggml_cuda_should_pad_weight(tensor)) {
const size_t padded_nb1 = ggml_cuda_padded_nb1(tensor);
const int64_t nrows = tensor->ne[1] * tensor->ne[2] * tensor->ne[3];
const size_t padded_sz = padded_nb1 * nrows;
if (padded_sz > size) {
size = padded_sz;
}
}

return size;
}

Expand Down Expand Up @@ -1642,6 +1698,15 @@ static void ggml_cuda_op_mul_mat_cublas(

const int64_t row_diff = row_high - row_low;

// src0 leading dimension: honor a padded row stride when src0 is passed through un-converted
// (row-contiguous). When src0 is converted/dequantized into a packed temp buffer, or copied
// to contiguous upstream, lda is the packed ne00.
const bool src0_row_contig = ggml_cuda_is_contiguous_rows(src0);
const int64_t s01 = (int64_t) (src0->nb[1] / ggml_type_size(src0->type));
const int64_t lda_f16 = (src0->type == GGML_TYPE_F16 && src0_row_contig) ? s01 : ne00;
const int64_t lda_f32 = (src0->type == GGML_TYPE_F32 && src0_row_contig) ? s01 : ne00;
const int64_t lda_bf16 = (src0->type == GGML_TYPE_BF16 && src0_row_contig) ? s01 : ne00;

int id = ggml_cuda_get_device();

// the main device has a larger memory buffer to hold the results from all GPUs
Expand All @@ -1656,11 +1721,11 @@ static void ggml_cuda_op_mul_mat_cublas(
const bool use_fp16 =
src0->type != GGML_TYPE_NVFP4 &&
(src0->type == GGML_TYPE_F16 || ggml_is_quantized(src0->type)) &&
ggml_is_contiguous(src0) &&
(ggml_is_contiguous(src0) || src0_row_contig) &&
row_diff == src0->ne[1] &&
dst->op_params[0] == GGML_PREC_DEFAULT;

if (supports_bf16 && src0->type == GGML_TYPE_BF16 && ggml_is_contiguous(src0) && row_diff == src0->ne[1]) {
if (supports_bf16 && src0->type == GGML_TYPE_BF16 && (ggml_is_contiguous(src0) || src0_row_contig) && row_diff == src0->ne[1]) {
ggml_cuda_pool_alloc<nv_bfloat16> src1_as_bf16(ctx.pool(id));
if (src1->type != GGML_TYPE_BF16) {
const to_bf16_cuda_t to_bf16_cuda = ggml_get_to_bf16_cuda(src1->type);
Expand All @@ -1680,7 +1745,7 @@ static void ggml_cuda_op_mul_mat_cublas(
CUBLAS_CHECK(
cublasGemmEx(ctx.cublas_handle(id), CUBLAS_OP_T, CUBLAS_OP_N,
row_diff, src1_ncols, ne10,
&alpha_f32, src0_ptr, CUDA_R_16BF, ne00,
&alpha_f32, src0_ptr, CUDA_R_16BF, lda_bf16,
src1_ptr, CUDA_R_16BF, ne10,
&beta_f32, dst_bf16.get(), CUDA_R_16BF, ldc,
CUBLAS_COMPUTE_32F,
Expand Down Expand Up @@ -1724,7 +1789,7 @@ static void ggml_cuda_op_mul_mat_cublas(
CUBLAS_CHECK(
cublasGemmEx(ctx.cublas_handle(id), CUBLAS_OP_T, CUBLAS_OP_N,
row_diff, src1_ncols, ne10,
&alpha, src0_ptr, CUDA_R_16F, ne00,
&alpha, src0_ptr, CUDA_R_16F, lda_f16,
src1_ptr, CUDA_R_16F, ne10,
&beta, dst_dd_i, CUDA_R_32F, ldc,
CUBLAS_COMPUTE_32F,
Expand All @@ -1738,7 +1803,7 @@ static void ggml_cuda_op_mul_mat_cublas(
CUBLAS_CHECK(
cublasGemmEx(ctx.cublas_handle(id), CUBLAS_OP_T, CUBLAS_OP_N,
row_diff, src1_ncols, ne10,
&alpha_f16, src0_ptr, CUDA_R_16F, ne00,
&alpha_f16, src0_ptr, CUDA_R_16F, lda_f16,
src1_ptr, CUDA_R_16F, ne10,
&beta_f16, dst_f16.get(), CUDA_R_16F, ldc,
CUBLAS_COMPUTE_16F,
Expand Down Expand Up @@ -1774,7 +1839,7 @@ static void ggml_cuda_op_mul_mat_cublas(
CUBLAS_CHECK(
cublasSgemm(ctx.cublas_handle(id), CUBLAS_OP_T, CUBLAS_OP_N,
row_diff, src1_ncols, ne10,
&alpha, src0_ddf_i, ne00,
&alpha, src0_ddf_i, lda_f32,
src1_ddf1_i, ne10,
&beta, dst_dd_i, ldc));
}
Expand Down Expand Up @@ -1845,7 +1910,10 @@ static void ggml_cuda_op_mul_mat(
const size_t q8_1_ts = sizeof(block_q8_1);
const size_t q8_1_bs = QK8_1;

const bool src0_is_contiguous = ggml_is_contiguous(src0);
// A non-quantized weight that is row-contiguous but has a padded row stride can be fed to
// cuBLAS directly (via lda); no need to copy it into a packed buffer.
const bool src0_is_contiguous = ggml_is_contiguous(src0)
|| (!ggml_is_quantized(src0->type) && ggml_cuda_is_contiguous_rows(src0));
const bool src1_is_contiguous = ggml_is_contiguous(src1);

const int64_t src1_padded_col_size = GGML_PAD(ne10, MATRIX_ROW_PADDING);
Expand Down Expand Up @@ -2018,7 +2086,9 @@ static void ggml_cuda_op_mul_mat(
}

// for split tensors the data begins at i0 == i0_offset_low
const size_t nbytes_src0_matrix = ne01*ne00*src0_ts / src0_bs;
// use the real per-matrix stride (nb[2]) so a padded row stride is honored;
// when src0 was copied to a packed temp buffer, fall back to the packed size
const size_t nbytes_src0_matrix = src0_is_contiguous ? src0->nb[2] : (ne01*ne00*src0_ts / src0_bs);
char * src0_dd_i = dev[id].src0_dd + ((i03/i03_divisor)*ne02 + (i02/i02_divisor)) * nbytes_src0_matrix;
float * src1_ddf_i = dev[id].src1_ddf + (i0*ne11 + src1_col_0) * ne10;
char * src1_ddq_i = dev[id].src1_ddq + src1_ddq_i_offset;
Expand Down
6 changes: 4 additions & 2 deletions ggml/src/ggml-cuda/mmvf.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -740,8 +740,10 @@ void ggml_cuda_op_mul_mat_vec_f(
const int cc = ggml_cuda_info().devices[id].cc;
const enum ggml_prec prec = fast_fp16_available(cc) ? ggml_prec(dst->op_params[0]) : GGML_PREC_F32;

// ggml_cuda_op provides single, contiguous matrices
const int64_t stride_row = ne00;
// honor a padded row stride (nb[1]) when src0 is passed through un-copied; equals ne00 for
// the packed case
const int64_t stride_row = ggml_cuda_is_contiguous_rows(src0)
? (int64_t) (src0->nb[1] / ggml_type_size(src0->type)) : ne00;
const int64_t stride_col_y = ne10;
const int64_t stride_col_dst = id == ctx.device ? ne0 : row_diff; // main device has larger memory buffer
const int64_t nchannels_x = 1;
Expand Down
34 changes: 32 additions & 2 deletions src/llama-model-loader.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1277,6 +1277,20 @@ struct ggml_tensor * llama_model_loader::create_tensor(
struct ggml_tensor * tensor = ggml_dup_tensor(ctx, cur);
ggml_set_name(tensor, ggml_get_name(cur));

// Mark GEMM weights (used as mul_mat / mul_mat_id src0) so backends may pad their row
// stride to avoid cache-set aliasing. Only real weights ("weight" suffix), not bias/scale.
{
llm_tensor weight_tensor = tn.tensor;
if (tn.tensor == LLM_TENSOR_TOKEN_EMBD && (flags & TENSOR_DUPLICATED)) {
weight_tensor = LLM_TENSOR_OUTPUT;
}
const bool is_weight = tn.suffix == nullptr || strcmp(tn.suffix, "weight") == 0;
const ggml_op weight_op = llm_tensor_info_for(weight_tensor).op;
if (is_weight && (weight_op == GGML_OP_MUL_MAT || weight_op == GGML_OP_MUL_MAT_ID)) {
tensor->flags |= GGML_TENSOR_FLAG_PAD_ROWS;
}
}

if (duplicated) {
size_data += ggml_nbytes(cur);
} else {
Expand Down Expand Up @@ -1535,6 +1549,16 @@ bool llama_model_loader::load_all_data(

size_t n_size = ggml_nbytes(cur);

// K-dim cache-line padding: the on-disk data is packed, but the destination tensor
// may carry a padded per-row stride (nb[1]). Read the packed size and upload the
// rows with a strided 2D copy.
const size_t packed_row = ggml_row_size(cur->type, cur->ne[0]);
const int64_t n_rows = ggml_nelements(cur) / cur->ne[0];
const bool row_padded = (size_t) cur->nb[1] != packed_row;
if (row_padded) {
n_size = packed_row * n_rows;
}

if (use_mmap) {
const auto & mapping = mappings.at(weight->idx);
ggml_backend_buffer_t buf_mmap = nullptr;
Expand All @@ -1560,6 +1584,8 @@ bool llama_model_loader::load_all_data(
auto & mmap_used = mmaps_used[weight->idx];
mmap_used.first = std::min(mmap_used.first, weight->offs);
mmap_used.second = std::max(mmap_used.second, weight->offs + n_size);
} else if (row_padded) {
ggml_backend_tensor_set_2d(cur, data, 0, packed_row, n_rows, cur->nb[1], packed_row);
} else {
ggml_backend_tensor_set(cur, data, 0, n_size);
}
Expand All @@ -1576,7 +1602,7 @@ bool llama_model_loader::load_all_data(
}
} else {
// If upload_backend is valid load the tensor in chunks to pinned memory and upload the buffers asynchronously to the GPU.
if (upload_backend) {
if (upload_backend && !row_padded) {
size_t offset = weight->offs;
alignment = file->read_alignment();
size_t aligned_offset = offset & ~(alignment - 1);
Expand Down Expand Up @@ -1632,7 +1658,11 @@ bool llama_model_loader::load_all_data(
read_buf.resize(n_size);
file->seek(weight->offs, SEEK_SET);
file->read_raw(read_buf.data(), n_size);
ggml_backend_tensor_set(cur, read_buf.data(), 0, n_size);
if (row_padded) {
ggml_backend_tensor_set_2d(cur, read_buf.data(), 0, packed_row, n_rows, cur->nb[1], packed_row);
} else {
ggml_backend_tensor_set(cur, read_buf.data(), 0, n_size);
}
if (check_tensors && !ggml_validate_row_data(cur->type, read_buf.data(), n_size)) {
throw std::runtime_error(format("tensor '%s' has invalid data", ggml_get_name(cur)));
}
Expand Down
Loading