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

FMHA SM100: Remove CUTLASS includes, write raw PTX inline asm

Browse Source 
CUTLASS headers transitively include cuda_bf16.h which has a CUDA 13.2
in_place_from bug. Writing tcgen05 PTX directly via inline asm instead.
No dependencies on CUTLASS C++ — pure PTX + CUDA runtime.
This commit is contained in:
biondizzle
2026-05-28 05:15:07 +00:00
parent 88d5995ec9
commit 48baea7728
1 changed files with 10 additions and 8 deletions
Download Patch File Download Diff File Expand all files Collapse all files

18
dsv4/kernels/attention/fmha_sm100.cuh
View File

@@ -60,14 +60,16 @@ typedef uint16_t bf16_t;
#endif
// CUTLASS C++ includes (CUDA device code only)
#if defined(__CUDA_ARCH__)
#include <cutlass/cutlass.h>
#include <cute/arch/mma_sm100_umma.hpp>
#include <cute/arch/copy_sm100.hpp>
#include <cute/arch/copy_sm100_tma.hpp>
#include <cute/arch/tmem_allocator_sm100.hpp>
#include <cute/arch/cluster_sm100.hpp>
#endif
// DISABLED: Transitive includes pull in cuda_bf16.h which has a CUDA 13.2 bug.
// We write tcgen05 PTX directly via inline asm instead.
// #if defined(__CUDA_ARCH__)
// #include <cutlass/cutlass.h>
// #include <cute/arch/mma_sm100_umma.hpp>
// #include <cute/arch/copy_sm100.hpp>
// #include <cute/arch/copy_sm100_tma.hpp>
// #include <cute/arch/tmem_allocator_sm100.hpp>
// #include <cute/arch/cluster_sm100.hpp>
// #endif
namespace dsv4 {
namespace kernels {