Skip to content

Commit 0b17416

Browse files
[MLIR] Replace LLVM_Type in bar.warp.sync and cp.async ops with I32
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]>
1 parent 13251f5 commit 0b17416

File tree

1 file changed

+2
-2
lines changed

1 file changed

+2
-2
lines changed

mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1387,7 +1387,7 @@ def NVVM_VoteSyncOp
13871387

13881388
def NVVM_SyncWarpOp :
13891389
NVVM_Op<"bar.warp.sync">,
1390-
Arguments<(ins LLVM_Type:$mask)> {
1390+
Arguments<(ins I32:$mask)> {
13911391
let summary = "Warp Barrier Synchronization Op";
13921392
let description = [{
13931393
The `nvvm.bar.warp.sync` operation performs barrier synchronization for threads
@@ -1473,7 +1473,7 @@ def NVVM_CpAsyncOp : NVVM_Op<"cp.async.shared.global">,
14731473
LLVM_PointerGlobal:$src,
14741474
I32Attr:$size,
14751475
LoadCacheModifierAttr:$modifier,
1476-
Optional<LLVM_Type>:$cpSize)> {
1476+
Optional<I32>:$cpSize)> {
14771477
let assemblyFormat = "$dst `,` $src `,` $size `,` `cache` `=` $modifier (`,` $cpSize^)? attr-dict `:` type(operands)";
14781478
let hasVerifier = 1;
14791479
let extraClassDeclaration = [{

0 commit comments

Comments
 (0)