From 48baea7728a0b8a47b640b94f48f6552d2c7cff7 Mon Sep 17 00:00:00 2001
From: biondizzle <biondizzle@gmail.com>
Date: Thu, 28 May 2026 05:15:07 +0000
Subject: [PATCH] FMHA SM100: Remove CUTLASS includes, write raw PTX inline asm
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

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.
---
 dsv4/kernels/attention/fmha_sm100.cuh | 18 ++++++++++--------
 1 file changed, 10 insertions(+), 8 deletions(-)

diff --git a/dsv4/kernels/attention/fmha_sm100.cuh b/dsv4/kernels/attention/fmha_sm100.cuh
index acfc7a49..0c289c5a 100644
--- a/dsv4/kernels/attention/fmha_sm100.cuh
+++ b/dsv4/kernels/attention/fmha_sm100.cuh
@@ -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 {