[MLIR] Replace LLVM_Type in bar.warp.sync and cp.async ops with I32 #167826
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
This patch replaces generic LLVM_Type with specific I32 type in NVVM operations. NVVM_SyncWarpOp: Change mask parameter from LLVM_Type to I32 NVVM_CpAsyncOp: Change cpSize parameter from Optional<LLVM_Type> to Optional<I32> Signed-off-by: Dharuni R Acharya <[email protected]>
Member
|
@llvm/pr-subscribers-mlir @llvm/pr-subscribers-mlir-llvm Author: Dharuni R Acharya (DharuniRAcharya) ChangesThis patch replaces generic
Full diff: https://github.com/llvm/llvm-project/pull/167826.diff 1 Files Affected:
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 4c13c5ddb2886..1c30d754a1792 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -1387,7 +1387,7 @@ def NVVM_VoteSyncOp
def NVVM_SyncWarpOp :
NVVM_Op<"bar.warp.sync">,
- Arguments<(ins LLVM_Type:$mask)> {
+ Arguments<(ins I32:$mask)> {
let summary = "Warp Barrier Synchronization Op";
let description = [{
The `nvvm.bar.warp.sync` operation performs barrier synchronization for threads
@@ -1473,7 +1473,7 @@ def NVVM_CpAsyncOp : NVVM_Op<"cp.async.shared.global">,
LLVM_PointerGlobal:$src,
I32Attr:$size,
LoadCacheModifierAttr:$modifier,
- Optional<LLVM_Type>:$cpSize)> {
+ Optional<I32>:$cpSize)> {
let assemblyFormat = "$dst `,` $src `,` $size `,` `cache` `=` $modifier (`,` $cpSize^)? attr-dict `:` type(operands)";
let hasVerifier = 1;
let extraClassDeclaration = [{
|
durga4github
approved these changes
Nov 13, 2025
Contributor
durga4github
left a comment
durga4github
left a comment
There was a problem hiding this comment.
LGTM. Thanks for the fix!
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.
This patch replaces generic
LLVM_Typewith specificI32type in NVVM operations.NVVM_SyncWarpOp: Change mask parameter fromLLVM_TypetoI32.NVVM_CpAsyncOp: Change cpSize parameter fromOptional<LLVM_Type>toOptional<I32>.