Skip to content

16. Testing the GPU Path and the Loop Transforms

Mountain goats on the mountain at night

What you will understand: how to bring Chapter 7's loop transforms and Chapters 10 through 12's GPU pipeline under the assertion-based suite Chapter 15 built, with no GPU required. Every test is shown here exactly as it exists in the repository, next to the real output it matches, so you can see what each CHECK line is checking. Along the way the chapter explains the GPU concepts those checks depend on: what a launch grid is, why one matrix becomes seven kernel parameters, and how to read the PTX the compiler produces. It also checks, again, that the new tests can fail, and records two mistakes made while writing them.

What you need to know first: Chapter 15's suite and its lit/FileCheck basics (this chapter re-explains only the directives it uses); Chapter 7's tiling and unrolling; Chapters 10 through 12's GPU recipe.

Compile and run

Every command, listing and output on this page is reproduced by the commands in Reproducing this chapter at the bottom of the page, which also lists what must be built first. The chain of builds is in Getting Started.

Why this page shows FAIL lines, and why that is good

Parts of this page come from negative controls: the same tests run against an older build, or against deliberately broken code (a "mutation"). In those runs a FAIL is the expected, wanted result: it means a test noticed the problem, which is how we know the tests are worth anything. What would be wrong is the opposite, a deliberately broken build that passes everything. The baseline (the unmodified current build) must always show every test passing, and it does. Each script that does this says so in its own header and prints a reminder at the top of its output.

Where the suite stands

Chapter 15 left 22 tests and a list of what was not covered. This chapter adds 12, for 34:

Group New tests What they pin down
loops/ tiling, unroll-full Chapter 7's tiling and full unrolling
gpu/ kernel kernel-outlining, nvvm-lowering, ptx-default Chapter 10: mg.add to a gpu.module kernel, to the nvvm dialect, to real PTX
gpu/ options ptx-index32, ptx-bare-pointers index-bitwidth=32; kernel-bare-ptr-calling-convention=1
gpu/ host host-runtime-calls, host-runtime-calls-bare Chapter 11/12: the mgpu* call sequence in the translated LLVM IR
gpu/ toolchain cuda-comparison, pass-order-crash the CUDA-style kernel compiles; a known LLVM 18.1.3 crash is pinned
execution/ gpu-host-with-stub-runtime the host glue run against a CPU stub runtime

Real run (code/run_out_34.txt):

-- Testing: 34 tests, 4 workers --
PASS: IR Forge / Mountain Goat :: execution/mismatch-columns-aborts.mlir (1 of 34)
PASS: IR Forge / Mountain Goat :: execution/mismatch-aborts.mlir (2 of 34)
PASS: IR Forge / Mountain Goat :: execution/static-add-runs.mlir (3 of 34)
PASS: IR Forge / Mountain Goat :: canonicalize/transpose-twice-dynamic-exact.mlir (4 of 34)
PASS: IR Forge / Mountain Goat :: execution/dynamic-shapes-run.mlir (5 of 34)
PASS: IR Forge / Mountain Goat :: canonicalize/transpose-twice.mlir (6 of 34)
PASS: IR Forge / Mountain Goat :: canonicalize/transpose-twice-dynamic-mixed.mlir (7 of 34)
PASS: IR Forge / Mountain Goat :: canonicalize/fold-constant-add.mlir (8 of 34)
PASS: IR Forge / Mountain Goat :: verifier/transpose-rank3.mlir (9 of 34)
PASS: IR Forge / Mountain Goat :: loops/fusion.mlir (10 of 34)
PASS: IR Forge / Mountain Goat :: execution/mixed-static-dynamic.mlir (11 of 34)
PASS: IR Forge / Mountain Goat :: lowering/dynamic-add-affine.mlir (12 of 34)
PASS: IR Forge / Mountain Goat :: verifier/unranked-rejected.mlir (13 of 34)
PASS: IR Forge / Mountain Goat :: execution/strided-view.mlir (14 of 34)
PASS: IR Forge / Mountain Goat :: verifier/add-static-mismatch.mlir (15 of 34)
PASS: IR Forge / Mountain Goat :: verifier/dynamic-accepted.mlir (16 of 34)
PASS: IR Forge / Mountain Goat :: lowering/mixed-add-one-check.mlir (17 of 34)
PASS: IR Forge / Mountain Goat :: lowering/bufferize-dynamic-add.mlir (18 of 34)
PASS: IR Forge / Mountain Goat :: lowering/static-add-affine.mlir (19 of 34)
PASS: IR Forge / Mountain Goat :: verifier/transpose-bad-shape.mlir (20 of 34)
PASS: IR Forge / Mountain Goat :: lowering/dynamic-transpose-no-check.mlir (21 of 34)
PASS: IR Forge / Mountain Goat :: lowering/bufferize-static-add.mlir (22 of 34)
PASS: IR Forge / Mountain Goat :: gpu/cuda-comparison.mlir (23 of 34)
PASS: IR Forge / Mountain Goat :: gpu/kernel-outlining.mlir (24 of 34)
PASS: IR Forge / Mountain Goat :: gpu/host-runtime-calls-bare.mlir (25 of 34)
PASS: IR Forge / Mountain Goat :: gpu/host-runtime-calls.mlir (26 of 34)
PASS: IR Forge / Mountain Goat :: gpu/nvvm-lowering.mlir (27 of 34)
PASS: IR Forge / Mountain Goat :: gpu/ptx-bare-pointers.mlir (28 of 34)
PASS: IR Forge / Mountain Goat :: gpu/pass-order-crash.mlir (29 of 34)
PASS: IR Forge / Mountain Goat :: gpu/ptx-default.mlir (30 of 34)
PASS: IR Forge / Mountain Goat :: loops/unroll-full.mlir (31 of 34)
PASS: IR Forge / Mountain Goat :: loops/tiling.mlir (32 of 34)
PASS: IR Forge / Mountain Goat :: gpu/ptx-index32.mlir (33 of 34)
PASS: IR Forge / Mountain Goat :: execution/gpu-host-with-stub-runtime.mlir (34 of 34)

Testing Time: 0.71s

Total Discovered Tests: 34
  Passed: 34 (100.00%)

A short primer: FileCheck directives used here

Each test is an MLIR file whose // RUN: lines say what to execute (with %-substitutions defined in lit.cfg.py) and whose // CHECK: lines are patterns that the output piped into FileCheck must satisfy, in order:

Directive Meaning
CHECK: text text must appear somewhere after the previous match
CHECK-NEXT: text text must appear on the very next line
CHECK-SAME: text text must appear on the same line as the previous match
CHECK-NOT: text text must not appear between the surrounding matches
CHECK-COUNT-n: text text must appear n times in a row
CHECK-LABEL: text an anchor that resets the search, so one function's checks can't match another's
{{regex}} a regular expression inside a pattern
%[[NAME:regex]] / %[[NAME]] capture something once, require the same text later

Registers and SSA names (%0, %rd7) differ between compiler versions, so the tests below match them with {{.*}} or captures and pin only what matters.

Part 1: the loop transforms (Chapter 7)

The input

Both loop tests start from Chapter 7's fused program: one 2x2 loop nest that adds two matrices element by element (with the transposed access Chapter 7 found) through a 1x1 temporary buffer.

module {
  func.func private @printMemrefF64(memref<*xf64>) attributes {llvm.emit_c_interface}
  func.func @compute(%arg0: memref<2x2xf64>, %arg1: memref<2x2xf64>) {
    %alloc = memref.alloc() : memref<1x1xf64>
    %alloc_0 = memref.alloc() : memref<2x2xf64>
    affine.for %arg2 = 0 to 2 {
      affine.for %arg3 = 0 to 2 {
        %0 = affine.load %arg0[%arg3, %arg2] : memref<2x2xf64>
        %1 = affine.load %arg1[%arg3, %arg2] : memref<2x2xf64>
        %2 = arith.addf %0, %1 : f64
        affine.store %2, %alloc[0, 0] : memref<1x1xf64>
        %3 = affine.load %alloc[0, 0] : memref<1x1xf64>
        affine.store %3, %alloc_0[%arg2, %arg3] : memref<2x2xf64>
      }
    }
    %cast = memref.cast %alloc_0 : memref<2x2xf64> to memref<*xf64>
    call @printMemrefF64(%cast) : (memref<*xf64>) -> ()
    return
  }
}

Tiling

The test:

// Chapter 7: --affine-loop-tile="tile-size=1" turns the fused 2-deep nest into 4 loops (tile loops, then point loops).
// RUN: %mg-opt %s --affine-loop-tile="tile-size=1" | %FileCheck %s
// CHECK: affine.for %[[I:.*]] = 0 to 2 {
// CHECK-NEXT: affine.for %[[J:.*]] = 0 to 2 {
// CHECK-NEXT: affine.for %[[PI:.*]] = #map(%[[I]]) to #map1(%[[I]]) {
// CHECK-NEXT: affine.for %[[PJ:.*]] = #map(%[[J]]) to #map1(%[[J]]) {
// CHECK: affine.load %arg0[%[[PJ]], %[[PI]]]
// CHECK: affine.store %{{.*}}, %{{.*}}[%[[PI]], %[[PJ]]]

The real output of mg-opt fused.mlir --affine-loop-tile="tile-size=1" (show/tile1.mlir):

    affine.for %arg2 = 0 to 2 {
      affine.for %arg3 = 0 to 2 {
        affine.for %arg4 = #map(%arg2) to #map1(%arg2) {
          affine.for %arg5 = #map(%arg3) to #map1(%arg3) {
            %0 = affine.load %arg0[%arg5, %arg4] : memref<2x2xf64>
            %1 = affine.load %arg1[%arg5, %arg4] : memref<2x2xf64>
            %2 = arith.addf %0, %1 : f64
            affine.store %2, %alloc[0, 0] : memref<1x1xf64>
            %3 = affine.load %alloc[0, 0] : memref<1x1xf64>
            affine.store %3, %alloc_0[%arg4, %arg5] : memref<2x2xf64>
          }
        }
      }
    }

Reading it: tiling splits each of the two original loops into an outer tile loop (%arg2, %arg3, stepping through tiles) and an inner point loop (%arg4, %arg5) that walks the elements inside one tile. #map(%arg2) is d0 and #map1(%arg2) is d0 + 1, so with a tile size of 1 each point loop runs over exactly [i, i+1): a one-iteration loop. That is the point of using size 1 on a 2x2 example: the structure is visible without changing what is computed. The CHECK-NEXT chain asserts all four loops nest directly, and the captures (%[[PJ:.*]]) assert that the load still reads [%arg5, %arg4] (point-j, point-i) while the store writes [%arg4, %arg5], so the tiled code keeps Chapter 7's index swap.

Full unrolling

// Chapter 7: fully unrolling both loops, then canonicalizing, leaves NO loops, four add-and-store groups
// with constant indices, and no leftover affine.apply index arithmetic.
// RUN: %mg-opt %s --affine-loop-unroll="unroll-full" --affine-loop-unroll="unroll-full" --canonicalize | %FileCheck %s
// CHECK-NOT: affine.for
// CHECK-NOT: affine.apply
// CHECK: affine.load %arg0[0, 0]
// CHECK: affine.store %{{.*}}, %{{.*}}[0, 0] : memref<2x2xf64>
// CHECK: affine.load %arg0[1, 0]
// CHECK: affine.store %{{.*}}, %{{.*}}[0, 1] : memref<2x2xf64>
// CHECK: affine.load %arg0[0, 1]
// CHECK: affine.store %{{.*}}, %{{.*}}[1, 0] : memref<2x2xf64>
// CHECK: affine.load %arg0[1, 1]
// CHECK: affine.store %{{.*}}, %{{.*}}[1, 1] : memref<2x2xf64>
// CHECK-NOT: affine.for

Real output after --affine-loop-unroll="unroll-full" run twice (each run unrolls only the innermost loops, as Chapter 7 found) and --canonicalize (show/unroll_full.mlir):

module {
  func.func private @printMemrefF64(memref<*xf64>) attributes {llvm.emit_c_interface}
  func.func @compute(%arg0: memref<2x2xf64>, %arg1: memref<2x2xf64>) {
    %alloc = memref.alloc() : memref<1x1xf64>
    %alloc_0 = memref.alloc() : memref<2x2xf64>
    %0 = affine.load %arg0[0, 0] : memref<2x2xf64>
    %1 = affine.load %arg1[0, 0] : memref<2x2xf64>
    %2 = arith.addf %0, %1 : f64
    affine.store %2, %alloc[0, 0] : memref<1x1xf64>
    %3 = affine.load %alloc[0, 0] : memref<1x1xf64>
    affine.store %3, %alloc_0[0, 0] : memref<2x2xf64>
    %4 = affine.load %arg0[1, 0] : memref<2x2xf64>
    %5 = affine.load %arg1[1, 0] : memref<2x2xf64>
    %6 = arith.addf %4, %5 : f64
    affine.store %6, %alloc[0, 0] : memref<1x1xf64>
    %7 = affine.load %alloc[0, 0] : memref<1x1xf64>
    affine.store %7, %alloc_0[0, 1] : memref<2x2xf64>
    %8 = affine.load %arg0[0, 1] : memref<2x2xf64>
    %9 = affine.load %arg1[0, 1] : memref<2x2xf64>
    %10 = arith.addf %8, %9 : f64
    affine.store %10, %alloc[0, 0] : memref<1x1xf64>
    %11 = affine.load %alloc[0, 0] : memref<1x1xf64>
    affine.store %11, %alloc_0[1, 0] : memref<2x2xf64>
    %12 = affine.load %arg0[1, 1] : memref<2x2xf64>
    %13 = affine.load %arg1[1, 1] : memref<2x2xf64>
    %14 = arith.addf %12, %13 : f64
    affine.store %14, %alloc[0, 0] : memref<1x1xf64>
    %15 = affine.load %alloc[0, 0] : memref<1x1xf64>
    affine.store %15, %alloc_0[1, 1] : memref<2x2xf64>
    %cast = memref.cast %alloc_0 : memref<2x2xf64> to memref<*xf64>
    call @printMemrefF64(%cast) : (memref<*xf64>) -> ()
    return
  }
}

The first CHECK-NOT: affine.for asserts no loop survives; CHECK-NOT: affine.apply asserts canonicalization removed the index arithmetic that unrolling introduced; then the four groups are checked in the order the output has them, with the exact constant indices. The final CHECK-NOT: affine.for after the last group guards against a loop reappearing at the end.

Part 2: the GPU kernel (Chapter 10)

Background: what the kernel is

A GPU launch runs the same function once per thread. Threads are grouped into blocks, and blocks into a grid. Here the grid is 2 blocks of 2 threads, and each thread computes one element of a 2x2 result: the block index picks the row and the thread index picks the column. The recipe that builds this from Chapter 4's loops is: prove the loops independent (affine-parallelize), turn them into scf.parallel (lower-affine), map the outer loop to blocks and the inner to threads (gpu-map-parallel-loops), build a launch (convert-parallel-loops-to-gpu), and move the body into its own module (gpu-kernel-outlining).

kernel-outlining

// Chapter 10: mg.add -> affine -> a real gpu.module kernel, 2 blocks x 2 threads, one element per thread.
// RUN: %mg-opt %inputs/static_add.mlir --convert-mg-to-affine | %mlir-opt %to-gpu | %FileCheck %s
// CHECK: gpu.launch_func @add_tensors_kernel::@add_tensors_kernel blocks in (%{{.*}}, %{{.*}}, %{{.*}}) threads in (%{{.*}}, %{{.*}}, %{{.*}})
// CHECK: gpu.module @add_tensors_kernel
// CHECK: gpu.func @add_tensors_kernel(%{{.*}}: index, %{{.*}}: index, %{{.*}}: memref<2x2xf64>, %{{.*}}: memref<2x2xf64>, %{{.*}}: memref<2x2xf64>) kernel
// CHECK: gpu.block_id x
// CHECK: gpu.thread_id x
// CHECK: memref.load
// CHECK: memref.load
// CHECK: arith.addf
// CHECK: memref.store
// CHECK: gpu.return

The real output it matches (show/kernel.mlir, host function then kernel module):

module attributes {gpu.container_module} {
  func.func @add_tensors(%arg0: memref<2x2xf64>, %arg1: memref<2x2xf64>) -> memref<2x2xf64> {
    %alloc = memref.alloc() : memref<2x2xf64>
    %c0 = arith.constant 0 : index
    %c2 = arith.constant 2 : index
    %c1 = arith.constant 1 : index
    %c1_0 = arith.constant 1 : index
    %0 = affine.apply #map(%c2)[%c0, %c1]
    %c2_1 = arith.constant 2 : index
    %c0_2 = arith.constant 0 : index
    %c1_3 = arith.constant 1 : index
    %1 = affine.apply #map(%c2_1)[%c0_2, %c1_3]
    gpu.launch_func  @add_tensors_kernel::@add_tensors_kernel blocks in (%0, %c1_0, %c1_0) threads in (%1, %c1_0, %c1_0)  args(%c1 : index, %c0 : index, %arg0 : memref<2x2xf64>, %arg1 : memref<2x2xf64>, %alloc : memref<2x2xf64>)
    return %alloc : memref<2x2xf64>
  }
  gpu.module @add_tensors_kernel {
    gpu.func @add_tensors_kernel(%arg0: index, %arg1: index, %arg2: memref<2x2xf64>, %arg3: memref<2x2xf64>, %arg4: memref<2x2xf64>) kernel {
      %0 = gpu.block_id  x
      %1 = gpu.block_id  y
      %2 = gpu.block_id  z
      %3 = gpu.thread_id  x
      %4 = gpu.thread_id  y
      %5 = gpu.thread_id  z
      %6 = gpu.grid_dim  x
      %7 = gpu.grid_dim  y
      %8 = gpu.grid_dim  z
      %9 = gpu.block_dim  x
      %10 = gpu.block_dim  y
      %11 = gpu.block_dim  z
      cf.br ^bb1
    ^bb1:  // pred: ^bb0
      %12 = affine.apply #map1(%0)[%arg0, %arg1]
      %c0 = arith.constant 0 : index
      %c2 = arith.constant 2 : index
      %c1 = arith.constant 1 : index
      %c1_0 = arith.constant 1 : index
      %c0_1 = arith.constant 0 : index
      %13 = affine.apply #map1(%3)[%c1_0, %c0_1]
      %14 = memref.load %arg2[%12, %13] : memref<2x2xf64>
      %15 = memref.load %arg3[%12, %13] : memref<2x2xf64>
      %16 = arith.addf %14, %15 : f64
      memref.store %16, %arg4[%12, %13] : memref<2x2xf64>
      gpu.return
    }
  }
}

The first CHECK pins the host side: a gpu.launch_func naming the kernel with a 3-D grid (blocks in) and a 3-D block (threads in); the unused y and z dimensions are all 1. The second pins the kernel module and the third its signature: two index arguments (the loop's step and lower bound, which the generated index arithmetic block_id * step + lb needs) and three memref<2x2xf64> operands. gpu.block_id x and gpu.thread_id x are the row and column. The y/z ids and the grid_dim/block_dim reads are generated but unused; they disappear later (see the PTX below).

nvvm-lowering

// Chapter 10: the kernel lowered to the nvvm dialect: special-register reads, an nvvm.kernel llvm.func, a chip target.
// RUN: %mg-opt %inputs/static_add.mlir --convert-mg-to-affine | %mlir-opt %to-gpu | %mlir-opt --pass-pipeline='builtin.module(lower-affine,nvvm-attach-target{chip=sm_70 features=+ptx60},gpu.module(convert-gpu-to-nvvm{index-bitwidth=64},reconcile-unrealized-casts))' | %FileCheck %s
// CHECK: gpu.module @add_tensors_kernel [#nvvm.target<chip = "sm_70">]
// CHECK: llvm.func @add_tensors_kernel(
// CHECK-SAME: attributes {gpu.kernel, nvvm.kernel}
// CHECK: nvvm.read.ptx.sreg.ctaid.x
// CHECK: nvvm.read.ptx.sreg.tid.x
// CHECK: llvm.fadd

nvvm-attach-target stamps the module with the target chip (sm_70: an arbitrary Volta-class choice, since nothing here runs); convert-gpu-to-nvvm rewrites gpu.block_id x into the NVVM special-register read nvvm.read.ptx.sreg.ctaid.x (and thread_id into tid.x) and the body into the llvm dialect. Real excerpt (show/nvvm.mlir):

  gpu.module @add_tensors_kernel [#nvvm.target<chip = "sm_70">]  {
    llvm.func @add_tensors_kernel(%arg0: i64, %arg1: i64, %arg2: !llvm.ptr, %arg3: !llvm.ptr, %arg4: i64, %arg5: i64, %arg6: i64, %arg7: i64, %arg8: i64, %arg9: !llvm.ptr, %arg10: !llvm.ptr, %arg11: i64, %arg12: i64, %arg13: i64, %arg14: i64, %arg15: i64, %arg16: !llvm.ptr, %arg17: !llvm.ptr, %arg18: i64, %arg19: i64, %arg20: i64, %arg21: i64, %arg22: i64) attributes {gpu.kernel, nvvm.kernel} {
      %0 = llvm.mlir.undef : !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)>

      %24 = nvvm.read.ptx.sreg.ctaid.x : i32
      %25 = llvm.sext %24 : i32 to i64
      %26 = nvvm.read.ptx.sreg.tid.x : i32
      %27 = llvm.sext %26 : i32 to i64
      llvm.br ^bb1

      %38 = llvm.getelementptr %arg10[%37] : (!llvm.ptr, i64) -> !llvm.ptr, f64
      %39 = llvm.load %38 : !llvm.ptr -> f64
      %40 = llvm.fadd %34, %39  : f64
      %41 = llvm.mlir.constant(2 : index) : i64
      %42 = llvm.mul %29, %41  : i64
      %43 = llvm.add %42, %27  : i64
      %44 = llvm.getelementptr %arg17[%43] : (!llvm.ptr, i64) -> !llvm.ptr, f64
      llvm.store %40, %44 : f64, !llvm.ptr
      llvm.return

Look at the kernel signature in the first excerpt: 23 parameters, which is where the next test comes in.

Why 23 parameters

A memref<2x2xf64> is not a bare pointer: at the calling-convention level (Chapter 1) it is a descriptor of an allocated pointer, an aligned pointer, an offset, two sizes and two strides, which is seven scalars. The kernel takes three memrefs plus two index arguments: 2 + 3 x 7 = 23. The kernel then spends its first instructions rebuilding descriptor structs it barely uses.

ptx-default

// Chapter 10: real PTX. 23 kernel parameters (three memrefs expanded into 7 scalars each, plus 2 indices),
// 64-bit index arithmetic, two global f64 loads, one add, one store.
// RUN: %mg-opt %inputs/static_add.mlir --convert-mg-to-affine | %mlir-opt %to-gpu | %mlir-opt --pass-pipeline='builtin.module(lower-affine,nvvm-attach-target{chip=sm_70 features=+ptx60},gpu.module(convert-gpu-to-nvvm{index-bitwidth=64},reconcile-unrealized-casts))' | %mlir-opt --gpu-module-to-binary="format=isa" -o %t.bin.mlir
// RUN: %decode-ptx %t.bin.mlir | %FileCheck %s
// CHECK: .target sm_70
// CHECK: .visible .entry add_tensors_kernel(
// CHECK-COUNT-23: .param .u64 add_tensors_kernel_param_{{[0-9]+}}
// CHECK-NOT: .param .u64 add_tensors_kernel_param_
// CHECK: mov.u32 %{{.*}}, %ctaid.x;
// CHECK: mov.u32 %{{.*}}, %tid.x;
// CHECK: mul.lo.s64
// CHECK: ld.global.f64
// CHECK: ld.global.f64
// CHECK: add.rn.f64
// CHECK: st.global.f64
// CHECK: ret;

--gpu-module-to-binary="format=isa" runs LLVM's NVPTX backend and stores the PTX as a hex-escaped string attribute; %decode-ptx (decode_ptx.py, from Chapter 10) turns that back into text. The real PTX the test matches (show/ptx64.ptx), header and first parameters:

//
// Generated by LLVM NVPTX Back-End
//

.version 6.0
.target sm_70
.address_size 64

    // .globl   add_tensors_kernel

.visible .entry add_tensors_kernel(
    .param .u64 add_tensors_kernel_param_0,
    .param .u64 add_tensors_kernel_param_1,
    .param .u64 add_tensors_kernel_param_2,
    .param .u64 add_tensors_kernel_param_3,
    .param .u64 add_tensors_kernel_param_4,

and the body:

    .reg .b64   %rd<19>;
    .reg .f64   %fd<4>;

    ld.param.u64    %rd1, [add_tensors_kernel_param_0];
    ld.param.u64    %rd2, [add_tensors_kernel_param_17];
    cvta.to.global.u64  %rd3, %rd2;
    ld.param.u64    %rd4, [add_tensors_kernel_param_1];
    ld.param.u64    %rd5, [add_tensors_kernel_param_10];
    cvta.to.global.u64  %rd6, %rd5;
    ld.param.u64    %rd7, [add_tensors_kernel_param_3];
    cvta.to.global.u64  %rd8, %rd7;
    mov.u32     %r1, %ctaid.x;
    cvt.s64.s32     %rd9, %r1;
    mov.u32     %r2, %tid.x;
    cvt.s64.s32     %rd10, %r2;
    mul.lo.s64  %rd11, %rd9, %rd1;
    add.s64     %rd12, %rd11, %rd4;
    shl.b64     %rd13, %rd12, 1;
    add.s64     %rd14, %rd13, %rd10;
    shl.b64     %rd15, %rd14, 3;
    add.s64     %rd16, %rd8, %rd15;
    ld.global.f64   %fd1, [%rd16];
    add.s64     %rd17, %rd6, %rd15;
    ld.global.f64   %fd2, [%rd17];
    add.rn.f64  %fd3, %fd1, %fd2;
    add.s64     %rd18, %rd3, %rd15;
    st.global.f64   [%rd18], %fd3;
    ret;

}

How to read the body (the instruction meanings below are PTX's, stated from the instruction names and this program's behavior; this chapter did not consult NVIDIA's PTX manual):

  • Lines 41 to 48: load the kernel parameters it actually needs: param_0 (the loop step), param_1 (the lower bound), and the three aligned data pointers (param_3, param_10, param_17, one per matrix: the second slot of each seven-field group). cvta.to.global converts each generic address to a global-memory address.
  • Lines 49 to 52: read %ctaid.x (block index) and %tid.x (thread index), widening each to 64 bits (cvt.s64.s32).
  • Lines 53 to 58: the index arithmetic: (block * step + lb) is the row; shl 1 multiplies by 2 (the row stride of a 2-column matrix); adding the thread index gives the element number; shl 3 multiplies by 8 (bytes per f64); adding the base pointer gives the address.
  • Lines 59 to 64: two ld.global.f64 loads, one add.rn.f64, one st.global.f64. This is the whole computation.

The test's CHECK-COUNT-23 followed by CHECK-NOT pins exactly 23 declared parameters (a pattern that matches the .param declarations but not the ld.param instructions in the body, which have a different spelling); mul.lo.s64 pins the 64-bit index arithmetic; the load/add/store sequence pins the computation. Notice what the backend did unasked: of the 23 declared parameters the body loads only five, and the unused block_id y/z, grid_dim and block_dim reads are gone.

The two options

// Chapter 10: index-bitwidth=32 turns the 64-bit index math into 32-bit (mad.lo.s32), without changing the parameter count.
// RUN: %mg-opt %inputs/static_add.mlir --convert-mg-to-affine | %mlir-opt %to-gpu | %mlir-opt --pass-pipeline='builtin.module(lower-affine,canonicalize,nvvm-attach-target{chip=sm_70 features=+ptx60},gpu.module(convert-gpu-to-nvvm{index-bitwidth=32},reconcile-unrealized-casts),gpu-module-to-binary{format=isa})' -o %t.bin.mlir
// RUN: %decode-ptx %t.bin.mlir | %FileCheck %s
// CHECK-COUNT-23: .param .u{{(32|64)}} add_tensors_kernel_param_{{[0-9]+}}
// CHECK-NOT: .param .u{{(32|64)}} add_tensors_kernel_param_
// CHECK: mad.lo.s32
// CHECK-NOT: mul.lo.s64
// CHECK: ld.global.f64

With index-bitwidth=32 the index arithmetic is 32-bit (show/ptx32.ptx):

    ld.param.u32    %r1, [add_tensors_kernel_param_0];
    ld.param.u64    %rd1, [add_tensors_kernel_param_17];
    cvta.to.global.u64  %rd2, %rd1;
    ld.param.u32    %r2, [add_tensors_kernel_param_1];
    ld.param.u64    %rd3, [add_tensors_kernel_param_10];
    cvta.to.global.u64  %rd4, %rd3;
    ld.param.u64    %rd5, [add_tensors_kernel_param_3];
    cvta.to.global.u64  %rd6, %rd5;
    mov.u32     %r3, %ctaid.x;
    mov.u32     %r4, %tid.x;
    mad.lo.s32  %r5, %r3, %r1, %r2;
    shl.b32     %r6, %r5, 1;
    add.s32     %r7, %r6, %r4;
    mul.wide.s32    %rd7, %r7, 8;
    add.s64     %rd8, %rd6, %rd7;
    ld.global.f64   %fd1, [%rd8];
    add.s64     %rd9, %rd4, %rd7;
    ld.global.f64   %fd2, [%rd9];
    add.rn.f64  %fd3, %fd1, %fd2;
    add.s64     %rd10, %rd2, %rd7;
    st.global.f64   [%rd10], %fd3;
    ret;

mad.lo.s32 (multiply-add) replaces the mul.lo.s64/add.s64 pair, and the two index parameters are now declared .u32, which is why the test's parameter pattern is .u{{(32|64)}}. The parameter count stays 23: narrowing the index does not touch the descriptor ABI.

// Chapter 12: kernel-bare-ptr-calling-convention=1 cuts the kernel from 23 parameters to 5 (static shapes only).
// RUN: %mlir-opt %inputs/add_host_device.mlir --pass-pipeline='builtin.module(func.func(gpu-async-region))' -o %t.async.mlir
// RUN: %mlir-opt %t.async.mlir --gpu-lower-to-nvvm-pipeline="cubin-chip=sm_70 cubin-features=+ptx60 cubin-format=isa kernel-bare-ptr-calling-convention=1" -o %t.bin.mlir
// RUN: %decode-ptx %t.bin.mlir | %FileCheck %s
// CHECK-COUNT-5: .param .u64 add_tensors_kernel_param_{{[0-9]+}}
// CHECK-NOT: .param .u64 add_tensors_kernel_param_
// CHECK: ld.global.f64
// CHECK: add.rn.f64
// CHECK: st.global.f64

kernel-bare-ptr-calling-convention=1 (Chapter 12's finding) passes each memref as one raw pointer, which works only because every shape here is static. Real PTX header (show/ptx_bare.ptx):

.visible .entry add_tensors_kernel(
    .param .u64 add_tensors_kernel_param_0,
    .param .u64 add_tensors_kernel_param_1,
    .param .u64 add_tensors_kernel_param_2,
    .param .u64 add_tensors_kernel_param_3,
    .param .u64 add_tensors_kernel_param_4
)
{

Five parameters: the step, the lower bound and three pointers. For scale, the CUDA-style comparison kernel below has three, because it hard-codes the loop bounds.

The CUDA comparison

// Chapter 10: the CUDA-style comparison kernel compiles with clang-18 (no CUDA SDK) to PTX with 3 parameters.
// RUN: %clang -x cuda --cuda-device-only --cuda-gpu-arch=sm_70 -nocudainc -nocudalib -S -O1 %inputs/cuda_add.cu -o %t.ptx
// RUN: %FileCheck %s < %t.ptx
// CHECK: .visible .entry add_cuda(
// CHECK-COUNT-3: .param .u64 add_cuda_param_{{[0-9]+}}
// CHECK-NOT: .param .u64 add_cuda_param_
// CHECK: ld.global.f64
// CHECK: add.f64
// CHECK: st.global.f64

The source it compiles (Inputs/cuda_add.cu; raw NVVM builtins because no CUDA headers are installed) and the PTX clang-18 produced (show/cuda.ptx):

// Equivalent hand-written CUDA-style kernel, compiled without the CUDA SDK.
// 'threadIdx.x'/'blockIdx.x' are NVVM builtins in disguise; we call them directly.
extern "C" __attribute__((global)) void add_cuda(const double *a, const double *b, double *c) {
  int row = __nvvm_read_ptx_sreg_ctaid_x();
  int col = __nvvm_read_ptx_sreg_tid_x();
  int i = row * 2 + col;
  c[i] = a[i] + b[i];
}

.visible .entry add_cuda(
    .param .u64 add_cuda_param_0,
    .param .u64 add_cuda_param_1,
    .param .u64 add_cuda_param_2
)
{
    .reg .b32   %r<5>;
    .reg .b64   %rd<11>;
    .reg .f64   %fd<4>;

// %bb.0:
    ld.param.u64    %rd1, [add_cuda_param_0];
    ld.param.u64    %rd2, [add_cuda_param_2];
    cvta.to.global.u64  %rd3, %rd2;
    ld.param.u64    %rd4, [add_cuda_param_1];
    cvta.to.global.u64  %rd5, %rd4;
    cvta.to.global.u64  %rd6, %rd1;
    mov.u32     %r1, %ctaid.x;
    mov.u32     %r2, %tid.x;
    shl.b32     %r3, %r1, 1;
    add.s32     %r4, %r3, %r2;
    mul.wide.s32    %rd7, %r4, 8;
    add.s64     %rd8, %rd6, %rd7;
    ld.global.f64   %fd1, [%rd8];
    add.s64     %rd9, %rd5, %rd7;
    ld.global.f64   %fd2, [%rd9];
    add.f64     %fd3, %fd1, %fd2;
    add.s64     %rd10, %rd3, %rd7;
    st.global.f64   [%rd10], %fd3;
    ret;
                                        // -- End function
}

Same memory traffic as the MLIR kernel (two global loads, one add, one store), three parameters, 32-bit index arithmetic.

Part 3: the host side (Chapters 11 and 12)

Background: the runtime-call sequence

A launch needs more than a kernel: device buffers are allocated and filled, the kernel module is loaded, the kernel is launched on a stream with a grid and block size and a parameter array, and everything is released. MLIR 18 lowers these to calls into a GPU runtime library. The input is the hand-written host function Chapter 11 introduced (gpu.alloc, gpu.memcpy, gpu.launch_func, gpu.dealloc), made async with gpu-async-region first because the lowering patterns only match async ops (Chapter 12 read that in MLIR's source):

#map = affine_map<(d0)[s0, s1] -> ((d0 - s0) ceildiv s1)>
#map1 = affine_map<(d0)[s0, s1] -> (d0 * s0 + s1)>
module attributes {gpu.container_module} {
  func.func @add_tensors(%a: memref<2x2xf64>, %b: memref<2x2xf64>) -> memref<2x2xf64> {
    %da = gpu.alloc () : memref<2x2xf64>
    %db = gpu.alloc () : memref<2x2xf64>
    %dc = gpu.alloc () : memref<2x2xf64>
    gpu.memcpy %da, %a : memref<2x2xf64>, memref<2x2xf64>
    gpu.memcpy %db, %b : memref<2x2xf64>, memref<2x2xf64>
    %c0 = arith.constant 0 : index
    %c1 = arith.constant 1 : index
    %c2 = arith.constant 2 : index
    gpu.launch_func @add_tensors_kernel::@add_tensors_kernel blocks in (%c2, %c1, %c1) threads in (%c2, %c1, %c1) args(%c1 : index, %c0 : index, %da : memref<2x2xf64>, %db : memref<2x2xf64>, %dc : memref<2x2xf64>)
    %r = memref.alloc() : memref<2x2xf64>
    gpu.memcpy %r, %dc : memref<2x2xf64>, memref<2x2xf64>
    gpu.dealloc %da : memref<2x2xf64>
    gpu.dealloc %db : memref<2x2xf64>
    gpu.dealloc %dc : memref<2x2xf64>
    return %r : memref<2x2xf64>
  }
  gpu.module @add_tensors_kernel {
    gpu.func @add_tensors_kernel(%arg0: index, %arg1: index, %arg2: memref<2x2xf64>, %arg3: memref<2x2xf64>, %arg4: memref<2x2xf64>) kernel {
      %0 = gpu.block_id  x
      %1 = gpu.block_id  y
      %2 = gpu.block_id  z
      %3 = gpu.thread_id  x
      %4 = gpu.thread_id  y
      %5 = gpu.thread_id  z
      %6 = gpu.grid_dim  x
      %7 = gpu.grid_dim  y
      %8 = gpu.grid_dim  z
      %9 = gpu.block_dim  x
      %10 = gpu.block_dim  y
      %11 = gpu.block_dim  z
      cf.br ^bb1
    ^bb1:  // pred: ^bb0
      %12 = affine.apply #map1(%0)[%arg0, %arg1]
      %c0 = arith.constant 0 : index
      %c2 = arith.constant 2 : index
      %c1 = arith.constant 1 : index
      %c1_0 = arith.constant 1 : index
      %c0_1 = arith.constant 0 : index
      %13 = affine.apply #map1(%3)[%c1_0, %c0_1]
      %14 = memref.load %arg2[%12, %13] : memref<2x2xf64>
      %15 = memref.load %arg3[%12, %13] : memref<2x2xf64>
      %16 = arith.addf %14, %15 : f64
      memref.store %16, %arg4[%12, %13] : memref<2x2xf64>
      gpu.return
    }
  }
}

host-runtime-calls

// Chapter 11: gpu.launch_func becomes mgpu* runtime calls at mlir-translate time; grid 2x1x1, block 2x1x1, 23 params.
// RUN: %mlir-opt %inputs/add_host_device.mlir --pass-pipeline='builtin.module(func.func(gpu-async-region))' -o %t.async.mlir
// RUN: %mlir-opt %t.async.mlir --gpu-lower-to-nvvm-pipeline="cubin-chip=sm_70 cubin-features=+ptx60 cubin-format=isa" -o %t.pipe.mlir
// RUN: %mlir-translate --mlir-to-llvmir %t.pipe.mlir -o %t.ll
// RUN: %FileCheck %s < %t.ll
// CHECK: call ptr @mgpuStreamCreate
// CHECK-COUNT-3: call ptr @mgpuMemAlloc
// CHECK-COUNT-2: call void @mgpuMemcpy
// CHECK: call ptr @mgpuModuleLoadJIT(ptr @add_tensors_kernel_bin_cst, i32 2)
// CHECK: call ptr @mgpuModuleGetFunction
// CHECK: call void @mgpuLaunchKernel(ptr %{{.*}}, i64 2, i64 1, i64 1, i64 2, i64 1, i64 1, i32 0, ptr %{{.*}}, ptr %{{.*}}, ptr null, i64 23)
// CHECK: call void @mgpuModuleUnload
// CHECK: call void @mgpuMemcpy
// CHECK-COUNT-3: call void @mgpuMemFree

The matching lines of the translated LLVM IR (show/host.ll): allocation and copies in, then the launch, then copy out and release:

  %15 = call ptr @mgpuStreamCreate()
  %16 = call ptr @mgpuMemAlloc(i64 ptrtoint (ptr getelementptr (double, ptr null, i32 4) to i64), ptr %15, i8 0)
  %17 = call ptr @mgpuMemAlloc(i64 ptrtoint (ptr getelementptr (double, ptr null, i32 4) to i64), ptr %15, i8 0)
  %18 = call ptr @mgpuMemAlloc(i64 ptrtoint (ptr getelementptr (double, ptr null, i32 4) to i64), ptr %15, i8 0)
  call void @mgpuMemcpy(ptr %16, ptr %1, i64 ptrtoint (ptr getelementptr (double, ptr null, i32 4) to i64), ptr %15)
  call void @mgpuMemcpy(ptr %17, ptr %8, i64 ptrtoint (ptr getelementptr (double, ptr null, i32 4) to i64), ptr %15)

  %67 = call ptr @mgpuModuleLoadJIT(ptr @add_tensors_kernel_bin_cst, i32 2)
  %68 = call ptr @mgpuModuleGetFunction(ptr %67, ptr @add_tensors_kernel_add_tensors_kernel_kernel_name)
  call void @mgpuLaunchKernel(ptr %68, i64 2, i64 1, i64 1, i64 2, i64 1, i64 1, i32 0, ptr %15, ptr %20, ptr null, i64 23)
  call void @mgpuModuleUnload(ptr %67)
  call void @mgpuStreamSynchronize(ptr %15)
  call void @mgpuStreamDestroy(ptr %15)

  %77 = call ptr @mgpuStreamCreate()
  call void @mgpuMemcpy(ptr %69, ptr %18, i64 ptrtoint (ptr getelementptr (double, ptr null, i32 4) to i64), ptr %77)
  call void @mgpuMemFree(ptr %16, ptr %77)
  call void @mgpuMemFree(ptr %17, ptr %77)
  call void @mgpuMemFree(ptr %18, ptr %77)
  call void @mgpuStreamSynchronize(ptr %77)
  call void @mgpuStreamDestroy(ptr %77)

The launch call's arguments, in order: the function handle; grid x, y, z = 2, 1, 1; block x, y, z = 2, 1, 1; shared memory 0; the stream; a pointer to an array of kernel parameter pointers; a null extra; and the parameter count, 23. The test pins the grid, block and count exactly. host-runtime-calls-bare is the same test with the bare-pointer option and a final i64 5:

// Chapter 12: with bare pointers the same launch passes 5 parameters instead of 23.
// RUN: %mlir-opt %inputs/add_host_device.mlir --pass-pipeline='builtin.module(func.func(gpu-async-region))' -o %t.async.mlir
// RUN: %mlir-opt %t.async.mlir --gpu-lower-to-nvvm-pipeline="cubin-chip=sm_70 cubin-features=+ptx60 cubin-format=isa kernel-bare-ptr-calling-convention=1" -o %t.pipe.mlir
// RUN: %mlir-translate --mlir-to-llvmir %t.pipe.mlir -o %t.ll
// RUN: %FileCheck %s < %t.ll
// CHECK: call void @mgpuLaunchKernel(ptr %{{.*}}, i64 2, i64 1, i64 1, i64 2, i64 1, i64 1, i32 0, ptr %{{.*}}, ptr %{{.*}}, ptr null, i64 5)

pass-order-crash: a test that pins a bug

// Chapter 12: PINS A KNOWN LLVM 18.1.3 BUG. gpu-to-llvm after gpu-module-to-binary segfaults mlir-opt (the launch
// pattern looks the kernel module up as a gpu::GPUModuleOp, which has already become a gpu.binary). If a newer
// toolchain fixes this, this test starts failing: that is the signal to revisit Chapters 11 and 12, not a regression.
// RUN: %mlir-opt %inputs/add_host_device.mlir --pass-pipeline='builtin.module(func.func(gpu-async-region))' -o %t.async.mlir
// RUN: %not --crash %mlir-opt %t.async.mlir --pass-pipeline='builtin.module(lower-affine,nvvm-attach-target{chip=sm_70 features=+ptx60},gpu.module(convert-gpu-to-nvvm{index-bitwidth=64},reconcile-unrealized-casts),gpu-module-to-binary{format=isa},gpu-to-llvm)' -o /dev/null

Chapter 12 established, from MLIR's source, why running gpu-to-llvm after gpu-module-to-binary segfaults mlir-opt-18: the launch pattern looks the kernel module up by type, and gpu-module-to-binary has already replaced it with a gpu.binary. This test asserts the crash with %not --crash. It is deliberately unusual: it is designed to start failing. If a newer LLVM fixes the bug the test turns red, and that is the signal to revisit Chapters 11 and 12, not a regression in this book. The test's header comment says so for whoever sees it fail. It documents a toolchain property rather than a book property, which is accepted here because those chapters rest on it.

gpu-host-with-stub-runtime

There is no GPU runtime library on this machine, so Chapter 11 wrote a CPU stand-in. Its source, in full:

// A CPU stand-in for MLIR's GPU runtime wrappers (the mgpu* entry points the lowered host code calls).
// There is no GPU here, so "device memory" is malloc and the "kernel" is a C function that reads the
// SAME launch parameters the PTX reads (params 0, 1, 3, 10, 17) and does the same arithmetic.
// This checks the HOST side (allocation, copies, launch geometry, 23-parameter packing). It does NOT
// execute the PTX, so it says nothing about whether the PTX itself is correct.
#include <stdint.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>

void *mgpuStreamCreate(void) { fprintf(stderr, "[stub] stream create\n"); return malloc(1); }
void mgpuStreamDestroy(void *s) { free(s); }
void mgpuStreamSynchronize(void *s) { (void)s; }
void *mgpuMemAlloc(uint64_t n, void *s, int shared) { (void)s; (void)shared; fprintf(stderr, "[stub] alloc %llu bytes\n", (unsigned long long)n); return malloc(n); }
void mgpuMemFree(void *p, void *s) { (void)s; free(p); }
void mgpuMemcpy(void *dst, void *src, uint64_t n, void *s) { (void)s; fprintf(stderr, "[stub] memcpy %llu bytes\n", (unsigned long long)n); memcpy(dst, src, n); }

static const char *kernel_name_loaded;
void *mgpuModuleLoadJIT(void *ptx, int optLevel) {
  const char *p = (const char *)ptx;
  fprintf(stderr, "[stub] module load: optLevel=%d, PTX has .entry add_tensors_kernel: %s\n", optLevel,
          strstr(p, ".entry add_tensors_kernel") ? "yes" : "NO");
  return (void *)p;
}
void *mgpuModuleGetFunction(void *mod, const char *name) { (void)mod; kernel_name_loaded = name; return (void *)name; }
void mgpuModuleUnload(void *mod) { (void)mod; }

void mgpuLaunchKernel(void *fn, intptr_t gx, intptr_t gy, intptr_t gz, intptr_t bx, intptr_t by, intptr_t bz,
                      int32_t smem, void *stream, void **params, void **extra, size_t nparams) {
  (void)gy; (void)gz; (void)by; (void)bz; (void)smem; (void)stream; (void)extra;
  fprintf(stderr, "[stub] launch %s: grid=(%ld,1,1) block=(%ld,1,1) nparams=%zu\n", (const char *)fn, (long)gx, (long)bx, nparams);
  // Two parameter layouts, depending on the kernel calling convention the host was compiled with:
  //   23 = memref descriptors expanded (default); 5 = bare pointers (kernel-bare-ptr-calling-convention).
  if (nparams != 23 && nparams != 5) { fprintf(stderr, "[stub] unexpected parameter count\n"); exit(2); }
  int64_t step = *(int64_t *)params[0], lb = *(int64_t *)params[1];
  double *a, *b, *c;
  if (nparams == 23) { a = *(double **)params[3]; b = *(double **)params[10]; c = *(double **)params[17]; }
  else               { a = *(double **)params[2]; b = *(double **)params[3];  c = *(double **)params[4]; }
  for (intptr_t blk = 0; blk < gx; blk++)
    for (intptr_t t = 0; t < bx; t++) {
      int64_t i = (blk * step + lb) * 2 + t;
      c[i] = a[i] + b[i];
    }
}

and the test that links the lowered host code with it, once with the default 23-parameter launch and once with 5:

// Chapter 11/12: the lowered host code, linked with a CPU STUB runtime, prints the right answer, with both the
// 23-parameter and the 5-parameter launch. This tests the host glue; it does NOT execute the PTX.
// RUN: %mlir-opt %inputs/add_host_device.mlir --pass-pipeline='builtin.module(func.func(gpu-async-region))' -o %t.async.mlir
// RUN: %mlir-opt %t.async.mlir --gpu-lower-to-nvvm-pipeline="cubin-chip=sm_70 cubin-features=+ptx60 cubin-format=isa" -o %t.p23.mlir
// RUN: %mlir-translate --mlir-to-llvmir %t.p23.mlir -o %t.p23.ll
// RUN: %clang -c %t.p23.ll -o %t.p23.o
// RUN: %clang %inputs/static_harness.c %t.p23.o %inputs/stub_gpu_runtime.c -o %t.p23.exe
// RUN: %t.p23.exe 2>%t.p23.err | %FileCheck %s
// RUN: %FileCheck %s --check-prefix=LAUNCH23 < %t.p23.err
// RUN: %mlir-opt %t.async.mlir --gpu-lower-to-nvvm-pipeline="cubin-chip=sm_70 cubin-features=+ptx60 cubin-format=isa kernel-bare-ptr-calling-convention=1" -o %t.p5.mlir
// RUN: %mlir-translate --mlir-to-llvmir %t.p5.mlir -o %t.p5.ll
// RUN: %clang -c %t.p5.ll -o %t.p5.o
// RUN: %clang %inputs/static_harness.c %t.p5.o %inputs/stub_gpu_runtime.c -o %t.p5.exe
// RUN: %t.p5.exe 2>%t.p5.err | %FileCheck %s
// RUN: %FileCheck %s --check-prefix=LAUNCH5 < %t.p5.err
// CHECK: 6 8
// CHECK-NEXT: 10 12
// LAUNCH23: [stub] module load: optLevel=2, PTX has .entry add_tensors_kernel: yes
// LAUNCH23: [stub] launch add_tensors_kernel: grid=(2,1,1) block=(2,1,1) nparams=23
// LAUNCH5: [stub] launch add_tensors_kernel: grid=(2,1,1) block=(2,1,1) nparams=5

The test checks both the printed answer (6 8 / 10 12) and the stub's own trace on stderr (grid=(2,1,1) block=(2,1,1) nparams=23 or 5). This tests the host glue only. The stub's "kernel" is C that mirrors the PTX, not an execution of it.

Two mistakes along the way

1. A wrong expectation, not a compiler bug. The first full run had 33 passes and 1 failure, gpu/ptx-index32. The test counted 23 parameters all declared .u64; with index-bitwidth=32 the two index parameters are .u32, and only seven lines matched:

ptx-index32.mlir:4:20: error: CHECK-COUNT: expected string not found in input (7 out of 23)

The compiler was right and the test was wrong. The pattern now accepts either width (.u{{(32|64)}}).

2. A bug in the mutation tool itself. Writing the mutation script below, mutation 1 (change index-bitwidth=64 to 32) was reported as passing: the test had not noticed. The cause was a sed expression written as 'index-bitwidth=64/s/64/32/', with no slashes around the address. sed parsed it as an insert command, so instead of substituting it added a junk line of text after every line of the test file. Diffing the "mutated" copy against the original showed:

0a1
> ndex-bitwidth=64/s/64/32/
1a3
> ndex-bitwidth=64/s/64/32/
...

The recipe was never changed, and the tool had counted "the file differs" as "the mutation applied". Two fixes: the expression was corrected ('/index-bitwidth=64/s/64/32/'), and the script now rejects any mutation that changes the line count, since a substitution never should. After the fix the same mutation fails the test and its diff is exactly one line. The lesson is about the whole method: a mutation check that can silently apply the wrong mutation is itself an untested test.

Can the new tests fail?

recipe_mutation.sh copies the suite, changes one thing in the copy, and runs only the affected test. These tests pin recipes and a stub runtime rather than mg source, so the mutations edit those:

#!/bin/sh
# Chapter 16: do the loop and GPU tests fail when the recipe or runtime they pin is changed?
# Each mutation edits ONE file in a COPY of the suite, then runs only the affected test(s) there.
# READ THIS FIRST: this script injects deliberate bugs into the GPU recipe and stub runtime and re-runs the tests.
#   FAIL lines (and non-zero exit statuses) in this script's output are the EXPECTED result: each injected bug SHOULD make at least one test fail.
#   The unmodified, current build is the baseline and must show no failures. A mutation or old build that makes NO test fail would be the problem.
echo "NOTE: this script deliberately runs broken or older code. FAIL lines below are EXPECTED: they show the tests can detect the problem. The baseline (unmodified current build) must show none."
HERE=$(cd "$(dirname "$0")" && pwd)
export MG_OPT=${MG_OPT:-$HERE/../../part14/code/build/mg-opt}
export DECODE_PTX=$(cd "$HERE/../../part10/code" && pwd)/decode_ptx.py
n=0
mut() {  # description | file under test/ | sed expression | test(s) to run
  n=$((n+1)); W=$HERE/work/rm-$n; rm -rf "$W"; mkdir -p "$W"; cp -r "$HERE/test" "$W/test"
  sed -i "$3" "$W/test/$2"
  if cmp -s "$HERE/test/$2" "$W/test/$2"; then echo "MUTATION $n NOT APPLIED ($1)"; return; fi
  # a substitution edits lines in place; if the line count changed, the sed expression did something else
  if [ "$(wc -l < "$HERE/test/$2")" != "$(wc -l < "$W/test/$2")" ]; then echo "MUTATION $n MALFORMED ($1): line count changed"; return; fi
  out=$(MG_TEST_TMP="$W/out" lit "$W/test/$4" 2>&1 | grep -E "^(PASS|FAIL)" | sed 's/IR Forge \/ Mountain Goat :: //;s/ ([0-9]* of [0-9]*)//')
  echo "MUTATION $n: $1"; echo "   $out"
}
mut "index-bitwidth=64 -> 32"                       gpu/ptx-default.mlir                '/index-bitwidth=64/s/64/32/' gpu/ptx-default.mlir
mut "target chip sm_70 -> sm_80"                     gpu/ptx-default.mlir                's/chip=sm_70/chip=sm_80/' gpu/ptx-default.mlir
mut "bare-pointer option dropped (5 -> 23 params)"   gpu/ptx-bare-pointers.mlir          's/ kernel-bare-ptr-calling-convention=1//' gpu/ptx-bare-pointers.mlir
mut "bare-pointer option added to the host test"     gpu/host-runtime-calls.mlir         '/RUN: %mlir-opt %t.async.mlir/s/cubin-format=isa"/cubin-format=isa kernel-bare-ptr-calling-convention=1"/' gpu/host-runtime-calls.mlir
mut "gpu-map-parallel-loops left out of the recipe"  gpu/kernel-outlining.mlir           's#%to-gpu#--pass-pipeline="builtin.module(func.func(affine-parallelize,lower-affine,convert-parallel-loops-to-gpu),gpu-kernel-outlining)"#' gpu/kernel-outlining.mlir
mut "tile-size 1 -> 2"                               loops/tiling.mlir                   's/tile-size=1/tile-size=2/' loops/tiling.mlir
mut "second full-unroll pass removed"                loops/unroll-full.mlir              's/ --affine-loop-unroll="unroll-full" --canonicalize/ --canonicalize/' loops/unroll-full.mlir
mut "stub runtime's kernel computes a-b, not a+b"    Inputs/stub_gpu_runtime.c           's/c\[i\] = a\[i\] + b\[i\];/c[i] = a[i] - b[i];/' execution/gpu-host-with-stub-runtime.mlir

Real output:

MUTATION 1: index-bitwidth=64 -> 32
   FAIL: gpu/ptx-default.mlir
MUTATION 2: target chip sm_70 -> sm_80
   FAIL: gpu/ptx-default.mlir
MUTATION 3: bare-pointer option dropped (5 -> 23 params)
   FAIL: gpu/ptx-bare-pointers.mlir
MUTATION 4: bare-pointer option added to the host test
   FAIL: gpu/host-runtime-calls.mlir
MUTATION 5: gpu-map-parallel-loops left out of the recipe
   FAIL: gpu/kernel-outlining.mlir
MUTATION 6: tile-size 1 -> 2
   FAIL: loops/tiling.mlir
MUTATION 7: second full-unroll pass removed
   FAIL: loops/unroll-full.mlir
MUTATION 8: stub runtime's kernel computes a-b, not a+b
   FAIL: execution/gpu-host-with-stub-runtime.mlir

All eight are caught, but they are not equally informative. For mutations 2 (a different chip) and 5 (a recipe missing gpu-map-parallel-loops) the broken recipe produced no output at all: FileCheck reported its first pattern missing from an empty input. Those tests failed because the pipeline errored, which proves only "this recipe still works", not "this specific line of output is still right". For mutations 1, 3, 4, 6, 7 and 8 the pipeline ran and a specific value changed (the index width, the parameter count, the tile structure, the unrolled loops, the printed numbers), which is the stronger evidence.

Chapter 15's checks, re-run with 34 tests

The older-build comparison and the three source mutations from Chapter 15 were re-run unchanged against the larger suite. Older builds (older_out_34.txt):

[100%] Built target mg-opt
=================== Chapter 7 build (no dynamic shapes)
  Failed: 13 (38.24%)
  Passed: 21 (61.76%)
FAIL: canonicalize/transpose-twice-dynamic-mixed.mlir
FAIL: execution/dynamic-shapes-run.mlir
FAIL: execution/mismatch-aborts.mlir
FAIL: execution/mismatch-columns-aborts.mlir
FAIL: execution/mixed-static-dynamic.mlir
FAIL: execution/strided-view.mlir
FAIL: lowering/bufferize-dynamic-add.mlir
FAIL: lowering/dynamic-add-affine.mlir
FAIL: lowering/dynamic-transpose-no-check.mlir
FAIL: lowering/mixed-add-one-check.mlir
FAIL: verifier/add-static-mismatch.mlir
FAIL: verifier/dynamic-accepted.mlir
FAIL: verifier/unranked-rejected.mlir
=================== Chapter 13 build (dynamic shapes, no runtime check)
  Failed:  6 (17.65%)
  Passed: 28 (82.35%)
FAIL: execution/mismatch-aborts.mlir
FAIL: execution/mismatch-columns-aborts.mlir
FAIL: execution/mixed-static-dynamic.mlir
FAIL: lowering/bufferize-dynamic-add.mlir
FAIL: lowering/dynamic-add-affine.mlir
FAIL: lowering/mixed-add-one-check.mlir

Source mutations (mutation_out_34.txt):

##### MUTATION 1: canonicalizer's exact-type guard removed
  Failed:  1 (2.94%)
  Passed: 33 (97.06%)
FAIL: canonicalize/transpose-twice-dynamic-mixed.mlir

##### MUTATION 2: runtime check only covers dimension 0
  Failed:  3 (8.82%)
  Passed: 31 (91.18%)
FAIL: execution/mismatch-columns-aborts.mlir
FAIL: lowering/bufferize-dynamic-add.mlir
FAIL: lowering/dynamic-add-affine.mlir

##### MUTATION 3: bufferization path forgets to emit the check
  Failed:  3 (8.82%)
  Passed: 31 (91.18%)
FAIL: execution/mismatch-aborts.mlir
FAIL: execution/mismatch-columns-aborts.mlir
FAIL: lowering/bufferize-dynamic-add.mlir

The same tests fail as in Chapter 15 (13 on Chapter 7's build, 6 on Chapter 13's, and the same three mutation results), and all 12 new tests pass on both older builds. That is expected, and also a limit worth stating: the new tests exercise the GPU and loop pipelines, which those builds share, so they cannot tell those builds apart. They guard recipes and toolchain behavior, not the dynamic-shape features of Chapters 13 and 14.

What this does and does not establish

  • The PTX is checked as text. The tests confirm that the compiler emits PTX with particular instructions, parameter counts and targets. The PTX was never assembled or run: there is no ptxas, driver or GPU here, so no test can say the kernel computes the right answer on hardware.
  • The stub-runtime test checks the host glue only. Mutation 8 shows it catches a wrong stub (a - b instead of a + b); it says nothing about the real PTX.
  • PTX text is a brittle thing to pin. Register numbers are not matched, but instruction choice (mul.lo.s64, mad.lo.s32) and parameter counts are. A different LLVM could emit different, equally correct instructions and break these tests with nothing wrong.
  • cuda-comparison needs clang-18's CUDA frontend. It ran here without the CUDA SDK, using raw NVVM builtins, as in Chapter 10.
  • No CI. The suite still runs only when someone runs it.
  • Chapters 1, 2 and 9 remain untested: the original scf.for demonstration, the dialect's own build, and the ExecutionEngine host program.
  • Dynamic-bound loop transforms (Chapter 7 on ? extents) are unexamined; the new loop tests use the static fused example. (Update, added after Chapter 17: now examined and tested there, taking the suite to 42 tests.)

Reproducing this chapter

cd docs/part15/code
./run_lit.sh              # 34 tests (needs ../../part14/code/build.sh first, and: pip install lit)
./show.sh                 # regenerates every output shown on this page under show/
./recipe_mutation.sh      # the eight recipe/stub mutations
./older_builds.sh         # the 34-test suite against Chapter 7's and 13's builds
./mutation.sh             # Chapter 15's three source mutations

Every file and output shown above is embedded from the repository when the site is built, so this page cannot drift from the code it describes (a missing file fails the build). As in Chapters 10 through 15, no documentation was consulted for any claim; each result comes from running the commands.

Chapter summary

This chapter extended the suite from 22 to 34 tests, covering Chapter 7's tiling and full unrolling and the whole GPU path of Chapters 10 through 12 without a GPU: kernel outlining, NVVM lowering, real PTX (23 parameters by default, 5 with bare pointers, 3 for the CUDA-style kernel), the index-bitwidth option, the host runtime-call sequence, the CUDA comparison compile, and a run of the lowered host code against a CPU stub runtime. One test deliberately pins Chapter 12's LLVM 18.1.3 crash and is designed to start failing when a newer LLVM fixes it. Eight recipe and stub mutations were all caught, though two only because the broken recipe produced no output. Two mistakes are recorded: a wrong expectation (the index parameters become .u32 at 32-bit width) and a bug in the new mutation script itself (a sed expression that inserted lines instead of substituting, so a mutation had never been applied).

Deliberately out of scope, stated explicitly: the PTX is checked only as text and has never been assembled or run, the stub test checks host glue only, pinning PTX instruction choices is brittle across LLVM versions, no CI runs the suite, and Chapters 1, 2 and 9 are still untested.

Self-check questions

1. gpu/pass-order-crash.mlir is designed to start failing. Why write such a test, and what should someone do when it fails?

Worked answer: it pins a property of the toolchain, the LLVM 18.1.3 segmentation fault when gpu-to-llvm runs after gpu-module-to-binary, that Chapters 11 and 12 depend on. If a newer LLVM fixes the bug, the test goes red. That is the signal to revisit those two chapters and the recipe, not a regression in this book, and the test's header says so for whoever sees it fail.

2. Eight recipe mutations were all caught, but the chapter calls two of the failures weaker evidence. Which two, and why?

Worked answer: mutation 2 (a different chip) and mutation 5 (a recipe missing gpu-map-parallel-loops). In both the broken recipe produced no output at all, so FileCheck reported its first pattern missing from an empty input. Those tests failed because the pipeline errored, which shows only that the recipe still works. For the other six (index width, parameter count in two tests, tile structure, unrolled loops, printed numbers) the pipeline ran and a specific value changed, which is the stronger evidence.

3. What went wrong the first time mutation 1 ran, and what guard was added?

Worked answer: the sed expression 'index-bitwidth=64/s/64/32/' had no slashes around its address, so sed read it as an insert command and added a junk line after every line of the test file instead of substituting. The recipe was never changed, so the unchanged test passed, and the script had treated "the file differs" as "the mutation applied". The expression was corrected and the script now rejects any mutation that changes the file's line count, since a substitution never should.

4. All 12 new tests pass on the Chapter 7 and Chapter 13 builds. What does that tell you about what they test?

Worked answer: they pin the GPU and loop pipelines and the toolchain's behavior, which those builds share, not the dynamic-shape features added in Chapters 13 and 14. That is why they cannot tell those builds apart, and it is expected: the earlier 22 tests are the ones that distinguish them.

5. What does the stub-runtime execution test tell you, and what does it not?

Worked answer: it checks the host glue: the program prints 6 8 / 10 12 with both the 23-parameter and the 5-parameter launch, and the stub's trace confirms the grid, block and parameter count. It does not execute the PTX, because the stub's kernel is C code written to mirror it. Mutation 8, changing the stub's a + b to a - b, shows the test catches a wrong stub, which is a statement about the test, not about the real kernel.