[MLIR][NVVM] Add `inline_ptx` op (#139923)

This op allows using PTX directly within the NVVM dialect, while greatly
simplifying llvm.inline_asm generation.

**Example 1: Read-only Parameters**

Sets `"l,r"` automatically. 
```
nvvm.inline_ptx "mbarrier.init.b64 [$0], $1;" (%barrier_gen, %count) : !llvm.ptr, i32

// Lowers to:
llvm.inline_asm has_side_effects asm_dialect = att
      "mbarrier.init.b64 [$0], $1;", "l,r" %arg0, %arg2 : (!llvm.ptr, i32) -> ()
```

**Example 2: Read-only and Write-only Parameters**

Sets `=f,f"` automatically. `=` is set because there is store. 
```
%0 = nvvm.inline_ptx "ex2.approx.ftz.f32 $0, $1;" (%input) : f32 -> f32

// Lowers to:
%0 = llvm.inline_asm has_side_effects asm_dialect = att "ex2.approx.ftz.f32 $0, $1;", "=f,f" %arg0 : (f32) -> f32
```

**Example 3: Predicate Usage**

Now `@$2` is set automatically for predication. 
```
nvvm.inline_ptx "mbarrier.init.b64 [$0], $1;" (%barrier_gen, %count), predicate = %pred : !llvm.ptr, i32, i1

// Lowers to:
llvm.inline_asm has_side_effects asm_dialect = att "@$2 mbarrier.init.b64 [$0], $1;", "l,r,b" %arg0, %arg2, %arg3 : (!llvm.ptr, i32, i1) -> ()
```

---------

Co-authored-by: Mehdi Amini <joker.eph@gmail.com>

0 files changed, 0 insertions, 0 deletions