We read every piece of feedback, and take your input very seriously.
To see all available qualifiers, see our documentation.
There was an error while loading. Please reload this page.
1 parent 494165f commit c8771abCopy full SHA for c8771ab
ggml-cuda/mma.cuh
@@ -23,7 +23,7 @@ struct mma_int_A_I16K4 {
23
24
__device__ __forceinline__ void load(const int * __restrict__ xs0, const int & stride) {
25
#if defined(INT8_MMA_AVAILABLE)
26
- const int * xs = xs0 + (threadIdx.x%I)*stride + (threadIdx.x/I)*(K/2);
+ const int * xs = xs0 + (threadIdx.x%I)*stride;
27
asm("ldmatrix.sync.aligned.m8n8.x2.b16 {%0, %1}, [%2];"
28
: "+r"(x[0]), "+r"(x[1])
29
: "l"(xs));
0 commit comments