From: Johannes Gäßler <redacted>
Date: Wed, 26 Jun 2024 06:28:02 +0000 (+0200)
Subject: CUDA: fix misaligned shared memory read (#8123)
X-Git-Tag: upstream/0.0.4488~1253
X-Git-Url: https://git.djapps.eu/?a=commitdiff_plain;h=c8771ab5f89387cdd7d9a8a69280dac46b45e02f;p=pkg%2Fggml%2Fsources%2Fllama.cpp

CUDA: fix misaligned shared memory read (#8123)
---

diff --git a/ggml-cuda/mma.cuh b/ggml-cuda/mma.cuh
index 0301a52f..5d87dd8e 100644
--- a/ggml-cuda/mma.cuh
+++ b/ggml-cuda/mma.cuh
@@ -23,7 +23,7 @@ struct mma_int_A_I16K4 {
 
     __device__ __forceinline__ void load(const int * __restrict__ xs0, const int & stride) {
 #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];"
             : "+r"(x[0]), "+r"(x[1])
             : "l"(xs));