mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-21 06:40:23 +08:00
Lei Wang
8bd0293740
* [Enhancement] Add global memory load/store intrinsics for various vector widths (#1710) * Introduced new intrinsics: ldg32, ldg64, ldg128, ldg256 for loading data from global memory. * Added corresponding store intrinsics: stg32, stg64, stg128, stg256 for storing data to global memory. * Updated built-in functions and templates to support these new operations. * Enhanced documentation for usage examples and type checks in the language module. * Updated CUDA code generation to utilize new intrinsics for improved performance. * lint fix * test fix * Refactor buffer type annotations in global memory load/store intrinsics Updated the type annotations for the `ldg` and `stg` functions to use `tir.Buffer` instead of the generic `Buffer` type. This change enhances type specificity and aligns with the TVM type system for better clarity and correctness in memory operations. * lint fix * [Enhancement] Introduce LDG/STG lowering options and transform pass * Added configuration options to enable/disable lowering of non-predicated and predicated global memory loads/stores to ldg/stg intrinsics. * Implemented the LowerLDGSTG transformation pass to optimize global memory access for better performance. * Updated tests to require a minimum CUDA compute version for relevant functionality. * Enhanced documentation for new configuration options and transformation behavior. * lint fix * Refactor LDG/STG intrinsic selection logic for improved readability * Replaced switch-case statements with if-else conditions for determining the appropriate LDG operation based on byte width. * Enhanced error handling for unsupported byte widths with a clear fatal log message. * lint fix * fix * testfix * Enhance LDGSTG pass to skip processing for CuTeDSL backend * Added a check to skip the transformation if the target has the "cutedsl" key. * Updated tests to require a minimum CUDA compute version of 10.0 for LDG and STG code generation tests. * Refactor global memory load/store functions to include conditional variants * Updated the CUDA code generation to support conditional versions of LDG and STG functions, enhancing flexibility in memory operations. * Modified existing functions to include new conditional variants (e.g., `ldg32_conditional`, `stg32_conditional`) for better control over memory access based on predicates. * Adjusted comments to reflect the changes in function signatures and usage. * Update comments and input handling for LDG intrinsics to include address and optional predicate * Modified comments for `ldg32`, `ldg64`, `ldg128`, and `ldg256` to clarify the function signatures, indicating the use of 'address' and an optional 'predicate'. * Changed the number of inputs for these LDG functions from 2 to -1 to accommodate the new signature format. * Add correctness verification to LDG code generation tests * Introduced assertions to verify the correctness of generated CUDA source for `ldg32`, `ldg64`, `ldg128`, and `ldg256` functions. * Added checks to ensure the output matches the expected reference values within specified tolerances. * Updated test cases to enhance validation of LDG intrinsic functionality. * test fix * Refactor LDG and STG intrinsics to use unified naming convention * Updated all instances of `ldg` and `stg` functions to their new names: `load_global` and `store_global`, respectively, across the codebase. * Modified related test cases to reflect the new function names, ensuring consistency in the testing framework. * Enhanced comments and documentation to clarify the changes in function signatures and usage. * Update reference values in LDG256 predicated code generation test to improve accuracy * lint fix * Rename test function for clarity and add new test for predicated load/store global intrinsics. Update assertions to reflect new naming convention in generated kernel source. * lint fix * subtype cast fix