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

P6: Try shared::cluster instead of shared::cta for TMA store

Browse Source
This commit is contained in:
biondizzle
2026-05-30 17:05:27 +00:00
parent 829a5f93ce
commit 2de300e281
1 changed files with 2 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
dsv4/kernels/attention/fmha_tma.cuh
View File

@@ -213,9 +213,9 @@ __device__ __forceinline__ void tma_store_2d(
int coord_x,
int coord_y
) {
// cp.async.bulk.tensor.2d.global.shared::cta.bulk_group [tensorMap, {coord}], [srcMem];
// cp.async.bulk.tensor.2d.global.shared::cluster.bulk_group [tensorMap, {coord}], [srcMem];
asm volatile(
"cp.async.bulk.tensor.2d.global.shared::cta.bulk_group "
"cp.async.bulk.tensor.2d.global.shared::cluster.bulk_group "
"[%1, {%2, %3}], [%0];"
:: "r"(smem_src),
"l"(tma_desc),