[all-commits] [llvm/llvm-project] d08b17: [MLIR][NVVM] Add `inline_ptx` op (#139923)

[Guray Ozen via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: d08b176edc82b8fbd417e0ead7b77abc5f45fce2
      https://github.com/llvm/llvm-project/commit/d08b176edc82b8fbd417e0ead7b77abc5f45fce2
  Author: Guray Ozen <guray.ozen at gmail.com>
  Date:   2025-05-15 (Thu, 15 May 2025)

  Changed paths:
    M mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
    M mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir

  Log Message:
  -----------
  [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 at gmail.com>



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications