diff --git a/CHANGELOG.md b/CHANGELOG.md
new file mode 100644
--- /dev/null
+++ b/CHANGELOG.md
@@ -0,0 +1,12 @@
+# Changelog
+
+## 0.1.0.0
+
+First release.
+
+- `runSdlApp` and `runSdlAppReduce` run nano-ui views in an SDL3 window.
+- Text shaped with SDL_ttf and HarfBuzz, drawn from the bundled Inter font
+  or installed fonts looked up by family name, with fallback fonts for
+  scripts the UI font lacks.
+- Clipboard, cursors, display scaling, images, and native file dialogs.
+- The `simd` flag compiles the draw-batch culler with AVX2 on x86-64.
diff --git a/LICENSE b/LICENSE
new file mode 100644
--- /dev/null
+++ b/LICENSE
@@ -0,0 +1,20 @@
+Copyright (c) 2026 goolord
+
+Permission is hereby granted, free of charge, to any person obtaining
+a copy of this software and associated documentation files (the
+"Software"), to deal in the Software without restriction, including
+without limitation the rights to use, copy, modify, merge, publish,
+distribute, sublicense, and/or sell copies of the Software, and to
+permit persons to whom the Software is furnished to do so, subject to
+the following conditions:
+
+The above copyright notice and this permission notice shall be included
+in all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
diff --git a/README.md b/README.md
new file mode 100644
--- /dev/null
+++ b/README.md
@@ -0,0 +1,35 @@
+# nano-ui-sdl
+
+SDL3 window backend for [nano-ui](https://github.com/goolord/nano-ui). Text is
+shaped with SDL_ttf and HarfBuzz, drawn from the bundled Inter font or installed
+fonts looked up by family name, with fallback fonts for other scripts. The
+backend also opens native file dialogs.
+
+```haskell
+{-# LANGUAGE OverloadedStrings #-}
+
+import NanoUI
+import NanoUI.Backend.Sdl (defaultSdlOptions, runSdlApp)
+
+main :: IO ()
+main = runSdlApp defaultSdlOptions (label "Hello")
+```
+
+`runSdlAppReduce` runs a view against a model and an update function, for use
+with `NanoUI.Emit`. `SdlOptions` sets the window, fonts, font size, theme, and
+vsync.
+
+## Running
+
+```sh
+cabal run nano-ui-sdl-anim     # tween and spring animations
+cabal bench nano-ui-sdl-bench  # runFrame and SDL drawing timings
+```
+
+## Requirements
+
+SDL3 and SDL3_ttf 3.2 or later, and pkg-config. The backend is behind the `sdl`
+flag, which is on by default.
+
+On x86-64, `-f simd` compiles the draw-batch culler with AVX2. A binary built
+that way needs an AVX2 CPU.
diff --git a/benchmark/SdlBench.hs b/benchmark/SdlBench.hs
new file mode 100644
--- /dev/null
+++ b/benchmark/SdlBench.hs
@@ -0,0 +1,155 @@
+{-# LANGUAGE OverloadedStrings #-}
+
+module Main (main) where
+
+import Control.Monad (replicateM_, void, when)
+import GHC.IO.Encoding (setLocaleEncoding, utf8)
+import GHC.Conc (getAllocationCounter)
+import NanoUI
+import NanoUI.Context (Context (..))
+import NanoUI.Testing (newPixelContext, runFrame)
+import NanoUI.Backend.Sdl (SdlEnv (..), sdlDrawFrame, syncDisplay, withSdlBench)
+import System.Exit (exitFailure)
+import System.IO (hSetEncoding, stderr, stdout)
+import System.Mem (performGC)
+import Test.Tasty.Bench
+import Text.Printf (printf)
+#if defined(mingw32_HOST_OS)
+import System.Win32 (setConsoleCP, setConsoleOutputCP)
+#endif
+
+benchWindowSize :: Size
+benchWindowSize = Size 800 600
+
+benchInput :: Input
+benchInput =
+  emptyInput
+    { inputWindowSize = benchWindowSize
+    , inputMousePos = V2 400 300
+    , inputMouseDown = True
+    }
+
+smallUi, mediumUi, largeUi :: NanoUI ()
+smallUi =
+  columnWith (gap 8) $ do
+    void (button "OK")
+    label "Hello"
+
+mediumUi =
+  columnWith
+    (grow . gap 8)
+    ( do
+        replicateM_ 12 $
+          gridWith 8 (gap 8) $
+            replicateM_ 8 (void (button "OK"))
+        label "nano-ui SDL bench"
+    )
+
+largeUi =
+  columnWith
+    (grow . gap 6)
+    ( do
+        replicateM_ 20 $
+          gridWith 10 (gap 6) $
+            replicateM_ 10 (void (button "Item"))
+        replicateM_ 8 (label "Status line with a bit of text")
+    )
+
+configureBenchIO :: IO ()
+configureBenchIO = do
+  setLocaleEncoding utf8
+  hSetEncoding stdout utf8
+  hSetEncoding stderr utf8
+#if defined(mingw32_HOST_OS)
+  void $ setConsoleCP 65001
+  void $ setConsoleOutputCP 65001
+#endif
+
+-- | Warm ASCII glyph lookups must not allocate: the atlas UV/bearing record
+-- is cached and shared per font, so a steady-state 'fmGlyph' hit is array
+-- reads and a pointer return. This gate catches reintroducing a
+-- per-character 'GlyphQuad' / 'Just' allocation on the text hot path.
+--
+-- The probe walks a shared 'Char' list rather than 'T.index', because
+-- 'T.index' allocates in this context and would mask the lookup cost.
+glyphLookupAlloc :: Context -> IO Integer
+glyphLookupAlloc ctx = do
+  (fm, _) <- ctxResolveFont ctx 16 WeightNormal FontStyleNormal FontRegular
+  let sample = "The quick brown fox jumps over the lazy dog 0123456789!?.,;:"
+      chars = sample
+      len = length chars
+      lookups = 20000 :: Int
+      step :: Int -> Float -> IO Float
+      step !n !acc =
+        if n <= 0
+           then pure acc
+          else
+            -- Force selection before the indirect glyph call; otherwise the
+            -- benchmark allocates a character-selection thunk per lookup.
+            let !c = chars !! (n `mod` len)
+              in drawGlyph fm c >>= \case
+                  Just gq -> step (n - 1) (acc + gqW gq)
+                  Nothing -> step (n - 1) acc
+  -- Warm every character so every lookup shares a cached 'Maybe'.
+  mapM_ (drawGlyph fm) chars
+  performGC
+  -- The thread allocation counter is current even if this probe never fills
+  -- the nursery. RTSStats.allocated_bytes only catches up at a GC.
+  before <- getAllocationCounter
+  _ <- step lookups 0
+  after <- getAllocationCounter
+  pure (fromIntegral before - fromIntegral after)
+
+-- | Bytes per warm lookup tolerated before the gate trips. The cached path
+-- should be zero; a reintroduced per-hit record would cost tens of bytes.
+glyphLookupAllocBudget :: Double
+glyphLookupAllocBudget = 1.0
+
+glyphLookupGate :: Context -> IO ()
+glyphLookupGate ctx = do
+  bytes <- glyphLookupAlloc ctx
+  let lookups = 20000 :: Int
+      perLookup = fromIntegral bytes / fromIntegral lookups :: Double
+  printf "glyph-lookup: %.3f B/lookup (budget %.1f)\n" perLookup glyphLookupAllocBudget
+  when (perLookup > glyphLookupAllocBudget) $ do
+    putStrLn "FAIL: warm glyph lookups allocate; expected the cached quad to be shared"
+    exitFailure
+
+main :: IO ()
+main = do
+  configureBenchIO
+  ctx0 <- newPixelContext
+  withSdlBench ctx0 $ \ctx sdlEnv -> do
+    (ctx', inp) <- syncDisplay ctx sdlEnv benchInput
+    warmup ctx' sdlEnv inp
+    glyphLookupGate ctx'
+    configureBenchIO
+    defaultMain
+      [ bgroup
+          "ui/runFrame"
+          [ benchRunFrame ctx' inp smallUi "small"
+          , benchRunFrame ctx' inp mediumUi "medium"
+          , benchRunFrame ctx' inp largeUi "large"
+          ]
+      , bgroup
+          "sdl3/draw"
+          [ benchDraw ctx' sdlEnv inp smallUi "small"
+          , benchDraw ctx' sdlEnv inp mediumUi "medium"
+          , benchDraw ctx' sdlEnv inp largeUi "large"
+          ]
+      ]
+
+warmup :: Context -> SdlEnv -> Input -> IO ()
+warmup ctx sdlEnv inp = do
+  void (runFrame ctx inp mediumUi)
+  void (sdlDrawFrame ctx mediumUi sdlEnv inp False)
+  void (runFrame ctx inp mediumUi)
+  void (sdlDrawFrame ctx mediumUi sdlEnv inp False)
+
+benchRunFrame :: Context -> Input -> NanoUI () -> String -> Benchmark
+benchRunFrame ctx inp ui name =
+  bench name $ whnfIO (void . runFrame ctx inp $ ui)
+
+benchDraw :: Context -> SdlEnv -> Input -> NanoUI () -> String -> Benchmark
+benchDraw ctx sdlEnv inp ui name =
+  bench name $ whnfIO (void . sdlDrawFrame ctx ui sdlEnv inp $ False)
diff --git a/cbits/nano_ui_batch.c b/cbits/nano_ui_batch.c
new file mode 100644
--- /dev/null
+++ b/cbits/nano_ui_batch.c
@@ -0,0 +1,176 @@
+#include "nano_ui_batch.h"
+#include "nano_ui_simd.h"
+
+#include <stdlib.h>
+
+struct NanoUiBatch {
+    SDL_Renderer *renderer;
+    const uint8_t *verts;
+    int vert_count;
+    const uint8_t *indices;
+    SDL_Texture *pending_texture;
+    int pending_start;
+    int pending_n;
+};
+
+NanoUiBatch *nano_ui_batch_create(SDL_Renderer *renderer)
+{
+    if (!renderer) {
+        return NULL;
+    }
+    NanoUiBatch *batch = (NanoUiBatch *)calloc(1, sizeof(NanoUiBatch));
+    if (!batch) {
+        return NULL;
+    }
+    batch->renderer = renderer;
+    return batch;
+}
+
+void nano_ui_batch_destroy(NanoUiBatch *batch)
+{
+    if (batch) {
+        nano_ui_batch_flush(batch);
+        free(batch);
+    }
+}
+
+void nano_ui_batch_flush(NanoUiBatch *batch)
+{
+    if (!batch || !batch->renderer || batch->pending_n < 3) {
+        if (batch) {
+            batch->pending_n = 0;
+            batch->pending_start = 0;
+            batch->pending_texture = NULL;
+        }
+        return;
+    }
+    const SDL_Vertex *sdl_verts = (const SDL_Vertex *)batch->verts;
+    const int *idx = (const int *)batch->indices + batch->pending_start;
+    SDL_RenderGeometry(batch->renderer, batch->pending_texture, sdl_verts, batch->vert_count, idx, batch->pending_n);
+    batch->pending_n = 0;
+    batch->pending_start = 0;
+    batch->pending_texture = NULL;
+}
+
+void nano_ui_batch_draw_range(
+    NanoUiBatch *batch,
+    const uint8_t *verts,
+    int vert_count,
+    const uint8_t *indices,
+    int index_start,
+    int index_n,
+    SDL_Texture *texture,
+    int has_damage,
+    float dmg_x,
+    float dmg_y,
+    float dmg_w,
+    float dmg_h)
+{
+    if (!batch || !verts || !indices || vert_count <= 0 || index_n < 3) {
+        return;
+    }
+    if (index_start < 0) {
+        index_start = 0;
+    }
+
+    if (has_damage && dmg_w > 0.f && dmg_h > 0.f && index_n >= 6) {
+        const SDL_Vertex *sdl_verts = (const SDL_Vertex *)verts;
+        const int *idx = (const int *)indices + index_start;
+        float dx0 = dmg_x;
+        float dy0 = dmg_y;
+        float dx1 = dmg_x + dmg_w;
+        float dy1 = dmg_y + dmg_h;
+
+        bool any_visible = false;
+        int q = 0;
+#if defined(NANO_UI_HAS_AVX2)
+        // 8 quads per iteration: the 8-wide AABB test amortizes the
+        // gather + compares, which matters on partial-redraw frames where
+        // long runs of quads fall entirely outside the damage rect.
+        {
+            const __m256 vdx0 = _mm256_set1_ps(dx0);
+            const __m256 vdy0 = _mm256_set1_ps(dy0);
+            const __m256 vdx1 = _mm256_set1_ps(dx1);
+            const __m256 vdy1 = _mm256_set1_ps(dy1);
+            float qx0[8], qy0[8], qx1[8], qy1[8];
+            for (; q + 48 <= index_n; q += 48) {
+                bool all_valid = true;
+                for (int k = 0; k < 8; k++) {
+                    int i0 = idx[q + k * 6];
+                    int i2 = idx[q + k * 6 + 2];
+                    if (i0 < 0 || i0 >= vert_count || i2 < 0 || i2 >= vert_count) {
+                        all_valid = false;
+                        break;
+                    }
+                    float x0 = sdl_verts[i0].position.x;
+                    float y0 = sdl_verts[i0].position.y;
+                    float x1 = sdl_verts[i2].position.x;
+                    float y1 = sdl_verts[i2].position.y;
+                    qx0[k] = x0 < x1 ? x0 : x1;
+                    qx1[k] = x0 > x1 ? x0 : x1;
+                    qy0[k] = y0 < y1 ? y0 : y1;
+                    qy1[k] = y0 > y1 ? y0 : y1;
+                }
+                if (!all_valid) {
+                    any_visible = true;
+                    break;
+                }
+                uint32_t mask = nano_ui_cull_8_quads_avx2(
+                    _mm256_loadu_ps(qx0), _mm256_loadu_ps(qy0),
+                    _mm256_loadu_ps(qx1), _mm256_loadu_ps(qy1),
+                    vdx0, vdy0, vdx1, vdy1);
+                if (mask != 0) {
+                    any_visible = true;
+                    break;
+                }
+            }
+        }
+#endif
+        for (; q + 6 <= index_n; q += 6) {
+            int i0 = idx[q];
+            int i2 = idx[q + 2];
+            if (i0 >= 0 && i0 < vert_count && i2 >= 0 && i2 < vert_count) {
+                float x0 = sdl_verts[i0].position.x;
+                float y0 = sdl_verts[i0].position.y;
+                float x1 = sdl_verts[i2].position.x;
+                float y1 = sdl_verts[i2].position.y;
+                float qx0 = x0 < x1 ? x0 : x1;
+                float qx1 = x0 > x1 ? x0 : x1;
+                float qy0 = y0 < y1 ? y0 : y1;
+                float qy1 = y0 > y1 ? y0 : y1;
+                if (nano_ui_aabb_intersects(qx0, qy0, qx1, qy1, dx0, dy0, dx1, dy1)) {
+                    any_visible = true;
+                    break;
+                }
+            } else {
+                any_visible = true;
+                break;
+            }
+        }
+        if (!any_visible && q > 0) {
+            return;
+        }
+    }
+
+    if (batch->pending_n > 0 &&
+        batch->verts == verts &&
+        batch->indices == indices &&
+        batch->pending_texture == texture &&
+        batch->pending_start + batch->pending_n == index_start)
+    {
+        batch->pending_n += index_n;
+        if (vert_count > batch->vert_count) {
+            batch->vert_count = vert_count;
+        }
+        return;
+    }
+
+    nano_ui_batch_flush(batch);
+
+    batch->verts = verts;
+    batch->vert_count = vert_count;
+    batch->indices = indices;
+    batch->pending_texture = texture;
+    batch->pending_start = index_start;
+    batch->pending_n = index_n;
+}
diff --git a/cbits/nano_ui_batch.h b/cbits/nano_ui_batch.h
new file mode 100644
--- /dev/null
+++ b/cbits/nano_ui_batch.h
@@ -0,0 +1,27 @@
+#ifndef NANO_UI_BATCH_H
+#define NANO_UI_BATCH_H
+
+#include <SDL3/SDL.h>
+#include <stdint.h>
+
+typedef struct NanoUiBatch NanoUiBatch;
+
+NanoUiBatch *nano_ui_batch_create(SDL_Renderer *renderer);
+void nano_ui_batch_destroy(NanoUiBatch *batch);
+void nano_ui_batch_flush(NanoUiBatch *batch);
+
+void nano_ui_batch_draw_range(
+    NanoUiBatch *batch,
+    const uint8_t *verts,
+    int vert_count,
+    const uint8_t *indices,
+    int index_start,
+    int index_n,
+    SDL_Texture *texture,
+    int has_damage,
+    float dmg_x,
+    float dmg_y,
+    float dmg_w,
+    float dmg_h);
+
+#endif
diff --git a/cbits/nano_ui_display.c b/cbits/nano_ui_display.c
new file mode 100644
--- /dev/null
+++ b/cbits/nano_ui_display.c
@@ -0,0 +1,73 @@
+#include <SDL3/SDL.h>
+#include <stddef.h>
+#include <stdbool.h>
+
+int nano_ui_window_refresh_rate(SDL_Window *window)
+{
+    if (!window) {
+        return 0;
+    }
+    SDL_DisplayID id = SDL_GetDisplayForWindow(window);
+    if (!id) {
+        return 0;
+    }
+    const SDL_DisplayMode *current = SDL_GetCurrentDisplayMode(id);
+    if (current && current->refresh_rate > 0) {
+        return current->refresh_rate;
+    }
+    /* Some drivers leave the current mode's refresh at 0 (variable-refresh
+     * panels, compositors that report a base rate). Fall back to the highest
+     * refresh among display modes at the current size so pacing still
+     * targets the panel's real cadence instead of the 60 Hz default. */
+    int best = 0;
+    int cw = current ? current->w : 0;
+    int ch = current ? current->h : 0;
+    if (cw <= 0 || ch <= 0) {
+        SDL_GetWindowSize(window, &cw, &ch);
+    }
+    int count = 0;
+    SDL_DisplayMode **modes = SDL_GetFullscreenDisplayModes(id, &count);
+    for (int i = 0; i < count && modes; i++) {
+        if (modes[i] && modes[i]->w == cw && modes[i]->h == ch && modes[i]->refresh_rate > best) {
+            best = modes[i]->refresh_rate;
+        }
+    }
+    if (modes) {
+        SDL_free(modes);
+    }
+    return best;
+}
+
+typedef void (*nano_ui_resize_cb)(void);
+
+static nano_ui_resize_cb g_resize_cb = NULL;
+
+static bool nano_ui_resize_watch(void *userdata, SDL_Event *event)
+{
+    (void)userdata;
+    if (!g_resize_cb || !event) {
+        return true;
+    }
+    if (event->type == SDL_EVENT_WINDOW_PIXEL_SIZE_CHANGED
+        || event->type == SDL_EVENT_WINDOW_RESIZED
+        || event->type == SDL_EVENT_WINDOW_DISPLAY_SCALE_CHANGED) {
+        g_resize_cb();
+    }
+    return true;
+}
+
+bool nano_ui_install_resize_watch(nano_ui_resize_cb cb)
+{
+    if (!SDL_AddEventWatch(nano_ui_resize_watch, NULL)) {
+        g_resize_cb = NULL;
+        return false;
+    }
+    g_resize_cb = cb;
+    return true;
+}
+
+void nano_ui_remove_resize_watch(void)
+{
+    SDL_RemoveEventWatch(nano_ui_resize_watch, NULL);
+    g_resize_cb = NULL;
+}
diff --git a/cbits/nano_ui_simd.h b/cbits/nano_ui_simd.h
new file mode 100644
--- /dev/null
+++ b/cbits/nano_ui_simd.h
@@ -0,0 +1,44 @@
+#ifndef NANO_UI_SIMD_H
+#define NANO_UI_SIMD_H
+
+#include <stdint.h>
+#include <stdbool.h>
+
+#if defined(__AVX2__)
+#include <immintrin.h>
+#define NANO_UI_HAS_AVX2 1
+#endif
+
+// Returns true if quad AABB [qx0, qy0, qx1, qy1] intersects [dx0, dy0, dx1, dy1]
+static inline bool nano_ui_aabb_intersects(
+    float qx0, float qy0, float qx1, float qy1,
+    float dx0, float dy0, float dx1, float dy1)
+{
+    return !(qx1 < dx0 || qx0 > dx1 || qy1 < dy0 || qy0 > dy1);
+}
+
+#if defined(NANO_UI_HAS_AVX2)
+
+// Tests 8 quad bounding boxes against damage rect [dx0, dy0, dx1, dy1].
+// Returns an 8-bit mask where bit i is 1 if quad i intersects the damage rectangle.
+static inline uint32_t nano_ui_cull_8_quads_avx2(
+    __m256 qx0, __m256 qy0, __m256 qx1, __m256 qy1,
+    __m256 dx0, __m256 dy0, __m256 dx1, __m256 dy1)
+{
+    __m256 outside = _mm256_or_ps(
+        _mm256_cmp_ps(qx1, dx0, _CMP_LT_OQ),
+        _mm256_or_ps(
+            _mm256_cmp_ps(qx0, dx1, _CMP_GT_OQ),
+            _mm256_or_ps(
+                _mm256_cmp_ps(qy1, dy0, _CMP_LT_OQ),
+                _mm256_cmp_ps(qy0, dy1, _CMP_GT_OQ)
+            )
+        )
+    );
+    uint32_t out_mask = (uint32_t)_mm256_movemask_ps(outside);
+    return (~out_mask) & 0xFF;
+}
+
+#endif
+
+#endif // NANO_UI_SIMD_H
diff --git a/cbits/nano_ui_text_atlas.c b/cbits/nano_ui_text_atlas.c
new file mode 100644
--- /dev/null
+++ b/cbits/nano_ui_text_atlas.c
@@ -0,0 +1,205 @@
+#include "nano_ui_text_atlas.h"
+
+#include <SDL3/SDL.h>
+#include <stdlib.h>
+#include <string.h>
+
+enum {
+    NANO_UI_TEXT_ATLAS_SIZE = 2048,
+    NANO_UI_TEXT_ATLAS_PAD = 1,
+    NANO_UI_WHITE_PATCH_SIZE = 4
+};
+
+struct NanoUiTextAtlas {
+    SDL_Renderer *renderer;
+    SDL_Texture *tex;
+    Uint8 *pixels;
+    int w;
+    int h;
+    int x;
+    int y;
+    int row_h;
+};
+
+static void init_white_pixel(NanoUiTextAtlas *atlas)
+{
+    if (!atlas || !atlas->pixels || atlas->w <= 0 || atlas->h <= 0) {
+        return;
+    }
+    for (int y = 0; y < NANO_UI_WHITE_PATCH_SIZE && y < atlas->h; y++) {
+        for (int x = 0; x < NANO_UI_WHITE_PATCH_SIZE && x < atlas->w; x++) {
+            size_t off = ((size_t)y * (size_t)atlas->w + (size_t)x) * 4;
+            atlas->pixels[off + 0] = 255;
+            atlas->pixels[off + 1] = 255;
+            atlas->pixels[off + 2] = 255;
+            atlas->pixels[off + 3] = 255;
+        }
+    }
+}
+
+static bool upload_all(NanoUiTextAtlas *atlas)
+{
+    if (!atlas->tex || !atlas->pixels || atlas->w <= 0 || atlas->h <= 0) {
+        return false;
+    }
+    return SDL_UpdateTexture(atlas->tex, NULL, atlas->pixels, atlas->w * 4);
+}
+
+static bool create_texture(NanoUiTextAtlas *atlas, int w, int h)
+{
+    SDL_Texture *tex =
+        SDL_CreateTexture(atlas->renderer, SDL_PIXELFORMAT_RGBA32, SDL_TEXTUREACCESS_STREAMING, w, h);
+    if (!tex) {
+        return false;
+    }
+    SDL_SetTextureBlendMode(tex, SDL_BLENDMODE_BLEND);
+    /* Bilinear filtering keeps glyph quads smooth when a quad lands off a
+     * whole texel boundary (fractional display scale, sub-pixel pen nudge,
+     * shaped-run placement). NEAREST snaps to the closest texel and makes
+     * scaled/slightly-misaligned text look blocky and pixelated. */
+    SDL_SetTextureScaleMode(tex, SDL_SCALEMODE_LINEAR);
+    Uint8 *px = (Uint8 *)calloc((size_t)w * (size_t)h, 4);
+    if (!px) {
+        SDL_DestroyTexture(tex);
+        return false;
+    }
+    atlas->tex = tex;
+    atlas->pixels = px;
+    atlas->w = w;
+    atlas->h = h;
+    init_white_pixel(atlas);
+    return upload_all(atlas);
+}
+
+static bool slot_for(NanoUiTextAtlas *atlas, int gw, int gh, int *out_x, int *out_y)
+{
+    int pad = NANO_UI_TEXT_ATLAS_PAD;
+    if (gw + 2 * pad > NANO_UI_TEXT_ATLAS_SIZE || gh + 2 * pad > NANO_UI_TEXT_ATLAS_SIZE) {
+        return false;
+    }
+    if (!atlas->tex) {
+        if (!create_texture(atlas, NANO_UI_TEXT_ATLAS_SIZE, NANO_UI_TEXT_ATLAS_SIZE)) {
+            return false;
+        }
+        atlas->x = NANO_UI_WHITE_PATCH_SIZE + pad;
+        atlas->y = pad;
+        atlas->row_h = NANO_UI_WHITE_PATCH_SIZE;
+    }
+    if (atlas->x + gw + pad <= atlas->w && atlas->y + gh + pad <= atlas->h) {
+        *out_x = atlas->x;
+        *out_y = atlas->y;
+        return true;
+    }
+    int next_y = atlas->y + (atlas->row_h > 0 ? atlas->row_h + pad : pad);
+    if (next_y + gh + pad <= atlas->h && gw + 2 * pad <= atlas->w) {
+        atlas->y = next_y;
+        atlas->x = pad;
+        atlas->row_h = 0;
+        *out_x = atlas->x;
+        *out_y = atlas->y;
+        return true;
+    }
+    return false;
+}
+
+/* Glyph surfaces arrive as RGBA32 (glyph_image_to_rgba in nano_ui_ttf.c). */
+static bool blit_surface(NanoUiTextAtlas *atlas, SDL_Surface *surface, int x, int y)
+{
+    const Uint8 *src = (const Uint8 *)surface->pixels;
+    int w = surface->w;
+    int h = surface->h;
+    for (int row = 0; row < h; row++) {
+        Uint8 *dst = atlas->pixels + ((y + row) * atlas->w + x) * 4;
+        memcpy(dst, src + (size_t)row * (size_t)surface->pitch, (size_t)w * 4);
+    }
+    SDL_Rect rect = {x, y, w, h};
+    return SDL_UpdateTexture(atlas->tex, &rect, atlas->pixels + (y * atlas->w + x) * 4, atlas->w * 4);
+}
+
+NanoUiTextAtlas *nano_ui_text_atlas_create(SDL_Renderer *renderer)
+{
+    if (!renderer) {
+        return NULL;
+    }
+    NanoUiTextAtlas *atlas = (NanoUiTextAtlas *)calloc(1, sizeof(NanoUiTextAtlas));
+    if (!atlas) {
+        return NULL;
+    }
+    atlas->renderer = renderer;
+    return atlas;
+}
+
+void nano_ui_text_atlas_destroy(NanoUiTextAtlas *atlas)
+{
+    if (!atlas) {
+        return;
+    }
+    if (atlas->tex) {
+        SDL_DestroyTexture(atlas->tex);
+    }
+    free(atlas->pixels);
+    free(atlas);
+}
+
+SDL_Texture *nano_ui_text_atlas_texture(NanoUiTextAtlas *atlas)
+{
+    return atlas ? atlas->tex : NULL;
+}
+
+bool nano_ui_text_atlas_insert_surface(
+    NanoUiTextAtlas *atlas,
+    SDL_Surface *surface,
+    float *out_x,
+    float *out_y,
+    float *out_w,
+    float *out_h)
+{
+    if (!atlas || !surface) {
+        return false;
+    }
+    int gw = surface->w;
+    int gh = surface->h;
+    if (gw <= 0 || gh <= 0) {
+        return false;
+    }
+    int x = 0;
+    int y = 0;
+    if (!slot_for(atlas, gw, gh, &x, &y)) {
+        return false;
+    }
+    if (!blit_surface(atlas, surface, x, y)) {
+        return false;
+    }
+    atlas->x = x + gw + NANO_UI_TEXT_ATLAS_PAD;
+    if (gh > atlas->row_h) {
+        atlas->row_h = gh;
+    }
+    if (out_x) {
+        *out_x = (float)x;
+    }
+    if (out_y) {
+        *out_y = (float)y;
+    }
+    if (out_w) {
+        *out_w = (float)gw;
+    }
+    if (out_h) {
+        *out_h = (float)gh;
+    }
+    return true;
+}
+
+void nano_ui_text_atlas_reset(NanoUiTextAtlas *atlas)
+{
+    if (!atlas) {
+        return;
+    }
+    atlas->x = NANO_UI_WHITE_PATCH_SIZE + NANO_UI_TEXT_ATLAS_PAD;
+    atlas->y = NANO_UI_TEXT_ATLAS_PAD;
+    atlas->row_h = NANO_UI_WHITE_PATCH_SIZE;
+    if (atlas->pixels && atlas->w > 0 && atlas->h > 0) {
+        memset(atlas->pixels, 0, (size_t)atlas->w * (size_t)atlas->h * 4);
+        init_white_pixel(atlas);
+        upload_all(atlas);
+    }
+}
diff --git a/cbits/nano_ui_text_atlas.h b/cbits/nano_ui_text_atlas.h
new file mode 100644
--- /dev/null
+++ b/cbits/nano_ui_text_atlas.h
@@ -0,0 +1,24 @@
+#ifndef NANO_UI_TEXT_ATLAS_H
+#define NANO_UI_TEXT_ATLAS_H
+
+#include <SDL3/SDL.h>
+#include <stdbool.h>
+
+typedef struct NanoUiTextAtlas NanoUiTextAtlas;
+
+NanoUiTextAtlas *nano_ui_text_atlas_create(SDL_Renderer *renderer);
+void nano_ui_text_atlas_destroy(NanoUiTextAtlas *atlas);
+
+SDL_Texture *nano_ui_text_atlas_texture(NanoUiTextAtlas *atlas);
+
+bool nano_ui_text_atlas_insert_surface(
+    NanoUiTextAtlas *atlas,
+    SDL_Surface *surface,
+    float *out_x,
+    float *out_y,
+    float *out_w,
+    float *out_h);
+
+void nano_ui_text_atlas_reset(NanoUiTextAtlas *atlas);
+
+#endif
diff --git a/cbits/nano_ui_ttf.c b/cbits/nano_ui_ttf.c
new file mode 100644
--- /dev/null
+++ b/cbits/nano_ui_ttf.c
@@ -0,0 +1,475 @@
+#include <SDL3/SDL.h>
+#include <SDL3_ttf/SDL_ttf.h>
+#include <SDL3_ttf/SDL_textengine.h>
+#include <stddef.h>
+#include <stdbool.h>
+#include <stdint.h>
+
+bool nano_ui_ttf_init(void)
+{
+    return TTF_Init();
+}
+
+void nano_ui_ttf_quit(void)
+{
+    TTF_Quit();
+}
+
+/* Kerning is on by default; pin it so shaping survives defaults changing.
+ * Direction and script stay unset: shaping detects them per run, so
+ * right-to-left and complex scripts shape as themselves. Light grid-fitting
+ * snaps stems to whole pixels the way terminals (alacritty/kitty) rasterize,
+ * giving crisp edges instead of soft, grey antialiased outlines. NORMAL keeps
+ * fractional outlines. */
+static TTF_Font *pin_font_rendering(TTF_Font *font)
+{
+    if (font) {
+        TTF_SetFontKerning(font, true);
+        TTF_SetFontHinting(font, TTF_HINTING_LIGHT);
+    }
+    return font;
+}
+
+TTF_Font *nano_ui_ttf_open_font(const char *path, float ptsize)
+{
+    return pin_font_rendering(TTF_OpenFont(path, ptsize));
+}
+
+/* A font at another size over the same source as @font@: the copy shares
+ * its file or memory stream, so any number of sizes read one source. */
+TTF_Font *nano_ui_ttf_copy_font(TTF_Font *font, float ptsize)
+{
+    TTF_Font *copy = font ? TTF_CopyFont(font) : NULL;
+    if (copy) {
+        TTF_SetFontSize(copy, ptsize);
+    }
+    return pin_font_rendering(copy);
+}
+
+void nano_ui_ttf_remove_fallback(TTF_Font *font, TTF_Font *fallback)
+{
+    if (font && fallback) {
+        TTF_RemoveFallbackFont(font, fallback);
+    }
+}
+
+TTF_Font *nano_ui_ttf_open_font_memory(const void *data, size_t size, float ptsize)
+{
+    /* The font reads this stream after the Haskell ByteString callback ends.
+     * Give the stream its own storage, released by its autoclose lifetime. */
+    SDL_IOStream *stream = SDL_IOFromDynamicMem();
+    if (!stream) {
+        return NULL;
+    }
+    if (SDL_WriteIO(stream, data, size) != size ||
+        SDL_SeekIO(stream, 0, SDL_IO_SEEK_SET) < 0) {
+        SDL_CloseIO(stream);
+        return NULL;
+    }
+    SDL_PropertiesID props = SDL_CreateProperties();
+    if (!props) {
+        SDL_CloseIO(stream);
+        return NULL;
+    }
+    SDL_SetPointerProperty(props, TTF_PROP_FONT_CREATE_IOSTREAM_POINTER, stream);
+    SDL_SetBooleanProperty(props, TTF_PROP_FONT_CREATE_IOSTREAM_AUTOCLOSE_BOOLEAN, true);
+    SDL_SetFloatProperty(props, TTF_PROP_FONT_CREATE_SIZE_FLOAT, ptsize);
+    TTF_Font *font = TTF_OpenFontWithProperties(props);
+    SDL_DestroyProperties(props);
+    return pin_font_rendering(font);
+}
+
+void nano_ui_ttf_close_font(TTF_Font *font)
+{
+    if (font) {
+        TTF_CloseFont(font);
+    }
+}
+
+float nano_ui_ttf_line_skip(TTF_Font *font)
+{
+    return (float)TTF_GetFontLineSkip(font);
+}
+
+float nano_ui_ttf_ascent(TTF_Font *font)
+{
+    return (float)TTF_GetFontAscent(font);
+}
+
+float nano_ui_ttf_space_advance(TTF_Font *font)
+{
+    int w = 0;
+    int h = 0;
+    if (TTF_GetStringSize(font, " ", 1, &w, &h) && w > 0) {
+        return (float)w;
+    }
+    int advance = 0;
+    if (TTF_GetGlyphMetrics(font, ' ', NULL, NULL, NULL, NULL, &advance) && advance > 0) {
+        return (float)advance;
+    }
+    return 0.f;
+}
+
+bool nano_ui_ttf_glyph_metrics(
+    TTF_Font *font,
+    Uint32 codepoint,
+    int *out_minx,
+    int *out_maxx,
+    int *out_miny,
+    int *out_maxy,
+    int *out_advance)
+{
+    if (!font) {
+        return false;
+    }
+    int minx = 0, maxx = 0, miny = 0, maxy = 0, advance = 0;
+    if (!TTF_GetGlyphMetrics(font, codepoint, &minx, &maxx, &miny, &maxy, &advance)) {
+        return false;
+    }
+    if (out_minx)   *out_minx   = minx;
+    if (out_maxx)   *out_maxx   = maxx;
+    if (out_miny)   *out_miny   = miny;
+    if (out_maxy)   *out_maxy   = maxy;
+    if (out_advance) *out_advance = advance;
+    return true;
+}
+
+static Uint8 glyph_channel_max(Uint8 r, Uint8 g, Uint8 b)
+{
+    Uint8 m = r;
+    if (g > m) {
+        m = g;
+    }
+    if (b > m) {
+        m = b;
+    }
+    return m;
+}
+
+static void force_white_rgb(SDL_Surface *surf)
+{
+    if (!surf || !surf->pixels || surf->format != SDL_PIXELFORMAT_RGBA32) {
+        return;
+    }
+    Uint8 *base = (Uint8 *)surf->pixels;
+    int pitch = surf->pitch;
+    for (int y = 0; y < surf->h; y++) {
+        Uint8 *row = base + y * pitch;
+        for (int x = 0; x < surf->w; x++) {
+            Uint8 *px = row + x * 4;
+            px[0] = 255;
+            px[1] = 255;
+            px[2] = 255;
+        }
+    }
+}
+
+static void invert_glyph_alpha(SDL_Surface *surf)
+{
+    if (!surf || !surf->pixels || surf->format != SDL_PIXELFORMAT_RGBA32) {
+        return;
+    }
+    Uint8 *base = (Uint8 *)surf->pixels;
+    int pitch = surf->pitch;
+    for (int y = 0; y < surf->h; y++) {
+        Uint8 *row = base + y * pitch;
+        for (int x = 0; x < surf->w; x++) {
+            row[x * 4 + 3] = (Uint8)(255 - row[x * 4 + 3]);
+        }
+    }
+}
+
+static SDL_Surface *glyph_image_to_rgba(SDL_Surface *raw, TTF_ImageType image_type)
+{
+    if (!raw) {
+        return NULL;
+    }
+
+    SDL_Surface *out = SDL_ConvertSurface(raw, SDL_PIXELFORMAT_RGBA32);
+    if (!out) {
+        return NULL;
+    }
+
+    if (image_type == TTF_IMAGE_ALPHA || image_type == TTF_IMAGE_SDF) {
+        /* Spec: color channels are white, alpha is coverage. Transparent
+         * white (a=0, rgb=255) must stay transparent. Do not use luma. */
+        force_white_rgb(out);
+        return out;
+    }
+
+    Uint8 *base = (Uint8 *)out->pixels;
+    int pitch = out->pitch;
+    Uint32 opaque_sum = 0;
+    Uint32 count = 0;
+    for (int y = 0; y < out->h; y++) {
+        Uint8 *row = base + y * pitch;
+        for (int x = 0; x < out->w; x++) {
+            Uint8 *px = row + x * 4;
+            Uint8 luma = glyph_channel_max(px[0], px[1], px[2]);
+            Uint8 a = px[3];
+            Uint8 cov = a > luma ? a : luma;
+            px[0] = 255;
+            px[1] = 255;
+            px[2] = 255;
+            px[3] = cov;
+            opaque_sum += cov;
+            count += 1;
+        }
+    }
+    if (count > 0 && opaque_sum > (255u * count) / 2u) {
+        invert_glyph_alpha(out);
+    }
+    return out;
+}
+
+bool nano_ui_ttf_render_glyph_surface(
+    TTF_Font *font,
+    Uint32 codepoint,
+    SDL_Surface **out_surface)
+{
+    if (!font || !out_surface) {
+        return false;
+    }
+
+    TTF_ImageType image_type = TTF_IMAGE_INVALID;
+    SDL_Surface *raw = TTF_GetGlyphImage(font, codepoint, &image_type);
+    if (!raw) {
+        SDL_Color white = {255, 255, 255, 255};
+        raw = TTF_RenderGlyph_Blended(font, codepoint, white);
+        image_type = TTF_IMAGE_ALPHA;
+    }
+    if (!raw) {
+        return false;
+    }
+
+    SDL_Surface *converted = glyph_image_to_rgba(raw, image_type);
+    SDL_DestroySurface(raw);
+    if (!converted) {
+        return false;
+    }
+
+    *out_surface = converted;
+    return true;
+}
+
+int nano_ui_ttf_get_kerning(TTF_Font *font, Uint32 prev_cp, Uint32 cp)
+{
+    if (!font) {
+        return 0;
+    }
+    int k = 0;
+    TTF_GetGlyphKerning(font, prev_cp, cp, &k);
+    return k;
+}
+
+/* ------------------------------------------------------------------------ */
+/* Shaping                                                                  */
+/* ------------------------------------------------------------------------ */
+
+/* One shaped line, copied out of SDL_ttf's text layout: per glyph
+ * (text offset, glyph index, dst x y w h, src x y w h) and its font; per
+ * cluster (byte offset, byte length, x, width, flags). Coordinates are
+ * pixels from the top left of the line. */
+typedef struct NanoUIShaped {
+    int w;
+    int h;
+    int num_glyphs;
+    int *glyphs;
+    TTF_Font **glyph_fonts;
+    int num_clusters;
+    int *clusters;
+    /* Bytes of the caller's text; anything after is the sentinel. */
+    int text_len;
+} NanoUIShaped;
+
+static bool SDLCALL nano_ui_capture_text(void *userdata, TTF_Text *text)
+{
+    NanoUIShaped *out = (NanoUIShaped *)userdata;
+    TTF_TextData *d = text->internal;
+    int copies = 0;
+    for (int i = 0; i < d->num_ops; i++) {
+        if (d->ops[i].cmd == TTF_DRAW_COMMAND_COPY) {
+            copies++;
+        }
+    }
+    out->w = d->w;
+    out->h = d->h;
+    out->glyphs = (int *)SDL_calloc(copies > 0 ? copies : 1, 10 * sizeof(int));
+    out->glyph_fonts = (TTF_Font **)SDL_calloc(copies > 0 ? copies : 1, sizeof(TTF_Font *));
+    out->clusters = (int *)SDL_calloc(d->num_clusters > 0 ? d->num_clusters : 1, 5 * sizeof(int));
+    if (!out->glyphs || !out->glyph_fonts || !out->clusters) {
+        return false;
+    }
+    /* The sentinel ends a left-to-right line and starts a right-to-left
+     * one, which then shifts back by its advance. */
+    int shift = 0;
+    for (int i = 0; i < d->num_clusters; i++) {
+        TTF_SubString *c = &d->clusters[i];
+        if (c->offset >= out->text_len && c->length > 0) {
+            int advance = 0;
+            TTF_GetGlyphMetrics(d->font, '|', NULL, NULL, NULL, NULL, &advance);
+            out->w = d->w - advance;
+            if ((c->flags & TTF_SUBSTRING_DIRECTION_MASK) == TTF_DIRECTION_RTL) {
+                shift = advance;
+            }
+        }
+    }
+    int g = 0;
+    for (int i = 0; i < d->num_ops; i++) {
+        TTF_DrawOperation *op = &d->ops[i];
+        if (op->cmd != TTF_DRAW_COMMAND_COPY || op->copy.text_offset >= out->text_len) {
+            continue;
+        }
+        int *slot = out->glyphs + g * 10;
+        slot[0] = op->copy.text_offset;
+        slot[1] = (int)op->copy.glyph_index;
+        slot[2] = op->copy.dst.x - shift;
+        /* SDL_ttf places a fallback font's glyphs from that font's own
+         * ascent; align them to the line's font's baseline instead. */
+        slot[3] = op->copy.dst.y - (TTF_GetFontAscent(op->copy.glyph_font) - TTF_GetFontAscent(d->font));
+        slot[4] = op->copy.dst.w;
+        slot[5] = op->copy.dst.h;
+        slot[6] = op->copy.src.x;
+        slot[7] = op->copy.src.y;
+        slot[8] = op->copy.src.w;
+        slot[9] = op->copy.src.h;
+        out->glyph_fonts[g] = op->copy.glyph_font;
+        g++;
+    }
+    out->num_glyphs = g;
+    int k = 0;
+    for (int i = 0; i < d->num_clusters; i++) {
+        TTF_SubString *c = &d->clusters[i];
+        if (c->offset >= out->text_len) {
+            continue;
+        }
+        int *slot = out->clusters + k * 5;
+        slot[0] = c->offset;
+        slot[1] = c->length;
+        slot[2] = c->rect.x - shift;
+        slot[3] = c->rect.w;
+        slot[4] = (int)c->flags;
+        k++;
+    }
+    out->num_clusters = k;
+    d->engine_text = out;
+    return true;
+}
+
+static void SDLCALL nano_ui_release_text(void *userdata, TTF_Text *text)
+{
+    (void)userdata;
+    (void)text;
+}
+
+/* Shape one line with the font and its fallbacks, in a direction (0 lets
+ * SDL_ttf pick, TTF_DIRECTION_LTR or TTF_DIRECTION_RTL). The caller frees
+ * the result with nano_ui_ttf_shaped_free.
+ *
+ * The line is shaped with a '|' after it, whose glyph and cluster are then
+ * dropped. SDL_ttf 3.2 cuts a fallback span that ends the text one character
+ * past its last cluster, losing a vowel sign merged into that cluster, and
+ * trims the width of trailing spaces. */
+bool nano_ui_ttf_shape(TTF_Font *font, const char *text, size_t len, int direction, NanoUIShaped *out)
+{
+    SDL_zerop(out);
+    if (!font || !text || len == 0) {
+        return font != NULL;
+    }
+    TTF_TextEngine engine;
+    SDL_INIT_INTERFACE(&engine);
+    engine.userdata = out;
+    engine.CreateText = nano_ui_capture_text;
+    engine.DestroyText = nano_ui_release_text;
+    out->text_len = (int)len;
+    bool sentinel = TTF_FontHasGlyph(font, '|');
+    char *padded = SDL_malloc(len + 1);
+    if (!padded) {
+        return false;
+    }
+    SDL_memcpy(padded, text, len);
+    padded[len] = '|';
+    TTF_Text *t = TTF_CreateText(&engine, font, padded, sentinel ? len + 1 : len);
+    SDL_free(padded);
+    if (!t) {
+        return false;
+    }
+    if (direction != 0) {
+        TTF_SetTextDirection(t, (TTF_Direction)direction);
+    }
+    /* Laying the text out hands it to the engine, which copies it. */
+    int w = 0, h = 0;
+    bool ok = TTF_GetTextSize(t, &w, &h);
+    if (ok && out->glyphs == NULL) {
+        ok = TTF_UpdateText(t) && out->glyphs != NULL;
+    }
+    TTF_DestroyText(t);
+    return ok;
+}
+
+void nano_ui_ttf_shaped_free(NanoUIShaped *shaped)
+{
+    if (shaped) {
+        SDL_free(shaped->glyphs);
+        SDL_free(shaped->glyph_fonts);
+        SDL_free(shaped->clusters);
+        SDL_zerop(shaped);
+    }
+}
+
+size_t nano_ui_ttf_shaped_size(void)
+{
+    return sizeof(NanoUIShaped);
+}
+
+/* Fields of a shaped line: 0 width, 1 height, 2 glyph count, 3 cluster count. */
+int nano_ui_ttf_shaped_int(const NanoUIShaped *shaped, int field)
+{
+    switch (field) {
+    case 0: return shaped->w;
+    case 1: return shaped->h;
+    case 2: return shaped->num_glyphs;
+    case 3: return shaped->num_clusters;
+    default: return 0;
+    }
+}
+
+/* 0 glyphs, 1 glyph fonts, 2 clusters. */
+void *nano_ui_ttf_shaped_ptr(const NanoUIShaped *shaped, int field)
+{
+    switch (field) {
+    case 0: return shaped->glyphs;
+    case 1: return (void *)shaped->glyph_fonts;
+    case 2: return shaped->clusters;
+    default: return NULL;
+    }
+}
+
+bool nano_ui_ttf_render_glyph_index_surface(TTF_Font *font, Uint32 glyph_index, SDL_Surface **out_surface)
+{
+    if (!font || !out_surface) {
+        return false;
+    }
+    TTF_ImageType image_type = TTF_IMAGE_INVALID;
+    SDL_Surface *raw = TTF_GetGlyphImageForIndex(font, glyph_index, &image_type);
+    if (!raw) {
+        return false;
+    }
+    SDL_Surface *converted = glyph_image_to_rgba(raw, image_type);
+    SDL_DestroySurface(raw);
+    if (!converted) {
+        return false;
+    }
+    *out_surface = converted;
+    return true;
+}
+
+bool nano_ui_ttf_has_glyph(TTF_Font *font, Uint32 ch)
+{
+    return font && TTF_FontHasGlyph(font, ch);
+}
+
+bool nano_ui_ttf_add_fallback(TTF_Font *font, TTF_Font *fallback)
+{
+    return font && fallback && TTF_AddFallbackFont(font, fallback);
+}
diff --git a/data/inter-LICENSE.txt b/data/inter-LICENSE.txt
new file mode 100644
--- /dev/null
+++ b/data/inter-LICENSE.txt
@@ -0,0 +1,92 @@
+Copyright (c) 2016 The Inter Project Authors (https://github.com/rsms/inter)
+
+This Font Software is licensed under the SIL Open Font License, Version 1.1.
+This license is copied below, and is also available with a FAQ at:
+http://scripts.sil.org/OFL
+
+-----------------------------------------------------------
+SIL OPEN FONT LICENSE Version 1.1 - 26 February 2007
+-----------------------------------------------------------
+
+PREAMBLE
+The goals of the Open Font License (OFL) are to stimulate worldwide
+development of collaborative font projects, to support the font creation
+efforts of academic and linguistic communities, and to provide a free and
+open framework in which fonts may be shared and improved in partnership
+with others.
+
+The OFL allows the licensed fonts to be used, studied, modified and
+redistributed freely as long as they are not sold by themselves. The
+fonts, including any derivative works, can be bundled, embedded,
+redistributed and/or sold with any software provided that any reserved
+names are not used by derivative works. The fonts and derivatives,
+however, cannot be released under any other type of license. The
+requirement for fonts to remain under this license does not apply
+to any document created using the fonts or their derivatives.
+
+DEFINITIONS
+"Font Software" refers to the set of files released by the Copyright
+Holder(s) under this license and clearly marked as such. This may
+include source files, build scripts and documentation.
+
+"Reserved Font Name" refers to any names specified as such after the
+copyright statement(s).
+
+"Original Version" refers to the collection of Font Software components as
+distributed by the Copyright Holder(s).
+
+"Modified Version" refers to any derivative made by adding to, deleting,
+or substituting -- in part or in whole -- any of the components of the
+Original Version, by changing formats or by porting the Font Software to a
+new environment.
+
+"Author" refers to any designer, engineer, programmer, technical
+writer or other person who contributed to the Font Software.
+
+PERMISSION AND CONDITIONS
+Permission is hereby granted, free of charge, to any person obtaining
+a copy of the Font Software, to use, study, copy, merge, embed, modify,
+redistribute, and sell modified and unmodified copies of the Font
+Software, subject to the following conditions:
+
+1) Neither the Font Software nor any of its individual components,
+in Original or Modified Versions, may be sold by itself.
+
+2) Original or Modified Versions of the Font Software may be bundled,
+redistributed and/or sold with any software, provided that each copy
+contains the above copyright notice and this license. These can be
+included either as stand-alone text files, human-readable headers or
+in the appropriate machine-readable metadata fields within text or
+binary files as long as those fields can be easily viewed by the user.
+
+3) No Modified Version of the Font Software may use the Reserved Font
+Name(s) unless explicit written permission is granted by the corresponding
+Copyright Holder. This restriction only applies to the primary font name as
+presented to the users.
+
+4) The name(s) of the Copyright Holder(s) or the Author(s) of the Font
+Software shall not be used to promote, endorse or advertise any
+Modified Version, except to acknowledge the contribution(s) of the
+Copyright Holder(s) and the Author(s) or with their explicit written
+permission.
+
+5) The Font Software, modified or unmodified, in part or in whole,
+must be distributed entirely under this license, and must not be
+distributed under any other license. The requirement for fonts to
+remain under this license does not apply to any document created
+using the Font Software.
+
+TERMINATION
+This license becomes null and void if any of the above conditions are
+not met.
+
+DISCLAIMER
+THE FONT SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO ANY WARRANTIES OF
+MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT
+OF COPYRIGHT, PATENT, TRADEMARK, OR OTHER RIGHT. IN NO EVENT SHALL THE
+COPYRIGHT HOLDER BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
+INCLUDING ANY GENERAL, SPECIAL, INDIRECT, INCIDENTAL, OR CONSEQUENTIAL
+DAMAGES, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+FROM, OUT OF THE USE OR INABILITY TO USE THE FONT SOFTWARE OR FROM
+OTHER DEALINGS IN THE FONT SOFTWARE.
diff --git a/data/inter.ttf b/data/inter.ttf
new file mode 100644
Binary files /dev/null and b/data/inter.ttf differ
diff --git a/examples/SdlAnim.hs b/examples/SdlAnim.hs
new file mode 100644
--- /dev/null
+++ b/examples/SdlAnim.hs
@@ -0,0 +1,234 @@
+module Main (main) where
+
+import Control.Monad (forM_, when)
+import NanoUI
+import NanoUI.Backend.Sdl (SdlOptions (..), defaultSdlOptions, runSdlApp)
+import qualified Data.Text as T
+import Text.Printf (printf)
+
+main :: IO ()
+main =
+  runSdlApp
+    defaultSdlOptions
+      { sdlAppTheme = Just benchTheme
+      , sdlAppShouldQuit = \inp -> inputKeysElem KeyEscape (inputKeys inp)
+      }
+    animUi
+
+-- Steenbeck flatbed: putty Formica, black 16mm path, cream frame, ruby rec.
+benchTheme :: Theme
+benchTheme =
+  let plate =
+        Style
+          { styleBg = colorRGBA 214 208 196 255
+          , styleFg = colorRGBA 28 26 22 255
+          , styleBorder = colorRGBA 92 86 76 255
+          , styleBorderWidth = 1
+          , styleCornerRadius = 1
+          , styleHoverBg = colorRGBA 222 216 204 255
+          , styleActiveBg = colorRGBA 196 190 178 255
+          }
+   in defaultTheme
+        { themeWindow = formica
+        , themePanel = plate
+        , themeFloatingWindow = plate
+        , themeButton =
+            Style
+              { styleBg = colorRGBA 196 190 176 255
+              , styleFg = colorRGBA 28 26 22 255
+              , styleBorder = colorRGBA 74 68 58 255
+              , styleBorderWidth = 1
+              , styleCornerRadius = 1
+              , styleHoverBg = colorRGBA 228 222 208 255
+              , styleActiveBg = colorRGBA 168 162 148 255
+              }
+        , themeInput =
+            Style
+              { styleBg = colorRGBA 176 170 156 255
+              , styleFg = colorRGBA 28 26 22 255
+              , styleBorder = colorRGBA 74 68 58 255
+              , styleBorderWidth = 1
+              , styleCornerRadius = 1
+              , styleHoverBg = colorRGBA 186 180 166 255
+              , styleActiveBg = colorRGBA 158 152 138 255
+              }
+        , themeSeparator = colorRGBA 92 86 76 255
+        , themeAccent = ruby
+        , themeMuted = colorRGBA 90 84 74 255
+        , themeOverlayDim = colorRGBA 40 36 30 160
+        }
+
+formica, film, paper, ruby, lamp, punch :: Color
+formica = colorRGBA 184 176 162 255
+film = colorRGBA 18 16 14 255
+paper = colorRGBA 244 236 220 255
+ruby = colorRGBA 154 42 36 255
+lamp = colorRGBA 255 196 92 255
+punch = colorRGBA 232 224 208 255
+
+animUi :: NanoUI ()
+animUi = do
+  (exposed, setExposed) <- useFlag False
+  (rewinding, setRewinding) <- useFlag False
+  (lampOn, setLamp) <- useFlag False
+  (bellowsOpen, setBellows) <- useFlag False
+  (tossed, setTossed) <- useFlag False
+  (stiffSpring, setStiffSpring) <- useFlag False
+  (throwRaw, setThrowRaw) <- useFloat 75
+  tossT <-
+    withKey ("toss" :: String)
+      ( animateTo
+          (Spring (if stiffSpring then presetStiff else presetBouncy))
+          (if tossed then 1 else 0)
+      )
+  lampT <-
+    if lampOn
+      then withKey ("lamp" :: String) (animate (Tween EaseInOutCubic 1.6 0) 0 1)
+      else pure 0
+  wash <-
+    withKey ("wash" :: String) (animateTo (Tween EaseInOutCubic 0.85 0) (if lampOn then 1 else 0))
+  bellowsT <-
+    withKey ("bellows" :: String) (animateTo (Tween EaseOutCubic 0.45 0) (if bellowsOpen then 1 else 0))
+  clock <-
+    if exposed
+      then withKey ("clock" :: String) (animate (Tween EaseLinear 3.2 0) 0 1)
+      else
+        if rewinding
+          then withKey ("clock" :: String) (animateTo (Tween EaseLinear 0.35 0) 0)
+          else pure 0
+  let frames = floor (clock * 128) :: Int
+      footage = T.pack (printf "%d+%02d" (frames `div` 16) (frames `mod` 16))
+      lampGlow = sin (lampT * pi)
+  scrollWith (tight . grow) $
+    columnWith (padAll 28 . gap 18 . fillW) $ do
+      rowWith (tight . gap 10 . alignMid . fillW) $ do
+        heading "16mm"
+        flex
+        withKey ("footage" :: String) (muted footage)
+      rowWith (tight . gap 8 . alignMid . fillW) $ do
+        whenM (button "Expose") (setExposed True >> setRewinding False)
+        whenM (button "Rewind") (setExposed False >> setRewinding True)
+        whenM (button (if lampOn then "Lamp off" else "Lamp on")) (setLamp (not lampOn))
+        flex
+        whenM (button (if tossed then "Catch" else "Toss")) (setTossed (not tossed))
+        whenM (button (if stiffSpring then "Stiff" else "Bouncy")) (setStiffSpring (not stiffSpring))
+      throwSec <- do
+        label "Throw"
+        newThrow <- slider 35 140 throwRaw
+        setThrowRaw newThrow
+        pure (newThrow / 100)
+      cycleThrow <- lockThrow exposed throwSec
+      let cycleLen = pullCycleLen cycleThrow
+      pullPhase <-
+        if exposed
+          then withKey ("pulldown" :: String) (animate (Tween EaseLinear cycleLen 0) 0 1)
+          else pure 0
+      let time = pullPhase * cycleLen
+      tossRail tossT
+      laneTs <- transport exposed rewinding cycleThrow time wash lampGlow
+      when
+        (rewinding && not exposed && abs clock < 0.001 && all settled laneTs)
+        (setRewinding False)
+      panelWith (padXY 12 10 . gap 8 . fixedW (220 + 180 * bellowsT)) $ do
+        rowWith (tight . gap 10 . alignMid . fillW) $ do
+          whenM
+            (button (if bellowsOpen then "Collapse" else "Extend"))
+            (setBellows (not bellowsOpen))
+          flex
+          let iris = 12 + 22 * bellowsT
+              irisCol = lerpColor film (lerpColor paper ruby 0.18) bellowsT
+          box (fixedWH iris iris) irisCol
+
+settled :: Float -> Bool
+settled x = abs x < 0.001
+
+lockThrow :: Bool -> Float -> NanoUI Float
+lockThrow exposed throwSec =
+  withKey ("cycleThrow" :: String) $
+    if exposed
+      then pure throwSec
+      else animateTo (Tween EaseLinear 0.35 0) 0 >> pure throwSec
+
+transport :: Bool -> Bool -> Float -> Float -> Float -> Float -> NanoUI [Float]
+transport exposed rewinding throwSec time wash glow = do
+  let washCol = lerpColor film lamp (wash * (0.45 + 0.55 * glow))
+  columnWith (tight . gap 0 . fillW) $ do
+    perfs
+    withKey ("washTop" :: String) (box (fixedH 10 . fillW) washCol)
+    ts <-
+      columnWith (padXY 0 10 . gap 10 . fillW) $
+        sequence
+          [ lane exposed rewinding throwSec time "Leader" EaseLinear
+          , lane exposed rewinding throwSec time "Gate" EaseInCubic
+          , lane exposed rewinding throwSec time "Shuttle" EaseOutCubic
+          , lane exposed rewinding throwSec time "Reg" EaseInOutCubic
+          , lane exposed rewinding throwSec time "Claw" EaseOutBack
+          , lane exposed rewinding throwSec time "Bezier" (EaseCubicBezier 0.33 0 0.2 1)
+          ]
+    withKey ("washBot" :: String) (box (fixedH 10 . fillW) washCol)
+    perfs
+    pure ts
+
+perfs :: NanoUI ()
+perfs =
+  rowWith (tight . gap 0 . alignMid . fillW) $ do
+    withKey ("perfL" :: String) (box (fixedWH 12 18) film)
+    forM_ [0 .. 16 :: Int] $ \i ->
+      withKey i $ do
+        box (fixedWH 6 18) film
+        box (fixedWH 7 6) punch
+        box (fixedWH 3 18) film
+    withKey ("perfR" :: String) (box (fillW . fixedH 18) film)
+
+pullHoldSec :: Float
+pullHoldSec = 1
+
+pullCycleLen :: Float -> Float
+pullCycleLen throwSec = 3 * pullHoldSec + 2 * throwSec
+
+laneT :: Ease -> Float -> Float -> Float
+laneT ease throwSec time =
+  let hold = pullHoldSec
+      out0 = hold
+      out1 = out0 + throwSec
+      topHold1 = out1 + hold
+      in0 = topHold1
+      in1 = in0 + throwSec
+   in if time < out0
+        then 0
+        else if time < out1
+          then applyEase ease ((time - out0) / throwSec)
+          else if time < in0
+            then 1
+            else if time < in1
+              then 1 - applyEase ease ((time - in0) / throwSec)
+              else 0
+
+tossRail :: Float -> NanoUI ()
+tossRail t = trackRow "Spring" t ruby
+
+lane :: Bool -> Bool -> Float -> Float -> T.Text -> Ease -> NanoUI Float
+lane exposed rewinding throwSec time name ease = do
+  t <-
+    if exposed
+      then pure (laneT ease throwSec time)
+      else
+        if rewinding
+          then withKey name (animateTo (Tween ease throwSec 0) 0)
+          else pure 0
+  trackRow name t paper
+  pure t
+
+trackRow :: T.Text -> Float -> Color -> NanoUI ()
+trackRow name t shuttle =
+  withKey name $
+    rowWith (tight . gap 12 . alignMid . fillW) $ do
+      labelWith (tight . fixedW 72) name
+      columnWith (tight . gap 0 . fillW) $ do
+        let travel = 394.0 :: Float
+        rowWith (tight . alignMid . fillW) $ do
+          box (fixedWH 4 22) film
+          spacer (Fixed (max 0 (t * travel))) Fit
+          box (fixedWH 16 16) shuttle
+          flex
+          box (fixedWH 4 22) film
diff --git a/lib/NanoUI/Backend/Sdl.hs b/lib/NanoUI/Backend/Sdl.hs
new file mode 100644
--- /dev/null
+++ b/lib/NanoUI/Backend/Sdl.hs
@@ -0,0 +1,103 @@
+-- | SDL3 backend: event loop, rendering, and application runners.
+module NanoUI.Backend.Sdl
+  ( RgbaImage (..)
+  , SdlDebugSnapshot (..)
+  , SdlEnv (..)
+  , SdlOptions (..)
+  , askSdlDebug
+  , setSdlUiFont
+  , defaultSdlOptions
+  , FileFilter (..)
+  , FileDialogOptions (..)
+  , FileDialogId (..)
+  , FileDialogResult (..)
+  , defaultFileDialogOptions
+  , openFileDialog
+  , saveFileDialog
+  , openFolderDialog
+  , pollFileDialog
+  , cancelFileDialog
+  , askOpenFileDialog
+  , askSaveFileDialog
+  , askOpenFolderDialog
+  , pollFileDialogUi
+  , NanoUIFont (..)
+  , listFontFamilies
+  , runSdlApp
+  , runSdlAppReduce
+  , sdlDrawFrame
+  , syncDisplay
+  , withSdl
+  , withSdlBench
+  , saveScreenshot
+  ) where
+
+import Control.Monad (unless)
+import Data.Foldable (foldlM)
+import Data.IORef (newIORef)
+import Data.Primitive.SmallArray (SmallArray)
+import Data.Typeable (Typeable)
+import NanoUI (NanoUI)
+import NanoUI.Sdl.Runner (askSdlDebug, drawReduceEff, sdlDrawFrame, setSdlUiFont)
+import NanoUI.Sdl.Session (runSdlSession)
+import NanoUI.Sdl.Debug (SdlDebugSnapshot (..))
+import NanoUI.Sdl.Window (RgbaImage (..), SdlEnv (..), SdlOptions (..), defaultSdlOptions, saveScreenshot, syncDisplay, withSdl, withSdlBench)
+import NanoUI.Sdl.Dialog
+  ( FileDialogId (..)
+  , FileDialogOptions (..)
+  , FileDialogResult (..)
+  , FileFilter (..)
+  , askOpenFileDialog
+  , askOpenFolderDialog
+  , askSaveFileDialog
+  , defaultFileDialogOptions
+  , openFileDialog
+  , openFolderDialog
+  , pollFileDialog
+  , cancelFileDialog
+  , pollFileDialogUi
+  , saveFileDialog
+  )
+import NanoUI.Sdl.NanoUIFont (NanoUIFont (..))
+import NanoUI.Sdl.Font.Search (listFontFamilies)
+import NanoUI.Testing (Context, newPixelContext, registerImage, runEff, withTheme)
+
+runSdlApp :: SdlOptions -> NanoUI () -> IO ()
+runSdlApp options ui = do
+  ctx <- sdlContext options
+  runSdlSession options ctx (const (pure ())) (sdlAppShouldQuit options) $ \c ->
+    sdlDrawFrame c ui
+
+runSdlAppReduce ::
+  (Typeable msg, Eq model) =>
+  SdlOptions ->
+  (msg -> model -> model) ->
+  model ->
+  (model -> NanoUI ()) ->
+  IO ()
+runSdlAppReduce options update model view = do
+  ctx <- sdlContext options
+  modelRef <- newIORef model
+  runSdlSession options ctx (const (pure ())) (sdlAppShouldQuit options) $
+    drawReduceEff runEff update modelRef view
+
+sdlContext :: SdlOptions -> IO Context
+sdlContext options = do
+  ctx0 <- newPixelContext
+  themed <- maybe (pure ctx0) (withTheme ctx0) (sdlAppTheme options)
+  ok <- registerImages themed (sdlAppImages options)
+  unless ok $ fail "registerImage failed"
+  pure themed
+
+registerImages :: Context -> SmallArray RgbaImage -> IO Bool
+registerImages ctx images =
+  foldlM (\ok img -> if ok then registerRgbaImage ctx img else pure False) True images
+
+registerRgbaImage :: Context -> RgbaImage -> IO Bool
+registerRgbaImage ctx img =
+  registerImage
+    ctx
+    (rgbaImageId img)
+    (rgbaImageWidth img)
+    (rgbaImageHeight img)
+    (rgbaImagePixels img)
diff --git a/lib/NanoUI/Sdl/Clipboard.hs b/lib/NanoUI/Sdl/Clipboard.hs
new file mode 100644
--- /dev/null
+++ b/lib/NanoUI/Sdl/Clipboard.hs
@@ -0,0 +1,25 @@
+module NanoUI.Sdl.Clipboard
+  ( withSdlClipboard
+  ) where
+
+import qualified Data.Text as T
+import Data.Text.Foreign (peekCString, withCString)
+import Foreign.Ptr (castPtr, nullPtr)
+import NanoUI.Testing (Context, withClipboard)
+import SDL3.Sys.Bindgen.Runtime.PtrConst qualified as PtrConst
+import SDL3.Sys.Clipboard (getClipboardText, setClipboardText)
+import SDL3.Sys.Stdinc (free)
+
+-- | Route the context's clipboard through SDL (UTF-8 text both ways).
+withSdlClipboard :: Context -> Context
+withSdlClipboard ctx = withClipboard ctx readClipboard writeClipboard
+  where
+    writeClipboard txt = withCString txt (setClipboardText . PtrConst.unsafeFromPtr)
+    readClipboard = do
+      ptr <- getClipboardText
+      if ptr == nullPtr
+        then pure Nothing
+        else do
+          txt <- peekCString ptr
+          free (castPtr ptr)
+          pure (if T.null txt then Nothing else Just txt)
diff --git a/lib/NanoUI/Sdl/Cursor.hs b/lib/NanoUI/Sdl/Cursor.hs
new file mode 100644
--- /dev/null
+++ b/lib/NanoUI/Sdl/Cursor.hs
@@ -0,0 +1,114 @@
+module NanoUI.Sdl.Cursor
+  ( SdlCursors (..)
+  , initCursors
+  , destroyCursors
+  , syncPointerCursor
+  ) where
+
+import Control.Monad (void, when)
+import Data.IORef (IORef, newIORef, readIORef, writeIORef)
+import Foreign.Ptr (Ptr, nullPtr)
+import NanoUI (Input (..))
+import NanoUI.Testing (Context, UiCursorKind (..), uiCursorKind)
+import System.Environment (lookupEnv)
+import System.IO (hPutStrLn, stderr)
+import qualified SDL3.Sys.Bindgen.Mouse as Mouse
+import SDL3.Sys.Mouse
+  ( createSystemCursorSafe
+  , destroyCursorSafe
+  , getDefaultCursorSafe
+  , setCursorSafe
+  )
+
+data SdlCursors = SdlCursors
+  { scDefault :: Ptr Mouse.SDL_Cursor
+  , scPointer :: Ptr Mouse.SDL_Cursor
+  , scText :: Ptr Mouse.SDL_Cursor
+  , scMoveFallback :: Ptr Mouse.SDL_Cursor
+  , scGrab :: Ptr Mouse.SDL_Cursor
+  , scGrabbing :: Ptr Mouse.SDL_Cursor
+  , scNsResize :: Ptr Mouse.SDL_Cursor
+  , scEwResize :: Ptr Mouse.SDL_Cursor
+  , scNwseResize :: Ptr Mouse.SDL_Cursor
+  , scNeswResize :: Ptr Mouse.SDL_Cursor
+  , scCurrent :: IORef UiCursorKind
+  , scTrace :: Bool
+  }
+
+initCursors :: IO SdlCursors
+initCursors = do
+  def <- getDefaultCursorSafe
+  ptr <- createSystemCursorSafe Mouse.SDL_SYSTEM_CURSOR_POINTER
+  text <- createSystemCursorSafe Mouse.SDL_SYSTEM_CURSOR_TEXT
+  moveFallback <- createSystemCursorSafe Mouse.SDL_SYSTEM_CURSOR_MOVE
+  ns <- createSystemCursorSafe Mouse.SDL_SYSTEM_CURSOR_NS_RESIZE
+  ew <- createSystemCursorSafe Mouse.SDL_SYSTEM_CURSOR_EW_RESIZE
+  nwse <- createSystemCursorSafe Mouse.SDL_SYSTEM_CURSOR_NWSE_RESIZE
+  nesw <- createSystemCursorSafe Mouse.SDL_SYSTEM_CURSOR_NESW_RESIZE
+  -- SDL_SYSTEM_CURSOR_GRAB (27) and GRABBING (28) have no bindgen patterns.
+  -- Where SDL or the platform lacks them creation returns NULL, and the move
+  -- cursor stands in.
+  grab <- createSystemCursorSafe (Mouse.SDL_SystemCursor 27)
+  grabbing <- createSystemCursorSafe (Mouse.SDL_SystemCursor 28)
+  current <- newIORef UiCursorDefault
+  -- Debug aid, read once here so cursor changes stay allocation-free:
+  -- NANO_CURSOR_TRACE=1 logs every cursor change to stderr.
+  trace <- (== Just "1") <$> lookupEnv "NANO_CURSOR_TRACE"
+  -- NULL cursors are tolerated: SDL_SetCursor(NULL) selects the platform
+  -- default arrow, which keeps us running on headless/dummy video drivers
+  -- where system cursor shapes are unavailable.
+  pure
+    SdlCursors
+      { scDefault = def
+      , scPointer = ptr
+      , scText = text
+      , scMoveFallback = moveFallback
+      , scGrab = if grab == nullPtr then moveFallback else grab
+      , scGrabbing = if grabbing == nullPtr then moveFallback else grabbing
+      , scNsResize = ns
+      , scEwResize = ew
+      , scNwseResize = nwse
+      , scNeswResize = nesw
+      , scCurrent = current
+      , scTrace = trace
+      }
+
+destroyOwnedCursor :: Ptr Mouse.SDL_Cursor -> Ptr Mouse.SDL_Cursor -> IO ()
+destroyOwnedCursor cur shared =
+  when (cur /= nullPtr && cur /= shared) $
+    destroyCursorSafe cur
+
+destroyCursors :: SdlCursors -> IO ()
+destroyCursors cursors = do
+  let fb = scMoveFallback cursors
+  destroyOwnedCursor (scPointer cursors) nullPtr
+  destroyOwnedCursor (scText cursors) nullPtr
+  destroyOwnedCursor (scNsResize cursors) nullPtr
+  destroyOwnedCursor (scEwResize cursors) nullPtr
+  destroyOwnedCursor (scNwseResize cursors) nullPtr
+  destroyOwnedCursor (scNeswResize cursors) nullPtr
+  destroyOwnedCursor (scGrab cursors) fb
+  destroyOwnedCursor (scGrabbing cursors) fb
+  destroyCursorSafe fb
+
+cursorPtr :: SdlCursors -> UiCursorKind -> Ptr Mouse.SDL_Cursor
+cursorPtr cursors = \case
+  UiCursorDefault -> scDefault cursors
+  UiCursorPointer -> scPointer cursors
+  UiCursorText -> scText cursors
+  UiCursorGrab -> scGrab cursors
+  UiCursorGrabbing -> scGrabbing cursors
+  UiCursorNsResize -> scNsResize cursors
+  UiCursorEwResize -> scEwResize cursors
+  UiCursorNwseResize -> scNwseResize cursors
+  UiCursorNeswResize -> scNeswResize cursors
+
+syncPointerCursor :: SdlCursors -> Context -> Input -> IO ()
+syncPointerCursor cursors ctx inp = do
+  want <- uiCursorKind ctx inp
+  cur <- readIORef (scCurrent cursors)
+  when (want /= cur) $ do
+    when (scTrace cursors) $
+      hPutStrLn stderr ("cursor: " ++ show want ++ " at " ++ show (inputMousePos inp))
+    void $ setCursorSafe (cursorPtr cursors want)
+    writeIORef (scCurrent cursors) want
diff --git a/lib/NanoUI/Sdl/Debug.hs b/lib/NanoUI/Sdl/Debug.hs
new file mode 100644
--- /dev/null
+++ b/lib/NanoUI/Sdl/Debug.hs
@@ -0,0 +1,77 @@
+module NanoUI.Sdl.Debug
+  ( SdlDebugSnapshot (..)
+  , SdlDebugSampler (..)
+  , newSdlDebugSampler
+  , emptySdlDebug
+  , traceFrame
+  ) where
+
+import Data.IORef (IORef, newIORef)
+import Data.Maybe (isJust)
+import Data.Text (Text, unpack)
+import NanoUI.Debug
+  ( CoreDebugSnapshot (..)
+  , DebugSamplerRef
+  , emptyCoreDebugSnapshot
+  , newDebugSampler
+  )
+import System.Environment (lookupEnv)
+import Text.Printf (printf)
+
+data SdlDebugSnapshot = SdlDebugSnapshot
+  { dbgCore      :: !CoreDebugSnapshot
+  , dbgScale     :: !Float
+  , dbgFontPath  :: !FilePath
+  , dbgRenderer  :: !Text
+  , dbgVsync     :: !Bool
+  , dbgRefreshHz :: !Int
+  }
+  deriving (Eq, Show)
+
+-- | The core sampler, the last published snapshot, and whether
+-- NANO_FRAME_TRACE (read once at creation) is set.
+data SdlDebugSampler = SdlDebugSampler
+  { sdsSampler  :: !DebugSamplerRef
+  , sdsSnapshot :: !(IORef SdlDebugSnapshot)
+  , sdsTrace    :: !Bool
+  }
+
+newSdlDebugSampler :: IO SdlDebugSampler
+newSdlDebugSampler =
+  SdlDebugSampler
+    <$> newDebugSampler
+    <*> newIORef emptySdlDebug
+    <*> (isJust <$> lookupEnv "NANO_FRAME_TRACE")
+
+emptySdlDebug :: SdlDebugSnapshot
+emptySdlDebug =
+  SdlDebugSnapshot
+    { dbgCore     = emptyCoreDebugSnapshot
+    , dbgScale    = 1
+    , dbgFontPath = ""
+    , dbgRenderer = ""
+    , dbgVsync    = True
+    , dbgRefreshHz = 0
+    }
+
+-- | Per-refresh timing trace (NANO_FRAME_TRACE). Prints the snapshot's phase
+-- EMAs so live-loop costs can be compared across builds.
+traceFrame :: SdlDebugSnapshot -> IO ()
+traceFrame s =
+  printf
+    "TRACE refreshHz=%3d rend=%s vsync=%d presentFps=%6.0f loopFps=%6.0f frameMs=%6.3f uiMs=%6.3f renderMs=%6.3f presentMs=%6.3f verts=%5d cmds=%2d presents=%d skips=%d\n"
+    (dbgRefreshHz s)
+    (unpack (dbgRenderer s))
+    (if dbgVsync s then 1 else 0 :: Int)
+    (dbgPresentFps c)
+    (dbgLoopFps c)
+    (dbgFrameMs c)
+    (dbgUiMs c)
+    (dbgRenderMs c)
+    (dbgPresentMs c)
+    (dbgVerts c)
+    (dbgCmds c)
+    (dbgPresents c)
+    (dbgSkips c)
+  where
+    c = dbgCore s
diff --git a/lib/NanoUI/Sdl/Dialog.hs b/lib/NanoUI/Sdl/Dialog.hs
new file mode 100644
--- /dev/null
+++ b/lib/NanoUI/Sdl/Dialog.hs
@@ -0,0 +1,285 @@
+-- | SDL3 native file dialogs.
+--
+-- These wrap SDL3's asynchronous dialog API ('SDL_ShowOpenFileDialog',
+-- 'SDL_ShowSaveFileDialog', and 'SDL_ShowOpenFolderDialog') into a
+-- non-blocking, poll-based interface. Launching a dialog returns a
+-- 'FileDialogId' immediately and the app keeps running its normal event loop;
+-- poll the handle on later frames to observe completion.
+--
+-- Threading: SDL3 may invoke the dialog callback on a background thread, so
+-- the callback here does little: it decodes the result, frees the
+-- FFI buffers it owned, wakes the event loop, and records the outcome. All
+-- UI-affecting work ('markDirty', releasing the callback 'FunPtr') is deferred
+-- to the thread that polls the result.
+module NanoUI.Sdl.Dialog
+  ( FileFilter (..)
+  , FileDialogOptions (..)
+  , defaultFileDialogOptions
+  , FileDialogId (..)
+  , FileDialogResult (..)
+  , openFileDialog
+  , saveFileDialog
+  , openFolderDialog
+  , pollFileDialog
+  , cancelFileDialog
+  , clearDialogState
+  , askOpenFileDialog
+  , askSaveFileDialog
+  , askOpenFolderDialog
+  , pollFileDialogUi
+  ) where
+
+import Control.Monad (forM, forM_, unless, void)
+import Data.Int (Int32)
+import Data.IntMap.Strict qualified as IM
+import Data.IORef (atomicModifyIORef', readIORef)
+import Data.Text (Text)
+import qualified Data.Text as T
+import Effectful (Eff, type (:>))
+import Foreign.C.String (CString, newCString, peekCString)
+import Foreign.C.Types (CChar)
+import Foreign.Marshal.Alloc (free)
+import Foreign.Marshal.Array (mallocArray, peekArray0)
+import Foreign.Ptr (FunPtr, Ptr, castFunPtr, castPtr, nullPtr)
+import Foreign.Storable (pokeElemOff)
+import NanoUI.Sdl.Dialog.Types
+  ( DialogCallback
+  , DialogCallbackFunPtr
+  , DialogState (..)
+  , FileDialogId (..)
+  , FileDialogResult (..)
+  , PendingDialog (..)
+  , clearDialogState
+  , drainRetired
+  , retireDialogCallback
+  )
+import NanoUI.Sdl.Display (pushRefreshEvent)
+import NanoUI.Sdl.Window (SdlEnv (..))
+import NanoUI.Testing (Ui, askHost, markDirty, uiIO)
+import SDL3.Sys.Bindgen.Dialog
+  ( SDL_DialogFileCallback (..)
+  , SDL_DialogFileCallback_Aux
+  , SDL_DialogFileFilter (..)
+  )
+import SDL3.Sys.Bindgen.Runtime.PtrConst qualified as PtrConst
+import SDL3.Sys.Dialog
+  ( showOpenFileDialogSafe
+  , showOpenFolderDialogSafe
+  , showSaveFileDialogSafe
+  )
+import SDL3.Sys.Video (raiseWindowSafe, restoreWindowSafe)
+import System.IO (hPutStrLn, stderr)
+
+-- | A file type filter shown in open/save dialogs.
+data FileFilter = FileFilter
+  { filterName :: !Text
+  -- ^ Human-readable label, e.g. @"Haskell source"@.
+  , filterPattern :: !Text
+  -- ^ Semicolon-separated extension list, e.g. @"hs;lhs"@, or @"*"@.
+  }
+  deriving (Eq, Show)
+
+-- | Common options for native file dialogs.
+data FileDialogOptions = FileDialogOptions
+  { dialogFilters :: ![FileFilter]
+  -- ^ File filters (ignored by folder dialogs).
+  , dialogDefaultLocation :: !(Maybe FilePath)
+  -- ^ Starting folder or file.
+  , dialogAllowMany :: !Bool
+  -- ^ Allow selecting more than one entry (ignored by save dialogs).
+  }
+  deriving (Eq, Show)
+
+-- | Sensible defaults: no filters, no default location, single selection.
+defaultFileDialogOptions :: FileDialogOptions
+defaultFileDialogOptions = FileDialogOptions [] Nothing False
+
+-- | Launch an open-file dialog. Returns a handle to poll for completion.
+openFileDialog :: SdlEnv -> FileDialogOptions -> IO FileDialogId
+openFileDialog env opts =
+  launchDialog env OpenDialog (dialogFilters opts) (dialogDefaultLocation opts) (dialogAllowMany opts)
+
+-- | Launch a save-file dialog. Returns a handle to poll for completion.
+saveFileDialog :: SdlEnv -> FileDialogOptions -> IO FileDialogId
+saveFileDialog env opts =
+  launchDialog env SaveDialog (dialogFilters opts) (dialogDefaultLocation opts) (dialogAllowMany opts)
+
+-- | Launch a folder-selection dialog. Returns a handle to poll for completion.
+openFolderDialog :: SdlEnv -> FileDialogOptions -> IO FileDialogId
+openFolderDialog env opts =
+  launchDialog env FolderDialog [] (dialogDefaultLocation opts) (dialogAllowMany opts)
+
+-- | Poll a previously launched dialog without blocking.
+--
+-- Each result is delivered exactly once: after the dialog completes, the
+-- first poll that observes the finished state returns it and forgets the
+-- handle, so later polls return 'FileDialogUnknown'.
+pollFileDialog :: SdlEnv -> FileDialogId -> IO FileDialogResult
+pollFileDialog env (FileDialogId did) = do
+  let st = sdlDialogState env
+  -- Callbacks retired by an earlier poll are now certainly returned.
+  drainRetired st
+  (mcb, result) <-
+    atomicModifyIORef' (dsPending st) $ \pending ->
+      case IM.lookup did pending of
+        Nothing -> (pending, (Nothing, FileDialogUnknown))
+        Just (PendingDialog FileDialogPending _) -> (pending, (Nothing, FileDialogPending))
+        Just (PendingDialog status cb) -> (IM.delete did pending, (Just cb, status))
+  case result of
+    FileDialogPending -> pure ()
+    FileDialogUnknown -> pure ()
+    _ -> do
+      -- Retire the callback for a later poll to free; the callback thread may
+      -- still be unwinding right now, and freeing a running wrapper is unsafe.
+      forM_ mcb (retireDialogCallback st)
+      -- The native dialog stole window focus; reclaim it so the app keeps
+      -- receiving hover/motion/wheel events without an extra click.
+      -- Restoration is a best-effort no-op when the window was never
+      -- minimized (its result is platform-dependent, so it is not a reliable
+      -- failure signal); only a failed raise means the window may still lack
+      -- focus and worth an audible warning.
+      void (restoreWindowSafe (sdlWindow env))
+      raised <- raiseWindowSafe (sdlWindow env)
+      unless raised $
+        hPutStrLn stderr "nano-ui: dialog completed but window raise failed; input may need a click"
+      -- The dialog finished; request a redraw so the caller can reflect the
+      -- result. Safe here: this runs on the polling (UI) thread.
+      ctx <- readIORef (sdlCachedCtx env)
+      markDirty ctx
+  pure result
+
+-- | Stop tracking a dialog handle without waiting for the native dialog to
+-- finish. The handle returns 'FileDialogUnknown' if polled afterwards.
+-- The native dialog keeps running until the user dismisses it; its result is
+-- discarded. If a dialog is abandoned while still open, its small FFI
+-- callback is left to be reclaimed at teardown or process exit.
+cancelFileDialog :: SdlEnv -> FileDialogId -> IO ()
+cancelFileDialog env (FileDialogId did) =
+  atomicModifyIORef' (dsPending (sdlDialogState env)) $ \m -> (IM.delete did m, ())
+
+-- | Open-file dialog, usable from within 'NanoUI' widget code. Returns
+-- 'Nothing' when there is no SDL host to launch a dialog.
+askOpenFileDialog :: Ui :> es => FileDialogOptions -> Eff es (Maybe FileDialogId)
+askOpenFileDialog opts = askHost >>= traverse (uiIO . (`openFileDialog` opts))
+
+-- | Save-file dialog, usable from within 'NanoUI' widget code. Returns
+-- 'Nothing' when there is no SDL host to launch a dialog.
+askSaveFileDialog :: Ui :> es => FileDialogOptions -> Eff es (Maybe FileDialogId)
+askSaveFileDialog opts = askHost >>= traverse (uiIO . (`saveFileDialog` opts))
+
+-- | Folder dialog, usable from within 'NanoUI' widget code. Returns
+-- 'Nothing' when there is no SDL host to launch a dialog.
+askOpenFolderDialog :: Ui :> es => FileDialogOptions -> Eff es (Maybe FileDialogId)
+askOpenFolderDialog opts = askHost >>= traverse (uiIO . (`openFolderDialog` opts))
+
+-- | Poll a dialog from within 'NanoUI' widget code.
+pollFileDialogUi :: Ui :> es => FileDialogId -> Eff es FileDialogResult
+pollFileDialogUi did = do
+  menv <- askHost
+  case menv of
+    Nothing -> pure FileDialogUnknown
+    Just env -> uiIO (pollFileDialog env did)
+
+data DialogKind = OpenDialog | SaveDialog | FolderDialog
+
+launchDialog ::
+  SdlEnv ->
+  DialogKind ->
+  [FileFilter] ->
+  Maybe FilePath ->
+  Bool ->
+  IO FileDialogId
+launchDialog env kind filters mDefault allowMany = do
+  (filtersPtr, filterStrs) <- allocFilters filters
+  (defaultPtr, defaultStr) <- allocDefault mDefault
+  let st = sdlDialogState env
+  -- Free callbacks from dialogs that finished earlier.
+  drainRetired st
+  did <- nextDialogId st
+  rawFp <- mkDialogCallback (onResult did st filterStrs filtersPtr defaultStr)
+  -- Register the handle before showing: the callback may fire before this
+  -- function returns, and it must find its entry.
+  atomicModifyIORef' (dsPending st) $ \m ->
+    (IM.insert did (PendingDialog FileDialogPending rawFp) m, ())
+  let cb = SDL_DialogFileCallback (castFunPtr rawFp :: FunPtr SDL_DialogFileCallback_Aux)
+      filtersConst = PtrConst.unsafeFromPtr filtersPtr
+      nfilters = fromIntegral (length filters)
+  case kind of
+    OpenDialog ->
+      showOpenFileDialogSafe cb nullPtr (sdlWindow env) filtersConst nfilters defaultPtr allowMany
+    SaveDialog ->
+      showSaveFileDialogSafe cb nullPtr (sdlWindow env) filtersConst nfilters defaultPtr
+    FolderDialog ->
+      showOpenFolderDialogSafe cb nullPtr (sdlWindow env) defaultPtr allowMany
+  pure (FileDialogId did)
+
+nextDialogId :: DialogState -> IO Int
+nextDialogId st = atomicModifyIORef' (dsNextId st) $ \n -> (n + 1, n + 1)
+
+-- | SDL invoked the callback: decode the file list, release the FFI buffers
+-- this launch owned, record the outcome, and only then wake the (possibly idle)
+-- event loop. The status must be visible before the wake, or the woken frame
+-- polls 'FileDialogPending', skips, and the result waits for an unrelated
+-- event.
+onResult ::
+  Int ->
+  DialogState ->
+  [CString] ->
+  Ptr SDL_DialogFileFilter ->
+  Maybe CString ->
+  Ptr () ->
+  Ptr () ->
+  Int32 ->
+  IO ()
+onResult did st filterStrs filtersPtr defaultStr _userdata filelistRaw _filterIdx = do
+  paths <- peekFileListRaw filelistRaw
+  let outcome =
+        case paths of
+          Nothing -> FileDialogFailed
+          Just [] -> FileDialogCancelled
+          Just ps -> FileDialogSelected ps
+  forM_ filterStrs free
+  free filtersPtr
+  forM_ defaultStr free
+  atomicModifyIORef' (dsPending st) $ \pending ->
+    (IM.adjust (\pl -> pl {pendingStatus = outcome}) did pending, ())
+  pushRefreshEvent
+
+-- | Decode SDL's null-terminated file list into a plain list of paths.
+--
+-- A null list pointer means SDL hit an error; a null first entry means the
+-- user canceled.
+peekFileListRaw :: Ptr () -> IO (Maybe [FilePath])
+peekFileListRaw filelistRaw
+  | filelistRaw == nullPtr = pure Nothing
+  | otherwise = Just <$> (peekArray0 nullPtr (castPtr filelistRaw) >>= traverse peekCString)
+
+allocFilters :: [FileFilter] -> IO (Ptr SDL_DialogFileFilter, [CString])
+allocFilters [] = pure (nullPtr, [])
+allocFilters fs = do
+  arr <- mallocArray (length fs)
+  strs <-
+    fmap concat $
+      forM (zip [0 ..] fs) $ \(i, FileFilter name pattern_) -> do
+        namePtr <- newCString (T.unpack name)
+        patternPtr <- newCString (T.unpack pattern_)
+        pokeElemOff
+          arr
+          i
+          ( SDL_DialogFileFilter
+              (PtrConst.unsafeFromPtr namePtr)
+              (PtrConst.unsafeFromPtr patternPtr)
+          )
+        pure [namePtr, patternPtr]
+  pure (arr, strs)
+
+allocDefault :: Maybe FilePath -> IO (PtrConst.PtrConst CChar, Maybe CString)
+allocDefault Nothing = pure (PtrConst.unsafeFromPtr nullPtr, Nothing)
+allocDefault (Just path) = do
+  cstr <- newCString path
+  pure (PtrConst.unsafeFromPtr cstr, Just cstr)
+
+-- SDL_DialogFileCallback is `void (*)(void *, const char * const *, int)`,
+-- flattened to `void *` pointers at the FFI boundary.
+foreign import ccall "wrapper"
+  mkDialogCallback :: DialogCallback -> IO DialogCallbackFunPtr
diff --git a/lib/NanoUI/Sdl/Dialog/Types.hs b/lib/NanoUI/Sdl/Dialog/Types.hs
new file mode 100644
--- /dev/null
+++ b/lib/NanoUI/Sdl/Dialog/Types.hs
@@ -0,0 +1,101 @@
+-- | Types and pending-dialog state shared between the SDL dialog backend and
+-- the SDL window lifecycle.
+module NanoUI.Sdl.Dialog.Types
+  ( FileDialogId (..)
+  , FileDialogResult (..)
+  , PendingDialog (..)
+  , DialogCallback
+  , DialogCallbackFunPtr
+  , DialogState (..)
+  , newDialogState
+  , clearDialogState
+  , drainRetired
+  , retireDialogCallback
+  ) where
+
+import Data.Int (Int32)
+import Data.IntMap.Strict (IntMap)
+import Data.IntMap.Strict qualified as IM
+import Data.IORef (IORef, atomicModifyIORef', newIORef, writeIORef)
+import Foreign.Ptr (FunPtr, Ptr, freeHaskellFunPtr)
+
+-- | Opaque handle returned by a non-blocking dialog launch. @0@ is never a
+-- valid handle.
+newtype FileDialogId = FileDialogId Int
+  deriving (Eq, Ord, Show)
+
+-- | Lifecycle state of a launched file dialog.
+data FileDialogResult
+  = FileDialogPending
+  -- ^ Still waiting for the user.
+  | FileDialogCancelled
+  -- ^ The user dismissed the dialog without choosing.
+  | FileDialogFailed
+  -- ^ SDL reported an error.
+  | FileDialogSelected [FilePath]
+  -- ^ The user chose one or more paths.
+  | FileDialogUnknown
+  -- ^ No dialog with this handle is being tracked. A handle becomes unknown
+  -- once its result has been delivered and consumed by 'pollFileDialog', or
+  -- after the dialog was abandoned via 'cancelFileDialog'. Never poll a
+  -- handle that returns 'FileDialogUnknown' again.
+  deriving (Eq, Show)
+
+-- | Shape of the SDL3 dialog callback, flattened to 'Ptr' at the FFI
+-- boundary.
+type DialogCallback = Ptr () -> Ptr () -> Int32 -> IO ()
+
+-- | A marshalled 'DialogCallback' allocated once per dialog launch.
+type DialogCallbackFunPtr = FunPtr DialogCallback
+
+-- | A tracked dialog: its current status plus the FFI callback that owns its
+-- completion. The callback is released only after the dialog completes and a
+-- poll consumes the result, so it is never freed while SDL could still invoke
+-- it.
+data PendingDialog = PendingDialog
+  { pendingStatus :: !FileDialogResult
+  , pendingCallback :: !DialogCallbackFunPtr
+  }
+
+-- | Pending dialogs, keyed by 'FileDialogId'.
+data DialogState = DialogState
+  { dsNextId :: !(IORef Int)
+  , dsPending :: !(IORef (IntMap PendingDialog))
+  , dsRetiredCur :: !(IORef [DialogCallbackFunPtr])
+  , dsRetiredPrev :: !(IORef [DialogCallbackFunPtr])
+  }
+
+-- | Create an empty dialog state.
+newDialogState :: IO DialogState
+newDialogState =
+  DialogState
+    <$> newIORef 0
+    <*> newIORef IM.empty
+    <*> newIORef []
+    <*> newIORef []
+
+-- | Forget every pending dialog. Used during SDL teardown: dialogs still
+-- open on the OS side keep running and their callbacks are left to the
+-- process, but all handles become 'FileDialogUnknown'. Entries still in the
+-- current retirement batch are not freed here: SDL may still be
+-- unwinding their wrappers during teardown; they leak to process exit.
+clearDialogState :: DialogState -> IO ()
+clearDialogState st = do
+  writeIORef (dsPending st) IM.empty
+  drainRetired st
+
+-- | Retire a consumed dialog callback for freeing on a later poll.
+retireDialogCallback :: DialogState -> DialogCallbackFunPtr -> IO ()
+retireDialogCallback st cb =
+  atomicModifyIORef' (dsRetiredCur st) (\cbs -> (cb : cbs, ()))
+
+-- | Free callback 'FunPtr's retired before the previous poll. Retiring parks
+-- them for one full poll first so the dialog callback thread has certainly
+-- returned before 'freeHaskellFunPtr' runs (freeing a wrapper while it
+-- executes is unsafe).
+drainRetired :: DialogState -> IO ()
+drainRetired st = do
+  cbs <- atomicModifyIORef' (dsRetiredPrev st) (\cbs -> ([], cbs))
+  mapM_ freeHaskellFunPtr cbs
+  cur <- atomicModifyIORef' (dsRetiredCur st) (\cur -> ([], cur))
+  writeIORef (dsRetiredPrev st) cur
diff --git a/lib/NanoUI/Sdl/Display.hs b/lib/NanoUI/Sdl/Display.hs
new file mode 100644
--- /dev/null
+++ b/lib/NanoUI/Sdl/Display.hs
@@ -0,0 +1,124 @@
+{-# LANGUAGE OverloadedRecordDot #-}
+
+module NanoUI.Sdl.Display
+  ( defaultFontSize
+  , queryWindowDisplayScale
+  , queryWindowRefreshHz
+  , queryWindowLogicalSize
+  , queryMouseWindowPos
+  , installResizeWatch
+  , refreshEventType
+  , initRefreshEvent
+  , pushRefreshEvent
+  ) where
+
+import Control.Monad (unless, void)
+import Data.IORef (IORef, newIORef, readIORef, writeIORef)
+import Foreign.C.Types (CInt (..))
+import Foreign.Marshal.Alloc (alloca, callocBytes)
+import Foreign.Ptr (FunPtr, Ptr, freeHaskellFunPtr)
+import Foreign.Storable (peek, poke, sizeOf)
+import Data.Word (Word32)
+import NanoUI (Size (..), V2 (..))
+import SDL3.Sys.Bindgen.Events (SDL_Event)
+import SDL3.Sys.Bindgen.Stdinc (Uint32 (..))
+import SDL3.Sys.Bindgen.Video (SDL_Window)
+import SDL3.Sys.Events (pushEvent, registerEvents)
+import SDL3.Sys.Mouse (getMouseState)
+import SDL3.Sys.Video (getWindowDisplayScale, getWindowSize)
+import System.IO.Unsafe (unsafePerformIO)
+
+defaultFontSize :: Float
+defaultFontSize = 16
+
+queryWindowDisplayScale :: Ptr SDL_Window -> IO Float
+queryWindowDisplayScale win = do
+  s <- getWindowDisplayScale win
+  pure (if s > 0 then s else 1)
+
+-- | Vertical refresh rate of the window's current display mode, in Hz
+-- (0 when unavailable).
+queryWindowRefreshHz :: Ptr SDL_Window -> IO Int
+queryWindowRefreshHz win = do
+  hz <- windowRefreshRateC win
+  pure (max 0 (fromIntegral hz))
+
+-- | Window size in window (logical) coordinates; 0x0 when SDL cannot say.
+-- SDL_GetWindowSize already returns the window-coordinate size, not pixels.
+-- Dividing by the display scale would shrink the logical size on DPI-scaled
+-- displays, making the retained framebuffer too small.
+queryWindowLogicalSize :: Ptr SDL_Window -> IO Size
+queryWindowLogicalSize win =
+  alloca $ \wp ->
+    alloca $ \hp -> do
+      ok <- getWindowSize win wp hp
+      if ok
+        then do
+          w <- peek wp
+          h <- peek hp
+          pure (Size (fromIntegral w) (fromIntegral h))
+        else pure (Size 0 0)
+
+-- | Pointer position relative to the window with mouse focus, in window
+-- coordinates. Uses 'SDL_GetMouseState' rather than the global pointer +
+-- window position: the latter is unreliable on Wayland (window position is not
+-- exposed) and breaks hover/wheel targeting.
+queryMouseWindowPos :: IO V2
+queryMouseWindowPos =
+  alloca $ \xp ->
+    alloca $ \yp -> do
+      void (getMouseState xp yp)
+      x <- peek xp
+      y <- peek yp
+      pure (V2 (realToFrac x) (realToFrac y))
+
+-- Windows runs a modal loop while the user drags the border, so the app
+-- event watch does not run. SDL still delivers resize events to this watch.
+installResizeWatch :: IO () -> IO (IO ())
+installResizeWatch act = do
+  fp <- mkResizeCb act
+  ok <- installResizeWatchC fp
+  unless ok $ fail "SDL_AddEventWatch failed"
+  pure $ do
+    removeResizeWatchC
+    freeHaskellFunPtr fp
+
+-- | The user event type that wakes the event loop, registered once per
+-- process by 'initRefreshEvent'; 0 until then.
+{-# NOINLINE refreshEventType #-}
+refreshEventType :: IORef Word32
+refreshEventType = unsafePerformIO (newIORef 0)
+
+-- | The event 'pushRefreshEvent' sends, filled in once by 'initRefreshEvent'.
+-- The core wakes the loop on every 'markDirty', so a push must not allocate.
+{-# NOINLINE refreshEvent #-}
+refreshEvent :: Ptr SDL_Event
+refreshEvent = unsafePerformIO (callocBytes (sizeOf (undefined :: SDL_Event)))
+
+initRefreshEvent :: IO Bool
+initRefreshEvent = do
+  registered <- readIORef refreshEventType
+  if registered /= 0
+    then pure True
+    else do
+      ty <- registerEvents 1
+      poke refreshEvent.type' (Uint32 ty)
+      writeIORef refreshEventType ty
+      pure (ty /= 0)
+
+pushRefreshEvent :: IO ()
+pushRefreshEvent = do
+  ty <- readIORef refreshEventType
+  unless (ty == 0) $ void (pushEvent refreshEvent)
+
+foreign import ccall unsafe "nano_ui_window_refresh_rate"
+  windowRefreshRateC :: Ptr SDL_Window -> IO CInt
+
+foreign import ccall "wrapper"
+  mkResizeCb :: IO () -> IO (FunPtr (IO ()))
+
+foreign import ccall safe "nano_ui_install_resize_watch"
+  installResizeWatchC :: FunPtr (IO ()) -> IO Bool
+
+foreign import ccall safe "nano_ui_remove_resize_watch"
+  removeResizeWatchC :: IO ()
diff --git a/lib/NanoUI/Sdl/Font.hs b/lib/NanoUI/Sdl/Font.hs
new file mode 100644
--- /dev/null
+++ b/lib/NanoUI/Sdl/Font.hs
@@ -0,0 +1,1446 @@
+module NanoUI.Sdl.Font
+  ( FontSource (..)
+  , GlyphAtlas
+  , withTtf
+  , fontSourceLabel
+  , newGlyphAtlas
+  , destroyGlyphAtlas
+  , prepareGlyphAtlasForFrame
+  , takeGlyphAtlasResetFlag
+  , glyphAtlasTexture
+  , SdlFontCache
+  , newSdlFontCache
+  , destroySdlFontCache
+  , reloadSdlFontCache
+  , sdlFontCacheSource
+  , withSdlFontCache
+  ) where
+
+import Control.Exception (SomeException, bracket, catch, throwIO)
+import Control.Monad (forM, forM_, unless, void, when)
+import Data.Bits ((.&.), (.|.), shiftL)
+import Foreign.Marshal.Alloc (alloca, allocaBytes)
+import Foreign.Marshal.Array (advancePtr, allocaArray)
+import Data.Char (isPrint, isSpace, ord)
+import NanoUI.Bidi (BidiRun (..), bidiRuns, needsBidi)
+import NanoUI.Sdl.Font.Search (searchFontFamilies)
+import System.IO.Unsafe (unsafePerformIO)
+import Data.ByteString (ByteString)
+import Data.ByteString.Unsafe (unsafeUseAsCStringLen)
+import Data.IORef (IORef, atomicModifyIORef', modifyIORef', newIORef, readIORef, writeIORef)
+import qualified Data.HashMap.Strict as HM
+import Data.Hashable (Hashable (..))
+import qualified Data.IntSet as IS
+import Data.Primitive.SmallArray
+  ( SmallArray
+  , indexSmallArray
+  , newSmallArray
+  , readSmallArray
+  , sizeofSmallArray
+  , smallArrayFromList
+  , writeSmallArray
+  )
+import Data.Primitive.PrimArray (PrimArray, indexPrimArray, newPrimArray, primArrayFromList, readPrimArray, setPrimArray, sizeofPrimArray, unsafeFreezePrimArray, writePrimArray)
+import Data.Int (Int32)
+import Data.Word (Word64)
+import Data.Text (Text)
+import qualified Data.Sequence as Seq
+import qualified Data.Text as T
+import qualified Data.ByteString.Short as SBS
+import qualified GHC.Foreign as GHC
+import GHC.IO.Encoding (getFileSystemEncoding)
+import Data.Text.Unsafe (lengthWord8)
+import Foreign.C.String (CString, withCString)
+import Foreign.C.Types (CFloat (..), CInt (..), CSize (..), CUInt (..))
+import Foreign.Ptr (IntPtr (..), Ptr, castPtr, intPtrToPtr, nullPtr, plusPtr, ptrToIntPtr)
+import Foreign.Storable (peek, peekElemOff, poke, sizeOf)
+import Data.Unique (hashUnique, newUnique)
+import qualified Data.ByteString as BS
+import System.Directory (getTemporaryDirectory, removeFile)
+import System.IO (hClose, openTempFile)
+import NanoUI
+  ( FontMetrics (..)
+  , FontBackend (..)
+  , FontStyle (..)
+  , FontVariant (..)
+  , FontWeight (..)
+  , GlyphQuad (..)
+  , ShapedGlyphs (..)
+  , ShapedText (..)
+  , monospaceMetrics
+  )
+import NanoUI.Testing
+  ( Context
+  , withExternalText
+  , withFontMetrics
+  , withFontResolver
+  , withMeasureText
+  , withMonoFontMetrics
+  , wrapMeasureCache
+  )
+import SDL3.Sys.Bindgen.Render (SDL_Renderer, SDL_Texture)
+import qualified Data.IntMap.Strict as IM
+import qualified Data.Text.Foreign as TF
+
+data SdlFont = SdlFont
+  { sfId :: !Word64
+  , sfFont :: Ptr ()
+  , sfLineSkip :: Float
+  , sfAscent :: Float
+  , sfSpaceAdvance :: Float
+  , sfTempPath :: !(Maybe FilePath)
+  , sfAlive :: !(IORef Bool)
+  , sfPointSize :: !Float
+  -- ^ The size the font was opened at, which its fallbacks open at too.
+  , sfFallbacks :: !(IORef (IM.IntMap SdlFont))
+  -- ^ Fonts shaping falls back to for characters this one lacks, by their
+  -- place in 'coverageFamilies': each is attached, at this font's size, the
+  -- first time a text needs a character only it covers.
+  }
+
+data FontSource
+  = FontFromPath !FilePath
+  | FontFromMemory !ByteString !FilePath
+  deriving (Show)
+
+fontSourceLabel :: FontSource -> FilePath
+fontSourceLabel (FontFromPath p) = p
+fontSourceLabel (FontFromMemory _ label) = label
+
+-- | Per-glyph atlas slot. UVs are normalised to [0,1] within the atlas texture.
+data GlyphSlot = GlyphSlot
+  { gsW :: {-# UNPACK #-} !Float -- pixel width of glyph image
+  , gsH :: {-# UNPACK #-} !Float -- pixel height of glyph image
+  , gsU0 :: {-# UNPACK #-} !Float
+  , gsV0 :: {-# UNPACK #-} !Float
+  , gsU1 :: {-# UNPACK #-} !Float
+  , gsV1 :: {-# UNPACK #-} !Float
+  , gsOffX :: {-# UNPACK #-} !Float -- bearing x (pixels, at font scale)
+  , gsOffY :: {-# UNPACK #-} !Float -- bearing y (pixels, at font scale)
+  , gsAdvX :: {-# UNPACK #-} !Float -- horizontal advance (pixels, at font scale)
+  }
+
+data GlyphAtlas = GlyphAtlas
+  { gaAtlas :: !(Ptr ())
+  , -- | Glyph slots by font id, then codepoint. Closing a font drops its inner
+    -- map instead of scanning every glyph.
+    gaEntries :: !(IORef (IM.IntMap (IM.IntMap (Maybe GlyphSlot))))
+  , gaEpoch :: !(IORef Word64)
+  , -- | An insertion failed (atlas out of space) during the last frame; the
+    -- atlas must be reset at the next frame start, before any quad is
+    -- recorded, so the reset can never wipe the texture underneath
+    -- already-recorded text.
+    gaNeedsReset :: !(IORef Bool)
+  , -- | The atlas was reset (or ran out of space) since the flag was last
+    -- cleared at frame start. Observed by the runner after the UI pass: a
+    -- set flag means the frame being built holds stale-UV or unplaceable
+    -- text quads and must not be presented.
+    gaResetFlag :: !(IORef Bool)
+  , -- | Glyph slots by font id, then glyph index: what shaped text draws.
+    gaIndexEntries :: !(IORef (IM.IntMap (IM.IntMap (Maybe GlyphSlot))))
+  , -- | Actions to run after every reset (re-warming the base fonts).
+    gaRewarmHooks :: !(IORef [IO ()])
+  , gaAlive :: !(IORef Bool)
+  }
+
+newFontId :: IO Word64
+newFontId = fromIntegral . hashUnique <$> newUnique
+
+-- Backend effects are confined to the owning SDL thread. Retained snapshots
+-- may outlive a window or a cache entry, but must never query a freed handle.
+ensureFontAlive :: SdlFont -> IO ()
+ensureFontAlive sf = do
+  alive <- readIORef (sfAlive sf)
+  unless alive (fail "font backend used after closeFont")
+
+ensureAtlasAlive :: GlyphAtlas -> IO ()
+ensureAtlasAlive ga = do
+  alive <- readIORef (gaAlive ga)
+  unless alive (fail "font backend used after destroyGlyphAtlas")
+
+-- | Maximum number of shaped-run cache entries per 'FontMetrics'. Dynamic,
+-- ever-changing text (FPS counters, timers, percentages, mouse positions)
+-- generates unique strings over time; without a bound the run cache (and
+-- the atlas rectangles its renders occupy) would grow without limit. The
+-- cap sits well above a realistic frame's string working set so steady
+-- static text is never evicted (re-rendering an evicted run leaks its old
+-- atlas rectangle); atlas exhaustion itself is recovered by the deferred
+-- reset in 'prepareGlyphAtlasForFrame', which also clears the whole cache.
+runCacheCap :: Int
+runCacheCap = 1024
+
+-- Entry-count limits alone do not bound retained text: edited oversized lines
+-- can otherwise keep thousands of full-document versions alive per font.
+cacheableText :: Text -> Bool
+cacheableText txt = T.compareLength txt 4096 /= GT
+
+-- | A hash map bounded by entry count: 'insertBounded' into a full cache evicts
+-- the first key of 'bcOrder' and returns its value so the owner can release
+-- it. Hashing a text key once beats comparing it at every level of a tree.
+data BoundedCache k v = BoundedCache
+  { bcEntries :: !(HM.HashMap k v)
+  , bcOrder :: !(Seq.Seq k)
+  -- ^ Each key once, oldest first. Its length is the entry count, which
+  -- 'HM.size' would have to count.
+  }
+
+emptyBounded :: BoundedCache k v
+emptyBounded = BoundedCache HM.empty Seq.empty
+
+insertBounded :: Hashable k => Int -> k -> v -> BoundedCache k v -> (BoundedCache k v, Maybe v)
+insertBounded cap k v (BoundedCache m order) =
+  case HM.alterF (\old -> (old, Just v)) k m of
+    (Just _, m') -> (BoundedCache m' order, Nothing)
+    (Nothing, m')
+      | Seq.length order >= cap
+      , victim Seq.:<| rest <- order ->
+          let (evicted, m'') = HM.alterF (\old -> (old, Nothing)) victim m'
+           in (BoundedCache m'' (rest Seq.|> k), evicted)
+      | otherwise -> (BoundedCache m' (order Seq.|> k), Nothing)
+
+-- Native glyph measurements have one representation, shared by metric-only
+-- preparation and atlas placement. Pixel bearings are unscaled here.
+data GlyphMetrics = GlyphMetrics
+  { gmMinX :: !Float
+  , gmMaxX :: !Float
+  , gmMinY :: !Float
+  , gmMaxY :: !Float
+  , gmAdvance :: !Float
+  }
+
+getGlyphMetrics :: SdlFont -> CUInt -> IO (Maybe GlyphMetrics)
+getGlyphMetrics sf cp = allocaArray 5 $ \p -> do
+  ok <-
+    ttfGlyphMetrics
+      (sfFont sf)
+      cp
+      p
+      (p `advancePtr` 1)
+      (p `advancePtr` 2)
+      (p `advancePtr` 3)
+      (p `advancePtr` 4)
+  let
+    metric i = fromIntegral <$> peekElemOff p i
+  if ok
+    then
+      Just
+        <$> (GlyphMetrics <$> metric 0 <*> metric 1 <*> metric 2 <*> metric 3 <*> metric 4)
+    else pure Nothing
+
+getGlyphAdvance :: SdlFont -> CUInt -> IO (Maybe Float)
+getGlyphAdvance sf cp = fmap gmAdvance <$> getGlyphMetrics sf cp
+
+-- Metric-only geometry has no atlas lifetime and never rasterises a surface.
+getGlyphGeometry :: SdlFont -> Float -> Char -> IO (Maybe GlyphQuad)
+getGlyphGeometry sf inv c =
+  fmap (metricsGlyphQuad sf inv) <$> getGlyphMetrics sf (fromIntegral (ord c))
+
+metricsGlyphQuad :: SdlFont -> Float -> GlyphMetrics -> GlyphQuad
+metricsGlyphQuad sf inv metrics =
+  GlyphQuad
+    (gmMinX metrics / inv)
+    ((sfAscent sf - gmMaxY metrics) / inv)
+    ((gmMaxX metrics - gmMinX metrics) / inv)
+    ((gmMaxY metrics - gmMinY metrics) / inv)
+    0
+    0
+    0
+    0
+
+newGlyphAtlas :: Ptr SDL_Renderer -> IO GlyphAtlas
+newGlyphAtlas ren = do
+  atlas <- textAtlasCreate ren
+  when (atlas == nullPtr) $ fail "nano_ui_text_atlas_create failed (glyph)"
+  entries <- newIORef IM.empty
+  indexEntries <- newIORef IM.empty
+  epoch <- newIORef 0
+  needsReset <- newIORef False
+  resetFlag <- newIORef False
+  rewarms <- newIORef []
+  alive <- newIORef True
+  pure
+    GlyphAtlas
+      { gaAtlas = atlas
+      , gaEntries = entries
+      , gaIndexEntries = indexEntries
+      , gaEpoch = epoch
+      , gaNeedsReset = needsReset
+      , gaResetFlag = resetFlag
+      , gaRewarmHooks = rewarms
+      , gaAlive = alive
+      }
+
+destroyGlyphAtlas :: GlyphAtlas -> IO ()
+destroyGlyphAtlas ga = do
+  alive <- atomicModifyIORef' (gaAlive ga) (\open -> (False, open))
+  when alive $ textAtlasDestroy (gaAtlas ga)
+
+-- | Register an action to run after every atlas reset (DPI change, font
+-- switch, exhaustion recovery). 'newSdlFontCache' registers one that re-warms
+-- the base fonts' ASCII glyphs, so the next frame pays no cold glyph misses.
+registerGlyphAtlasRewarm :: GlyphAtlas -> IO () -> IO ()
+registerGlyphAtlasRewarm ga hook = modifyIORef' (gaRewarmHooks ga) (hook :)
+
+resetGlyphAtlas :: GlyphAtlas -> IO ()
+resetGlyphAtlas ga = do
+  modifyIORef' (gaEpoch ga) (+1)
+  writeIORef (gaEntries ga) IM.empty
+  writeIORef (gaIndexEntries ga) IM.empty
+  writeIORef (gaNeedsReset ga) False
+  textAtlasReset (gaAtlas ga)
+  hooks <- readIORef (gaRewarmHooks ga)
+  mapM_ id hooks
+  writeIORef (gaResetFlag ga) True
+
+-- | An atlas insertion failed: the atlas is out of space. The reset is
+-- deferred to the next frame start ('prepareGlyphAtlasForFrame') so quads
+-- already recorded this frame keep sampling valid pixels, and the frame
+-- itself is marked invalid so the runner drops it instead of presenting
+-- text that could not be placed.
+markAtlasExhausted :: GlyphAtlas -> IO ()
+markAtlasExhausted ga = do
+  writeIORef (gaNeedsReset ga) True
+  writeIORef (gaResetFlag ga) True
+
+-- | Frame-start atlas maintenance: reset the atlas if an insertion failed
+-- during the previous frame, then clear the mid-frame reset flag. Must run
+-- before the frame's UI pass records any quads.
+prepareGlyphAtlasForFrame :: GlyphAtlas -> IO ()
+prepareGlyphAtlasForFrame ga = do
+  needs <- readIORef (gaNeedsReset ga)
+  when needs $ resetGlyphAtlas ga
+  writeIORef (gaResetFlag ga) False
+
+-- | Test-and-clear the mid-frame reset flag. 'True' means the atlas was
+-- reset (or ran out of space) while the frame was being built, so quads
+-- recorded before that point may hold stale UVs; the caller must not
+-- present that frame.
+takeGlyphAtlasResetFlag :: GlyphAtlas -> IO Bool
+takeGlyphAtlasResetFlag ga = atomicModifyIORef' (gaResetFlag ga) (\v -> (False, v))
+
+-- | Pre-rasterise printable ASCII into the glyph atlas to avoid cold misses
+-- on the first rendered frame.
+warmGlyphAtlas :: GlyphAtlas -> SdlFont -> IO ()
+warmGlyphAtlas ga sf =
+  mapM_ (\c -> lookupOrInsertGlyph ga sf c) [' ' .. '~']
+
+-- | Look up or insert a glyph into the atlas.  Returns 'Nothing' for
+-- characters that have no glyph (e.g. control characters).
+lookupOrInsertGlyph :: GlyphAtlas -> SdlFont -> Char -> IO (Maybe GlyphSlot)
+lookupOrInsertGlyph ga sf c = do
+  entries <- readIORef (gaEntries ga)
+  case IM.lookup (fromIntegral (sfId sf)) entries >>= IM.lookup (ord c) of
+    Just mSlot -> pure mSlot
+    Nothing -> do
+      let !cp = fromIntegral (ord c) :: CUInt
+      mMetrics <- getGlyphMetrics sf cp
+      mSlot <- case mMetrics of
+        Nothing -> pure Nothing
+        Just metrics ->
+          placeGlyphImage ga (ttfRenderGlyphSurface (sfFont sf) cp) >>= \case
+            Nothing -> pure Nothing
+            Just slot -> do
+              -- TTF_GetGlyphImage is a tight bitmap. Place it with the font
+              -- bearings: pen + minX, lineTop + (ascent - maxY). Do not clamp
+              -- minX; monospace glyphs are often centered (minX > 0).
+              let !placed = slot {gsOffX = gmMinX metrics, gsOffY = sfAscent sf - gmMaxY metrics, gsAdvX = gmAdvance metrics}
+              pure (Just placed)
+      modifyIORef' (gaEntries ga) (IM.insertWith IM.union (fromIntegral (sfId sf)) (IM.singleton (ord c) mSlot))
+      pure mSlot
+
+-- | Look up or insert a glyph by font and glyph index, the way shaped text
+-- names glyphs. Glyphs are keyed by the font's id, which is never reused, and
+-- rendered through its handle.
+lookupOrInsertGlyphIndex :: GlyphAtlas -> Int -> Int -> Int -> IO (Maybe GlyphSlot)
+lookupOrInsertGlyphIndex ga fontKey handle gi = do
+  entries <- readIORef (gaIndexEntries ga)
+  case IM.lookup fontKey entries >>= IM.lookup gi of
+    Just mSlot -> pure mSlot
+    Nothing -> do
+      mSlot <- placeGlyphImage ga (ttfRenderGlyphIndexSurface (intPtrToPtr (IntPtr handle)) (fromIntegral gi))
+      modifyIORef' (gaIndexEntries ga) (IM.insertWith IM.union fontKey (IM.singleton gi mSlot))
+      pure mSlot
+
+-- | Render a glyph image into a surface and copy it into the atlas, as a slot
+-- with no bearings or advance. 'Nothing' when there is no image or no room.
+-- A full atlas is reset at the next frame start (see 'markAtlasExhausted'):
+-- wiping the texture here would leave quads already recorded this frame
+-- sampling blank pixels. The glyph is unavailable for the rest of the frame,
+-- which is dropped.
+placeGlyphImage :: GlyphAtlas -> (Ptr (Ptr ()) -> IO Bool) -> IO (Maybe GlyphSlot)
+placeGlyphImage ga render = do
+  surf <- alloca $ \sp -> do
+    poke sp nullPtr
+    ok <- render sp
+    if ok then peek sp else pure nullPtr
+  if surf == nullPtr
+    then pure Nothing
+    else do
+      mPos <- tryInsert (gaAtlas ga) surf
+      freeSurface surf
+      case mPos of
+        Nothing -> do
+          markAtlasExhausted ga
+          pure Nothing
+        Just (px, py, tw, th) -> do
+          let !slot =
+                GlyphSlot
+                  { gsW = tw
+                  , gsH = th
+                  , gsU0 = px / glyphAtlasSize
+                  , gsV0 = py / glyphAtlasSize
+                  , gsU1 = (px + tw) / glyphAtlasSize
+                  , gsV1 = (py + th) / glyphAtlasSize
+                  , gsOffX = 0
+                  , gsOffY = 0
+                  , gsAdvX = 0
+                  }
+          pure (Just slot)
+
+-- | Width and height of the glyph atlas texture; mirrors
+-- NANO_UI_TEXT_ATLAS_SIZE in nano_ui_text_atlas.c.
+glyphAtlasSize :: Float
+glyphAtlasSize = 2048
+
+-- | A line shaped by SDL_ttf: its layout for measuring and caret placement,
+-- its measured width and height, and per glyph nine numbers (glyph index, destination x y w h,
+-- source x y w h, in raster pixels) with the index of the font that has it
+-- among the line's font and its fallbacks.
+data Shaped = Shaped
+  { shapedText :: !ShapedText
+  , shapedSize :: !(Float, Float)
+  -- ^ Kept whole so measuring a cached line allocates nothing.
+  , _shapedGlyphs :: !(PrimArray Int32)
+  , _shapedFontIndices :: !(PrimArray Int32)
+  , _shapedFonts :: !(SmallArray SdlFont)
+  }
+
+-- | The pieces of a line shaped one at a time, in visual order: character
+-- start, end, and the SDL_ttf direction (0 for a line in one direction).
+-- SDL_ttf keeps a right-to-left text's edge spaces on the side they are
+-- stored, so they are shaped apart and placed on the side they read.
+shapingRuns :: Text -> [(Int, Int, CInt)]
+shapingRuns txt
+  | needsBidi txt = concatMap directed (bidiRuns txt)
+  | otherwise = [(0, T.length txt, 0)]
+  where
+    directed r
+      | not (runRightToLeft r) = [(runStart r, runEnd r, 4)]
+      | otherwise =
+          let run = T.take (runEnd r - runStart r) (T.drop (runStart r) txt)
+              lead = T.length (T.takeWhile isSpace run)
+              trail = T.length (T.takeWhileEnd isSpace run)
+              coreStart = runStart r + lead
+              coreEnd = runEnd r - trail
+           in if coreStart >= coreEnd
+                then [(runStart r, runEnd r, 4)]
+                else
+                  [(coreEnd, runEnd r, 4) | trail > 0]
+                    ++ [(coreStart, coreEnd, 5)]
+                    ++ [(runStart r, coreStart, 4) | lead > 0]
+
+-- | Shape one line with the font and its fallbacks. A line mixing
+-- directions is split into direction runs, each shaped on its own and placed
+-- in visual order, since SDL_ttf shapes a text in one direction. Caret
+-- positions come from the clusters: a cluster's characters share its width,
+-- from its left edge in a left-to-right run and from its right edge in a
+-- right-to-left one.
+shapeLine :: SdlFont -> Float -> Text -> IO Shaped
+shapeLine sf inv txt = do
+  ensureFontAlive sf
+  fallbacks <- IM.elems <$> readIORef (sfFallbacks sf)
+  let fontList = sf : fallbacks
+      n = T.length txt
+      totalBytes = lengthWord8 txt
+      runs = shapingRuns txt
+      pieceCount = length runs
+      ascii = totalBytes == n
+  -- The byte each character starts at, and the character starting at each
+  -- byte (n inside a character and at the end). ASCII needs neither.
+  byteOfChar <- newPrimArray (if ascii then 0 else n + 1)
+  charOfByte <- newPrimArray (if ascii then 0 else totalBytes + 1)
+  let byteAt i = if ascii then pure i else readPrimArray byteOfChar i
+      charAt b = if ascii then pure (min n b) else readPrimArray charOfByte (min totalBytes b)
+      indexChars !i !b t = do
+        writePrimArray byteOfChar i b
+        case T.uncons t of
+          Nothing -> pure ()
+          Just (c, rest) -> do
+            writePrimArray charOfByte b i
+            indexChars (i + 1) (b + utf8Length c) rest
+  unless ascii $ do
+    setPrimArray charOfByte 0 (totalBytes + 1) n
+    indexChars 0 0 txt
+  size <- fromIntegral <$> ttfShapedSize
+  allocaBytes (size * max 1 pieceCount) $ \outs -> do
+    let resultOf p = outs `plusPtr` (p * size)
+    -- Shape every piece first, so the output arrays are sized once.
+    let shapeAll !_ [] = pure ()
+        shapeAll !p ((start, end, dir) : rest) = do
+          withUtf8 (T.take (end - start) (T.drop start txt)) $ \cstr len -> do
+            ok <- ttfShape (sfFont sf) cstr len dir (resultOf p)
+            unless ok $ ttfShapedFree (resultOf p)
+          shapeAll (p + 1) rest
+    shapeAll 0 runs
+    let countGlyphs !p !acc
+          | p >= pieceCount = pure acc
+          | otherwise = do
+              g <- ttfShapedInt (resultOf p) 2
+              countGlyphs (p + 1) (acc + fromIntegral g)
+    glyphCount <- countGlyphs 0 0
+    glyphs <- newPrimArray (glyphCount * 9)
+    fontIndices <- newPrimArray glyphCount
+    -- Caret stops by character in raster pixels, NaN where no cluster
+    -- starts; a later cluster covering a character wins.
+    stops <- newPrimArray n
+    setPrimArray stops 0 n (0 / 0 :: Float)
+    let fillPieces !p !pen !height !g0 !inkEnd !endStop pieces = case pieces of
+          [] -> pure (pen, height, inkEnd, endStop)
+          (start, _, _) : rest -> do
+            let result = resultOf p
+            byteStart <- byteAt start
+            w <- fromIntegral <$> ttfShapedInt result 0
+            h <- fromIntegral <$> ttfShapedInt result 1
+            nGlyphs <- fromIntegral <$> ttfShapedInt result 2
+            nClusters <- fromIntegral <$> ttfShapedInt result 3
+            glyphPtr <- castPtr <$> ttfShapedPtr result 0
+            fontPtr <- castPtr <$> ttfShapedPtr result 1
+            clusterPtr <- castPtr <$> ttfShapedPtr result 2
+            let glyphInt :: Int -> IO Int32
+                glyphInt k = fromIntegral <$> peekElemOff (glyphPtr :: Ptr CInt) k
+                fontIndex ptr = go 0 fontList
+                  where
+                    go !k (f : fs) = if sfFont f == ptr then k else go (k + 1) fs
+                    go !_ [] = 0
+                -- SDL_ttf's ten numbers a glyph start with its text offset,
+                -- which carets take from the clusters instead.
+                copyGlyphs !i !ink
+                  | i >= nGlyphs = pure ink
+                  | otherwise = do
+                      let o = (g0 + i) * 9
+                          field !k
+                            | k > 9 = pure ()
+                            | otherwise = do
+                                v <- glyphInt (i * 10 + k)
+                                writePrimArray glyphs (o + k - 1) (if k == 2 then v + fromIntegral pen else v)
+                                field (k + 1)
+                      field 1
+                      x <- glyphInt (i * 10 + 2)
+                      gw <- glyphInt (i * 10 + 4)
+                      ptr <- peekElemOff (fontPtr :: Ptr (Ptr ())) i
+                      writePrimArray fontIndices (g0 + i) (fontIndex ptr)
+                      copyGlyphs (i + 1) (max ink (fromIntegral x + pen + fromIntegral gw))
+                clusterInt :: Int -> IO Int
+                clusterInt k = fromIntegral <$> peekElemOff (clusterPtr :: Ptr CInt) k
+                -- A cluster's characters share its width, from its left edge
+                -- left to right and from its right edge right to left.
+                placeClusters !i !end
+                  | i >= nClusters = pure end
+                  | otherwise = do
+                      off0 <- clusterInt (i * 5)
+                      len <- clusterInt (i * 5 + 1)
+                      x0 <- clusterInt (i * 5 + 2)
+                      cw <- clusterInt (i * 5 + 3)
+                      flags <- clusterInt (i * 5 + 4)
+                      if len <= 0
+                        then placeClusters (i + 1) end
+                        else do
+                          let off = off0 + byteStart
+                              x = fromIntegral (x0 + pen) :: Float
+                              fw = fromIntegral cw
+                              rtl = flags .&. 0xFF == 5
+                          c0 <- charAt off
+                          c1 <- charAt (off + len)
+                          let k = max 1 (c1 - c0)
+                              stop j
+                                | c0 + j >= n = pure ()
+                                | rtl = writePrimArray stops (c0 + j) (x + fw - fw * fromIntegral j / fromIntegral k)
+                                | otherwise = writePrimArray stops (c0 + j) (x + fw * fromIntegral j / fromIntegral k)
+                          forM_ [0 .. k - 1] stop
+                          let end'
+                                | isNaN end && c1 == n = if rtl then x else x + fw
+                                | otherwise = end
+                          placeClusters (i + 1) end'
+            ink <- if glyphPtr == nullPtr then pure inkEnd else copyGlyphs 0 inkEnd
+            end <- if clusterPtr == nullPtr then pure endStop else placeClusters 0 endStop
+            ttfShapedFree result
+            fillPieces (p + 1) (pen + w) (max height h) (g0 + nGlyphs) ink end rest
+    (total, height, inkEnd, endStop) <- fillPieces 0 0 (0 :: Int) 0 0 (0 / 0) runs
+    carets <- newPrimArray (n + 1)
+    let fillCarets !i !prev
+          | i >= n = pure ()
+          | otherwise = do
+              v <- readPrimArray stops i
+              let v' = if isNaN v then prev else v / inv
+              writePrimArray carets i v'
+              fillCarets (i + 1) v'
+    fillCarets 0 0
+    writePrimArray carets n ((if isNaN endStop then fromIntegral total else endStop) / inv)
+    caretArr <- unsafeFreezePrimArray carets
+    glyphArr <- unsafeFreezePrimArray glyphs
+    indexArr <- unsafeFreezePrimArray fontIndices
+    let !width = fromIntegral total / inv
+        !measured = (width, fromIntegral height / inv)
+    pure (Shaped (ShapedText width (fromIntegral inkEnd / inv) caretArr) measured glyphArr indexArr (smallArrayFromList fontList))
+  where
+    utf8Length c
+      | ord c < 0x80 = 1
+      | ord c < 0x800 = 2
+      | ord c < 0x10000 = 3
+      | otherwise = 4 :: Int
+
+-- | Open the fonts that cover what this one lacks, the first time a text
+-- has a character it cannot draw.
+ensureCoverage :: SdlFont -> Text -> IO ()
+ensureCoverage sf txt =
+  unless (T.all (\c -> ord c < 128) txt) $
+    forM_ (T.unpack txt) $ \c ->
+      when (ord c >= 128 && isPrint c) $ do
+        -- Whether the font or a fallback it already has draws the character.
+        has <- ttfHasGlyph (sfFont sf) (fromIntegral (ord c))
+        unless has $
+          coverageSourceFor c >>= \case
+            Just (source, probe) -> attachFallback sf source probe
+            Nothing -> pure ()
+
+-- | Open the coverage source @source@ at the font's size, sharing the
+-- source's stream with its probe, and attach it. Fallbacks stay in
+-- 'coverageFamilies' order, so a character two of them draw comes from the
+-- one listed first whichever was attached first.
+attachFallback :: SdlFont -> Int -> Ptr () -> IO ()
+attachFallback sf source probe = do
+  attached <- readIORef (sfFallbacks sf)
+  unless (IM.member source attached) $ do
+    copy <- ttfCopyFont probe (realToFrac (sfPointSize sf))
+    unless (copy == nullPtr) $ do
+      fallback <- readSdlFont (sfPointSize sf) Nothing copy
+      let attached' = IM.insert source fallback attached
+      case IM.lookupMax attached of
+        Just (lastSource, _) | lastSource > source -> do
+          forM_ attached $ \f -> ttfRemoveFallback (sfFont sf) (sfFont f)
+          forM_ attached' $ \f -> ttfAddFallback (sfFont sf) (sfFont f)
+        _ -> void (ttfAddFallback (sfFont sf) copy)
+      writeIORef (sfFallbacks sf) attached'
+
+-- | What the process knows about coverage fonts: the installed files, found
+-- once; a probe font for each opened so far (null when it failed to open),
+-- which every size copies from; and the first source drawing each character
+-- asked about (-1 for none).
+data Coverage = Coverage
+  { covSources :: !(SmallArray SBS.ShortByteString)
+  -- ^ Paths as the file system's bytes, kept compactly for the session and
+  -- handed to SDL_ttf as they are.
+  , covProbes :: !(IM.IntMap (Ptr ()))
+  , covChars :: !(IM.IntMap Int)
+  }
+
+{-# NOINLINE coverageRef #-}
+coverageRef :: IORef (Maybe Coverage)
+coverageRef = unsafePerformIO (newIORef Nothing)
+
+-- | The first coverage source, in 'coverageFamilies' order, that draws the
+-- character, and its probe. Probes open only as far down the list as a
+-- search goes, once a session, whatever the number of font sizes.
+coverageSourceFor :: Char -> IO (Maybe (Int, Ptr ()))
+coverageSourceFor c = do
+  cov0 <-
+    readIORef coverageRef >>= \case
+      Just cov -> pure cov
+      Nothing -> do
+        files <- searchFontFamilies coverageFamilies `catch` \(_ :: SomeException) -> pure []
+        -- The file system encoding turns a path back into the bytes it was
+        -- read from, including bytes that are not valid in that encoding.
+        enc <- getFileSystemEncoding
+        sources <- forM files $ \path -> GHC.withCStringLen enc path $ \cstr -> SBS.toShort <$> BS.packCStringLen cstr
+        pure (Coverage (smallArrayFromList sources) IM.empty IM.empty)
+  let cp = ord c
+      probeOf cov i = case IM.lookup i (covProbes cov) of
+        Just probe -> pure (probe, cov)
+        Nothing -> do
+          probe <- SBS.useAsCString (indexSmallArray (covSources cov) i) $ \cpath -> ttfOpenFont cpath 12
+          pure (probe, cov {covProbes = IM.insert i probe (covProbes cov)})
+      search cov i
+        | i >= sizeofSmallArray (covSources cov) = pure (-1, cov)
+        | otherwise = do
+            (probe, cov') <- probeOf cov i
+            has <- if probe == nullPtr then pure False else ttfHasGlyph probe (fromIntegral cp)
+            if has then pure (i, cov') else search cov' (i + 1)
+  (source, cov1) <- case IM.lookup cp (covChars cov0) of
+    Just known -> pure (known, cov0)
+    Nothing -> do
+      (found, cov') <- search cov0 0
+      pure (found, cov' {covChars = IM.insert cp found (covChars cov')})
+  writeIORef coverageRef (Just cov1)
+  pure $ case IM.lookup source (covProbes cov1) of
+    Just probe | source >= 0 -> Just (source, probe)
+    _ -> Nothing
+
+-- | Close the coverage probes, before SDL_ttf shuts down. Fallbacks copied
+-- from them keep their shared streams open until they close themselves.
+closeCoverageProbes :: IO ()
+closeCoverageProbes =
+  readIORef coverageRef >>= \case
+    Nothing -> pure ()
+    Just cov -> do
+      forM_ (covProbes cov) $ \probe -> unless (probe == nullPtr) (ttfCloseFont probe)
+      writeIORef coverageRef (Just cov {covProbes = IM.empty})
+
+-- | Fallback families in the order shaping tries them: broad Latin, Greek
+-- and Cyrillic first, then scripts, then symbols, across Linux, Windows and
+-- macOS names.
+coverageFamilies :: [String]
+coverageFamilies =
+  [ "Noto Sans", "DejaVu Sans"
+  , "Noto Sans Arabic", "Noto Sans Hebrew", "Noto Sans Devanagari", "Noto Sans Bengali"
+  , "Noto Sans Tamil", "Noto Sans Telugu", "Noto Sans Gujarati", "Noto Sans Gurmukhi"
+  , "Noto Sans Kannada", "Noto Sans Malayalam", "Noto Sans Sinhala", "Noto Sans Thai"
+  , "Noto Sans Lao", "Noto Sans Khmer", "Noto Sans Myanmar", "Noto Sans Armenian"
+  , "Noto Sans Georgian", "Noto Sans Ethiopic", "Noto Sans CJK", "Noto Sans CJK SC", "Noto Sans CJK JP"
+  , "Noto Sans Symbols", "Noto Sans Symbols 2", "Noto Sans Math"
+  , "Segoe UI", "Segoe UI Symbol", "Nirmala UI", "Leelawadee UI", "Microsoft YaHei"
+  , "Yu Gothic", "Malgun Gothic", "Arial Unicode MS"
+  , "Geeza Pro", "Kohinoor Devanagari", "Thonburi", "PingFang SC", "Hiragino Sans"
+  , "Apple SD Gothic Neo", "Apple Symbols"
+  ]
+
+-- | ASCII glyph-cache slot. The cached 'Maybe' is shared on every hit, so a
+-- warm lookup returns the same heap object instead of rebuilding
+-- @Just GlyphQuad@ on each character.
+data CachedQuad
+  = UncachedQuad
+  -- Preserve the cached object even with -funbox-strict-fields: unpacking it
+  -- defeats sharing and reconstructs the lookup result on every hit.
+  | Cached {-# NOUNPACK #-} !(Maybe GlyphQuad)
+
+-- | Build immutable metric snapshots and explicit IO rasterisation callbacks.
+-- Font queries happen in 'fbPrepare'; atlas insertion happens in 'fbDrawShaped'
+-- and 'fbDrawGlyph'. All coordinates are logical (unscaled).
+--
+-- Standard ASCII (0..127) lookups are backed by a 'SmallMutableArray'
+-- fast path for branchless O(1) in-memory indexing, with automatic cache invalidation
+-- whenever the underlying glyph atlas is reset.
+{-# NOINLINE buildGlyphFontMetrics #-}
+buildGlyphFontMetrics :: GlyphAtlas -> SdlFont -> Float -> IO (FontMetrics, Text -> IO (Float, Float))
+buildGlyphFontMetrics ga sf scale = do
+  let !inv = if scale > 0 then scale else 1
+      baseFm = ttfFontMetricsScaled sf scale
+
+  -- Query ASCII metrics once for both advances and geometry, without atlas rasterization.
+  asciiMetrics <- mapM (getGlyphMetrics sf) [0 .. 127 :: CUInt]
+  -- Reuse fixed ASCII geometry across dynamic labels, so preparing a fresh
+  -- counter string does not query native glyph metrics for each character.
+  let !asciiGeometry = smallArrayFromList (map (fmap (metricsGlyphQuad sf inv)) asciiMetrics)
+      !asciiAdvances =
+        primArrayFromList
+          [ case metrics of
+              Nothing  -> sfSpaceAdvance sf / inv
+              Just m -> gmAdvance m / inv
+          | metrics <- asciiMetrics
+          ]
+
+  -- Cache of ASCII 0..127 glyph quads with epoch-based invalidation
+  asciiCacheArr <- newSmallArray 128 UncachedQuad
+  initEpoch <- readIORef (gaEpoch ga)
+  asciiEpochRef <- newIORef initEpoch
+
+  -- Non-ASCII glyph quads are memoised per font here (keyed by codepoint) so
+  -- repeated text still hits a shared value instead of rebuilding the record
+  -- on every character. Invalidated with the atlas epoch.
+  nonAsciiCacheRef <- newIORef IM.empty
+  nonAsciiEpochRef <- newIORef initEpoch
+
+  -- Kerning pairs are sparse and each miss costs a shaped 2-glyph
+  -- layout, so a pair cache keeps the hot pen loops off the FFI
+  -- boundary after first contact. Keyed by packed codepoint pair on this
+  -- 'FontMetrics' (the font id is implicit).
+  kernCacheRef <- newIORef emptyBounded
+
+  -- Shaped lines: SDL3_ttf lays each string out with its kerning,
+  -- ligatures, contextual forms, fallback fonts and right-to-left runs. A
+  -- line's layout is kept with its metric snapshot, which survives atlas
+  -- resets; the glyph quads drawn from it hold atlas UVs, so their cache is
+  -- dropped with the atlas epoch. Both caches are bounded by 'runCacheCap'.
+  preparedRef <- newIORef emptyBounded
+  shapedRef <- newIORef emptyBounded
+  quadCacheRef <- newIORef emptyBounded
+  initQuadEpoch <- readIORef (gaEpoch ga)
+  quadEpochRef <- newIORef initQuadEpoch
+
+  let
+    slotToQuad !gs =
+      GlyphQuad
+        { gqX  = gsOffX gs / inv
+        , gqY  = gsOffY gs / inv
+        , gqW  = gsW    gs / inv
+        , gqH  = gsH    gs / inv
+        , gqU0 = gsU0   gs
+        , gqV0 = gsV0   gs
+        , gqU1 = gsU1   gs
+        , gqV1 = gsV1   gs
+        }
+
+    resetAsciiCache !epoch = do
+      writeIORef asciiEpochRef epoch
+      mapM_ (\i -> writeSmallArray asciiCacheArr i UncachedQuad) [0 .. 127 :: Int]
+
+    lookupAsciiQuad !cp = do
+      curEpoch <- readIORef (gaEpoch ga)
+      lastEpoch <- readIORef asciiEpochRef
+      when (curEpoch /= lastEpoch) $ resetAsciiCache curEpoch
+      cached <- readSmallArray asciiCacheArr cp
+      case cached of
+        Cached mq -> pure mq
+        UncachedQuad -> do
+          mSlot <- lookupOrInsertGlyph ga sf (toEnum cp)
+          newEpoch <- readIORef (gaEpoch ga)
+          if newEpoch /= curEpoch
+            then do
+              -- The atlas was reset during insertion: this slot's UVs are
+              -- already stale, so do not cache them.
+              resetAsciiCache newEpoch
+              pure (fmap slotToQuad mSlot)
+            else do
+              let !mq = fmap slotToQuad mSlot
+              writeSmallArray asciiCacheArr cp (Cached mq)
+              pure mq
+
+    lookupNonAsciiQuad !c = do
+      curEpoch <- readIORef (gaEpoch ga)
+      lastEpoch <- readIORef nonAsciiEpochRef
+      when (curEpoch /= lastEpoch) $ do
+        writeIORef nonAsciiEpochRef curEpoch
+        writeIORef nonAsciiCacheRef IM.empty
+      m <- readIORef nonAsciiCacheRef
+      case IM.lookup (ord c) m of
+        Just mq -> pure mq
+        Nothing -> do
+          mSlot <- lookupOrInsertGlyph ga sf c
+          newEpoch <- readIORef (gaEpoch ga)
+          if newEpoch /= curEpoch
+            then pure (fmap slotToQuad mSlot)
+            else do
+              let !mq = fmap slotToQuad mSlot
+              modifyIORef' nonAsciiCacheRef (IM.insert (ord c) mq)
+              pure mq
+
+    {-# NOINLINE glyphLookup #-}
+    glyphLookup !c = do
+      ensureFontAlive sf
+      ensureAtlasAlive ga
+      let !cp = ord c
+      if (fromIntegral cp :: Word) < 128
+             then lookupAsciiQuad cp
+             else lookupNonAsciiQuad c
+
+    {-# NOINLINE advanceLookup #-}
+    advanceLookup !c =
+      let !cp = ord c
+       in if (fromIntegral cp :: Word) < 128
+             then pure (indexPrimArray asciiAdvances cp)
+             else do
+              mAdv <- getGlyphAdvance sf (fromIntegral cp)
+              pure $! case mAdv of
+                Nothing  -> sfSpaceAdvance sf / inv
+                Just adv -> adv / inv
+
+    {-# NOINLINE kernLookup #-}
+    kernLookup !prev !c = do
+      -- The cache lives on this 'FontMetrics', so the font id is constant and
+      -- the pair can be packed into a single Int key: no tuple on the hot path.
+      let !pk = (ord prev `shiftL` 21) .|. ord c
+      cache <- readIORef kernCacheRef
+      case HM.lookup pk (bcEntries cache) of
+        Just k -> pure k
+        Nothing -> do
+          raw <- ttfGetKerning (sfFont sf) (fromIntegral (ord prev) :: CUInt) (fromIntegral (ord c) :: CUInt)
+          let !k = fromIntegral raw / inv
+          -- At most 4096 pairs per font.
+          writeIORef kernCacheRef $! fst (insertBounded 4096 pk k cache)
+          pure k
+
+    -- The glyph quads of a shaped line, from the atlas. Quads are cached per
+    -- text and dropped with the atlas epoch, when their UVs go stale.
+    {-# NOINLINE shapedLookup #-}
+    shapedLookup !txt
+      | T.null txt = pure Nothing
+      | otherwise = do
+          ensureFontAlive sf
+          ensureAtlasAlive ga
+          ep <- readIORef (gaEpoch ga)
+          quadEp <- readIORef quadEpochRef
+          when (quadEp /= ep) $ do
+            writeIORef quadEpochRef ep
+            writeIORef quadCacheRef emptyBounded
+          cache <- readIORef quadCacheRef
+          -- Entries are kept wrapped so a hit returns them without allocating.
+          case HM.lookup txt (bcEntries cache) of
+            Just quads -> pure quads
+            Nothing -> do
+              shaped <- shapeOf txt
+              quads <- Just <$> placeGlyphs shaped
+              epAfter <- readIORef (gaEpoch ga)
+              when (cacheableText txt && epAfter == ep) $
+                modifyIORef' quadCacheRef (fst . insertBounded runCacheCap txt quads)
+              pure quads
+
+    -- Put a shaped line's glyphs in the atlas. A glyph the atlas has no room
+    -- for draws nothing, and the atlas resets before the next frame.
+    placeGlyphs (Shaped _ _ glyphs fontIndices fonts) = do
+      let !count = sizeofPrimArray fontIndices
+      out <- newPrimArray (count * 8)
+      let go !i
+            | i >= count = pure ()
+            | otherwise = do
+                let g k = fromIntegral (indexPrimArray glyphs (i * 9 + k)) :: Float
+                    o = i * 8
+                    write k v = writePrimArray out (o + k) v
+                    font = indexSmallArray fonts (fromIntegral (indexPrimArray fontIndices i))
+                    IntPtr handle = ptrToIntPtr (sfFont font)
+                mSlot <- lookupOrInsertGlyphIndex ga (fromIntegral (sfId font)) handle (fromIntegral (indexPrimArray glyphs (i * 9)))
+                write 0 (g 1 / inv)
+                write 1 (g 2 / inv)
+                write 2 (g 3 / inv)
+                write 3 (g 4 / inv)
+                case mSlot of
+                  Just slot -> do
+                    -- A glyph drawn in part samples only its source rect.
+                    let sx = g 5
+                        sy = g 6
+                        sw = g 7
+                        sh = g 8
+                        u0 = gsU0 slot + sx / glyphAtlasSize
+                        v0 = gsV0 slot + sy / glyphAtlasSize
+                        u1 = if sw > 0 then u0 + sw / glyphAtlasSize else gsU1 slot
+                        v1 = if sh > 0 then v0 + sh / glyphAtlasSize else gsV1 slot
+                    write 4 u0
+                    write 5 v0
+                    write 6 u1
+                    write 7 v1
+                  Nothing -> do
+                    let (u, v, _, _) = deadUv
+                    write 4 u
+                    write 5 v
+                    write 6 u
+                    write 7 v
+                go (i + 1)
+      go 0
+      ShapedGlyphs <$> unsafeFreezePrimArray out
+
+    -- A UV rect that always samples transparent pixels: column 4 sits
+    -- right of the 4px white patch (columns 0..3) and left of the first
+    -- slot (allocations start at x = 5), and the final row is never
+    -- written because every slot keeps 1px of padding.
+    deadUv :: (Float, Float, Float, Float)
+    deadUv =
+      let !u = 4.5 / glyphAtlasSize
+          !v = (glyphAtlasSize - 0.5) / glyphAtlasSize
+       in (u, v, u, v)
+
+    -- The shaped layout of a line, shared by measuring, preparing and
+    -- drawing it. Fonts that cover characters this one lacks join it before
+    -- the line is shaped.
+    shapeOf !txt = do
+      shapedCache <- readIORef shapedRef
+      case HM.lookup txt (bcEntries shapedCache) of
+        Just shaped -> pure shaped
+        Nothing -> do
+          ensureCoverage sf txt
+          shaped <- shapeLine sf inv txt
+          when (cacheableText txt) $
+            writeIORef shapedRef $! fst (insertBounded runCacheCap txt shaped shapedCache)
+          pure shaped
+
+    -- The width shaping draws with, so layout and drawing agree.
+    measure !txt
+      | T.null txt = pure emptySize
+      | otherwise = shapedSize <$> shapeOf txt
+    !emptySize = (0, sfLineSkip sf / inv)
+
+    glyphGeometry c
+      | ord c < 128 = pure (indexSmallArray asciiGeometry (ord c))
+      | otherwise = getGlyphGeometry sf inv c
+
+    backend = FontBackend prepareText shapedLookup glyphLookup
+
+    prepareText txt = do
+      ensureFontAlive sf
+      prepared <- readIORef preparedRef
+      case HM.lookup txt (bcEntries prepared) of
+        Just fm -> pure fm
+        Nothing -> do
+          let insertChar m c = IM.insert (ord c) c m
+              chars = T.foldl' insertChar (T.foldl' insertChar IM.empty " HxM") txt
+          advances <- traverse advanceLookup chars
+          geometry <- traverse glyphGeometry chars
+          let gather !pairs !previous remaining = case T.uncons remaining of
+                Nothing -> pure pairs
+                Just (c, rest) -> do
+                  let key = (ord previous `shiftL` 21) .|. ord c
+                  pairs' <- if IM.member key pairs then pure pairs else do
+                    k <- kernLookup previous c
+                    pure $! IM.insert key k pairs
+                  gather pairs' c rest
+          seedKerns <- gather IM.empty ' ' "xM"
+          kerns <- gather seedKerns 'M' txt
+          shaped <- if T.null txt then pure Nothing else Just <$> shapeOf txt
+          let !layout = fmap shapedText shaped
+          let !fm = baseFm
+                { fmAdvance = \c ->
+                    let cp = ord c
+                     in if cp < 128 then indexPrimArray asciiAdvances cp
+                          else IM.findWithDefault (sfSpaceAdvance sf / inv) cp advances
+                , fmKerning = \a b -> IM.findWithDefault 0 ((ord a `shiftL` 21) .|. ord b) kerns
+                , fmGlyph = \c -> IM.findWithDefault Nothing (ord c) geometry
+                , fmShape = \t -> if t == txt then layout else Nothing
+                , fmBackend = Just backend
+                , fmSnapScale = inv
+                }
+          when (cacheableText txt) $
+            writeIORef preparedRef $! fst (insertBounded runCacheCap txt fm prepared)
+          pure fm
+
+  fm <- prepareText ""
+  pure (fm, measure)
+
+-- | Return the SDL_Texture backing the glyph atlas, for passing to the renderer.
+glyphAtlasTexture :: GlyphAtlas -> IO (Ptr SDL_Texture)
+glyphAtlasTexture ga = textAtlasTexture (gaAtlas ga)
+
+-- ---------------------------------------------------------------------------
+withTtf :: IO a -> IO a
+withTtf act =
+  bracket startup shutdown $ \_ -> act
+  where
+    startup = do
+      ok <- ttfInit
+      when (not ok) $ fail "TTF_Init failed"
+    shutdown _ = closeCoverageProbes >> ttfQuit
+
+openFont :: FilePath -> Float -> IO SdlFont
+openFont path ptsize =
+  withCString path $ \cpath -> do
+    font <- ttfOpenFont cpath (realToFrac ptsize)
+    when (font == nullPtr) $
+      fail ("TTF_OpenFont failed for " ++ path)
+    readSdlFont ptsize Nothing font
+
+openFontFromMemory :: ByteString -> FilePath -> Float -> IO SdlFont
+openFontFromMemory bs label ptsize =
+  unsafeUseAsCStringLen bs $ \(ptr, len) -> do
+    (fontPtr, mTemp) <-
+      ttfOpenFontMemory (castPtr ptr) (fromIntegral len) (realToFrac ptsize) >>= \f ->
+        if f /= nullPtr
+          then pure (f, Nothing)
+          else openFontFromMemoryTemp bs ptsize
+    when (fontPtr == nullPtr) $
+      fail ("TTF_OpenFont failed for in-memory font " ++ label)
+    readSdlFont ptsize mTemp fontPtr
+
+openFontFromMemoryTemp :: ByteString -> Float -> IO (Ptr (), Maybe FilePath)
+openFontFromMemoryTemp bs openPt = do
+  tmpDir <- getTemporaryDirectory
+  (path, h) <- openTempFile tmpDir "nano-ui-font-"
+  BS.hPut h bs
+  hClose h
+  withCString path $ \cpath -> do
+    font <- ttfOpenFont cpath (realToFrac openPt)
+    if font == nullPtr
+      then removeFile path >> pure (nullPtr, Nothing)
+      else pure (font, Just path)
+
+-- | Wrap an open TTF font; @mTemp@ is a temp file to delete on close.
+readSdlFont :: Float -> Maybe FilePath -> Ptr () -> IO SdlFont
+readSdlFont ptsize mTemp font = do
+  fid <- newFontId
+  alive <- newIORef True
+  fallbacks <- newIORef IM.empty
+  lineSkip <- ttfLineSkip font
+  ascent <- ttfAscent font
+  spaceAdv <- ttfSpaceAdvance font
+  pure
+    SdlFont
+      { sfId = fid
+      , sfAlive = alive
+      , sfFont = font
+      , sfLineSkip = realToFrac lineSkip
+      , sfAscent = realToFrac ascent
+      , sfSpaceAdvance = realToFrac spaceAdv
+      , sfTempPath = mTemp
+      , sfPointSize = ptsize
+      , sfFallbacks = fallbacks
+      }
+
+openFontSource :: FontSource -> Float -> IO SdlFont
+openFontSource (FontFromPath path) ptsize = openFont path ptsize
+openFontSource (FontFromMemory bs label) ptsize =
+  openFontFromMemory bs label ptsize
+
+openFontSourceWithFallback :: FontSource -> FontSource -> Float -> IO SdlFont
+openFontSourceWithFallback primary fallback ptsize =
+  openFontSource primary ptsize
+    `catch` \(e :: SomeException) ->
+      if fontSourcesSame primary fallback
+        then throwIO e
+        else openFontSource fallback ptsize
+          `catch` \(_ :: SomeException) -> throwIO e
+
+fontSourcesSame :: FontSource -> FontSource -> Bool
+fontSourcesSame (FontFromPath a) (FontFromPath b) = a == b
+fontSourcesSame (FontFromMemory _ la) (FontFromMemory _ lb) = la == lb
+fontSourcesSame _ _ = False
+
+closeFont :: SdlFont -> IO ()
+closeFont sf = do
+  alive <- atomicModifyIORef' (sfAlive sf) (\open -> (False, open))
+  when alive $ do
+    ttfCloseFont (sfFont sf)
+    mapM_ removeFile (sfTempPath sf)
+    readIORef (sfFallbacks sf) >>= mapM_ closeFont
+
+-- | Install glyph-atlas-backed 'FontMetrics' (from 'buildGlyphFontMetrics')
+-- so that 'pushText' emits per-glyph textured quads into the draw arena.
+-- Text measurement uses the primary font's shaped lines.
+withTtfMeasureGlyph ::
+  Context ->
+  (Text -> IO (Float, Float)) -> -- ^ primary font measurement
+  FontMetrics -> -- ^ glyph-atlas fm for primary font
+  FontMetrics -> -- ^ glyph-atlas fm for mono font
+  Float ->
+  Context
+withTtfMeasureGlyph ctx measure fm monoFm scale =
+  let ctx1 =
+        withExternalText
+          ( withMeasureText
+              (withMonoFontMetrics (withFontMetrics ctx fm) monoFm)
+              measure
+          )
+          False
+   in wrapMeasureCache scale ctx1 measure
+
+ttfFontMetricsScaled :: SdlFont -> Float -> FontMetrics
+ttfFontMetricsScaled sf scale =
+  let inv = if scale > 0 then scale else 1
+   in (monospaceMetrics (sfLineSkip sf / inv))
+        { fmAscent = sfAscent sf / inv
+        , fmAdvance = const (sfSpaceAdvance sf / inv)
+        }
+
+tryInsert :: Ptr () -> Ptr () -> IO (Maybe (Float, Float, Float, Float))
+tryInsert atlas surf =
+  allocaBytes (4 * sizeOf (0 :: CFloat)) $ \px -> do
+    let py = plusPtr px (sizeOf (0 :: CFloat))
+        tw = plusPtr py (sizeOf (0 :: CFloat))
+        th = plusPtr tw (sizeOf (0 :: CFloat))
+    ok <- textAtlasInsertSurface atlas surf px py tw th
+    if ok
+      then do
+        x <- realToFrac <$> peek px
+        y <- realToFrac <$> peek py
+        w <- realToFrac <$> peek tw
+        h <- realToFrac <$> peek th
+        pure (Just (x, y, w, h))
+      else pure Nothing
+
+withUtf8 :: Text -> (CString -> CSize -> IO a) -> IO a
+withUtf8 txt act =
+  TF.useAsPtr txt $ \ptr len ->
+    act (castPtr ptr) (fromIntegral len)
+
+foreign import ccall unsafe "nano_ui_ttf_init"
+  ttfInit :: IO Bool
+
+foreign import ccall unsafe "nano_ui_ttf_quit"
+  ttfQuit :: IO ()
+
+foreign import ccall unsafe "nano_ui_ttf_open_font"
+  ttfOpenFont :: CString -> CFloat -> IO (Ptr ())
+
+foreign import ccall unsafe "nano_ui_ttf_open_font_memory"
+  ttfOpenFontMemory :: Ptr () -> CSize -> CFloat -> IO (Ptr ())
+
+foreign import ccall unsafe "nano_ui_ttf_close_font"
+  ttfCloseFont :: Ptr () -> IO ()
+
+foreign import ccall unsafe "nano_ui_ttf_line_skip"
+  ttfLineSkip :: Ptr () -> IO CFloat
+
+foreign import ccall unsafe "nano_ui_ttf_ascent"
+  ttfAscent :: Ptr () -> IO CFloat
+
+foreign import ccall unsafe "nano_ui_ttf_space_advance"
+  ttfSpaceAdvance :: Ptr () -> IO CFloat
+
+foreign import ccall unsafe "nano_ui_text_atlas_create"
+  textAtlasCreate :: Ptr SDL_Renderer -> IO (Ptr ())
+
+foreign import ccall unsafe "nano_ui_text_atlas_destroy"
+  textAtlasDestroy :: Ptr () -> IO ()
+
+foreign import ccall unsafe "nano_ui_text_atlas_reset"
+  textAtlasReset :: Ptr () -> IO ()
+
+foreign import ccall unsafe "nano_ui_text_atlas_texture"
+  textAtlasTexture :: Ptr () -> IO (Ptr SDL_Texture)
+
+foreign import ccall unsafe "nano_ui_text_atlas_insert_surface"
+  textAtlasInsertSurface ::
+    Ptr () ->
+    Ptr () ->
+    Ptr CFloat ->
+    Ptr CFloat ->
+    Ptr CFloat ->
+    Ptr CFloat ->
+    IO Bool
+
+foreign import ccall unsafe "SDL_DestroySurface"
+  freeSurface :: Ptr () -> IO ()
+
+foreign import ccall unsafe "nano_ui_ttf_glyph_metrics"
+  ttfGlyphMetrics ::
+    Ptr () ->   -- font
+    CUInt ->    -- codepoint
+    Ptr CInt -> -- out_minx
+    Ptr CInt -> -- out_maxx
+    Ptr CInt -> -- out_miny
+    Ptr CInt -> -- out_maxy
+    Ptr CInt -> -- out_advance
+    IO Bool
+
+foreign import ccall unsafe "nano_ui_ttf_render_glyph_surface"
+  ttfRenderGlyphSurface ::
+    Ptr () ->        -- font
+    CUInt ->         -- codepoint
+    Ptr (Ptr ()) ->  -- out_surface
+    IO Bool
+
+foreign import ccall unsafe "nano_ui_ttf_shape"
+  ttfShape :: Ptr () -> CString -> CSize -> CInt -> Ptr () -> IO Bool
+
+foreign import ccall unsafe "nano_ui_ttf_shaped_free"
+  ttfShapedFree :: Ptr () -> IO ()
+
+foreign import ccall unsafe "nano_ui_ttf_shaped_size"
+  ttfShapedSize :: IO CSize
+
+foreign import ccall unsafe "nano_ui_ttf_shaped_int"
+  ttfShapedInt :: Ptr () -> CInt -> IO CInt
+
+foreign import ccall unsafe "nano_ui_ttf_shaped_ptr"
+  ttfShapedPtr :: Ptr () -> CInt -> IO (Ptr ())
+
+foreign import ccall unsafe "nano_ui_ttf_render_glyph_index_surface"
+  ttfRenderGlyphIndexSurface :: Ptr () -> CUInt -> Ptr (Ptr ()) -> IO Bool
+
+foreign import ccall unsafe "nano_ui_ttf_has_glyph"
+  ttfHasGlyph :: Ptr () -> CUInt -> IO Bool
+
+foreign import ccall unsafe "nano_ui_ttf_add_fallback"
+  ttfAddFallback :: Ptr () -> Ptr () -> IO Bool
+
+foreign import ccall unsafe "nano_ui_ttf_remove_fallback"
+  ttfRemoveFallback :: Ptr () -> Ptr () -> IO ()
+
+foreign import ccall unsafe "nano_ui_ttf_copy_font"
+  ttfCopyFont :: Ptr () -> CFloat -> IO (Ptr ())
+
+foreign import ccall unsafe "nano_ui_ttf_get_kerning"
+  ttfGetKerning :: Ptr () -> CUInt -> CUInt -> IO CInt
+
+-- ---------------------------------------------------------------------------
+-- Font cache: the base sans and mono faces, plus fonts opened per size and
+-- variant on demand
+
+-- | A font variant and its point size key, @round (targetPt * 2)@.
+data FontCacheKey = FontCacheKey !FontVariant !Int
+  deriving (Eq)
+
+instance Hashable FontCacheKey where
+  hashWithSalt s (FontCacheKey variant ptKey) =
+    s `hashWithSalt` fromEnum variant `hashWithSalt` ptKey
+
+data CachedFontEntry = CachedFontEntry
+  { cfeFont    :: !SdlFont
+  , cfeFm      :: !FontMetrics
+  , cfeMeasure :: !(Text -> IO (Float, Float))
+  }
+
+data SdlFontCache = SdlFontCache
+  { sfcPrimarySourceRef :: !(IORef FontSource)
+  , sfcFallbackSource :: !FontSource
+  , sfcMonoSource     :: !FontSource
+  , sfcMonoFallback   :: !FontSource
+  , sfcGlyphAtlas     :: !GlyphAtlas
+  , sfcBasePt         :: !Float
+  , sfcScaleRef       :: !(IORef Float)
+  -- ^ The display scale, owned by the window and read here.
+  , sfcBaseEntries    :: !(IORef (CachedFontEntry, CachedFontEntry))
+  , sfcDynamicCache   :: !(IORef (BoundedCache FontCacheKey CachedFontEntry))
+  }
+
+-- | Open the base sans and mono fonts at the display scale, and re-warm them
+-- into the glyph atlas after every atlas reset.
+newSdlFontCache ::
+  FontSource -> -- ^ primary font source
+  FontSource -> -- ^ fallback font source
+  FontSource -> -- ^ mono font source
+  FontSource -> -- ^ mono fallback font source
+  GlyphAtlas ->
+  Float ->      -- ^ base font size (pt)
+  IORef Float -> -- ^ display scale
+  IO SdlFontCache
+newSdlFontCache primary fallback mono monoFb ga basePt scaleRef = do
+  scale <- readIORef scaleRef
+  primaryRef <- newIORef primary
+  sansEntry <- openCachedFont ga scale primary fallback basePt
+  monoEntry <- openCachedFont ga scale mono monoFb basePt
+  baseEntriesRef <- newIORef (sansEntry, monoEntry)
+  cacheRef <- newIORef emptyBounded
+  -- The hook reads the base entries when it runs, so a reset always warms the
+  -- live fonts, never ones already closed.
+  let rewarm = do
+        (sans, monoBase) <- readIORef baseEntriesRef
+        warmGlyphAtlas ga (cfeFont sans)
+        warmGlyphAtlas ga (cfeFont monoBase)
+  registerGlyphAtlasRewarm ga rewarm
+  rewarm
+  pure
+    SdlFontCache
+      { sfcPrimarySourceRef = primaryRef
+      , sfcFallbackSource = fallback
+      , sfcMonoSource     = mono
+      , sfcMonoFallback   = monoFb
+      , sfcGlyphAtlas     = ga
+      , sfcBasePt         = basePt
+      , sfcScaleRef       = scaleRef
+      , sfcBaseEntries    = baseEntriesRef
+      , sfcDynamicCache   = cacheRef
+      }
+
+-- | A font from a source (or its fallback) at a point size, rasterised at
+-- the display scale, with its glyph metrics.
+openCachedFont :: GlyphAtlas -> Float -> FontSource -> FontSource -> Float -> IO CachedFontEntry
+openCachedFont ga scale primary fallback pt = do
+  font <- openFontSourceWithFallback primary fallback (pt * scale)
+  (fm, measure) <- buildGlyphFontMetrics ga font scale
+  pure (CachedFontEntry font fm measure)
+
+-- | The primary (sans) family's source, for the debug readout.
+sdlFontCacheSource :: SdlFontCache -> IO FontSource
+sdlFontCacheSource cache = readIORef (sfcPrimarySourceRef cache)
+
+-- | Close fonts and drop their glyphs from the shared atlas index.
+closeCachedFonts :: GlyphAtlas -> [SdlFont] -> IO ()
+closeCachedFonts ga fonts = do
+  fallbacks <- concat <$> mapM (fmap IM.elems . readIORef . sfFallbacks) fonts
+  let handles = IS.fromList [fromIntegral (sfId f) | f <- fonts ++ fallbacks]
+  mapM_ closeFont fonts
+  modifyIORef' (gaEntries ga) (`IM.withoutKeys` IS.fromList (map (fromIntegral . sfId) fonts))
+  modifyIORef' (gaIndexEntries ga) (`IM.withoutKeys` handles)
+
+-- | Close every open font, base and dynamic.
+destroySdlFontCache :: SdlFontCache -> IO ()
+destroySdlFontCache cache = do
+  dynamic <- atomicModifyIORef' (sfcDynamicCache cache) (\c -> (emptyBounded, c))
+  (sans, mono) <- readIORef (sfcBaseEntries cache)
+  closeCachedFonts (sfcGlyphAtlas cache) (cfeFont sans : cfeFont mono : map cfeFont (HM.elems (bcEntries dynamic)))
+
+-- | Reopen the base fonts from @source@ at the current display scale, close
+-- every dynamic size, and reset the glyph atlas, which re-warms the new base
+-- fonts.
+reloadSdlFontCache :: SdlFontCache -> FontSource -> IO ()
+reloadSdlFontCache cache source = do
+  destroySdlFontCache cache
+  writeIORef (sfcPrimarySourceRef cache) source
+  scale <- readIORef (sfcScaleRef cache)
+  let ga = sfcGlyphAtlas cache
+  sansEntry <- openCachedFont ga scale source (sfcFallbackSource cache) (sfcBasePt cache)
+  monoEntry <- openCachedFont ga scale (sfcMonoSource cache) (sfcMonoFallback cache) (sfcBasePt cache)
+  writeIORef (sfcBaseEntries cache) (sansEntry, monoEntry)
+  resetGlyphAtlas ga
+
+-- | Install the cache's base fonts as the context's measurement and glyph
+-- metrics, and its sizes and variants as the font resolver.
+withSdlFontCache :: SdlFontCache -> Context -> IO Context
+withSdlFontCache cache ctx = do
+  scale <- readIORef (sfcScaleRef cache)
+  (sans, mono) <- readIORef (sfcBaseEntries cache)
+  pure (withFontResolver (withTtfMeasureGlyph ctx (cfeMeasure sans) (cfeFm sans) (cfeFm mono) scale) (resolveSdlFont cache) (resolveSdlMeasure cache))
+
+-- | The open font for a size and variant. Weight and style pick nothing
+-- here: they are drawn synthetically over the regular face, because SDL_ttf's
+-- style flags change its layout boxes but not the glyph images shaped text
+-- draws, so a styled face would not line up.
+getOrLoadCachedFont ::
+  SdlFontCache ->
+  Float ->
+  FontWeight ->
+  FontStyle ->
+  FontVariant ->
+  IO CachedFontEntry
+getOrLoadCachedFont cache sz _weight _style var = do
+  let basePt = sfcBasePt cache
+      rawPt = if sz > 0 then sz else basePt
+      -- Quantize dynamic sizes to 0.5 pt increments so dragging sliders
+      -- doesn't create hundreds of redundant TTF_Font instances.
+      targetPt = fromIntegral (round (rawPt * 2.0) :: Int) / 2.0
+      ptKey = round (targetPt * 2.0) :: Int
+      basePtKey = round (basePt * 2.0) :: Int
+      isBase =
+        ptKey == basePtKey
+  if isBase
+    then do
+      (sansEntry, monoEntry) <- readIORef (sfcBaseEntries cache)
+      pure (if var == FontMono then monoEntry else sansEntry)
+    else do
+      let key = FontCacheKey var ptKey
+      dynamic <- readIORef (sfcDynamicCache cache)
+      case HM.lookup key (bcEntries dynamic) of
+        Just entry -> do
+          -- Least recently used goes first: move a hit to the back of the
+          -- eviction order, so fonts drawn every frame are never closed.
+          case Seq.viewr (bcOrder dynamic) of
+            _ Seq.:> newest | newest == key -> pure ()
+            _ ->
+              -- Each key appears in the order once, and a hot key sits near
+              -- the back, so search from the right and delete that one entry.
+              let order = bcOrder dynamic
+               in writeIORef (sfcDynamicCache cache) $!
+                    dynamic {bcOrder = maybe order (`Seq.deleteAt` order) (Seq.elemIndexR key order) Seq.|> key}
+          pure entry
+        Nothing -> do
+          scale <- readIORef (sfcScaleRef cache)
+          primarySans <- readIORef (sfcPrimarySourceRef cache)
+          let (primary, fallback) =
+                if var == FontMono
+                  then (sfcMonoSource cache, sfcMonoFallback cache)
+                  else (primarySans, sfcFallbackSource cache)
+          -- Dynamic fonts are not warmed: they insert only glyphs drawn.
+          entry <- openCachedFont (sfcGlyphAtlas cache) scale primary fallback targetPt
+          -- At most 48 dynamic sizes stay open.
+          let (dynamic', evicted) = insertBounded 48 key entry dynamic
+          writeIORef (sfcDynamicCache cache) $! dynamic'
+          mapM_ (closeCachedFonts (sfcGlyphAtlas cache) . pure . cfeFont) evicted
+          pure entry
+
+resolveSdlFont ::
+  SdlFontCache ->
+  Float ->
+  FontWeight ->
+  FontStyle ->
+  FontVariant ->
+  IO (FontMetrics, Bool)
+resolveSdlFont cache sz weight style var = do
+  entry <- getOrLoadCachedFont cache sz weight style var
+  pure (cfeFm entry, False)
+
+resolveSdlMeasure ::
+  SdlFontCache ->
+  Float ->
+  FontWeight ->
+  FontStyle ->
+  FontVariant ->
+  Text ->
+  IO (Float, Float)
+resolveSdlMeasure cache sz weight style var txt = do
+  entry <- getOrLoadCachedFont cache sz weight style var
+  cfeMeasure entry txt
diff --git a/lib/NanoUI/Sdl/Font/Inter.hs b/lib/NanoUI/Sdl/Font/Inter.hs
new file mode 100644
--- /dev/null
+++ b/lib/NanoUI/Sdl/Font/Inter.hs
@@ -0,0 +1,17 @@
+{-# LANGUAGE TemplateHaskell #-}
+
+-- | Embedded default UI font. Import only from font bootstrap code.
+module NanoUI.Sdl.Font.Inter
+  ( fontInterBytes
+  , fontInterLabel
+  ) where
+
+import Data.ByteString (ByteString)
+import Data.FileEmbed (embedFileRelative)
+
+fontInterLabel :: FilePath
+fontInterLabel = "inter.ttf"
+
+{-# NOINLINE fontInterBytes #-}
+fontInterBytes :: ByteString
+fontInterBytes = $(embedFileRelative "data/inter.ttf")
diff --git a/lib/NanoUI/Sdl/Font/Resolve.hs b/lib/NanoUI/Sdl/Font/Resolve.hs
new file mode 100644
--- /dev/null
+++ b/lib/NanoUI/Sdl/Font/Resolve.hs
@@ -0,0 +1,43 @@
+module NanoUI.Sdl.Font.Resolve
+  ( embeddedFontSource
+  , resolveNanoUIFont
+  , defaultFontSearch
+  , defaultFontSearchMono
+  ) where
+
+import NanoUI.Sdl.Font (FontSource (..))
+import NanoUI.Sdl.Font.Inter (fontInterBytes, fontInterLabel)
+import NanoUI.Sdl.Font.Search (searchFonts)
+import NanoUI.Sdl.NanoUIFont (NanoUIFont (..))
+
+embeddedFontSource :: FontSource
+embeddedFontSource = FontFromMemory fontInterBytes fontInterLabel
+
+defaultFontSearch :: NanoUIFont
+defaultFontSearch =
+  FontSearch
+    [ "Inter"
+    , "Montserrat"
+    , "Work Sans"
+    , "Roboto"
+    , "Open Sans"
+    , "Helvetica Neue"
+    ]
+
+defaultFontSearchMono :: NanoUIFont
+defaultFontSearchMono =
+  FontSearch
+    [ "Consolas"
+    , "Courier New"
+    , "Liberation Mono"
+    , "DejaVu Sans Mono"
+    , "monospace"
+    ]
+
+resolveNanoUIFont :: NanoUIFont -> IO FontSource
+resolveNanoUIFont DefaultFont = pure embeddedFontSource
+resolveNanoUIFont (FontFilePath path) = pure (FontFromPath path)
+resolveNanoUIFont (FontSearch names) =
+  searchFonts names >>= \case
+    Just path -> pure (FontFromPath path)
+    Nothing -> pure embeddedFontSource
diff --git a/lib/NanoUI/Sdl/Font/Search.hs b/lib/NanoUI/Sdl/Font/Search.hs
new file mode 100644
--- /dev/null
+++ b/lib/NanoUI/Sdl/Font/Search.hs
@@ -0,0 +1,231 @@
+-- | Locate system font files by walking the standard font directories for the
+-- current platform, without fontconfig. Each candidate file is matched
+-- against the requested family name using a normalised-filename heuristic.
+module NanoUI.Sdl.Font.Search
+  ( searchFonts
+  , searchFontFamilies
+  , listFontFamilies
+  ) where
+
+import Control.Exception (IOException, catch)
+import Data.Containers.ListUtils (nubOrd)
+import Data.IORef (IORef, newIORef, readIORef, writeIORef)
+import Data.Char (isDigit, isLower, isSpace, isUpper, toLower)
+import Data.List (isInfixOf, minimumBy, sort, stripPrefix)
+import Data.Maybe (fromMaybe, listToMaybe, mapMaybe)
+import Data.Ord (Down (..), comparing)
+import qualified Data.Set as Set
+import System.Directory (getHomeDirectory)
+import System.Directory.Recursive (getFilesRecursive)
+import System.Environment (lookupEnv)
+import System.FilePath (takeBaseName, takeExtension, (</>))
+import System.Info (os)
+import System.IO.Unsafe (unsafePerformIO)
+
+-- | Try each family name in order, returning the first font file that
+-- matches.  Generic families like @monospace@ are expanded to a list of
+-- concrete families first.
+searchFonts :: [String] -> IO (Maybe FilePath)
+searchFonts names = case concatMap families names of
+  [] -> pure Nothing
+  candidates -> do
+    files <- fontStems
+    pure (listToMaybe (mapMaybe (`bestMatch` files) candidates))
+  where
+    families name =
+      let norm = normalize name
+       in if null norm then [] else maybe [norm] (concatMap families) (expandGeneric norm)
+
+-- | The file for each family that is installed, in the order asked, reading
+-- the font directories once.
+searchFontFamilies :: [String] -> IO [FilePath]
+searchFontFamilies names = do
+  files <- fontStems
+  pure (nubOrd (mapMaybe (\name -> bestMatch (normalize name) files) names))
+
+-- | Human-readable names for every installed font family, deduped and sorted.
+-- Each name is a usable 'searchFonts' token: the same normalization is applied
+-- to both the requested family and the file stem, so a listed name always
+-- resolves back to (at least) the file it came from. Non-text faces (icons,
+-- colour emoji) are included; callers that need sans families can filter the
+-- result themselves.
+listFontFamilies :: IO [String]
+listFontFamilies = do
+  files <- fontStems
+  pure (Set.toAscList (Set.fromList (map (prettyFamily . takeBaseName . snd) files)))
+
+-- | Filename stem -> display family. Everything from the first @-@ is treated
+-- as style (\"Regular\", \"Bold Italic\", ...); camel case is split so
+-- @NotoSansArabic@ reads as @Noto Sans Arabic@. Kept case-insensitively
+-- compatible with 'normalize'.
+prettyFamily :: String -> String
+prettyFamily = separateCamel . stripStyle
+  where
+    stripStyle s = case break (== '-') s of
+      (base, _) -> base
+    separateCamel = go
+      where
+        go [] = []
+        go (c : cs) = c : goTail c cs
+        goTail _ [] = []
+        goTail prev (c : cs)
+          | isUpper c && (isLower prev || isDigit prev) = ' ' : c : goTail c cs
+          | otherwise = c : goTail c cs
+
+-- ---------------------------------------------------------------------------
+-- Directory traversal
+
+-- | Every font file under every standard font directory for this platform,
+-- with its normalised name, normalised once for all the families matched
+-- against it. User directories come first so that user-installed fonts win
+-- over system ones; missing or unreadable roots are skipped. The directories
+-- are walked once per process.
+fontStems :: IO [(String, FilePath)]
+fontStems =
+  readIORef fontStemsRef >>= \case
+    Just stems -> pure stems
+    Nothing -> do
+      roots <- defaultFontDirs
+      files <- concat <$> mapM (fmap (sort . filter isFontFile) . filesBelow) roots
+      let stems = map (\path -> (normalize (takeBaseName path), path)) files
+      writeIORef fontStemsRef (Just stems)
+      pure stems
+
+{-# NOINLINE fontStemsRef #-}
+fontStemsRef :: IORef (Maybe [(String, FilePath)])
+fontStemsRef = unsafePerformIO (newIORef Nothing)
+
+filesBelow :: FilePath -> IO [FilePath]
+filesBelow root =
+  getFilesRecursive root `catch` \(_ :: IOException) -> pure []
+
+defaultFontDirs :: IO [FilePath]
+defaultFontDirs =
+  case os of
+    "darwin" -> macDirs
+    "mingw32" -> winDirs
+    _ -> linuxDirs
+  where
+    linuxDirs :: IO [FilePath]
+    linuxDirs = do
+      home <- getHomeDirectory
+      pure
+        [ home </> ".local/share/fonts"
+        , "/usr/local/share/fonts"
+        , "/usr/share/fonts"
+        ]
+
+    macDirs :: IO [FilePath]
+    macDirs = do
+      home <- getHomeDirectory
+      pure
+        [ home </> "Library/Fonts"
+        , "/Library/Fonts"
+        , "/System/Library/Fonts"
+        ]
+
+    winDirs :: IO [FilePath]
+    winDirs = do
+      mRoot <- lookupEnv "SystemRoot"
+      let systemDir = fromMaybe "C:\\Windows" mRoot </> "Fonts"
+      mLocal <- lookupEnv "LOCALAPPDATA"
+      let userDirs =
+            maybe [] (\l -> [l </> "Microsoft" </> "Windows" </> "Fonts"]) mLocal
+      pure (userDirs ++ [systemDir])
+
+isFontFile :: FilePath -> Bool
+isFontFile path =
+  map toLower (takeExtension path) `elem` [".ttf", ".otf", ".ttc", ".otc"]
+
+-- ---------------------------------------------------------------------------
+-- Family matching
+
+-- | Pick the highest-scoring file for @norm@ (a normalised family name).
+bestMatch :: String -> [(String, FilePath)] -> Maybe FilePath
+bestMatch norm files =
+  case [(score, path) | (stem, path) <- files, Just score <- [maximum (Nothing : map (`matchScore` stem) candidates)]] of
+    [] -> Nothing
+    scored ->
+      -- minimumBy keeps the first tie; descending scores prefer the best face.
+      let (_, best) = minimumBy (comparing (Down . fst)) scored
+       in Just best
+  where
+    candidates = norm : familyAliases norm
+
+matchScore :: String -> String -> Maybe Int
+matchScore "" _ = Nothing
+matchScore norm stem
+  | stem == norm = Just 100
+  | otherwise =
+      case stripPrefix norm stem of
+        Nothing -> Nothing
+        Just t
+          | t `elem` regularTails -> Just 90
+          | t `elem` otherTails -> Just 70
+          | "variable" `isInfixOf` t -> Just 20
+          | otherwise -> Just 60
+
+-- | Style tails that indicate the regular weight of a family.
+regularTails :: [String]
+regularTails = ["regular", "r", "normal", "medium", "text"]
+
+-- | Style tails for non-regular weights; still worth preferring over an
+-- unrelated font, but a regular (or exact) match wins.
+otherTails :: [String]
+otherTails =
+  [ "bold"
+  , "italic"
+  , "light"
+  , "semibold"
+  , "semibolditalic"
+  , "extrabold"
+  , "thin"
+  , "black"
+  , "regularitalic"
+  , "bolditalic"
+  , "mediumitalic"
+  , "oblique"
+  ]
+
+-- | Windows ships @Consolas@ as @consola.ttf@ and @Courier New@ as
+-- @cour.ttf@, so those families need filename aliases that prefix-match
+-- differently than their family names.
+familyAliases :: String -> [String]
+familyAliases "consolas" = ["consola"]
+familyAliases "couriernew" = ["cour"]
+familyAliases _ = []
+
+-- | Expand a generic CSS family into concrete families to search in order.
+expandGeneric :: String -> Maybe [String]
+expandGeneric "monospace" =
+  Just
+    [ "DejaVu Sans Mono"
+    , "Liberation Mono"
+    , "Ubuntu Mono"
+    , "Noto Sans Mono"
+    , "Consolas"
+    , "Courier New"
+    ]
+expandGeneric "sansserif" =
+  Just
+    [ "DejaVu Sans"
+    , "Liberation Sans"
+    , "Noto Sans"
+    , "Open Sans"
+    , "Helvetica Neue"
+    ]
+expandGeneric "serif" =
+  Just
+    [ "DejaVu Serif"
+    , "Liberation Serif"
+    , "Noto Serif"
+    , "Times New Roman"
+    ]
+expandGeneric _ = Nothing
+
+-- | Fold case, drop whitespace and separators, so family \/ file names can be
+-- compared loosely ("DejaVu Sans Mono" vs @DejaVuSansMono.ttf@).
+normalize :: String -> String
+normalize =
+  map toLower
+    . filter (\c -> not (isSpace c) && c /= '-' && c /= '_')
diff --git a/lib/NanoUI/Sdl/Image.hs b/lib/NanoUI/Sdl/Image.hs
new file mode 100644
--- /dev/null
+++ b/lib/NanoUI/Sdl/Image.hs
@@ -0,0 +1,71 @@
+module NanoUI.Sdl.Image
+  ( ImageAtlas
+  , newImageAtlas
+  , destroyImageAtlas
+  , syncImageAtlas
+  , lookupImage
+  )
+where
+
+import Control.Exception (mask_)
+import Control.Monad (when)
+import Data.IORef (IORef, newIORef, readIORef, writeIORef)
+import Data.Word (Word8)
+import Foreign.ForeignPtr (ForeignPtr, withForeignPtr)
+import Foreign.Ptr (Ptr, castPtr, nullPtr)
+import NanoUI.Testing (Context, atlasSnapshot, atlasTextureId)
+import SDL3.Sys.Bindgen.Blendmode (sDL_BLENDMODE_BLEND)
+import SDL3.Sys.Bindgen.Pixels qualified as Pixels
+import SDL3.Sys.Bindgen.Render (SDL_Renderer, SDL_Texture)
+import SDL3.Sys.Bindgen.Render qualified as Render
+import SDL3.Sys.Bindgen.Runtime.PtrConst qualified as PtrConst
+import SDL3.Sys.Render (createTextureSafe, destroyTexture, setTextureBlendMode, updateTextureSafe)
+
+-- A texture and its metadata have one lifetime; publish them together.
+newtype ImageAtlas = ImageAtlas (IORef (Maybe AtlasTexture))
+
+data AtlasTexture = AtlasTexture
+  { atTexture :: !(Ptr SDL_Texture)
+  , atGeneration :: !Int
+  }
+
+newImageAtlas :: IO ImageAtlas
+newImageAtlas = ImageAtlas <$> newIORef Nothing
+
+destroyImageAtlas :: ImageAtlas -> IO ()
+destroyImageAtlas (ImageAtlas ref) = mask_ $ do
+  old <- readIORef ref
+  writeIORef ref Nothing
+  mapM_ (destroyTexture . atTexture) old
+
+syncImageAtlas :: Ptr SDL_Renderer -> ImageAtlas -> Context -> IO ()
+syncImageAtlas ren atlas@(ImageAtlas ref) ctx = do
+  snap <- atlasSnapshot ctx
+  case snap of
+    Nothing -> pure ()
+    Just (w, h, pixels, gen) -> do
+      old <- readIORef ref
+      when (Just gen /= fmap atGeneration old) $
+        uploadAtlas ren atlas w h pixels gen
+
+uploadAtlas ::
+  Ptr SDL_Renderer -> ImageAtlas -> Int -> Int -> ForeignPtr Word8 -> Int -> IO ()
+uploadAtlas ren (ImageAtlas ref) w h pixels gen = mask_ $
+  withForeignPtr pixels $ \ptr -> do
+    tex <- createTextureSafe ren Pixels.SDL_PIXELFORMAT_RGBA32 Render.SDL_TEXTUREACCESS_STATIC (fromIntegral w) (fromIntegral h)
+    ok <-
+      if tex == nullPtr
+        then pure False
+        else do
+          _ <- setTextureBlendMode tex (fromIntegral sDL_BLENDMODE_BLEND)
+          uploaded <- updateTextureSafe tex (PtrConst.unsafeFromPtr nullPtr) (PtrConst.unsafeFromPtr (castPtr ptr)) (fromIntegral (w * 4))
+          if uploaded then pure True else destroyTexture tex >> pure False
+    when ok $ do
+      old <- readIORef ref
+      writeIORef ref (Just (AtlasTexture tex gen))
+      mapM_ (destroyTexture . atTexture) old
+
+lookupImage :: ImageAtlas -> Int -> IO (Maybe (Ptr SDL_Texture))
+lookupImage (ImageAtlas ref) tid
+  | tid == atlasTextureId = fmap atTexture <$> readIORef ref
+  | otherwise = pure Nothing
diff --git a/lib/NanoUI/Sdl/Input.hs b/lib/NanoUI/Sdl/Input.hs
new file mode 100644
--- /dev/null
+++ b/lib/NanoUI/Sdl/Input.hs
@@ -0,0 +1,309 @@
+{-# LANGUAGE OverloadedRecordDot #-}
+{-# LANGUAGE NoFieldSelectors #-}
+
+-- | SDL3 event polling and waiting, and translation of SDL events into
+-- 'NanoUI.Input.Input'.
+module NanoUI.Sdl.Input
+  ( SdlEvent (..)
+  , pollEvents
+  , waitEvent
+  , waitEventTimeout
+  , applyEvent
+  , isHardQuit
+  , isButtonEdge
+  ) where
+
+import Data.Bits ((.&.))
+import Data.IORef (readIORef)
+import qualified Data.Text as T
+import Data.Text (Text)
+import qualified Data.Text.Foreign as TF
+import Data.Word (Word32)
+import Foreign.C.Types (CFloat)
+import Foreign.Marshal.Alloc (alloca)
+import Foreign.Ptr (Ptr, nullPtr)
+import Foreign.Storable (Storable (..))
+import GHC.Records.Compat (getField)
+import SDL3.Sys.Bindgen.Runtime.CBool qualified as CBool
+import SDL3.Sys.Bindgen.Runtime.PtrConst qualified as PtrConst
+import NanoUI
+  ( Input (..)
+  , Key (..)
+  , Modifiers (..)
+  , DropEvent (..)
+  , DropType (..)
+  , V2 (..)
+  , appendInputKey
+  , v2Add
+  )
+import NanoUI.Input (MouseButton (..), appendDropEvent, applyMouseButton)
+import NanoUI.Sdl.Display (refreshEventType)
+import SDL3.Sys.Bindgen.Events
+  ( SDL_Event (..)
+  , SDL_EventType (..)
+  , SDL_KeyboardEvent
+  )
+import SDL3.Sys.Bindgen.Events qualified as Events
+import SDL3.Sys.Events (pollEventSafe, waitEventSafe, waitEventTimeoutSafe)
+import SDL3.Sys.Bindgen.Keycode
+  ( SDL_Keycode (..)
+  , SDL_Keymod (..)
+  , sDLK_BACKSPACE
+  , sDLK_DELETE
+  , sDLK_DOWN
+  , sDLK_END
+  , sDLK_ESCAPE
+  , sDLK_HOME
+  , sDLK_LEFT
+  , sDLK_RETURN
+  , sDLK_RIGHT
+  , sDLK_TAB
+  , sDLK_UP
+  , sDL_KMOD_ALT
+  , sDL_KMOD_CTRL
+  , sDL_KMOD_SHIFT
+  )
+import SDL3.Sys.Bindgen.Mouse (sDL_BUTTON_LEFT, sDL_BUTTON_RIGHT)
+import SDL3.Sys.Bindgen.Stdinc (Sint32 (..), Uint32 (..))
+import SDL3.Sys.Keyboard (getModState)
+
+data SdlEvent
+  = EvQuit
+  | EvResize Int Int
+  | EvDisplayScale
+  | EvKey Key Modifiers
+  | EvText Text Modifiers
+  | EvMouseMotion V2 Modifiers
+  | EvMousePress V2 Modifiers
+  | EvMouseRelease V2 Modifiers
+  | EvMouseRightPress V2 Modifiers
+  | EvMouseRightRelease V2 Modifiers
+  | EvScroll V2
+  | EvDrop DropEvent
+  | EvRefresh
+  | EvWindowRedraw
+  deriving (Eq, Show)
+
+-- | Drain every pending event, oldest first.
+pollEvents :: IO [SdlEvent]
+pollEvents =
+  alloca $ \(p :: Ptr SDL_Event) -> do
+    refreshTy <- readIORef refreshEventType
+    let drain acc = do
+          got <- pollEventSafe p
+          if got
+            then decodeEvent refreshTy p >>= \ev -> drain (maybe acc (: acc) ev)
+            else pure (reverse acc)
+    drain []
+
+waitEvent :: IO (Maybe SdlEvent)
+waitEvent =
+  alloca $ \p -> do
+    got <- waitEventSafe p
+    if got then readIORef refreshEventType >>= \ty -> decodeEvent ty p else pure Nothing
+
+waitEventTimeout :: Int -> IO (Maybe SdlEvent)
+waitEventTimeout ms =
+  alloca $ \p -> do
+    got <- waitEventTimeoutSafe p (fromIntegral ms)
+    if got then readIORef refreshEventType >>= \ty -> decodeEvent ty p else pure Nothing
+
+-- | Translate one SDL event, given the refresh event type; 'Nothing' for
+-- events the UI ignores.
+decodeEvent :: Word32 -> Ptr SDL_Event -> IO (Maybe SdlEvent)
+decodeEvent refreshTy p = do
+  Uint32 w <- peek p.type'
+  if refreshTy /= 0 && w == refreshTy
+    then pure (Just EvRefresh)
+    else case SDL_EventType (fromIntegral w) of
+      Events.SDL_EVENT_QUIT -> pure (Just EvQuit)
+      Events.SDL_EVENT_WINDOW_RESIZED -> Just <$> windowResized p
+      -- Pixel size changes are ignored here; syncDisplay re-queries logical size.
+      Events.SDL_EVENT_WINDOW_PIXEL_SIZE_CHANGED -> pure (Just EvDisplayScale)
+      Events.SDL_EVENT_WINDOW_DISPLAY_SCALE_CHANGED -> pure (Just EvDisplayScale)
+      -- The window manager damaged our window surface (occlusion, compositor
+      -- effects, restore). The backbuffer contents are gone; the next present
+      -- must be full or stale regions flash.
+      Events.SDL_EVENT_WINDOW_EXPOSED -> pure (Just EvWindowRedraw)
+      Events.SDL_EVENT_WINDOW_RESTORED -> pure (Just EvWindowRedraw)
+      Events.SDL_EVENT_KEY_DOWN -> keyDown p
+      Events.SDL_EVENT_TEXT_INPUT -> textInput p
+      Events.SDL_EVENT_MOUSE_MOTION -> Just <$> mouseMotion p
+      Events.SDL_EVENT_MOUSE_BUTTON_DOWN -> mouseButton p True
+      Events.SDL_EVENT_MOUSE_BUTTON_UP -> mouseButton p False
+      Events.SDL_EVENT_MOUSE_WHEEL -> Just <$> mouseWheel p
+      Events.SDL_EVENT_DROP_FILE -> Just <$> dropEvent p DropFile
+      Events.SDL_EVENT_DROP_TEXT -> Just <$> dropEvent p DropText
+      Events.SDL_EVENT_DROP_BEGIN -> Just <$> dropEvent p DropBegin
+      Events.SDL_EVENT_DROP_COMPLETE -> Just <$> dropEvent p DropComplete
+      Events.SDL_EVENT_DROP_POSITION -> Just <$> dropEvent p DropPosition
+      _ -> pure Nothing
+
+keyDown :: Ptr SDL_Event -> IO (Maybe SdlEvent)
+keyDown p = do
+  ke <- peek p.key
+  let mods = keyModifiers ke
+      code = fromIntegral (getField @"key" ke :: SDL_Keycode) :: Word32
+      repeating = CBool.toBool (getField @"repeat" ke)
+  pure $
+    case mapSpecialKey code of
+      Just k
+        | not repeating || isRepeatableKey k -> Just (EvKey k mods)
+        | otherwise -> Nothing
+      Nothing
+        -- Ctrl chords produce no text-input event; report the printable key
+        -- symbol (SDL folds Shift into it, so Ctrl+Shift+= arrives as '+').
+        | modCtrl mods && code >= 32 && code <= 126 ->
+            Just (EvText (T.singleton (toEnum (fromIntegral code))) mods)
+        | otherwise -> Nothing
+
+textInput :: Ptr SDL_Event -> IO (Maybe SdlEvent)
+textInput p = do
+  te <- peek p.text
+  mods <- peekModifiers
+  let textPtr = PtrConst.unsafeToPtr (getField @"text" te)
+  if textPtr == nullPtr
+    then pure Nothing
+    else do
+      txt <- TF.peekCString textPtr
+      pure (if T.null txt then Nothing else Just (EvText txt mods))
+
+mouseMotion :: Ptr SDL_Event -> IO SdlEvent
+mouseMotion p = do
+  me <- peek p.motion
+  mods <- peekModifiers
+  let x = getField @"x" me :: CFloat
+      y = getField @"y" me :: CFloat
+  pure (EvMouseMotion (V2 (realToFrac x) (realToFrac y)) mods)
+
+mouseButton :: Ptr SDL_Event -> Bool -> IO (Maybe SdlEvent)
+mouseButton p down = do
+  be <- peek p.button
+  mods <- peekModifiers
+  let x = getField @"x" be :: CFloat
+      y = getField @"y" be :: CFloat
+      pos = V2 (realToFrac x) (realToFrac y)
+      btn = getField @"button" be
+  pure $
+    if btn == fromIntegral sDL_BUTTON_LEFT
+      then Just (if down then EvMousePress pos mods else EvMouseRelease pos mods)
+      else
+        if btn == fromIntegral sDL_BUTTON_RIGHT
+          then Just (if down then EvMouseRightPress pos mods else EvMouseRightRelease pos mods)
+          else Nothing
+
+mouseWheel :: Ptr SDL_Event -> IO SdlEvent
+mouseWheel p = do
+  we <- peek p.wheel
+  let x = getField @"x" we :: CFloat
+      y = getField @"y" we :: CFloat
+  pure (EvScroll (V2 (realToFrac x) (negate (realToFrac y))))
+
+windowResized :: Ptr SDL_Event -> IO SdlEvent
+windowResized p = do
+  we <- peek p.window
+  let Sint32 w = getField @"data1" we
+      Sint32 h = getField @"data2" we
+  pure (EvResize (fromIntegral w) (fromIntegral h))
+
+dropEvent :: Ptr SDL_Event -> DropType -> IO SdlEvent
+dropEvent p ty = do
+  de <- peek p.drop
+  let x = getField @"x" de :: CFloat
+      y = getField @"y" de :: CFloat
+      at = Just (V2 (realToFrac x) (realToFrac y))
+      pos =
+        case ty of
+          DropPosition -> at
+          DropFile -> at
+          DropText -> at
+          _ -> Nothing
+      dataPtr = PtrConst.unsafeToPtr (getField @"data'" de)
+  payload <-
+    if dataPtr == nullPtr
+      then pure ""
+      else TF.peekCString dataPtr
+  pure (EvDrop (DropEvent ty pos payload))
+
+peekModifiers :: IO Modifiers
+peekModifiers = modFromKeymod <$> getModState
+
+modFromKeymod :: SDL_Keymod -> Modifiers
+modFromKeymod km =
+  let m = word32 km
+   in Modifiers
+        { modShift = m .&. word32 sDL_KMOD_SHIFT /= 0
+        , modCtrl = m .&. word32 sDL_KMOD_CTRL /= 0
+        , modAlt = m .&. word32 sDL_KMOD_ALT /= 0
+        }
+
+keyModifiers :: SDL_KeyboardEvent -> Modifiers
+keyModifiers ke = modFromKeymod (getField @"mod" ke)
+
+word32 :: Integral a => a -> Word32
+word32 = fromIntegral
+
+mapSpecialKey :: Word32 -> Maybe Key
+mapSpecialKey k
+  | k == word32 sDLK_ESCAPE = Just KeyEscape
+  | k == word32 sDLK_RETURN = Just KeyEnter
+  | k == word32 sDLK_TAB = Just KeyTab
+  | k == word32 sDLK_BACKSPACE = Just KeyBackspace
+  | k == word32 sDLK_DELETE = Just KeyDelete
+  | k == word32 sDLK_LEFT = Just KeyLeft
+  | k == word32 sDLK_RIGHT = Just KeyRight
+  | k == word32 sDLK_UP = Just KeyUp
+  | k == word32 sDLK_DOWN = Just KeyDown
+  | k == word32 sDLK_HOME = Just KeyHome
+  | k == word32 sDLK_END = Just KeyEnd
+  | otherwise = Nothing
+
+isRepeatableKey :: Key -> Bool
+isRepeatableKey KeyBackspace = True
+isRepeatableKey KeyDelete = True
+isRepeatableKey KeyLeft = True
+isRepeatableKey KeyRight = True
+isRepeatableKey KeyUp = True
+isRepeatableKey KeyDown = True
+isRepeatableKey KeyHome = True
+isRepeatableKey KeyEnd = True
+isRepeatableKey _ = False
+
+applyEvent :: Input -> SdlEvent -> Input
+applyEvent inp ev =
+  case ev of
+    EvQuit -> inp
+    EvDisplayScale -> inp
+    EvResize _ _ -> inp
+    EvKey k mods -> inp {inputKeys = appendInputKey k (inputKeys inp), inputModifiers = mods}
+    EvText txt mods ->
+      inp {inputChars = inputChars inp <> txt, inputModifiers = mods}
+    EvMouseMotion pos mods ->
+      inp {inputMousePos = pos, inputModifiers = mods}
+    EvMousePress pos mods ->
+      (applyMouseButton MouseLeft True inp) {inputMousePos = pos, inputModifiers = mods}
+    EvMouseRelease pos mods ->
+      (applyMouseButton MouseLeft False inp) {inputMousePos = pos, inputModifiers = mods}
+    EvMouseRightPress pos mods ->
+      (applyMouseButton MouseRight True inp) {inputMousePos = pos, inputModifiers = mods}
+    EvMouseRightRelease pos mods ->
+      (applyMouseButton MouseRight False inp) {inputMousePos = pos, inputModifiers = mods}
+    EvScroll delta -> inp {inputScroll = v2Add (inputScroll inp) delta}
+    EvDrop dropEv -> inp {inputDrops = appendDropEvent dropEv (inputDrops inp)}
+    EvRefresh -> inp {inputWindowRedraw = True}
+    EvWindowRedraw -> inp {inputWindowRedraw = True}
+
+isButtonEdge :: SdlEvent -> Bool
+isButtonEdge ev =
+  case ev of
+    EvMousePress _ _ -> True
+    EvMouseRelease _ _ -> True
+    EvMouseRightPress _ _ -> True
+    EvMouseRightRelease _ _ -> True
+    _ -> False
+
+isHardQuit :: SdlEvent -> Bool
+isHardQuit ev =
+  case ev of
+    EvText txt mods -> (txt == "c" && modCtrl mods) || txt == "\ETX"
+    _ -> False
diff --git a/lib/NanoUI/Sdl/NanoUIFont.hs b/lib/NanoUI/Sdl/NanoUIFont.hs
new file mode 100644
--- /dev/null
+++ b/lib/NanoUI/Sdl/NanoUIFont.hs
@@ -0,0 +1,16 @@
+-- | Font selection for the SDL backend.
+module NanoUI.Sdl.NanoUIFont
+  ( NanoUIFont (..)
+  ) where
+
+-- | SDL text font selection.
+-- 'DefaultFont' uses the embedded Inter subset.
+-- 'FontSearch' tries each family name against the platform font directories
+-- (walked recursively at runtime). When nothing matches, the embedded font
+-- is used as fallback.
+-- 'FontFilePath' loads that file and does not search.
+data NanoUIFont
+  = DefaultFont
+  | FontSearch [String]
+  | FontFilePath FilePath
+  deriving (Eq, Show)
diff --git a/lib/NanoUI/Sdl/Render.hs b/lib/NanoUI/Sdl/Render.hs
new file mode 100644
--- /dev/null
+++ b/lib/NanoUI/Sdl/Render.hs
@@ -0,0 +1,243 @@
+module NanoUI.Sdl.Render
+  ( RenderBatch
+  , newRenderBatch
+  , destroyRenderBatch
+  , flushRenderBatch
+  , renderDrawDataPass
+  , snapDamage
+  ) where
+
+import NanoUI.Sdl.Image (ImageAtlas, lookupImage)
+
+import Control.Monad (void, when)
+import Data.Bits (shiftR, (.&.))
+import Data.IORef (IORef, newIORef, readIORef, writeIORef)
+import Data.Primitive.PrimArray (indexPrimArray, sizeofPrimArray)
+import Data.Word (Word8)
+import Foreign.C.Types (CFloat (..), CInt (..))
+import Foreign.ForeignPtr (withForeignPtr)
+import Foreign.Marshal.Alloc (free, malloc)
+import Foreign.Ptr (Ptr, nullPtr)
+import Foreign.Storable (poke)
+import NanoUI (Color (..), Rect (..), rectIntersect)
+import NanoUI.Testing
+  ( Damage (..)
+  , DrawCmd (..)
+  , DrawData (..)
+  , LayerSlice (..)
+  , damageIsEmpty
+  , glyphAtlasTextureId
+  )
+import SDL3.Sys.Bindgen.Rect (SDL_Rect (..))
+import SDL3.Sys.Bindgen.Render (SDL_Renderer, SDL_Texture)
+import SDL3.Sys.Bindgen.Runtime.PtrConst qualified as PtrConst
+import SDL3.Sys.Render
+  ( renderClearSafe
+  , setRenderClipRect
+  , setRenderDrawColorSafe
+  )
+
+data ClipState
+  = ClipNone
+  | ClipKey {-# UNPACK #-} !Int {-# UNPACK #-} !Int {-# UNPACK #-} !Int {-# UNPACK #-} !Int
+  deriving (Eq)
+
+{-# INLINE snapDamage #-}
+snapDamage :: Float -> Damage -> Damage
+snapDamage _ DamageFull = DamageFull
+snapDamage scale (DamageClip (Rect x y w h)) =
+  let px = fromIntegral (floor (x * scale) :: Int) / scale
+      py = fromIntegral (floor (y * scale) :: Int) / scale
+      pw = fromIntegral (ceiling ((x + w) * scale) :: Int) / scale - px
+      ph = fromIntegral (ceiling ((y + h) * scale) :: Int) / scale - py
+   in DamageClip (Rect px py pw ph)
+
+{-# INLINE toClipKey #-}
+toClipKey :: Rect -> ClipState
+toClipKey (Rect x y w h) =
+  let px = floor x :: Int
+      py = floor y :: Int
+      x1 = ceiling (x + w) :: Int
+      y1 = ceiling (y + h) :: Int
+   in ClipKey px py (max 1 (x1 - px)) (max 1 (y1 - py))
+
+applyClipState :: RenderBatch -> IORef ClipState -> Ptr SDL_Renderer -> ClipState -> IO ()
+applyClipState batch ref ren next = do
+  prev <- readIORef ref
+  when (prev /= next) $ do
+    flushRenderBatch batch
+    writeIORef ref next
+    void $ case next of
+      ClipNone -> setRenderClipRect ren (PtrConst.unsafeFromPtr nullPtr)
+      ClipKey px py pw ph -> do
+        let rect = rbClipRect batch
+        poke rect (SDL_Rect (fromIntegral px) (fromIntegral py) (fromIntegral pw) (fromIntegral ph))
+        setRenderClipRect ren (PtrConst.unsafeFromPtr rect)
+
+-- | Draw every command in layer-slice order, clipped to its own rect and to
+-- the damage. A full repaint with a clear colour clears the target first.
+renderDrawDataPass :: RenderBatch -> Ptr SDL_Renderer -> Maybe Color -> DrawData -> ImageAtlas -> Ptr SDL_Texture -> Damage -> IO ()
+renderDrawDataPass batch ren mClear drawData images glyphTex damage =
+  when (not (damageIsEmpty damage)) $ do
+    clipRef <- newIORef ClipNone
+    void $ setRenderClipRect ren (PtrConst.unsafeFromPtr nullPtr)
+    case (mClear, damage) of
+      (Just clearColor, DamageFull) -> do
+        let (cr, cg, cb, ca) = unpackColor clearColor
+        void $ setRenderDrawColorSafe ren cr cg cb ca
+        void $ renderClearSafe ren
+      (Just _clearColor, DamageClip r) ->
+        -- The draw list starts a clip frame with its own window backdrop,
+        -- so the clip needs no clear here.
+        applyClipState batch clipRef ren (toClipKey r)
+      (Nothing, DamageClip r) -> applyClipState batch clipRef ren (toClipKey r)
+      (Nothing, DamageFull) -> pure ()
+    let clip = case damage of
+          DamageFull -> Nothing
+          DamageClip r -> Just r
+        vc = drawVertexCount drawData
+        cmds = drawCommands drawData
+        slices = drawLayerSlices drawData
+    withForeignPtr (drawVertices drawData) $ \vp ->
+      withForeignPtr (drawIndices drawData) $ \ip ->
+        let goLy !li
+              | li >= sizeofPrimArray slices = pure ()
+              | otherwise = do
+                  let LayerSlice off cnt = indexPrimArray slices li
+                      goCmd !j
+                        | j >= cnt = pure ()
+                        | otherwise = do
+                            drawCmd batch ren vp vc ip images glyphTex clip clipRef (indexPrimArray cmds (off + j))
+                            goCmd (j + 1)
+                  goCmd 0
+                  goLy (li + 1)
+         in goLy 0
+    applyClipState batch clipRef ren ClipNone
+
+{-# INLINE drawCmd #-}
+drawCmd ::
+  RenderBatch ->
+  Ptr SDL_Renderer ->
+  Ptr Word8 ->
+  Int ->
+  Ptr Word8 ->
+  ImageAtlas ->
+  Ptr SDL_Texture ->
+  Maybe Rect ->
+  IORef ClipState ->
+  DrawCmd ->
+  IO ()
+drawCmd batch ren vp vc ip images glyphTex mDamage clipRef cmd = do
+  let !count = fromIntegral (cmdIndexCount cmd)
+      !cmdRect = Rect (cmdClipX cmd) (cmdClipY cmd) (cmdClipW cmd) (cmdClipH cmd)
+      !cmdOpen = cmdClipW cmd >= 1e8 || cmdClipH cmd >= 1e8
+      live = case (mDamage, cmdOpen) of
+        (Nothing, _) -> Just cmdRect
+        (Just dmg, True) -> Just dmg
+        (Just dmg, False) -> rectIntersect dmg cmdRect
+  when (count >= 3) $
+    case live of
+      Nothing -> pure ()
+      Just clip -> do
+        if cmdOpen && mDamage == Nothing
+          then applyClipState batch clipRef ren ClipNone
+          else applyClipState batch clipRef ren (toClipKey clip)
+        let !start = fromIntegral (cmdIndexOffset cmd)
+            !texId = cmdTextureId cmd
+        tex <-
+          if texId == glyphAtlasTextureId
+            then pure glyphTex
+            else if texId > 0
+              then maybe nullPtr id <$> lookupImage images texId
+              else pure nullPtr
+        batchDrawRange batch vp vc ip start count tex mDamage
+
+{-# INLINE unpackColor #-}
+unpackColor :: Color -> (Word8, Word8, Word8, Word8)
+unpackColor (Color w) =
+  ( fromIntegral ((w `shiftR` 24) .&. 0xFF)
+  , fromIntegral ((w `shiftR` 16) .&. 0xFF)
+  , fromIntegral ((w `shiftR` 8) .&. 0xFF)
+  , fromIntegral (w .&. 0xFF)
+  )
+
+-- | The C batch and a clip rect it passes to SDL, both owned for the session.
+data RenderBatch = RenderBatch
+  { rbBatch :: !(Ptr ())
+  , rbClipRect :: !(Ptr SDL_Rect)
+  }
+
+-- | Create a persistent render batch. Reusing one batch across frames avoids
+-- a C calloc/free pair per presented frame; flush after each render pass.
+newRenderBatch :: Ptr SDL_Renderer -> IO RenderBatch
+newRenderBatch ren = do
+  p <- batchCreate ren
+  if p == nullPtr
+    then fail "nano_ui_batch_create failed"
+    else RenderBatch p <$> malloc
+
+destroyRenderBatch :: RenderBatch -> IO ()
+destroyRenderBatch batch = do
+  batchDestroy (rbBatch batch)
+  free (rbClipRect batch)
+
+flushRenderBatch :: RenderBatch -> IO ()
+flushRenderBatch batch = batchFlush (rbBatch batch)
+
+batchDrawRange ::
+  RenderBatch ->
+  Ptr Word8 ->
+  Int ->
+  Ptr Word8 ->
+  Int ->
+  Int ->
+  Ptr SDL_Texture ->
+  Maybe Rect ->
+  IO ()
+batchDrawRange batch verts vc indices start n tex mDmg =
+  batchDrawRangeC
+    (rbBatch batch)
+    verts
+    (ci vc)
+    indices
+    (ci start)
+    (ci n)
+    tex
+    hasDmg
+    (cf dx)
+    (cf dy)
+    (cf dw)
+    (cf dh)
+  where
+    ci = fromIntegral
+    (hasDmg, dx, dy, dw, dh) = case mDmg of
+      Nothing -> (0, 0, 0, 0, 0)
+      Just (Rect x y w h) -> (1, x, y, w, h)
+
+cf :: Float -> CFloat
+cf = realToFrac
+
+foreign import ccall unsafe "nano_ui_batch_create"
+  batchCreate :: Ptr SDL_Renderer -> IO (Ptr ())
+
+foreign import ccall unsafe "nano_ui_batch_destroy"
+  batchDestroy :: Ptr () -> IO ()
+
+foreign import ccall unsafe "nano_ui_batch_flush"
+  batchFlush :: Ptr () -> IO ()
+
+foreign import ccall unsafe "nano_ui_batch_draw_range"
+  batchDrawRangeC ::
+    Ptr () ->
+    Ptr Word8 ->
+    CInt ->
+    Ptr Word8 ->
+    CInt ->
+    CInt ->
+    Ptr SDL_Texture ->
+    CInt ->
+    CFloat ->
+    CFloat ->
+    CFloat ->
+    CFloat ->
+    IO ()
diff --git a/lib/NanoUI/Sdl/Runner.hs b/lib/NanoUI/Sdl/Runner.hs
new file mode 100644
--- /dev/null
+++ b/lib/NanoUI/Sdl/Runner.hs
@@ -0,0 +1,283 @@
+{-# LANGUAGE DataKinds #-}
+
+-- | SDL3 draw path: retained damage updates or direct continuous presentation.
+module NanoUI.Sdl.Runner
+  ( sdlDrawFrame
+  , drawReduceEff
+  , askSdlDebug
+  , setSdlUiFont
+  ) where
+
+import Control.Exception (finally, mask_)
+import Control.Monad (unless, void, when)
+import Data.IORef (IORef, readIORef, writeIORef)
+import Data.Typeable (Typeable)
+import GHC.Clock (getMonotonicTime)
+import NanoUI
+  ( Input (..)
+  , NanoUI
+  , Size (..)
+  , V2 (..)
+  , themeWindow
+  )
+import Effectful (Eff, IOE, type (:>))
+import NanoUI.Testing
+  ( Context
+  , Damage (..)
+  , DrawData (..)
+  , Ui
+  , askHost
+  , ctxPaintFull
+  , ctxTheme
+  , damageFull
+  , damageIsEmpty
+  , drawCmdCount
+  , markDirty
+  , runEff
+  , runFrameEff
+  , runFrameReduceEff
+  , takeDamage
+  , uiIO
+  )
+import NanoUI.Debug (CoreDebugSnapshot (..), noteDebugPresent, noteDebugSkip, refreshDebugSnapshot)
+import NanoUI.Sdl.Debug
+  ( SdlDebugSampler (..)
+  , SdlDebugSnapshot (..)
+  , emptySdlDebug
+  , traceFrame
+  )
+import NanoUI.Sdl.Display (queryMouseWindowPos, queryWindowLogicalSize)
+import NanoUI.Sdl.Font
+  ( fontSourceLabel
+  , glyphAtlasTexture
+  , sdlFontCacheSource
+  , prepareGlyphAtlasForFrame
+  , takeGlyphAtlasResetFlag
+  )
+import NanoUI.Sdl.NanoUIFont (NanoUIFont)
+import NanoUI.Sdl.Render (flushRenderBatch, renderDrawDataPass, snapDamage)
+import NanoUI.Sdl.Window (SdlEnv (..))
+import Foreign.Marshal.Alloc (alloca)
+import Foreign.Ptr (Ptr, nullPtr)
+import Foreign.Storable (peek)
+import qualified NanoUI.Sdl.Image as SdlImage
+import SDL3.Sys.Bindgen.Blendmode (sDL_BLENDMODE_NONE)
+import SDL3.Sys.Bindgen.Pixels qualified as Pixels
+import SDL3.Sys.Bindgen.Render (SDL_Texture)
+import SDL3.Sys.Bindgen.Render qualified as Render
+import SDL3.Sys.Bindgen.Runtime.PtrConst qualified as PtrConst
+import SDL3.Sys.Render
+  ( createTexture
+  , destroyTexture
+  , getRenderOutputSize
+  , renderPresentSafe
+  , renderTexture
+  , setRenderClipRect
+  , setRenderScale
+  , setRenderTarget
+  , setTextureBlendMode
+  )
+
+sdlDrawFrame :: Context -> NanoUI () -> SdlEnv -> Input -> Bool -> IO (Bool, Input)
+sdlDrawFrame ctx ui env inp forceFull =
+  drawFrameWith ctx env inp forceFull $ do
+    (_, _, drawData, dirtyAfterUi) <- runFrameEff runEff ctx inp ui
+    pure (drawData, dirtyAfterUi)
+
+-- | Both application styles share atlas maintenance, retain preparation,
+-- timing, and presentation. Only evaluation of the UI differs.
+drawFrameWith :: Context -> SdlEnv -> Input -> Bool -> IO (DrawData, Bool) -> IO (Bool, Input)
+drawFrameWith ctx env inp forceFull evaluateUi = do
+  SdlImage.syncImageAtlas (sdlRenderer env) (sdlImages env) ctx
+  (tex, presentFull) <- prepareRetain ctx env inp forceFull
+  t0 <- getMonotonicTime
+  (drawData, dirtyAfterUi) <- evaluateUi
+  t1 <- getMonotonicTime
+  finishDraw ctx env inp tex presentFull t0 t1 drawData dirtyAfterUi
+
+-- | Choose the render target and whether to repaint everything. Must run
+-- before the frame so paint can cull to damage for retained partial updates.
+prepareRetain :: Context -> SdlEnv -> Input -> Bool -> IO (Ptr SDL_Texture, Bool)
+prepareRetain ctx env inp forceFull = do
+  -- Glyph-atlas maintenance before any quad is recorded: if the atlas ran
+  -- out of space during the previous frame, reset it now (re-warming the
+  -- base fonts) so a reset can never wipe the texture underneath
+  -- already-recorded text mid-frame.
+  prepareGlyphAtlasForFrame (sdlGlyphAtlas env)
+  scale <- readIORef (sdlScaleRef env)
+  let Size lw lh = inputWindowSize inp
+      pw = max 1 (round (lw * scale))
+      ph = max 1 (round (lh * scale))
+  -- Continuous sessions repaint every pixel, so retaining and copying a
+  -- second framebuffer only adds a target switch and a full-window blit.
+  -- A null target selects the window backbuffer directly. Direct drawing is
+  -- equivalent to the retained blit only at the same pixel dimensions;
+  -- content scale and window pixel density can differ.
+  direct <-
+    if sdlContinuous env
+      then
+        alloca $ \wp ->
+          alloca $ \hp -> do
+            ok <- getRenderOutputSize (sdlRenderer env) wp hp
+            ow <- peek wp
+            oh <- peek hp
+            pure (ok && fromIntegral ow == pw && fromIntegral oh == ph)
+      else pure False
+  (tex, retainNew) <-
+    if direct
+      then pure (nullPtr, False)
+      else ensureRetain env pw ph scale
+  let presentFull = forceFull || retainNew || sdlContinuous env || inputWindowRedraw inp
+  writeIORef (ctxPaintFull ctx) presentFull
+  pure (tex, presentFull)
+
+drawReduceEff ::
+  (IOE :> es, Typeable msg, Eq model) =>
+  (forall x. Eff es x -> IO x) ->
+  (msg -> model -> model) ->
+  IORef model ->
+  (model -> Eff (Ui : es) ()) ->
+  Context ->
+  SdlEnv ->
+  Input ->
+  Bool ->
+  IO (Bool, Input)
+drawReduceEff unlift update modelRef view ctx env inp forceFull =
+  drawFrameWith ctx env inp forceFull $ do
+    m <- readIORef modelRef
+    (_, m', _, drawData, dirtyAfterUi) <- runFrameReduceEff unlift update ctx inp m view
+    writeIORef modelRef m'
+    pure (drawData, dirtyAfterUi)
+
+finishDraw :: Context -> SdlEnv -> Input -> Ptr SDL_Texture -> Bool -> Double -> Double -> DrawData -> Bool -> IO (Bool, Input)
+finishDraw ctx env inp tex presentFull t0 t1 drawData dirtyAfterUi = do
+  let uiMs = (t1 - t0) * 1000
+  scale <- readIORef (sdlScaleRef env)
+  dmg0 <- takeDamage ctx
+  let Size lw lh = inputWindowSize inp
+  -- Frame damage from writeDamage is authoritative: a live animation whose
+  -- key is out of view or scroll-clipped produces empty damage, and forcing
+  -- DamageFull here would turn every skip frame into a full present. A
+  -- window redraw event (expose/restore) is the exception: the backbuffer
+  -- is gone, so the next present must be full.
+  let damage =
+        if presentFull
+          then DamageFull
+          else snapDamage scale dmg0
+  writeIORef (sdlLastPresented env) False
+  -- A glyph-atlas reset or exhaustion during the UI pass means quads
+  -- recorded before that point hold stale (or unplaceable) UVs. Drop the
+  -- frame instead of presenting it: the screen keeps the previous valid
+  -- frame, 'damageFull' forces a full repaint, and
+  -- 'prepareGlyphAtlasForFrame' resets the atlas before the next frame
+  -- records any quads, so text never flickers or vanishes for a frame.
+  atlasReset <- takeGlyphAtlasResetFlag (sdlGlyphAtlas env)
+  if atlasReset || damageIsEmpty damage || lw <= 0 || lh <= 0
+    then do
+      when atlasReset $ do
+        damageFull ctx
+        markDirty ctx
+      noteDebugSkip (sdsSampler (sdlDebug env))
+      pure (atlasReset || dirtyAfterUi, inp)
+    else do
+      -- A null texture draws full-repaint sessions straight to the window.
+      okBegin <- setRenderTarget (sdlRenderer env) tex
+      okScale <- setRenderScale (sdlRenderer env) scale scale
+      unless (okBegin && okScale) $ fail "SDL_SetRenderTarget/Scale failed"
+      theme <- readIORef (ctxTheme ctx)
+      glyphTex <- glyphAtlasTexture (sdlGlyphAtlas env)
+      -- Persistent batch created once per session (sdlBatch): no C
+      -- calloc/free pair per presented frame. Flush unconditionally so an
+      -- aborted pass cannot leak pending geometry into the next frame.
+      --
+      -- Full repaints clear the target, including bare backdrop regions.
+      -- Partial updates preserve the undamaged part of the retained texture.
+      let batch = sdlBatch env
+      renderDrawDataPass
+        batch
+        (sdlRenderer env)
+        (if damage == DamageFull then Just (themeWindow theme) else Nothing)
+        drawData
+        (sdlImages env)
+        glyphTex
+        damage
+        `finally` flushRenderBatch batch
+      t2 <- getMonotonicTime
+      -- Damage limits updates to the retained texture, not the final copy:
+      -- SDL leaves the window backbuffer undefined after each present.
+      -- Restore the window's pixel coordinate system before polling events.
+      -- Retained sessions do this as part of their final texture copy.
+      okBlit <-
+        if tex == nullPtr
+          then setRenderScale (sdlRenderer env) 1 1
+          else do
+            okTarget <- setRenderTarget (sdlRenderer env) nullPtr
+            okClip <- setRenderClipRect (sdlRenderer env) (PtrConst.unsafeFromPtr nullPtr)
+            void $ setRenderScale (sdlRenderer env) 1 1
+            okCopy <- renderTexture (sdlRenderer env) tex (PtrConst.unsafeFromPtr nullPtr) (PtrConst.unsafeFromPtr nullPtr)
+            pure (okTarget && okClip && okCopy)
+      unless okBlit $ fail "SDL window presentation preparation failed"
+      void $ renderPresentSafe (sdlRenderer env)
+      t3 <- getMonotonicTime
+      let renderMs = (t2 - t1) * 1000
+          presentMs = (t3 - t2) * 1000
+          frameMs = (t3 - t0) * 1000
+      noteDebugPresent (sdsSampler (sdlDebug env)) uiMs renderMs presentMs frameMs
+        (drawVertexCount drawData) (drawIndexCount drawData) (drawCmdCount drawData)
+      writeIORef (sdlLastPresented env) True
+      pure (dirtyAfterUi, inp)
+
+ensureRetain :: SdlEnv -> Int -> Int -> Float -> IO (Ptr SDL_Texture, Bool)
+ensureRetain env w h scale = do
+  (tex, ow, oh, oldScale) <- readIORef (sdlRetain env)
+  let scaleChanged = abs (oldScale - scale) > 0.001
+  if tex /= nullPtr && ow == w && oh == h
+    then do
+      when scaleChanged $ writeIORef (sdlRetain env) (tex, w, h, scale)
+      -- Same pixel size after a DPI change still holds the old present.
+      pure (tex, scaleChanged)
+    else mask_ $ do
+      -- Allocate before replacing: failure leaves the owned texture valid.
+      tex' <- createTexture (sdlRenderer env) Pixels.SDL_PIXELFORMAT_RGBA32 Render.SDL_TEXTUREACCESS_TARGET (fromIntegral w) (fromIntegral h)
+      when (tex' == nullPtr) $ fail "SDL_CreateTexture(retain) failed"
+      void $ setTextureBlendMode tex' (fromIntegral sDL_BLENDMODE_NONE)
+      writeIORef (sdlRetain env) (tex', w, h, scale)
+      unless (tex == nullPtr) $ destroyTexture tex
+      pure (tex', True)
+
+askSdlDebug :: Ui :> es => Eff es SdlDebugSnapshot
+askSdlDebug = do
+  menv <- askHost @SdlEnv
+  case menv of
+    Nothing -> pure emptySdlDebug
+    Just env -> uiIO $ do
+      let sampler = sdlDebug env
+      -- The display is queried only when the snapshot refreshes.
+      refreshDebugSnapshot (sdsSampler sampler) (sdsSnapshot sampler) $ \core -> do
+        scale <- readIORef (sdlScaleRef env)
+        fontSource <- sdlFontCacheSource (sdlFontCache env)
+        Size ww wh <- queryWindowLogicalSize (sdlWindow env)
+        V2 mx my <- queryMouseWindowPos
+        let snap =
+              SdlDebugSnapshot
+                { dbgCore = core {dbgWinW = ww, dbgWinH = wh, dbgMouseX = mx, dbgMouseY = my}
+                , dbgScale = scale
+                , dbgFontPath = fontSourceLabel fontSource
+                , dbgRenderer = sdlRendererName env
+                , dbgVsync = sdlVsync env
+                , dbgRefreshHz = round (1 / sdlRefreshPeriod env)
+                }
+        when (sdsTrace sampler) (traceFrame snap)
+        pure snap
+
+-- | Request a UI font family. The SDL display thread resolves and applies it
+-- before the next frame (see 'NanoUI.Sdl.Window.syncDisplay'), rebuilding the
+-- glyph atlas and text resolver. A no-op on non-SDL hosts.
+setSdlUiFont :: Ui :> es => NanoUIFont -> Eff es ()
+setSdlUiFont font = do
+  menv <- askHost @SdlEnv
+  case menv of
+    Nothing -> pure ()
+    Just env -> uiIO $ do
+      cur <- readIORef (sdlFontRequestRef env)
+      when (cur /= font) $ writeIORef (sdlFontRequestRef env) font
diff --git a/lib/NanoUI/Sdl/Session.hs b/lib/NanoUI/Sdl/Session.hs
new file mode 100644
--- /dev/null
+++ b/lib/NanoUI/Sdl/Session.hs
@@ -0,0 +1,187 @@
+{-# LANGUAGE DataKinds #-}
+
+-- | SDL window session loop: event poll, resize sync, frame present.
+module NanoUI.Sdl.Session
+  ( runSdlSession
+  ) where
+
+import Control.Exception (bracket)
+import Control.Monad (void, when)
+import Data.IORef (newIORef, readIORef, writeIORef)
+import NanoUI (Input (..), emptyInput)
+import NanoUI.Sdl.Debug (SdlDebugSampler (..))
+import NanoUI.Input (clearEphemeral)
+import NanoUI.Runner
+  ( SessionDriver (..)
+  , newDrawingLock
+  , runSessionLoop
+  , shouldRedrawFrame
+  , tryWithDrawingLock
+  )
+import NanoUI.Testing
+  ( Context
+  , clearDirty
+  )
+import NanoUI.Sdl.Cursor (syncPointerCursor)
+import NanoUI.Sdl.Input
+  ( SdlEvent (..)
+  , applyEvent
+  , isButtonEdge
+  , isHardQuit
+  , pollEvents
+  , waitEvent
+  , waitEventTimeout
+  )
+import NanoUI.Sdl.Display (installResizeWatch)
+import NanoUI.Sdl.Window (SdlEnv (..), SdlOptions (..), syncDisplay, withSdl)
+import SDL3.Sys.Bindgen.Blendmode (sDL_BLENDMODE_BLEND)
+import SDL3.Sys.Render (setRenderDrawBlendModeSafe, setRenderVSync)
+
+
+runSdlSession ::
+  SdlOptions ->
+  Context ->
+  (SdlEnv -> IO ()) ->
+  (Input -> Bool) ->
+  (Context -> SdlEnv -> Input -> Bool -> IO (Bool, Input)) ->
+  IO ()
+runSdlSession options ctx setup shouldQuit drawFn =
+  withSdl options ctx $ \ctx0 env -> do
+    setup env
+    void $ setRenderDrawBlendModeSafe (sdlRenderer env) (fromIntegral sDL_BLENDMODE_BLEND)
+    ctxRef <- newIORef ctx0
+    prev <- newIORef emptyInput
+    drawing <- newDrawingLock
+    startupDone <- newIORef False
+    startupCatchup <- newIORef False
+    -- The resize watch presents with vsync off: Windows' modal size loop
+    -- cannot take the next drag step while a present waits for vblank. The
+    -- main loop turns vsync back on before its own frames.
+    vsyncPaused <- newIORef False
+    -- Window size the resize watch presented since the main loop last
+    -- decided whether to draw. That frame already covers the size change and
+    -- expose events the loop is about to see.
+    resizePresented <- newIORef Nothing
+    let onResize = do
+          void $
+            tryWithDrawingLock drawing $ do
+              liveCtx <- readIORef ctxRef
+              inp <- readIORef prev
+              scale0 <- readIORef (sdlScaleRef env)
+              (ctx', inpSynced) <- syncDisplay liveCtx env (clearEphemeral inp)
+              writeIORef ctxRef ctx'
+              done <- readIORef startupDone
+              if not done
+                then do
+                  writeIORef prev inpSynced
+                  writeIORef startupCatchup True
+                else do
+                  scale1 <- readIORef (sdlScaleRef env)
+                  if inputWindowSize inpSynced == inputWindowSize inp && scale1 == scale0
+                    then writeIORef prev inpSynced
+                    else do
+                      paused <- readIORef vsyncPaused
+                      when (sdlVsync env && not paused) $ do
+                        void $ setRenderVSync (sdlRenderer env) 0
+                        writeIORef vsyncPaused True
+                      (_, s) <- drawFn ctx' env inpSynced True
+                      writeIORef prev s
+                      writeIORef resizePresented (Just (inputWindowSize s))
+    -- A refresh wake (a finished file dialog) may postdate the watch's frame,
+    -- so it voids that frame's cover.
+    let noteWake evs = do
+          when (EvRefresh `elem` evs) $ writeIORef resizePresented Nothing
+          pure evs
+    let drainUntilQuiet c inp = do
+          pending <- pollEvents
+          (c', inp') <- syncDisplay c env (foldl' applyEvent inp pending)
+          if null pending
+            then pure (c', inp')
+            else drainUntilQuiet c' inp'
+    let inpSeed = emptyInput {inputWindowSize = sdlWindowSize options}
+    (ctx1, inp0) <- drainUntilQuiet ctx0 inpSeed
+    writeIORef ctxRef ctx1
+    scale0 <- readIORef (sdlScaleRef env)
+    (_, synced0) <- drawFn ctx1 env inp0 True
+    -- First present can apply DPI. Prev rects are empty on that frame.
+    -- Draw once more before idle or the Controls page stays stretched
+    -- until the first mouse move.
+    (ctx1b, inp0b) <- drainUntilQuiet ctx1 synced0
+    writeIORef ctxRef ctx1b
+    scaleSettle <- readIORef (sdlScaleRef env)
+    let paintedSize = inputWindowSize inp0b
+    (_, synced0b) <- drawFn ctx1b env inp0b True
+    clearDirty ctx1b
+    (ctx2, inp1) <- drainUntilQuiet ctx1b synced0b
+    writeIORef ctxRef ctx2
+    scale1 <- readIORef (sdlScaleRef env)
+    catchup <- readIORef startupCatchup
+    synced1 <-
+      if catchup || inputWindowSize inp1 /= paintedSize || abs (scale1 - scaleSettle) > 0.001 || abs (scaleSettle - scale0) > 0.001
+        then do
+          (_, s) <- drawFn ctx2 env inp1 True
+          clearDirty ctx2
+          pure s
+        else pure inp1
+    writeIORef startupCatchup False
+    writeIORef startupDone True
+    writeIORef prev synced1
+    let drv =
+          SessionDriver
+            { sdPollEvents    = pollEvents >>= noteWake
+            , sdWaitEvents    = \t -> do
+                -- Take the rest of the queue with the event that ended the
+                -- wait, so one pass sees a whole burst (a resize queues
+                -- several window events at once).
+                woke <- if t < 0 then waitEvent else waitEventTimeout t
+                case woke of
+                  Nothing -> pure []
+                  Just ev -> noteWake . (ev :) =<< pollEvents
+            , sdApplyEvent    = applyEvent
+            , sdIsButtonEdge  = isButtonEdge
+            , sdIsHardQuit    = isHardQuit
+            , sdIsSessionQuit = (== EvQuit)
+            , sdSyncDisplay   = \c inp -> do
+                paused <- readIORef vsyncPaused
+                when paused $ do
+                  void $ setRenderVSync (sdlRenderer env) 1
+                  writeIORef vsyncPaused False
+                (c', inp') <- syncDisplay c env inp
+                writeIORef ctxRef c'
+                writeIORef prev inp'
+                pure (c', inp')
+            , sdDebug         = sdsSampler (sdlDebug env)
+            , sdContinuous    = sdlContinuous env
+              -- With vsync on, presents throttle the loop. With vsync off a
+              -- live animation would spin at max speed, so wait ~2 ms short
+              -- of the frame period, leaving the slack for alignFrameStart:
+              -- SDL_WaitEventTimeout overruns by ~1 ms, and a wait that
+              -- returns past the boundary makes the frame late.
+            , sdPacingMs      = if sdlVsync env then 16 else max 1 (floor (sdlRefreshPeriod env * 1000) - 2)
+              -- A frame that skipped (empty damage, e.g. an animation scrolled
+              -- out of view) did not wait for vblank, so it must not loop
+              -- without waiting.
+            , sdPresentPaces  = if sdlVsync env then readIORef (sdlLastPresented env) else pure False
+            , sdShouldDraw    = \c prevInp inpSynced wasAnim wantDebug -> do
+                presented <- readIORef resizePresented
+                writeIORef resizePresented Nothing
+                let (prevInp', inpSynced') = case presented of
+                      Just size
+                        | size == inputWindowSize inpSynced ->
+                            (prevInp {inputWindowSize = size}, inpSynced {inputWindowRedraw = False})
+                      _ -> (prevInp, inpSynced)
+                shouldRedrawFrame c prevInp' inpSynced' wasAnim (sdlContinuous env) wantDebug
+            , sdDraw          = \c inpSynced forceFull -> do
+                writeIORef resizePresented Nothing
+                ms <- tryWithDrawingLock drawing (drawFn c env inpSynced (forceFull || sdlContinuous env))
+                case ms of
+                  Just (dirtyOut, s) -> do
+                    writeIORef prev s
+                    pure (dirtyOut, s)
+                  Nothing -> pure (False, inpSynced)
+            , sdOnCursor      = syncPointerCursor (sdlCursors env)
+            , sdAlignSec      = sdlRefreshPeriod env
+            , sdShouldQuit    = shouldQuit
+            }
+    bracket (installResizeWatch onResize) id $ \_ ->
+      runSessionLoop drv ctx2 synced1
diff --git a/lib/NanoUI/Sdl/Window.hs b/lib/NanoUI/Sdl/Window.hs
new file mode 100644
--- /dev/null
+++ b/lib/NanoUI/Sdl/Window.hs
@@ -0,0 +1,434 @@
+module NanoUI.Sdl.Window
+  ( RgbaImage (..)
+  , SdlEnv (..)
+  , SdlOptions (..)
+  , defaultSdlOptions
+  , withSdl
+  , withSdlBench
+  , syncDisplay
+  , saveScreenshot
+  ) where
+
+import Control.Exception (bracket)
+import Control.Monad (unless, void, when)
+import Data.Bits ((.|.))
+import Data.ByteString (ByteString)
+import Data.ByteString qualified as BS
+import Data.Maybe (isJust, isNothing)
+import Data.IORef (IORef, newIORef, readIORef, writeIORef)
+import System.Environment (lookupEnv)
+import Text.Read (readMaybe)
+import Data.Primitive.SmallArray (SmallArray)
+import Data.Text (Text)
+import Data.Text.Foreign qualified as TextForeign
+import Foreign.C.String (withCString)
+import Foreign.Marshal.Alloc (alloca)
+import Foreign.Ptr (Ptr, nullPtr)
+import Foreign.Storable (peek)
+import NanoUI (ImageId, Input (..), Size (..), Theme)
+import NanoUI.Context (Context (..), setDrawSnapScale)
+import NanoUI.Testing (clearMeasureCache, markDirty, setHost, setWakeLoop)
+import NanoUI.Sdl.Display
+  ( defaultFontSize
+  , initRefreshEvent
+  , pushRefreshEvent
+  , queryMouseWindowPos
+  , queryWindowDisplayScale
+  , queryWindowLogicalSize
+  , queryWindowRefreshHz
+  )
+import NanoUI.Sdl.Clipboard (withSdlClipboard)
+import NanoUI.Sdl.Cursor (SdlCursors (..), destroyCursors, initCursors)
+import NanoUI.Sdl.Font
+  ( FontSource (..)
+  , GlyphAtlas
+  , SdlFontCache
+  , destroyGlyphAtlas
+  , destroySdlFontCache
+  , newGlyphAtlas
+  , newSdlFontCache
+  , reloadSdlFontCache
+  , sdlFontCacheSource
+  , withSdlFontCache
+  , withTtf
+  )
+import NanoUI.Sdl.Font.Resolve
+  ( embeddedFontSource
+  , resolveNanoUIFont
+  , defaultFontSearch
+  , defaultFontSearchMono
+  )
+import NanoUI.Sdl.NanoUIFont (NanoUIFont (..))
+import NanoUI.Sdl.Debug (SdlDebugSampler, newSdlDebugSampler)
+import NanoUI.Sdl.Dialog.Types (DialogState (..), clearDialogState, newDialogState)
+import NanoUI.Sdl.Image (ImageAtlas, destroyImageAtlas, newImageAtlas)
+import NanoUI.Sdl.Render (RenderBatch, destroyRenderBatch, newRenderBatch)
+import SDL3.Sys.Bindgen.Hints (sDL_HINT_ASSERT, sDL_HINT_RENDER_VSYNC, sDL_HINT_VIDEO_DRIVER)
+import SDL3.Sys.Bindgen.Render (SDL_Renderer, SDL_Texture)
+import SDL3.Sys.Bindgen.Runtime.PtrConst qualified as PtrConst
+import SDL3.Sys.Bindgen.Video (SDL_Window, SDL_WindowFlags (..))
+import SDL3.Sys.Bindgen.Init (SDL_InitFlags (..), sDL_INIT_VIDEO)
+import SDL3.Sys.Hints (setHint)
+import SDL3.Sys.Init (initSafe, quitSafe)
+import SDL3.Sys.Keyboard (startTextInputSafe, stopTextInputSafe)
+import SDL3.Sys.Render
+  ( createWindowAndRendererSafe
+  , destroyRendererSafe
+  , destroyTexture
+  , getRendererName
+  , renderReadPixels
+  , setRenderScale
+  , setRenderVSync
+  )
+import SDL3.Sys.Surface (destroySurface, saveBMP)
+import SDL3.Sys.Video (destroyWindowSafe)
+
+-- | Initial RGBA asset uploaded before the first frame.
+data RgbaImage = RgbaImage
+  { rgbaImageId :: !ImageId
+  , rgbaImageWidth :: !Int
+  , rgbaImageHeight :: !Int
+  , rgbaImagePixels :: !ByteString
+  }
+
+-- | Application-owned SDL settings.
+data SdlOptions = SdlOptions
+  { sdlWindowTitle :: !Text
+  -- ^ Window title (default: @"nano-ui"@).
+  , sdlWindowSize :: !Size
+  -- ^ Initial window size in logical units (default: 1280x800).
+  , sdlWindowResizable :: !Bool
+  -- ^ Allow the window to be resized (default: 'True').
+  , sdlWindowFullscreen :: !Bool
+  -- ^ Open the window in fullscreen mode (default: 'False').
+  , sdlWindowBorderless :: !Bool
+  -- ^ Create a borderless window (default: 'False').
+  , sdlWindowAlwaysOnTop :: !Bool
+  -- ^ Keep the window on top of other windows (default: 'False').
+  , sdlWindowHidden :: !Bool
+  -- ^ Start the window hidden (default: 'False').
+  , sdlAppVsync :: !Bool
+  -- ^ Enable vertical synchronization (default: 'True').
+  , sdlAppContinuous :: !Bool
+  -- ^ Continuous unthrottled rendering without waiting for events (default: 'False').
+  , sdlAppFont :: !NanoUIFont
+  -- ^ UI font (default: embedded Inter).
+  , sdlAppMonoFont :: !NanoUIFont
+  -- ^ Monospace font (default: embedded Inter).
+  , sdlAppFontSize :: !Float
+  -- ^ Base font size in points (default: 16).
+  , sdlAppTheme :: !(Maybe Theme)
+  -- ^ Initial UI theme override (default: 'Nothing').
+  , sdlAppShouldQuit :: !(Input -> Bool)
+  -- ^ Predicate on user input to trigger application exit (default: @const False@).\
+  , sdlAppImages :: !(SmallArray RgbaImage)
+  -- ^ Initial RGBA textures registered before the first frame.
+  }
+
+defaultSdlOptions :: SdlOptions
+defaultSdlOptions =
+  SdlOptions
+    { sdlWindowTitle = "nano-ui"
+    , sdlWindowSize = defaultWindowSize
+    , sdlWindowResizable = True
+    , sdlWindowFullscreen = False
+    , sdlWindowBorderless = False
+    , sdlWindowAlwaysOnTop = False
+    , sdlWindowHidden = False
+    , sdlAppVsync = True
+    , sdlAppContinuous = False
+    , sdlAppFont = defaultFontSearch
+    , sdlAppMonoFont = defaultFontSearchMono
+    , sdlAppFontSize = defaultFontSize
+    , sdlAppTheme = Nothing
+    , sdlAppShouldQuit = const False
+    , sdlAppImages = mempty
+    }
+
+-- | SDL_WINDOW_HIGH_PIXEL_DENSITY (0x2000): without it the window's surface
+-- gets scale 1.0 even on a 2x / HiDPI output, so the compositor upscales the
+-- whole window (blurry "looks upscaled"). With it, SDL_GetWindowDisplayScale
+-- returns the real output scale, the window keeps its logical size, and the
+-- pixel buffer (and therefore the retain texture, glyph atlas, and fonts)
+-- rasterizes at the native pixel density.
+windowFlags :: SdlOptions -> SDL_WindowFlags
+windowFlags opts =
+  SDL_WindowFlags $
+    0x0000000000002000
+      .|. flag sdlWindowResizable 0x0000000000000020
+      .|. flag sdlWindowFullscreen 0x0000000000000001
+      .|. flag sdlWindowBorderless 0x0000000000000010
+      .|. flag sdlWindowAlwaysOnTop 0x0000000000010000
+      .|. flag sdlWindowHidden 0x0000000000000008
+  where
+    flag field bit = if field opts then bit else 0
+
+-- Hidden only. Do not combine with resizable for bench windows on Windows.
+sdlWindowHiddenFlag :: SDL_WindowFlags
+sdlWindowHiddenFlag = SDL_WindowFlags 0x0000000000000008
+
+scaleEpsilon :: Float
+scaleEpsilon = 0.001
+
+data SdlEnv = SdlEnv
+  { sdlWindow :: Ptr SDL_Window
+  , sdlRenderer :: Ptr SDL_Renderer
+  , sdlRendererName :: !Text
+  , sdlBatch :: RenderBatch
+  , sdlFontRequestRef :: !(IORef NanoUIFont)
+  , sdlFontAppliedRef :: !(IORef NanoUIFont)
+  , sdlForcedScale :: !(Maybe Float)
+  -- ^ NANO_FORCE_SCALE override of the display scale, read at startup.
+  , sdlScaleRef :: IORef Float
+  , sdlGlyphAtlas :: GlyphAtlas
+  , sdlImages :: ImageAtlas
+  , sdlCursors :: SdlCursors
+  , sdlDebug :: SdlDebugSampler
+  , sdlRetain :: IORef (Ptr SDL_Texture, Int, Int, Float)
+  , sdlLastPresented :: IORef Bool
+  , sdlVsync :: !Bool
+  , sdlRefreshPeriod :: !Double
+  , sdlContinuous :: !Bool
+  , sdlCachedCtx :: !(IORef Context)
+  , sdlFontCache :: !SdlFontCache
+  , sdlDialogState :: !DialogState
+  }
+
+defaultWindowSize :: Size
+defaultWindowSize = Size 1280 800
+
+-- Layout in logical coordinates; draw/text rasterize at native pixel density.
+syncDisplay :: Context -> SdlEnv -> Input -> IO (Context, Input)
+syncDisplay ctx env inp = do
+  scale <- maybe (queryWindowDisplayScale (sdlWindow env)) pure (sdlForcedScale env)
+  oldScale <- readIORef (sdlScaleRef env)
+  let scaleChanged = abs (scale - oldScale) > scaleEpsilon
+  when scaleChanged $ do
+    -- Presents leave the renderer at 1:1 pixels; re-assert it only when the
+    -- display scale moves.
+    ok <- setRenderScale (sdlRenderer env) 1 1
+    unless ok $ fail "SDL_SetRenderScale failed"
+    writeIORef (sdlScaleRef env) scale
+    setDrawSnapScale ctx scale
+  -- Runtime font-family switch: the app publishes its requested family through
+  -- 'setSdlUiFont'; resolve and apply it here, on the display thread, before
+  -- the next frame so the atlas, metrics, and text resolver agree.
+  requested <- readIORef (sdlFontRequestRef env)
+  applied <- readIORef (sdlFontAppliedRef env)
+  let fontChanged = requested /= applied
+  when (scaleChanged || fontChanged) $ do
+    source <-
+      if fontChanged
+        then resolveNanoUIFont requested
+        else sdlFontCacheSource (sdlFontCache env)
+    writeIORef (sdlFontAppliedRef env) requested
+    reloadSdlFontCache (sdlFontCache env) source
+    writeIORef (sdlCachedCtx env) . withSdlClipboard =<< withSdlFontCache (sdlFontCache env) ctx
+    clearMeasureCache ctx
+    markDirty ctx
+  queried <- queryWindowLogicalSize (sdlWindow env)
+  let winSize =
+        case queried of
+          Size 0 0 ->
+            case inputWindowSize inp of
+              Size 0 0 -> defaultWindowSize
+              s -> s
+          s -> s
+  mouse <- queryMouseWindowPos
+  ctxMeasured <- readIORef (sdlCachedCtx env)
+  pure (ctxMeasured, inp {inputWindowSize = winSize, inputMousePos = mouse})
+
+-- | Everything a window session is opened with, besides the context.
+data WindowConfig = WindowConfig
+  { wcTitle :: !Text
+  , wcSize :: !Size
+  , wcFlags :: !SDL_WindowFlags
+  , wcBench :: !Bool
+  -- ^ Hidden benchmark window: bench hints, no vsync setup or text input.
+  , wcVsync :: !Bool
+  , wcContinuous :: !Bool
+  , wcUiFont :: !NanoUIFont
+  , wcMonoFont :: !NanoUIFont
+  , wcFontSize :: !Float
+  }
+
+withSdl :: SdlOptions -> Context -> (Context -> SdlEnv -> IO a) -> IO a
+withSdl opts ctx =
+  withSdlWindow
+    ctx
+    WindowConfig
+      { wcTitle = sdlWindowTitle opts
+      , wcSize = sdlWindowSize opts
+      , wcFlags = windowFlags opts
+      , wcBench = False
+      , wcVsync = sdlAppVsync opts
+      , wcContinuous = sdlAppContinuous opts
+      , wcUiFont = sdlAppFont opts
+      , wcMonoFont = sdlAppMonoFont opts
+      , wcFontSize = sdlAppFontSize opts
+      }
+
+withSdlBench :: Context -> (Context -> SdlEnv -> IO a) -> IO a
+withSdlBench ctx =
+  withSdlWindow
+    ctx
+    WindowConfig
+      { wcTitle = "nano-ui-bench"
+      , wcSize = Size 800 600
+      , wcFlags = sdlWindowHiddenFlag
+      , wcBench = True
+      , wcVsync = False
+      , wcContinuous = True
+      , wcUiFont = DefaultFont
+      , wcMonoFont = DefaultFont
+      , wcFontSize = defaultFontSize
+      }
+
+withSdlWindow :: Context -> WindowConfig -> (Context -> SdlEnv -> IO a) -> IO a
+withSdlWindow ctx cfg act =
+  withTtf $ do
+    let hint name value =
+          BS.useAsCString name $ \cname ->
+            BS.useAsCString value $ \cvalue ->
+              void $ setHint (PtrConst.unsafeFromPtr cname) (PtrConst.unsafeFromPtr cvalue)
+    if wcBench cfg
+      then do
+        hint sDL_HINT_ASSERT "always_ignore"
+        hint sDL_HINT_RENDER_VSYNC "0"
+      else do
+        hint sDL_HINT_RENDER_VSYNC (if wcVsync cfg then "1" else "0")
+        -- SDL3 only auto-picks Wayland when the compositor has the fifo-v1 /
+        -- commit-timing-v1 protocols. Without them (sway, wlroots, many
+        -- others) it selects X11/XWayland, giving a scale-1 window on a
+        -- scale-2 (or fractional) output that the compositor upscales, so
+        -- text looks blurred. Native Wayland with
+        -- SDL_WINDOW_HIGH_PIXEL_DENSITY rasterizes at the real output scale.
+        -- An explicit SDL_VIDEO_DRIVER wins, and pure X11 sessions are left
+        -- alone.
+        wayland <- lookupEnv "WAYLAND_DISPLAY"
+        driver <- lookupEnv "SDL_VIDEO_DRIVER"
+        when (isJust wayland && isNothing driver) $
+          hint sDL_HINT_VIDEO_DRIVER "wayland"
+    fontSource <- resolveNanoUIFont (wcUiFont cfg)
+    monoSource <- resolveNanoUIFont (wcMonoFont cfg)
+    bracket
+      (startSdlWindow ctx cfg fontSource monoSource)
+      (\(_, env) -> stopSdlWindow (wcBench cfg) env)
+      (uncurry act)
+
+startSdlWindow :: Context -> WindowConfig -> FontSource -> FontSource -> IO (Context, SdlEnv)
+startSdlWindow ctx cfg fontSource monoSource = do
+  videoOk <- initSafe (SDL_InitFlags (fromIntegral sDL_INIT_VIDEO))
+  unless videoOk $ fail "SDL_Init(SDL_INIT_VIDEO) failed"
+  refreshOk <- initRefreshEvent
+  unless refreshOk $ fail "SDL_RegisterEvents failed for refresh wake"
+  let Size w h = wcSize cfg
+      bench = wcBench cfg
+  -- NANO_FORCE_SCALE: debug override of the display scale.
+  forcedEnv <- lookupEnv "NANO_FORCE_SCALE"
+  let forcedScale = case forcedEnv >>= readMaybe of
+        Just s | s > 0 -> Just s
+        _ -> Nothing
+  env <-
+    TextForeign.withCString (wcTitle cfg) $ \titlePtr ->
+      alloca $ \winPtr ->
+        alloca $ \renPtr -> do
+          ok <-
+            createWindowAndRendererSafe
+              (PtrConst.unsafeFromPtr titlePtr)
+              (round w)
+              (round h)
+              (wcFlags cfg)
+              winPtr
+              renPtr
+          unless ok $ fail "SDL_CreateWindowAndRenderer failed"
+          win <- peek winPtr
+          ren <- peek renPtr
+          scale <- queryWindowDisplayScale win
+          setDrawSnapScale ctx scale
+          refreshHz <- queryWindowRefreshHz win
+          rendererName <- getRendererName ren >>= \name ->
+            if PtrConst.unsafeToPtr name == nullPtr then pure "unknown" else TextForeign.peekCString (PtrConst.unsafeToPtr name)
+          scaleRef <- newIORef scale
+          fontRequestRef <- newIORef (wcUiFont cfg)
+          fontAppliedRef <- newIORef (wcUiFont cfg)
+          glyphAtlas <- newGlyphAtlas ren
+          images <- newImageAtlas
+          cursors <- initCursors
+          debug <- newSdlDebugSampler
+          retain <- newIORef (nullPtr, 0, 0, 0)
+          fontCache <-
+            newSdlFontCache
+              fontSource
+              embeddedFontSource
+              monoSource
+              embeddedFontSource
+              glyphAtlas
+              (wcFontSize cfg)
+              scaleRef
+          cachedCtx <- newIORef . withSdlClipboard =<< withSdlFontCache fontCache ctx
+          let refreshPeriod =
+                if refreshHz > 0
+                  then 1 / fromIntegral refreshHz
+                  else 1 / 60
+          scaleOk <- setRenderScale ren 1 1
+          unless scaleOk $ fail "SDL_SetRenderScale failed"
+          unless bench $ do
+            void $ setRenderVSync ren (if wcVsync cfg then 1 else 0)
+            void $ startTextInputSafe win
+          dialogState <- newDialogState
+          lastPresented <- newIORef False
+          batch <- newRenderBatch ren
+          pure
+            SdlEnv
+              { sdlWindow = win
+              , sdlRenderer = ren
+              , sdlRendererName = rendererName
+              , sdlBatch = batch
+              , sdlFontRequestRef = fontRequestRef
+              , sdlFontAppliedRef = fontAppliedRef
+              , sdlForcedScale = forcedScale
+              , sdlScaleRef = scaleRef
+              , sdlGlyphAtlas = glyphAtlas
+              , sdlImages = images
+              , sdlCursors = cursors
+              , sdlDebug = debug
+              , sdlRetain = retain
+              , sdlLastPresented = lastPresented
+              , sdlVsync = wcVsync cfg
+              , sdlRefreshPeriod = refreshPeriod
+              , sdlContinuous = wcContinuous cfg
+              , sdlCachedCtx = cachedCtx
+              , sdlFontCache = fontCache
+              , sdlDialogState = dialogState
+              }
+  ctx' <- readIORef (sdlCachedCtx env)
+  setHost ctx' env
+  setWakeLoop ctx' pushRefreshEvent
+  pure (ctx', env)
+
+stopSdlWindow :: Bool -> SdlEnv -> IO ()
+stopSdlWindow bench env = do
+  clearDialogState (sdlDialogState env)
+  (tex, _, _, _) <- readIORef (sdlRetain env)
+  unless (tex == nullPtr) $ destroyTexture tex
+  destroyRenderBatch (sdlBatch env)
+  destroyCursors (sdlCursors env)
+  destroyImageAtlas (sdlImages env)
+  destroySdlFontCache (sdlFontCache env)
+  destroyGlyphAtlas (sdlGlyphAtlas env)
+  unless bench $ void $ stopTextInputSafe (sdlWindow env)
+  void $ setRenderScale (sdlRenderer env) 1 1
+  destroyRendererSafe (sdlRenderer env)
+  destroyWindowSafe (sdlWindow env)
+  quitSafe
+
+saveScreenshot :: SdlEnv -> FilePath -> IO Bool
+saveScreenshot env path = do
+  surface <- renderReadPixels (sdlRenderer env) (PtrConst.unsafeFromPtr nullPtr)
+  if surface == nullPtr
+    then pure False
+    else withCString path $ \cpath -> do
+      ok <- saveBMP surface (PtrConst.unsafeFromPtr cpath)
+      destroySurface surface
+      pure ok
diff --git a/nano-ui-sdl.cabal b/nano-ui-sdl.cabal
new file mode 100644
--- /dev/null
+++ b/nano-ui-sdl.cabal
@@ -0,0 +1,185 @@
+cabal-version:      3.4
+name:               nano-ui-sdl
+version:            0.1.0.0
+synopsis:           SDL3 window backend for nano-ui
+description:
+    Runs nano-ui views in an SDL3 window with TrueType text and native file
+    dialogs. Requires SDL3, SDL3_ttf, and pkg-config.
+license:            MIT AND OFL-1.1
+license-file:       LICENSE
+author:             goolord
+maintainer:         zacharyachurchill@gmail.com
+category:           Graphics
+homepage:           https://github.com/goolord/nano-ui
+bug-reports:        https://github.com/goolord/nano-ui/issues
+build-type:         Simple
+tested-with:        GHC ==9.10.3 || ==9.14.1
+extra-doc-files:
+    CHANGELOG.md
+    README.md
+    data/inter-LICENSE.txt
+
+extra-source-files:
+    data/inter.ttf
+    cbits/*.h
+
+source-repository head
+    type:     git
+    location: https://github.com/goolord/nano-ui.git
+    subdir:   packages/nano-ui-sdl
+
+flag sdl
+    description: Build the SDL3 backend. Needs SDL3, SDL3_ttf, and pkg-config.
+    manual: True
+    default: True
+
+flag simd
+    description: Compile the batch culler with AVX2 (x86-64 only; the binary then requires an AVX2 CPU)
+    manual: True
+    default: False
+
+common extensions
+    default-language: GHC2024
+    default-extensions:
+        OverloadedStrings
+
+common rts-options
+    ghc-options:
+        -rtsopts
+        -threaded
+        "-with-rtsopts=-N1 -A64m -T -I0"
+
+-- -N1 keeps the threaded runtime on one capability. Do not pass tasty's -j1
+-- via -with-rtsopts; use: cabal bench ... --benchmark-options=-j1
+common bench-rts-options
+    ghc-options: -rtsopts -threaded "-with-rtsopts=-N1 -T"
+
+common warnings
+  ghc-options:
+    -Wall
+    -Wextra
+    -Wcompat
+    -Widentities
+    -Wincomplete-record-updates
+    -Wincomplete-uni-patterns
+    -Wmissing-export-lists
+    -Wmissing-home-modules
+    -Wpartial-fields
+    -Wredundant-constraints
+    -Wunused-packages
+
+library
+    import:           extensions
+    import:           warnings
+    exposed-modules:
+        NanoUI.Backend.Sdl
+        NanoUI.Sdl.Input
+        NanoUI.Sdl.NanoUIFont
+    other-modules:
+        NanoUI.Sdl.Session
+        NanoUI.Sdl.Runner
+        NanoUI.Sdl.Clipboard
+        NanoUI.Sdl.Cursor
+        NanoUI.Sdl.Debug
+        NanoUI.Sdl.Dialog
+        NanoUI.Sdl.Dialog.Types
+        NanoUI.Sdl.Display
+        NanoUI.Sdl.Font
+        NanoUI.Sdl.Font.Inter
+        NanoUI.Sdl.Font.Resolve
+        NanoUI.Sdl.Font.Search
+        NanoUI.Sdl.Image
+        NanoUI.Sdl.Render
+        NanoUI.Sdl.Window
+    build-depends:
+        base >=4.20 && <4.23,
+        bytestring >=0.11 && <0.13,
+        containers >=0.6.7 && <0.9,
+        directory >=1.3.7 && <1.4,
+        dir-traverse >=0.2.3 && <0.3,
+        effectful-core >=2.5 && <2.8,
+        filepath >=1.4.100 && <1.6,
+        file-embed >=0.0.16 && <0.1,
+        hashable >=1.4 && <1.6,
+        nano-ui ^>=0.1,
+        primitive >=0.8 && <0.10,
+        record-hasfield >=1.0 && <1.1,
+        sdl3-bindgen-sys >=0.0.0.3 && <0.0.1,
+        text >=2.1.2 && <2.2,
+        unordered-containers >=0.2.19 && <0.3
+    hs-source-dirs:   lib
+    if !flag(sdl)
+        buildable: False
+    if flag(sdl)
+        c-sources:        cbits/nano_ui_ttf.c
+                          cbits/nano_ui_display.c
+                          cbits/nano_ui_batch.c
+                          cbits/nano_ui_text_atlas.c
+        include-dirs:     cbits
+        pkgconfig-depends: sdl3 >= 3.2, sdl3-ttf >= 3.2
+        -- nano_ui_batch.c falls back to a scalar cull loop without __AVX2__.
+        if flag(simd) && arch(x86_64)
+            cc-options:   -mavx2
+
+test-suite nano-ui-font-search-test
+    import:           extensions, warnings
+    type:             exitcode-stdio-1.0
+    main-is:          FontSearch.hs
+    hs-source-dirs:   test, lib
+    other-modules:    NanoUI.Sdl.Font.Search
+    build-depends:
+        base >=4.20 && <4.23,
+        containers >=0.6.7 && <0.9,
+        directory >=1.3.7 && <1.4,
+        dir-traverse >=0.2.3 && <0.3,
+        filepath >=1.4.100 && <1.6
+
+test-suite nano-ui-font-effects-test
+    import:           extensions, warnings, rts-options
+    type:             exitcode-stdio-1.0
+    main-is:          FontEffects.hs
+    hs-source-dirs:   test
+    build-depends:
+        base >=4.20 && <4.23,
+        nano-ui,
+        nano-ui-sdl,
+        primitive >=0.8 && <0.10,
+        text >=2.1.2 && <2.2
+    if !flag(sdl)
+        buildable: False
+
+executable nano-ui-sdl-anim
+    import:           extensions
+    import:           rts-options
+    import:           warnings
+    main-is:          SdlAnim.hs
+    build-depends:
+        base >=4.20 && <4.23,
+        nano-ui,
+        nano-ui-sdl,
+        text >=2.1.2 && <2.2
+    hs-source-dirs:   examples
+    if os(windows)
+        ghc-options: -optl-mconsole
+    if !flag(sdl)
+        buildable: False
+
+benchmark nano-ui-sdl-bench
+    import:           extensions
+    import:           bench-rts-options
+    import:           warnings
+    default-extensions: CPP
+    type:             exitcode-stdio-1.0
+    main-is:          SdlBench.hs
+    build-depends:
+        base >=4.20 && <4.23,
+        nano-ui,
+        nano-ui-sdl,
+        tasty-bench >=0.3 && <0.6
+    if os(windows)
+        build-depends: Win32 >=2.13 && <2.15
+    hs-source-dirs:   benchmark
+    if os(windows)
+        ghc-options: -optl-mconsole
+    if !flag(sdl)
+        buildable: False
diff --git a/test/FontEffects.hs b/test/FontEffects.hs
new file mode 100644
--- /dev/null
+++ b/test/FontEffects.hs
@@ -0,0 +1,107 @@
+module Main (main) where
+
+import Control.Exception (IOException, evaluate, try)
+import Control.Monad (forM_, unless, void)
+import Data.List (isInfixOf)
+import Data.IORef (writeIORef)
+import qualified Data.Text as T
+import Data.Primitive.PrimArray (indexPrimArray, sizeofPrimArray)
+import NanoUI
+import NanoUI.Testing (newPixelContext, textIndexAtX)
+import NanoUI.Backend.Sdl (NanoUIFont (..), SdlEnv (..), syncDisplay, withSdlBench)
+import NanoUI.Context (ctxResolveFont, ctxResolveMeasure)
+import System.Environment (setEnv)
+import System.Mem (performGC)
+
+main :: IO ()
+main = do
+  setEnv "SDL_VIDEODRIVER" "dummy"
+  setEnv "SDL_RENDER_DRIVER" "software"
+  ctx0 <- newPixelContext
+  (font, snapshot, width, quad) <- withSdlBench ctx0 $ \ctx env -> do
+    (fm, _) <- ctxResolveFont ctx 16 WeightNormal FontStyleNormal FontRegular
+    let text = "AV To fi café λ"
+    prepared <- prepareFontMetrics fm text
+    (hostWidth, _) <- ctxResolveMeasure ctx 16 WeightNormal FontStyleNormal FontRegular text
+    measured <- evaluate (lineWidth prepared text)
+    unless (abs (hostWidth - measured) < 0.01) $
+      fail ("prepared shaped width differs from SDL measurement: " ++ show (hostWidth, measured))
+    scaled <- lineWidthIO (scaleFontMetrics 1.5 fm) text
+    unless (abs (scaled - measured * 1.5) < 0.01) $
+      fail "effectful font scaling lost its metric scale"
+    before <- drawShaped fm text
+    unless (maybe False (\(ShapedGlyphs q) -> sizeofPrimArray q >= 8 * 10) before) $
+      fail "shaped drawing returned no glyph quads"
+    -- Exceed both cache caps through measurement only. This must not fill or
+    -- reset the atlas, nor alter an already-rasterised run.
+    forM_ [1 .. 1100 :: Int] $ \n -> do
+      let labelText = "counter " <> T.pack (show n)
+      p <- prepareFontMetrics fm labelText
+      void (evaluate (lineWidth p labelText))
+    after <- drawShaped fm text
+    unless (before == after) $
+      fail "metric preparation mutated the atlas or raster cache"
+    let oversized = T.replicate 1000 "W"
+    large <- prepareFontMetrics fm oversized
+    largeGlyphs <- drawShaped fm oversized
+    unless (maybe False (\(ShapedGlyphs q) -> sizeofPrimArray q == 8 * 1000) largeGlyphs && lineWidth large oversized > 0) $
+      fail "text wider than the atlas lost its glyphs"
+    shapingChecks fm
+    performGC
+    again <- lineWidthIO fm text
+    unless (abs (again - measured) < 0.01) $
+      fail "font metrics changed after cache eviction and GC"
+    -- Exercise the actual font/atlas replacement path. Old pure snapshots
+    -- remain valid while old native callbacks must reject their closed font.
+    writeIORef (sdlFontRequestRef env) (FontSearch [])
+    (replacement, _) <- syncDisplay ctx env emptyInput
+    expectClosed (drawShaped fm text)
+    (fresh, _) <- ctxResolveFont replacement 16 WeightNormal FontStyleNormal FontRegular
+    freshWidth <- lineWidthIO fresh text
+    freshQuad <- drawShaped fresh text
+    unless (abs (freshWidth - measured) < 0.01 && freshQuad /= Nothing) $
+      fail "font/atlas replacement failed to restore shaped text"
+    pure (fresh, prepared, measured, freshQuad)
+  -- Pure metric and quad values are safe to evaluate after the native font,
+  -- atlas and SDL session are closed; native effects fail before dereferencing.
+  afterClose <- evaluate (lineWidth snapshot "AV To fi café λ")
+  unless (afterClose == width) $ fail "immutable snapshot changed after shutdown"
+  void (evaluate quad)
+  expectClosed (prepareFontMetrics font "new text")
+  expectClosed (drawShaped font "AV To fi café λ")
+  expectClosed (drawGlyph font 'A')
+  putStrLn "font effects: ok"
+
+-- | Shaping reorders right-to-left text and gives every character a caret:
+-- an Arabic word's carets run right to left, a mixed line keeps its Latin
+-- carets increasing, and a click lands on the nearest caret.
+shapingChecks :: FontMetrics -> IO ()
+shapingChecks fm = do
+  let arabic = "مرحبا"
+      mixed = "Hi مرحبا 12"
+  pArabic <- prepareFontMetrics fm arabic
+  case fmShape pArabic arabic of
+    Nothing -> fail "Arabic text was not shaped"
+    Just st -> do
+      let carets = [indexPrimArray (stCarets st) i | i <- [0 .. sizeofPrimArray (stCarets st) - 1]]
+      unless (length carets == T.length arabic + 1) $ fail "Arabic carets do not cover every character"
+      -- The script sets the direction, so the word runs right to left
+      -- whichever font draws it.
+      unless (and (zipWith (>=) carets (drop 1 carets))) $
+        fail ("Arabic carets do not decrease: " <> show carets)
+      unless (textIndexAtX pArabic arabic (maximum carets + 1) == 0) $
+        fail "a click at the right edge of Arabic text did not land before its first character"
+  pMixed <- prepareFontMetrics fm mixed
+  case fmShape pMixed mixed of
+    Nothing -> fail "mixed text was not shaped"
+    Just st -> do
+      let caret i = indexPrimArray (stCarets st) i
+      unless (caret 0 < caret 1 && caret 1 < caret 2) $ fail "Latin carets in a mixed line do not increase"
+      unless (caret 9 < caret 10 && caret 10 < caret 11) $ fail "digits after right-to-left text do not run left to right"
+
+expectClosed :: IO a -> IO ()
+expectClosed action = do
+  result <- try (void action) :: IO (Either IOException ())
+  case result of
+    Left err | "used after" `isInfixOf` show err -> pure ()
+    _ -> fail "retained font callback did not reject its closed native handle"
diff --git a/test/FontSearch.hs b/test/FontSearch.hs
new file mode 100644
--- /dev/null
+++ b/test/FontSearch.hs
@@ -0,0 +1,69 @@
+module Main (main) where
+
+import Control.Exception (bracket)
+import Control.Monad (unless)
+import NanoUI.Sdl.Font.Search (listFontFamilies, searchFonts)
+import System.Directory
+  ( createDirectory
+  , createDirectoryIfMissing
+  , getTemporaryDirectory
+  , removeFile
+  , removePathForcibly
+  )
+import System.Environment (lookupEnv, setEnv, unsetEnv)
+import System.FilePath ((</>))
+import System.IO (hClose, openTempFile)
+import System.Info (os)
+
+-- Discovery only inspects filenames; no native display or valid font data is needed.
+main :: IO ()
+main = bracket temporaryRoot removePathForcibly $ \root ->
+  bracket (lookupEnv homeVar) restoreHome $ \_ -> do
+    setEnv homeVar root
+    let
+      fonts =
+        root </> case os of
+          "mingw32" -> "Microsoft/Windows/Fonts"
+          "darwin" -> "Library/Fonts"
+          _ -> ".local/share/fonts"
+      regular = fonts </> "NanoSearchFixture-Regular.ttf"
+      bold = fonts </> "NanoSearchFixture-Bold.ttf"
+      fallback = fonts </> "NanoFallbackFixture.otf"
+      boldOnly = fonts </> "NanoBoldOnlyFixture-Bold.ttf"
+    -- The font directories are walked once per process, so every fixture
+    -- exists before the first search.
+    createDirectoryIfMissing True fonts
+    mapM_ (`writeFile` "") [regular, bold, fallback, boldOnly]
+    expect "regular face" (Just regular) =<< searchFonts ["Nano Search Fixture"]
+    expect "ordered fallback" (Just fallback)
+      =<< searchFonts
+        ["", "NanoMissingFixture", "NanoFallbackFixture", "NanoSearchFixture"]
+    expect "empty request" Nothing =<< searchFonts []
+    expect "missing family" Nothing =<< searchFonts ["NanoMissingFixture"]
+    families <- listFontFamilies
+    expect
+      "deduplicated family"
+      ["Nano Search Fixture"]
+      (filter (== "Nano Search Fixture") families)
+    expect "non-regular fallback" (Just boldOnly) =<< searchFonts ["NanoBoldOnlyFixture"]
+    putStrLn "font search: ok"
+
+expect :: (Eq a, Show a) => String -> a -> a -> IO ()
+expect label expected actual =
+  unless (actual == expected) $
+    fail (label ++ ": expected " ++ show expected ++ ", got " ++ show actual)
+
+homeVar :: String
+homeVar = if os == "mingw32" then "LOCALAPPDATA" else "HOME"
+
+restoreHome :: Maybe String -> IO ()
+restoreHome = maybe (unsetEnv homeVar) (setEnv homeVar)
+
+temporaryRoot :: IO FilePath
+temporaryRoot = do
+  tmp <- getTemporaryDirectory
+  (path, handle) <- openTempFile tmp "nano-ui-font-search"
+  hClose handle
+  removeFile path
+  createDirectory path
+  pure path
