[Mlir-commits] [mlir] [MLIR][NVVM] Update mbarrier.arrive.expect_tx Op (PR #169922)

[Sohaib Iftikhar]

sohaibiftikhar wrote:

This seems to break the TMA tests for sm90.
https://github.com/llvm/llvm-project/tree/main/mlir/test/Integration/GPU/CUDA/sm90

with a stack trace:
```
F0000 00:00:1764743226.477247    2267 logging.cc:51] assert.h assertion failed at mlir/include/mlir/Target/LLVMIR/ModuleTranslation.h:91 in llvm::Value *&mlir::LLVM::ModuleTranslation::mapValue(Value): llvm == nullptr && "attempting to map a value that is already mapped"
*** Check failure stack trace: ***
    @     0x563fe59f2ee4  __assert_fail
    @     0x563fe09bfccc  mlir::LLVM::ModuleTranslation::mapValue()
    @     0x563fe0cd6342  (anonymous namespace)::NVVMDialectLLVMIRTranslationInterface::convertOperation()
    @     0x563fe3fa2d04  mlir::LLVM::ModuleTranslation::convertOperation()
    @     0x563fe3fa35b3  mlir::LLVM::ModuleTranslation::convertBlockImpl()
    @     0x563fe3fa80f0  mlir::LLVM::ModuleTranslation::convertOneFunction()
    @     0x563fe3faac62  mlir::LLVM::ModuleTranslation::convertFunctions()
    @     0x563fe3fadfe5  mlir::translateModuleToLLVMIR()
    @     0x563fe0ca95e1  mlir::LLVM::ModuleToObject::run()
    @     0x563fe0996bc3  mlir::gpu::detail::TargetAttrInterfaceInterfaceTraits::FallbackModel<>::serializeToObject()
    @     0x563fe08cae63  mlir::gpu::transformGpuModulesToBinaries()
    @     0x563fe08cbf38  (anonymous namespace)::GpuModuleToBinaryPass::runOnOperation()
    @     0x563fe47b4f8a  mlir::detail::OpToOpPassAdaptor::run()
    @     0x563fe47b5c4d  mlir::detail::OpToOpPassAdaptor::runPipeline()
    @     0x563fe47bccb1  mlir::PassManager::runPasses()
    @     0x563fe47bc1ca  mlir::PassManager::run()
    @     0x563fde80999c  performActions()
    @     0x563fde808db6  llvm::function_ref<>::callback_fn<>()
    @     0x563fe492c99f  mlir::splitAndProcessBuffer()
    @     0x563fde7fd783  mlir::MlirOptMain()
```

All 4 tma_load tests are affected.

https://github.com/llvm/llvm-project/pull/169922