Add intrinsics related to DMA operations #3
Conversation
This file contains hidden or 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
add mesh_tensor functions: annotate_mesh_tensor_info, mesh_tensor_copy
For many NMP, memory is often distributed. To support distributed memory, we add a set of helper functions for working with mesh-distributed tensors. With this, we can let users to control the layout of tiles over distributed memory (cores) and add annotations for the compiling passes about the layout. Then, compiler can optimize the computing performance by utilizing the distributed properties of memory and the layout tiles.
…), dma_store(), and create_dma_descriptor(). The relevant files are language/dma.py, src/op/builtin_dma.cc, and src/op/builtin_dma.h. There are also two test files, examples/mytest/testDMAOp.py and examples/mytest/testDMDDescriptor.py, to test whether the registration of the three intrinsics is successful.
Collaborator
JiaqiGuoSunlune
left a comment
JiaqiGuoSunlune
left a comment
There was a problem hiding this comment.
- As discussed on Monday, we may better move test cases to folder "testing/python"
- Please also enable pre-commit by 1)
pip install pre-commit2)pre-commit install. It will automatically correct the coding format for us. - If possible, please add testing oracle (expect) to the test cases (Here is an example from tilelang https://github.com/tile-ai/tilelang/blob/6aaf3c7ae25ebbd8c5f6ad1f45081289234a7943/testing/python/transform/test_tilelang_transform_let_inline.py#L16)
- If I understand the usage correctly, in tilelang DSL, developers can directly use T.dma_load , after lowering to TIR, an dma descriptor will be created automatically?
1 task
Collaborator
Author
yes. since our DMA related ops haven't been implemented, wo can refer to the lower function in the original copy.cc file. |
…g expected results in the test files.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
4 participants
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.
Three intrinsics related to DMA operations have been added: dma_load(), dma_store(), and create_dma_descriptor().
The relevant files are
language/dma.py,src/op/builtin_dma.cc, andsrc/op/builtin_dma.h. This update did not modify the original Tilelang's files.There are also two test files,
examples/mytest/testDMAOp.pyandexamples/mytest/testDMDDescriptor.py, to test whether the registration of the three intrinsics is successful.All files were compiled successfully.