diff --git a/dsv4/kernels/attention/fmha_sm100.cuh b/dsv4/kernels/attention/fmha_sm100.cuh index 7d8cc947..20535464 100644 --- 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(__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(addr >> 4); // leading_byte_offset: bits [16,30), shifted left by 4