From c2cbb164cace436b9a779f5ab8ef3145a8573df5 Mon Sep 17 00:00:00 2001 From: leejet Date: Thu, 23 Oct 2025 22:57:27 +0800 Subject: [PATCH 1/4] fix k_compute_batched_ptrs --- ggml/src/ggml-cuda/ggml-cuda.cu | 11 +++++++++-- 1 file changed, 9 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index f5a6a751acfd5..4d5c14a88daf4 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1957,8 +1957,15 @@ static void ggml_cuda_mul_mat_batched_cublas_impl(ggml_backend_cuda_context & ct size_t src1_stride_size = sizeof(cuda_t); - dim3 block_dims(ne13, ne12); - k_compute_batched_ptrs<<<1, block_dims, 0, main_stream>>>( + const int threads_x = 32; + const int threads_y = 32; + dim3 block_dims(threads_x, threads_y); + + dim3 grid_dims( + (ne13 + threads_x - 1) / threads_x, + (ne12 + threads_y - 1) / threads_y + ); + k_compute_batched_ptrs<<>>( src0_ptr, src1_ptr, dst_t, ptrs_src.get(), ptrs_dst.get(), ne12, ne13, From e707662f25bf9d574b3d565c6476d4ec38178003 Mon Sep 17 00:00:00 2001 From: leejet Date: Thu, 23 Oct 2025 23:47:00 +0800 Subject: [PATCH 2/4] add backend ops test --- tests/test-backend-ops.cpp | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 991c62597962d..0a16069a47bcc 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -6562,6 +6562,10 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_mul_mat(type_a, type_b, 16, 1, 1024, {3, 2}, {1, 1})); test_cases.emplace_back(new test_mul_mat(type_a, type_b, 16, 8, 1024, {3, 2}, {1, 1})); test_cases.emplace_back(new test_mul_mat(type_a, type_b, 16, 16, 1024, {3, 2}, {1, 1})); + + // test cases with large batch size + test_cases.emplace_back(new test_mul_mat(type_a, type_b, 16, 8, 256, {1024, 2}, {1, 1})); + test_cases.emplace_back(new test_mul_mat(type_a, type_b, 16, 8, 256, {4096, 1}, {1, 1})); } } for (ggml_type type_a : other_types) { From 96e2320ce8714867792d83c302b2f2ce09c15862 Mon Sep 17 00:00:00 2001 From: leejet Date: Sat, 25 Oct 2025 00:50:57 +0800 Subject: [PATCH 3/4] Update ggml/src/ggml-cuda/ggml-cuda.cu MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Co-authored-by: Johannes Gäßler --- ggml/src/ggml-cuda/ggml-cuda.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 4d5c14a88daf4..d6be18183021f 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1957,8 +1957,8 @@ static void ggml_cuda_mul_mat_batched_cublas_impl(ggml_backend_cuda_context & ct size_t src1_stride_size = sizeof(cuda_t); - const int threads_x = 32; - const int threads_y = 32; + const int threads_x = 16; + const int threads_y = 16; dim3 block_dims(threads_x, threads_y); dim3 grid_dims( From 1da98548623dfe1176e5b2d72693f7bde1026530 Mon Sep 17 00:00:00 2001 From: leejet Date: Sat, 25 Oct 2025 00:52:35 +0800 Subject: [PATCH 4/4] reduce the batch size --- tests/test-backend-ops.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 0a16069a47bcc..be1ed7cf29ba3 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -6564,8 +6564,7 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_mul_mat(type_a, type_b, 16, 16, 1024, {3, 2}, {1, 1})); // test cases with large batch size - test_cases.emplace_back(new test_mul_mat(type_a, type_b, 16, 8, 256, {1024, 2}, {1, 1})); - test_cases.emplace_back(new test_mul_mat(type_a, type_b, 16, 8, 256, {4096, 1}, {1, 1})); + test_cases.emplace_back(new test_mul_mat(type_a, type_b, 16, 8, 256, {1536, 1}, {1, 1})); } } for (ggml_type type_a : other_types) {