v0.1.4 · tile-ai tilelang · Discussion #406 · GitHub

LeiWang1999
Apr 18, 2025
Maintainer

What's Changed

[Bugfix] Support T.clear for let binding by @LeiWang1999 in [Bugfix] Support T.clear for let binding #268
[Bugfix] Add TMA and Producer Buffer Analysis in Warp Specialized Rewriter by @chengyupku in [Bugfix] Add TMA and Producer Buffer Analysis in Warp Specialized Rewriter #269
[Refactor] Improve flash attention example and layout comparison logic by @LeiWang1999 in [Refactor] Improve flash attention example and layout comparison logic #270
[Bugfix]Add CUDA availability check in CtypesKernelAdapter by @XueSongTap in [Bugfix]Add CUDA availability check in CtypesKernelAdapter #267
[CI] Add gemm performance test by @xwhzz in [CI] Add gemm performance test #274
[Language] Introduce T.ptr and T.Tensor by @LeiWang1999 in [Language] Introduce T.ptr and T.Tensor #276
[Refactor] Enhance Autotune by @yyttt6 in [Refactor] Enhance Autotune #266
[Refactor] Update cache key generation in KernelCache by @LeiWang1999 in [Refactor] Update cache key generation in KernelCache #283
[Docs][Tutorial] Add tutorial for auto-tuning by @yyttt6 in [Docs][Tutorial] Add tutorial for auto-tuning #285
[Refactor] Deprecated T.Buffer as arguments and rename related calls into T.Tensor by @LeiWang1999 in [Refactor] Deprecated T.Buffer as arguments and rename related calls into T.Tensor #281
[Doc] Update README.md to correct documentation link for TileLang debug tools by @chengyupku in [Doc] Update README.md to correct documentation link for TileLang debug tools #286
[Feature] Introduce NoSetMaxNReg for warp specialization by @chengyupku in [Feature] Introduce NoSetMaxNReg for warp specialization #289
[Language] Proxy tvm ir to make linter happy by @LeiWang1999 in [Language] Proxy tvm ir to make linter happy #287
[Bugfix] Enable bfloat16 atomic operations only for CUDA architectures greater than 7.5 by @LeiWang1999 in [Bugfix] Enable bfloat16 atomic operations only for CUDA architectures greater than 7.5 #291
[Doc] Update Python API docs generation by @xwhzz in [Doc] Update Python API docs generation #278
[Doc] Remove citation page by @LeiWang1999 in [Doc] Remove citation page #292
[Dev] Correcting cxx compiler by @penguin-wwy in [Dev] Correcting cxx compiler #294
[doc/example] add gemv doc and examples by @botbw in [doc/example] add gemv doc and examples #293
[Feature] Implement ParallelLoopTransformer for enhanced loop analysis by @LeiWang1999 in [Feature] Implement ParallelLoopTransformer for enhanced loop analysis #295
[Enhancement] Update AtomicAdd functions for BFLOAT16 in common.h by @LeiWang1999 in [Enhancement] Update AtomicAdd functions for BFLOAT16 in common.h #297
[Refactor] Improve documentation and add detailed docstrings across multiple modules by @LeiWang1999 in [Refactor] Improve documentation and add detailed docstrings across multiple modules #298
[Bugfix] Correct method call for block reduction check when analyzing memory footprint by @NaOHCC in [Bugfix] Correct method call for block reduction check when analyzing memory footprint #299
[Dynamic Symbolic] Refactor passes with dynamic symbolic and check shape bound precisely by @tzj-fxz in [Dynamic Symbolic] Refactor passes with dynamic symbolic and check shape bound precisely #302
Add autotune to conv example by @yyttt6 in Add autotune to conv example #301
[Bugfix] Resolve autotuner bugs for blocksparse GEMM example by @tth37 in [Bugfix] Resolve autotuner bugs for blocksparse GEMM example #300
[Bugfix] Replace profiler.mod with profiler.adapter to fix AttributeError by @LeslinD in [Bugfix] Replace profiler.mod with profiler.adapter to fix AttributeError #305
[Enhancement] Add support for CUDA architecture 8.9 in GEMM template by @LeiWang1999 in [Enhancement] Add support for CUDA architecture 8.9 in GEMM template #304
[BugFix] Fix unintended Git config overrides in CI runners by @xwhzz in [BugFix] Fix unintended Git config overrides in CI runners #306
[Cache] Implement in-memory cache by @LeiWang1999 in [Cache] Implement in-memory cache #308
[Bugfix] Updated autotune usage in the examples to align with the latest changes by @LeiWang1999 in [Bugfix] Updated autotune usage in the examples to align with the latest changes #309
[Bugfix] Fix dynamic axis with variable extent by @LeiWang1999 in [Bugfix] Fix dynamic axis with variable extent #311
[Bugfix] Fix layout conflict issue for gqa decoding examples by @LeiWang1999 in [Bugfix] Fix layout conflict issue for gqa decoding examples #314
[Bugfix] Fixed the handling logic of IfThenElseNode in if_stmt_binding by @chengyupku in [Bugfix] Fixed the handling logic of IfThenElseNode in if_stmt_binding #315
[Bugfix] Fix logic error in ReduceOp when handling CUDA architecture by @chengyupku in [Bugfix] Fix logic error in ReduceOp when handling CUDA architecture #316
[CostModel] Introduce cuda driver api to get precise shared memory capacity by @LeiWang1999 in [CostModel] Introduce cuda driver api to get precise shared memory capacity #317
[Dev] Add FP8 Quantization Examples and Absolute Maximum Reduction Operation Support by @chengyupku in [Dev] Add FP8 Quantization Examples and Absolute Maximum Reduction Operation Support #320
[Tools] Summarize TFLOPS Information from a tilelang program by @yyttt6 in [Tools] Summarize TFLOPS Information from a tilelang program #321
Support block_N sizes that are 2^n in deepgemm example by @zcnrex in Support block_N sizes that are 2^n in deepgemm example #319
[Feat] Enhance CUDA Property Handling by @LeiWang1999 in [Feat] Enhance CUDA Property Handling #322
[Bugfix] add a patch to fix T.abs on float16 by @botbw in [Bugfix] add a patch to fix T.abs on float16 #325
[AMD] Adapt rocm and support T.gemm with transpose_b=False for amd backend by @LeiWang1999 in [AMD] Adapt rocm and support T.gemm with transpose_b=False for amd backend #327
[Dynamic Symbolic] Adaptively vectorize with different condition expressions by @tzj-fxz in [Dynamic Symbolic] Adaptively vectorize with different condition expressions #326
[Bugfix] Fix fragment layout annotation in example gqa decode by @LeiWang1999 in [Bugfix] Fix fragment layout annotation in example gqa decode #329
[AMD] Support Transpose_A=True and GEMM_RS for hip backend by @LeiWang1999 in [AMD] Support Transpose_A=True and GEMM_RS for hip backend #331
[Refactor] Optimize RMS normalization kernel in rms_norm.py by @chengyupku in [Refactor] Optimize RMS normalization kernel in rms_norm.py #333
[AMD] Fix for missing composable kernel include path when compile kernels on amd gpus by @LeiWang1999 in [AMD] Fix for missing composable kernel include path when compile kernels on amd gpus #334
[Example] Add sparse gqa decode example by @xiayuqing0622 in [Example] Add sparse gqa decode example #332
[Enhancement] Enhance FP8/FP4 type handling in CUDA codegen by @LeiWang1999 in [Enhancement] Enhance FP8/FP4 type handling in CUDA codegen #323
[Doc] Fix typo and heading level in GEMV tutorial by @yeh-sudo in [Doc] Fix typo and heading level in GEMV tutorial #337
[Dev] Add Group Cast FP8 Example by @chengyupku in [Dev] Add Group Cast FP8 Example #338
[Enhancement] Support region padding when convert buffer load to buffer region by @LeiWang1999 in [Enhancement] Support region padding when convert buffer load to buffer region #342
[Example] Add triton block sparse gqa decode by @YizhaoGao in [Example] Add triton block sparse gqa decode #341
[Enhancement] Support index bit width configuration by @LeiWang1999 in [Enhancement] Support index bit width configuration #343
[Bugfix] Fix X_amax Correctness Issue in Group Cast FP8 by @chengyupku in [Bugfix] Fix X_amax Correctness Issue in Group Cast FP8 #345
[Bugfix] Fix Transposed Fragment Layout for amd GEMM_RS matrix core by @LeiWang1999 in [Bugfix] Fix Transposed Fragment Layout for amd GEMM_RS matrix core #346
[AutoTune] Refactor AutoTuneArtifact to utilize kernel as context instead of profiler by @LeiWang1999 in [AutoTune] Refactor AutoTuneArtifact to utilize kernel as context instead of profiler #344
[Bugfix] Compile/"cached" still not loading cached kernel for example in example_mha_bwd by @Alex4210987 in [Bugfix] Compile/"cached" still not loading cached kernel for example in example_mha_bwd #339
[Refactor] Implement thread-local storage for FrameStack in frame.py and kernel.py by @LeiWang1999 in [Refactor] Implement thread-local storage for FrameStack in frame.py and kernel.py #352
[Typo] Replace kernel.func with kernel in mla benchmark scripts by @LeiWang1999 in [Typo] Replace kernel.func with kernel in mla benchmark scripts #354
[AMD][Docker] Create Dockerfile for ROCm environment setup by @LeiWang1999 in [AMD][Docker] Create Dockerfile for ROCm environment setup #355
[Enhancement] Update group_per_split_token_cast_to_fp8 to support multiple data types by @chengyupku in [Enhancement] Update group_per_split_token_cast_to_fp8 to support multiple data types #356
[Enhancement] Support pass config disable_warp_specialize to disable auto specialization on hopper by @LeiWang1999 in [Enhancement] Support pass config disable_warp_specialize to disable auto specialization on hopper #357
[Example] Introduce autotuning example for GEMM with enhanced configuration options by @chengyupku in [Example] Introduce autotuning example for GEMM with enhanced configuration options #360
[Example] Handle Scenarios in Which a Threadblock is Assigned Only Invalid Block Indices for Sparse Attention by @xiayuqing0622 in [Example] Handle Scenarios in Which a Threadblock is Assigned Only Invalid Block Indices for Sparse Attention #361
[Bugfix] Correct dynamic shared memory size error handling in HIP by @LeiWang1999 in [Bugfix] Correct dynamic shared memory size error handling in HIP #362
[AMD] Implement Deepseek MLA for AMD by @LeiWang1999 in [AMD] Implement Deepseek MLA for AMD #363
[Bugfix] Fix compilation issues for amd cdna element size check by @LeiWang1999 in [Bugfix] Fix compilation issues for amd cdna element size check #364
[AMD] Support FlashMLA with num split template for AMD gpus by @LeiWang1999 in [AMD] Support FlashMLA with num split template for AMD gpus #366
[MLA][AMD] Add amd mla benchmarking by @LeiWang1999 in [MLA][AMD] Add amd mla benchmarking #367
[Bugfix] Adjust Autotuner threadpool max_workers limit to available CPUs by @tth37 in [Bugfix] Adjust Autotuner threadpool max_workers limit to available CPUs #368
[Language] Introduce T.any_of and T.all_of to reduce a bool arrary by @LeiWang1999 in [Language] Introduce T.any_of and T.all_of to reduce a bool arrary #371
[AMD][Setup] Support HIP in setup.py by @zhhangBian in [AMD][Setup] Support HIP in setup.py #369
[Typo] Remove debug print by @LeiWang1999 in [Typo] Remove debug print #373
[Docs] Add AMD Flash MLA Documentation to Tutorials Section by @LeiWang1999 in [Docs] Add AMD Flash MLA Documentation to Tutorials Section #376
[Bugfix] Add filelock for cython build by @LeiWang1999 in [Bugfix] Add filelock for cython build #377
[Typo] Remove unused comments generated by copilot by @LeiWang1999 in [Typo] Remove unused comments generated by copilot #379
[Doc] Add deepseek_mla to documentation index by @LeiWang1999 in [Doc] Add deepseek_mla to documentation index #380
[Refactor] Remove debug message in pass legalize_safe_memory_access by @LeiWang1999 in [Refactor] Remove debug message in pass legalize_safe_memory_access #381
[Enhancement][Pipeline] More precise copy code block detection in pipeline by @LeiWang1999 in [Enhancement][Pipeline] More precise copy code block detection in pipeline #384
[Revert] Revert modifications for pass FlattenBuffer by @LeiWang1999 in [Revert] Revert modifications for pass FlattenBuffer #385
[Dynamic Symbolic] Add pass_config to customize vectorization and tail split by @tzj-fxz in [Dynamic Symbolic] Add pass_config to customize vectorization and tail split #383
[Pytest Fix] Wrap tests in dynamic benchmark by @tzj-fxz in [Pytest Fix] Wrap tests in dynamic benchmark #387
[Doc] Update README.md for deepseek_mla on AMD by @LeiWang1999 in [Doc] Update README.md for deepseek_mla on AMD #389
[Pipeline][Enhancement] Add copy_prepare stage to support mask and index caching by @LeiWang1999 in [Pipeline][Enhancement] Add copy_prepare stage to support mask and index caching #392
[Refactor] Refactor warp_specialized_rewriter to support multiple acquire/release patterns by @chengyupku in [Refactor] Refactor warp_specialized_rewriter to support multiple acquire/release patterns #391
[Enhancement] Report Error Body in ParallelOp Layout Inference by @chengyupku in [Enhancement] Report Error Body in ParallelOp Layout Inference #394
[Bugfix] Support T.Parallel with local register assignment by @LeiWang1999 in [Bugfix] Support T.Parallel with local register assignment #395
[Enhancement] Introduce a smarter warp partition strategy by @LeiWang1999 in [Enhancement] Introduce a smarter warp partition strategy #396
[Example] Add bitnet-1.58b examples by @LeiWang1999 in [Example] Add bitnet-1.58b examples #399
Bump transformers from 4.40 to 4.48.0 in /examples/bitnet-1.58b by @dependabot in Bump transformers from 4.40 to 4.48.0 in /examples/bitnet-1.58b #400
[BugFix] Address should aligned with access size in tail split by @tzj-fxz in [BugFix] Address should aligned with access size in tail split #401
[Enhancement] Move T.any_of and T.all_of op registration from python into cpp by @Cunxiao2002 in [Enhancement] Move T.any_of and T.all_of op registration from python into cpp #398
Add preliminary support for bf16 for AMD by @OscarSavolainen in Add preliminary support for bf16 for AMD #388
[BugFix] Conditions Robustness in dynamic vectorize by @tzj-fxz in [BugFix] Conditions Robustness in dynamic vectorize #404
[CI] Update CI configuration to run pytest with automatic parallelization by @LeiWang1999 in [CI] Update CI configuration to run pytest with automatic parallelization #393
[Documentation] Fix Installation Documentation by @andyluo03 in [Documentation] Fix Installation Documentation #405

New Contributors

@XueSongTap made their first contribution in [Bugfix]Add CUDA availability check in CtypesKernelAdapter #267
@botbw made their first contribution in [doc/example] add gemv doc and examples #293
@NaOHCC made their first contribution in [Bugfix] Correct method call for block reduction check when analyzing memory footprint #299
@tth37 made their first contribution in [Bugfix] Resolve autotuner bugs for blocksparse GEMM example #300
@LeslinD made their first contribution in [Bugfix] Replace profiler.mod with profiler.adapter to fix AttributeError #305
@zcnrex made their first contribution in Support block_N sizes that are 2^n in deepgemm example #319
@xiayuqing0622 made their first contribution in [Example] Add sparse gqa decode example #332
@yeh-sudo made their first contribution in [Doc] Fix typo and heading level in GEMV tutorial #337
@YizhaoGao made their first contribution in [Example] Add triton block sparse gqa decode #341
@zhhangBian made their first contribution in [AMD][Setup] Support HIP in setup.py #369
@dependabot made their first contribution in Bump transformers from 4.40 to 4.48.0 in /examples/bitnet-1.58b #400
@OscarSavolainen made their first contribution in Add preliminary support for bf16 for AMD #388
@andyluo03 made their first contribution in [Documentation] Fix Installation Documentation #405

Full Changelog: v0.1.3...v0.1.4

This discussion was created from the release v0.1.4.

Replies: 1 comment

LxinJ05
Jun 6, 2025

TileLang 感觉太猛了，看到最近北京有一个 AI 编译器活动分享→https://hdxu.cn/1cmhk，阵容很硬啊，AMD 都请来了，一整个期待住了

0 replies

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