* refactor code
* enhance tutorial
* Enhance error handling and code generation in CUDA and TileLang components
This commit introduces several improvements across multiple files:
- Added more informative error messages in GEMM layout checks
- Updated CUDA codegen to support more flexible function signature generation
- Improved TMA descriptor initialization and kernel dispatch logic
- Refined library generation and source code parsing utilities
- Enhanced error handling in various adapter and wrapper classes
* Add thread tag validation for warp specialization
Introduce a ThreadTagChecker to validate that a PrimFunc only uses threadIdx.x before applying warp specialization. This prevents unintended transformations on kernels with complex thread binding and provides a clear warning to users about potential issues with warp specialization.
* Update TileLang Profiling and Compilation in Flash Decoding Examples
Refactor the profiling and compilation workflow in two flash decoding example scripts:
- Replace `tilelang.lower()` and `tilelang.Profiler()` with `tilelang.compile()`
- Simplify profiler initialization using `get_profiler()`
- Update method calls to use the new profiler and compiled kernel objects
- Maintain existing performance benchmarking and validation logic
* Refactor and clean up code formatting in TileLang testing and adapter modules
This commit includes several code style and formatting improvements:
- Adjust whitespace and line breaks in test files
- Improve code formatting in CUDA source wrapper and adapter utilities
- Enhance readability of function calls and argument handling
- Remove unnecessary whitespace and standardize indentation
- Simplify function signatures and argument parsing
* Refactor CUDA codegen and improve code formatting
This commit includes several improvements to CUDA code generation and formatting:
- Enhance function signature generation in CodeGenTileLangCUDA
- Improve code formatting and readability in CUDA-related files
- Simplify parameter handling and type annotations
- Clean up whitespace and line breaks in codegen and layout files
---------
Co-authored-by: Ubuntu <dlisuser@h100testl730RPS.xu5snccwrbtejcqqalluoku5hb.xx.internal.cloudapp.net>
* [Feature] Add CTypes JIT kernel support for dynamic shapes and multi-stream execution
- Enhance CtypesKernelAdapter to handle dynamic symbolic shapes
- Add support for multi-stream kernel execution in CTypes backend
- Implement dynamic shape handling in test_tilelang_jit_gemm_ctypes.py
- Add symbolic shape utility function in tilelang.language
- Update profiler to improve flexibility in benchmark selection
* Remove redundant thread binding in GEMM kernel implementations
- Remove unnecessary `thread_binding` line in GEMM kernel functions
- Clean up code in `examples/gemm/README.md` and `testing/python/kernel/test_tilelang_kernel_int4_gemm_mma.py`
- Enhance code readability by removing redundant thread binding annotation
* Fix indentation in int4 GEMM kernel test file
- Correct indentation for function calls in `test_tilelang_kernel_int4_gemm_mma.py`
- Remove extra indentation in `mma_emitter.ldmatrix_a()` and `mma_emitter.ldmatrix_b()` calls
- Improve code formatting for better readability
* [CI][Test] Add test cases for tilelang transform MultiVersionBuffer and WarpSpecialized
* Relax the mismatch ratio restrictions in the flash_linear_attention and mha tests
* [Enhancement] Add VectorizeLoop function and update imports for compatibility
* [CI][Test] Improve test cases for vectorization and fix typos in parser comments
* lint fix
* Fix incorrect module reference for VectorizeLoop transformation
* Refactor vectorize_loop transformation by removing unused extent mutation logic
* [Enhancement] Add support for FP8 data types and global barriers in CUDA codegen
* Fix formatting in CUDA FP8 header file for consistency
* Refactor CI workflow to use 'tilelang_ci' virtual environment and update CUDA type printing for better clarity
* Update submodule 'tvm' to latest commit for improved functionality
* Refactor execution backend references from 'dl_pack' to 'dlpack' for consistency and clarity; add apply_simplify function to simplify PrimFunc or IRModule.
* Refactor CUDA code for improved readability; clean up formatting and remove unnecessary whitespace in multiple files.
* Refactor import statement in test_tilelang_kernel_dequantize_gemm.py to use 'tilelang.language' for consistency
* Add CUDA requirements to FP8 test cases and update references for clarity
* Add a blank line for improved readability in test_tilelang_kernel_fp8_gemm_mma.py
* Fix data type in reference result calculation for consistency in test_tilelang_kernel_gemm_mma_intrinsic.py
* Add CUDA requirements and FP8 test cases for matmul and gemv simulations
* Remove debug print statements and use tilelang's testing assertion for result validation in test_tilelang_kernel_gemm_mma_intrinsic.py
* Remove outdated comment regarding FP8 tests in test_tilelang_kernel_gemv_simt.py
* Add BF16 support to matrix multiplication and introduce corresponding test cases
* Add a blank line for improved readability in BF16 GEMM test
* Update acknowledgements in README to include supervision by Zhi Yang at Peking University
* [Enhancement] Add VectorizeLoop function and update imports for compatibility
* [CI][Test] Improve test cases for vectorization and fix typos in parser comments
* lint fix
* Fix incorrect module reference for VectorizeLoop transformation
* Refactor vectorize_loop transformation by removing unused extent mutation logic
* [Enhancement] Add support for FP8 data types and global barriers in CUDA codegen
* Fix formatting in CUDA FP8 header file for consistency
* Refactor CI workflow to use 'tilelang_ci' virtual environment and update CUDA type printing for better clarity
* Update submodule 'tvm' to latest commit for improved functionality
* Refactor execution backend references from 'dl_pack' to 'dlpack' for consistency and clarity; add apply_simplify function to simplify PrimFunc or IRModule.
* Refactor CUDA code for improved readability; clean up formatting and remove unnecessary whitespace in multiple files.
* Refactor import statement in test_tilelang_kernel_dequantize_gemm.py to use 'tilelang.language' for consistency
* Add CUDA requirements to FP8 test cases and update references for clarity
* Add a blank line for improved readability in test_tilelang_kernel_fp8_gemm_mma.py
* Fix data type in reference result calculation for consistency in test_tilelang_kernel_gemm_mma_intrinsic.py
* Add CUDA requirements and FP8 test cases for matmul and gemv simulations
* Remove debug print statements and use tilelang's testing assertion for result validation in test_tilelang_kernel_gemm_mma_intrinsic.py
* Remove outdated comment regarding FP8 tests in test_tilelang_kernel_gemv_simt.py
* implement jit test case
* [Dev] implement auto tune test case for matrix multiplication
* Implement test for legalize memory access and vectorized loop
* lint fix
* introduce run_once
* Refactor callback function names for consistency and improve code readability
* enhance documentations
* lint fix
* lint fix
* lint fix
* lint fix
* fix formatting issues in rt_mod_hip.cc
* add random seed initialization for deterministic testing
* [Dev] Add FlashDecoding example
* [CI][Test] Add test cases for tilelang kernel convolution
* [CI][Test] Add test cases for tilelang kernel FlashAttention
* Reduce the number of stages to ensure the shared memory allocation is valid
* Temporarily remove the dim128 case
* lint
* update einops in requirements-dev.txt
* update einops in requirements-test.txt
* remove einops in requirements-dev.txt