From 829a5f93ce9408da58e681a3e2c9ba4575f8c2c5 Mon Sep 17 00:00:00 2001 From: biondizzle Date: Sat, 30 May 2026 17:04:38 +0000 Subject: [PATCH] =?UTF-8?q?P6:=20Fix=20TMA=20store=20PTX=20=E2=80=94=20rem?= =?UTF-8?q?ove=20.tile=20modifier,=20fix=20wait=5Fgroup=20syntax?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- dsv4/kernels/attention/fmha_tma.cuh | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/dsv4/kernels/attention/fmha_tma.cuh b/dsv4/kernels/attention/fmha_tma.cuh index 362ca91a..09ef2060 100644 --- a/dsv4/kernels/attention/fmha_tma.cuh +++ b/dsv4/kernels/attention/fmha_tma.cuh @@ -213,8 +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]; asm volatile( - "cp.async.bulk.tensor.2d.global.shared::cta.tile.bulk_group " + "cp.async.bulk.tensor.2d.global.shared::cta.bulk_group " "[%1, {%2, %3}], [%0];" :: "r"(smem_src), "l"(tma_desc), @@ -237,14 +238,14 @@ __device__ __forceinline__ void tma_store_commit_group() { * N=0 waits for all groups. N=1 waits for the most recent group. */ __device__ __forceinline__ void tma_store_wait_group(int n) { - asm volatile("cp.async.bulk.wait_group.read %0;" :: "r"(n) : "memory"); + asm volatile("cp.async.bulk.wait_group %0;" :: "r"(n) : "memory"); } /** * Wait for all pending TMA store groups to complete. */ __device__ __forceinline__ void tma_store_wait_all() { - asm volatile("cp.async.bulk.wait_group.read 0;" ::: "memory"); + asm volatile("cp.async.bulk.wait_group 0;" ::: "memory"); } } // namespace dsv4::kernels::attention