This was my first time probing an instruction encoding instead of only reading an existing encoder. I did not discover the five-instruction sequence: Faith Ekstrand had already put CUDA's output in the work item. My part started where the issue stopped: making NAK represent PLOP3 sign-bit sources, emit the observed encoding, and survive real compiler paths.

The names used below:

PTX    NVIDIA's portable assembly language
SASS   the machine instructions executed by an NVIDIA GPU
NIR    Mesa's shared compiler intermediate representation
NAK    Mesa's NVIDIA shader backend: NIR → SASS
GPR    a 32-bit general-purpose register, one value per GPU lane
Pred   a one-bit predicate register used for conditions

The part that was already known

The issue contained this SASS:

IADD3 R0, R2, R5, RZ
PLOP3.LUT P0, PT, R5.SIGN, R2.SIGN, R0.SIGN, 0x2,  0x0
PLOP3.LUT P1, PT, R5.SIGN, R2.SIGN, R0.SIGN, 0x40, 0x0
SEL R9, R0, 0x7fffffff, !P0
SEL R9, R9, 0x80000000, !P1

That answered “what does NVIDIA's compiler select?” It did not answer how NAK should type those sources, which SM70 encoding form to use, or what legalization and uniformization may do to them.

Where the operation disappeared

NIR already had nir_op_iadd_sat. NAK requested generic lowering unconditionally:

lower_uadd_sat: dev.sm < 70,
lower_usub_sat: dev.sm < 70,
lower_iadd_sat: true, // TODO

So this was not a missing-correctness bug. NIR expanded the operation before NAK saw it. The native path required changing that option for SM70+, adding the NIR translation, building the sequence, and carrying the new source form through NAK IR to the encoder.

Why the LUTs are 0x02 and 0x40

There is no single signed IADD.SAT instruction here. Start with a wrapping sum. Signed addition overflows only when both inputs have the same sign and the wrapped result has the opposite sign:

x.sign  y.sign  sum.sign  result
   0       0        1     positive overflow → INT_MAX
   1       1        0     negative overflow → INT_MIN

PLOP3 is a three-input boolean lookup table. Its LUT byte has one bit for each of the eight input rows. With NAK's (x, y, sum) ordering, 001 is bit 1, hence 1 << 1 = 0x02. The row 110 is bit 6, hence 1 << 6 = 0x40.

sum          = IADD3(x, y, 0)
pos_overflow = PLOP3(x.sign, y.sign, sum.sign, LUT=0x02)
neg_overflow = PLOP3(x.sign, y.sign, sum.sign, LUT=0x40)
tmp          = SEL(pos_overflow, INT_MAX, sum)
dst          = SEL(neg_overflow, INT_MIN, tmp)

Asking CUDA for the machine code

The sequence from the issue can be produced from a small SM89 PTX kernel containing one add.sat.s32:

.version 8.8
.target sm_89
.address_size 64

.visible .entry test_iadd_sat(
    .param .u64 out_ptr,
    .param .u64 in_ptr)
{
    .reg .b64 %rd<4>;
    .reg .s32 %r<4>;

    ld.param.u64 %rd1, [out_ptr];
    ld.param.u64 %rd2, [in_ptr];
    add.u64 %rd3, %rd2, 4;
    ld.global.s32 %r1, [%rd2];
    ld.global.s32 %r2, [%rd3];
    add.sat.s32 %r3, %r1, %r2;
    st.global.u32 [%rd1], %r3;
    ret;
}

The corresponding ptxas and nvdisasm commands are:

ptxas -arch=sm_89 iadd_sat.ptx -o iadd_sat.cubin
nvdisasm --print-code --print-instruction-encoding iadd_sat.cubin |
  rg 'IADD3|PLOP3|SEL'

The relevant output was the same five-instruction core shown in the issue:

IADD3 R2, R0, R3, RZ
PLOP3.LUT P0, PT, R0.SIGN, R3.SIGN, R2.SIGN, 0x2,  0x0
PLOP3.LUT P1, PT, R0.SIGN, R3.SIGN, R2.SIGN, 0x40, 0x0
SEL R7, R2, 0x7fffffff, !P0
SEL R7, R7, 0x80000000, !P1

ptxas gave me NVIDIA's instruction selection. nvdisasm printed both the readable SASS and each 128-bit instruction encoding. I did not try to reverse every bit of PLOP3. NAK already knew the ordinary predicate form, common ALU register fields, LUT field, and predicate destinations. I compared that known encoder with the all-.SIGN instruction emitted by CUDA and added the missing GPR-source form.

The first encoder version reused NAK's common three-source ALU encoding and filled only the fields established by the comparison:

127                                                   0
┌──────────────┬──────┬──────┬──────────┬─────────────┐
│ common ALU,  │ dst1 │ dst0 │ LUT      │ opcode/form │
│ source and   │84–86 │81–83 │72–79     │0–11         │
│ control bits │      │      │          │             │
└──────────────┴──────┴──────┴──────────┴─────────────┘
e.encode_alu(0x01f, None,
             Some(&self.srcs[0]),
             Some(&self.srcs[1]),
             Some(&self.srcs[2]));
e.set_field(72..80, self.ops[0].lut);
e.set_pred_dst(81..84, &self.dsts[0]);
e.set_pred_dst(84..87, &self.dsts[1]);

The diagram is an overlay of the fields this code writes, not a claim that I recovered the full 128-bit format.

I then made NAK's disassembly test encode the IR and ask NVIDIA's decoder for the expected text:

export NAK_TEST=/path/to/mesa-build/src/nouveau/compiler/nak

PATH=/path/to/cuda/bin:$PATH \
  "$NAK_TEST" \
  --test --exact nvdisasm_tests::test_plop3 --test-threads 1

This proves that NVIDIA's decoder agrees with the fields NAK emitted. It does not prove that the instruction computes the right value. For that I needed the GPU.

My first model was deliberately narrow

Before reading Máté Pinczel's commits or the later discussion, I modeled .SIGN as one instruction-wide mode: either all three PLOP3 sources were predicates or all three were GPR sign bits. That was enough for iadd_sat, where both LUTs use x.sign, y.sign, and sum.sign.

The patch connected the whole vertical slice:

NIR iadd_sat
  → NAK builder: IADD3 + two PLOP3 + two SEL
  → OpPLop3 with sign-mode GPR sources
  → legalization
  → SM70+ encoding
  → AD107

The local hardware test stores the GPU result and compares it with Rust's saturating_add for boundary inputs and 100 deterministic random pairs:

PATH=/path/to/cuda/bin:$PATH \
  "$NAK_TEST" \
  --test --exact hw_tests::test_iadd_sat --test-threads 1

The useful failure

The focused hardware and disassembly tests passed. I then ran the Vulkan Conformance Test Suite (CTS) dot-accumulate-saturate cases. These are not named iadd_sat tests, but NIR lowers a signed saturating dot product to a normal dot product followed by iadd_sat with the accumulator. That made them a wider test of the new path.

844 pass
152 fail
failure: NAK assertion during pipeline compilation

A normal GPR holds a separate value for each GPU lane. A uniform register holds one value shared by every lane. Some CTS variants used values that NAK could move into uniform registers, and NAK considered every PLOP3 uniform-capable. My new .SIGN encoding supported only normal GPR sources. The compiler therefore reached an unsupported uniform form and stopped at the encoder assertion; the GPU never ran a wrong result.

I kept the assertion. The fix belonged earlier: sign-source PLOP3 is non-uniform, non-register inputs must be copied into GPRs, and predicate optimizations must not treat its ALU sources as ordinary predicates. After that, the full filter covered signed, unsigned, and mixed-signedness dot-accumulate-saturate cases:

// Ordinary predicate PLOP3 may be uniform; .SIGN PLOP3 may not.
Op::PLop3(op) => !op.src_sign,

// The encoder's .SIGN form consumes normal GPRs.
if self.src_sign {
    for src in &mut self.srcs {
        b.copy_alu_src_if_not_reg(src, RegFile::GPR, SrcType::GPR);
    }
}
uniform value
    → copy to a normal GPR in each lane
    → warp PLOP3 R?.SIGN, R?.SIGN, R?.SIGN
export MESA_PREFIX=/path/to/patched-mesa-install
export VK_CTS=/path/to/VK-GL-CTS/build/external/vulkancts/modules/vulkan/deqp-vk

env VK_DRIVER_FILES="$MESA_PREFIX/share/vulkan/icd.d/nouveau_icd.x86_64.json" \
  LD_LIBRARY_PATH="$MESA_PREFIX/lib" \
  MESA_SHADER_CACHE_DISABLE=true \
  "$VK_CTS" -n \
  'dEQP-VK.spirv_assembly.instruction.compute.op*dotaccsatkhr.*' \
  --deqp-log-filename=iadd_sat.qpa \
  --deqp-log-images=disable
Passed: 996/996
Failed: 0

Then I found the broader model

In the issue I wrote explicitly that my all-.SIGN implementation predated the comment and Máté's commits. Faith's proposed direction was to make .SIGN a source modifier, assuming PLOP3 supports arbitrary mixtures of rN.sign and predicates. Máté's WIP already modeled that per source and included mixed-source nvdisasm coverage. I built his commit locally; that test passed with CUDA 13.3.73.

That changed the IR question. With P for a predicate source and S for a register sign bit, three sources have eight logical layouts:

PPP  PPS  PSP  PSS  SPP  SPS  SSP  SSS

I did not assume there must be eight physical encodings. I took a known SM89 PLOP3 binary, varied all 4,096 values of its 12-bit opcode field, and ran every candidate through CUDA 13.3 nvdisasm. The decoder exposed these relevant forms:

0x81c        PPP: predicate-only PLOP3
0x21d        PSP: one .SIGN source
0x21e        PSS: two .SIGN sources
0x21f        SSS: three .SIGN sources
0xa1d–0xa1f  constant-buffer source variants
0x89c        uniform UPLOP3

This was automated, not a visual search through 4,096 binary strings. The scanner kept the other 116 bits of a known instruction fixed, replaced bits 0–11 with each candidate, wrote the four 32-bit words to a temporary file, and retained outputs containing PLOP3:

let base = [0x0200721f_u32, 0x00000003,
            0x00700204,     0x000fc000];

for opcode in 0_u32..0x1000 {
    let code = [
        (base[0] & !0xfff) | opcode,
        base[1], base[2], base[3],
    ];
    // Write code as little-endian bytes, then:
    nvdisasm -b SM89 --print-raw candidate.bin
}

For example, the controlled template changes like this:

logical form   four 32-bit words                         low 12 bits
PPP            0200781c 00000003 00700204 000fc000       0x81c
SSS            0200721f 00000003 00700204 000fc000       0x21f
                       ^ the scanned field is here

The word comparison tells me which controlled field changed. nvdisasm tells me how NVIDIA's decoder interprets that candidate. Neither alone proves execution semantics, which is why I followed the scan with the AD107 truth-table test.

There was no separate decoded opcode for PPS, SPP, SPS, or SSP. The useful model is four canonical physical layouts. Legalization can reorder a logical layout into one of them, but it must permute the LUT inputs at the same time. Swapping sources without swapping the LUT axes silently changes the boolean function.

For a concrete reorder, suppose the logical operation is SPP(a, b, c, lut): a is a GPR sign bit, while b and c are predicates. The hardware form puts the one sign source in the middle slot:

logical:   SPP(a, b, c, lut)
physical: PSP(b, a, c, permute_xy(lut))

Only swapping a and b would make LUT input x receive the old y. Applying the same permutation to the LUT preserves all eight rows of the original boolean function.

Checking all eight logical layouts

I added an AD107 truth-table test. It uses 0 and INT_MIN, so one boolean can be supplied either as value != 0 through a predicate or as value.sign through a GPR. For every P/S layout, it runs all eight input rows with two asymmetric LUTs, 0x1b and 0xe4. These are test functions, not the 0x02 and 0x40 overflow LUTs above; their asymmetry makes a bad source order or LUT permutation visible.

8 source layouts × 8 truth-table rows × 2 LUTs = 128 result checks

All eight logical layouts pass after source canonicalization and the matching LUT permutation. This is the part that convinced me the per-source modifier model is better than my original instruction-wide flag.

I also tried to make ptxas select mixed-source PLOP3 directly. Schematically, I asked for a boolean function combining one sign test with predicates that already represented boolean conditions:

sx = (x < 0)
out = LUT(sx, p, q)

ptxas folded straightforward versions into ISETP instead of emitting a mixed PLOP3. I then used PTX vote operations to stop the predicates from being trivially folded. That produced a PLOP3, but ptxas first materialized sx as a predicate, so all three final sources were predicates:

sign test + predicate logic  → ISETP
opaque predicates via vote  → predicate materialization → PPP PLOP3

These are schematic versions of the probes; I did not preserve the exact mixed-probe PTX files. The result only describes CUDA's instruction selection for those inputs. It does not disprove mixed encodings; the decoder scan and hardware test answer that question more directly.

Results and limits

The narrow native iadd_sat path passes its AD107 hardware test, external disassembly test, and 996 CTS cases. A corpus of 110 Fossilize databases and 1,334 comparable shader executables produced no output changes because none contained a surviving nir_op_iadd_sat. That is corpus coverage, not a performance result.

The broader mixed-source branch is still research code. Its focused truth-table test passes, but the full NAK run is not green: 60 tests pass and two fail, including modifier legalization in Máté's original iadd_sat test and an existing plop2 path. The opcode scanner and debug printing also do not belong in an MR.

I have not opened an MR. I posted the evidence on the work item and asked whether the Vulkan accelerated properties belong in the same series, how maintainers want the commits split, and whether the per-source model matches the intended direction. The code remains local while I wait for that direction.

References