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;