5 Commits

Author	SHA1	Message	Date
biondizzle	ce4c4b6fcb	debug empty output	2026-05-14 22:13:32 +00:00
biondizzle	c016e66e23	Add CUDA sync + NaN/Inf check after each expert GEMM in grouped kernel	2026-05-14 11:27:58 +00:00
biondizzle	869151d211	Fix kernel.py: remove broken expand on scale factors (was expanding sf to weight size)	2026-05-14 10:36:16 +00:00
biondizzle	25cbc85afe	Replace kernel.py with thin wrapper around pre-compiled _C extension	2026-05-14 09:25:56 +00:00
biondizzle	f375c80bfe	feat: CUTLASS NVFP4 block-scaled GEMM kernel (native SM100 Blackwell) ... - Native NVFP4 block-scaled MMA using CUTLASS MainloopSm100TmaUmmaWarpSpecializedBlockScaled - Invokes mxf8f6f4.block_scale tensor core instructions (tcgen05.mma) - E2M1 (packed int8) + UE4M3 (float8_e4m3fn) block-16 scales → BF16 output - No dequantization: hardware block-scaled MMA avoids costly dequantize+BF16 path - PyTorch CUDA extension with CollectiveBuilder auto-deduction - Grouped expert GEMM for MoE dispatch (32 experts/rank, top-6 routing) - Integrated into nvfp4_mega_moe.py as primary path with TileLang fallback - Standalone C API (cutlass_nvfp4_gemm.cu) for direct B200 compilation - Build script, setup.py, and test script for B200 deployment Files: cutlass_nvfp4_gemm/ — Kernel source, PyTorch binding, build/test scripts nvfp4_mega_moe.py — Updated to use CUTLASS kernel when available	2026-05-13 23:11:15 +00:00