v0.1.7 · tile-ai tilelang · Discussion #1383 · GitHub

LeiWang1999
Dec 7, 2025
Maintainer

What's Changed

[PATCH] Static libg++ linking fix by @LeiWang1999 in [PATCH] Static libg++ linking fix #854
[Analyzer] Enhance ConstIntBoundAnalyzer and IntervalSet with modular set analysis by @LeiWang1999 in [Analyzer] Enhance ConstIntBoundAnalyzer and IntervalSet with modular set analysis #856
[Doc] Optimize the quickstart guide for clarity and not just for CUDA by @LeiWang1999 in [Doc] Optimize the quickstart guide for clarity and not just for CUDA #858
[TMA] Bugfix when a shared buffer is both issued with tma store and tma load by @LeiWang1999 in [TMA] Bugfix when a shared buffer is both issued with tma store and tma load #857
[AMD][MLA] Fix mla autotune for rocm by @LeiWang1999 in [AMD][MLA] Fix mla autotune for rocm #861
[Bugfix] Ensure correct handling for cases where seq_q<seq_kv in flash attention examples by @Rachmanino in [Bugfix] Ensure correct handling for cases where seq_q<seq_kv in flash attention examples #864
[AMD] refactor MatrixCoreIntrinEmitter by @Paran0idy in [AMD] refactor MatrixCoreIntrinEmitter #860
[Feat] Add fast sine and cosine definitions in CUDA templates by @Rachmanino in [Feat] Add fast sine and cosine definitions in CUDA templates #865
[Layout] Support layout forward with multi dimension by @LeiWang1999 in [Layout] Support layout forward with multi dimension #867
[Autotune][Conv] optimize convolution examples to use autotune by @LeiWang1999 in [Autotune][Conv] optimize convolution examples to use autotune #866
[Example] Add examples to support efficient attention sink forward process by @Rachmanino in [Example] Add examples to support efficient attention sink forward process #853
[Parser] Adapt Parser to work with Python 3.8 in some cases by @LeiWang1999 in [Parser] Adapt Parser to work with Python 3.8 in some cases #869
[Fix] Fix bug 0905: tilelang doesn't vectorize B[i,j] = c[i] + A[i,j] by @kurisu6912 in [Fix] Fix bug 0905: tilelang doesn't vectorize B[i,j] = c[i] + A[i,j] #798
[Language] Support sequence comparisons by @LeiWang1999 in [Language] Support sequence comparisons #872
[Language] Support loop_break primitive by @chengyupku in [Language] Support loop_break primitive #873
[Bugfix] Use ExprDeepEqual instead of StructuralEqual when merge consecutive If stmt by @LeiWang1999 in [Bugfix] Use ExprDeepEqual instead of StructuralEqual when merge consecutive If stmt #876
[Language] Support atomic add with ret by @LeiWang1999 in [Language] Support atomic add with ret #870
[Cython] Remove an incorrect check by @LJC00118 in [Cython] Remove an incorrect check #880
Update amd_ci.yml by @Alex4210987 in Update amd_ci.yml #881
[FastMath] Disable default TVM fastmath intrinsic dispatch and add explicit fastmath op to invoke by @LeiWang1999 in [FastMath] Disable default TVM fastmath intrinsic dispatch and add explicit fastmath op to invoke #875
[Example] Add efficient attention sink backward implementations and tests by @Rachmanino in [Example] Add efficient attention sink backward implementations and tests #877
[Precision] Introduce T.ieee_rsqrt and related high precision op by @LeiWang1999 in [Precision] Introduce T.ieee_rsqrt and related high precision op #882
[Dist] Provide an option to include commit ID in version by @LeiWang1999 in [Dist] Provide an option to include commit ID in version #884
[Example] Optimize sink attention forward via swizzled layout and report benchmark results by @Rachmanino in [Example] Optimize sink attention forward via swizzled layout and report benchmark results #885
[Layout] Introduce Flexible Parallel to Support T.serial and local buffers inside T.Parallel loop by @LeiWang1999 in [Layout] Introduce Flexible Parallel to Support T.serial and local buffers inside T.Parallel loop #844
[Bugfix][Enhancement] Fix a bug in previous commit and enhance cuda backend by @Hamerlate in [Bugfix][Enhancement] Fix a bug in previous commit and enhance cuda backend #887
[Bugfix] Fix CopyNode Lower method to include disable_tma flag in GetCopyInst by @Rachmanino in [Bugfix] Fix CopyNode Lower method to include disable_tma flag in GetCopyInst #888
[Layout] Fix plot layout by @Paran0idy in [Layout] Fix plot layout #890
[Example] Add example by @LeiWang1999 in [Example] Add example #894
[News] Add announcement of support for Huawei Ascend chips by @xwhzz in [News] Add announcement of support for Huawei Ascend chips #895
[Example] Add sparse mla examples by @LeiWang1999 in [Example] Add sparse mla examples #896
[Typo] Fix backend name for Huawei Ascend by @xwhzz in [Typo] Fix backend name for Huawei Ascend #898
[CI] Legalize math related test by @LeiWang1999 in [CI] Legalize math related test #899
[Bugfix] Fix flops comp and softmax scale in mla by @Edenzzzz in [Bugfix] Fix flops comp and softmax scale in mla #900
[Example] Specify a fixed commit for the flash-linear-attention repository and optimize nsa examples by @LeiWang1999 in [Example] Specify a fixed commit for the flash-linear-attention repository and optimize nsa examples #913
[CI] optimize CI time for sparse gemm by @botbw in [CI] optimize CI time for sparse gemm #906
[Enhancement] Include compile flags into the hash key of cached kernels by @Rachmanino in [Enhancement] Include compile flags into the hash key of cached kernels #911
[Bugfix] Fix saving kernel source code where JITKernel.artifact is None by @zjudmd1015 in [Bugfix] Fix saving kernel source code where JITKernel.artifact is None #921
[CI] Refactor import paths in dequantization examples to use dequantize_utils by @LeiWang1999 in [CI] Refactor import paths in dequantization examples to use dequantize_utils #914
[Example] Add MLA decode ws example by @chengyupku in [Example] Add MLA decode ws example #928
[CI] Fix documentation runner by adding 'nvidia' tag by @xwhzz in [CI] Fix documentation runner by adding 'nvidia' tag #927
[Layout] Strict annotate completed replicated layout for fragment with constant index by @LeiWang1999 in [Layout] Strict annotate completed replicated layout for fragment with constant index #929
[Bugfix] Fix tensor memory copy layout by @Hamerlate in [Bugfix] Fix tensor memory copy layout #933
[Example] Optimize online_softmax example by @lijinpei in [Example] Optimize online_softmax example #934
[Example] Add correctness assert into dsa example by @LeiWang1999 in [Example] Add correctness assert into dsa example #937
[Enhancement] Enhance and add new GQA backward examples for Hopper by @Rachmanino in [Enhancement] Enhance and add new GQA backward examples for Hopper #930
[Enhancement] Fix lint to improve grouped GEMM performance with TMA by @Cunxiao2002 in [Enhancement] Fix lint to improve grouped GEMM performance with TMA #938
[Example] Introduce split+sum template, and optimize atomic_add performance for bwd examples by @LeiWang1999 in [Example] Introduce split+sum template, and optimize atomic_add performance for bwd examples #940
[Example] Disable TMA and enable FastMath for NSA Examples ([Example] Disable TMA and enable FastMath for NSA Examples (#941) #941) by @LeiWang1999 in [Example] Disable TMA and enable FastMath for NSA Examples (#941) #941
[Example] Revert the atomic/split&sum templates in MHA backward examples by @Rachmanino in [Example] Revert the atomic/split&sum templates in MHA backward examples #943
[Example] Add sparse mla bwd example for deepseek_v32 by @Zhichenzzz in [Example] Add sparse mla bwd example for deepseek_v32 #919
[Profiler]Adds CUPTI profiler support by @Cunxiao2002 in [Profiler]Adds CUPTI profiler support #936
[Enhancement] Support Copy for Buffer Load witih scalar indices by @LeiWang1999 in [Enhancement] Support Copy for Buffer Load witih scalar indices #946
[Code Style] Refine nvrtc compile related check style by @BBuf in [Code Style] Refine nvrtc compile related check style #945
[Backend] Add metal backend by @oraluben in [Backend] Add metal backend #799
[CI] enable dependabot for GHA workflows by @XuehaiPan in [CI] enable dependabot for GHA workflows #950
Modify the SM architecture number to support Thor’s sm110. by @iloveai8086 in Modify the SM architecture number to support Thor’s sm110. #957
[CI] auto-cancel in-progress PR CI when new commits are pushed by @XuehaiPan in [CI] auto-cancel in-progress PR CI when new commits are pushed #956
[bug] fix type object is not subscriptable in py38 by @BBuf in [bug] fix type object is not subscriptable in py38 #959
[Bugfix][Doc] Add astroid version constraint to requirements.txt by @xwhzz in [Bugfix][Doc] Add astroid version constraint to requirements.txt #958
[CI]: Bump actions/setup-python from 2 to 6 by @dependabot[bot] in [CI]: Bump actions/setup-python from 2 to 6 #951
[CI]: Bump astral-sh/setup-uv from 6 to 7 by @dependabot[bot] in [CI]: Bump astral-sh/setup-uv from 6 to 7 #952
[CI]: Bump actions/github-script from 7 to 8 by @dependabot[bot] in [CI]: Bump actions/github-script from 7 to 8 #954
[CI]: Bump actions/checkout from 2 to 5 by @dependabot[bot] in [CI]: Bump actions/checkout from 2 to 5 #953
[TileOp] Implement WGMMA for T.gemm_v2 by @LeiWang1999 in [TileOp] Implement WGMMA for T.gemm_v2 #813
[Docs] add CODE_OF_CONDUCT.md by @XuehaiPan in [Docs] add CODE_OF_CONDUCT.md #965
[Example] Add support for bfloat16 and user-defined sm_scale in attention sink examples by @Rachmanino in [Example] Add support for bfloat16 and user-defined sm_scale in attention sink examples #924
[Bugfix] Do not force inline let stmt by @LeiWang1999 in [Bugfix] Do not force inline let stmt #947
[CI] add pre-commit integration by @XuehaiPan in [CI] add pre-commit integration #955
[Doc] Install docs add docker install method by @BBuf in [Doc] Install docs add docker install method #961
[Bugfix] Fix dummy kernel compliation by @SiriusNEO in [Bugfix] Fix dummy kernel compliation #962
[CI][Refactor] Refactor non-test CI workflow files by @XuehaiPan in [CI][Refactor] Refactor non-test CI workflow files #971
[TileOp] Implememt CumSum1D by @LeiWang1999 in [TileOp] Implememt CumSum1D #978
[Language] Enhance T.alloc_var for AugAssign and AnnAsign by @LeiWang1999 in [Language] Enhance T.alloc_var for AugAssign and AnnAsign #979
[Refactor] Refactor Pass InjectFenceProxy and expose some warp group primitives in frontend by @LeiWang1999 in [Refactor] Refactor Pass InjectFenceProxy and expose some warp group primitives in frontend #977
[Typo] Remove debug print by @LeiWang1999 in [Typo] Remove debug print #980
[Bugfix] Use access_ptr("r") instead of access_ptr("w") for correct pipeline analysis by @LeiWang1999 in [Bugfix] Use access_ptr("r") instead of access_ptr("w") for correct pipeline analysis #983
[Feature][Example] Support TMA reduce operation and update GQA bwd example by @chengyupku in [Feature][Example] Support TMA reduce operation and update GQA bwd example #969
[Bugfix] Add NVIDIA HPC SDK support in CUDA detection (Add support for NVIDIA HPC SDK (nvhpc) installation paths #974) by @Degeneracy-Evil in [Bugfix] Add NVIDIA HPC SDK support in CUDA detection (#974) #976
[BugFix] Robust gemm policy for sparse_mla_fwd in Hopper and Ada Lovelace architectures by @tzj-fxz in [BugFix] Robust gemm policy for sparse_mla_fwd in Hopper and Ada Lovelace architectures #984
[Bugfix] Fallback torch.accelerator.synchronize() to torch.cuda.synchronize() by @yyttt6 in [Bugfix] Fallback torch.accelerator.synchronize() to torch.cuda.synchronize() #987
[Bugfix]:Fix atomicadd auto vectorize identify var error by @yyttt6 in [Bugfix]:Fix atomicadd auto vectorize identify var error #883
[CI] Speed up sparse tensor core test via vectorized generating sparse data by @LeiWang1999 in [CI] Speed up sparse tensor core test via vectorized generating sparse data #1009
[Build] Migrate to scikit-build-core by @oraluben in [Build] Migrate to scikit-build-core #939
[CI] Removes redundant environment variable by @Cunxiao2002 in [CI] Removes redundant environment variable #1020
[Transform] Migrate LowerIntrin from tvm into tilelang by @LeiWang1999 in [Transform] Migrate LowerIntrin from tvm into tilelang #999
[Lint] Prefer American English spelling by @XuehaiPan in [Lint] Prefer American English spelling #1022
[Build] Prefer libs from local build dir by @oraluben in [Build] Prefer libs from local build dir #1027
[Language] Support Consequential assignments like 'a = b = c = 1' by @LeiWang1999 in [Language] Support Consequential assignments like 'a = b = c = 1' #992
[CI] Removes debug print statements from the example. by @Cunxiao2002 in [CI] Removes debug print statements from the example. #1030
[Enhancement] Update abs function for half_t and bfloat_t to use cutlass implementation by @Rachmanino in [Enhancement] Update abs function for half_t and bfloat_t to use cutlass implementation #1023
[Bugfix] Recover code for flexible parallel by @LeiWang1999 in [Bugfix] Recover code for flexible parallel #1032
[CI] Disable buggy(maybe) warp specialized kernel ci test for H20 by @LeiWang1999 in [CI] Disable buggy(maybe) warp specialized kernel ci test for H20 #1033
[TIR] Revert some changes of Pass LowerIntrin by @LeiWang1999 in [TIR] Revert some changes of Pass LowerIntrin #1035
[Env] Optimize the mechanism for locating TL_LIBS by @LeiWang1999 in [Env] Optimize the mechanism for locating TL_LIBS #1038
[CUDA] Add pack functions for FP8 types by @LJC00118 in [CUDA] Add pack functions for FP8 types #967
[Language] Expose T.get_warp_idx_sync and T.shuffle_elect for efficient thread election by @LeiWang1999 in [Language] Expose T.get_warp_idx_sync and T.shuffle_elect for efficient thread election #989
[AMD] fix bug&add amd fp8 examples by @Alex4210987 in [AMD] fix bug&add amd fp8 examples #966
[CI][Refactor] Merge test CI workflow files into one by @XuehaiPan in [CI][Refactor] Merge test CI workflow files into one #973
[BugFix] Phaseout dependency of Triton in sink examples to make CI happy by @Rachmanino in [BugFix] Phaseout dependency of Triton in sink examples to make CI happy #1045
[Refactor] Use has_simt_copy to decide whether to insert set_max_nreg by @chengyupku in [Refactor] Use has_simt_copy to decide whether to insert set_max_nreg #982
[Feature]: Add test for atomicadd auto vectorize and remove useless code by @yyttt6 in [Feature]: Add test for atomicadd auto vectorize and remove useless code #1019
Allow mma gemm for all cuda arch by @oraluben in Allow mma gemm for all cuda arch #1047
[Bugfix] Improves compatibility when checking for MPS availability in different PyTorch builds. by @LeiWang1999 in [Bugfix] Improves compatibility when checking for MPS availability in different PyTorch builds. #1051
[CI] Fix ROCm CI by @XuehaiPan in [CI] Fix ROCm CI #1043
[Enhancement] Add support for symbolic dimensions in Cython kernel adapter and improve static shape validation in wrapper by @Rachmanino in [Enhancement] Add support for symbolic dimensions in Cython kernel adapter and improve static shape validation in wrapper #1024
Automatically initialize submodule if missing by @LeiWang1999 in Automatically initialize submodule if missing #1052
[Enhancement] Remove constraint requiring last dimension stride to be 1 by @LJC00118 in [Enhancement] Remove constraint requiring last dimension stride to be 1 #1040
[CI] Disable autofix for pre-commit CI by @LeiWang1999 in [CI] Disable autofix for pre-commit CI #1053
[Enhancement] Improve CUDA compiler detection in CMake by @LJC00118 in [Enhancement] Improve CUDA compiler detection in CMake #1054
[Enhancement] Introduce a workaround for layout inference for local buffer store by @LeiWang1999 in [Enhancement] Introduce a workaround for layout inference for local buffer store #1055
[Refactor] Refactor Pass LegalizeSafeMemoryAccess to support recursive load/store rewrite by @SiriusNEO in [Refactor] Refactor Pass LegalizeSafeMemoryAccess to support recursive load/store rewrite #1050
Making version parser more robust against missing or unavailable metadata by @LeiWang1999 in Making version parser more robust against missing or unavailable metadata #1061
[DOC] Add document for develop with PYTHONPATH by @LeiWang1999 in [DOC] Add document for develop with PYTHONPATH #1062
[CI]:Reduce test shapes to avoid OOM errors during CI. by @yyttt6 in [CI]:Reduce test shapes to avoid OOM errors during CI. #1060
[Benchmark] Add H800 SXM Benchmark results by @LeiWang1999 in [Benchmark] Add H800 SXM Benchmark results #1063
[Misc] Add GitHub issue templates by @XuehaiPan in [Misc] Add GitHub issue templates #1057
[Refactor][Example] Update linear attention examples and add tests by @Rachmanino in [Refactor][Example] Update linear attention examples and add tests #1010
[Enhancement] Deprecate split&sum in attn bwd examples on Hopper by @Rachmanino in [Enhancement] Deprecate split&sum in attn bwd examples on Hopper #1065
[Benchmark] Add matmul FP16 benchmark results by @LeiWang1999 in [Benchmark] Add matmul FP16 benchmark results #1067
[CI]: Bump actions/checkout from 4 to 5 by @dependabot[bot] in [CI]: Bump actions/checkout from 4 to 5 #1070
[Example] Update GQA varlen fwd and MHA varlen fwd by @chengyupku in [Example] Update GQA varlen fwd and MHA varlen fwd #1071
[Parallel] Support T.Parallel with dynamic extents by @LeiWang1999 in [Parallel] Support T.Parallel with dynamic extents #990
[Layout] Utilizing IsEqual instead of StructuralEqual by @LeiWang1999 in [Layout] Utilizing IsEqual instead of StructuralEqual #1073
[Cache] raise errors for tileang.clear_cache() by @LeiWang1999 in [Cache] raise errors for tileang.clear_cache() #1077
[Feature] Support Reduce operators for bitwise and/or/xor by @tzj-fxz in [Feature] Support Reduce operators for bitwise and/or/xor #1074
[Autotune] Add autotune coverage for symbolic M and normalize cache key by @LeiWang1999 in [Autotune] Add autotune coverage for symbolic M and normalize cache key #1075
[Language] Recommend using T.dynamic instead of T.symbolic by @LeiWang1999 in [Language] Recommend using T.dynamic instead of T.symbolic #1076
[Language] Efficient T.reduce_ with shared memory input/output by @LeiWang1999 in [Language] Efficient T.reduce_ with shared memory input/output #1080
[Bugfix] Fix missing reg alloc in custom warp specialization by @chengyupku in [Bugfix] Fix missing reg alloc in custom warp specialization #1084
[Enhancement] Update async intrinsic handling in inject_fence_proxy by @Rachmanino in [Enhancement] Update async intrinsic handling in inject_fence_proxy #1068
[Feature] Add GQA backward kernel with varlen input by @tzj-fxz in [Feature] Add GQA backward kernel with varlen input #1082
[BugFix] Add memory order argument for non-vectorized atomic add by @tzj-fxz in [BugFix] Add memory order argument for non-vectorized atomic add #1081
[Refactor] Rename cython output to tilelang_cython and relocate its path by @LeiWang1999 in [Refactor] Rename cython output to tilelang_cython and relocate its path #1086
[Target] Enhance target selection helpers and documentation by @LeiWang1999 in [Target] Enhance target selection helpers and documentation #1085
[Cleanup] Remove tilelang.disable_cache() calls from examples and tests by @Rachmanino in [Cleanup] Remove tilelang.disable_cache() calls from examples and tests #1088
[PassConfig] Introduce PassConfig TL_STORAGE_REWRITE_DETECT_INPLACE by @LeiWang1999 in [PassConfig] Introduce PassConfig TL_STORAGE_REWRITE_DETECT_INPLACE #1089
[Language] Support tilelang alloc_var(dtype, init=x) by @LeiWang1999 in [Language] Support tilelang alloc_var(dtype, init=x) #1092
[Bugfix] Fix missing host cuTensorMapEncodeIm2col call by @chengyupku in [Bugfix] Fix missing host cuTensorMapEncodeIm2col call #1094
[GQA] Add regional atomic add to slightly boost performance by @tzj-fxz in [GQA] Add regional atomic add to slightly boost performance #1093
[Example] Add block level high performance gemv example by @LeiWang1999 in [Example] Add block level high performance gemv example #1097
[Refactor] Optimize debug message for parallel inference by @LeiWang1999 in [Refactor] Optimize debug message for parallel inference #1096
[CI][Lint] Retire format.sh and add clang-tidy to GHA workflow by @XuehaiPan in [CI][Lint] Retire format.sh and add clang-tidy to GHA workflow #1044
[Refactor] Use forceinline in ldmatrix and update mamba scan kernel by @chengyupku in [Refactor] Use forceinline in ldmatrix and update mamba scan kernel #1104
[Maint] Update uncommitted change detection command in format.sh by @XuehaiPan in [Maint] Update uncommitted change detection command in format.sh #1102
[Benchmark] Add Mamba2_chunk_scan benchmark by @chengyupku in [Benchmark] Add Mamba2_chunk_scan benchmark #1109
[Benchmark] Update Mamba2_chunk_scan benchmark by @chengyupku in [Benchmark] Update Mamba2_chunk_scan benchmark #1110
[Lint] Enable pyupgrade linter in ruff by @oraluben in [Lint] Enable pyupgrade linter in ruff #963
[Refactor] Improve scalar handling in CopyNode and update loop partition dtype logi by @LeiWang1999 in [Refactor] Improve scalar handling in CopyNode and update loop partition dtype logi #1111
[Feature] Enhance vectorized conversion support in CUDA codegen by @Rachmanino in [Feature] Enhance vectorized conversion support in CUDA codegen #1095
[Feature] Support None type as input for T.ptr and T.Tensor by @xwhzz in [Feature] Support None type as input for T.ptr and T.Tensor #1114
[Bugfix] Resolve mixed stride dtype issue (inconsistent int32/int64 values) by @LeiWang1999 in [Bugfix] Resolve mixed stride dtype issue (inconsistent int32/int64 values) #1119
[Feature] Add memory_order PTX for vectorized atomic add by @tzj-fxz in [Feature] Add memory_order PTX for vectorized atomic add #1112
[CI]: Bump actions/upload-artifact from 4 to 5 by @dependabot[bot] in [CI]: Bump actions/upload-artifact from 4 to 5 #1128
[CI]: Bump actions/download-artifact from 5 to 6 by @dependabot[bot] in [CI]: Bump actions/download-artifact from 5 to 6 #1127
[Enhancement] Add missing fence_barrier_init primitive after mbarrier init by @chengyupku in [Enhancement] Add missing fence_barrier_init primitive after mbarrier init #1121
[Feature]:Add device assert by @yyttt6 in [Feature]:Add device assert #1116
[Build][CI] Build and test SDist in release CI by @XuehaiPan in [Build][CI] Build and test SDist in release CI #1098
[Benchmark] Update triton and helion baselines in mamba-chuk-scan by @chengyupku in [Benchmark] Update triton and helion baselines in mamba-chuk-scan #1131
Add int2 and longlong4 pack functions by @LJC00118 in Add int2 and longlong4 pack functions #1129
[BugFix] Add memory order and testing script for split version GQA bwd kernel by @tzj-fxz in [BugFix] Add memory order and testing script for split version GQA bwd kernel #1100
[Bugfix] Correctly construct the argument list for atomic add based on the vector size by @LeiWang1999 in [Bugfix] Correctly construct the argument list for atomic add based on the vector size #1137
[AMD] Supoort T.gemm_v2 for AMD Backend by @Paran0idy in [AMD] Supoort T.gemm_v2 for AMD Backend #1136
[BugFix] alloc_var init failed to handle complex expression by @kurisu6912 in [BugFix] alloc_var init failed to handle complex expression #1144
[Refactor] Remove amd gemm_v2 tests by @LeiWang1999 in [Refactor] Remove amd gemm_v2 tests #1149
[BugFix] Implement bfloat16 support in CUDA code generation with min/max functions and inf/nan values by @Rachmanino in [BugFix] Implement bfloat16 support in CUDA code generation with min/max functions and inf/nan values #1143
[Bugfix] Implement classic arena algorithm for shmem merge and WAW conflict detection by @LeiWang1999 in [Bugfix] Implement classic arena algorithm for shmem merge and WAW conflict detection #1146
[CI] allow dirty workspace for format.sh and introduce loop carry thread sync unit test by @LeiWang1999 in [CI] allow dirty workspace for format.sh and introduce loop carry thread sync unit test #1153
[CI] use Python urllib to download file instead of Wget by @XuehaiPan in [CI] use Python urllib to download file instead of Wget #1154
[BugFix] Correct direct copy from bf16 to fp8 by @Cunxiao2002 in [BugFix] Correct direct copy from bf16 to fp8 #1090
[Refactor]:Move device_assert from extern_call to intrin_call by @yyttt6 in [Refactor]:Move device_assert from extern_call to intrin_call #1134
[Enhancement] Enhance Cast operations Vectorization by @LJC00118 in [Enhancement] Enhance Cast operations Vectorization #1156
[Bugfix] Enhance LetStmt handling in Vectorize Loop Pass by @LeiWang1999 in [Bugfix] Enhance LetStmt handling in Vectorize Loop Pass #1159
[Release] Bump version to v0.1.6.post2 by @LeiWang1999 in [Release] Bump version to v0.1.6.post2 #1160
[FFI] Rebase tvm to v0.22.0 to utilize tvm-ffi by @LeiWang1999 in [FFI] Rebase tvm to v0.22.0 to utilize tvm-ffi #1108
[Bugfix] Enable code lowering with producer‑copy‑only program by @LeiWang1999 in [Bugfix] Enable code lowering with producer‑copy‑only program #1168
[Bugfix] Support 16bits shfl_sync by @LeiWang1999 in [Bugfix] Support 16bits shfl_sync #1169
[Testing] Move TMA 1D and test for its functionality by @tzj-fxz in [Testing] Move TMA 1D and test for its functionality #1167
[Refactor]: Change the params in pytest to avoid oom error during ci by @yyttt6 in [Refactor]: Change the params in pytest to avoid oom error during ci #1170
[Bugfix] Fix tvm import path for editable build by @LeiWang1999 in [Bugfix] Fix tvm import path for editable build #1172
[Language] Expose T.warpgroup_fence_operand for nvcc code motion by @LeiWang1999 in [Language] Expose T.warpgroup_fence_operand for nvcc code motion #986
[Language] Add Correctness and performance check scripts for V2 by @LeiWang1999 in [Language] Add Correctness and performance check scripts for V2 #1174
[Bugfix] Legalize Datatype for mma intrinisc codegen by @LeiWang1999 in [Bugfix] Legalize Datatype for mma intrinisc codegen #1179
[CI]: Bump actions/download-artifact from 5 to 6 by @dependabot[bot] in [CI]: Bump actions/download-artifact from 5 to 6 #1177
[CI]: Bump actions/upload-artifact from 4 to 5 by @dependabot[bot] in [CI]: Bump actions/upload-artifact from 4 to 5 #1178
[Language] Initial version of tilelang frontend v2 by @kurisu6912 in [Language] Initial version of tilelang frontend v2 #1120
[Fix] fix type imcompatible error in [BUG] index of T.Tensor cannot be int64 #1115 by @kurisu6912 in [Fix] fix type imcompatible error in #1115 #1180
[CI] [pre-commit.ci] autoupdate by @pre-commit-ci[bot] in [CI] [pre-commit.ci] autoupdate #1183
[Fix] Remove unsupported type params by @kurisu6912 in [Fix] Remove unsupported type params #1186
[Feature] Enhance fill operation to support various buffer types by @LeiWang1999 in [Feature] Enhance fill operation to support various buffer types #1189
[Refactor] Improve Python3.9 compatibility for ParamSpec and Self by @LeiWang1999 in [Refactor] Improve Python3.9 compatibility for ParamSpec and Self #1190
[Feat] Add swap like grammar in tuple assignment by @kurisu6912 in [Feat] Add swap like grammar in tuple assignment #1185
[Release] Unify local build scripts to use cibuildwheel and reduce size of sdist by @oraluben in [Release] Unify local build scripts to use cibuildwheel and reduce size of sdist #1171
[Langauge] Support n>256 for v2 by @LeiWang1999 in [Langauge] Support n>256 for v2 #1182
[GQA] Use TMA in GQA bwd kernel to boost performance by @tzj-fxz in [GQA] Use TMA in GQA bwd kernel to boost performance #1176
[Example] Update GQA varlen fwd by @chengyupku in [Example] Update GQA varlen fwd #1173
[Refactor] Dynamic registration of FP8 data type for compatibility with older PyTorch versions by @LeiWang1999 in [Refactor] Dynamic registration of FP8 data type for compatibility with older PyTorch versions #1197
[Feature] Add tl.infinity operator for infinity handling of bfloat16 by @Rachmanino in [Feature] Add tl.infinity operator for infinity handling of bfloat16 #1175
[SM70] Refactor and minor fix for SM70 by @LeiWang1999 in [SM70] Refactor and minor fix for SM70 #1195
[CI] Enable ccache for CIBW on Linux by @oraluben in [CI] Enable ccache for CIBW on Linux #1184
[Feat] Add support for T.serial with step and negative step by @kurisu6912 in [Feat] Add support for T.serial with step and negative step #1188
[Feat] Add A Pass to Handle Negative Index by @kurisu6912 in [Feat] Add A Pass to Handle Negative Index #1192
Fix type errors in reduce.h by @LJC00118 in Fix type errors in reduce.h #1204
[Bugfix] Improves the accuracy of dependency analysis in the storage access by @LeiWang1999 in [Bugfix] Improves the accuracy of dependency analysis in the storage access #1205
[Bugfix][Language V2] Capture closure variables from program by @LeiWang1999 in [Bugfix][Language V2] Capture closure variables from program #1206
Fix Dockerfile.cu128 by @createthis in Fix Dockerfile.cu128 #1208
[Enhancement] Improve handling of negative indices for ramp and broadcast node by @LeiWang1999 in [Enhancement] Improve handling of negative indices for ramp and broadcast node #1207
[Bugfix] Enhane LetStmt Handling in Pipeline Transform by @LeiWang1999 in [Bugfix] Enhane LetStmt Handling in Pipeline Transform #1212
[Fix] Fix buffer re-import typo in tilelang.languge by @kurisu6912 in [Fix] Fix buffer re-import typo in tilelang.languge #1214
[Build] Explicitly add libtvm as a dep of libtilelang by @oraluben in [Build] Explicitly add libtvm as a dep of libtilelang #1215
[Utils] Add source export, NVCC-based PTX/SASS dump, logging by @LeiWang1999 in [Utils] Add source export, NVCC-based PTX/SASS dump, logging #1216
[Bugfix] Improve error handling in LayoutNode InverseWithLevel by @LeiWang1999 in [Bugfix] Improve error handling in LayoutNode InverseWithLevel #1220
[Enhancement] Improve iterator handling in layout utilities and parallel operations by @LeiWang1999 in [Enhancement] Improve iterator handling in layout utilities and parallel operations #1221
[Language] Refactor reduce and support shared memory as its in/out by @LeiWang1999 in [Language] Refactor reduce and support shared memory as its in/out #1219
[GQA] Add varlen decoding kernel with logits saving by @tzj-fxz in [GQA] Add varlen decoding kernel with logits saving #1223
[Enhancement] Add thread count validation for ReduceOp fragment layout inference by @LeiWang1999 in [Enhancement] Add thread count validation for ReduceOp fragment layout inference #1225
[Refactor] Simplify logic in the CompleteBufferFragment by @LeiWang1999 in [Refactor] Simplify logic in the CompleteBufferFragment #1226
[Refactor] Refactor version retrieval logic in tilelang package by @LeiWang1999 in [Refactor] Refactor version retrieval logic in tilelang package #1227
[CPU] Minor fix for cpu backend by @LeiWang1999 in [CPU] Minor fix for cpu backend #1230
[Feature] Add Release Plan issue template by @LeiWang1999 in [Feature] Add Release Plan issue template #1231
[Fix] Fix a type that make wrong T.macro backtrace by @kurisu6912 in [Fix] Fix a type that make wrong T.macro backtrace #1234
[Refactor] Add kernel selection option for GEMM v1 in environment settings by @LeiWang1999 in [Refactor] Add kernel selection option for GEMM v1 in environment settings #1200
[Bugfix] Minor fix in builder.py by @LJC00118 in [Bugfix] Minor fix in builder.py #1235
[Language] Add type stubs for tir op by @kurisu6912 in [Language] Add type stubs for tir op #1239
[Enhancement] Support Layout/Fragment Reshape by @LeiWang1999 in [Enhancement] Support Layout/Fragment Reshape #1241
[Bugfix] Minor fix for tcgen05 by @LeiWang1999 in [Bugfix] Minor fix for tcgen05 #1242
RMSNorm epsilon refine in the example by @pengxin99 in RMSNorm epsilon refine in the example #1243
[AMD] enable amd ci test & fix bug & fix dockerfile by @Paran0idy in [AMD] enable amd ci test & fix bug & fix dockerfile #1244
[Refactor] Phaseout legacy loop vectorize dynamic pass by @LeiWang1999 in [Refactor] Phaseout legacy loop vectorize dynamic pass #1245
[Bugfix] Fix fp8 dtype for some cases by @LeiWang1999 in [Bugfix] Fix fp8 dtype for some cases #1246
[Minor] Remove git_commit.txt by @SiriusNEO in [Minor] Remove git_commit.txt #1249
[Language][Reshape] Improve variable handling and ensure correctness during Layout Reshape by @LeiWang1999 in [Language][Reshape] Improve variable handling and ensure correctness during Layout Reshape #1248
[Refactor] Update buffer handling in copy and atomic operations by @LeiWang1999 in [Refactor] Update buffer handling in copy and atomic operations #1247
[Language] Add missing while statement by @kurisu6912 in [Language] Add missing while statement #1254
[BugFix] Add autotune and exp2 for GDN kernel by @tzj-fxz in [BugFix] Add autotune and exp2 for GDN kernel #1258
[BugFix] Refactor attention kernel to handle OOB positions by filling with -inf instead of clearing accumulators. by @Rachmanino in [BugFix] Refactor attention kernel to handle OOB positions by filling with -inf instead of clearing accumulators. #1222
[fix] NVRTC execution backend by @lucifer1004 in [fix] NVRTC execution backend #1256
[AMD] Update CK for ROCm7 by @Paran0idy in [AMD] Update CK for ROCm7 #1262
[BugFix] Remove memory_order in atomic constexpr and fix NSA bwd by @KevinZeng08 in [BugFix] Remove memory_order in atomic constexpr and fix NSA bwd #1260
[Example] Add GQA decoding kernel with varlen page table by @tzj-fxz in [Example] Add GQA decoding kernel with varlen page table #1265
[Refactor] add support for numpy dtype conversion by @kurisu6912 in [Refactor] add support for numpy dtype conversion #1255
[EXAMPLE] In the flash attention example keep the max of all blocks seen in scores_max numerical stability by @vpj in [EXAMPLE] In the flash attention example keep the max of all blocks seen in scores_max numerical stability #1148
[Docs] Improve Installation Guide by @SiriusNEO in [Docs] Improve Installation Guide #1270
[Enhancement] Keep max score attention across blocks in FlashAttention for better numerical stablity by @Rachmanino in [Enhancement] Keep max score attention across blocks in FlashAttention for better numerical stablity #1269
[Bugfix] Fix multiple cg defination when using T.sync_grid by @chengyupku in [Bugfix] Fix multiple cg defination when using T.sync_grid #1272
[Minor] Remove from __future__ import annotations for python 3.8 by @oraluben in [Minor] Remove from __future__ import annotations for python 3.8 #1273
[BugFix] Adding extra parameters into autotune hashkey by @SiriusNEO in [BugFix] Adding extra parameters into autotune hashkey #1274
Fix various issues under int64_t static and dynamic shape. by @Elevator14B in Fix various issues under int64_t static and dynamic shape. #1218
Bug fix for Gated Delta Net benchmark script by @learning-chip in Bug fix for Gated Delta Net benchmark script #1267
[Bugfix] Minor fix for some cases by @LeiWang1999 in [Bugfix] Minor fix for some cases #1278
[Language] Add shape check in T.view/reshape by @SiriusNEO in [Language] Add shape check in T.view/reshape #1277
[FFI] Use tvm ffi as the default execution backend by @LeiWang1999 in [FFI] Use tvm ffi as the default execution backend #1259
[Bugfix] Supply missing T.print for bool type by @LeiWang1999 in [Bugfix] Supply missing T.print for bool type #1279
[Fix] Fix memory leak bug by @kurisu6912 in [Fix] Fix memory leak bug #1281
[Enhancement] Enhance CUDA compilation by integrating pass context configuration by @LeiWang1999 in [Enhancement] Enhance CUDA compilation by integrating pass context configuration #1283
Fix the bug in issue [BUG] example_tilelang_nsa_fwd.py is inconsistent with the reference implementation and specific parameters . #1266 by @sea-with-sakura in Fix the bug in issue #1266 #1284
[Language][UX] Nested loop checker in pre-lowering stage by @SiriusNEO in [Language][UX] Nested loop checker in pre-lowering stage #1288
[Compatibility] Support CUDA 11.3 by @LeiWang1999 in [Compatibility] Support CUDA 11.3 #1290
[Feat] Add support for using T.Tensor(n * 2 + 1) in function annotation by @kurisu6912 in [Feat] Add support for using T.Tensor(n * 2 + 1) in function annotation #1285
[Feat] Add missing support to pass reference by T.Var annotation by @kurisu6912 in [Feat] Add missing support to pass reference by T.Var annotation #1291
[Enhancement] Shared Memory Size Can be Dynamic by @LeiWang1999 in [Enhancement] Shared Memory Size Can be Dynamic #1294
[Fix] Remove unused let_bindings_ in CodeGenC to fix [BUG] Weird TVM internal Error #1300 by @kurisu6912 in [Fix] Remove unused let_bindings_ in CodeGenC to fix #1300 #1305
[Bugfix] Fallback to the old AtomicAdd implementation for legacy architectures by @LeiWang1999 in [Bugfix] Fallback to the old AtomicAdd implementation for legacy architectures #1306
[Fix] Fix frame scope error in T.macro by @kurisu6912 in [Fix] Fix frame scope error in T.macro #1308
[WIP] support more dtypes for tcgen05 by @PannenetsF in [WIP] support more dtypes for tcgen05 #1229
Improve memory access safety and T.assume handling by @LJC00118 in Improve memory access safety and T.assume handling #1292
[Bugfix] Fix autotune cache by @LeiWang1999 in [Bugfix] Fix autotune cache #1315
[Refactor] Backup Analyzer to get the appropriate arith informations by @LeiWang1999 in [Refactor] Backup Analyzer to get the appropriate arith informations #1311
Revert "[WIP] support more dtypes for tcgen05 ([WIP] support more dtypes for tcgen05 #1229)" by @LeiWang1999 in Revert "[WIP] support more dtypes for tcgen05 (#1229)" #1323
[CI]: Bump actions/checkout from 5 to 6 by @dependabot[bot] in [CI]: Bump actions/checkout from 5 to 6 #1319
[CI]: Bump pypa/cibuildwheel from 3.2 to 3.3 by @dependabot[bot] in [CI]: Bump pypa/cibuildwheel from 3.2 to 3.3 #1318
[Installation] Fix building using customized TVM path by @SiriusNEO in [Installation] Fix building using customized TVM path #1326
[Release] Allow developer with write permission to trigger wheel release by @oraluben in [Release] Allow developer with write permission to trigger wheel release #1322
[Feat] Support warp reduce by @Rachmanino in [Feat] Support warp reduce #1316
[Enhancement] Support more dtype in T.print by @xwhzz in [Enhancement] Support more dtype in T.print #1329
[BugFix] Use BufferRegion in tl.cumsum to infer buffer shape by @SiriusNEO in [BugFix] Use BufferRegion in tl.cumsum to infer buffer shape #1321
[Fix] Fix uint narrowing bug in [BUG] Wrong type recognition in T.serial #1310 by @kurisu6912 in [Fix] Fix uint narrowing bug in #1310 #1320
[Refactor] Disable strided buffer load inside tvm ([BUG] strided index cause silent bug #1301) by @kurisu6912 in [Refactor] Disable strided buffer load inside tvm (#1301) #1332
[Refactor] Moving NormalizeToBufferRegion and MakeAccessPtrFromRegion to utils by @LeiWang1999 in [Refactor] Moving NormalizeToBufferRegion and MakeAccessPtrFromRegion to utils #1333
[Fix] Fix bug copying from or to local buffer ([BUG] T.copy has bad behavior from global memory to local memory #1304) by @kurisu6912 in [Fix] Fix bug copying from or to local buffer (#1304) #1324
[Language][UX] Semantic check for parallel fragment access by @SiriusNEO in [Language][UX] Semantic check for parallel fragment access #1338
Add unit tests for T.assume by @LJC00118 in Add unit tests for T.assume #1341
[Feat] Extend LegalizeNegativeIndex to support buffer store stmts by @ConvolutedDog in [Feat] Extend LegalizeNegativeIndex to support buffer store stmts #1339
[Refactor] Phaseout vmap for Tile Operators by @LeiWang1999 in [Refactor] Phaseout vmap for Tile Operators #1334
[Enhancement] add more dtype and fix mma.ws for fp16 for tcgen05 by @PannenetsF in [Enhancement] add more dtype and fix mma.ws for fp16 for tcgen05 #1327
[Refactor] Enhance CopyNode's IterVar Creation and Range Handling by @LeiWang1999 in [Refactor] Enhance CopyNode's IterVar Creation and Range Handling #1346
[Fix] Fix missing not operator in frontend ([BUG] ValueError: Cannot use and / or / not operator to Expr, hint: use tvm.tir.all / tvm.tir.any instead #1347) by @kurisu6912 in [Fix] Fix missing not operator in frontend (#1347) #1348
[Enhancement] Add support for k_pack in gemm_mfma by @Gongen-Ali in [Enhancement] Add support for k_pack in gemm_mfma #1344
Add sparse fine-tuning kernel for deepseek sparse attention to example by @hyx1999 in Add sparse fine-tuning kernel for deepseek sparse attention to example #1296
[Refactor] Improve assertion handling in CodeGenCHost and ArgBinder by @LeiWang1999 in [Refactor] Improve assertion handling in CodeGenCHost and ArgBinder #1352
[Refactor] Simplify index sign state handling in LegalizeNegativeIndex by @LeiWang1999 in [Refactor] Simplify index sign state handling in LegalizeNegativeIndex #1354
[Enhancement] Improve error handling and assertion messages across runtime and argument binding by @LeiWang1999 in [Enhancement] Improve error handling and assertion messages across runtime and argument binding #1356
[Bugfix] Disable floordiv optimization due to integer overflow risk by @LJC00118 in [Bugfix] Disable floordiv optimization due to integer overflow risk #1355
[Bugfix] Fix the jit_kernel issue by @gfvvz in [Bugfix] Fix the jit_kernel issue #1357
[Bugfix] Bind thread range for fragment inference in Parallel strict layout inference stage. by @LeiWang1999 in [Bugfix] Bind thread range for fragment inference in Parallel strict layout inference stage. #1359
[Analysis] Enhance NestedLoopChecker with tile op cases by @SiriusNEO in [Analysis] Enhance NestedLoopChecker with tile op cases #1358
[Language] support T.gemm_sp_v2 on sm80 and sm89 by @botbw in [Language] support T.gemm_sp_v2 on sm80 and sm89 #1056
[Bugfix] Update TIR registration for GemmSPPy to use tile operation by @LeiWang1999 in [Bugfix] Update TIR registration for GemmSPPy to use tile operation #1361
[Enhancement] Implement dynamic unroll factor in CUDA code generation by @LeiWang1999 in [Enhancement] Implement dynamic unroll factor in CUDA code generation #1360
[CI] [pre-commit.ci] autoupdate by @pre-commit-ci[bot] in [CI] [pre-commit.ci] autoupdate #1362
[Bugfix] Remove debug print in PyStmtFunctionVisitor by @LeiWang1999 in [Bugfix] Remove debug print in PyStmtFunctionVisitor #1363
[Debug] Always include line info in NVCC command for improved profiling by @LeiWang1999 in [Debug] Always include line info in NVCC command for improved profiling #1364
[Enhancemnet] Minor fix to speed up testing by @LeiWang1999 in [Enhancemnet] Minor fix to speed up testing #1365
[Enhancement] Add DISABLE_CACHE environment variables by @SiriusNEO in [Enhancement] Add DISABLE_CACHE environment variables #1368
[Refactor]: Remove useless include in atomicadd_vectorize.h by @yyttt6 in [Refactor]: Remove useless include in atomicadd_vectorize.h #1371
[Refactor] Generalize fp8 process by @LeiWang1999 in [Refactor] Generalize fp8 process #1372
[Layout] Enhance Free Layout Inference by @LeiWang1999 in [Layout] Enhance Free Layout Inference #1375
[Enhancement] Introduce buffer var lca analysis for pass plan buffer allocations by @LeiWang1999 in [Enhancement] Introduce buffer var lca analysis for pass plan buffer allocations #1376
[Tool] Provide layout visualization tool by @Cunxiao2002 in [Tool] Provide layout visualization tool #1353
[Release] Relax constraint of tvm-ffi to compatible version by @oraluben in [Release] Relax constraint of tvm-ffi to compatible version #1373
[Language] Tilelang LazyJIT Experimental Version by @kurisu6912 in [Language] Tilelang LazyJIT Experimental Version #1337
[Builder] Enhance variable name binding and scope management by @LeiWang1999 in [Builder] Enhance variable name binding and scope management #1378
[Bugfix] make cuda driver api compat with cuda12/13, along with tests by @PannenetsF in [Bugfix] make cuda driver api compat with cuda12/13, along with tests #1379
[Fix] typo in cuda attr by @PannenetsF in [Fix] typo in cuda attr #1380
[Language V2] Minor fix for complex annotations by @LeiWang1999 in [Language V2] Minor fix for complex annotations #1381
[Release] Bump Version into 0.1.7 by @LeiWang1999 in [Release] Bump Version into 0.1.7 #1377
[Typing] Enhance compatibility for advanced typing features for Py39 by @LeiWang1999 in [Typing] Enhance compatibility for advanced typing features for Py39 #1382

New Contributors

@LJC00118 made their first contribution in [Cython] Remove an incorrect check #880
@Edenzzzz made their first contribution in [Bugfix] Fix flops comp and softmax scale in mla #900
@zjudmd1015 made their first contribution in [Bugfix] Fix saving kernel source code where JITKernel.artifact is None #921
@lijinpei made their first contribution in [Example] Optimize online_softmax example #934
@Zhichenzzz made their first contribution in [Example] Add sparse mla bwd example for deepseek_v32 #919
@BBuf made their first contribution in [Code Style] Refine nvrtc compile related check style #945
@XuehaiPan made their first contribution in [CI] enable dependabot for GHA workflows #950
@iloveai8086 made their first contribution in Modify the SM architecture number to support Thor’s sm110. #957
@Degeneracy-Evil made their first contribution in [Bugfix] Add NVIDIA HPC SDK support in CUDA detection (#974) #976
@pre-commit-ci[bot] made their first contribution in [CI] [pre-commit.ci] autoupdate #1183
@createthis made their first contribution in Fix Dockerfile.cu128 #1208
@pengxin99 made their first contribution in RMSNorm epsilon refine in the example #1243
@KevinZeng08 made their first contribution in [BugFix] Remove memory_order in atomic constexpr and fix NSA bwd #1260
@vpj made their first contribution in [EXAMPLE] In the flash attention example keep the max of all blocks seen in scores_max numerical stability #1148
@Elevator14B made their first contribution in Fix various issues under int64_t static and dynamic shape. #1218
@learning-chip made their first contribution in Bug fix for Gated Delta Net benchmark script #1267
@sea-with-sakura made their first contribution in Fix the bug in issue #1266 #1284
@PannenetsF made their first contribution in [WIP] support more dtypes for tcgen05 #1229
@ConvolutedDog made their first contribution in [Feat] Extend LegalizeNegativeIndex to support buffer store stmts #1339
@Gongen-Ali made their first contribution in [Enhancement] Add support for k_pack in gemm_mfma #1344

Full Changelog: 0.1.6...v0.1.7

This discussion was created from the release v0.1.7.

Replies: 0 comments

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment