GitLab

* Introduced a new function `get_nvcc_compiler` in nvcc.py to obtain the path to the nvcc compiler.
* Updated LibraryGenerator to use `get_nvcc_compiler` instead of hardcoding the nvcc command, improving maintainability and flexibility.

* [New Feature] Add FP8 Flash Attention Implementation (#412)

* Introduce a new example script for FP8 Flash Attention in `example_mla_decode_kv_fp8.py`, showcasing the use of tilelang for efficient attention computation.
* Implement the `flashattn` function with optimized memory management and kernel execution.
* Include a reference program for comparison and performance evaluation.
* Add command-line argument parsing for batch size, number of heads, and dimensions to facilitate testing and experimentation.
* Enhance the overall structure and readability of the code.

This addition aims to improve the performance of attention mechanisms in deep learning models by leveraging FP8 precision and optimized kernel execution.

* lint fix

* optimize quick start

* lint fix

* [Language] make linter and type checker happy with mocking

* Apply suggestions from code review
Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>

* Refactor BaseTensor class in proxy.py to implement __getitem__ and __setitem__ methods, enhancing type checking and linting compliance. Added method stubs for from_ptr and other subclasses for improved clarity and maintainability.

* Refactor type imports in proxy.py to enhance clarity and maintainability by replacing the built-in Self type with the Self type from typing_extensions.

---------
Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>

* Phase out attr

* Remove unused dependencies from requirements files

* Update TVM submodule to latest commit 4776d31

* Update Installation.md

* Update installation prerequisites in documentation

* cu128 Docker Image

---------
Co-authored-by: Lei Wang <34334180+LeiWang1999@users.noreply.github.com>
Co-authored-by: Andy <andyluo918@gmail.com>

* Update TVM submodule and enhance vectorization logic in loop_vectorize.cc

- Updated the TVM submodule to the latest commit.
- Simplified the vectorization process by ensuring that the vectorized expression is simplified after vectorization, improving expression handling.
- Added checks in loop_fusion_utils.h to prevent fusion of loops with non-power-of-2 extents, enhancing robustness in loop transformations.

* lint fix

* Update Installation.md

* Update installation prerequisites in documentation

---------
Co-authored-by: Lei Wang <34334180+LeiWang1999@users.noreply.github.com>

* Update CI configuration to run pytest with automatic parallelization using the '-n auto' option.

* Enhance Cython JIT Adapter Compilation Logic

- Improved the locking mechanism during the compilation of the Cython JIT adapter to prevent race conditions.
- Added checks to determine if another process has already compiled the library, reducing unnecessary recompilation.
- Cleaned up the code by removing redundant imports and ensuring proper handling of temporary files during compilation failures.
- Updated vectorization logic in loop_vectorize.cc to allow optional simplification of vectorized expressions.

This update enhances performance and reliability in the JIT compilation process.

* lint fix

* Update CI configuration to run pytest with 4 parallel jobs instead of auto-detection

* Add pytest markers for serial execution in MHA tests

- Added @pytest.mark.serial to multiple MHA test functions to ensure they run sequentially.
- This change improves test reliability by preventing potential race conditions during execution.

* Update TVM submodule and enhance vectorization logic in loop_vectorize.cc

- Updated the TVM submodule to the latest commit.
- Modified the vectorization logic to include optional simplification of vectorized expressions and added checks to ensure the usage of vectorized variables, improving performance and reliability in expression handling.

* Remove @pytest.mark.serial from multiple MHA test functions to allow parallel execution. This change enhances test performance by enabling concurrent test runs while maintaining reliability.

* Remove tvm_simplify_test.py file, eliminating the test for expression simplification in TVM. This cleanup helps streamline the codebase by removing unused test cases.

* Remove unused pytest import from test_tilelang_kernel_mha.py to streamline the test file.

* lint fix

* Update TVM submodule and refine vectorization logic in loop_vectorize.cc

- Updated the TVM submodule to the latest commit.
- Adjusted the return statements in loop_vectorize.cc to improve expression handling and ensure consistency in the visitor pattern.

* Refactor vectorization logic in loop_vectorize.cc

- Removed the check for the usage of the vectorized variable in the vectorization logic, simplifying the expression handling.
- This change enhances the clarity and efficiency of the vectorization process.

* Enhance vectorization checks in loop_vectorize.cc

- Added a check to ensure the vectorized expression uses the vectorized variable, improving the robustness of the vectorization logic.
- This change refines the expression handling and ensures that only valid vectorized expressions are processed.

* Implement non-local buffer checks for loop vectorization in layout_inference.cc

- Added logic to check for non-local buffer loads and stores before applying vectorization to loops. This enhancement ensures that vectorization is only applied when appropriate, improving the correctness of the loop transformations.

* Refactor buffer handling in pipeline planning and layout inference

- Renamed GlobalCopyPatternDetector to BufferRegionCollector for clarity and updated its logic to collect buffer read/write regions.
- Enhanced the handling of conditional expressions in pipeline planning, allowing for better management of stages related to conditional statements.
- Improved the processing of buffer regions during read/write operations, ensuring accurate tracking of buffer usage across different stages.

* Refactor vectorization checks in loop_vectorize.cc

- Removed the check for the usage of the vectorized variable in the vectorization logic, simplifying the expression handling.
- This change enhances the clarity and efficiency of the vectorization process, ensuring that valid vectorized expressions are processed without unnecessary checks.

* Add bf16 support for AMD in quickstart example

* Reduced git diff

* Move bf16 vector definition into common.h

* Added unit tests for basic AMD bf16 matmul

* lint fix

---------
Co-authored-by: OscarSavNS <oscar.savolainen@nscale.com>
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>

* [Enhancement] Move T.any_of and T.all_of op registration from python into cpp

* format

* add license

* [BugFix] Address should aligned with access size in tail split

* Lint

* Lint

Bumps [transformers](https://github.com/huggingface/transformers) from 4.40 to 4.48.0.
- [Release notes](https://github.com/huggingface/transformers/releases)
- [Commits](https://github.com/huggingface/transformers/compare/v4.40.0...v4.48.0

)

---
updated-dependencies:
- dependency-name: transformers
  dependency-version: 4.48.0
  dependency-type: direct:production
...
Signed-off-by: dependabot[bot] <support@github.com>
Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>

* Update copyright notice in example_mha_bwd_wgmma_pipelined.py to reflect Tile-AI Corporation ownership.

* lint fix

* make it python 3.8- happy

* [Enhancement] Improve loop partitioning and vectorization logic in layout inference and loop vectorization

- Enhanced the VisitStmt_ method to support local buffer handling in parallel loops, allowing for register usage without explicit thread binding.
- Updated loop vectorization logic to simplify expressions and ensure accurate vector size calculations, improving performance and clarity in the vectorization process.

* lint fix

* [Refactor] Update warp size checks and enhance warp partitioning logic in GEMM

- Changed warp_n size check from 16 to 8 in gemm_layouts.cc to improve compatibility with specific configurations.
- Refactored warp partitioning logic in gemm.cc to prioritize N dimension for better performance based on aspect ratio.
- Introduced a new CompileArgs dataclass in autotuner to streamline compile argument management and improve code clarity.

* lint fix

* [Enhancement] Initialize jit_compile in AutoTuner class

- Added initialization for jit_compile attribute in the AutoTuner class to ensure it is set to None by default.
- Updated the assignment logic for jit_compile to prevent overwriting an existing compile function, enhancing the flexibility of the AutoTuner's compilation process.

* make it python 3.8- happy

* [Enhancement] Improve loop partitioning and vectorization logic in layout inference and loop vectorization

- Enhanced the VisitStmt_ method to support local buffer handling in parallel loops, allowing for register usage without explicit thread binding.
- Updated loop vectorization logic to simplify expressions and ensure accurate vector size calculations, improving performance and clarity in the vectorization process.

* lint fix

Added detailed error messages in the InferLayout method to provide better context when layout conflicts occur. This includes the body of the operation that triggered the error, aiding in debugging and layout validation.

Updated SyncPatternMap to use vectors for acquire and release, enhancing flexibility in handling synchronization patterns. Improved barrier handling logic for both producer and consumer cases, ensuring accurate synchronization in the pipeline.

* [Enhancement][Pipeline] Improve pipeline stage information handling and copy stage detection

- Added detailed documentation for the PipelineStageInfo structure to clarify its parameters.
- Enhanced the VisitStmt_ method to handle annotations for pipeline order and stage more effectively.
- Implemented logic to determine if a stage is used by a copy operation, adjusting the stage assignment accordingly.
- Processed the tail copy stage to ensure correct ordering and stage assignment in the pipeline planning process.

* lint fix

* Update README.md for deepseek_mla: Refine performance comparison details and add acknowledgment section. Adjusted performance metrics for TileLang, highlighting its efficiency over Triton and assembly kernels. Included gratitude to the AMD ROCm team for their contributions.

* Update README.md for deepseek_mla: Clarify performance metrics for TileLang, specifying the range of performance parity with hand-optimized assembly kernels. This adjustment enhances the accuracy of the comparative analysis against Triton implementations.

[Dynamic Symbolic] Add pass_config to customize vectorization and tail split
[Pytest Fix] Wrap tests in dynamic benchmark

* [Dynamic Symbolic] Add pass_config to customize vectorization and tail split

* Lint

* Only check for vectorized dimension. Add docs.

* Lint

* Update comment for cache directory in .gitignore

* Use CUTLASS convention to represent dynamic alignment. Fix bugs

* Add benchmark examples

* Add more benchmarks. Fix accumulate type bug.

* Lint

* Lint

* Test Lint

* Lint

* Test Lint

* Lint

* Fix typo

* Lint

* Lint

---------
Co-authored-by: Lei Wang <34334180+LeiWang1999@users.noreply.github.com>

* fix

* Update submodule TVM to latest commit and enhance FlattenBuffer pass in TileLang engine. Added boolean handling in buffer loading and improved address_of detection in flattening logic.

* lint fix

* Update legalize_safe_memory_access.cc

* Add cache path handling and file locking in Cython adapter

- Introduced a new cache path based on the code hash for the Cython JIT adapter, enhancing cache management.
- Added a lock file mechanism to ensure safe access during cache operations, improving concurrency handling.
- These changes aim to optimize the compilation process and prevent race conditions during library loading.

* lint fix

* refactor

* refactor

* Add GlobalCopyPatternDetector to identify global memory copy patterns

- Introduced a new class, GlobalCopyPatternDetector, to detect specific memory copy patterns in statements.
- Enhanced the PipelinePlanner to utilize this detector for determining copy stages based on global and local memory scopes.
- Improved code clarity and maintainability by encapsulating detection logic within the new class.

* Refactor copy stage detection logic in pipeline planning

- Simplified the determination of copy stages by directly assigning the result of GlobalCopyPatternDetector to pinfo.copy_stage.
- Removed redundant checks for read and write scopes, enhancing code clarity and maintainability.

* lint fix

* Update legalize_safe_memory_access.cc

* Add cache path handling and file locking in Cython adapter

- Introduced a new cache path based on the code hash for the Cython JIT adapter, enhancing cache management.
- Added a lock file mechanism to ensure safe access during cache operations, improving concurrency handling.
- These changes aim to optimize the compilation process and prevent race conditions during library loading.

* lint fix

* Add deepseek_mla to documentation index (#380)

* lint fix

* Remove debug print statement from OptimizeForTarget function and enhance library loading mechanism in Cython adapter. Implemented file locking during cache access and added checks for library size before loading. Introduced temporary file handling for safer compilation of Cython JIT adapter.

* Update comments in Cython adapter for clarity and consistency. Changed Chinese comments to English for better readability and understanding of the code's functionality, specifically regarding file handling and compilation processes.

* Refactor comments in Cython adapter for improved clarity. Updated comment on cache file deletion for consistency and removed unnecessary whitespace in file handling section.

Remove debug print statement from OptimizeForTarget function and enhance library loading mechanism in Cython adapter. Implemented file locking during cache access and added checks for library size before loading. Introduced temporary file handling for safer compilation of Cython JIT adapter. (#377)

* [Add] Introduce deepseek_mla documentation for high-performance FlashMLA with TileLang

- Added a comprehensive guide on writing high-performance kernels using TileLang, focusing on the Multi-Head Latent Attention (MLA) mechanism.
- Included benchmark results comparing FlashMLA, TileLang, Torch, Triton, and FlashInfer, highlighting TileLang's efficiency and ease of use.
- Detailed implementation strategies, including layout inference, threadblock swizzling, shared memory swizzling, and warp specialization.
- Provided examples and explanations of optimization techniques to enhance performance in GPU kernel programming.

* doc update

* [Add] Enhance AMD FlashMLA implementation and documentation

- Refactored variable names in `benchmark_mla_decode_amd_tilelang.py` for clarity, changing `Q_shared` and `Q_pe_shared` to `Q_local` and `Q_pe_local` to reflect their usage in register allocation.
- Added a new `README.md` detailing the high-performance FlashMLA implementation on AMD MI300X accelerators, including architectural considerations, optimization strategies, and performance evaluation.
- Introduced a performance comparison figure to illustrate the efficiency of the TileLang implementation against other frameworks.

* lint fix

* [Add] Expand deepseek_mla documentation for AMD MI300X optimization strategies

- Introduced a new section detailing architectural differences and optimization strategies for implementing FlashMLA on AMD MI300X accelerators.
- Highlighted key considerations such as instruction set variations, shared memory constraints, tile size flexibility, and memory bank conflict swizzling.
- Included performance evaluation results demonstrating TileLang's efficiency compared to other frameworks.
- Discussed future optimization opportunities for memory bank conflict mitigation and dimension parallelization.

* [Enhancement] Add variable check in GlobalMemChecker for safe memory access validation

- Introduced a check in the GlobalMemChecker to determine if the index used in memory access has any variable components, enhancing the safety of memory access validation.
- Updated the condition handling in store operations to ensure that only boolean conditions are processed, improving type safety and error handling in memory operations.

* [Refactor] Rename VecAllocAccess to TLVecAllocAccess and enhance buffer access handling

- Renamed the `VecAllocAccess` class to `TLVecAllocAccess` for clarity in its purpose.
- Improved the handling of buffer access by mutating extents and rewriting access in the body, ensuring compatibility with vectorized operations.
- Added a TODO comment to suggest moving this pass to occur before StorageFlatten/FlattenBuffer for better optimization.
- Introduced a print statement in the phase optimization process for debugging purposes.

* lint fix

* add hip setup support

* add env.find_hip func

* Delete install_hip.sh as we already have install_rocm.sh

* modify hip to rocm

---------
Co-authored-by: Lei Wang <34334180+LeiWang1999@users.noreply.github.com>

* [Enhancement] Introduce logical operations `any_of` and `all_of` for buffer checks

- Added new logical operations `any_of` and `all_of` to the TileLang language interface, allowing users to check conditions across buffer elements.
- Implemented corresponding intrinsic calls for CUDA, enhancing the functionality of the TileLang framework.
- Updated the `allocate.py` to handle boolean types correctly in shared memory allocations.
- Introduced tests for the new logical operations to ensure correctness and performance.
Co-authored-by: Zhiwen Mo <zhiwen.mo25@ic.ac.uk>

* lint fix

---------
Co-authored-by: Zhiwen Mo <zhiwen.mo25@ic.ac.uk>

* [Bugfix] Adjust Autotuner threadpool `max_workers` limit to available CPUs

* [Example] Small fix on example_blocksparse_gemm.py

* [Add] Introduce benchmark scripts for MLA decoding with AMD support

- Added three new benchmark scripts: `benchmark_mla_decode_amd_tilelang.py`, `benchmark_mla_decode_amd_torch.py`, and `benchmark_mla_decode_amd_triton.py` to evaluate the performance of the MLA decoding mechanism across different frameworks.
- Each script includes implementations for attention calculation, performance profiling, and output validation against reference implementations.
- Enhanced command-line argument parsing for customizable input parameters, including batch size, number of heads, and dimensions.
- Integrated performance comparison functionality to facilitate benchmarking between different implementations.

* lint fix

* lint fix

---------
Co-authored-by: Zhiwen Mo <zhiwen.mo25@ic.ac.uk>

[Enhancement] Update kernel declaration pattern to support launch bounds in match_declare_kernel function (#366)

- Modified the regex pattern in `match_declare_kernel` to accommodate optional `__launch_bounds__` specifications, enhancing the function's ability to match kernel declarations accurately.
- This change improves the flexibility of kernel matching in the source code, allowing for more complex kernel definitions.
Co-authored-by: Zhiwen Mo <zhiwen.mo25@ic.ac.uk>

* [Refactor] Update AutoTuner run method and timeout handling

- Modified the `run` method to reduce the default timeout from 100 to 30 seconds for improved responsiveness.
- Changed the `get_input_tensors_supply` call to disable output generation, enhancing performance during tensor supply retrieval.
- Refactored the latency measurement to streamline the benchmarking process, ensuring proper timeout handling with `ThreadPoolExecutor`.
- Added logging for timeout occurrences to aid in debugging and performance analysis.

* bug fix

* lint fix

* [Bugfix] Correct dynamic shared memory size error handling in HIP wrapper

- Updated the error handling logic in `PREDEF_ATTRIBUTE_SET_DYNAMIC_MEMORY_HIP` to check if the dynamic shared memory size exceeds the maximum limit of 65536.
- Improved error message clarity by specifying the function name and the attempted size, ensuring better debugging information.
- Ensured the function returns 0 upon successful setting of the dynamic shared memory size.

* [Add] Implement example for MLA decoding with AMD support

- Introduced a new example script `example_mla_decode_amd.py` demonstrating the use of the flash attention mechanism with AMD hardware.
- Implemented functions for attention calculation, including support for split processing and combining outputs.
- Added command-line argument parsing for customizable input parameters such as batch size, number of heads, and dimensions.
- Included a reference implementation for validation against the Tile-AI output, ensuring correctness of the implementation.
- Enhanced performance profiling and output comparison for debugging and optimization purposes.

* lint fix

- Updated the error handling logic in `PREDEF_ATTRIBUTE_SET_DYNAMIC_MEMORY_HIP` to check if the dynamic shared memory size exceeds the maximum limit of 65536.
- Improved error message clarity by specifying the function name and the attempted size, ensuring better debugging information.
- Ensured the function returns 0 upon successful setting of the dynamic shared memory size.

[Example] Handle Scenarios in Which a Threadblock is Assigned Only Invalid Block Indices for Sparse Attention  (#361)

* Fix issue where threadblock with only invalid blocks produces incorrect output.

* fix score scale

* format