From 0fd3e12a521b87176bd5998f58a0de1d13e5c358 Mon Sep 17 00:00:00 2001
From: biondizzle <biondizzle@gmail.com>
Date: Thu, 28 May 2026 16:46:56 +0000
Subject: [PATCH] Fix TMA test: globalStrides in bytes not elements

---
 tests/unit/test_tma_driver.cu | 4 +++-
 1 file changed, 3 insertions(+), 1 deletion(-)

diff --git a/tests/unit/test_tma_driver.cu b/tests/unit/test_tma_driver.cu
index 5f49dca3..3a48a7a1 100644
--- a/tests/unit/test_tma_driver.cu
+++ b/tests/unit/test_tma_driver.cu
@@ -83,8 +83,10 @@ int main() {
     cudaMemcpy(d_src, h_src, DATA_BYTES, cudaMemcpyHostToDevice);
 
     // Create CUtensorMap
+    // globalDim: sizes in ELEMENTS (x=cols, y=rows)
+    // globalStrides: strides in BYTES
     uint64_t gdim[] = {(uint64_t)COLS, (uint64_t)ROWS};
-    uint64_t gstr[] = {1, (uint64_t)COLS};
+    uint64_t gstr[] = {2, (uint64_t)COLS * 2};  // BF16=2 bytes/element, row stride = COLS*2
     uint32_t tdim[] = {TILE_COLS, TILE_ROWS};
     uint32_t tstr[] = {1, TILE_COLS};