diff --git a/.github/workflows/cmake.yml b/.github/workflows/cmake.yml
new file mode 100644
index 0000000..b738127
--- /dev/null
+++ b/.github/workflows/cmake.yml
@@ -0,0 +1,82 @@
+# Booth has no CMake build of its own; these jobs check the other direction,
+# that a CMake project can still consume an installed Booth.
+name: CMake package
+
+on:
+  push:
+    branches: [master]
+  pull_request:
+    branches: [master]
+
+jobs:
+  consumer:
+    name: Example consumer builds
+    runs-on: ubuntu-latest
+    steps:
+      - uses: actions/checkout@v4
+
+      - name: Build Booth
+        run: make
+
+      - name: Install to a prefix
+        run: make install PREFIX="$PWD/_prefix"
+
+      - name: Configure the example
+        run: cmake -S examples/cmake -B _example -DCMAKE_PREFIX_PATH="$PWD/_prefix"
+
+      - name: Build the example
+        run: cmake --build _example
+
+      - name: Check the artefacts
+        run: |
+          cd _example
+          for f in vadd_ptx.ptx vadd_gfx942.hsaco vadd_asm.s vadd_tt.cpp vadd_host.o; do
+            test -s "$f" || { echo "missing or empty: $f"; exit 1; }
+          done
+          # --tensix emits a whole Metalium program, so the siblings the
+          # package config declares as byproducts have to be there too.
+          for f in vadd_tt_host.cpp vadd_tt_reader.cpp vadd_tt_writer.cpp; do
+            test -s "$f" || { echo "missing or empty tensix sibling: $f"; exit 1; }
+          done
+          # --amdgpu once ignored -o and wrote to stdout with the register-plan
+          # line mixed in. Assembly on line one is what says it still honours it.
+          head -1 vadd_asm.s | grep -q amdgcn_target || {
+            echo "vadd_asm.s does not start with assembly:"; head -3 vadd_asm.s; exit 1; }
+          echo "all artefacts present"
+
+      - name: Rebuild does no work
+        run: |
+          out=$(cmake --build _example)
+          echo "$out"
+          # if, not &&: under bash -e a failing grep in an && list would fail
+          # the step in exactly the case we want to pass.
+          if echo "$out" | grep -q "Booth "; then
+            echo "second build recompiled a kernel"
+            exit 1
+          fi
+          echo "incremental build was a no-op"
+
+  arch-sync:
+    name: Arch list matches the compiler
+    runs-on: ubuntu-latest
+    steps:
+      - uses: actions/checkout@v4
+
+      # The package config carries its own copy of the target list so it can
+      # reject a typo at configure time. That copy has to keep up with main.c,
+      # and nothing else would notice if it did not.
+      - name: Compare target lists
+        run: |
+          grep -oE '"--(gfx[0-9]+[a-z]?|xe[0-9a-z-]*)"' src/main.c \
+            | tr -d '"' | sed 's/^--//' | sort -u > /tmp/kath_arches
+          sed -n '/set(_arches/,/)/p' cmake/BoothConfig.cmake.in \
+            | sed -e 's/set(_arches//' -e 's/)//' | tr -s ' \t' '\n' \
+            | grep -E '^(gfx|xe)' | sort -u > /tmp/cmake_arches
+          if ! diff -u /tmp/kath_arches /tmp/cmake_arches; then
+            echo
+            echo "The target list in cmake/BoothConfig.cmake.in has drifted from"
+            echo "the flags main.c accepts. Lines marked - are in the compiler but"
+            echo "not the package config; + is the other way round."
+            exit 1
+          fi
+          echo "target lists agree ($(wc -l < /tmp/kath_arches) targets)"
diff --git a/.gitignore b/.gitignore
index 8b909eb..0d7c183 100644
--- a/.gitignore
+++ b/.gitignore
@@ -155,3 +155,8 @@ tdf_flag_out.txt
 
 # per-host object trees
 /build/
+
+# Where the CMake package job stages an install and builds the example
+# consumer. Same paths locally if you follow examples/cmake/CMakeLists.txt.
+/_prefix/
+/_example/
diff --git a/CHANGELOG.md b/CHANGELOG.md
index c26c795..8c3c4f9 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -39,6 +39,11 @@ Booth — Changelog
   no longer lose the lanes that did not take the branch
   (Zane Hambly, 2026-07-25)
 
+- `--amdgpu` honours `-o`, and the register-plan line goes to stderr rather
+  than into the middle of the assembly on stdout, where it stopped the result
+  assembling
+  (Zane Hambly, 2026-08-06)
+
 ### Tensix
 
 - `--tt-chip` selects wormhole or blackhole, so L1 and text limits follow the
@@ -82,6 +87,11 @@ Booth — Changelog
   the linker a mix of COFF and ELF
   (Zane Hambly, 2026-07-28)
 
+- `make install`, honouring `PREFIX` and `DESTDIR`, and a CMake package config
+  alongside it, so a downstream project can `find_package(Booth)` and build
+  kernels with `booth_add_kernel()`
+  (Zane Hambly, 2026-08-06)
+
 ### CI and tests
 
 - #140: numeric regression against SLATEC known-good values, across cpu,
@@ -92,6 +102,11 @@ Booth — Changelog
   block from the kernel rather than a fixed 64 bytes
   (Zane Hambly, 2026-07-27)
 
+- build an example CMake consumer against a staged install, and check the
+  target list in the package config has not drifted from the flags the
+  compiler accepts
+  (Zane Hambly, 2026-08-06)
+
 - validate Tensix ELFs against tt-metal's loader and run RV64 under QEMU
   (Zane Hambly, 2026-07-23)
 
@@ -100,6 +115,10 @@ Booth — Changelog
 
 ### Documentation
 
+- document consuming Booth from CMake in `docs/cmake.md`, and link it from the
+  README
+  (Zane Hambly, 2026-08-06)
+
 - drop the LLVM requirement from the usage documentation
   (Zane Hambly, 2026-07-27)
 
diff --git a/Makefile b/Makefile
index 83b0528..30994a1 100644
--- a/Makefile
+++ b/Makefile
@@ -140,6 +140,45 @@ $(OBJDIR)/runtime/%.o: runtime/%.c
 	@mkdir -p $(dir $@)
 	$(CC) $(TCFLAGS) -c $< -o $@
 
+# ---- Install ----
+# Booth is consumed as a compiler rather than a library, so an install is the
+# binary, the message catalogues --lang reads, and a CMake package config so
+# downstream CMake projects can find_package(Booth) and run kath as a build
+# step. Version comes out of the header so there is one place to bump it.
+PREFIX  ?= /usr/local
+BINDIR   = $(DESTDIR)$(PREFIX)/bin
+SHAREDIR = $(DESTDIR)$(PREFIX)/share/booth
+CMAKEDIR = $(DESTDIR)$(PREFIX)/lib/cmake/Booth
+
+# Matching on the macro name rather than the whole "#define" line: make 3.81,
+# which is what macOS ships, takes a # inside $(shell) as the start of a
+# comment and swallows the rest of the call. Quoting does not save it, and
+# neither does awk over sed, so the # simply has to go.
+VER_MAJOR := $(shell awk '$$2 == "BC_VERSION_MAJOR" {print $$3}' src/barracuda.h)
+VER_MINOR := $(shell awk '$$2 == "BC_VERSION_MINOR" {print $$3}' src/barracuda.h)
+VER_PATCH := $(shell awk '$$2 == "BC_VERSION_PATCH" {print $$3}' src/barracuda.h)
+VERSION   := $(VER_MAJOR).$(VER_MINOR).$(VER_PATCH)
+
+# MinGW gcc appends .exe when -o names no suffix, so the built file is not
+# always $(TARGET); clean has always known this, install needs to as well.
+EXE :=
+ifneq (,$(findstring MINGW,$(UNAME_S)))
+  EXE := .exe
+endif
+
+install: $(TARGET)
+	install -d $(BINDIR) $(SHAREDIR)/lang $(CMAKEDIR)
+	install -m 755 $(TARGET)$(EXE) $(BINDIR)/$(TARGET)$(EXE)
+	install -m 644 lang/en.txt lang/mi.txt $(SHAREDIR)/lang/
+	sed -e 's/@BOOTH_VERSION@/$(VERSION)/g' -e 's/@BOOTH_VERSION_MAJOR@/$(VER_MAJOR)/g' \
+	    cmake/BoothConfig.cmake.in > $(CMAKEDIR)/BoothConfig.cmake
+	sed -e 's/@BOOTH_VERSION@/$(VERSION)/g' -e 's/@BOOTH_VERSION_MAJOR@/$(VER_MAJOR)/g' \
+	    cmake/BoothConfigVersion.cmake.in > $(CMAKEDIR)/BoothConfigVersion.cmake
+
+uninstall:
+	rm -f $(BINDIR)/$(TARGET)$(EXE)
+	rm -rf $(SHAREDIR) $(CMAKEDIR)
+
 clean:
 	rm -rf $(OBJDIR)
 	rm -f $(TARGET) $(TARGET).exe trunner trunner.exe
@@ -148,4 +187,4 @@ clean:
 # linked in and the build silently disagrees with the source.
 -include $(OBJECTS:.o=.d) $(TOBJS:.o=.d) $(HOSTRT:.o=.d)
 
-.PHONY: all clean test
+.PHONY: all clean test install uninstall
diff --git a/README.md b/README.md
index 2b46bfd..4ddc808 100644
--- a/README.md
+++ b/README.md
@@ -36,6 +36,7 @@ The binary is `kath` (after Kathleen but if she picked Australia or New Zealand
 ## Documentation
 
 - **[Usage](docs/usage.md)** — every backend, every flag, and the runtime launcher
+- **[CMake](docs/cmake.md)** — installing Booth and compiling kernels from a CMake project
 - **[Feature status](docs/features.md)** — what compiles today, and what doesn't yet
 - **[Mainframe curios](docs/mainframe.md)** — ABEND dumps, SNAP, SYSPRINT, TDF
 - **[Validated hardware](docs/hardware.md)** — the silicon it's been tested on, and the test suite
diff --git a/cmake/BoothConfig.cmake.in b/cmake/BoothConfig.cmake.in
new file mode 100644
index 0000000..65a81aa
--- /dev/null
+++ b/cmake/BoothConfig.cmake.in
@@ -0,0 +1,181 @@
+# Booth package config, for find_package(Booth CONFIG).
+#
+# Booth is a compiler, not a library, so there is nothing here to link
+# against. What you get is the kath binary as an imported target and a
+# booth_add_kernel() rule that runs it over a source file at build time.
+
+set(Booth_VERSION "@BOOTH_VERSION@")
+
+# This file installs to <prefix>/lib/cmake/Booth, so the prefix is three
+# levels up. Deriving it rather than baking it in keeps a relocated or
+# staged install working.
+get_filename_component(_booth_prefix "${CMAKE_CURRENT_LIST_DIR}/../../.." ABSOLUTE)
+
+find_program(Booth_EXECUTABLE
+    NAMES kath
+    HINTS "${_booth_prefix}/bin")
+
+if(NOT Booth_EXECUTABLE)
+    set(Booth_FOUND FALSE)
+    set(Booth_NOT_FOUND_MESSAGE
+        "Booth config was found but the kath binary was not; expected it in ${_booth_prefix}/bin")
+    return()
+endif()
+
+if(NOT TARGET Booth::kath)
+    add_executable(Booth::kath IMPORTED)
+    set_target_properties(Booth::kath PROPERTIES IMPORTED_LOCATION "${Booth_EXECUTABLE}")
+endif()
+
+# Message catalogues for --lang. Absent on a binary-only install, so callers
+# should treat this as a hint rather than a guarantee.
+set(Booth_LANG_DIR "${_booth_prefix}/share/booth/lang")
+
+# booth_add_kernel(<name>
+#     SOURCE <file>
+#     BACKEND amdgpu|amdgpu-bin|nvidia-ptx|tensix|rv-elf|cpu|rv64|metal|intel-spirv
+#     [LANGUAGE CUDA|HIP|TRITON]
+#     [ARCH <gfx…|xe…>]
+#     [TT_CHIP blackhole|wormhole]
+#     [OUTPUT <file>]
+#     [EXCLUDE_FROM_ALL]
+#     [INCLUDE_DIRECTORIES <dir>…]
+#     [COMPILE_DEFINITIONS <name[=val]>…]
+#     [OPTIONS <flag>…]
+#     [DEPENDS <file>…])
+#
+# Adds a custom target <name> that compiles SOURCE with kath. The output path
+# lands in the BOOTH_KERNEL_OUTPUT property on that target, and in a
+# <name>_OUTPUT variable in the caller's scope.
+#
+# The tables below live inside the function on purpose: a package config is
+# included in whatever scope called find_package, so anything left at
+# directory scope would go missing when the function is called from a sibling
+# directory that never called it.
+function(booth_add_kernel name)
+    set(_backends amdgpu amdgpu-bin nvidia-ptx tensix rv-elf cpu rv64 metal intel-spirv)
+    # Default suffix per backend, index-matched to _backends. Text for the
+    # assembly and source-emitting backends, loadable images for the rest.
+    set(_exts     ".s"   ".hsaco"   ".ptx"      ".cpp" ".elf"  ".o" ".o"  ".metal" ".spv")
+    set(_arches
+        gfx90a gfx942
+        gfx1030 gfx1031 gfx1032 gfx1033 gfx1034 gfx1035 gfx1036
+        gfx1100 gfx1101 gfx1102 gfx1103
+        gfx1150 gfx1151 gfx1152 gfx1153
+        gfx1200 gfx1201
+        xe-lpg xe-hpg xe-hpc xe2)
+
+    cmake_parse_arguments(PARSE_ARGV 1 BK
+        "EXCLUDE_FROM_ALL"
+        "SOURCE;BACKEND;LANGUAGE;ARCH;TT_CHIP;OUTPUT;LANG"
+        "OPTIONS;INCLUDE_DIRECTORIES;COMPILE_DEFINITIONS;DEPENDS")
+
+    if(BK_UNPARSED_ARGUMENTS)
+        message(FATAL_ERROR "booth_add_kernel(${name}): unrecognised arguments: ${BK_UNPARSED_ARGUMENTS}")
+    endif()
+    if(NOT BK_SOURCE)
+        message(FATAL_ERROR "booth_add_kernel(${name}): SOURCE is required")
+    endif()
+    if(NOT BK_BACKEND)
+        message(FATAL_ERROR "booth_add_kernel(${name}): BACKEND is required, one of: ${_backends}")
+    endif()
+
+    # list(FIND) rather than IN_LIST, which would rely on the consumer's
+    # CMP0057 being NEW.
+    list(FIND _backends "${BK_BACKEND}" _bi)
+    if(_bi EQUAL -1)
+        message(FATAL_ERROR "booth_add_kernel(${name}): unknown BACKEND '${BK_BACKEND}', want one of: ${_backends}")
+    endif()
+
+    # kath takes one input file per invocation, so a kernel is one source.
+    get_filename_component(_src "${BK_SOURCE}" ABSOLUTE BASE_DIR "${CMAKE_CURRENT_SOURCE_DIR}")
+
+    set(_args "--${BK_BACKEND}")
+
+    if(BK_LANGUAGE)
+        string(TOUPPER "${BK_LANGUAGE}" _lang)
+        if(_lang STREQUAL "HIP")
+            list(APPEND _args --hip)
+        elseif(_lang STREQUAL "TRITON")
+            list(APPEND _args --triton)
+        elseif(NOT _lang STREQUAL "CUDA")
+            message(FATAL_ERROR "booth_add_kernel(${name}): LANGUAGE must be CUDA, HIP or TRITON")
+        endif()
+    endif()
+
+    if(BK_ARCH)
+        list(FIND _arches "${BK_ARCH}" _ai)
+        if(_ai EQUAL -1)
+            message(FATAL_ERROR "booth_add_kernel(${name}): unknown ARCH '${BK_ARCH}', want one of: ${_arches}")
+        endif()
+        list(APPEND _args "--${BK_ARCH}")
+    endif()
+
+    if(BK_TT_CHIP)
+        if(NOT BK_TT_CHIP MATCHES "^(blackhole|wormhole)$")
+            message(FATAL_ERROR "booth_add_kernel(${name}): TT_CHIP must be blackhole or wormhole")
+        endif()
+        list(APPEND _args --tt-chip "${BK_TT_CHIP}")
+    endif()
+
+    foreach(_d IN LISTS BK_INCLUDE_DIRECTORIES)
+        list(APPEND _args -I "${_d}")
+    endforeach()
+    foreach(_d IN LISTS BK_COMPILE_DEFINITIONS)
+        list(APPEND _args -D "${_d}")
+    endforeach()
+    if(BK_LANG)
+        list(APPEND _args --lang "${BK_LANG}")
+    endif()
+    list(APPEND _args ${BK_OPTIONS})
+
+    if(BK_OUTPUT)
+        get_filename_component(_out "${BK_OUTPUT}" ABSOLUTE BASE_DIR "${CMAKE_CURRENT_BINARY_DIR}")
+    else()
+        list(GET _exts ${_bi} _ext)
+        set(_out "${CMAKE_CURRENT_BINARY_DIR}/${name}${_ext}")
+    endif()
+
+    get_filename_component(_outdir "${_out}" DIRECTORY)
+    file(MAKE_DIRECTORY "${_outdir}")
+
+    # --tensix emits a whole Metalium program, not one file: the compute kernel
+    # at OUTPUT plus reader/writer/host sources and their ELFs. Two naming rules
+    # are in play, one off the stem and one off the full OUTPUT name, so both
+    # are spelled out. Declaring them as byproducts is what lets `clean` reach
+    # them and stops Ninja treating them as unexplained.
+    set(_byproducts "")
+    if(BK_BACKEND STREQUAL "tensix")
+        get_filename_component(_outname "${_out}" NAME)
+        # Booth strips only the final extension, so a regex rather than NAME_WE
+        # (strips from the first dot) or NAME_WLE (needs CMake 3.14).
+        string(REGEX REPLACE "\\.[^./\\\\]*$" "" _stem "${_outname}")
+        foreach(_s _host.cpp _reader.cpp _writer.cpp _compute.elf _reader.elf _writer.elf)
+            list(APPEND _byproducts "${_outdir}/${_stem}${_s}")
+        endforeach()
+        foreach(_s _compute.bin _compute.ttinsn)
+            list(APPEND _byproducts "${_outdir}/${_outname}${_s}")
+        endforeach()
+    endif()
+
+    # Depending on the binary rebuilds kernels when the compiler itself
+    # changes. kath emits no depfile, so an edited header will not retrigger
+    # this on its own; list headers in DEPENDS if you need that.
+    add_custom_command(
+        OUTPUT "${_out}"
+        BYPRODUCTS ${_byproducts}
+        COMMAND "${Booth_EXECUTABLE}" ${_args} -o "${_out}" "${_src}"
+        DEPENDS "${_src}" "${Booth_EXECUTABLE}" ${BK_DEPENDS}
+        COMMENT "Booth ${BK_BACKEND}: ${name}"
+        VERBATIM)
+
+    if(BK_EXCLUDE_FROM_ALL)
+        add_custom_target(${name} DEPENDS "${_out}")
+    else()
+        add_custom_target(${name} ALL DEPENDS "${_out}")
+    endif()
+    set_target_properties(${name} PROPERTIES BOOTH_KERNEL_OUTPUT "${_out}")
+    set(${name}_OUTPUT "${_out}" PARENT_SCOPE)
+endfunction()
+
+set(Booth_FOUND TRUE)
diff --git a/cmake/BoothConfigVersion.cmake.in b/cmake/BoothConfigVersion.cmake.in
new file mode 100644
index 0000000..ce63aef
--- /dev/null
+++ b/cmake/BoothConfigVersion.cmake.in
@@ -0,0 +1,18 @@
+# Version check for find_package(Booth <version>). Same-major compatibility:
+# a consumer asking for 5.1 is satisfied by any later 5.x, never by 6.x.
+
+set(PACKAGE_VERSION "@BOOTH_VERSION@")
+
+if(PACKAGE_VERSION VERSION_LESS PACKAGE_FIND_VERSION)
+    set(PACKAGE_VERSION_COMPATIBLE FALSE)
+else()
+    if(PACKAGE_FIND_VERSION_MAJOR STREQUAL "@BOOTH_VERSION_MAJOR@")
+        set(PACKAGE_VERSION_COMPATIBLE TRUE)
+    else()
+        set(PACKAGE_VERSION_COMPATIBLE FALSE)
+    endif()
+
+    if(PACKAGE_FIND_VERSION STREQUAL PACKAGE_VERSION)
+        set(PACKAGE_VERSION_EXACT TRUE)
+    endif()
+endif()
diff --git a/docs/cmake.md b/docs/cmake.md
new file mode 100644
index 0000000..d560bb0
--- /dev/null
+++ b/docs/cmake.md
@@ -0,0 +1,84 @@
+# Using Booth from CMake
+
+Booth builds with its own Makefile; there is no CMake build of the compiler
+itself. What this gives you is the other direction, a CMake project consuming
+an installed Booth to compile kernels as part of its build.
+
+## Installing
+
+```sh
+make
+sudo make install                 # /usr/local by default
+make install PREFIX=$HOME/.local  # or wherever
+```
+
+That puts `kath` in `<prefix>/bin`, the `--lang` catalogues in
+`<prefix>/share/booth/lang`, and the package config in
+`<prefix>/lib/cmake/Booth`. `DESTDIR` is honoured for staged installs.
+
+## Finding it
+
+```cmake
+find_package(Booth 5.1 REQUIRED)
+```
+
+If the prefix is not one CMake already searches, point it there with
+`-DCMAKE_PREFIX_PATH=<prefix>`. Version matching is same-major: 5.1 is
+satisfied by any later 5.x, never by 6.x.
+
+You get `Booth_VERSION`, `Booth_EXECUTABLE`, `Booth_LANG_DIR`, and a
+`Booth::kath` imported target.
+
+## Compiling a kernel
+
+```cmake
+booth_add_kernel(vadd_amd
+    SOURCE  vadd.cu
+    BACKEND amdgpu-bin
+    ARCH    gfx942)
+```
+
+That adds a target `vadd_amd` that runs `kath` at build time. The output path
+comes back two ways, as a `BOOTH_KERNEL_OUTPUT` target property and as a
+`vadd_amd_OUTPUT` variable in the calling scope:
+
+```cmake
+add_custom_command(TARGET app POST_BUILD
+    COMMAND ${CMAKE_COMMAND} -E copy ${vadd_amd_OUTPUT} $<TARGET_FILE_DIR:app>)
+```
+
+The `cpu` and `rv64` backends emit ordinary relocatable objects, so those can
+go straight into a target:
+
+```cmake
+booth_add_kernel(vadd_host SOURCE vadd.cu BACKEND cpu)
+add_executable(app main.c)
+add_dependencies(app vadd_host)
+target_link_libraries(app PRIVATE ${vadd_host_OUTPUT})
+```
+
+### Arguments
+
+| | |
+|---|---|
+| `SOURCE` | The one input file. `kath` compiles a single file per run, so a kernel is one source. |
+| `BACKEND` | `amdgpu`, `amdgpu-bin`, `nvidia-ptx`, `tensix`, `rv-elf`, `cpu`, `rv64`, `metal`, `intel-spirv`. |
+| `LANGUAGE` | `CUDA` (default), `HIP`, or `TRITON`. `.hip` files pick HIP up on their own. |
+| `ARCH` | `gfx90a`, `gfx942`, `gfx1030`…`gfx1201`, or `xe-lpg`/`xe-hpg`/`xe-hpc`/`xe2`. |
+| `TT_CHIP` | `blackhole` or `wormhole`. |
+| `OUTPUT` | Output path. Defaults to `<name>` in the build dir with a suffix picked from the backend. |
+| `INCLUDE_DIRECTORIES` | Passed through as `-I`. |
+| `COMPILE_DEFINITIONS` | Passed through as `-D`. |
+| `OPTIONS` | Any other flags, handed to `kath` untouched. |
+| `DEPENDS` | Extra files to rebuild on. |
+| `EXCLUDE_FROM_ALL` | Build only when something asks for it. |
+
+### Two things to know
+
+`kath` writes no depfile, so editing a header the kernel includes will not
+retrigger a build on its own. List those headers in `DEPENDS` if you need it.
+
+The `tensix` backend emits a whole Metalium program rather than one file: the
+compute kernel at `OUTPUT`, plus reader, writer and host sources and their
+ELFs alongside it. Those are declared as byproducts, so `clean` reaches them,
+but `BOOTH_KERNEL_OUTPUT` names only the compute kernel.
diff --git a/examples/cmake/CMakeLists.txt b/examples/cmake/CMakeLists.txt
new file mode 100644
index 0000000..17e9a92
--- /dev/null
+++ b/examples/cmake/CMakeLists.txt
@@ -0,0 +1,28 @@
+# A project consuming an installed Booth. Build it against a prefix you
+# installed to:
+#
+#     make install PREFIX=/tmp/booth
+#     cmake -S examples/cmake -B build -DCMAKE_PREFIX_PATH=/tmp/booth
+#     cmake --build build
+#
+# LANGUAGES NONE because nothing here needs a host compiler; every target is
+# a kernel Booth produces. See docs/cmake.md for linking the cpu and rv64
+# objects into a real executable.
+
+cmake_minimum_required(VERSION 3.10)
+project(booth_cmake_example LANGUAGES NONE)
+
+find_package(Booth 5.1 REQUIRED)
+message(STATUS "Booth ${Booth_VERSION} at ${Booth_EXECUTABLE}")
+
+booth_add_kernel(vadd_ptx    SOURCE vadd.cu BACKEND nvidia-ptx)
+booth_add_kernel(vadd_gfx942 SOURCE vadd.cu BACKEND amdgpu-bin ARCH gfx942)
+booth_add_kernel(vadd_asm    SOURCE vadd.cu BACKEND amdgpu     ARCH gfx1100)
+booth_add_kernel(vadd_tt     SOURCE vadd.cu BACKEND tensix     TT_CHIP blackhole)
+booth_add_kernel(vadd_host   SOURCE vadd.cu BACKEND cpu)
+
+# Every kernel reports where it landed, as a target property and as a
+# variable in this scope.
+get_target_property(_ptx vadd_ptx BOOTH_KERNEL_OUTPUT)
+message(STATUS "PTX at ${_ptx}")
+message(STATUS "host object at ${vadd_host_OUTPUT}")
diff --git a/examples/cmake/vadd.cu b/examples/cmake/vadd.cu
new file mode 100644
index 0000000..2055682
--- /dev/null
+++ b/examples/cmake/vadd.cu
@@ -0,0 +1,5 @@
+__global__ void vadd(const float *a, const float *b, float *c, int n)
+{
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < n) c[i] = a[i] + b[i];
+}
diff --git a/src/amdgpu/amd_rplan.c b/src/amdgpu/amd_rplan.c
index 51db9b4..51ca544 100644
--- a/src/amdgpu/amd_rplan.c
+++ b/src/amdgpu/amd_rplan.c
@@ -135,10 +135,12 @@ void amd_rplan(amd_module_t *A)
         rp_alloc(MF, A->target, &st);
 
         const char *name = A->bir->strings + MF->name;
-        printf("  rplan %s: wave%u, sgp_imp=%u, scratch=%s, "
-               "%u loads, %u stores\n",
-               name, MF->wavefront_size, MF->imp_sgp,
-               st.n_alloca ? "yes" : "no",
-               st.n_loads, st.n_stores);
+        /* stderr, not stdout: --amdgpu writes assembly to stdout when no -o is
+         * given, and a diagnostic in that stream is a .s that won't assemble. */
+        fprintf(stderr, "  rplan %s: wave%u, sgp_imp=%u, scratch=%s, "
+                "%u loads, %u stores\n",
+                name, MF->wavefront_size, MF->imp_sgp,
+                st.n_alloca ? "yes" : "no",
+                st.n_loads, st.n_stores);
     }
 }
diff --git a/src/main.c b/src/main.c
index 91c3639..70a5f3b 100644
--- a/src/main.c
+++ b/src/main.c
@@ -195,7 +195,21 @@ static int run_bir_backends(bir_module_t *bir, const backend_cfg_t *cfg)
             if (cfg->mode_amdgpu_bin)
                 amdgpu_emit_elf(amd,
                     cfg->output_file ? cfg->output_file : "a.hsaco");
-            else
+            else if (cfg->output_file) {
+                /* -o used to be silently ignored here, so a build system had
+                 * no way to ask for the assembly other than capturing stdout.
+                 * No -o still means stdout, which is what anyone piping it
+                 * today already relies on. */
+                FILE *af = fopen(cfg->output_file, "w");
+                if (!af) {
+                    fprintf(stderr, "error: cannot open %s for writing\n",
+                            cfg->output_file);
+                    rc = BC_ERR_IO;
+                } else {
+                    amdgpu_emit_asm(amd, af);
+                    fclose(af);
+                }
+            } else
                 amdgpu_emit_asm(amd, stdout);
         } else {
             if (arc != BC_ERR_VERIFY)
@@ -459,7 +473,8 @@ static void usage(const char *prog)
         "  --xe-hpg      Target Xe-HPG (Alchemist, Battlemage) [default]\n"
         "  --xe-hpc      Target Xe-HPC (Ponte Vecchio)\n"
         "  --xe2         Target Xe2 (Lunar Lake, next-gen Arc)\n"
-        "  -o <file>     Output file (for --amdgpu-bin, --tensix, --nvidia-ptx, --metal, --intel-spirv)\n"
+        "  -o <file>     Output file (for --amdgpu, --amdgpu-bin, --tensix, --nvidia-ptx,\n"
+        "                --metal, --intel-spirv). --amdgpu writes to stdout without it.\n"
         "  --lang <file> Load translated error messages\n"
         "  --version     Print version and exit\n"
         "  --help        Show this message\n"