Commit b1c2863e2 for llama.cpp

commit b1c2863e2c2c861ab3009aad8622826472100750
Author: lingyezhixing <144504450+lingyezhixing@users.noreply.github.com>
Date:   Tue Sep 22 00:11:29 2026 +0800

    cuda: fix sm_70 tile compilation error (#29224)

    The 5-argument load_ldmatrix added in 1884824fd only defines tile<16,8>, so the Volta tile<8,4> does not match. See https://github.com/ggml-org/llama.cpp/issues/29222 for details. Building on 1884824fd, generalize the tile shape of the 5-argument load_ldmatrix from <16,8> to <I,J>, so the non-swizzle branch forwards to the 3-argument loader for any shape. Local compilation and testing passed.

    Assisted-by: DeepSeek V4.1 Flash (OpenCode)

diff --git a/ggml/src/ggml-cuda/mma.cuh b/ggml/src/ggml-cuda/mma.cuh
index 6af2b6a14..3583ba5e1 100644
--- a/ggml/src/ggml-cuda/mma.cuh
+++ b/ggml/src/ggml-cuda/mma.cuh
@@ -873,14 +873,16 @@ namespace ggml_cuda_mma {
     }

     // Load from tile element (i0, j0), swz tells if the tile is stored swizzled.
-    template <bool swz, typename T, data_layout dl>
+    template <bool swz, int I, int J, typename T, data_layout dl>
     static __device__ __forceinline__ void load_ldmatrix(
-            tile<16, 8, T, dl> & t, const T * __restrict__ tile_base, const int i0, const int j0, const int stride) {
+            tile<I, J, T, dl> & t, const T * __restrict__ tile_base, const int i0, const int j0, const int stride) {
         if constexpr (!swz) {
             load_ldmatrix(t, tile_base + i0*stride + j0, stride);
             return;
         }
 #if defined(TURING_MMA_AVAILABLE)
+        static_assert(I == 16, "bad tile width");
+        static_assert(J ==  8, "bad tile height");
         const int i = i0 + threadIdx.x % t.I;
         const int j = j0 + (threadIdx.x / t.I) * (t.J / 2);
         int * xi = (int *) t.x;