biondizzle/nvfp4-megamoe-kernel
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

test: try 16x256b.x1 with column step=4 (4 cols per read)

Browse Source
This commit is contained in:
biondizzle
2026-05-28 23:03:51 +00:00
parent 95003eced2
commit e1a708a187
1 changed files with 2 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
tests/unit/test_tmem_4warp_read.cu
View File

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