spv/lift: canonicalize loop shortcut before phi materialization

[niklasha]

While playing around with rust-gpu I stepped upon a miscompilation, which I debuggged to be a problem in spirt where loops could clobber values in non-deterministic ways. After quite some bughunt, spirt's "lift" logic seems to be at error. The included commit-set adds regression tests matching the problem I faced, and a patch for spirt, that fixes them (and my original miscompilation problem).

Disclosure: This change was developed with LLM assistance (gpt-5.3-codex); I reviewed, tested, and take responsibility for the final patch.

spv/lift: canonicalize loop shortcut to preserve loop edge value mapping

Summary

This MR fixes a spv::lift loop-shape issue where a specific structured loop encoding can preserve a fragile shortcut CFG pattern through lifting.

The fix adds a guarded canonicalization in FuncLifting::from_func_decl (after linear branch fusion, before dead-block pruning and phi case collection) and adds fixture-backed structural + semantic regression tests.

Problem

In a narrow loop pattern:

• loop header branches to a body SelectBranch(BoolCond)
• one body arm is an empty pass-through to the body merge
• body merge branches to loop_continue
• loop_continue conditionally branches to {loop_header, loop_merge} on the same condition (or negation)

this preserved shape can mis-handle edge-carried values in later phi materialization-sensitive paths.

Changes

1) spv::lift canonicalization

File: src/spv/lift.rs

Added helper methods in FuncLifting:

• recompute_use_counts
• is_passthrough_branch_to
• is_const_opcode
• continue_cond_relation
• rewrite_continue_as_unconditional_backedge
• canonicalize_loop_continue_shortcuts

Canonicalization behavior (only when strict guards match):

• retarget the pass-through body arm directly to loop_merge
• remove the body Selection merge marker for that case
• rewrite loop_continue into unconditional backedge to loop header
• keep only header-target phi payloads on rewritten continue edge
• recompute predecessor/use counts before dead-block pruning

2) Regression coverage

Files:

• tests/regression_loop_continue_shortcut.rs
• tests/data/loop-continue-shortcut*.spvbin
• tests/data/basic.frag.glsl.dbg.spvbin

Test strategy:

• Run lower -> structurize -> link -> lift on fixed SPIR-V fixtures
• Parse lifted blocks and check both:
  • structural shortcut shape detector
  • semantic loop-carried OpUndef detector on the shortcut path
• Assert:
  • repro fixture set does not retain the shortcut shape
  • repro fixture set does not contain loop-carried OpUndef through shortcut phis
  • control fixture set does not produce false positives

Named tests:

• no_loop_continue_shortcut_shape_after_lift
• no_loop_carried_undef_from_shortcut_after_lift
• detector_does_not_trigger_on_non_loop_fixture

Fixture sets include:

• pre-test while variants (const, len)
• nested pre-test while variants (const, len)
• controls (post-test loop, no-loop, basic fragment)

Safety/Correctness guards

The canonicalization is conservative and only fires when all of the following hold:

• exact terminator kinds/attrs for header, body, body-merge, and continue blocks
• expected merge topology (Loop + body Selection)
• continue target ordering and uniqueness constraints
• predecessor count constraints for body-merge and continue
• phi payload arity compatibility for both body-merge and loop-header edges
• loop-merge has no phis in rewritten shape

No rewrite occurs when checks fail.

Validation

Executed locally:

• cargo test --test regression_loop_continue_shortcut
• cargo test -q
• cargo test --release --all-targets
• cargo build -q --examples

All pass.

Scope

Semantics-affecting code changes are isolated to:

• src/spv/lift.rs

Test-only additions are isolated to:

• tests/regression_loop_continue_shortcut.rs
• tests/data/*.spvbin fixture files

Tracking

• Issue: spv::lift: loop shortcut preserved with undef-fed loop-carried phi (same body/continue condition) #31

[eddyb]

Problem

In a narrow loop pattern:

• loop header branches to a body SelectBranch(BoolCond)
• one body arm is an empty pass-through to the body merge
• body merge branches to loop_continue
• loop_continue conditionally branches to {loop_header, loop_merge} on the same condition (or negation)

It would help a lot to have a GLSL/WGSL example, I believe you mean that:

while(f()) {
    g();
}

ends up encoded like this in SPIR-T (for reasons that go back to RVSDG):

do {
    bool cond = f();
    if(cond) {
        g();
    }
} while(cond);

And if you were to use SPIRV-Cross or Naga to get GLSL/WGSL after SPIR-V -> SPIR-T -> SPIR-V, you'd see the latter do-while with the bool value being passed through.

---

Whether or not we improve this aspect, could you first open an issue about it?
(You should also try to see if spirv-opt cleans it up, and what passes do that in it)

Any amount of cleverness comes at a cost, and SPIR-T may be suboptimal in a few places, to maximize the likelihood of correctness (as we don't have the tools for formally verified translation etc.).

[niklasha]

Thanks, this is a helpful framing.

Yes, your understanding is exactly the family I was targeting, with one extra detail: the problematic case is when a value carried around the loop gets a merge input from the pass-through arm that is OpUndef in SPIR-V.

A clear source-level shape (GLSL/WGSL) is:

#version 450
layout(local_size_x = 1) in;
layout(set = 0, binding = 0, std430) buffer OutBuf { uint out_data[]; } out_buf;

void main() {
    uint i = 0u;
    while (i < 1u) {
        if (i < 1u) {
            i = i + 1u;
        }
    }
    out_buf.out_data[0] = i; // expected: 1
}

@group(0) @binding(0)
var<storage, read_write> out_data: array<u32>;

@compute @workgroup_size(1)
fn main() {
  var i: u32 = 0u;
  while (i < 1u) {
    if (i < 1u) {
      i = i + 1u;
    }
  }
  out_data[0] = i; // expected: 1u
}

The problematic lowered form (in pseudo-GLSL) is:

uint i = 0u;
do {
    uint next;              // corresponds to OpUndef on one incoming path
    if (i < 1u) {
        next = i + 1u;
    }
    i = next;               // value merge can carry undef from pass-through arm
} while (i < 1u);

In the concrete failing fixture, the key pattern is:

%16 = OpUndef %uint
...
%25 = OpPhi %uint %32 %53 %16 %54
...
OpBranchConditional %31 %53 %54      ; body cond
...
OpBranchConditional %31 %51 %57      ; continue cond (same condition)

So the issue is not just while -> do/while structurally, but that this specific shortcut shape can preserve an undef-fed loop-carried value in a way that later phi materialization paths mishandle.

Also per your suggestion: I tested spirv-opt -O and spirv-opt -Os on the concrete failing .link.spv; in this case both retained the shortcut shape and the undef-fed phi (i.e. they did not clean it up).

I can open a dedicated issue with:

• the failing fixture,
• a minimal script to reproduce pre-fix vs post-fix behavior,
• and the spirv-opt observations.