fix: cvta.to.shared PTX for SMEM address

2026-05-28 05:09:50 +00:00
parent e12a81ae36
commit db17d8db9a
1 changed files with 3 additions and 1 deletions
--- a/dsv4/kernels/attention/fmha_sm100.cuh
+++ b/dsv4/kernels/attention/fmha_sm100.cuh
@@ -128,7 +128,9 @@ __device__ __forceinline__ uint64_t make_umma_smem_desc_bf16(
    uint64_t desc = 0;
    
    // start_address: bits [0,14), shifted left by 4 (16B-aligned)
-    uint32_t addr = static_cast<uint32_t>(__cvta_generic_to_shared(smem_ptr));
+    // Use __cvta_to_smem intrinsic or manual SMEM address computation
+    uint32_t addr = 0;
+    asm volatile("cvta.to.shared.u32 %0, %1;" : "=r"(addr) : "l"(smem_ptr));
    uint16_t start_addr = static_cast<uint16_t>(addr >> 4);
    
    // leading_byte_offset: bits [16,30), shifted left by 4