docs: cover the MLIR frontend

README.md

@@ -6,16 +6,21 @@ 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
 
 Takes CUDA C, HIP, or Triton source (the same files you'd hand to `nvcc`, `ROCm`, or Triton's JIT) and turns them into AMD RDNA 2/3/4 binaries, NVIDIA PTX, Tenstorrent Metalium C++ or native RV32IM, or just plain x86-64 you can run on a laptop with no GPU in it.
 
 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

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.

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.