Public/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-09-20 01:31:31 +02:00
Code Issues Packages Projects Releases Wiki Activity

revert lds128 for wmma loading

Browse Source
This commit is contained in:
zhang hui
2025-11-25 17:06:40 +08:00
parent ae1c500aae
commit db9ae8b6b4
1 changed files with 2 additions and 7 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/mma.cuh
+2 -7
View File
@@ -441,13 +441,8 @@ namespace ggml_cuda_mma {
int64_t * xi = (int64_t *) t.x;
const int64_t * xs = (int64_t *) ((const int *) xs0 + (threadIdx.x % t.I) * stride + 2 * (threadIdx.x / t.I));
xi[0] = xs[0];
}else if constexpr (I == 16 && J == 8) {
int64_t * xi = (int64_t *) t.x;
const int64_t * xs = (int64_t *) ((const int *) xs0 + (threadIdx.x % t.I) * stride + 4 * (threadIdx.x / t.I));
xi[0] = xs[0];
const int64_t * xs1 = (int64_t *) ((const int *) xs0 + (threadIdx.x % t.I) * stride + 4 * (threadIdx.x / t.I) + 2);
xi[1] = xs1[0];
}else if constexpr (I == 16 && J == 8 && (std::is_same_v<T, half2> || std::is_same_v<T, nv_bfloat162>)) {
ggml_cuda_memcpy_1<sizeof(t.x)>(t.x, xs0 + t.get_i(0) * stride + t.get_j(0));
}else{
NO_DEVICE_CODE;
}