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

CUDA: fix misaligned shared memory read (#8123)

Browse source
This commit is contained in:
Johannes Gäßler 2024-06-26 08:28:02 +02:00 • committed by GitHub
parent 494165f3b6
commit c8771ab5f8
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
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));