CFG, SSA & liveness — sass2mlir

Walkthrough

The analysis stack, end to end

Everything on this page runs on a lifted module — the out/saxpy.mlir from Capturing & lifting kernels, or any other .mlir the lifter produced. saxpy is straight-line code with a single branch, though, so to see the interesting output we use a small reduction kernel with a real loop: a per-block dot product. Lifted, it’s 33 instructions of sm_90a SASS in three blocks.

nvcc -arch=sm_90a -cubin dot.cu -o dot.cubin
sass2mlir harvest dot.cubin -o out/
ls out/
# dot.code   dot.mlir

The kernel under analysis

// dot.cu — one warp per block accumulates a strided partial dot product
extern "C" __global__ void dot(const float* x, const float* y, float* out, int n) {
  float acc = 0.f;
  int iters = n / blockDim.x;                    // uniform across the warp
  for (int t = 0; t < iters; ++t) {
    int i = t * blockDim.x + threadIdx.x;        // lane-dependent
    acc += x[i] * y[i];
  }
  if (threadIdx.x == 0) out[blockIdx.x] = acc;   // simplified: no cross-lane reduce
}

The loop is what we want: it produces a preheader, a body with a back-edge, and an epilogue — the smallest CFG that exercises every analysis below.

The four analyses — cfg, defuse, liveness, uniform — all take the same shape: a lifted module, a --fn glob, output on stdout. They run on the recovered SSA of the 770-op sass dialect, and none of them can touch a byte of machine code. The last section proves that.

Step 1 · cfg

Recovering the control-flow graph

CFG recovery is the first analysis — SSA, liveness, and register allocation are all built on the block structure it produces.

sass2mlir cfg out/dot.mlir --fn dot
# dot (sm_90a): 3 blocks, 4 edges, 1 back-edge
# bb0 (entry): 0x0000–0x00e0  → succ: bb1 (unconditional)
# bb1 (loop):  0x00f0–0x01b0  → succ: bb1 (taken, back-edge), bb2 (fall-through)
# bb2:         0x01c0–0x0200  → succ: (exit)

bb0 is the preheader: parameter loads, the threadIdx.x read, the accumulator zeroing. bb1 is the loop body — its conditional branch back to 0x00f0 is the back-edge, with the not-taken path falling through to the epilogue bb2, which stores and EXITs.

Every terminator falls into one of five kinds:

Terminator	Class	Successors
BRA / BRX / JMP, no guard or PT guard	unconditional	1 (target)
BRA under any other guard	conditional	2 (taken + fall-through)
RET	return	none
EXIT	thread exit	none
BSSY / BSYNC	convergence pair	BSSY pushes a reconvergence point; BSYNC branches to it

BSSY/BSYNC exist because a GPU warp diverges and reconverges: BSSY pushes a convergence point, and the matching BSYNC branches to it when the warp reconverges. dot has no divergent region, so they don’t appear here; any kernel with an if both sides of which execute will show the pair.

Branch targets are absolute, resolved per instruction from target = addr + 16 + (rel << scale), where addr is the branching instruction’s own offset and rel/scale come out of its decoded operand. The back-edge above: the BRA sits at 0x01b0 and decodes rel = -0xd0, scale = 0, so target = 0x01b0 + 0x10 + (-0xd0 << 0) = 0x00f0 — exactly bb1’s leader. Leaders are the entry instruction, every branch target, and the instruction after every terminator; leaders partition the stream, and each block’s successor edges are read off its terminator kind.

For a picture, emit Graphviz and render it:

sass2mlir cfg out/dot.mlir --fn dot --dot -o cfg.dot
dot -Tsvg cfg.dot -o cfg.svg

Recovery is cross-checked against nvdisasm -bbcfg as an independent oracle: the recovered block boundaries have to agree with NVIDIA’s own basic-block disassembly of the same cubin.

nvdisasm -bbcfg dot.cubin > nvdisasm.cfg
sass2mlir cfg out/dot.mlir --fn dot --dot -o sass2mlir.cfg
# normalized to (block-start, successors) pairs and diffed: no output, 3/3 blocks agree

Step 2 · SSA

Block arguments, def-use, reaching definitions

Cross-block dataflow is native MLIR: φ-nodes at joins become block arguments (Braun-style SSA). The loop’s φ-nodes are bb1’s arguments — the running accumulator, the index, and the two load cursors — and the back-edge supplies their next values:

^bb1(%acc: !sass.reg, %t: !sass.reg, %pa: !sass.reg, %pb: !sass.reg):
  %21 = sass.LDG_0x381.E guard(%7) [%pa]  : (!sass.reg) -> (!sass.reg)
  %22 = sass.LDG_0x381.E guard(%7) [%pb]  : (!sass.reg) -> (!sass.reg)
  // acc += x[i] * y[i] — the loop-carried value:
  %23 = sass.FFMA_0x223.FTZ guard(%7) %21, %22, %acc
          : (!sass.reg, !sass.reg, !sass.reg) -> (!sass.reg)
  %25 = sass.IADD3_0x210 guard(%7) %t, 0x1 : (!sass.reg) -> (!sass.reg)      // t += 1
  %26 = sass.ISETP_0x40c.LT.AND guard(%7) %25, %14
          : (!sass.reg, !sass.reg) -> (!sass.pred)                            // t < iters
  // taken: back-edge feeding the phis; not-taken: fall through to bb2
  sass.BRA_0x942 guard(%26) ^bb1(%23, %25, %27, %28
          : !sass.reg, !sass.reg, !sass.reg, !sass.reg) : () -> ()

%acc is defined by the preheader on the first iteration and by %23 on every later one — the block argument is exactly the φ that merges those two definitions, with no explicit φ instruction anywhere.

defuse walks this SSA and reports each op’s definitions and uses, plus the reaching definitions at each block entry:

sass2mlir defuse out/dot.mlir --fn dot
# %21 = sass.LDG_0x381.E ...        defs: {%21}   uses: {%pa (bb1 arg)}
# %23 = sass.FFMA_0x223.FTZ ...     defs: {%23}   uses: {%21, %22, %acc (bb1 arg)}
# %25 = sass.IADD3_0x210 ...        defs: {%25}   uses: {%t (bb1 arg)}
# reaching defs into bb1: {%acc: {%9 (bb0), %23 (bb1)}, %t: {%10 (bb0), %25 (bb1)}, ...}
# reaching defs into bb2: {%23 (bb1), %16 (bb0)}

The reaching sets are the loop structure stated as dataflow: bb1’s %acc has exactly two reaching definitions, one from each incoming edge. Anything a transform pass needs to ask — “is this value written anywhere in the loop?” — is a query over this output, not a re-walk of the instruction stream.

Step 3 · liveness

Per-block live-in and live-out

Liveness runs the standard backward dataflow over the recovered blocks and prints register sets — physical R/UR/P names, since this is machine code:

sass2mlir liveness out/dot.mlir --fn dot
# bb0: live-in  {R2, R3, UR4, UR6}          live-out {R0, R1, R2, R3, UR4, UR6, P0}
# bb1: live-in  {R0, R1, R2, R3, UR4, UR6}  live-out {R0, R2, R3, UR4, UR6, P1}
# bb2: live-in  {R0, R2, R3}                live-out {}

bb1’s live-in and live-out overlap in {R0, R2, R3, UR4, UR6} — those are the loop-carried registers: the accumulator, the cursors, the bound, and the trip count on the uniform file. P0 dies at bb0’s exit branch; P1 is produced and consumed inside bb1 and never leaves it.

This is the analysis register allocation stands on: the Chaitin-Briggs allocator builds its interference graph from exactly these sets, and its --pinned mode — reproduce the allocation the hardware was given, byte-identically — is the corpus-wide oracle that proves the liveness model right, because any error shows up as a wrong register in the output.

Step 4 · uniform

Which values are warp-uniform

Volta and later ship a scalar uniform datapath — per-warp UR/UP registers and ops — next to the per-lane vector path. Values that are provably identical across all 32 lanes of a warp can live there, freeing vector registers and issue slots. uniform is the analysis that decides which values qualify:

sass2mlir uniform out/dot.mlir --fn dot
# uniform:     UR4 (param base &x), UR6 (param base &y), UR8 (blockDim.x),
#              UR10 (trip count n/blockDim), %t (loop counter), UP0 (t < iters check)
# non-uniform: %16 (lane id, S2R SR_TID.X), %18 (element index i),
#              %21, %22 (LDG results), %23 (FFMA result)

The uniform set is exactly what you’d hope: the kernel parameter bases read from the constant bank, blockDim.x and the trip count derived from it, the loop counter, and the loop’s bounds check — all provably identical in every lane of the warp. Anything touched by SR_TID.X — the lane id itself, the element index i, the gathered loads, the accumulated result — is non-uniform. Note that in the lifted module the uniform values are still ordinary !sass.reg SSA values; migrating the proven-uniform web onto !sass.ureg is the pass’s job, not the analysis’s.

The rule that makes this usable is that the analysis is conservative by construction: any uncertainty marks a value non-uniform. A spurious “uniform” is not a missed optimization, it’s a wrong-code bug — a lane-varying value placed in a UR register silently broadcasts one lane’s value to the whole warp. So the analysis only says yes when it can prove it, and the scalar-to-uniform pass trusts that answer directly.

Verification

Read-only by construction

None of the four analyses mutates the module. Each one only reads the lifted IR — no op’s $raw attribute can be invalidated — so the byte-exact round trip survives any sequence of them not by discipline but by construction. Run the whole stack and then the safety net:

sass2mlir cfg      out/dot.mlir --fn dot > /dev/null
sass2mlir defuse   out/dot.mlir --fn dot > /dev/null
sass2mlir liveness out/dot.mlir --fn dot > /dev/null
sass2mlir uniform  out/dot.mlir --fn dot > /dev/null

sass2mlir roundtrip dot.cubin
# roundtrip: sm_90a dot: 33/33 instructions byte-identical (0 mismatches)

33 of 33 instructions byte-identical, after every analysis has run over the module. That’s the property the transform passes then deliberately spend: once a pass edits an op, its $raw is invalidated and nvisa_encode re-encodes it structurally — and crosscheck is the audit that the encoder gets that right.