mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-24 07:04:00 +08:00
KurisandClaude Opus 4.6
f309d8147e
* [Bugfix][Subtype] Fix scalar fp4 load/store codegen and dynamic stride inference for sub-byte types The scalar fp4 buffer access path in codegen_cuda.cc used GetBufferRef which applies a /2 index division (mapping two fp4 elements to one byte), then performed a plain byte assignment — destroying the neighboring nibble. This caused incorrect results whenever fp4 stores couldn't be vectorized (e.g. with dynamic/non-constant strides in StridedTensor). Fix: use tl_fp4_packed_load/tl_fp4_packed_store with (fp4_e2_2_t*) cast for scalar fp4 accesses to non-packed buffers, which correctly handles nibble-level read-modify-write via set_x/set_y. Also add stride_scale factor to _process_dynamic_symbolic in the cython and tvm_ffi adapters, so torch strides (in storage units) are correctly converted to logical element strides for sub-byte dtypes. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> * [Test][Subtype] Add correctness tests for fp4 dynamic-stride store codegen Add regression test that scatters fp4 rows into a non-contiguous StridedTensor with dynamic strides and verifies byte-for-byte match against a static-stride reference. Parametrized over block_size. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> * [Enhancement][Subtype] Inject stride divisibility assume for sub-byte dtypes and add scalar fp4 store codegen test - Extend InjectAssumes pass to emit `assume(stride % pack_factor == 0)` for non-last-dimension strides of sub-byte buffers (e.g. fp4 with pack_factor=2). This provides the TIR analyzer with divisibility information before LetStmt inlining. - Add `test_subtype_fp4_scalar_store_codegen`: scatter fp4 elements via indirection to force the scalar (non-vectorized) store codegen path and verify nibble-level correctness. - Convert new tests to eager-style @tilelang.jit. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> * style: apply clang-format and ruff formatting Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> * fix lint error * fix: unpack 4-tuple in cython_wrapper dynamic_symbolic_map lookup --------- Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>