diff --git a/tests/unit/test_tmem_4warp_read.cu b/tests/unit/test_tmem_4warp_read.cu index 6b4333c5..954735ab 100644 --- a/tests/unit/test_tmem_4warp_read.cu +++ b/tests/unit/test_tmem_4warp_read.cu @@ -73,7 +73,7 @@ test_16x256b_loads(float* results) { { uint32_t u0, u1, u2, u3; asm volatile("tcgen05.ld.sync.aligned.16x256b.x1.b32 {%0, %1, %2, %3}, [%4];" - : "=r"(u0), "=r"(u1), "=r"(u2), "=r"(u3) : "r"(tb + 1)); + : "=r"(u0), "=r"(u1), "=r"(u2), "=r"(u3) : "r"(tb + 4)); // next 4 columns asm volatile("tcgen05.wait::ld.sync.aligned;"); load_count++; @@ -88,7 +88,7 @@ test_16x256b_loads(float* results) { { uint32_t u0, u1, u2, u3; asm volatile("tcgen05.ld.sync.aligned.16x256b.x1.b32 {%0, %1, %2, %3}, [%4];" - : "=r"(u0), "=r"(u1), "=r"(u2), "=r"(u3) : "r"(tb + 8)); + : "=r"(u0), "=r"(u1), "=r"(u2), "=r"(u3) : "r"(tb + 32)); // columns 8-11 asm volatile("tcgen05.wait::ld.sync.aligned;"); load_count++;