-
Notifications
You must be signed in to change notification settings - Fork 25
gpu to gpu runtime pass #488
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
gpu to gpu runtime pass #488
Conversation
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Pull request overview
This PR introduces a new transformation pass that converts GPU dialect operations (specifically gpu.launch_func) into Wave runtime library calls. The pass expects IR containing gpu.launch_func ops and transforms them into explicit calls to wave_load_kernel and wave_launch_kernel functions.
Key changes:
- Adds the
WaterGPUtoGPURuntimepass that lowersgpu.launch_funcoperations to runtime calls - Registers the ROCDL dialect in water-opt to support
#rocdl.targetattributes in GPU binary objects - Includes comprehensive tests covering basic functionality, multiple kernel launches, and error validation
Reviewed changes
Copilot reviewed 5 out of 5 changed files in this pull request and generated 9 comments.
Show a summary per file
| File | Description |
|---|---|
| water/lib/Transforms/GPUToGPURuntime.cpp | Implements the GPU to runtime pass with kernel loading and launching logic |
| water/include/water/Transforms/Passes.td | Defines the new water-gpu-to-gpu-runtime pass |
| water/lib/Transforms/CMakeLists.txt | Adds build configuration for the new pass and GPU/LLVM dialect dependencies |
| water/tools/water-opt/water-opt.cpp | Registers ROCDLDialect to support rocdl.target attributes used in tests |
| water/test/Transforms/gpu-to-gpu-runtime.mlir | Provides comprehensive test coverage including error cases |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
| std::string strVal = str.str(); | ||
| strVal.append(std::string_view("\0", 1)); | ||
| return LLVM::createGlobalString( | ||
| loc, builder, getUniqueLLVMGlobalName(mod, symbolTable, varName), | ||
| strVal, LLVM::Linkage::Internal); |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
| std::string strVal = str.str(); | |
| strVal.append(std::string_view("\0", 1)); | |
| return LLVM::createGlobalString( | |
| loc, builder, getUniqueLLVMGlobalName(mod, symbolTable, varName), | |
| strVal, LLVM::Linkage::Internal); | |
| Twine strVal = str + "\0" | |
| return LLVM::createGlobalString( | |
| loc, builder, getUniqueLLVMGlobalName(mod, symbolTable, varName), | |
| strVal.str(), LLVM::Linkage::Internal); |
use Twine instead of std::string_view
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
It doesn't work this way as Twine interprets "\0" just as empty string, need to wrap into StringRef
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Pull request overview
Copilot reviewed 6 out of 6 changed files in this pull request and generated 1 comment.
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
021e0a2 to
6a25b3c
Compare
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Pull request overview
Copilot reviewed 6 out of 6 changed files in this pull request and generated 2 comments.
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
| mlir::LLVM::LLVMDialect, mlir::memref::MemRefDialect, | ||
| mlir::scf::SCFDialect, mlir::vector::VectorDialect, | ||
| wave::WaveDialect>(); | ||
| mlir::LLVM::LLVMDialect, mlir::ROCDL::ROCDLDialect, |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I suspect you also need to modify cmake to link the dialect.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I will add it, but water-opt is already linking with all dialects via get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS)
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Extracted from #452
This pass expects IR with
gpu.launch_funcops and converts it to wave runtime lib calls.