[Mlir-commits] [mlir] [MLIR][NVVM] Add tensormap.replace NVVM Op (PR #174926)

[Srinivasa Ravi]

================
@@ -4775,6 +4848,50 @@ PermuteOp::getIntrinsicIDAndArgs(Operation &op, LLVM::ModuleTranslation &mt,
   return {IDs[modeIndex], args};
 }
 
+mlir::NVVM::IDArgPair TensormapReplaceOp::getIntrinsicIDAndArgs(
+    Operation &op, LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder) {
+  auto thisOp = cast<NVVM::TensormapReplaceOp>(op);
+
+  llvm::SmallVector<llvm::Value *> args;
+  args.push_back(mt.lookupValue(thisOp.getAddr()));
+  if (thisOp.getOrd())
+    args.push_back(builder.getInt32(thisOp.getOrd().value()));
+  if (thisOp.getNewValue())
+    args.push_back(mt.lookupValue(thisOp.getNewValue()));
+  if (auto attr = thisOp.getNewValueAttr()) {
+    auto val =
+        llvm::TypeSwitch<mlir::Attribute, unsigned>(*attr)
+            .Case<TensormapElemtypeAttr, TensormapInterleaveLayoutAttr,
+                  TensormapSwizzleModeAttr, TensormapSwizzleAtomicityAttr,
+                  TensormapFillModeAttr>([](auto attr) {
+              return static_cast<unsigned>(attr.getValue());
+            })
+            .Default([](auto attr) {
+              llvm_unreachable("Invalid attribute type");
+              return 0;
+            });
+    args.push_back(builder.getInt32(val));
+  }
+
+  static constexpr llvm::Intrinsic::ID IDs[] = {
+      llvm::Intrinsic::nvvm_tensormap_replace_global_address,
+      llvm::Intrinsic::nvvm_tensormap_replace_rank,
+      llvm::Intrinsic::nvvm_tensormap_replace_box_dim,
+      llvm::Intrinsic::nvvm_tensormap_replace_global_dim,
+      llvm::Intrinsic::nvvm_tensormap_replace_global_stride,
+      llvm::Intrinsic::nvvm_tensormap_replace_element_stride,
+      llvm::Intrinsic::nvvm_tensormap_replace_elemtype,
+      llvm::Intrinsic::nvvm_tensormap_replace_interleave_layout,
+      llvm::Intrinsic::nvvm_tensormap_replace_swizzle_mode,
+      llvm::Intrinsic::nvvm_tensormap_replace_swizzle_atomicity,
+      llvm::Intrinsic::nvvm_tensormap_replace_fill_mode,
+  };
+
+  unsigned fieldIndex = static_cast<unsigned>(thisOp.getField());
----------------
Wolfram70 wrote:

+1 for looking into this. Sounds like it'd simplify the Op->intrinsic lowering for a lot of Ops.

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