llvm-project

5d300afa - [MLIR][NVVM] Add support for multiple return values in `inline_ptx` (#153774)

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