public_git/whisper.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/whisper.cpp.git synced 2026-09-30 11:36:38 +02:00
Code Issues Packages Projects Releases Wiki Activity

cuda: fix sm_70 tile compilation error (llama/29224)

Browse Source 
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)
This commit is contained in:
lingyezhixing
2026-09-23 20:46:47 +03:00
committed by Georgi Gerganov
parent 9a7d43df1e
commit 0e640a1801
1 changed files with 4 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/mma.cuh
+4 -2
View File
@@ -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;