From e1a708a187e2068fb1a8f32773453416a7a7564c Mon Sep 17 00:00:00 2001
From: biondizzle <biondizzle@gmail.com>
Date: Thu, 28 May 2026 23:03:51 +0000
Subject: [PATCH] test: try 16x256b.x1 with column step=4 (4 cols per read)

---
 tests/unit/test_tmem_4warp_read.cu | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

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++;