From: leejet <redacted>
Date: Sun, 26 Oct 2025 18:13:31 +0000 (+0800)
Subject: ggml: fix cuda kernel launch configuration for k_compute_batched_ptrs to support... 
X-Git-Tag: upstream/1.8.3~417
X-Git-Url: https://git.djapps.eu/?a=commitdiff_plain;h=4f4246dcb4657db3ba220502b237cddd7d68cf4c;p=pkg%2Fggml%2Fsources%2Fwhisper.cpp

ggml: fix cuda kernel launch configuration for k_compute_batched_ptrs to support large batch (llama/16744)

* fix k_compute_batched_ptrs

* add backend ops test

* Update ggml/src/ggml-cuda/ggml-cuda.cu

Co-authored-by: Johannes Gäßler <redacted>

* reduce the batch size

---------

Co-authored-by: Johannes Gäßler <redacted>
---

diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index 19f72975..6b688bfe 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 = 16;
+        const int threads_y = 16;
+        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<<<grid_dims, block_dims, 0, main_stream>>>(
                 src0_ptr, src1_ptr, dst_t,
                 ptrs_src.get(), ptrs_dst.get(),
                 ne12, ne13,