[Dev][TL] Integrate TL Dequant Implementation into BitBLAS OPs #214
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
The select_scheduler function in the dense/__init__.py module has been refactored to use a fine-grained interface. This change provides more flexibility and enables the implementation of high-performance kernels. Update MatmulScheduler class in matmul_tensorcore.py The MatmulScheduler class in the matmul_tensorcore.py module has been updated to calculate the number of threads based on the block size and warp size. This ensures optimal GPU warp configuration for NVIDIA GPUs. Improve test_general_matmul_tilelang_kernel.py The test_general_matmul_tilelang_kernel.py module has been improved to include additional test cases and assertions for correctness.
…_tilelang_kernel.py to use centered random values for input tensors
|
Though TL enables more program flexibility for us to write kernels, but it's hard to implement all dequant kernels within a simple template(as our tir schedule based template) |
This was referenced Oct 7, 2024
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 includes significant changes to the
bitblaslibrary, mainly focusing on the addition of new schedulers, improvements to matrix multiplication operations, and updates to testing and dependencies. The most important changes are grouped into themes below.New Schedulers and Enhancements:
float16andint8target data types inget_lop3_intrin_groupfunction inbitblas/gpu/intrin/lop3.py. ([bitblas/gpu/intrin/lop3.pyR1680-R1690](https://github.com/microsoft/BitBLAS/pull/214/files#diff-15fd74b90c3b956e9864e35778f26b27f6c9a7cfae35037967f420fda9a0bbe5R1680-R1690))bitblas/ops/general_matmul/__init__.pyand updated the_select_schedulermethod to return this new scheduler. ([[1]](https://github.com/microsoft/BitBLAS/pull/214/files#diff-74fe5dd2824cb03a0fb2b0a913a2fc5caeb9c08e5368c318cd32b3af7e6f52edR15),[[2]](https://github.com/microsoft/BitBLAS/pull/214/files#diff-74fe5dd2824cb03a0fb2b0a913a2fc5caeb9c08e5368c318cd32b3af7e6f52edL594-R614))MatmulDequantizeSchedulerand related functions inbitblas/ops/general_matmul/tilelang/dequantize/__init__.py. ([bitblas/ops/general_matmul/tilelang/dequantize/__init__.pyR3-R102](https://github.com/microsoft/BitBLAS/pull/214/files#diff-422bb6fd30915da2280e418fe97aab5bdf246b548321577b128cd1652bd68ec2R3-R102))Code Simplification and Refactoring:
mainfunction inbitblas/ops/general_matmul/tilelang/dense/matmul_tensorcore.pyfor better readability and maintainability. ([[1]](https://github.com/microsoft/BitBLAS/pull/214/files#diff-eacc57f40c9e810b3503e297685ed5eb8372c922201e18a4be1c1c2e20c93353L515-R522),[[2]](https://github.com/microsoft/BitBLAS/pull/214/files#diff-eacc57f40c9e810b3503e297685ed5eb8372c922201e18a4be1c1c2e20c93353L652-R662),[[3]](https://github.com/microsoft/BitBLAS/pull/214/files#diff-eacc57f40c9e810b3503e297685ed5eb8372c922201e18a4be1c1c2e20c93353L686-R700),[[4]](https://github.com/microsoft/BitBLAS/pull/214/files#diff-eacc57f40c9e810b3503e297685ed5eb8372c922201e18a4be1c1c2e20c93353L867-R884),[[5]](https://github.com/microsoft/BitBLAS/pull/214/files#diff-eacc57f40c9e810b3503e297685ed5eb8372c922201e18a4be1c1c2e20c93353L979-R999),[[6]](https://github.com/microsoft/BitBLAS/pull/214/files#diff-eacc57f40c9e810b3503e297685ed5eb8372c922201e18a4be1c1c2e20c93353L1009-R1033))__repr__method inbitblas/ops/general_matmul/tilelang/dense/matmul_tensorcore.pyto includewarp_Mandwarp_Ndetails. ([bitblas/ops/general_matmul/tilelang/dense/matmul_tensorcore.pyR320-R321](https://github.com/microsoft/BitBLAS/pull/214/files#diff-eacc57f40c9e810b3503e297685ed5eb8372c922201e18a4be1c1c2e20c93353R320-R321))Testing Enhancements:
matmul_torch_forwardandmatmul_torch_forward_dequanttotesting/python/operators/test_general_matmul_ops_backend_tl.pyfor validating matrix multiplication operations with and without dequantization. ([[1]](https://github.com/microsoft/BitBLAS/pull/214/files#diff-cb0d29b36116f888480a1c2a5cee67a69ad3d5434522ce51fb92731242adc2cfR78-R227),[[2]](https://github.com/microsoft/BitBLAS/pull/214/files#diff-cb0d29b36116f888480a1c2a5cee67a69ad3d5434522ce51fb92731242adc2cfR248-R259))MatmulDequantizeSchedulerin the imports oftesting/python/operators/test_general_matmul_tilelang_kernel.py. ([testing/python/operators/test_general_matmul_tilelang_kernel.pyR13-R15](https://github.com/microsoft/BitBLAS/pull/214/files#diff-052c67c47659338f3612e7c47384c54f88929e9b32c40d9e797b3f5307ff3896R13-R15))Dependency Updates:
yapfversion inrequirements-dev.txtandrequirements-test.txtto0.40.2. ([[1]](https://github.com/microsoft/BitBLAS/pull/214/files#diff-2b4945591edfeaa4cf4d3f155e66d4b43d1bda7a55d881d5cf3107f1b05abbbcL2-R2),[[2]](https://github.com/microsoft/BitBLAS/pull/214/files#diff-685da804fbcac569d75387e475e57d1de687a54c6c41b3aa4057694cfb5abc4bL2-R2))Miscellaneous:
interleave_weightfunction to the imports inbitblas/quantization/__init__.py. ([bitblas/quantization/__init__.pyL12-R16](https://github.com/microsoft/BitBLAS/pull/214/files#diff-aeb54b540a85cbc63bdf9e661a713906e58ddb8c69f56d090bd811e2ba9b4b97L12-R16))