tree 2c0f18ae8abcb3b9587f0ef0257e6d9240ad3cba
parent 6fc031291955d5d3c69f7df9b9f7460c473d114a
author Guray Ozen <guray.ozen@gmail.com> 1747314729 +0200
committer GitHub <noreply@github.com> 1747314729 +0200
gpgsig -----BEGIN PGP SIGNATURE-----
 
 wsFcBAABCAAQBQJoJegpCRC1aQ7uu5UhlAAAZCwQAKAwp71kx7lBxIPKUwwcE4Hn
 mQbXE8CrqRawKXqpqzkRvveLs7yakd7E4mUP7GWTY9bMw+wkg0wK2fIVmIpip83g
 RIqn+7oo20R6UmjcxGGPBRJ5fSpbpYCxVu4oJi8WbKVYMNwxL+k0he4ZQN5U2L8Y
 Cpsy/8KHCnhOmsnuwRgQErYTgp6MAO5UUkq4X4Xftz2qhtq4V07Wr+NpjkWswUNt
 8TSRu8oXwLwWpvlKmY1KlLVqy0GhCv+vSO4h/LFZCWLVGwJEeGkM9YRHs/YX1eZ5
 pnD4+HYFixmAJLPqY7wY+tJU7wjSuFk8I9wSKUcdqxlHnTKLdXj0QqOEyTDMjzmp
 SAKv+ccX73MoxQctt5+siTQDxykzRdoN4TCe5gtmJ9eY4j83xXq+J298GlYXjP/z
 C97xv1pm52/q2lv4xLqQOA42tWZQXVfiVGWgLmgMBLMl4+0EijMrE1IXOcaQMrq7
 I5X8l00KZjQ31RoChhDtaNAfrDvMi9hz/AckMReWAQIvhClqkkpG8rYLboUtvWFF
 F6XLe0vNg31raRgfolx7VJ1M0qnB9SesWMVi1FGwWcP/HKyN4rcXqrTbAVu4vHVl
 0RKtl6b+D2Qm1ssqn68B4Rb3lpNwZLdjyqxdqUtDjpwtAhr0rNb8wAyaKwpw5aJF
 sSINsqKCFiDIwWY8I4X2
 =KSzp
 -----END PGP SIGNATURE-----
 

[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>