v0.1.6

[LeiWang1999]

What's Changed

• [Bugfix] Added missing thread offsets and other information to reduce by @LeiWang1999 in [Bugfix] Added missing thread offsets and other information to reduce #646
• [Bugfix] Adjust role assignment in warp specialization based on read access by @chengyupku in [Bugfix] Adjust role assignment in warp specialization based on read access #647
• Fix/jit kernel use target by @meinie0826 in Fix/jit kernel use target #648
• [Bugfix] Remove small array reuse condition in shared memory allocation merging by @LeiWang1999 in [Bugfix] Remove small array reuse condition in shared memory allocation merging #654
• [Enhancement] Add role assignment for AllocateNode in warp specialization by @chengyupku in [Enhancement] Add role assignment for AllocateNode in warp specialization #657
• [Bugfix][CI] Bug fixing and migrate CI from ada to hopper by @xwhzz in [Bugfix][CI] Bug fixing and migrate CI from ada to hopper #652
• [CI] Enable cache for virtual env and parallelize pytest via xdist by @LeiWang1999 in [CI] Enable cache for virtual env and parallelize pytest via xdist #660
• [Cache] Support shared cache directories for multiple process by @LeiWang1999 in [Cache] Support shared cache directories for multiple process #649
• [Enhancement] Add compile_flags parameter to JIT kernel and adapter classes for improved compilation control by @xwhzz in [Enhancement] Add compile_flags parameter to JIT kernel and adapter classes for improved compilation control #656
• add the support of rocm arch detecting by @zhangnju in add the support of rocm arch detecting #661
• [BugFix] Do not modify strict layout in common or relax level of layout inference. More conditions on layout checking by @tzj-fxz in [BugFix] Do not modify strict layout in common or relax level of layout inference. More conditions on layout checking #653
• [Bugfix][Docs] Update documentation build process and configurations for autoapi support by @xwhzz in [Bugfix][Docs] Update documentation build process and configurations for autoapi support #663
• [Enhancement] Improve buffer conflict detection in thread storage synchronization by @LeiWang1999 in [Enhancement] Improve buffer conflict detection in thread storage synchronization #658
• [Bugfix] Consider buffer data type into indices provably disjoint analysis by @LeiWang1999 in [Bugfix] Consider buffer data type into indices provably disjoint analysis #664
• [Bugfix] Remove redundant T.fill to fix precision issue by @xuchangtolearn in [Bugfix] Remove redundant T.fill to fix precision issue #667
• [Enhancement] Refactor buffer index handling for improved precision a… by @Alex4210987 in [Enhancement] Refactor buffer index handling for improved precision a… #671
• Reverts [Enhancement] Refactor buffer index handling for improved precision a… #671 by @LeiWang1999 in Reverts tile-ai/tilelang#671 #672
• [Bugfix] Passing correct nvcc to cmake by @chenyang78 in [Bugfix] Passing correct nvcc to cmake #670
• [CI] Improve format check output and automate commit of changes by @xwhzz in [CI] Improve format check output and automate commit of changes #669
• [Bugfix][CI] Use valid runner labels in workflow by @xwhzz in [Bugfix][CI] Use valid runner labels in workflow #674
• [Enhancement] passing verbose to LibraryGenerator by @chenyang78 in [Enhancement] passing verbose to LibraryGenerator #673
• [Enhancement] Enhance lint error messaging in CI by @xwhzz in [Enhancement] Enhance lint error messaging in CI #675
• Refactor to support upstream tvm by @Hzfengsy in Refactor to support upstream tvm #595
• Do not check for short variables by @oraluben in Do not check for short variables #676
• [Refactor] Phaseout version with commit id in editable model by @LeiWang1999 in [Refactor] Phaseout version with commit id in editable model #677
• [CI] Update CI workflow to use Python 3.12 by @LeiWang1999 in [CI] Update CI workflow to use Python 3.12 #679
• [Enhancement] Output cache-file-related messages with verbose=True by @chenyang78 in [Enhancement] Output cache-file-related messages with verbose=True #683
• [Enhancement] Enhance warp specialization logic by @chengyupku in [Enhancement] Enhance warp specialization logic #680
• Add Flash Attn example on amd mi300 series by @Alex4210987 in Add Flash Attn example on amd mi300 series #682
• [Enhancement] Refactored buffer detection logic in warp_specialized_rewriter.cc by @chengyupku in [Enhancement] Refactored buffer detection logic in warp_specialized_rewriter.cc #685
• [Fix] fix some issues with JIT decorators existing in the examples by @Cunxiao2002 in [Fix] fix some issues with JIT decorators existing in the examples #681
• [Enhancement] Add --ptxas-options=--register-usage-level=10 option by @LeiWang1999 in [Enhancement] Add --ptxas-options=--register-usage-level=10 option #684
• [Feature]:Add auto vectorize for atomic add by @yyttt6 in [Feature]:Add auto vectorize for atomic add #686
• [Refactor] Rebase pipeline injector from upstream tvm by @LeiWang1999 in [Refactor] Rebase pipeline injector from upstream tvm #687
• [Refactor] Introduce GemmInst for different targets handling by @LeiWang1999 in [Refactor] Introduce GemmInst for different targets handling #688
• [Enhancement] Optimize BF16 casting performance by @xwhzz in [Enhancement] Optimize BF16 casting performance #689
• [Smem Reuse] Optimize to do memory alignment on identical buffers. by @LeiWang1999 in [Smem Reuse] Optimize to do memory alignment on identical buffers. #693
• [Version] Keep local commit id as it somehow help with debugging by @LeiWang1999 in [Version] Keep local commit id as it somehow help with debugging #697
• [Example] Optimize warp specialize flashmla example by @LeiWang1999 in [Example] Optimize warp specialize flashmla example #698
• Bump transformers from 4.52.1 to 4.53.0 in /examples/bitnet-1.58b by @dependabot[bot] in Bump transformers from 4.52.1 to 4.53.0 in /examples/bitnet-1.58b #700
• Gated Delta Net(GDN) kernel implementation in TileLang by @tzj-fxz in Gated Delta Net(GDN) kernel implementation in TileLang #695
• Trivial update to calculate target arch by @oraluben in Trivial update to calculate target arch #702
• [CI] Remove Flash Attention dependency by @LeiWang1999 in [CI] Remove Flash Attention dependency #705
• [Layout] Introduce a new layout inference mechanism by @LeiWang1999 in [Layout] Introduce a new layout inference mechanism #699
• [Pipeline] Optimize inject software pipeline and pipeline planing pass by @LeiWang1999 in [Pipeline] Optimize inject software pipeline and pipeline planing pass #706
• Low-bit kernels fix and implementation by @tzj-fxz in Low-bit kernels fix and implementation #704
• [Feat] Support gemm with stride by @smallscientist1 in [Feat] Support gemm with stride #701
• [Enhancement] Add eviction policy support for TMA operations, enhance CUDA codegen, and introduce new pass config by @xwhzz in [Enhancement] Add eviction policy support for TMA operations, enhance CUDA codegen, and introduce new pass config #690
• [Enhancement] Enhance the robustness and generality of MLA examples by @Rachmanino in [Enhancement] Enhance the robustness and generality of MLA examples #709
• [Refactor] MergeAnnotations function to accept Map<Any, Any> instead of Map<String, Any> by @LeiWang1999 in [Refactor] MergeAnnotations function to accept Map<Any, Any> instead of Map<String, Any> #710
• [Pipeline] Phaseout fragment and double buffer info from pipeline pass by @LeiWang1999 in [Pipeline] Phaseout fragment and double buffer info from pipeline pass #711
• [Pipeline] Skip condition expression analysis for global reading by @LeiWang1999 in [Pipeline] Skip condition expression analysis for global reading #713
• [Index] Relocate Int64 Auto Promoter to ConfigBitWidth Pass, removing it from FlattenBuffer by @LeiWang1999 in [Index] Relocate Int64 Auto Promoter to ConfigBitWidth Pass, removing it from FlattenBuffer #714
• [CI] Bind build-test CI to NVIDIA as AMD runners are being introduced by @LeiWang1999 in [CI] Bind build-test CI to NVIDIA as AMD runners are being introduced #718
• fix: NVRTC backend by @lucifer1004 in fix: NVRTC backend #717
• [CUDA] Init support for sm_120 by @oraluben in [CUDA] Init support for sm_120 #716
• [Bugfix] Correct git configuration in docs CI by @xwhzz in [Bugfix] Correct git configuration in docs CI #720
• [Chore] fix typos by @lucifer1004 in [Chore] fix typos #719
• [CI][AMD] Add AMD GPU CI and fix some related bugs by @Alex4210987 in [CI][AMD] Add AMD GPU CI and fix some related bugs #694
• [Carver][Bugfix] Correct score function for warp tile selection in tensorcore policy by @NaOHCC in [Carver][Bugfix] Correct score function for warp tile selection in tensorcore policy #724
• [Refactor] Refactor CUDA code generation to simplify eviction policy handling by @LeiWang1999 in [Refactor] Refactor CUDA code generation to simplify eviction policy handling #721
• [Language] Introduce StridedTensor to support non contigious torch inputs by @LeiWang1999 in [Language] Introduce StridedTensor to support non contigious torch inputs #722
• [Enhancement][Bugfix] Fix bug in warp specialized pass and add gemm_sr fallback support for Hopper by @xwhzz in [Enhancement][Bugfix] Fix bug in warp specialized pass and add gemm_sr fallback support for Hopper #712
• 📝 Add docstrings to fix by @coderabbitai[bot] in 📝 Add docstrings to fix #726
• fix amd ci&add examples by @Alex4210987 in fix amd ci&add examples #729
• [Feature] Low-bit twiddling dequantization and FP4 GEMM by @tzj-fxz in [Feature] Low-bit twiddling dequantization and FP4 GEMM #725
• 📝 Add docstrings to mxfp4 by @coderabbitai[bot] in 📝 Add docstrings to mxfp4 #732
• [Refactor] Refactor env into a more flexible version by @LeiWang1999 in [Refactor] Refactor env into a more flexible version #740
• [Bugfix] Align stride index validation with torch in CythonKernelWrapper by @LeiWang1999 in [Bugfix] Align stride index validation with torch in CythonKernelWrapper #743
• [Bugfix]:Fix atomic add auto vectorize memory access out of bound error by @yyttt6 in [Bugfix]:Fix atomic add auto vectorize memory access out of bound error #742
• 📝 Add docstrings to main by @coderabbitai[bot] in 📝 Add docstrings to main #745
• [Refactor] Refactor barrier management by @LeiWang1999 in [Refactor] Refactor barrier management #744
• [Refactor] Merge bulk copy into copy and improve layout inference for bulk copy by @LeiWang1999 in [Refactor] Merge bulk copy into copy and improve layout inference for bulk copy #746
• [Refactor] Merge ThreadPartialSync and ThreadStorageSync by @LeiWang1999 in [Refactor] Merge ThreadPartialSync and ThreadStorageSync #741
• [Enhancement] Optimize loop body handling in IR by @chengyupku in [Enhancement] Optimize loop body handling in IR #749
• [MXFP4] Fix bugs and optimize exponential operation by @tzj-fxz in [MXFP4] Fix bugs and optimize exponential operation #750
• [Enhancement] Add DispatchInstruction specialization for fp8 types in gemm_sm90.h by @LeiWang1999 in [Enhancement] Add DispatchInstruction specialization for fp8 types in gemm_sm90.h #751
• [Enhancement] Add shape checking for reduce options by @kurisu6912 in [Enhancement] Add shape checking for reduce options #748
• [Bugfix] Add missing FP8 header include by @LeiWang1999 in [Bugfix] Add missing FP8 header include #752
• [MXFP4] Add bias to MXFP4 GEMM kernel by @tzj-fxz in [MXFP4] Add bias to MXFP4 GEMM kernel #753
• [Bugfix][WS] Consider loop min extent when computing phase id by @LeiWang1999 in [Bugfix][WS] Consider loop min extent when computing phase id #754
• [Typo] Remove disable_cache in some tests by @LeiWang1999 in [Typo] Remove disable_cache in some tests #755
• [README] Update GDN README for clarity and add acknowledgements by @chengyupku in [README] Update GDN README for clarity and add acknowledgements #758
• cutlass v4.2.0 supporting cuda 13 by @johnnynunez in cutlass v4.2.0 supporting cuda 13 #760
• [Feature] Add 1D TMA support by @tzj-fxz in [Feature] Add 1D TMA support #761
• [Example] Add vertical slash sparse attention pattern by @xwhzz in [Example] Add vertical slash sparse attention pattern #762
• [Bugfix] Address PassContext contamination from CI and fix incorrect rewrites in warp specialized pass by @xwhzz in [Bugfix] Address PassContext contamination from CI and fix incorrect rewrites in warp specialized pass #767
• [MXFP4] Add 1D TMA copy for Scale tensor in MXFP4 GEMM by @tzj-fxz in [MXFP4] Add 1D TMA copy for Scale tensor in MXFP4 GEMM #766
• [CUTLASS] hot fix blackwell by @johnnynunez in [CUTLASS] hot fix blackwell #768
• [Refactor] Refactor Operator into TileOperator and with tvm reflection by @LeiWang1999 in [Refactor] Refactor Operator into TileOperator and with tvm reflection #763
• [Reducer] Introduce alloc_reducer to separate inter and intra warp reduction by @LeiWang1999 in [Reducer] Introduce alloc_reducer to separate inter and intra warp reduction #757
• 📝 Add docstrings to pytile_0826 by @coderabbitai[bot] in 📝 Add docstrings to pytile_0826 #770
• [Bugfix]:Fix atomic add auto vectorize negative optimization by @yyttt6 in [Bugfix]:Fix atomic add auto vectorize negative optimization #765
• 📝 Add docstrings to reducer_0825 by @coderabbitai[bot] in 📝 Add docstrings to reducer_0825 #772
• Allow fill global buffer by @kurisu6912 in Allow fill global buffer #774
• [BugFix] Refactor the op check in LowerTileOp pass using the member function instead of string match by @tzj-fxz in [BugFix] Refactor the op check in LowerTileOp pass using the member function instead of string match #771
• [Enhancement] Add exp fallback for bf16 by @xwhzz in [Enhancement] Add exp fallback for bf16 #776
• [Lint] Introduce clang-tidy into format.sh by @LeiWang1999 in [Lint] Introduce clang-tidy into format.sh #777
• [Cache] Introduce detailed target information for the disk kernel cache by @LeiWang1999 in [Cache] Introduce detailed target information for the disk kernel cache #780
• [Example]Adds example for top-k operation by @Cunxiao2002 in [Example]Adds example for top-k operation #775
• [Math] Dispatch T.rsqrt(x) into cuda intrin instead of 1 / T.sqrt(x) by @LeiWang1999 in [Math] Dispatch T.rsqrt(x) into cuda intrin instead of 1 / T.sqrt(x) #781
• [CI] Adds pytest-durations for test timing by @Cunxiao2002 in [CI] Adds pytest-durations for test timing #782
• [Refactor] Support python reflection for tile operators by @LeiWang1999 in [Refactor] Support python reflection for tile operators #783
• fix amd tir&add examples by @Alex4210987 in fix amd tir&add examples #784
• [Nvidia][SM121] Add intrin.h include to gemm_mma.h for sm120+ by @HaoKang-Timmy in [Nvidia][SM121] Add intrin.h include to gemm_mma.h for sm120+ #785
• [Feat] Add tilelang T.assume support and assume injection for buffer shapes by @kurisu6912 in [Feat] Add tilelang T.assume support and assume injection for buffer shapes #787
• [Bugfix] Fix incorrect synchronization bug in minference example by @xwhzz in [Bugfix] Fix incorrect synchronization bug in minference example #786
• [AMD] fix bugs in warp shuffle by @txs19991 in [AMD] fix bugs in warp shuffle #790
• [AMD] fix mfma op interface by @Paran0idy in [AMD] fix mfma op interface #791
• [TMA] Automatically lower 1d tma in appropriate cases by @LeiWang1999 in [TMA] Automatically lower 1d tma in appropriate cases #788
• [CI]Adds pytest timeout to CI by @Cunxiao2002 in [CI]Adds pytest timeout to CI #792
• [Enhancement] Resolve reference cycle by @LeiWang1999 in [Enhancement] Resolve reference cycle #795
• [Bugfix] Fix index handling to promote 64-bit integers by @LeiWang1999 in [Bugfix] Fix index handling to promote 64-bit integers #796
• [AMD] support mfma i32_16x16x32_i8 by @Paran0idy in [AMD] support mfma i32_16x16x32_i8 #800
• [TileOp] Introduce a experimental python defined T.gemm_v2 by @LeiWang1999 in [TileOp] Introduce a experimental python defined T.gemm_v2 #793
• [Bugfix] Expose alloc_reducer definition to the python side by @LeiWang1999 in [Bugfix] Expose alloc_reducer definition to the python side #802
• [Refactor] Use new namespace and enhance dispatch macros for mma by @LeiWang1999 in [Refactor] Use new namespace and enhance dispatch macros for mma #801
• [AMD] support fp8 T.gemm by @txs19991 in [AMD] support fp8 T.gemm #804
• [AMD] support preshuffle weight mfma by @Paran0idy in [AMD] support preshuffle weight mfma #806
• Add pytest-durations to requirements for ROCm by @Alex4210987 in Add pytest-durations to requirements for ROCm #810
• Add ruff config to check for useless spaces by @oraluben in Add ruff config to check for useless spaces #807
• [Feature] Add ptx_cp_async_barrier_noinc intrinsic and related functionality by @chengyupku in [Feature] Add ptx_cp_async_barrier_noinc intrinsic and related functionality #809
• [Fix] Fix lower bug when buffer store is not guarded by any tile op by @kurisu6912 in [Fix] Fix lower bug when buffer store is not guarded by any tile op #794
• [feat] support gemm_sp for ampere arch by @botbw in [feat] support gemm_sp for ampere arch #691
• [Refactor] Update TVM subproject and refactor BlockNode handling in warp_specialized_rewriter.cc by @chengyupku in [Refactor] Update TVM subproject and refactor BlockNode handling in warp_specialized_rewriter.cc #812
• [Refactor] Reopen [Fix] Fix lower bug when buffer store is not guarded by any tile op #794 Fix lower bug when buffer store is not guarded by any tile op by @kurisu6912 in [Refactor] Reopen #794 Fix lower bug when buffer store is not guarded by any tile op #817
• [Refactor] Update TVM subproject and streamline buffer store handling by @chengyupku in [Refactor] Update TVM subproject and streamline buffer store handling #816
• [Example] add w4a8 gemm kernel by @Cunxiao2002 in [Example] add w4a8 gemm kernel #815
• [CI] fix rocm ci by @Cunxiao2002 in [CI] fix rocm ci #819
• [example] fix unused param in mhs example by @botbw in [example] fix unused param in mhs example #821
• [DSL] Support python tenary if then else expression by @LeiWang1999 in [DSL] Support python tenary if then else expression #822
• [Bugfix] Bug fix when git command is not installed by @LeiWang1999 in [Bugfix] Bug fix when git command is not installed #823
• [Bugfix] Skip fp4 dtype binding when using older versions of ml_dtypes by @LeiWang1999 in [Bugfix] Skip fp4 dtype binding when using older versions of ml_dtypes #824
• [Enhancement] Add a MXFP4 grouped GEMM example for FusedMoE by @Rachmanino in [Enhancement] Add a MXFP4 grouped GEMM example for FusedMoE #811
• [CMake] Added support for statically linked system libc library by @LeiWang1999 in [CMake] Added support for statically linked system libc library #825
• [Refactor] Refactor some build related configurations by @LeiWang1999 in [Refactor] Refactor some build related configurations #827
• [CI] Test Fix: Handle BufferLoad nodes when T.gemm input has a stride by @LeiWang1999 in [CI] Test Fix: Handle BufferLoad nodes when T.gemm input has a stride #843
• [Refactor] Turn off ENABLE_FAST_MATH by default by @LeiWang1999 in [Refactor] Turn off ENABLE_FAST_MATH by default #846
• [AMD] fix bf16x2 dtype codegen by @Paran0idy in [AMD] fix bf16x2 dtype codegen #847
• [Typing] Fallback from Python 3.10+ type syntax for compatibility by @LeiWang1999 in [Typing] Fallback from Python 3.10+ type syntax for compatibility #848
• [TIR] Refactor division simplification in RewriteSimplifier by @LeiWang1999 in [TIR] Refactor division simplification in RewriteSimplifier #849
• [Py38] Revert typing and parser updates for Python 3.8 compatibility by @LeiWang1999 in [Py38] Revert typing and parser updates for Python 3.8 compatibility #850
• [Bugfix] Disable Memory Info Analysis for local.var by @LeiWang1999 in [Bugfix] Disable Memory Info Analysis for local.var #851
• [Release] Bump Version to 0.1.6 by @LeiWang1999 in [Release] Bump Version to 0.1.6 #818

New Contributors

• @meinie0826 made their first contribution in Fix/jit kernel use target #648
• @zhangnju made their first contribution in add the support of rocm arch detecting #661
• @chenyang78 made their first contribution in [Bugfix] Passing correct nvcc to cmake #670
• @Hzfengsy made their first contribution in Refactor to support upstream tvm #595
• @coderabbitai[bot] made their first contribution in 📝 Add docstrings to fix #726
• @kurisu6912 made their first contribution in [Enhancement] Add shape checking for reduce options #748
• @johnnynunez made their first contribution in cutlass v4.2.0 supporting cuda 13 #760
• @HaoKang-Timmy made their first contribution in [Nvidia][SM121] Add intrin.h include to gemm_mma.h for sm120+ #785
• @txs19991 made their first contribution in [AMD] fix bugs in warp shuffle #790
• @Paran0idy made their first contribution in [AMD] fix mfma op interface #791

Full Changelog: https://github.com/tile-ai/tilelang/commits/0.1.6

---

This discussion was created from the release v0.1.6.