vbatts/llama.cpp
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

CUDA: fix misaligned shared memory read

Browse source
This commit is contained in:
Johannes Gäßler 2024-06-26 00:09:56 +02:00
parent 6fcbf68235
commit 9bf8ec5788
1 changed files with 1 additions and 1 deletions
Download patch file Download diff file Expand all files Collapse all files

2
ggml-cuda/mma.cuh
View file

@ -23,7 +23,7 @@ struct mma_int_A_I16K4 {
__device__ __forceinline__ void load(const int * __restrict__ xs0, const int & stride) { __device__ __forceinline__ void load(const int * __restrict__ xs0, const int & stride) {
#if defined(INT8_MMA_AVAILABLE) #if defined(INT8_MMA_AVAILABLE)
const int * xs = xs0 + (threadIdx.x%I)*stride + (threadIdx.x/I)*(K/2); const int * xs = xs0 + (threadIdx.x%I)*stride;
asm("ldmatrix.sync.aligned.m8n8.x2.b16 {%0, %1}, [%2];" asm("ldmatrix.sync.aligned.m8n8.x2.b16 {%0, %1}, [%2];"
: "+r"(x[0]), "+r"(x[1]) : "+r"(x[0]), "+r"(x[1])
: "l"(xs)); : "l"(xs));