[Mlir-commits] [mlir] [MLIR][NVVM] Add `inline_ptx` op (PR #139923)

[Guray Ozen]

================
@@ -236,6 +236,76 @@ foreach index = !range(0, 32) in {
   def NVVM_EnvReg # index # Op : NVVM_SpecialRegisterOp<"read.ptx.sreg.envreg" # index>;
 }
 
+//===----------------------------------------------------------------------===//
+// Inline PTX op definition
+//===----------------------------------------------------------------------===//
+
+def NVVM_InlinePtxOP : NVVM_Op<"inline_ptx", 
+  [DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>, 
+    AttrSizedOperandSegments]>
+{
+  let summary = "Inline PTX Op";
+  let description = [{This op allows using PTX directly within the NVVM 
+    dialect, while greatly simplifying llvm.inline_asm generation. It 
+    automatically handles register size selection and sets the correct 
+    read/write access for each operand. The operation leverages the 
----------------
grypp wrote:

When I add ReadWrite, then I will add another operand. so it'll probably look like following:
```
nvvm.inline_ptx "mbarrier.init.b64 [$0], $1;" (%barrier_gen, %count), read_write = (%0,%1,%2) : !llvm.ptr, i32, types...
``


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