From 0b17416fdb07dbbefad2232943dfdf95fb95faf7 Mon Sep 17 00:00:00 2001
From: Dharuni R Acharya <dharunira@nvidia.com>
Date: Thu, 13 Nov 2025 06:48:08 +0000
Subject: [PATCH] [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 <dharunira@nvidia.com>
---
 mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

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 = [{