diff options
author | Guray Ozen <guray.ozen@gmail.com> | 2025-08-18 08:37:55 +0200 |
---|---|---|
committer | GitHub <noreply@github.com> | 2025-08-18 08:37:55 +0200 |
commit | 5d300afa8060c3a4d3c26b5eebc19550ab21e052 (patch) | |
tree | 182b33a3bc1c80d7170cf287c03250eee30c318a /llvm/lib/Target/WebAssembly/Disassembler/WebAssemblyDisassembler.cpp | |
parent | e6e874ce8f055f5b8c5d7f8c7fb0afe764d1d350 (diff) | |
download | llvm-5d300afa8060c3a4d3c26b5eebc19550ab21e052.zip llvm-5d300afa8060c3a4d3c26b5eebc19550ab21e052.tar.gz llvm-5d300afa8060c3a4d3c26b5eebc19550ab21e052.tar.bz2 |
[MLIR][NVVM] Add support for multiple return values in `inline_ptx` (#153774)
This PR adds the ability for `nvvm.inline_ptx` to return multiple
values, matching the expected semantics in PTX while respecting LLVM’s
constraints.
LLVM’s `inline_asm` op does not natively support multiple returns —
instead, it requires packing results into an LLVM `struct` and then
extracting them. This PR implements automatic packing/unpacking so that
multiple return values can be expressed naturally in MLIR without extra
user boilerplate.
**Example**
MLIR:
```
%r1, %r2 = nvvm.inline_ptx "{
.reg .pred p;
setp.ge.s32 p, $2, $3;
selp.s32 $0, $2, $3, p;
selp.s32 $1, $2, $3, !p;
}" (%a, %b) : i32, i32 -> i32, i32
%r3 = llvm.add %r1, %r2 : i32
```
Lowered LLVM IR:
```
%1 = llvm.inline_asm has_side_effects asm_dialect = att "{\0A\09 .reg .pred p;\0A\09 setp.ge.s32 p, $2, $3;\0A\09 selp.s32 $0, $2, $3, p;\0A\09 selp.s32 $1, $2, $3, !p;\0A\09}\0A", "=r,=r,r,r" %a, %b : (i32, i32) -> !llvm.struct<(i32, i32)>
%2 = llvm.extractvalue %1[0] : !llvm.struct<(i32, i32)>
%3 = llvm.extractvalue %1[1] : !llvm.struct<(i32, i32)>
%4 = llvm.add %2, %3 : i32
```
Diffstat (limited to 'llvm/lib/Target/WebAssembly/Disassembler/WebAssemblyDisassembler.cpp')
0 files changed, 0 insertions, 0 deletions