From e6436b1bcaf5f350dbde921b44bd9422c865e9d2 Mon Sep 17 00:00:00 2001 From: ZaneHam Date: Wed, 12 Aug 2026 22:16:55 +1200 Subject: [PATCH] docs: cover the MLIR frontend --- README.md | 11 ++++++++--- docs/features.md | 2 ++ docs/usage.md | 29 +++++++++++++++++++++++++++++ 3 files changed, 39 insertions(+), 3 deletions(-) diff --git a/README.md b/README.md index 4be63d4..044b01e 100644 --- a/README.md +++ b/README.md @@ -6,9 +6,10 @@ It is named to honour Kathleen Booth: creator of the first assembly language, co A running log of what's changed is in [CHANGELOG.md](CHANGELOG.md). -**Update:** Fortran `do concurrent` kernels now compile to AMD, NVIDIA and -x86-64 through LFortran, checked against SLATEC values in CI. See -[Using Fortran](docs/usage.md#fortran). +**Update:** `--mlir` reads MLIR text, with no LLVM anywhere in the path. +`func.func` and the `arith` dialect lower to BIR and go down the same pipeline +CUDA and Triton use. It is a small subset on purpose, and anything outside it +is named and refused rather than skipped. See [Using MLIR](docs/usage.md#mlir). ## What It Does @@ -16,6 +17,10 @@ Takes CUDA C, HIP, or Triton source (the same files you'd hand to `nvcc`, `ROCm` That last one still surprises me a bit. You can write a Triton kernel, matmul and all, and run it on a machine that's never seen a GPU, from scratch, no LLVM, straight to native. I haven't come across anyone else doing Triton like this, but I'd happily be proven wrong, so give me a yell if you've seen it somewhere. +Fortran `do concurrent` kernels go down the same path through +[LFortran](https://lfortran.org/), checked against SLATEC values in CI. See +[Using Fortran](docs/usage.md#fortran). + It also borrows a pile of operational discipline from the mainframe world: real crash dumps when a kernel faults, structured output routed by class, parameter snapshots on entry. See [docs/mainframe.md](docs/mainframe.md) if that sounds like your kind of thing. ## Getting it diff --git a/docs/features.md b/docs/features.md index 0262f22..3787808 100644 --- a/docs/features.md +++ b/docs/features.md @@ -39,6 +39,7 @@ The following CUDA features compile to working GFX9/GFX10/GFX11/GFX12 machine co - Struct pass-by-value - Triton tile shape inference: rank-0/1/2 shape annotation on every expression, constexpr default propagation (`BLOCK: tl.constexpr = 256` resolves to `vec[256]`), numpy-style broadcasting, `[:, None]` / `[None, :]` reshape patterns - Triton matmul on the CPU: `tl.dot` lowers and runs via `--cpu`, with a runtime K-loop so the contraction can be any size. Rank-2 tiles materialise and unroll +- MLIR frontend (`--mlir`): `func.func`, `return`, `arith.constant`, the `arith` binops, compares and conversions lower to BIR and run through the same pipeline as CUDA and Triton. `--mlir --pp` reprints what was read - x86-64 CPU backend (`--cpu`): CUDA and Triton kernels compile to a host object and run with no GPU. SIMT becomes a thread loop. Stack-everything codegen, no register allocator yet - TDF (Tile DataFlow) IR layer above BIR: regions / channels / NoC arcs as first-class compiler concepts, L1 placement, fission pass for multi-core kernels - SYSPRINT: class-tagged structured kernel output, pattern-routed sinks on the host. See [mainframe.md](mainframe.md) for the kernel/host workflow. @@ -54,6 +55,7 @@ Being honest about limitations is important. Here's what's missing: - Host code generation (only device code is compiled) - Rank-2 matrix codegen on the GPU backends (MFMA on AMD, mma.sync on NVIDIA). Triton `tl.dot` already runs on the CPU backend, materialised and unrolled with a K-loop, but the GPU matrix-instruction path is a separate job. On GPU targets rank-2 tiles still refuse cleanly with E099, no silent wrong code. - CPU backend is correct-first: stack-everything codegen, no register allocator yet, single block per call, and `tl.load` masks aren't honoured (so keep the launch's nthreads equal to the element count). It runs; it isn't fast. +- MLIR beyond `func`, `return` and `arith`. No `memref`, `scf` or `gpu` dialect, and a function body with more than one block is refused rather than flattened. Nothing marks an MLIR function as a kernel, so it lowers as a device function: `--cpu` and `--rv64` emit real code, `--metal` counts it as a kernel, and `--amdgpu-bin` and `--nvidia-ptx` report zero kernels and write an empty container - Soft-float for the Tenstorrent native RV32IM path. The runtime exists and validates against host FPU; wiring it into `--rv-elf` is a sitting's work away. None of these are architectural blockers. They're all "haven't got round to it yet" items. diff --git a/docs/usage.md b/docs/usage.md index 5ff8eb5..5f7db08 100644 --- a/docs/usage.md +++ b/docs/usage.md @@ -193,3 +193,32 @@ no FPU, so float has to go through a soft-float runtime, which is in progress. If something breaks and you cannot tell whether it is an LFortran gap or a Booth one, raise it here and it will get sorted out from this side. + +## MLIR + +`--mlir` reads MLIR text. The reader is Ondřej Čertík's, vendored under +`src/mlir/vendor`, and there is no LLVM in the path. + +```bash +# Lower MLIR and compile it, same as any other frontend +./kath --mlir kernel.mlir --cpu -o kernel.o +./kath --mlir kernel.mlir --rv64 -o kernel.o + +# Dump the BIR the lowering produced +./kath --mlir --ir kernel.mlir + +# Reprint what was read instead of lowering it, for telling a misreading +# from a bad file +./kath --mlir --pp kernel.mlir +``` + +Accepted today: `func.func` with named arguments, `return`, `arith.constant`, +every `arith` integer and float binop, `arith.cmpi` and `arith.cmpf`, and the +`arith` conversions. Anything else is named on stderr and the whole lowering +fails, because an op skipped quietly is a kernel that compiles and computes +something different. + +Nothing in MLIR marks a function as a kernel yet, so everything lowers as a +device function. `--cpu` and `--rv64` give you real code; the GPU backends will +take the file and then report zero kernels, which is an honest answer until +the `gpu` dialect arrives.