From fcf56f3152875a913824dfddcd4c74ee42433c04 Mon Sep 17 00:00:00 2001 From: Sohaib Iftikhar Date: Wed, 3 Dec 2025 19:49:31 +0000 Subject: [PATCH] [MLIR][NVVM] Fix lowering logic after fddf7b05 --- mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 7 ++----- 1 file changed, 2 insertions(+), 5 deletions(-) diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td index a96d65d3fcacd..cb83ec23bc76e 100644 --- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td +++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td @@ -948,8 +948,7 @@ def NVVM_MBarrierArriveExpectTxOp : NVVM_PTXBuilder_Op<"mbarrier.arrive.expect_t auto [id, args] = NVVM::MBarrierArriveExpectTxOp::getIntrinsicIDAndArgs( *op, moduleTranslation, builder); - int addrSpace = llvm::cast(op.getAddr().getType()).getAddressSpace(); - if (addrSpace != NVVM::NVVMMemorySpace::SharedCluster) + if (op.getNumResults() > 0) $res = createIntrinsicCall(builder, id, args); else createIntrinsicCall(builder, id, args); @@ -985,9 +984,7 @@ def NVVM_MBarrierArriveDropExpectTxOp : NVVM_Op<"mbarrier.arrive_drop.expect_tx" string llvmBuilder = [{ auto [id, args] = NVVM::MBarrierArriveDropExpectTxOp::getIntrinsicIDAndArgs( *op, moduleTranslation, builder); - - int addrSpace = llvm::cast(op.getAddr().getType()).getAddressSpace(); - if (addrSpace != NVVM::NVVMMemorySpace::SharedCluster) + if (op.getNumResults() > 0) $res = createIntrinsicCall(builder, id, args); else createIntrinsicCall(builder, id, args);