Quickstart: five minutes to a byte-exact round trip

This walks the core loop once, end to end: get a cubin, lift its kernels to MLIR, look at the IR, and lower it back to the exact bytes. Everything here is read-only — no kernel is modified. Assumes sass2mlir is built and on your PATH and nvcc is available.

Get a cubin

Any file with embedded SASS works — an existing app’s cubin extracted with cuobjdump -xelf all app, or one you compile on the spot. The smallest possible one:

// saxpy.cu
extern "C" __global__ void saxpy(float a, const float* x, float* y, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) y[i] = a * x[i] + y[i];
}

nvcc -arch=sm_90a -cubin saxpy.cu -o saxpy.cubin

Harvest the kernels

harvest parses the cubin’s ELF structure and writes the raw machine code plus the lifted MLIR for every kernel it contains:

sass2mlir harvest saxpy.cubin -o out/
ls out/
# saxpy.code   saxpy.mlir

.code is the raw 128-bit-per-instruction SASS stream; .mlir is the same kernel lifted into the sass dialect.

Read the lifted IR

cat out/saxpy.mlir

module {
  sass.func "saxpy" {
    %0 = builtin.unrealized_conversion_cast to !sass.reg   // live-in: kernel params
    ...
    // the a * x[i] + y[i] fused multiply-adds, predicated on the i < n guard:
    %41 = sass.FFMA_0x223.FTZ guard(%38) %27, %14, %40
            : (!sass.reg, !sass.reg, !sass.reg) -> (!sass.reg)
    ...
    sass.BRA_0x942 guard(%38) ^bb1 : () -> ()              // branch over the store
    sass.STG_0x382.E guard(%38) [%36], %41 : (!sass.reg, !sass.reg) -> ()
    sass.exit
  }
}

One op per instruction, named sass.<MNEMONIC>_<opcode>; SSA values for registers; the guard(%p) predicate inline. Every op also carries its original 128-bit encoding as a $raw attribute — that’s what makes the trip back exact.

Prove the round trip

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

This did lift → print → reparse → lower and compared bytes against the original cubin. lower(lift(bytes)) == bytes, per instruction — the property the whole framework is regression-gated on.

Look at the structure

sass2mlir cfg out/saxpy.mlir --fn saxpy
# bb0 (entry): 0x0000–0x0150  → succ: bb1 (taken), bb2 (fall-through)
# bb1:         0x0160–0x0190  → succ: bb3
# ...

sass2mlir liveness out/saxpy.mlir --fn saxpy
# bb0: live-in {R2, R3, UR4}  live-out {R4, R5, P0} ...

Control flow is recovered into MLIR blocks, φ-nodes become block arguments, and the standard analyses — def-use, liveness, uniformity — run directly on the recovered SSA.