From 2de300e281ac8c66895edd6f3e154f0144db08e3 Mon Sep 17 00:00:00 2001
From: biondizzle <biondizzle@gmail.com>
Date: Sat, 30 May 2026 17:05:27 +0000
Subject: [PATCH] P6: Try shared::cluster instead of shared::cta for TMA store

---
 dsv4/kernels/attention/fmha_tma.cuh | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/dsv4/kernels/attention/fmha_tma.cuh b/dsv4/kernels/attention/fmha_tma.cuh
index 09ef2060..cdd4da77 100644
--- a/dsv4/kernels/attention/fmha_tma.cuh
+++ b/dsv4/kernels/attention/fmha_tma.cuh
@@ -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),