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

auto: pre-test commit

Browse Source
This commit is contained in:
biondizzle
2026-05-28 16:42:24 +00:00
parent a4211559cf
commit d9df1e6486
1 changed files with 9 additions and 7 deletions
Download Patch File Download Diff File Expand all files Collapse all files

16
tests/unit/test_tma_proper.cu
View File

@@ -101,16 +101,18 @@ int main() {
// Create CUtensorMap for a (ROWS, COLS) BF16 tensor
// Layout: (rows, cols) = (128, 16) in row-major
// TMA tile: (128, 16) — one tile covers the whole matrix
uint64_t gdim[] = {(uint64_t)COLS, (uint64_t)ROWS};
uint64_t gstr[] = {1, (uint64_t)COLS};
uint32_t tdim[] = {TILE_COLS, TILE_ROWS};
uint32_t tstr[] = {1, TILE_COLS};
CUtensorMap tma_desc_host;
CUresult res = cuTensorMapEncodeTiled(
&tma_desc_host,
CU_TENSOR_MAP_DATA_TYPE_UINT16, // BF16 = uint16
2, // 2D tensor
d_src, // global address
(uint64_t[]){COLS, ROWS}, // global dims (x=cols, y=rows)
(uint64_t[]){1, COLS}, // global strides (in elements)
(uint32_t[]){TILE_COLS, TILE_ROWS}, // tile dims
(uint32_t[]){1, TILE_COLS}, // tile strides (in elements)
CU_TENSOR_MAP_DATA_TYPE_UINT16,
2,
d_src,
gdim, gstr, tdim, tstr,
CU_TENSOR_MAP_INTERLEAVE_NONE,
CU_TENSOR_MAP_SWIZZLE_NONE, // No swizzle for now
CU_TENSOR_MAP_L2_PROMOTION_NONE,