[MLIR][NVVM] Fix lowering logic after fddf7b0 (#170545)

Without this mapping fails when there is no result specified.

See:
#169922 (comment)

To reproduce error on `main`:

```bash
mkdir -p build && cd build
cmake -G Ninja ../llvm \
  -DLLVM_ENABLE_PROJECTS=mlir \
  -DLLVM_TARGETS_TO_BUILD="host;NVPTX" \
  -DMLIR_ENABLE_CUDA_RUNNER=ON \
  -DMLIR_RUN_CUDA_TENSOR_CORE_TESTS=ON \
  -DMLIR_RUN_CUDA_SM90_TESTS=ON \
  -DMLIR_GPU_COMPILATION_TEST_FORMAT=fatbin \
  -DMLIR_INCLUDE_INTEGRATION_TESTS=ON \
  -DLLVM_ENABLE_ASSERTIONS=ON \
  -DLLVM_INSTALL_UTILS=ON \
  -DCMAKE_BUILD_TYPE=Release
ninja

bin/mlir-opt ../mlir/test/Integration/GPU/CUDA/sm90/tma_load_128x64_swizzle128b.mlir -gpu-lower-to-nvvm-pipeline="cubin-chip=sm_90 cubin-features=+ptx80 opt-level=3"
```

Author: sohaibiftikhar

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<LLVMPointerType>(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<LLVMPointerType>(op.getAddr().getType()).getAddressSpace();
-    if (addrSpace != NVVM::NVVMMemorySpace::SharedCluster)
+    if (op.getNumResults() > 0)
       $res = createIntrinsicCall(builder, id, args);
     else
       createIntrinsicCall(builder, id, args);