furyhawk/ik_llama.cpp
1
0
Fork 0
mirror of https://github.com/ikawrakow/ik_llama.cpp.git synced 2026-07-21 02:05:35 +00:00
Code Issues Packages Projects Releases Wiki Activity

CUDA: fix MUL with non-contiguous src0 and scalar src1 (#2072)

Browse Source 
The scalar fast-path in ggml_cuda_op_mul routes to ggml_cuda_op_scale_tensor,
which reads src0 and writes dst as flat contiguous buffers. With a non-contiguous
src0 (for example a row-gapped view) this ignored the per-row strides: only the
first row was correct and later rows read the wrong memory. The CPU backend
respects the strides, so the two backends diverged.

Guard the fast-path on contiguous src0 and dst; non-contiguous inputs now fall
through to the general bin_bcast path, which honours the strides.

Add a non-contiguous test_bin_bcast variant covering both the scalar (scale) and
vector (general) paths.

Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
This commit is contained in:
Joel Farthing
2026-07-03 08:27:03 +02:00
committed by GitHub
co-authored by Joel Farthing
parent a5e41bc210
commit 86d8e9a13c
2 changed files with 27 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/binbcast.cu
+5 -1
View File
@@ -575,7 +575,11 @@ static __global__ void k_mul_fast(int ne0, int nelem, const float * x, const flo
}
void ggml_cuda_op_mul(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
if (ggml_nelements(dst->src[1]) == 1 && dst->src[1]->type == GGML_TYPE_F32 && dst->src[0]->type == GGML_TYPE_F32) {
// The scalar fast-path reads src0 and writes dst as flat/contiguous buffers, so it is only
// valid when both are contiguous. A non-contiguous (e.g. row-gapped view) src0 must fall
// through to the general bin_bcast path, which honours the per-dimension strides.
if (ggml_nelements(dst->src[1]) == 1 && dst->src[1]->type == GGML_TYPE_F32 && dst->src[0]->type == GGML_TYPE_F32 &&
ggml_is_contiguous(dst->src[0]) && ggml_is_contiguous(dst)) {
ggml_cuda_op_scale_tensor(ctx, dst);
return;
}
tests/test-backend-ops.cpp
+22 -4
View File
@@ -850,9 +850,10 @@ struct test_bin_bcast : public test_case {
const ggml_type type;
const std::array<int64_t, 4> ne;
const std::array<int, 4> nr;
const int nc; // 1: make src0 a non-contiguous view (row-gapped) instead of a fresh tensor
std::string vars() override {
return VARS_TO_STR3(type, ne, nr);
return VARS_TO_STR4(type, ne, nr, nc);
}
size_t op_size(ggml_tensor * t) override {
@@ -861,11 +862,23 @@ struct test_bin_bcast : public test_case {
test_bin_bcast(op_t op, ggml_type type = GGML_TYPE_F32,
std::array<int64_t, 4> ne = {10, 10, 1, 1},
std::array<int, 4> nr = {1, 2, 1, 1})
: op(op), type(type), ne(ne), nr(nr) {}
std::array<int, 4> nr = {1, 2, 1, 1},
int nc = 0)
: op(op), type(type), ne(ne), nr(nr), nc(nc) {}
ggml_tensor * build_graph(ggml_context * ctx) override {
ggml_tensor * a = ggml_new_tensor_4d(ctx, type, ne[0]*nr[0], ne[1]*nr[1], ne[2]*nr[2], ne[3]*nr[3]);
ggml_tensor * a;
if (nc) {
// non-contiguous src0: allocate padded in dim0 and view it back, so nb[1] > ne[0]*nb[0]
// (contiguous within a row, a gap between rows). This is the strided-view shape that
// openPangu's mHC produces and that the CUDA scalar-mul path used to misread.
ggml_tensor * a_full = ggml_new_tensor_4d(ctx, type,
2*ne[0]*nr[0], ne[1]*nr[1], ne[2]*nr[2], ne[3]*nr[3]);
a = ggml_view_4d(ctx, a_full, ne[0]*nr[0], ne[1]*nr[1], ne[2]*nr[2], ne[3]*nr[3],
a_full->nb[1], a_full->nb[2], a_full->nb[3], 0);
} else {
a = ggml_new_tensor_4d(ctx, type, ne[0]*nr[0], ne[1]*nr[1], ne[2]*nr[2], ne[3]*nr[3]);
}
ggml_tensor * b = ggml_new_tensor(ctx, type, 4, ne.data());
ggml_tensor * out = op(ctx, a, b);
return out;
@@ -2219,6 +2232,11 @@ static bool test_backend(ggml_backend_t backend, test_mode mode, const char * op
add_test_bin_bcast(GGML_TYPE_F32, {16, 10, 10, 10}, {1, 2, 2, 2});
add_test_bin_bcast(GGML_TYPE_F32, {16, 10, 10, 10}, {2, 2, 2, 2});
// non-contiguous src0 (row-gapped view), the openPangu mHC strided-view shape.
// The scalar-src1 case used to take a CUDA scale fast-path that read src0 as flat/contiguous.
test_cases.emplace_back(new test_bin_bcast(ggml_mul, GGML_TYPE_F32, {1, 1, 1, 1}, {4, 32, 1, 1}, /*nc=*/1)); // scalar src1 (scale path)
test_cases.emplace_back(new test_bin_bcast(ggml_mul, GGML_TYPE_F32, {1, 32, 1, 1}, {4, 1, 1, 1}, /*nc=*/1)); // vector src1 (general bin_bcast path)
// stable diffusion
add_test_bin_bcast(GGML_TYPE_F32, {1280, 1, 1, 1}, {1, 1, 1, 1});
add_test_bin_bcast(GGML_TYPE_F32, {1280, 1, 1, 1}, {1, 16, 16, 1});