mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-22 06:50:25 +08:00
Tong WU
550e25d493
* Add TMA copy for packed global FP4 to unpacked SMEM. Wire IsValidTMACopyDtypePair and ALIGN16B tensor map descriptors for float4_e2m1fn global to float4_e2m1_unpacked shared transfers, and extend language tests for TMA and SIMT copy restrictions. * Extend tcgen05 MMA metadata for FP4 f8f6f4 and dtype helpers. * Refactor tcgen05_gemm_blockscaled API and lowering. * Use separate A/B dtypes in GEMM tileop emitters across backends. * Add maint/gemm correctness evaluation scripts and update gemm_sm100 examples. * Fix profiler assert_allclose to compare outputs only. * lint * canonicalize data types * add dpgm fp8-fp4 kernel * resolve comments * Refine FP4 unpacked SMEM dtype handling Treat float4_e2m1_unpacked as a SMEM/TMA-only storage tag instead of a general vector/register dtype. Remove vectorized Python aliases and keep codegen uses constrained to shared-memory storage/reference paths that lower to uint8_t. Allow packed FP4 global tensors to TMA-load into unpacked shared memory with the ALIGN16B TensorMap format, while keeping unpacked storage out of TMA store paths. Use payload-bit accounting for global strides and mbarrier transactions. Add focused dtype and TMA descriptor tests, and validate the real tcgen05 fp4-unpacked GEMM paths. * Restrict FP4 unpack TMA to loads * Hide vectorized unpacked FP4 language aliases * Update TVM for unpacked FP4 symbol exports * Rebind TVM submodule to merged unpacked FP4 export fix * Enable cache for persistent tcgen5 GEMM example