5 Commits

Author	SHA1	Message	Date
biondizzle	1c39e21d87	fix: remove broken L1 weight interleave ... The interleave assumed gate/up were pre-interleaved in groups of 16 and that we needed 2CTA UMMA layout. Both wrong: 1. vLLM w13_weight is plain concat [gate; up] along output dim 2. Our CUTLASS kernel uses ClusterShape 1x1x1, not 2CTA The interleave was shuffling weights into nonsense, making L1 GEMM compute the wrong thing, and chunk(2) would split wrong halves.	2026-05-14 13:05:45 +00:00
biondizzle	d3f35c9465	cleanup: remove abandoned TileLang and Mojo files ... - Deleted: layout.mojo, mega_moe.mojo, quantize.mojo (Mojo attempt) - Deleted: nvfp4_blockscaled_gemm.py, staging.py, nvfp4_mega_moe.py (TileLang top-level) - Deleted: tilelang_nvfp4_gemm.py, tilelang_kernels.py, nvfp4_dequant.py (TileLang package) - Deleted: src/weight_transform.py (duplicate of package version) - Fixed nvfp4_mega_moe.py: inlined unpack_ue4m3_u32, removed TileLang fallback imports - Fixed weight_transform.py: renamed function, removed TileLang alias, updated docs - Fixed __init__.py: removed TileLang alias, updated docstring - CUTLASS is the only kernel path now	2026-05-14 12:44:47 +00:00
biondizzle	ebc0ab0cac	Fix: keep scales as float8_e4m3fn, don't pack to uint32 (min_all_cuda unsupported)	2026-05-13 21:54:39 +00:00
biondizzle	1a452ffabd	Fix weight_transform signature to match nightly vLLM finalize_weights call	2026-05-13 21:41:43 +00:00
biondizzle	c2b752c2fe	Initial: TileLang NVFP4 mega_moe kernel package ... - nvfp4_mega_moe_full: drop-in replacement for deep_gemm.mega.fp8_nvfp4_mega_moe - transform_nvfp4_weights_for_mega_moe: weight transformation (tested) - SymmBuffer + get_symm_buffer_for_nvfp4_mega_moe: API-matching stubs - MEGA_MOE_STATIC=1 support for pipeline testing - pyproject.toml for pip install	2026-05-13 15:44:51 +00:00