[Feature] Add CTypes JIT kernel support #100
Merged
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
…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 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
- 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
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This pull request introduces several changes to the
tilelangproject, focusing on improving thectypesadapter and adding new benchmarking and testing functionalities. The most important changes include adding new test functions, modifying theCtypesKernelAdapterclass to handle dynamic symbolic shapes, and updating thetilelang.languagemodule to support symbolic variables.Testing Enhancements:
test_ctypes_kernel_do_bench,test_ctypes_kernel_multi_stream, andtest_ctypes_dynamic_shapeto evaluate the performance and correctness of thectypeskernel under different conditions.Adapter Improvements:
CtypesKernelAdapterclass to handle dynamic symbolic shapes by introducing adynamic_symbolic_mapand the_process_dynamic_symbolicmethod.prim_func,srcpath,libpath,lib_code, andis_dynamicto theCtypesKernelAdapterclass to improve code readability and maintainability.Code Simplification:
TLCUDASourceWrapperWithDynamicclass and integrated its functionality directly into theTLCUDASourceWrapperclass, simplifying the codebase. [1] [2]Language Module Enhancements:
symbolicfunction to thetilelang.languagemodule to support the creation of symbolic variables, which can be used in dynamic shape handling.Profiler Enhancements:
determine_profilermethod to theProfilerclass to automatically select the appropriate profiler based on the provided function or module type.