15. A Real Test Suite: lit, FileCheck, and Proof That the Tests Can Fail¶
What you will understand: how to turn this book's "run it once and paste the output" scripts into an assertion-based regression suite that fails when behavior changes, and, more important, how to show that it can fail. Every one of the suite's 22 original tests is printed here exactly as it exists in the repository, next to the real output it checks, so you can see what each pattern is matching and why. Chapters 12 through 14 each flagged the same gap ("no lit/FileCheck tests, no CI"); this chapter closes the first half of it. Continuous integration is not done here.
What you need to know first: Chapters 3 through 14's behavior. The suite tests what those chapters built; it adds no compiler features.
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.
Why a test suite, and why this kind¶
Chapters 3 through 14 each ended with a pasted output, produced by a command run once. That proves the command worked then. It does not stop the next change from quietly breaking it, and two of the chapters already contain bugs found only later (Chapter 13 fixed a canonicalizer that produced invalid IR; Chapter 14 closed an out-of-bounds read). A test suite turns each pasted output into a check that runs again and fails loudly when the behavior changes.
MLIR's own project uses two tools for exactly this, and so does this suite:
litdiscovers test files, runs the shell commands written in each one's// RUN:lines, and reports pass or fail by exit status.FileCheckreads text on its standard input and checks it against// CHECK:patterns written in the same test file. A test is "run the compiler, pipe its output into FileCheck".
How lit finds and runs a test¶
lit.cfg.py tells lit what a test is (a .mlir file), where to run it, and which %-names to substitute inside RUN: lines. The suite's, in full:
# lit configuration for the Mountain Goat (mg) regression suite.
# Which mg-opt is tested is chosen by the MG_OPT environment variable, so the same suite can be pointed at
# different builds (Chapter 15 does exactly that to prove the tests can fail).
import os
import lit.formats
config.name = "IR Forge / Mountain Goat"
config.test_format = lit.formats.ShTest(execute_external=False)
config.suffixes = [".mlir"]
config.excludes = ["Inputs"]
config.test_source_root = os.path.dirname(__file__)
config.test_exec_root = os.environ.get("MG_TEST_TMP", "/tmp/mg-lit-tmp")
mg_opt = os.environ.get("MG_OPT")
if not mg_opt:
raise RuntimeError("set MG_OPT to the mg-opt binary under test")
inputs = os.path.join(config.test_source_root, "Inputs")
# Chapter 10's recipe: mg.add's affine loops -> a gpu.module kernel (blocks = rows, threads = columns).
to_gpu = ("--pass-pipeline='builtin.module(func.func(affine-parallelize,lower-affine,gpu-map-parallel-loops,"
"convert-parallel-loops-to-gpu),gpu-kernel-outlining)'")
lower = ("--lower-affine --convert-scf-to-cf --convert-arith-to-llvm --finalize-memref-to-llvm "
"--convert-func-to-llvm --reconcile-unrealized-casts")
config.substitutions += [
("%mg-opt", mg_opt),
("%mlir-translate", "mlir-translate-18"),
("%mlir-opt", "mlir-opt-18"),
("%decode-ptx", "python3 " + os.environ.get("DECODE_PTX", os.path.join(config.test_source_root, "..", "..", "..", "part10", "code", "decode_ptx.py"))),
("%clang", "clang-18"),
("%FileCheck", "FileCheck-18"),
("%not", "not-18"),
("%inputs", inputs),
("%lower-to-llvm", lower),
("%lower-with-stderr", "--lower-affine --convert-scf-to-cf --mg-lower-assert-to-stderr --convert-arith-to-llvm "
"--finalize-memref-to-llvm --convert-func-to-llvm --reconcile-unrealized-casts"),
("%to-gpu", to_gpu),
]
# The language tour (docs/tour): its examples are run in place, like the chapters' examples.
config.substitutions += [("%tour", os.path.join(config.test_source_root, "..", "..", "..", "tour", "code"))]
# Chapter 40: GA-1, a cycle-counting model of a matrix accelerator (pure Python; no compiler involved).
ch40 = os.path.join(config.test_source_root, "..", "..", "..", "part40", "code")
config.substitutions += [("%ch40", ch40)]
# Chapter 39: a differential profile (valgrind callgrind) of the lowering pass. No compiler change.
ch39 = os.path.join(config.test_source_root, "..", "..", "..", "part39", "code")
config.substitutions += [("%ch39", ch39)]
# Chapter 38: outlining loop nests (--mg-outline-loops, mgc --outline). '%mgc38' must come before the shorter '%mgc'.
ch38 = os.path.join(config.test_source_root, "..", "..", "..", "part38", "code")
config.substitutions += [
("%mgc38", "env MG_OPT=" + mg_opt + " " + os.path.join(ch38, "mgc")),
("%ex38", os.path.join(ch38, "examples")),
("%ch38", ch38),
]
# Chapter 37: inside the LLVM stage (reading the IR, the passes, the vectorizer's reasons, the backend). No new compiler operation; the tests run LLVM's own tools.
ch37 = os.path.join(config.test_source_root, "..", "..", "..", "part37", "code")
config.substitutions += [
("%ex37", os.path.join(ch37, "examples")),
("%ch37", ch37),
]
# Chapter 36: Chapter 30's architecture (two blocks, two heads each) trained with a hand-derived backward pass. No new compiler operation ('%mgc32' runs the examples).
ch36 = os.path.join(config.test_source_root, "..", "..", "..", "part36", "code")
config.substitutions += [
("%ex36", os.path.join(ch36, "examples")),
("%ch36", ch36),
]
# Chapter 35: a one-block causal transformer language model trained with a hand-derived backward pass. No new compiler operation ('%mgc32' runs the examples).
ch35 = os.path.join(config.test_source_root, "..", "..", "..", "part35", "code")
config.substitutions += [
("%ex35", os.path.join(ch35, "examples")),
("%ch35", ch35),
]
# Chapter 34: why a block has a feed-forward network (backward through layer norm, relu and residuals). No new compiler operation ('%mgc32' runs the examples).
ch34 = os.path.join(config.test_source_root, "..", "..", "..", "part34", "code")
config.substitutions += [
("%ex34", os.path.join(ch34, "examples")),
("%ch34", ch34),
]
# Chapter 33: an attention classifier trained by hand-derived backpropagation. No new compiler operation: the examples run with Chapter 32's driver ('%mgc32').
ch33 = os.path.join(config.test_source_root, "..", "..", "..", "part33", "code")
config.substitutions += [
("%ex33", os.path.join(ch33, "examples")),
("%ch33", ch33),
]
# Chapter 32: held-out data, weight decay and greedy generation ('%mgc32' etc. must come before the shorter '%mgc').
ch32 = os.path.join(config.test_source_root, "..", "..", "..", "part32", "code")
config.substitutions += [
("%mgc32", "env MG_OPT=" + mg_opt + " " + os.path.join(ch32, "mgc")),
("%ex32", os.path.join(ch32, "examples")),
("%ch32", ch32),
]
# Chapter 31: training a bigram model ('%mgc31' etc. must come before the shorter '%mgc').
ch31 = os.path.join(config.test_source_root, "..", "..", "..", "part31", "code")
config.substitutions += [
("%mgc31", "env MG_OPT=" + mg_opt + " " + os.path.join(ch31, "mgc")),
("%ex31", os.path.join(ch31, "examples")),
("%ch31", ch31),
]
# Chapter 30: a small transformer ('%mgc30' etc. must come before the shorter '%mgc').
ch30 = os.path.join(config.test_source_root, "..", "..", "..", "part30", "code")
config.substitutions += [
("%mgc30", "env MG_OPT=" + mg_opt + " " + os.path.join(ch30, "mgc")),
("%ex30", os.path.join(ch30, "examples")),
("%ch30", ch30),
]
# Chapter 29: exp, softmax and attention ('%mgc29' etc. must come before the shorter '%mgc').
ch29 = os.path.join(config.test_source_root, "..", "..", "..", "part29", "code")
config.substitutions += [
("%mgc29", "env MG_OPT=" + mg_opt + " " + os.path.join(ch29, "mgc")),
("%ex29", os.path.join(ch29, "examples")),
("%ch29", ch29),
]
# Chapter 28: tensors of any rank and contract(...) in Mountain Goat itself ('%mgc28' etc. must come before the shorter '%mgc').
ch28 = os.path.join(config.test_source_root, "..", "..", "..", "part28", "code")
config.substitutions += [
("%mgc28", "env MG_OPT=" + mg_opt + " " + os.path.join(ch28, "mgc")),
("%ex28", os.path.join(ch28, "examples")),
("%ch28", ch28),
]
# Chapter 27: tensor contractions through Mountain Goat's matrix product (files used in place).
config.substitutions += [("%cpp27", os.path.join(config.test_source_root, "..", "..", "..", "part27", "code", "cpp"))]
# Chapter 26: the gradient-descent trainer (its Mountain Goat file and C++ driver are used in place).
config.substitutions += [("%cpp26", os.path.join(config.test_source_root, "..", "..", "..", "part26", "code", "cpp"))]
# Chapter 25: tests for the early chapters' own files, run in place ('%docs' is the book's docs directory).
config.substitutions += [("%docs", os.path.join(config.test_source_root, "..", "..", ".."))]
# Chapter 24: the matmul loop-order option ('%mgc24' must come before '%mgc').
ch24 = os.path.join(config.test_source_root, "..", "..", "..", "part24", "code")
config.substitutions += [
("%mgc24", "env MG_OPT=" + mg_opt + " " + os.path.join(ch24, "mgc")),
("%ex24", os.path.join(ch24, "examples")),
("%cpp21b", os.path.join(config.test_source_root, "..", "..", "..", "part21", "code", "cpp")),
]
# Chapter 23: mgc's -O and --passes options, and the benchmark examples ('%mgc23' must come before '%mgc').
ch23 = os.path.join(config.test_source_root, "..", "..", "..", "part23", "code")
config.substitutions += [
("%mgc23", "env MG_OPT=" + mg_opt + " " + os.path.join(ch23, "mgc")),
("%ex23", os.path.join(ch23, "examples")),
]
# Chapter 22: broadcasting, reductions, relu. (Listed before Chapter 21's: lit replaces by prefix, so '%mgc22' must come before '%mgc'.)
ch22 = os.path.join(config.test_source_root, "..", "..", "..", "part22", "code")
config.substitutions += [
("%mgc22", "env MG_OPT=" + mg_opt + " " + os.path.join(ch22, "mgc")),
("%mgfront22", "python3 " + os.path.join(ch22, "mgfront.py")),
("%ex22", os.path.join(ch22, "examples")),
("%cpp22", os.path.join(ch22, "cpp")),
]
# Chapter 21: the extended driver/front end (more operations), also referenced in place.
ch21 = os.path.join(config.test_source_root, "..", "..", "..", "part21", "code")
config.substitutions += [
("%mgc21", "env MG_OPT=" + mg_opt + " " + os.path.join(ch21, "mgc")),
("%mgfront21", "python3 " + os.path.join(ch21, "mgfront.py")),
("%ex21", os.path.join(ch21, "examples")),
("%cpp21", os.path.join(ch21, "cpp")),
]
# Chapter 20: the surface-syntax driver and its examples (they live in part20/code, not copied, so tests cannot drift from the book).
ch20 = os.path.join(config.test_source_root, "..", "..", "..", "part20", "code")
config.substitutions += [
("%mgc", "env MG_OPT=" + mg_opt + " " + os.path.join(ch20, "mgc")),
("%mgfront", "python3 " + os.path.join(ch20, "mgfront.py")),
("%ex", os.path.join(ch20, "examples")),
("%cpp", os.path.join(ch20, "cpp")),
("%cxx", "clang++-18"),
]
config.environment["PATH"] = os.environ["PATH"]
# Chapter 41: a back end from the mg dialect to GA-1 (Python over `mgc mlir` text).
ch41 = os.path.join(config.test_source_root, "..", "..", "..", "part41", "code")
config.substitutions += [("%ch41", ch41)]
# Chapter 42: lowering scf to cf last to first (--mg-scf-to-cf-reverse, mgc --reverse-loops); mg-opt of this chapter is the newest build.
ch42 = os.path.join(config.test_source_root, "..", "..", "..", "part42", "code")
config.substitutions += [("%ch42", ch42)]
# Chapter 43: mgc's defaults changed (loops lowered last to first, big functions outlined); the driver is part43/code/mgc, its mg-opt is Chapter 42's.
ch43 = os.path.join(config.test_source_root, "..", "..", "..", "part43", "code")
config.substitutions += [("%ch43", ch43)]
# Chapter 44: fast-math flags (--mg-set-fastmath, mgc --fast-math); the newest mg-opt build is this chapter's.
ch44 = os.path.join(config.test_source_root, "..", "..", "..", "part44", "code")
config.substitutions += [("%ch44", ch44)]
# Chapter 45: automatic differentiation (autograd.py, a source-to-source transformation; it runs programs with mgc and needs mg-opt of Chapter 44's build).
ch45 = os.path.join(config.test_source_root, "..", "..", "..", "part45", "code")
config.substitutions += [("%ch45", ch45)]
# Chapter 46: Hessian-vector products (hvp.py differentiates the generated backward pass again).
ch46 = os.path.join(config.test_source_root, "..", "..", "..", "part46", "code")
config.substitutions += [("%ch46", ch46)]
# Appendices: the language reference (A) is checked against the compiler, the self-check answers (D) against the chapter pages.
apx = os.path.join(config.test_source_root, "..", "..", "..")
config.substitutions += [("%apxA", os.path.join(apx, "appendixA", "code")), ("%apxD", os.path.join(apx, "appendixD"))]
The one design decision worth noticing is MG_OPT: which mg-opt binary is tested is an environment variable, not a constant. That is what lets the same suite be pointed at older builds later in this chapter to show the tests can fail. Each test gets its own scratch name (%t), so tests cannot read each other's files.
The launcher script is three lines of real logic:
#!/bin/sh
# Run the regression suite against an mg-opt build. Default: the newest build that exists, Chapter 44, then 42, 38, 32, 31, 30, 29, 28, 24, 22, 21, 20, 19, 14 (run ../../part44/code/build.sh etc. first).
# Override with MG_OPT=/path/to/mg-opt. Needs: pip install lit; FileCheck-18, not-18, mlir-translate-18, clang-18.
HERE=$(cd "$(dirname "$0")" && pwd)
export MG_OPT=${MG_OPT:-$(for b in part44 part42 part38 part32 part31 part30 part29 part28 part24 part22 part21 part20 part19 part14; do [ -x $HERE/../../$b/code/build/mg-opt ] && echo $HERE/../../$b/code/build/mg-opt && break; done)}
[ -x "$MG_OPT" ] || { echo "no mg-opt at $MG_OPT (build it first)"; exit 2; }
export MG_TEST_TMP=${MG_TEST_TMP:-$HERE/work/lit-out}
mkdir -p "$MG_TEST_TMP"
exec lit "$@" "$HERE/test"
The tools, and one deviation¶
This sandbox has FileCheck-18 and not-18 from the llvm-18-tools package, but no lit for LLVM 18: apt-get install python3-lit and apt-get install lit both report "Unable to locate package". lit was installed from PyPI instead: version 23.1.2, which is not the toolchain's version (18.1.3). lit is a standalone Python test runner that only launches processes, so the mismatch is unlikely to matter, but it is a deviation, stated rather than hidden.
A FileCheck primer¶
Each CHECK pattern must be found in the piped output 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 cannot match another's |
{{regex}} |
a regular expression inside a pattern |
%[[NAME:regex]], %[[NAME]] |
capture text once, require the same text later |
SSA value names (%0, %dim_2) differ between compiler versions, so the tests below match them with {{.*}} and pin only what matters.
The suite: 22 tests in five groups¶
| Group | Tests | What each one pins down |
|---|---|---|
verifier/ |
5 | static shape mismatch rejected; bad transpose shape rejected; rank 3 rejected; unranked rejected with an explicit message; dynamic and mixed shapes accepted |
canonicalize/ |
4 | two constants fold; transpose-of-transpose folds; folds with dynamic extents when types match exactly; is left alone when types differ |
lowering/ |
6 | static add lowers with no runtime check; dynamic add gets extents, a check per dimension and a dynamic allocation; a mixed add gets exactly one check; transpose gets none; the One-Shot Bufferize path matches |
loops/ |
1 | --affine-loop-fusion merges Chapter 7's two loop nests |
execution/ |
6 | compile to native code, run, check output and exit status |
Real run (code/run_out.txt, recorded when the suite had 22 tests; Chapter 16 grows it to 34):
-- Testing: 22 tests, 4 workers --
PASS: IR Forge / Mountain Goat :: canonicalize/fold-constant-add.mlir (1 of 22)
PASS: IR Forge / Mountain Goat :: canonicalize/transpose-twice-dynamic-mixed.mlir (2 of 22)
PASS: IR Forge / Mountain Goat :: canonicalize/transpose-twice.mlir (3 of 22)
PASS: IR Forge / Mountain Goat :: canonicalize/transpose-twice-dynamic-exact.mlir (4 of 22)
PASS: IR Forge / Mountain Goat :: execution/mixed-static-dynamic.mlir (5 of 22)
PASS: IR Forge / Mountain Goat :: execution/dynamic-shapes-run.mlir (6 of 22)
PASS: IR Forge / Mountain Goat :: execution/mismatch-columns-aborts.mlir (7 of 22)
PASS: IR Forge / Mountain Goat :: execution/mismatch-aborts.mlir (8 of 22)
PASS: IR Forge / Mountain Goat :: loops/fusion.mlir (9 of 22)
PASS: IR Forge / Mountain Goat :: lowering/dynamic-add-affine.mlir (10 of 22)
PASS: IR Forge / Mountain Goat :: verifier/transpose-bad-shape.mlir (11 of 22)
PASS: IR Forge / Mountain Goat :: verifier/dynamic-accepted.mlir (12 of 22)
PASS: IR Forge / Mountain Goat :: lowering/static-add-affine.mlir (13 of 22)
PASS: IR Forge / Mountain Goat :: verifier/transpose-rank3.mlir (14 of 22)
PASS: IR Forge / Mountain Goat :: lowering/bufferize-dynamic-add.mlir (15 of 22)
PASS: IR Forge / Mountain Goat :: lowering/mixed-add-one-check.mlir (16 of 22)
PASS: IR Forge / Mountain Goat :: lowering/dynamic-transpose-no-check.mlir (17 of 22)
PASS: IR Forge / Mountain Goat :: verifier/unranked-rejected.mlir (18 of 22)
PASS: IR Forge / Mountain Goat :: lowering/bufferize-static-add.mlir (19 of 22)
PASS: IR Forge / Mountain Goat :: verifier/add-static-mismatch.mlir (20 of 22)
PASS: IR Forge / Mountain Goat :: execution/static-add-runs.mlir (21 of 22)
PASS: IR Forge / Mountain Goat :: execution/strided-view.mlir (22 of 22)
Testing Time: 0.42s
Total Discovered Tests: 22
Passed: 22 (100.00%)
Group 1: the verifier¶
A verifier is the dialect's own check that an operation is well formed (Chapter 3). Each test feeds a deliberately bad program to mg-opt and checks the diagnostic. The %not in front means "this command must fail", and 2>&1 sends the error text into FileCheck.
// A static shape mismatch must be rejected by the verifier, not at runtime.
// RUN: %not %mg-opt %s 2>&1 | %FileCheck %s
// CHECK: error: 'mg.add' op mg.add operands must have compatible shapes, got 'tensor<2x2xf64>' and 'tensor<2x3xf64>'
func.func @f(%a: tensor<2x2xf64>, %b: tensor<2x3xf64>) -> tensor<2x2xf64> {
%0 = mg.add %a, %b : tensor<2x2xf64>, tensor<2x3xf64> -> tensor<2x2xf64>
func.return %0 : tensor<2x2xf64>
}
The real output the pattern matches (show/verify_add-static-mismatch.txt):
verifier/add-static-mismatch.mlir:5:8: error: 'mg.add' op mg.add operands must have compatible shapes, got 'tensor<2x2xf64>' and 'tensor<2x3xf64>' -> 'tensor<2x2xf64>'
%0 = mg.add %a, %b : tensor<2x2xf64>, tensor<2x3xf64> -> tensor<2x2xf64>
^
verifier/add-static-mismatch.mlir:5:8: note: see current operation: %0 = "mg.add"(%arg0, %arg1) : (tensor<2x2xf64>, tensor<2x3xf64>) -> tensor<2x2xf64>
A static mismatch (2x2 plus 2x3) must die at compile time. The pattern pins the full message including both types, which makes the test strict about wording (see the end of this chapter). The other three rejection tests follow the same shape:
// RUN: %not %mg-opt %s 2>&1 | %FileCheck %s
// CHECK: error: 'mg.transpose' op mg.transpose result shape must be the input shape reversed
func.func @f(%a: tensor<2x3xf64>) -> tensor<2x3xf64> {
%0 = mg.transpose %a : tensor<2x3xf64> to tensor<2x3xf64>
func.return %0 : tensor<2x3xf64>
}
verifier/transpose-bad-shape.mlir:4:8: error: 'mg.transpose' op mg.transpose result shape must be the input shape reversed
%0 = mg.transpose %a : tensor<2x3xf64> to tensor<2x3xf64>
^
verifier/transpose-bad-shape.mlir:4:8: note: see current operation: %0 = "mg.transpose"(%arg0) : (tensor<2x3xf64>) -> tensor<2x3xf64>
// RUN: %not %mg-opt %s 2>&1 | %FileCheck %s
// CHECK: error: 'mg.transpose' op mg.transpose only supports rank-2 tensors
func.func @f(%a: tensor<2x2x2xf64>) -> tensor<2x2x2xf64> {
%0 = mg.transpose %a : tensor<2x2x2xf64> to tensor<2x2x2xf64>
func.return %0 : tensor<2x2x2xf64>
}
verifier/transpose-rank3.mlir:4:8: error: 'mg.transpose' op mg.transpose only supports rank-2 tensors
%0 = mg.transpose %a : tensor<2x2x2xf64> to tensor<2x2x2xf64>
^
verifier/transpose-rank3.mlir:4:8: note: see current operation: %0 = "mg.transpose"(%arg0) : (tensor<2x2x2xf64>) -> tensor<2x2x2xf64>
// Chapter 13: an explicit diagnostic instead of an unchecked cast.
// RUN: %not %mg-opt %s 2>&1 | %FileCheck %s
// CHECK: error: 'mg.transpose' op mg.transpose only supports ranked tensors
func.func @f(%a: tensor<*xf64>) -> tensor<*xf64> {
%0 = mg.transpose %a : tensor<*xf64> to tensor<*xf64>
func.return %0 : tensor<*xf64>
}
verifier/unranked-rejected.mlir:5:8: error: 'mg.transpose' op mg.transpose only supports ranked tensors
%0 = mg.transpose %a : tensor<*xf64> to tensor<*xf64>
^
verifier/unranked-rejected.mlir:5:8: note: see current operation: %0 = "mg.transpose"(%arg0) : (tensor<*xf64>) -> tensor<*xf64>
The last one matters for a reason Chapter 13 found: the original verifier cast an unranked tensor to a ranked type without checking, which in a release build happened to print a plausible message by luck. This test pins the explicit diagnostic that replaced it. The fifth test is the opposite direction: programs that must be accepted.
// Chapter 13: dynamic extents verify, including a mixed ?x2 + 3x2 add.
// RUN: %mg-opt %s | %FileCheck %s
// CHECK-LABEL: func.func @dyn
// CHECK: mg.add %arg0, %arg1 : tensor<?x?xf64>, tensor<?x?xf64> -> tensor<?x?xf64>
// CHECK-LABEL: func.func @mixed
// CHECK: mg.add %arg0, %arg1 : tensor<?x2xf64>, tensor<3x2xf64> -> tensor<?x2xf64>
// CHECK-LABEL: func.func @t
// CHECK: mg.transpose %arg0 : tensor<?x?xf64> to tensor<?x?xf64>
func.func @dyn(%a: tensor<?x?xf64>, %b: tensor<?x?xf64>) -> tensor<?x?xf64> {
%0 = mg.add %a, %b : tensor<?x?xf64>, tensor<?x?xf64> -> tensor<?x?xf64>
func.return %0 : tensor<?x?xf64>
}
func.func @mixed(%a: tensor<?x2xf64>, %b: tensor<3x2xf64>) -> tensor<?x2xf64> {
%0 = mg.add %a, %b : tensor<?x2xf64>, tensor<3x2xf64> -> tensor<?x2xf64>
func.return %0 : tensor<?x2xf64>
}
func.func @t(%a: tensor<?x?xf64>) -> tensor<?x?xf64> {
%0 = mg.transpose %a : tensor<?x?xf64> to tensor<?x?xf64>
func.return %0 : tensor<?x?xf64>
}
with CHECK-LABEL keeping each function's checks to its own function. The real output, printed back unchanged:
module {
func.func @dyn(%arg0: tensor<?x?xf64>, %arg1: tensor<?x?xf64>) -> tensor<?x?xf64> {
%0 = mg.add %arg0, %arg1 : tensor<?x?xf64>, tensor<?x?xf64> -> tensor<?x?xf64>
return %0 : tensor<?x?xf64>
}
func.func @mixed(%arg0: tensor<?x2xf64>, %arg1: tensor<3x2xf64>) -> tensor<?x2xf64> {
%0 = mg.add %arg0, %arg1 : tensor<?x2xf64>, tensor<3x2xf64> -> tensor<?x2xf64>
return %0 : tensor<?x2xf64>
}
func.func @t(%arg0: tensor<?x?xf64>) -> tensor<?x?xf64> {
%0 = mg.transpose %arg0 : tensor<?x?xf64> to tensor<?x?xf64>
return %0 : tensor<?x?xf64>
}
}
Group 2: the canonicalizer¶
Canonicalization (Chapter 3) rewrites a program into a simpler equivalent one. These tests pin what it must do and, equally, what it must not do.
// Chapter 3: AddOp::fold folds two constants into one.
// RUN: %mg-opt %s --canonicalize | %FileCheck %s
// CHECK: mg.constant dense<{{\[\[}}6.000000e+00, 8.000000e+00], [1.000000e+01, 1.200000e+01]]> : tensor<2x2xf64>
// CHECK-NOT: mg.add
func.func @main() {
%a = mg.constant dense<[[1.0, 2.0], [3.0, 4.0]]> : tensor<2x2xf64>
%b = mg.constant dense<[[5.0, 6.0], [7.0, 8.0]]> : tensor<2x2xf64>
%s = mg.add %a, %b : tensor<2x2xf64>, tensor<2x2xf64> -> tensor<2x2xf64>
mg.print %s : tensor<2x2xf64>
func.return
}
module {
func.func @main() {
%0 = mg.constant dense<[[6.000000e+00, 8.000000e+00], [1.000000e+01, 1.200000e+01]]> : tensor<2x2xf64>
mg.print %0 : tensor<2x2xf64>
return
}
}
The two constants and the add are gone, replaced by one constant holding the sum (CHECK-NOT: mg.add asserts the add did not survive). The regular expression {{\[\[}} is just an escaped literal [[. The transpose pair:
// Chapter 3: transpose(transpose(x)) -> x.
// RUN: %mg-opt %s --canonicalize | %FileCheck %s
// CHECK: func.func @tt
// CHECK-NOT: mg.transpose
// CHECK: return %arg0 : tensor<2x3xf64>
func.func @tt(%a: tensor<2x3xf64>) -> tensor<2x3xf64> {
%0 = mg.transpose %a : tensor<2x3xf64> to tensor<3x2xf64>
%1 = mg.transpose %0 : tensor<3x2xf64> to tensor<2x3xf64>
func.return %1 : tensor<2x3xf64>
}
module {
func.func @tt(%arg0: tensor<2x3xf64>) -> tensor<2x3xf64> {
return %arg0 : tensor<2x3xf64>
}
}
The next two are the pair that guard Chapter 13's bug. With dynamic extents the rewrite still fires when its replacement has exactly the right type:
// Chapter 13: with dynamic extents the rewrite still fires when the types match exactly.
// RUN: %mg-opt %s --canonicalize | %FileCheck %s
// CHECK-NOT: mg.transpose
// CHECK: return %arg0 : tensor<?x2xf64>
func.func @tt(%a: tensor<?x2xf64>) -> tensor<?x2xf64> {
%0 = mg.transpose %a : tensor<?x2xf64> to tensor<2x?xf64>
%1 = mg.transpose %0 : tensor<2x?xf64> to tensor<?x2xf64>
func.return %1 : tensor<?x2xf64>
}
module {
func.func @tt(%arg0: tensor<?x2xf64>) -> tensor<?x2xf64> {
return %arg0 : tensor<?x2xf64>
}
}
but it must be left alone when it would change a type (it once produced invalid IR here):
// Chapter 13: the rewrite must NOT fire when it would change the type (it used to emit invalid IR).
// RUN: %mg-opt %s --canonicalize | %FileCheck %s
// CHECK: mg.transpose %arg0 : tensor<?x2xf64> to tensor<2x?xf64>
// CHECK: mg.transpose %0 : tensor<2x?xf64> to tensor<?x?xf64>
// CHECK: return %1 : tensor<?x?xf64>
func.func @tt(%a: tensor<?x2xf64>) -> tensor<?x?xf64> {
%0 = mg.transpose %a : tensor<?x2xf64> to tensor<2x?xf64>
%1 = mg.transpose %0 : tensor<2x?xf64> to tensor<?x?xf64>
func.return %1 : tensor<?x?xf64>
}
module {
func.func @tt(%arg0: tensor<?x2xf64>) -> tensor<?x?xf64> {
%0 = mg.transpose %arg0 : tensor<?x2xf64> to tensor<2x?xf64>
%1 = mg.transpose %0 : tensor<2x?xf64> to tensor<?x?xf64>
return %1 : tensor<?x?xf64>
}
}
Group 3: the lowerings¶
These check the shape of the IR the two lowering paths produce. The first is the baseline: a static add must lower to constant-bound loops with no runtime machinery.
// Chapter 4/5: static mg.add -> affine loop nest with constant bounds, and NO runtime checks.
// RUN: %mg-opt %s --convert-mg-to-affine | %FileCheck %s
// CHECK: memref.alloc() : memref<2x2xf64>
// CHECK: affine.for %{{.*}} = 0 to 2 {
// CHECK: affine.for %{{.*}} = 0 to 2 {
// CHECK: affine.load
// CHECK: affine.load
// CHECK: arith.addf
// CHECK: affine.store
// CHECK-NOT: cf.assert
// CHECK-NOT: memref.dim
func.func @add_tensors(%a: tensor<2x2xf64>, %b: tensor<2x2xf64>) -> tensor<2x2xf64> {
%0 = mg.add %a, %b : tensor<2x2xf64>, tensor<2x2xf64> -> tensor<2x2xf64>
func.return %0 : tensor<2x2xf64>
}
module {
func.func @add_tensors(%arg0: memref<2x2xf64>, %arg1: memref<2x2xf64>) -> memref<2x2xf64> {
%alloc = memref.alloc() : memref<2x2xf64>
affine.for %arg2 = 0 to 2 {
affine.for %arg3 = 0 to 2 {
%0 = affine.load %arg0[%arg2, %arg3] : memref<2x2xf64>
%1 = affine.load %arg1[%arg2, %arg3] : memref<2x2xf64>
%2 = arith.addf %0, %1 : f64
affine.store %2, %alloc[%arg2, %arg3] : memref<2x2xf64>
}
}
return %alloc : memref<2x2xf64>
}
}
The two CHECK-NOT lines are the point: the dynamic-shape work of Chapters 13 and 14 must cost static programs nothing. The dynamic add, where each dimension's size is read at runtime and compared before the loop:
// Chapters 13/14: dynamic add -> memref.dim extents, a runtime check per dimension, dynamic alloc.
// RUN: %mg-opt %s --convert-mg-to-affine | %FileCheck %s
// CHECK-LABEL: func.func @add_dyn
// CHECK: arith.cmpi eq
// CHECK-NEXT: cf.assert {{.*}}, "mg.add: operand shapes differ at runtime in dimension 0"
// CHECK: arith.cmpi eq
// CHECK-NEXT: cf.assert {{.*}}, "mg.add: operand shapes differ at runtime in dimension 1"
// CHECK: memref.alloc(%{{.*}}, %{{.*}}) : memref<?x?xf64>
// CHECK: affine.for
// CHECK: affine.for
func.func @add_dyn(%a: tensor<?x?xf64>, %b: tensor<?x?xf64>) -> tensor<?x?xf64> {
%0 = mg.add %a, %b : tensor<?x?xf64>, tensor<?x?xf64> -> tensor<?x?xf64>
func.return %0 : tensor<?x?xf64>
}
#map = affine_map<(d0) -> (d0)>
module {
func.func @add_dyn(%arg0: memref<?x?xf64>, %arg1: memref<?x?xf64>) -> memref<?x?xf64> {
%c0 = arith.constant 0 : index
%dim = memref.dim %arg0, %c0 : memref<?x?xf64>
%c0_0 = arith.constant 0 : index
%dim_1 = memref.dim %arg1, %c0_0 : memref<?x?xf64>
%0 = arith.cmpi eq, %dim, %dim_1 : index
cf.assert %0, "mg.add: operand shapes differ at runtime in dimension 0"
%c1 = arith.constant 1 : index
%dim_2 = memref.dim %arg0, %c1 : memref<?x?xf64>
%c1_3 = arith.constant 1 : index
%dim_4 = memref.dim %arg1, %c1_3 : memref<?x?xf64>
%1 = arith.cmpi eq, %dim_2, %dim_4 : index
cf.assert %1, "mg.add: operand shapes differ at runtime in dimension 1"
%c0_5 = arith.constant 0 : index
%dim_6 = memref.dim %arg0, %c0_5 : memref<?x?xf64>
%c1_7 = arith.constant 1 : index
%dim_8 = memref.dim %arg0, %c1_7 : memref<?x?xf64>
%alloc = memref.alloc(%dim_6, %dim_8) : memref<?x?xf64>
%c0_9 = arith.constant 0 : index
affine.for %arg2 = #map(%c0_9) to #map(%dim_6) {
affine.for %arg3 = #map(%c0_9) to #map(%dim_8) {
%2 = affine.load %arg0[%arg2, %arg3] : memref<?x?xf64>
%3 = affine.load %arg1[%arg2, %arg3] : memref<?x?xf64>
%4 = arith.addf %2, %3 : f64
affine.store %4, %alloc[%arg2, %arg3] : memref<?x?xf64>
}
}
return %alloc : memref<?x?xf64>
}
}
CHECK-NEXT asserts each cf.assert directly follows its comparison. The mixed case, where only the dimension that is dynamic is checked:
// Chapter 14: ?x2 + 3x2 checks only dimension 0 (dimension 1 is static on both sides).
// RUN: %mg-opt %s --convert-mg-to-affine | %FileCheck %s
// CHECK: cf.assert {{.*}}, "mg.add: operand shapes differ at runtime in dimension 0"
// CHECK-NOT: cf.assert
func.func @mixed(%a: tensor<?x2xf64>, %b: tensor<3x2xf64>) -> tensor<?x2xf64> {
%0 = mg.add %a, %b : tensor<?x2xf64>, tensor<3x2xf64> -> tensor<?x2xf64>
func.return %0 : tensor<?x2xf64>
}
#map = affine_map<(d0) -> (d0)>
module {
func.func @mixed(%arg0: memref<?x2xf64>, %arg1: memref<3x2xf64>) -> memref<?x2xf64> {
%c0 = arith.constant 0 : index
%dim = memref.dim %arg0, %c0 : memref<?x2xf64>
%c0_0 = arith.constant 0 : index
%dim_1 = memref.dim %arg1, %c0_0 : memref<3x2xf64>
%0 = arith.cmpi eq, %dim, %dim_1 : index
cf.assert %0, "mg.add: operand shapes differ at runtime in dimension 0"
%c0_2 = arith.constant 0 : index
%dim_3 = memref.dim %arg0, %c0_2 : memref<?x2xf64>
%alloc = memref.alloc(%dim_3) : memref<?x2xf64>
%c0_4 = arith.constant 0 : index
%c2 = arith.constant 2 : index
affine.for %arg2 = #map(%c0_4) to #map(%dim_3) {
affine.for %arg3 = 0 to 2 {
%1 = affine.load %arg0[%arg2, %arg3] : memref<?x2xf64>
%2 = affine.load %arg1[%arg2, %arg3] : memref<3x2xf64>
%3 = arith.addf %1, %2 : f64
affine.store %3, %alloc[%arg2, %arg3] : memref<?x2xf64>
}
}
return %alloc : memref<?x2xf64>
}
}
and the transpose, which has one operand and therefore nothing to compare, with the result's extents swapped (%c1 then %c0):
// Chapter 13/14: transpose has one operand, so no shape check; the result extents are swapped dims.
// RUN: %mg-opt %s --convert-mg-to-affine | %FileCheck %s
// CHECK-NOT: cf.assert
// CHECK: memref.dim %arg0, %c1
// CHECK: memref.dim %arg0, %c0
// CHECK: memref.alloc(%{{.*}}, %{{.*}}) : memref<?x?xf64>
// CHECK: affine.load %arg0[%arg{{[0-9]+}}, %arg{{[0-9]+}}]
func.func @t(%a: tensor<?x?xf64>) -> tensor<?x?xf64> {
%0 = mg.transpose %a : tensor<?x?xf64> to tensor<?x?xf64>
func.return %0 : tensor<?x?xf64>
}
#map = affine_map<(d0) -> (d0)>
module {
func.func @t(%arg0: memref<?x?xf64>) -> memref<?x?xf64> {
%c1 = arith.constant 1 : index
%dim = memref.dim %arg0, %c1 : memref<?x?xf64>
%c0 = arith.constant 0 : index
%dim_0 = memref.dim %arg0, %c0 : memref<?x?xf64>
%alloc = memref.alloc(%dim, %dim_0) : memref<?x?xf64>
%c0_1 = arith.constant 0 : index
affine.for %arg1 = #map(%c0_1) to #map(%dim) {
affine.for %arg2 = #map(%c0_1) to #map(%dim_0) {
%0 = affine.load %arg0[%arg2, %arg1] : memref<?x?xf64>
affine.store %0, %alloc[%arg1, %arg2] : memref<?x?xf64>
}
}
return %alloc : memref<?x?xf64>
}
}
The last two cover Chapter 6's other path, One-Shot Bufferize, which must produce the same structure. Note its function arguments are strided memref types (Chapter 6's finding), and the test pins that:
// Chapter 6 + 13 + 14: the One-Shot Bufferize path emits the same dynamic structure and the same checks.
// RUN: %mg-opt %s --one-shot-bufferize="bufferize-function-boundaries" | %FileCheck %s
// CHECK: func.func @add_dyn(%arg0: memref<?x?xf64, strided<[?, ?], offset: ?>>, %arg1: memref<?x?xf64, strided<[?, ?], offset: ?>>) -> memref<?x?xf64>
// CHECK: cf.assert {{.*}}, "mg.add: operand shapes differ at runtime in dimension 0"
// CHECK: cf.assert {{.*}}, "mg.add: operand shapes differ at runtime in dimension 1"
// CHECK: memref.alloc(%{{.*}}, %{{.*}}) : memref<?x?xf64>
func.func @add_dyn(%a: tensor<?x?xf64>, %b: tensor<?x?xf64>) -> tensor<?x?xf64> {
%0 = mg.add %a, %b : tensor<?x?xf64>, tensor<?x?xf64> -> tensor<?x?xf64>
func.return %0 : tensor<?x?xf64>
}
#map = affine_map<(d0) -> (d0)>
module {
func.func @add_dyn(%arg0: memref<?x?xf64, strided<[?, ?], offset: ?>>, %arg1: memref<?x?xf64, strided<[?, ?], offset: ?>>) -> memref<?x?xf64> {
%c0 = arith.constant 0 : index
%dim = memref.dim %arg0, %c0 : memref<?x?xf64, strided<[?, ?], offset: ?>>
%c0_0 = arith.constant 0 : index
%dim_1 = memref.dim %arg1, %c0_0 : memref<?x?xf64, strided<[?, ?], offset: ?>>
%0 = arith.cmpi eq, %dim, %dim_1 : index
cf.assert %0, "mg.add: operand shapes differ at runtime in dimension 0"
%c1 = arith.constant 1 : index
%dim_2 = memref.dim %arg0, %c1 : memref<?x?xf64, strided<[?, ?], offset: ?>>
%c1_3 = arith.constant 1 : index
%dim_4 = memref.dim %arg1, %c1_3 : memref<?x?xf64, strided<[?, ?], offset: ?>>
%1 = arith.cmpi eq, %dim_2, %dim_4 : index
cf.assert %1, "mg.add: operand shapes differ at runtime in dimension 1"
%c0_5 = arith.constant 0 : index
%dim_6 = memref.dim %arg0, %c0_5 : memref<?x?xf64, strided<[?, ?], offset: ?>>
%c1_7 = arith.constant 1 : index
%dim_8 = memref.dim %arg0, %c1_7 : memref<?x?xf64, strided<[?, ?], offset: ?>>
%alloc = memref.alloc(%dim_6, %dim_8) : memref<?x?xf64>
%c0_9 = arith.constant 0 : index
affine.for %arg2 = #map(%c0_9) to #map(%dim_6) {
affine.for %arg3 = #map(%c0_9) to #map(%dim_8) {
%2 = affine.load %arg0[%arg2, %arg3] : memref<?x?xf64, strided<[?, ?], offset: ?>>
%3 = affine.load %arg1[%arg2, %arg3] : memref<?x?xf64, strided<[?, ?], offset: ?>>
%4 = arith.addf %2, %3 : f64
affine.store %4, %alloc[%arg2, %arg3] : memref<?x?xf64>
}
}
%cast = memref.cast %alloc : memref<?x?xf64> to memref<?x?xf64, strided<[?, ?], offset: ?>>
return %alloc : memref<?x?xf64>
}
}
// Chapter 6: static One-Shot Bufferize path: no dynamic machinery, no runtime check.
// RUN: %mg-opt %s --one-shot-bufferize="bufferize-function-boundaries" | %FileCheck %s
// CHECK: memref.alloc() : memref<2x2xf64>
// CHECK: affine.for %{{.*}} = 0 to 2
// CHECK-NOT: cf.assert
// CHECK-NOT: memref.dim
func.func @add_tensors(%a: tensor<2x2xf64>, %b: tensor<2x2xf64>) -> tensor<2x2xf64> {
%0 = mg.add %a, %b : tensor<2x2xf64>, tensor<2x2xf64> -> tensor<2x2xf64>
func.return %0 : tensor<2x2xf64>
}
module {
func.func @add_tensors(%arg0: memref<2x2xf64, strided<[?, ?], offset: ?>>, %arg1: memref<2x2xf64, strided<[?, ?], offset: ?>>) -> memref<2x2xf64> {
%alloc = memref.alloc() : memref<2x2xf64>
affine.for %arg2 = 0 to 2 {
affine.for %arg3 = 0 to 2 {
%0 = affine.load %arg0[%arg2, %arg3] : memref<2x2xf64, strided<[?, ?], offset: ?>>
%1 = affine.load %arg1[%arg2, %arg3] : memref<2x2xf64, strided<[?, ?], offset: ?>>
%2 = arith.addf %0, %1 : f64
affine.store %2, %alloc[%arg2, %arg3] : memref<2x2xf64>
}
}
%cast = memref.cast %alloc : memref<2x2xf64> to memref<2x2xf64, strided<[?, ?], offset: ?>>
return %alloc : memref<2x2xf64>
}
}
Group 4: loop transforms¶
One test, for Chapter 7's fusion (Chapter 16 adds tiling and unrolling). The header of the test, then the real output of --affine-loop-fusion on Chapter 7's program:
// Chapter 7: --affine-loop-fusion merges the add nest and the transpose nest (4 loops -> 2).
// RUN: %mg-opt %s --affine-loop-fusion | %FileCheck %s
// CHECK: func.func @compute
// CHECK: affine.for
// CHECK-NEXT: affine.for
// CHECK-NOT: affine.for
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
}
}
Four loops (two nests of two) became one nest of two: CHECK/CHECK-NEXT pin the nested pair and the trailing CHECK-NOT: affine.for pins that no further loop remains.
Group 5: running the compiled program¶
The strongest tests compile to a native executable and run it. This is a pipeline of RUN: lines, each one a step Chapters 4 through 8 introduced: lower with mg-opt, translate to LLVM IR, compile, link a C harness, run, and pipe the program's output into FileCheck. The static case, on both lowering paths:
// Chapter 8: static add, both paths, compiled to a native executable and run.
// RUN: %mg-opt %inputs/static_add.mlir --convert-mg-to-affine %lower-to-llvm -o %t.a.mlir
// RUN: %mlir-translate --mlir-to-llvmir %t.a.mlir -o %t.a.ll
// RUN: %clang -c %t.a.ll -o %t.a.o
// RUN: %clang %inputs/static_harness.c %t.a.o -o %t.a.exe
// RUN: %t.a.exe | %FileCheck %s
// RUN: %mg-opt %inputs/static_add.mlir --one-shot-bufferize="bufferize-function-boundaries" -o %t.b0.mlir
// RUN: %mg-opt %t.b0.mlir %lower-to-llvm -o %t.b.mlir
// RUN: %mlir-translate --mlir-to-llvmir %t.b.mlir -o %t.b.ll
// RUN: %clang -c %t.b.ll -o %t.b.o
// RUN: %clang %inputs/static_harness.c %t.b.o -o %t.b.exe
// RUN: %t.b.exe | %FileCheck %s
// CHECK: 6 8
// CHECK-NEXT: 10 12
// (input program lives in Inputs/static_add.mlir)
The two real outputs (show/exec_static-path-A.txt and -B.txt):
The dynamic case: one compiled function, five different non-square shapes, both paths.
// Chapters 13/14: ONE compiled function handles several non-square shapes, on both lowering paths.
// RUN: %mg-opt %inputs/dyn.mlir --convert-mg-to-affine %lower-to-llvm -o %t.a.mlir
// RUN: %mlir-translate --mlir-to-llvmir %t.a.mlir -o %t.a.ll
// RUN: %clang -c %t.a.ll -o %t.a.o
// RUN: %clang %inputs/dyn_harness.c %t.a.o -o %t.a.exe
// RUN: %t.a.exe | %FileCheck %s
// RUN: %mg-opt %inputs/dyn.mlir --one-shot-bufferize="bufferize-function-boundaries" -o %t.b0.mlir
// RUN: %mg-opt %t.b0.mlir %lower-to-llvm -o %t.b.mlir
// RUN: %mlir-translate --mlir-to-llvmir %t.b.mlir -o %t.b.ll
// RUN: %clang -c %t.b.ll -o %t.b.o
// RUN: %clang %inputs/dyn_harness.c %t.b.o -o %t.b.exe
// RUN: %t.b.exe | %FileCheck %s
// CHECK-LABEL: add 2x3 -> 2x3:
// CHECK-NEXT: 11 22 33
// CHECK-NEXT: 44 55 66
// CHECK-LABEL: add 3x2 -> 3x2:
// CHECK-NEXT: 11 22
// CHECK-NEXT: 33 44
// CHECK-NEXT: 55 66
// CHECK-LABEL: add 1x6 -> 1x6:
// CHECK-NEXT: 11 22 33 44 55 66
// CHECK-LABEL: transpose 2x3 -> 3x2:
// CHECK-NEXT: 1 4
// CHECK-NEXT: 2 5
// CHECK-NEXT: 3 6
// CHECK-LABEL: transpose 1x6 -> 6x1:
// CHECK-NEXT: 1
// CHECK-NEXT: 2
// CHECK-NEXT: 3
// CHECK-NEXT: 4
// CHECK-NEXT: 5
// CHECK-NEXT: 6
add 2x3 -> 2x3:
11 22 33
44 55 66
add 3x2 -> 3x2:
11 22
33 44
55 66
add 1x6 -> 1x6:
11 22 33 44 55 66
transpose 2x3 -> 3x2:
1 4
2 5
3 6
transpose 1x6 -> 6x1:
1
2
3
4
5
6
The harness that makes those calls (Inputs/dyn_harness.c) builds the rank-2 memref descriptor Chapter 8 introduced, with sizes chosen at runtime:
#include <stdint.h>
#include <stdio.h>
typedef struct { double *allocated; double *aligned; int64_t offset; int64_t sizes[2]; int64_t strides[2]; } Memref2D;
extern Memref2D add_dyn(double *a0, double *a1, int64_t off0, int64_t s0_0, int64_t s0_1, int64_t st0_0, int64_t st0_1,
double *b0, double *b1, int64_t off1, int64_t s1_0, int64_t s1_1, int64_t st1_0, int64_t st1_1);
extern Memref2D t_dyn(double *a0, double *a1, int64_t off0, int64_t s0_0, int64_t s0_1, int64_t st0_0, int64_t st0_1);
static void show(const char *name, Memref2D r) {
printf("%s -> %ldx%ld:\n", name, (long)r.sizes[0], (long)r.sizes[1]);
for (int64_t i = 0; i < r.sizes[0]; i++) {
for (int64_t j = 0; j < r.sizes[1]; j++) printf("%g ", r.aligned[r.offset + i * r.strides[0] + j * r.strides[1]]);
printf("\n");
}
}
int main(void) {
double a[6] = {1, 2, 3, 4, 5, 6}, b[6] = {10, 20, 30, 40, 50, 60};
show("add 2x3", add_dyn(a, a, 0, 2, 3, 3, 1, b, b, 0, 2, 3, 3, 1));
show("add 3x2", add_dyn(a, a, 0, 3, 2, 2, 1, b, b, 0, 3, 2, 2, 1));
show("add 1x6", add_dyn(a, a, 0, 1, 6, 6, 1, b, b, 0, 1, 6, 6, 1));
show("transpose 2x3", t_dyn(a, a, 0, 2, 3, 3, 1));
show("transpose 1x6", t_dyn(a, a, 0, 1, 6, 6, 1));
#ifdef STRIDED
// The same 6 numbers viewed as a 2x3 matrix in COLUMN-major order: strides (1, 2) instead of (3, 1).
// Logical matrix: [[1,3,5],[2,4,6]]. Only a callee whose type carries strides can honor this.
show("transpose 2x3 column-major view", t_dyn(a, a, 0, 2, 3, 1, 2));
#endif
return 0;
}
The column-major view only the bufferization path's strided types can honor:
// Chapter 13: only the One-Shot Bufferize path's strided argument type honors a column-major view.
// RUN: %mg-opt %inputs/dyn.mlir --one-shot-bufferize="bufferize-function-boundaries" -o %t.b0.mlir
// RUN: %mg-opt %t.b0.mlir %lower-to-llvm -o %t.b.mlir
// RUN: %mlir-translate --mlir-to-llvmir %t.b.mlir -o %t.b.ll
// RUN: %clang -c %t.b.ll -o %t.b.o
// RUN: %clang -DSTRIDED %inputs/dyn_harness.c %t.b.o -o %t.b.exe
// RUN: %t.b.exe | %FileCheck %s
// CHECK-LABEL: transpose 2x3 column-major view -> 3x2:
// CHECK-NEXT: 1 2
// CHECK-NEXT: 3 4
// CHECK-NEXT: 5 6
add 2x3 -> 2x3:
11 22 33
44 55 66
add 3x2 -> 3x2:
11 22
33 44
55 66
add 1x6 -> 1x6:
11 22 33 44 55 66
transpose 2x3 -> 3x2:
1 4
2 5
3 6
transpose 1x6 -> 6x1:
1
2
3
4
5
6
transpose 2x3 column-major view -> 3x2:
1 2
3 4
5 6
The two runtime-mismatch tests expect the program to abort. %not --crash means "must die from a signal", and the test's last RUN: shows the explanatory message by running under stdbuf -oL (line-buffered), because Chapter 14 found the message is lost when stdout is a pipe:
// Chapter 14: a runtime shape mismatch must abort (SIGABRT), on both paths, and never return a result.
// RUN: %mg-opt %inputs/dyn.mlir --convert-mg-to-affine %lower-to-llvm -o %t.a.mlir
// RUN: %mlir-translate --mlir-to-llvmir %t.a.mlir -o %t.a.ll
// RUN: %clang -c %t.a.ll -o %t.a.o
// RUN: %clang %inputs/mismatch_harness.c %t.a.o -o %t.a.exe
// RUN: %not --crash %t.a.exe
// RUN: %mg-opt %inputs/dyn.mlir --one-shot-bufferize="bufferize-function-boundaries" -o %t.b0.mlir
// RUN: %mg-opt %t.b0.mlir %lower-to-llvm -o %t.b.mlir
// RUN: %mlir-translate --mlir-to-llvmir %t.b.mlir -o %t.b.ll
// RUN: %clang -c %t.b.ll -o %t.b.o
// RUN: %clang %inputs/mismatch_harness.c %t.b.o -o %t.b.exe
// RUN: %not --crash %t.b.exe
// The message goes to stdout (puts) and abort() does not flush it, so only a line-buffered run shows it:
// RUN: %not --crash stdbuf -oL %t.a.exe 2>&1 | %FileCheck %s
// CHECK: mg.add: operand shapes differ at runtime in dimension 0
The second one mismatches only the columns (2x3 plus 2x2, rows equal). It was added after the failure-injection check below exposed that no native test covered that case:
// Chapter 15: rows agree, COLUMNS differ (2x3 + 2x2). Only a dimension-1 check can catch this, so this guards against a check that covers just dimension 1.
// RUN: %mg-opt %inputs/dyn.mlir --convert-mg-to-affine %lower-to-llvm -o %t.a.mlir
// RUN: %mlir-translate --mlir-to-llvmir %t.a.mlir -o %t.a.ll
// RUN: %clang -c %t.a.ll -o %t.a.o
// RUN: %clang %inputs/mismatch_cols_harness.c %t.a.o -o %t.a.exe
// RUN: %not --crash %t.a.exe
// RUN: %mg-opt %inputs/dyn.mlir --one-shot-bufferize="bufferize-function-boundaries" -o %t.b0.mlir
// RUN: %mg-opt %t.b0.mlir %lower-to-llvm -o %t.b.mlir
// RUN: %mlir-translate --mlir-to-llvmir %t.b.mlir -o %t.b.ll
// RUN: %clang -c %t.b.ll -o %t.b.o
// RUN: %clang %inputs/mismatch_cols_harness.c %t.b.o -o %t.b.exe
// RUN: %not --crash %t.b.exe
// The message goes to stdout (puts) and abort() does not flush it, so only a line-buffered run shows it:
// RUN: %not --crash stdbuf -oL %t.a.exe 2>&1 | %FileCheck %s
// CHECK: mg.add: operand shapes differ at runtime in dimension 1
#include <stdint.h>
#include <stdio.h>
typedef struct { double *allocated; double *aligned; int64_t offset; int64_t sizes[2]; int64_t strides[2]; } Memref2D;
extern Memref2D add_dyn(double*,double*,int64_t,int64_t,int64_t,int64_t,int64_t,double*,double*,int64_t,int64_t,int64_t,int64_t,int64_t);
int main(void) {
double a[6] = {1,2,3,4,5,6}, b[4] = {10,20,30,40};
// Rows agree (2 == 2), COLUMNS do not (3 vs 2): only a dimension-1 check can catch this.
Memref2D r = add_dyn(a,a,0,2,3,3,1, b,b,0,2,2,2,1);
printf("returned %ldx%ld (should be unreachable)\n", (long)r.sizes[0], (long)r.sizes[1]);
return 0;
}
And the mixed static/dynamic case, matching sizes then mismatched:
// Chapter 14: ?x2 + 3x2: matching sizes work, mismatched sizes abort.
// RUN: %mg-opt %inputs/mixed.mlir --convert-mg-to-affine %lower-to-llvm -o %t.m.mlir
// RUN: %mlir-translate --mlir-to-llvmir %t.m.mlir -o %t.m.ll
// RUN: %clang -c %t.m.ll -o %t.m.o
// RUN: %clang %inputs/mixed_harness.c %t.m.o -o %t.m.exe
// RUN: %t.m.exe 3 | %FileCheck %s
// RUN: %not --crash %t.m.exe 2
// CHECK: mixed with a=3x2, b=3x2 -> 3x2: 11 22 33 44 55 66
mixed with a=3x2, b=3x2 -> 3x2: 11 22 33 44 55 66
mg.add: operand shapes differ at runtime in dimension 0
Aborted
exit status: 134
A failure while writing it¶
The first full run had 19 passes and 2 failures, both tests using not --crash. lit's built-in not printed:
for that form. The verifier tests that wrote plain not had passed, but at that point a passing not test could not be told apart from a broken one, which is exactly why the next section exists. The fix was to stop relying on lit's builtin and route every not through LLVM's own binary: %not is defined as not-18 in lit.cfg.py above, and not-18 supports --crash. After that, 22 of 22.
Passing is not evidence: can the tests fail?¶
A suite that has only ever passed has not been shown to test anything. Two checks, both scripted.
Check 1: the same tests against older builds¶
older_builds.sh points the unchanged suite at two earlier mg-opt builds. Every failure should be a feature that build lacks.
#!/bin/sh
# Point the SAME suite at two older mg-opt builds to show the tests can fail:
# - Chapter 7's build (no dynamic shapes at all), assembled from Chapters 3, 5, 6, 7's own files
# - Chapter 13's build (dynamic shapes, no runtime check), built by Chapter 13's own build.sh
# READ THIS FIRST: this script runs the SAME test suite against two OLDER builds of the compiler on purpose.
# FAIL lines (and non-zero exit statuses) in this script's output are the EXPECTED result: each older build lacks features the tests check (Chapter 7: no dynamic shapes; Chapter 13: no run-time check), so the tests that need them SHOULD fail. That proves the tests can 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."
set -e
HERE=$(cd "$(dirname "$0")" && pwd); D=$HERE/../..; W=$HERE/work/ch7; T=$W/tree
mkdir -p $T/include/mg $T/lib $T/tools $W/build/include/mg
cp $D/part7/code/CMakeLists.txt $T/
cp $D/part3/code/MgDialect.td $D/part3/code/MgOps.td $D/part3/code/MgDialect.h $D/part3/code/MgOps.h $T/include/mg/
cp $D/part3/code/MgDialect.cpp $D/part5/code/LowerToAffine.cpp $D/part6/code/MgBufferizableOpInterfaceImpl.cpp $T/lib/
cp $D/part7/code/mg-opt.cpp $T/tools/
(cd $W/build && cmake $T -DMLIR_DIR=/usr/lib/llvm-18/lib/cmake/mlir -DLLVM_DIR=/usr/lib/llvm-18/lib/cmake/llvm \
-DCMAKE_BUILD_TYPE=Release -DLLVM_ENABLE_ASSERTIONS=OFF >/dev/null 2>&1 && make -j8 2>&1 | tail -1)
[ -x $D/part13/code/build/mg-opt ] || $D/part13/code/build.sh | tail -1
for b in "Chapter 7 build (no dynamic shapes):$W/build/mg-opt" "Chapter 13 build (dynamic shapes, no runtime check):$D/part13/code/build/mg-opt"; do
echo "=================== ${b%%:*}"
MG_OPT=${b#*:} MG_TEST_TMP=$HERE/work/old-out lit "$HERE/test" 2>&1 | grep -E "^FAIL|Passed|Failed:" | sed 's/IR Forge \/ Mountain Goat :: //;s/ ([0-9]* of [0-9]*)//' | sort -u
done
Real output (the 22-test suite; Chapter 16 re-runs this with 34):
[100%] Built target mg-opt
=================== Chapter 7 build (no dynamic shapes)
Failed: 13 (59.09%)
Passed: 9 (40.91%)
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 (27.27%)
Passed: 16 (72.73%)
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
The Chapter 7 build fails 13 and passes 9: every dynamic-shape test, the mixed-type canonicalizer test, the unranked test and every runtime-check test fail. The nine passes are the static ones plus the dynamic-extent transpose fold (Chapter 3's original pattern already handled the exact-type case), which is right: that behavior never changed. The Chapter 13 build, which has dynamic shapes but no runtime check, fails exactly the tests about the runtime check. One failure on the Chapter 7 build is not a missing feature: verifier/add-static-mismatch fails because the error wording changed in Chapter 13 ("operands must have the same shape" became "operands must have compatible shapes"), not because the behavior (rejecting the program) changed. Pinning diagnostic text makes a test brittle in that way. It was kept deliberately: if the message changes, a human should look.
Check 2: three deliberate bugs in the current source¶
mutation.sh copies Chapter 14's source tree three times, introduces one bug in each copy, builds it, and runs the suite:
#!/bin/sh
# Three deliberate source bugs, each applied to a COPY of Chapter 14's source tree, built, and tested.
# Needs ../../part14/code/build.sh to have been run (it creates the tree/ directory this copies).
# READ THIS FIRST: this script injects deliberate bugs ("mutations") into the compiler 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; that is how we know the tests would notice that bug.
# 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."
set -e
HERE=$(cd "$(dirname "$0")" && pwd); SRC=$HERE/../../part14/code/tree
mutate() { # name, python edit snippet (operates on files under $1)
name=$1; W=$HERE/work/mut-$name; rm -rf "$W"; mkdir -p "$W/build/include/mg"; cp -r "$SRC" "$W/tree"
python3 -c "$2" "$W/tree"
(cd "$W/build" && cmake "$W/tree" -DMLIR_DIR=/usr/lib/llvm-18/lib/cmake/mlir -DLLVM_DIR=/usr/lib/llvm-18/lib/cmake/llvm \
-DCMAKE_BUILD_TYPE=Release -DLLVM_ENABLE_ASSERTIONS=OFF >/dev/null 2>&1 && make -j8 2>&1 | grep -E " error " || true)
MG_OPT=$W/build/mg-opt MG_TEST_TMP=$W/out lit "$HERE/test" 2>&1 | grep -E "^FAIL|Passed|Failed:" | sed 's/IR Forge \/ Mountain Goat :: //;s/ ([0-9]* of [0-9]*)//' | sort -u
}
echo "##### MUTATION 1: canonicalizer's exact-type guard removed"
mutate 1 'import sys; p=sys.argv[1]+"/lib/MgDialect.cpp"; s=open(p).read(); a=" if (innerTranspose.getInput().getType() != op.getType())\n return mlir::failure();\n"; assert a in s; open(p,"w").write(s.replace(a,""))'
echo; echo "##### MUTATION 2: runtime check only covers dimension 0"
mutate 2 'import sys; p=sys.argv[1]+"/include/mg/DynamicShapes.h"; s=open(p).read(); a=" for (int64_t d = 0; d < lhsType.getRank(); ++d) {\n if (!lhsType.isDynamicDim(d) && !rhsType.isDynamicDim(d))"; assert a in s; open(p,"w").write(s.replace(a,a.replace("d < lhsType.getRank()","d < 1")))'
echo; echo "##### MUTATION 3: bufferization path forgets to emit the check"
mutate 3 'import sys; p=sys.argv[1]+"/lib/MgBufferizableOpInterfaceImpl.cpp"; s=open(p).read(); i=s.index(" mg::dyn::assertSameShape("); j=s.index("\n llvm::SmallVector<Value, 2> extents;", i); open(p,"w").write(s[:i]+s[j:])'
Real output:
##### MUTATION 1: canonicalizer's exact-type guard removed
Failed: 1 (4.55%)
Passed: 21 (95.45%)
FAIL: canonicalize/transpose-twice-dynamic-mixed.mlir
##### MUTATION 2: runtime check only covers dimension 0
Failed: 3 (13.64%)
Passed: 19 (86.36%)
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 (13.64%)
Passed: 19 (86.36%)
FAIL: execution/mismatch-aborts.mlir
FAIL: execution/mismatch-columns-aborts.mlir
FAIL: lowering/bufferize-dynamic-add.mlir
Every mutation is caught, by the tests aimed at it: the canonicalizer guard (Chapter 13's bug returning) by transpose-twice-dynamic-mixed; the missing dimension check by the lowering tests that expect an assert for dimension 1 and by the column-only mismatch test; the forgotten bufferization check by bufferize-dynamic-add and the abort tests.
The second mutation exposed a weakness in the suite as first written, which is the main thing this check is for. With 21 tests, "check only dimension 0" was caught only by tests of the IR's structure; no native test ever mismatched the columns while the rows agreed, because Chapter 13's mismatch case differs in rows first. A runtime bug that structural tests happened to anticipate would have slipped past every native test. The 22nd test, mismatch-columns-aborts (shown in Group 5), was added to close that, and the mutation was re-run: it now fails there too.
What the suite does not cover¶
- No CI. Nothing runs these tests automatically on a push; they run when someone executes
./run_lit.sh. A GitHub Actions workflow is the natural next step and is not written here, partly because it would need the MLIR 18 packages installed on the runner, which this sandbox cannot test. - Coverage at the end of this chapter: Chapters 1, 2, 9, 10, 11 and 12 had no tests. (Update, added after Chapter 16: Chapters 10 through 12 and Chapter 7's tiling and unrolling are now covered, taking the suite to 34 tests; Chapters 1, 2 and 9 are still untested.)
- Loop transforms: one test at the end of this chapter (fusion only).
- Tests encode current behavior. They show the compiler behaves as it does now, not that it is correct; the expected values were taken from outputs Chapters 3 through 14 had already validated by hand.
- The abort message is only checked under
stdbuf -oL, because testing it when piped would be testing the C library, notmg. - Diagnostics are pinned verbatim, with the brittleness described above.
Reproducing this chapter¶
pip install lit # lit 23.1.2 was used
../part14/code/build.sh # builds the mg-opt under test
cd docs/part15/code
./run_lit.sh # the suite
./show.sh # regenerates every output shown on this page under show/
./older_builds.sh # the same suite against Chapter 7's and Chapter 13's builds
./mutation.sh # the three source mutations
Every test, script 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 14, no documentation was consulted for any claim here; the not --crash behavior, the lit version mismatch and every result above come from running the commands.
Chapter summary¶
This chapter replaced "run once and paste the output" with an assertion-based suite of 22 lit/FileCheck tests across the verifier, the canonicalizer, both lowering paths, loop fusion and native execution, every one printed on the page beside the real output it matches. A first-run failure (lit's built-in not rejecting --crash) was fixed by routing not through LLVM's own not-18. Because a suite that has only passed has not been shown to test anything, the same tests were run against older builds (13 failures on Chapter 7's, 6 on Chapter 13's, each a feature that build lacks) and against three deliberate source bugs, each caught by the tests aimed at it. The mutation check exposed a real gap in the suite as first written, a check covering only dimension 0, closed by a column-only mismatch test.
Deliberately out of scope, stated explicitly: no CI runs the suite on a push, Chapters 1, 2, 9 and the GPU chapters had no tests at the end of this chapter (Chapter 16 covers the GPU path and two more loop transforms), lit came from PyPI at a different version from the toolchain, and diagnostics are pinned verbatim, which makes some tests brittle.
Self-check questions¶
1. What does CHECK-NOT assert, and which tests lean on it to show that the dynamic-shape work costs static programs nothing?
Worked answer: CHECK-NOT: text requires that text does not appear between the surrounding matches. The static lowering tests (static-add-affine and bufferize-static-add) use CHECK-NOT: cf.assert and CHECK-NOT: memref.dim after the loop checks, so if the dynamic machinery leaked into a static program the test would fail.
2. Why is the mg-opt under test chosen by an environment variable rather than hard-coded in lit.cfg.py?
Worked answer: so the identical suite can be pointed at other builds. That is how the chapter shows the tests can fail: older_builds.sh runs the unchanged tests with MG_OPT set to Chapter 7's and Chapter 13's builds, and mutation.sh sets it to builds with deliberate bugs. A suite tied to one binary could only ever report on that binary.
3. The two not --crash tests failed on the first run, while the verifier tests written with plain not passed. Why was that not good enough, and what was the fix?
Worked answer: lit's builtin not did not support --crash ("not: command not found" for that form), so those tests could not run. The plain-not tests passing was not reassuring either, because at that point a passing not test could not be told apart from a broken one. The fix was to define %not as LLVM's not-18, which supports --crash, and use it everywhere; the older-build runs then showed the verifier tests really do fail when the behavior is absent.
4. What did mutating the code to check only dimension 0 reveal about the suite as first written?
Worked answer: with 21 tests, that mutation was caught only by tests of the IR's structure (the lowering tests expecting an assert for dimension 1). No native test mismatched the columns while the rows agreed, because the only mismatch case differs in rows first, so a runtime bug in dimension 1 could have passed every executable test. The 22nd test, mismatch-columns-aborts (2x3 plus 2x2), was added and the mutation re-run: it now fails there too.
5. Why is pinning an exact diagnostic message a double-edged choice? Give the example from this chapter.
Worked answer: it makes the test strict about wording, so it also fails when only the text changes. On Chapter 7's build, verifier/add-static-mismatch fails not because the program is accepted but because the message then read "operands must have the same shape" and Chapter 13 changed it to "operands must have compatible shapes". The test was kept pinned deliberately, on the view that a changed message should make a human look, but it is a brittleness a reader should know about.