diff --git a/.github/workflows/modal-deploy.yml b/.github/workflows/modal-deploy.yml
index 35410a825..ad8bb19db 100644
--- a/.github/workflows/modal-deploy.yml
+++ b/.github/workflows/modal-deploy.yml
@@ -27,10 +27,13 @@ jobs:
       - name: Install Modal
         run: |
           python -m pip install --upgrade pip
-          pip install modal
+          pip install "modal>=1.5.0"
+
+      - name: Warm Torch header cache
+        run: |
+          PYTHONPATH=${{ github.workspace }}/src:${{ github.workspace }}/src/runners modal run src/runners/modal_runner.py::warm_pch
 
       - name: Deploy job
         run: |
           # TODO: Remove pythonpath here
-          PYTHONPATH=${{ github.workspace }}/src modal deploy src/runners/modal_runner_archs.py
-          
+          PYTHONPATH=${{ github.workspace }}/src:${{ github.workspace }}/src/runners modal deploy src/runners/modal_runner_archs.py
diff --git a/docs/benchmarks/modal-pch.json b/docs/benchmarks/modal-pch.json
new file mode 100644
index 000000000..0d4a08f92
--- /dev/null
+++ b/docs/benchmarks/modal-pch.json
@@ -0,0 +1,58 @@
+{
+  "date": "2026-09-07",
+  "gpu": "T4",
+  "cpu": 4,
+  "memory_mib": 16384,
+  "cuda": "13.3.0",
+  "torch": "2.12.0+cu130",
+  "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7",
+  "cutlass_gemm_version": "4.5.2",
+  "volume": "kernelbot-pch-test-20260907",
+  "ninja_parallelism": "default",
+  "baseline": "/usr/bin/g++; common timing shim in both modes; header tracing outside timed builds",
+  "sampling": "Two variants per case, each with/without PCH in separate containers. CUTLASS has two repetitions per variant/mode. Plain C++ uses the median of ten object compilations per container.",
+  "nvcc_timing": "Recorded for the CUTLASS GEMM runs. The earlier matrix instruments CXX only.",
+  "benchmark_sha256": "b22fbb07c7e78a5426544411f6e5dfbd394bfaa4364cf9d8dd0fe7f92f3a039c",
+  "compiler_sha256": "6d82acd795c3e8f69e63f1c1c62fb8afc52bc9c06a1a6a739a385907f26279d1",
+  "runs": {
+    "simple": "https://modal.com/apps/coreauto/main/ap-ZBKWqPp5sSkMJCFjKmWwPI",
+    "other_cases": "https://modal.com/apps/coreauto/main/ap-5aIt3temjLqpwFSfQhW6LI",
+    "cutlass_gemm": "https://modal.com/apps/coreauto/main/ap-6MXYSgeVomyahsGXOTKuG3",
+    "cutlass_gemm_implicit": "https://modal.com/apps/coreauto/main/ap-cMFfCst1YafIPb9GCly6AQ",
+    "all_warmup_profiles": "https://modal.com/apps/coreauto/main/ap-4ym5VrAiAuC127zircLT5b"
+  },
+  "measurements": [
+    {"case":"inline_simple","variant":0,"repeat":0,"cached":false,"build_s":18.800584544,"cpp_s":17.997382098,"nvcc_s":null,"pch_consumed":false,"pch_key":null,"checks":6,"max_error":null,"outputs":[2.0,2.0,2.0,2.0,2.0],"source_sha256":"4ec65ca0b08feea11ce48b7a0ebd322676aab622d35eba1da63b605643e46435","container_id":"ta-01M1YT654VFX478EV36GSNAGCR"},
+    {"case":"inline_simple","variant":0,"repeat":0,"cached":true,"build_s":9.449116172,"cpp_s":8.415279449,"nvcc_s":null,"pch_consumed":true,"pch_key":"cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38","checks":6,"max_error":null,"outputs":[2.0,2.0,2.0,2.0,2.0],"source_sha256":"4ec65ca0b08feea11ce48b7a0ebd322676aab622d35eba1da63b605643e46435","container_id":"ta-01M1YT65B8E9149HBHXBTA0NAR"},
+    {"case":"inline_simple","variant":1,"repeat":0,"cached":true,"build_s":10.546373582,"cpp_s":9.191087943000001,"nvcc_s":null,"pch_consumed":true,"pch_key":"cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38","checks":6,"max_error":null,"outputs":[3.0,3.0,3.0,3.0,3.0],"source_sha256":"7e064d69277762a33d2a65a4d6258e5422c7d2b3535694ce9afd6a6304f6993a","container_id":"ta-01M1YT654V66S3D05ZP1KSGBXR"},
+    {"case":"inline_simple","variant":1,"repeat":0,"cached":false,"build_s":18.893621677,"cpp_s":18.087305365,"nvcc_s":null,"pch_consumed":false,"pch_key":null,"checks":6,"max_error":null,"outputs":[3.0,3.0,3.0,3.0,3.0],"source_sha256":"7e064d69277762a33d2a65a4d6258e5422c7d2b3535694ce9afd6a6304f6993a","container_id":"ta-01M1YT65B8TXRZGSM820Z3Q16R"},
+    {"case":"load_files","variant":0,"repeat":0,"cached":false,"build_s":18.484804069,"cpp_s":17.716654833999996,"nvcc_s":null,"pch_consumed":false,"pch_key":null,"checks":6,"max_error":null,"outputs":[2.0,2.0,2.0,2.0,2.0],"source_sha256":"4ec65ca0b08feea11ce48b7a0ebd322676aab622d35eba1da63b605643e46435","container_id":"ta-01M1YT1MEHZQJTJ91K03EYFKJR"},
+    {"case":"load_files","variant":0,"repeat":0,"cached":true,"build_s":8.844153404,"cpp_s":7.813371175,"nvcc_s":null,"pch_consumed":true,"pch_key":"cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38","checks":6,"max_error":null,"outputs":[2.0,2.0,2.0,2.0,2.0],"source_sha256":"4ec65ca0b08feea11ce48b7a0ebd322676aab622d35eba1da63b605643e46435","container_id":"ta-01M1YT1MM8C9EBBGEZCWXZAW5R"},
+    {"case":"load_files","variant":1,"repeat":0,"cached":true,"build_s":8.704362362000001,"cpp_s":7.6940678259999995,"nvcc_s":null,"pch_consumed":true,"pch_key":"cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38","checks":6,"max_error":null,"outputs":[3.0,3.0,3.0,3.0,3.0],"source_sha256":"7e064d69277762a33d2a65a4d6258e5422c7d2b3535694ce9afd6a6304f6993a","container_id":"ta-01M1YT20M1E9G09T9PZBWNQ65R"},
+    {"case":"load_files","variant":1,"repeat":0,"cached":false,"build_s":18.118759269999998,"cpp_s":17.33960136,"nvcc_s":null,"pch_consumed":false,"pch_key":null,"checks":6,"max_error":null,"outputs":[3.0,3.0,3.0,3.0,3.0],"source_sha256":"7e064d69277762a33d2a65a4d6258e5422c7d2b3535694ce9afd6a6304f6993a","container_id":"ta-01M1YT276G72J91NSPNG9VWAWR"},
+    {"case":"plain_cpp","variant":0,"repeat":0,"cached":false,"build_s":0.19980638800000028,"cpp_s":0.12402935050000008,"nvcc_s":null,"pch_consumed":false,"pch_key":null,"checks":1,"max_error":null,"outputs":[3],"source_sha256":"f5fbb559817c22a648c4be4067de69dd806e6e9f2c9be71aa84f2371ecf8d781","container_id":"ta-01M1YT2FP2X908DV2VFHCJPT6R"},
+    {"case":"plain_cpp","variant":0,"repeat":0,"cached":true,"build_s":0.2823680209999999,"cpp_s":0.20570012650000002,"nvcc_s":null,"pch_consumed":false,"pch_key":null,"checks":1,"max_error":null,"outputs":[3],"source_sha256":"f5fbb559817c22a648c4be4067de69dd806e6e9f2c9be71aa84f2371ecf8d781","container_id":"ta-01M1YT2QSWZR9SN6VGQJC16VSR"},
+    {"case":"plain_cpp","variant":1,"repeat":0,"cached":true,"build_s":0.29172716450000014,"cpp_s":0.21366047899999963,"nvcc_s":null,"pch_consumed":false,"pch_key":null,"checks":1,"max_error":null,"outputs":[4],"source_sha256":"4f5b6c5b34f3f71e1e94ad89fb640e348e9bd0232d5cdfc75075c5e827674378","container_id":"ta-01M1YT2TY6A6XZNT3DJDM9AW6R"},
+    {"case":"plain_cpp","variant":1,"repeat":0,"cached":false,"build_s":0.20363044750000014,"cpp_s":0.12606111700000033,"nvcc_s":null,"pch_consumed":false,"pch_key":null,"checks":1,"max_error":null,"outputs":[4],"source_sha256":"4f5b6c5b34f3f71e1e94ad89fb640e348e9bd0232d5cdfc75075c5e827674378","container_id":"ta-01M1YT3294D1MHQ774CC40NP6R"},
+    {"case":"flag_miss","variant":0,"repeat":0,"cached":false,"build_s":18.314668328000003,"cpp_s":17.468810722999997,"nvcc_s":null,"pch_consumed":false,"pch_key":null,"checks":6,"max_error":null,"outputs":[2.0,2.0,2.0,2.0,2.0],"source_sha256":"4ec65ca0b08feea11ce48b7a0ebd322676aab622d35eba1da63b605643e46435","container_id":"ta-01M1YT32ZPHKD6NAA1RQHDFZNR"},
+    {"case":"flag_miss","variant":0,"repeat":0,"cached":true,"build_s":20.427356331,"cpp_s":19.146695692999998,"nvcc_s":null,"pch_consumed":false,"pch_key":null,"checks":6,"max_error":null,"outputs":[2.0,2.0,2.0,2.0,2.0],"source_sha256":"4ec65ca0b08feea11ce48b7a0ebd322676aab622d35eba1da63b605643e46435","container_id":"ta-01M1YT36X83KGN6Y6BSXPHJB5R"},
+    {"case":"flag_miss","variant":1,"repeat":0,"cached":true,"build_s":20.175504439,"cpp_s":19.005651779999997,"nvcc_s":null,"pch_consumed":false,"pch_key":null,"checks":6,"max_error":null,"outputs":[3.0,3.0,3.0,3.0,3.0],"source_sha256":"7e064d69277762a33d2a65a4d6258e5422c7d2b3535694ce9afd6a6304f6993a","container_id":"ta-01M1YT3AQERWZZTC7SBEKE1EFR"},
+    {"case":"flag_miss","variant":1,"repeat":0,"cached":false,"build_s":17.837603579,"cpp_s":17.036646379,"nvcc_s":null,"pch_consumed":false,"pch_key":null,"checks":6,"max_error":null,"outputs":[3.0,3.0,3.0,3.0,3.0],"source_sha256":"7e064d69277762a33d2a65a4d6258e5422c7d2b3535694ce9afd6a6304f6993a","container_id":"ta-01M1YT3MBVRB90BKSZW7DR6FHR"},
+    {"case":"cutlass_gemm","variant":0,"repeat":0,"cached":false,"build_s":19.794986437,"cpp_s":18.943877617,"nvcc_s":[12.304675192000001],"pch_consumed":false,"pch_key":null,"checks":10,"max_error":4.2646570364013314e-06,"outputs":[8.0],"source_sha256":"c0cc1e80088a516e7d3b0a21890de234e88d4cf0b2b278c6e8e74d7912d4a14d","container_id":"ta-01M1YTRK5PMETR512ZRYZAWEJR"},
+    {"case":"cutlass_gemm","variant":0,"repeat":0,"cached":true,"build_s":12.692048463999999,"cpp_s":8.008807361,"nvcc_s":[11.653260732],"pch_consumed":true,"pch_key":"cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38","checks":10,"max_error":4.2646570364013314e-06,"outputs":[8.0],"source_sha256":"c0cc1e80088a516e7d3b0a21890de234e88d4cf0b2b278c6e8e74d7912d4a14d","container_id":"ta-01M1YTRK5PF1P3622Q8VNYA6KR"},
+    {"case":"cutlass_gemm","variant":1,"repeat":0,"cached":true,"build_s":13.076952955000001,"cpp_s":7.930763236000001,"nvcc_s":[11.847514466],"pch_consumed":true,"pch_key":"cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38","checks":10,"max_error":8.529314072802663e-06,"outputs":[16.0],"source_sha256":"b926f64a92cfd59840083d417ec09e2db0dd93574c97203ad4ac0d08925ccaee","container_id":"ta-01M1YTRM3CV3B2HBFEF32S6QVR"},
+    {"case":"cutlass_gemm","variant":1,"repeat":0,"cached":false,"build_s":19.559281418,"cpp_s":18.747936609999996,"nvcc_s":[12.55567741],"pch_consumed":false,"pch_key":null,"checks":10,"max_error":8.529314072802663e-06,"outputs":[16.0],"source_sha256":"b926f64a92cfd59840083d417ec09e2db0dd93574c97203ad4ac0d08925ccaee","container_id":"ta-01M1YTRKCWZPE610GWKMMNAHFR"},
+    {"case":"cutlass_gemm","variant":0,"repeat":1,"cached":false,"build_s":21.135683762,"cpp_s":20.226324679999998,"nvcc_s":[12.684764970999998],"pch_consumed":false,"pch_key":null,"checks":10,"max_error":4.2646570364013314e-06,"outputs":[8.0],"source_sha256":"c0cc1e80088a516e7d3b0a21890de234e88d4cf0b2b278c6e8e74d7912d4a14d","container_id":"ta-01M1YTSM3BNSJWKC1QNYJ4A7NR"},
+    {"case":"cutlass_gemm","variant":0,"repeat":1,"cached":true,"build_s":12.152523803,"cpp_s":7.759680262,"nvcc_s":[11.084306554000001],"pch_consumed":true,"pch_key":"cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38","checks":10,"max_error":4.2646570364013314e-06,"outputs":[8.0],"source_sha256":"c0cc1e80088a516e7d3b0a21890de234e88d4cf0b2b278c6e8e74d7912d4a14d","container_id":"ta-01M1YTSN1GCNZ2DRWEXVY91GQR"},
+    {"case":"cutlass_gemm","variant":1,"repeat":1,"cached":true,"build_s":12.466685846,"cpp_s":7.833615612999999,"nvcc_s":[11.412195861],"pch_consumed":true,"pch_key":"cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38","checks":10,"max_error":8.529314072802663e-06,"outputs":[16.0],"source_sha256":"b926f64a92cfd59840083d417ec09e2db0dd93574c97203ad4ac0d08925ccaee","container_id":"ta-01M1YTT5X2Z3VGZN0ZXA8J8FFR"},
+    {"case":"cutlass_gemm","variant":1,"repeat":1,"cached":false,"build_s":20.099891596000003,"cpp_s":19.24853497,"nvcc_s":[12.522667313000001],"pch_consumed":false,"pch_key":null,"checks":10,"max_error":8.529314072802663e-06,"outputs":[16.0],"source_sha256":"b926f64a92cfd59840083d417ec09e2db0dd93574c97203ad4ac0d08925ccaee","container_id":"ta-01M1YTT4XD62ZGV4CJYD6AM7WR"},
+    {"case":"cutlass_gemm_implicit","variant":0,"repeat":0,"cached":false,"build_s":49.344148073,"cpp_s":20.027217276,"nvcc_s":[48.269329267],"pch_consumed":false,"pch_key":null,"checks":10,"max_error":4.2646570364013314e-06,"outputs":[8.0],"source_sha256":"c0cc1e80088a516e7d3b0a21890de234e88d4cf0b2b278c6e8e74d7912d4a14d","container_id":"ta-01M1YVFWZ4BD7M4B5P4EKGNRSR"},
+    {"case":"cutlass_gemm_implicit","variant":0,"repeat":0,"cached":true,"build_s":44.628304248000006,"cpp_s":8.165473682,"nvcc_s":[43.488331118],"pch_consumed":true,"pch_key":"cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38","checks":10,"max_error":4.2646570364013314e-06,"outputs":[8.0],"source_sha256":"c0cc1e80088a516e7d3b0a21890de234e88d4cf0b2b278c6e8e74d7912d4a14d","container_id":"ta-01M1YVFWZ4A8FTFDZ76TX6XT6R"},
+    {"case":"cutlass_gemm_implicit","variant":1,"repeat":0,"cached":true,"build_s":45.660435061,"cpp_s":8.445875664,"nvcc_s":[44.495367471],"pch_consumed":true,"pch_key":"cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38","checks":10,"max_error":8.529314072802663e-06,"outputs":[16.0],"source_sha256":"b926f64a92cfd59840083d417ec09e2db0dd93574c97203ad4ac0d08925ccaee","container_id":"ta-01M1YVFWSFP54HSM30AAW6BC0R"},
+    {"case":"cutlass_gemm_implicit","variant":1,"repeat":0,"cached":false,"build_s":44.895219539,"cpp_s":18.66360581,"nvcc_s":[44.011524792],"pch_consumed":false,"pch_key":null,"checks":10,"max_error":8.529314072802663e-06,"outputs":[16.0],"source_sha256":"b926f64a92cfd59840083d417ec09e2db0dd93574c97203ad4ac0d08925ccaee","container_id":"ta-01M1YVFWSFP2G8HR0D01V9F6JR"},
+    {"case":"cutlass_gemm_implicit","variant":0,"repeat":1,"cached":false,"build_s":46.818345875999995,"cpp_s":18.425574178,"nvcc_s":[45.895286961000004],"pch_consumed":false,"pch_key":null,"checks":10,"max_error":4.2646570364013314e-06,"outputs":[8.0],"source_sha256":"c0cc1e80088a516e7d3b0a21890de234e88d4cf0b2b278c6e8e74d7912d4a14d","container_id":"ta-01M1YVHXJN3XFWSW7KCQEPWSDR"},
+    {"case":"cutlass_gemm_implicit","variant":0,"repeat":1,"cached":true,"build_s":45.564655846,"cpp_s":8.125612553,"nvcc_s":[44.392049244999995],"pch_consumed":true,"pch_key":"cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38","checks":10,"max_error":4.2646570364013314e-06,"outputs":[8.0],"source_sha256":"c0cc1e80088a516e7d3b0a21890de234e88d4cf0b2b278c6e8e74d7912d4a14d","container_id":"ta-01M1YVHZH53EH13G4H5Y2PH07R"},
+    {"case":"cutlass_gemm_implicit","variant":1,"repeat":1,"cached":true,"build_s":50.934128441,"cpp_s":9.066257041,"nvcc_s":[49.628436083],"pch_consumed":true,"pch_key":"cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38","checks":10,"max_error":8.529314072802663e-06,"outputs":[16.0],"source_sha256":"b926f64a92cfd59840083d417ec09e2db0dd93574c97203ad4ac0d08925ccaee","container_id":"ta-01M1YVJ80Q2MBQKQXP74XRF2HR"},
+    {"case":"cutlass_gemm_implicit","variant":1,"repeat":1,"cached":false,"build_s":53.681691927,"cpp_s":21.277357138000003,"nvcc_s":[52.342306416],"pch_consumed":false,"pch_key":null,"checks":10,"max_error":8.529314072802663e-06,"outputs":[16.0],"source_sha256":"b926f64a92cfd59840083d417ec09e2db0dd93574c97203ad4ac0d08925ccaee","container_id":"ta-01M1YVJETDXTHV0W4J3E158HYR"}
+  ]
+}
diff --git a/docs/modal-pch-cache.md b/docs/modal-pch-cache.md
new file mode 100644
index 000000000..845900cba
--- /dev/null
+++ b/docs/modal-pch-cache.md
@@ -0,0 +1,102 @@
+# Modal Torch header cache
+
+The runner sets `CXX` to a GCC wrapper that loads a precompiled
+`torch/extension.h` from a Modal Volume. Ordinary `load_inline()` calls need no
+changes. Leave `use_pch=False`; PyTorch's separate native PCH option is unnecessary.
+The submission path assumes `load_inline()`.
+
+Only the Torch host header is cached. CUDA/CUTLASS templates, user headers, kernel
+objects, and shared libraries still compile in each submission container.
+Different sources and module names share a PCH when their host flags match;
+changing a user header after the Torch include still changes the compiled code.
+
+## Profiles and fallback
+
+The CPU warmer builds CPU-only and CUDA-linked profiles for these `extra_cflags`:
+`[]`, `["-O2"]`, `["-O3"]`, and `["-O3", "-ffast-math"]`.
+Keys include installed headers, compiler, Python/Torch versions, host flags, and
+include environment. CUDA-only flags do not change the host key.
+
+Unknown profiles compile normally; sources without the initial Torch include
+bypass PCH. With `no_implicit_headers=True`, explicitly include
+`<torch/extension.h>` first in C++ to use the cache. Python/Triton code that does
+not invoke `CXX` is unaffected.
+
+Misses and bypasses still pay wrapper startup cost. There is no guarantee of a
+speedup. Set `CXX=/usr/bin/g++` before building to bypass the wrapper entirely.
+`KERNELBOT_PCH_DISABLE=1` disables cache lookup but retains wrapper startup.
+
+## Deployment
+
+The main/dev workflow warms profiles before deploying GPU runners. Manually,
+run both commands in the same Modal profile/environment:
+
+```sh
+PYTHONPATH=src:src/runners uv run modal run src/runners/modal_runner.py::warm_pch
+PYTHONPATH=src:src/runners uv run modal deploy src/runners/modal_runner_archs.py
+```
+
+`KERNELBOT_PCH_VOLUME` overrides the default `kernelbot-torch-pch` Volume.
+`warm_pch --profiles cuda-default,cuda-O3` warms selected profiles.
+Rewarm after changing the image or installed headers, then replace existing
+runner containers to refresh their Volume snapshots.
+
+The CPU warmer accepts fixed profiles and commits completed headers. GPU runners
+mount the Volume read-only and use `restrict_modal_access=True`; otherwise the
+Modal upload API could bypass the read-only mount. Restricted runners cannot
+reload the Volume themselves. GCC validates a PCH before using it.
+
+## Benchmark
+
+```sh
+PYTHONPATH=src:src/runners KERNELBOT_PCH_VOLUME=kernelbot-pch-test \
+  uv run modal run scripts/modal_pch_benchmark.py
+# Actual tensor-core GEMM, two tile variants, two repetitions per mode:
+PYTHONPATH=src:src/runners KERNELBOT_PCH_VOLUME=kernelbot-pch-test \
+  uv run modal run scripts/modal_pch_benchmark.py --cases cutlass_gemm,cutlass_gemm_implicit --repeats 2
+```
+
+Every build runs in a fresh T4 container (4 CPU cores, 16 GiB), with normal Ninja
+parallelism. The baseline uses `/usr/bin/g++`; both modes use the same timing
+instrumentation. GCC header tracing verifies actual PCH use in a separate,
+untimed compilation. The tests check changed headers and source against distinct
+expected outputs. The script also retains file-based and plain C++ diagnostics;
+these are not representative submission paths.
+
+The CUTLASS case instantiates SM75 tensor-core GEMMs with FP16 inputs, FP32
+accumulation/output, 128×128×32 and 128×64×32 threadblock tiles. It checks
+`alpha * A @ B.T` for random and zero inputs, tile tails, and a non-default CUDA
+stream against a float64 PyTorch reference (`rtol=atol=2e-4`). This measures compile
+time, not GEMM execution speed.
+
+## Results
+
+Measured on CUDA 13.3 / PyTorch 2.12.0+cu130, T4. Values are medians from separate
+baseline and cached containers; two variants per case, with two repetitions for
+CUTLASS. The table covers `load_inline()` builds.
+
+| Build | Baseline | Cached |
+| --- | ---: | ---: |
+| Minimal CUDA headers | 18.85s | 10.00s |
+| CUTLASS GEMM, default implicit CUDA headers | 48.08s | 45.61s |
+| CUTLASS GEMM, `no_implicit_headers=True` | 19.95s | 12.58s |
+| Unwarmed host flags | 18.08s | 20.30s |
+
+With default implicit CUDA headers, the CUTLASS GEMM improved only **1.05×**,
+with overlapping baseline/cached timing ranges. With `no_implicit_headers=True`,
+it improved **1.59×** (37% less time). In that case, host compilation fell from
+19.10s to 7.88s, while NVCC took 12.54s and 11.53s and became the bottleneck.
+The PCH does not cache CUDA headers or templates. Keeping Torch headers out of
+the CUDA translation unit produces the larger gain for this GEMM; the fixture
+keeps Tensor handling in C++ and explicitly includes its required headers.
+
+All 160 GEMM checks passed, with maximum absolute error 8.53e-6. The remaining
+cases, including the supplemental diagnostics, passed another 76 checks.
+All eight warmup profiles also completed and
+committed on this image.
+
+The `load_inline()` case with unwarmed host flags was slower in this sample.
+Separate containers introduce timing variability, so the measured difference
+does not isolate wrapper overhead. Larger CUDA builds may hide the host-side
+savings entirely.
+[Per-container results and Modal runs](benchmarks/modal-pch.json).
diff --git a/pyproject.toml b/pyproject.toml
index 27038f2f6..2955ec938 100644
--- a/pyproject.toml
+++ b/pyproject.toml
@@ -15,7 +15,7 @@ dependencies = [
     "audioop-lts; python_version >= '3.13'",
     "python-dotenv",
     "requests",
-    "modal",
+    "modal>=1.5.0",
     "psycopg2-binary",
     "yoyo-migrations",
     "better_profanity",
diff --git a/scripts/modal_pch_benchmark.py b/scripts/modal_pch_benchmark.py
new file mode 100644
index 000000000..b4191ff16
--- /dev/null
+++ b/scripts/modal_pch_benchmark.py
@@ -0,0 +1,335 @@
+"""Benchmark Torch PCH reuse with CUDA, CUTLASS, file-based builds, and cache misses.
+
+PYTHONPATH=src:src/runners KERNELBOT_PCH_VOLUME=kernelbot-pch-test \
+    uv run modal run scripts/modal_pch_benchmark.py
+"""
+
+import json
+import statistics
+from pathlib import Path
+
+import modal
+from modal_runner import PCH_MOUNT, cuda_image, pch_volume, warm_pch
+
+app = modal.App("kernelbot-pch-matrix")
+app.include(warm_pch.app)
+CASES = (
+    "inline_simple", "load_files", "plain_cpp", "flag_miss",
+    "cutlass_gemm", "cutlass_gemm_implicit",
+)
+
+CACHED_CASES = set(CASES) - {"plain_cpp", "flag_miss"}
+
+# Instrument both modes identically; do not enable GCC's verbose header tracing
+# in timed builds. Baseline delegates directly to g++, cache mode to the wrapper.
+OBSERVER = r'''#!/usr/bin/env python3
+import json, os, subprocess, sys, time
+started = time.perf_counter()
+kind = "nvcc" if sys.argv[0].endswith("nvcc") else "cxx"
+compiler = "/usr/local/cuda/bin/nvcc" if kind == "nvcc" else os.environ["PCH_TEST_COMPILER"]
+code = subprocess.call([compiler, *sys.argv[1:]])
+if "-c" in sys.argv:
+    with open(os.environ["PCH_TEST_COMMANDS"] + "." + kind, "a") as f:
+        f.write(json.dumps({"args": sys.argv[1:], "seconds": time.perf_counter()-started}) + "\n")
+sys.exit(code)
+'''
+
+CPP = r'''
+#include <torch/extension.h>
+#include <c10/cuda/CUDAGuard.h>
+#include <c10/cuda/CUDAStream.h>
+#include "variant.h"
+
+void add_impl(const void*, const void*, void*, int, cudaStream_t);
+int version() { return kVariant; }
+torch::Tensor add_cuda(torch::Tensor a, torch::Tensor b) {
+    TORCH_CHECK(a.is_cuda() && b.is_cuda() && a.device() == b.device());
+    TORCH_CHECK(a.scalar_type() == torch::kFloat16 && b.scalar_type() == torch::kFloat16);
+    TORCH_CHECK(a.is_contiguous() && b.is_contiguous() && a.sizes() == b.sizes());
+    c10::cuda::CUDAGuard guard(a.device());
+    auto out = torch::empty_like(a);
+    add_impl(a.data_ptr(), b.data_ptr(), out.data_ptr(), a.numel(),
+             c10::cuda::getCurrentCUDAStream().stream());
+    TORCH_CHECK(cudaGetLastError() == cudaSuccess);
+    return out;
+}
+'''
+
+CUDA = r'''
+#include <cuda_runtime.h>
+#include <cuda_fp16.h>
+#include "variant.h"
+__global__ void add_kernel(const __half* a, const __half* b, __half* out, int n) {
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < n) out[i] = __hadd(__hadd(a[i], b[i]), __float2half(float(kVariant)));
+}
+void add_impl(const void* a, const void* b, void* out, int n, cudaStream_t stream) {
+    constexpr int threads = THREADS;
+    add_kernel<<<(n + threads - 1) / threads, threads, 0, stream>>>(
+        static_cast<const __half*>(a), static_cast<const __half*>(b), static_cast<__half*>(out), n);
+}
+'''
+
+GEMM_CPP = r'''
+#include <torch/extension.h>
+#include <c10/cuda/CUDAGuard.h>
+#include <c10/cuda/CUDAStream.h>
+#include "variant.h"
+
+int gemm_impl(const void*, const void*, float*, int, int, int, cudaStream_t);
+int version() { return kVariant; }
+torch::Tensor add_cuda(torch::Tensor a, torch::Tensor b) {
+    TORCH_CHECK(a.is_cuda() && b.is_cuda() && a.device() == b.device());
+    TORCH_CHECK(a.scalar_type() == torch::kFloat16 && b.scalar_type() == torch::kFloat16);
+    TORCH_CHECK(a.dim() == 2 && b.dim() == 2 && a.is_contiguous() && b.is_contiguous());
+    TORCH_CHECK(a.size(1) == b.size(1) && a.size(1) % 8 == 0 && b.size(0) % 4 == 0);
+    TORCH_CHECK(a.size(0) > 0 && b.size(0) > 0 && a.size(1) > 0);
+    TORCH_CHECK(a.size(0) < INT_MAX && b.size(0) < INT_MAX && a.size(1) < INT_MAX);
+    c10::cuda::CUDAGuard guard(a.device());
+    auto out = torch::empty({a.size(0), b.size(0)}, a.options().dtype(torch::kFloat32));
+    int status = gemm_impl(a.data_ptr(), b.data_ptr(), out.data_ptr<float>(),
+                          a.size(0), b.size(0), a.size(1),
+                          c10::cuda::getCurrentCUDAStream().stream());
+    TORCH_CHECK(status == 0, "CUTLASS status: ", status);
+    return out;
+}
+'''
+
+GEMM_CUDA = r'''
+#include <cuda_runtime.h>
+#include <cutlass/gemm/device/gemm.h>
+#include "variant.h"
+
+// A[M,K] and B[N,K] are contiguous: interpret B as column-major [K,N].
+using Gemm = cutlass::gemm::device::Gemm<
+    cutlass::half_t, cutlass::layout::RowMajor,
+    cutlass::half_t, cutlass::layout::ColumnMajor,
+    float, cutlass::layout::RowMajor, float,
+    cutlass::arch::OpClassTensorOp, cutlass::arch::Sm75,
+    cutlass::gemm::GemmShape<128, TILE_N, 32>,
+    cutlass::gemm::GemmShape<64, 64, 32>,
+    cutlass::gemm::GemmShape<16, 8, 8>,
+    cutlass::epilogue::thread::LinearCombination<float, 4, float, float>,
+    cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<>, 2>;
+
+int gemm_impl(const void* a, const void* b, float* out, int m, int n, int k, cudaStream_t stream) {
+    Gemm::Arguments args({m, n, k},
+        {static_cast<const cutlass::half_t*>(a), k},
+        {static_cast<const cutlass::half_t*>(b), k},
+        {out, n}, {out, n}, {float(kVariant), 0.0f});
+    auto status = Gemm::can_implement(args);
+    if (status != cutlass::Status::kSuccess) return int(status);
+    if (Gemm::get_workspace_size(args) != 0) return int(cutlass::Status::kErrorInvalidProblem);
+    return int(Gemm{}(args, nullptr, stream));
+}
+'''
+
+
+@app.function(
+    image=cuda_image,
+    gpu="T4",
+    cpu=4,
+    memory=16384,
+    timeout=900,
+    max_containers=4,
+    single_use_containers=True,
+    restrict_modal_access=True,
+    volumes={PCH_MOUNT: pch_volume.with_mount_options(read_only=True)},
+)
+def trial(case: str, variant: int, enabled: bool, repeat: int = 0):
+    import hashlib
+    import os
+    import re
+    import subprocess
+    import tempfile
+    import time
+
+    import torch
+
+    started = time.perf_counter()
+    os.environ.pop("MAX_JOBS", None)
+    os.environ.pop("KERNELBOT_PCH_DISABLE", None)
+    os.environ.pop("KERNELBOT_PCH_TRACE", None)
+    os.environ.pop("KERNELBOT_PCH_WRITE", None)
+    os.environ["TORCH_CUDA_ARCH_LIST"] = "7.5"
+    os.environ["PCH_TEST_COMPILER"] = "/opt/kernelbot-pch/compiler.py" if enabled else "/usr/bin/g++"
+    with tempfile.TemporaryDirectory() as directory:
+        root = Path(directory)
+        observer = root / "observe-cxx"
+        observer.write_text(OBSERVER)
+        observer.chmod(0o755)
+        nvcc = root / "observe-nvcc"
+        nvcc.write_text(OBSERVER)
+        nvcc.chmod(0o755)
+        os.environ["PYTORCH_NVCC"] = str(nvcc)
+        commands = root / "commands.jsonl"
+        os.environ["CXX"] = str(observer)
+        os.environ["PCH_TEST_COMMANDS"] = str(commands)
+        header = f"#pragma once\nconstexpr int kVariant = {variant + 1};\n"
+        (root / "variant.h").write_text(header)
+        cpp = CPP
+        cuda = CUDA.replace("THREADS", str(256 if variant == 0 else 128))
+        if case.startswith("cutlass_gemm"):
+            cpp = GEMM_CPP
+            cuda = GEMM_CUDA.replace("TILE_N", str(128 if variant == 0 else 64))
+        name = "pch_matrix_shared" if case == "load_files" else f"pch_matrix_{case}_{variant}_{int(enabled)}"
+        if case == "plain_cpp":
+            cpp = '#include <array>\n#include <cstdio>\n#include "variant.h"\nint main() {\n'
+            cpp += 'std::array<int, 3> a{1, 2, 3}; std::printf("%d\\n", a[1] + kVariant); }\n'
+            (root / "main.cpp").write_text(cpp)
+            # Multiple direct compilations resolve wrapper overhead at millisecond scale.
+            times = []
+            logs = []
+            for _ in range(10):
+                tick = time.perf_counter()
+                result = subprocess.run(
+                    [str(observer), "-std=c++20", "-c", str(root / "main.cpp"), "-o", str(root / "main.o")],
+                    capture_output=True, text=True, check=True,
+                )
+                times.append(time.perf_counter() - tick)
+                logs.append(result.stdout + result.stderr)
+            subprocess.run(["g++", str(root / "main.o"), "-o", str(root / "main")], check=True)
+            actual = int(subprocess.check_output([str(root / "main")]))
+            assert actual == variant + 3
+            wall = statistics.median(times)
+            log = "\n".join(logs)
+            correctness_checks = 1
+            device_outputs = [actual]
+        else:
+            (root / "binding.cpp").write_text(
+                cpp + '\nPYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {\n'
+                'm.def("add_cuda", &add_cuda); m.def("version", &version); }\n'
+            )
+            (root / "kernel.cu").write_text(cuda)
+            (root / "cpp.txt").write_text(cpp)
+            (root / "cuda.txt").write_text(cuda)
+            # A subprocess captures compiler output and limits timed work to load/load_inline.
+            loader = '''
+import json, os, time
+from pathlib import Path
+import torch
+from torch.utils.cpp_extension import load, load_inline
+root = Path(os.environ["PCH_TEST_ROOT"])
+case = os.environ["PCH_TEST_CASE"]
+variant = int(os.environ["PCH_TEST_VARIANT"])
+options = dict(name=os.environ["PCH_TEST_NAME"], build_directory=str(root), verbose=True)
+if case.startswith("cutlass_gemm"):
+    # Restore CUDA half operators for CUTLASS; host flags stay unchanged.
+    options["extra_cuda_cflags"] = [
+        "-U__CUDA_NO_HALF_OPERATORS__", "-U__CUDA_NO_HALF_CONVERSIONS__",
+        "-U__CUDA_NO_HALF2_OPERATORS__", "-U__CUDA_NO_BFLOAT16_CONVERSIONS__",
+        "--expt-relaxed-constexpr",
+        "-O3", "-I/opt/cutlass/include", "-I/opt/cutlass/tools/util/include",
+    ]
+started = time.perf_counter()
+if case == "load_files":
+    module = load(sources=[str(root / "binding.cpp"), str(root / "kernel.cu")], **options)
+else:
+    module = load_inline(cpp_sources=(root / "cpp.txt").read_text(), cuda_sources=(root / "cuda.txt").read_text(),
+        functions=["add_cuda", "version"], no_implicit_headers=case != "cutlass_gemm_implicit",
+        extra_cflags=[f"-DPCH_CASE={variant+1}"] if case == "flag_miss" else [], **options)
+seconds = time.perf_counter() - started
+assert module.version() == variant + 1
+outputs = []
+errors = []
+if case.startswith("cutlass_gemm"):
+    torch.manual_seed(17)
+    with torch.cuda.stream(torch.cuda.Stream()):
+        for m, n, k in ((128, 128, 64), (129, 136, 64), (256, 256, 128), (512, 512, 256)):
+            a = torch.randn(m, k, device="cuda", dtype=torch.float16) / 4
+            b = torch.randn(n, k, device="cuda", dtype=torch.float16) / 4
+            for zero in (False, True):
+                if zero:
+                    a.zero_()
+                actual = module.add_cuda(a, b)
+                reference = (a.double() @ b.double().T) * (variant + 1)
+                torch.testing.assert_close(actual.double(), reference, rtol=2e-4, atol=2e-4)
+                errors.append(float((actual.double() - reference).abs().max()))
+        a = torch.ones(8, 8, device="cuda", dtype=torch.float16)
+        actual = module.add_cuda(a, a)
+        torch.testing.assert_close(actual, torch.full_like(actual, 8 * (variant + 1)), rtol=0, atol=0)
+        outputs.append(float(actual[0, 0]))
+else:
+    for n in (1, 127, 128, 129, 1025):
+        a = torch.arange(n, device="cuda", dtype=torch.float16) % 8
+        b = torch.ones_like(a)
+        actual = module.add_cuda(a, b)
+        expected = a + b + (variant + 1)
+        torch.testing.assert_close(actual, expected, rtol=0, atol=0)
+        outputs.append(float(actual[0]))
+(root / "result.json").write_text(json.dumps(dict(seconds=seconds, outputs=outputs, errors=errors)))
+'''
+            env = os.environ | {
+                "PCH_TEST_ROOT": str(root), "PCH_TEST_CASE": case,
+                "PCH_TEST_VARIANT": str(variant), "PCH_TEST_NAME": name,
+            }
+            result = subprocess.run(["python3", "-c", loader], env=env, capture_output=True, text=True)
+            log = result.stdout + result.stderr
+            assert result.returncode == 0, log[-14000:]
+            measurement = json.loads((root / "result.json").read_text())
+            wall = measurement["seconds"]
+            device_outputs = measurement["outputs"]
+            correctness_checks = 10 if case.startswith("cutlass_gemm") else 6
+            times = [wall]
+        observed = [json.loads(line) for line in Path(str(commands) + ".cxx").read_text().splitlines()]
+        hit_keys = re.findall(r"\[kernelbot-pch\] hit (\w+)", log)
+        misses = "[kernelbot-pch] miss" in log
+        expected_hit = enabled and case in CACHED_CASES
+        assert bool(hit_keys) == expected_hit, log[-14000:]
+        assert misses == (enabled and case == "flag_miss"), log[-14000:]
+        # Independently verify actual GCC consumption after all timing/correctness checks.
+        # This extra compilation and its header tracing are excluded from measured times.
+        args = observed[0]["args"]
+        args[args.index("-o") + 1] = str(root / "diagnostic.o")
+        diagnostic = subprocess.run(
+            [os.environ["PCH_TEST_COMPILER"], *args, "-H"], capture_output=True, text=True,
+        )
+        assert diagnostic.returncode == 0, diagnostic.stderr[-10000:]
+        consumed = "! /kernelbot-pch/" in diagnostic.stderr
+        assert consumed == expected_hit, diagnostic.stderr[-10000:]
+        cuda_commands = Path(str(commands) + ".nvcc")
+        cuda_times = [json.loads(line)["seconds"] for line in cuda_commands.read_text().splitlines()] \
+            if cuda_commands.exists() else []
+        return {
+            "case": case, "variant": variant, "enabled": enabled, "repeat": repeat, "extension_name": name,
+            "seconds": wall, "timing_samples": times,
+            "host_compile_seconds": statistics.median(row["seconds"] for row in observed),
+            "pch_consumed": consumed, "hit_keys": hit_keys, "cache_miss": misses,
+            "correctness_checks": correctness_checks, "outputs": device_outputs,
+            "source_sha256": hashlib.sha256((cpp + cuda + header).encode()).hexdigest(),
+            "container_id": os.environ["MODAL_TASK_ID"], "image_id": os.environ["MODAL_IMAGE_ID"],
+            "gpu": torch.cuda.get_device_name(), "torch": str(torch.__version__),
+            "nvcc_seconds": cuda_times,
+            "max_absolute_error": max(measurement["errors"], default=0) if case.startswith("cutlass_gemm") else 0,
+            "remote_seconds": time.perf_counter() - started,
+        }
+
+
+@app.local_entrypoint()
+def main(output: str = "/tmp/kernelbot-pch-matrix.json", cases: str = "", repeats: int = 1):
+    selected = cases.split(",") if cases else CASES
+    if not set(selected) <= set(CASES):
+        raise ValueError(f"Choose from {CASES}")
+    print(warm_pch.remote(profiles="cuda-default"))
+    jobs = [
+        (case, variant, enabled, repeat)
+        for repeat in range(repeats)
+        for case in selected
+        for variant in (0, 1)
+        for enabled in ((False, True) if variant == 0 else (True, False))
+    ]
+    rows = []
+    for row in trial.starmap(jobs):
+        rows.append(row)
+        Path(output).write_text(json.dumps(rows, indent=2) + "\n")
+        print(json.dumps(row), flush=True)
+    assert len({row["container_id"] for row in rows}) == len(rows)
+    hit_keys = {key for row in rows for key in row["hit_keys"]}
+    assert len(hit_keys) == (1 if set(selected) & CACHED_CASES else 0)
+    for case in selected:
+        group = [row for row in rows if row["case"] == case]
+        assert len({row["source_sha256"] for row in group}) == 2
+        assert group[0]["outputs"] != group[2]["outputs"], "Expected distinct results from changed submissions"
+        cold = statistics.median(row["seconds"] for row in group if not row["enabled"])
+        warm = statistics.median(row["seconds"] for row in group if row["enabled"])
+        print(f"{case}: baseline={cold:.4f}s cache={warm:.4f}s ratio={cold / warm:.3f}x")
diff --git a/src/runners/modal_runner.py b/src/runners/modal_runner.py
index 5c07faf34..b177f99d1 100644
--- a/src/runners/modal_runner.py
+++ b/src/runners/modal_runner.py
@@ -1,8 +1,11 @@
+import os
 import signal
+import subprocess
 import traceback
 from contextlib import contextmanager
+from pathlib import Path
 
-from modal import App, Image
+from modal import App, Image, Volume
 
 from libkernelbot.run_eval import FullResult, SystemInfo, run_config
 
@@ -113,6 +116,24 @@
     )
 )
 
+# The writer is a trusted CPU-only warmup function. Submission containers mount
+# this Volume read-only; no submission objects or shared libraries are persisted.
+PCH_MOUNT = "/kernelbot-pch"
+PCH_VOLUME_NAME = os.environ.get("KERNELBOT_PCH_VOLUME", "kernelbot-torch-pch")
+pch_volume = Volume.from_name(PCH_VOLUME_NAME, create_if_missing=True)
+pch_files = Path(__file__).parent / "pch"
+cuda_image = (
+    cuda_image.add_local_file(
+        pch_files / "compiler.py", "/opt/kernelbot-pch/compiler.py", copy=True
+    )
+    .add_local_file(pch_files / "warm.py", "/opt/kernelbot-pch/warm.py", copy=True)
+    .run_commands(
+        "chmod +x /opt/kernelbot-pch/compiler.py",
+        "python /opt/kernelbot-pch/compiler.py --kernelbot-fingerprint",
+    )
+    .env({"CXX": "/opt/kernelbot-pch/compiler.py", "KERNELBOT_PCH_VOLUME": PCH_VOLUME_NAME})
+)
+
 cuda_image = cuda_image.add_local_python_source(
     "libkernelbot",
     "modal_runner",
@@ -168,3 +189,28 @@ def modal_run_config(  # noqa: C901
             runs={},
             system=SystemInfo(),
         )
+
+
+@app.function(
+    image=cuda_image,
+    cpu=4,
+    memory=16384,
+    timeout=1800,
+    max_containers=1,
+    volumes={PCH_MOUNT: pch_volume},
+)
+def warm_pch(profiles: str = ""):
+    # A comma-separated selector also works with the Modal CLI.
+    """Run before deploying GPU runners: modal run modal_runner.py::warm_pch."""
+    allowed = {
+        f"{device}-{flags}"
+        for device in ("cpu", "cuda")
+        for flags in ("default", "O2", "O3", "O3-fast-math")
+    }
+    selected = profiles.split(",") if profiles else []
+    if not set(selected) <= allowed:
+        raise ValueError(f"Unknown profiles; choose from {sorted(allowed)}")
+    pch_volume.reload()
+    subprocess.run(["python3", "/opt/kernelbot-pch/warm.py", *selected], check=True)
+    pch_volume.commit()
+    return "Torch PCH profiles committed"
diff --git a/src/runners/modal_runner_archs.py b/src/runners/modal_runner_archs.py
index 55f025a0a..c33813056 100644
--- a/src/runners/modal_runner_archs.py
+++ b/src/runners/modal_runner_archs.py
@@ -1,6 +1,6 @@
 # This file contains wrapper functions for running
 # Modal apps on specific devices. We will fix this later.
-from modal_runner import MODAL_RUN_TIMEOUT_SECONDS, app, cuda_image, modal_run_config
+from modal_runner import MODAL_RUN_TIMEOUT_SECONDS, PCH_MOUNT, app, cuda_image, modal_run_config, pch_volume
 
 from libkernelbot.validation_runtime import run_validation_config
 
@@ -12,6 +12,8 @@
         image=cuda_image,
         name=f"run_cuda_script_{gpu_slug}",
         serialized=True,
+        restrict_modal_access=True,
+        volumes={PCH_MOUNT: pch_volume.with_mount_options(read_only=True)},
         timeout=MODAL_RUN_TIMEOUT_SECONDS,
     )(modal_run_config)
     app.function(
@@ -19,6 +21,8 @@
         image=cuda_image,
         name=f"run_pytorch_script_{gpu_slug}",
         serialized=True,
+        restrict_modal_access=True,
+        volumes={PCH_MOUNT: pch_volume.with_mount_options(read_only=True)},
         timeout=MODAL_RUN_TIMEOUT_SECONDS,
     )(modal_run_config)
 
@@ -27,5 +31,7 @@
     image=cuda_image,
     name="run_validation_script_b200",
     serialized=True,
+    restrict_modal_access=True,
+    volumes={PCH_MOUNT: pch_volume.with_mount_options(read_only=True)},
     timeout=MODAL_RUN_TIMEOUT_SECONDS,
 )(run_validation_config)
diff --git a/src/runners/pch/compiler.py b/src/runners/pch/compiler.py
new file mode 100755
index 000000000..f7268b677
--- /dev/null
+++ b/src/runners/pch/compiler.py
@@ -0,0 +1,145 @@
+#!/usr/bin/env python3
+"""Transparent GCC wrapper for the trusted, read-only torch header PCH Volume.
+
+Only the first implicit load_inline header is cached. Kernel objects and shared
+libraries stay in the submission container's normal torch extension cache.
+"""
+
+import hashlib
+import json
+import os
+import re
+import subprocess
+import sys
+import time
+import uuid
+from pathlib import Path
+
+COMPILER = "/usr/bin/g++"
+ROOT = Path(os.environ.get("KERNELBOT_PCH_ROOT", "/kernelbot-pch"))
+FINGERPRINT = Path("/opt/kernelbot-pch/fingerprint")
+
+
+def compile_flags(args):
+    """Return preprocessing/compilation flags, excluding per-module build outputs."""
+    if args.count("-c") != 1 or "-o" not in args:
+        return None
+    if args.index("-c") + 1 == len(args):
+        return None
+    source = args[args.index("-c") + 1]
+    if not source.endswith(".cpp"):
+        return None
+    # Do not change header order or inject torch into no_implicit_headers users.
+    try:
+        prefix = Path(source).read_text()
+    except (OSError, UnicodeError):
+        return None
+    if not re.match(r"\s*#\s*include\s*<torch/extension\.h>", prefix):
+        return None
+    flags = []
+    index = 0
+    while index < len(args):
+        arg = args[index]
+        if arg in {"-c", "-o", "-MF", "-MT", "-MQ"}:
+            index += 2
+            continue
+        if arg not in {"-MMD", "-MD", "-H"} and not arg.startswith("-DTORCH_EXTENSION_NAME="):
+            flags.append(arg)
+        index += 1
+    return flags
+
+
+def cache_key(flags, fingerprint):
+    # Include header-search environment as well as argv. Never share incompatible
+    # macro, optimization, ABI, language-standard or include-path configurations.
+    payload = [
+        fingerprint,
+        flags,
+        {
+            k: os.environ.get(k)
+            for k in (
+                "CPATH",
+                "CPLUS_INCLUDE_PATH",
+                "C_INCLUDE_PATH",
+                "GCC_EXEC_PREFIX",
+                "COMPILER_PATH",
+            )
+        },
+    ]
+    return hashlib.sha256(json.dumps(payload, sort_keys=True).encode()).hexdigest()
+
+
+def make_fingerprint():
+    """Run once when building the immutable image, never on a submission's GPU."""
+    import sysconfig
+
+    import torch
+
+    digest = hashlib.sha256()
+    digest.update(Path(__file__).read_bytes())
+    digest.update(subprocess.check_output([COMPILER, "-v"], stderr=subprocess.STDOUT))
+    digest.update(sys.version.encode())
+    digest.update(torch.__version__.encode())
+    digest.update(Path("/etc/os-release").read_bytes())
+    # Hash installed headers, including CUDA and compiler headers, so rebuilding
+    # the image with changed header contents cannot reuse stale PCHs.
+    roots = [
+        Path(torch.__file__).parent / "include",
+        Path(sysconfig.get_path("include")),
+        Path("/usr/include"),
+        *Path("/usr/lib/gcc").glob("*/*/include*"),
+        Path("/usr/local/cuda/include"),
+        Path("/opt/cutlass/include"),
+        Path("/opt/cutlass/tools/util/include"),
+        Path("/opt/mathdx/include"),
+        Path("/opt/mathdx/external/cutlass/include"),
+    ]
+    for root in roots:
+        for path in sorted(root.rglob("*")):
+            if path.is_file():
+                digest.update(str(path).encode())
+                digest.update(path.read_bytes())
+    FINGERPRINT.write_text(digest.hexdigest())
+
+
+def run_compiler(args):
+    trace = os.environ.get("KERNELBOT_PCH_TRACE") == "1" and "-c" in args
+    started = time.perf_counter()
+    result = subprocess.call([COMPILER, *args, *(["-H"] if trace and "-include" in args else [])])
+    if trace:
+        print(f"[kernelbot-pch] compile_seconds={time.perf_counter() - started:.3f}", flush=True)
+    return result
+
+
+def main(args):
+    if args == ["--kernelbot-fingerprint"]:
+        make_fingerprint()
+        return 0
+    if os.environ.get("KERNELBOT_PCH_DISABLE") == "1":
+        return run_compiler(args)
+    flags = compile_flags(args)
+    if flags is None:
+        return run_compiler(args)
+    key = cache_key(flags, FINGERPRINT.read_text())
+    entry = ROOT / key
+    header = entry / "torch.h"
+    pch = entry / "torch.h.gch"
+    # Write mode exists only on the trusted CPU warmer, never the GPU runner.
+    if not pch.exists() and os.environ.get("KERNELBOT_PCH_WRITE") == "1":
+        entry.mkdir(parents=True, exist_ok=True)
+        header.write_text("#include <torch/extension.h>\n")
+        temporary = entry / f"torch.h.gch.{uuid.uuid4().hex}.tmp"
+        subprocess.run(
+            [COMPILER, *flags, "-x", "c++-header", str(header), "-o", str(temporary)], check=True
+        )
+        temporary.replace(pch)
+        print(f"[kernelbot-pch] built {key}", flush=True)
+    if pch.is_file():
+        print(f"[kernelbot-pch] hit {key}", flush=True)
+        return run_compiler([*args, "-include", str(header), "-Winvalid-pch"])
+    print(f"[kernelbot-pch] miss {key}; compiling normally", flush=True)
+    return run_compiler(args)
+
+
+if __name__ == "__main__":
+    sys.exit(main(sys.argv[1:]))
diff --git a/src/runners/pch/warm.py b/src/runners/pch/warm.py
new file mode 100644
index 000000000..fa369c794
--- /dev/null
+++ b/src/runners/pch/warm.py
@@ -0,0 +1,36 @@
+"""Trusted CPU-only header warmup. No submission source or flags are accepted."""
+
+import os
+import sys
+import tempfile
+
+# Keep the common profiles bounded. Unsupported flags simply compile normally.
+PROFILES = {"default": [], "O2": ["-O2"], "O3": ["-O3"], "O3-fast-math": ["-O3", "-ffast-math"]}
+
+
+def warm(profiles=()):
+    from torch.utils.cpp_extension import load_inline
+
+    os.environ["KERNELBOT_PCH_WRITE"] = "1"
+    os.environ["MAX_JOBS"] = "1"
+    # PyTorch computes CUDA flags even for a C++-only warmup with with_cuda=True.
+    # The host PCH is architecture-independent; no GPU or CUDA source is needed.
+    os.environ["TORCH_CUDA_ARCH_LIST"] = "7.5"
+    for with_cuda in (False, True):
+        for index, (name, flags) in enumerate(PROFILES.items()):
+            if profiles and f"{'cuda' if with_cuda else 'cpu'}-{name}" not in profiles:
+                continue
+            with tempfile.TemporaryDirectory() as build:
+                load_inline(
+                    name=f"kernelbot_pch_warm_{int(with_cuda)}_{index}",
+                    cpp_sources="int warm() { return 1; }",
+                    functions=["warm"],
+                    extra_cflags=flags,
+                    with_cuda=with_cuda,
+                    build_directory=build,
+                    verbose=True,
+                )
+
+
+if __name__ == "__main__":
+    warm(sys.argv[1:])
diff --git a/tests/test_pch_compiler.py b/tests/test_pch_compiler.py
new file mode 100644
index 000000000..09ab44227
--- /dev/null
+++ b/tests/test_pch_compiler.py
@@ -0,0 +1,117 @@
+"""PCH cache isolation and transparent fallback without a local torch install."""
+
+import importlib.util
+from pathlib import Path
+
+import pytest
+
+spec = importlib.util.spec_from_file_location(
+    "pch_compiler", Path(__file__).parents[1] / "src/runners/pch/compiler.py"
+)
+compiler = importlib.util.module_from_spec(spec)
+spec.loader.exec_module(compiler)
+
+
+def command(tmp_path, name="first", extra=()):
+    source = tmp_path / f"{name}.cpp"
+    source.write_text("#include <torch/extension.h>\nint value() { return 1; }\n")
+    return [
+        "-MMD",
+        "-MF",
+        f"{name}.d",
+        f"-DTORCH_EXTENSION_NAME={name}",
+        "-DTORCH_API_INCLUDE_EXTENSION_H",
+        "-isystem",
+        "/torch/include",
+        "-fPIC",
+        "-std=c++17",
+        *extra,
+        "-c",
+        str(source),
+        "-o",
+        f"{name}.o",
+    ]
+
+
+def test_different_kernels_share_header_key(tmp_path):
+    first = compiler.compile_flags(command(tmp_path, "first"))
+    second_args = command(tmp_path, "second")
+    # Source bodies and user headers following Torch are compiled anew, not cached.
+    (tmp_path / "second.cpp").write_text(
+        '#include <torch/extension.h>\n#include "custom.h"\nint value() { return CUSTOM + 2; }\n'
+    )
+    (tmp_path / "custom.h").write_text("#define CUSTOM 7\n")
+    second = compiler.compile_flags(second_args)
+    assert compiler.cache_key(first, "image") == compiler.cache_key(second, "image")
+
+
+@pytest.mark.parametrize(
+    "flags",
+    [
+        ["-O3"],
+        ["-D_GLIBCXX_USE_CXX11_ABI=0"],
+        ["-std=c++20"],
+        ["-I/other/headers"],
+        ["-DOTHER=1"],
+        ["-ffast-math"],
+    ],
+)
+def test_incompatible_flags_do_not_reuse(tmp_path, flags):
+    normal = compiler.compile_flags(command(tmp_path))
+    changed = compiler.compile_flags(command(tmp_path, extra=flags))
+    assert compiler.cache_key(normal, "image") != compiler.cache_key(changed, "image")
+
+
+def test_header_and_environment_invalidation(tmp_path, monkeypatch):
+    flags = compiler.compile_flags(command(tmp_path))
+    key = compiler.cache_key(flags, "image")
+    assert key != compiler.cache_key(flags, "new headers/compiler/torch")
+    monkeypatch.setenv("CPLUS_INCLUDE_PATH", "/new/headers")
+    assert key != compiler.cache_key(flags, "image")
+
+
+@pytest.mark.parametrize(
+    "source", ["int value() { return 1; }", "#define CUSTOM 1\n#include <torch/extension.h>"]
+)
+def test_does_not_inject_headers_or_reorder_macros(tmp_path, source):
+    args = command(tmp_path)
+    Path(args[args.index("-c") + 1]).write_text(source)
+    assert compiler.compile_flags(args) is None
+
+
+@pytest.mark.parametrize(
+    "args", [["--version"], ["-v"], ["a.o", "-shared", "-o", "a.so"], ["-o", "a.o", "-c"]]
+)
+def test_pass_through_non_compilation(args):
+    assert compiler.compile_flags(args) is None
+
+
+def test_miss_is_read_only_and_falls_back(tmp_path, monkeypatch):
+    args = command(tmp_path)
+    fingerprint = tmp_path / "fingerprint"
+    fingerprint.write_text("image")
+    cache = tmp_path / "cache"
+    monkeypatch.setattr(compiler, "FINGERPRINT", fingerprint)
+    monkeypatch.setattr(compiler, "ROOT", cache)
+    monkeypatch.delenv("KERNELBOT_PCH_WRITE", raising=False)
+    calls = []
+    monkeypatch.setattr(compiler, "run_compiler", lambda argv: calls.append(argv) or 0)
+    assert compiler.main(args) == 0
+    assert calls == [args]
+    assert not cache.exists()
+
+
+def test_hit_injects_pch_without_reusing_kernel_objects(tmp_path, monkeypatch):
+    args = command(tmp_path)
+    fingerprint = tmp_path / "fingerprint"
+    fingerprint.write_text("image")
+    monkeypatch.setattr(compiler, "FINGERPRINT", fingerprint)
+    monkeypatch.setattr(compiler, "ROOT", tmp_path)
+    key = compiler.cache_key(compiler.compile_flags(args), "image")
+    entry = tmp_path / key
+    entry.mkdir()
+    (entry / "torch.h.gch").write_bytes(b"cached pch")
+    calls = []
+    monkeypatch.setattr(compiler, "run_compiler", lambda argv: calls.append(argv) or 0)
+    assert compiler.main(args) == 0
+    assert calls == [[*args, "-include", str(entry / "torch.h"), "-Winvalid-pch"]]
diff --git a/tests/test_pch_warmup.py b/tests/test_pch_warmup.py
new file mode 100644
index 000000000..d9a235369
--- /dev/null
+++ b/tests/test_pch_warmup.py
@@ -0,0 +1,57 @@
+"""The shared writer only accepts fixed profiles and publishes successful builds."""
+
+import importlib
+import subprocess
+from pathlib import Path
+from types import SimpleNamespace
+
+import pytest
+
+
+@pytest.fixture
+def runner(monkeypatch):
+    monkeypatch.syspath_prepend(str(Path(__file__).parents[1] / "src/runners"))
+    module = importlib.import_module("modal_runner")
+    events = []
+    monkeypatch.setattr(
+        module,
+        "pch_volume",
+        SimpleNamespace(
+            reload=lambda: events.append("reload"), commit=lambda: events.append("commit")
+        ),
+    )
+    monkeypatch.setattr(module.subprocess, "run", lambda args, **kw: events.append(args))
+    return module, events
+
+
+def test_only_publishes_after_successful_warmup(runner):
+    module, events = runner
+    module.warm_pch.local("cuda-default,cuda-O3")
+    assert events == [
+        "reload",
+        ["python3", "/opt/kernelbot-pch/warm.py", "cuda-default", "cuda-O3"],
+        "commit",
+    ]
+
+
+@pytest.mark.parametrize(
+    "profiles", ["-O3", "cuda-default; touch /tmp/oops", "cuda-unknown", "cpu-default,"]
+)
+def test_rejects_arbitrary_compiler_arguments(runner, profiles):
+    module, events = runner
+    with pytest.raises(ValueError, match="Unknown profiles"):
+        module.warm_pch.local(profiles)
+    assert not events
+
+
+def test_failed_warmup_does_not_commit(runner, monkeypatch):
+    module, events = runner
+
+    def fail(*args, **kwargs):
+        raise subprocess.CalledProcessError(1, args[0])
+
+    monkeypatch.setattr(module.subprocess, "run", fail)
+    with pytest.raises(subprocess.CalledProcessError):
+        module.warm_pch.local("cuda-default")
+    assert events == ["reload"]
+
diff --git a/uv.lock b/uv.lock
index 75feb80d8..1d0db15cd 100644
--- a/uv.lock
+++ b/uv.lock
@@ -2,7 +2,8 @@ version = 1
 revision = 3
 requires-python = ">=3.11"
 resolution-markers = [
-    "python_full_version >= '3.13'",
+    "python_full_version >= '3.14'",
+    "python_full_version == '3.13.*'",
     "python_full_version < '3.13'",
 ]
 
@@ -236,6 +237,54 @@ wheels = [
     { url = "https://files.pythonhosted.org/packages/f3/dd/0b074d89e903cc771721cde2c4bf3d8c9d114b5bd791af5c62bcf5fb9459/better_profanity-0.7.0-py3-none-any.whl", hash = "sha256:bd4c529ea6aa2db1aaa50524be1ed14d0fe5c664f1fd88c8bc388c7e9f9f00e8", size = 46104, upload-time = "2020-11-02T10:49:56.066Z" },
 ]
 
+[[package]]
+name = "cbor2"
+version = "6.1.4"
+source = { registry = "https://pypi.org/simple" }
+sdist = { url = "https://files.pythonhosted.org/packages/c6/14/b02446bacfe44351b1689c04937ade007588f44570431880a6937e525e6c/cbor2-6.1.4.tar.gz", hash = "sha256:01ecc79a28f33d17331943ce508fc1e21f4b06553c73f874f4c77120d72b2ef9", size = 90840, upload-time = "2026-08-01T20:41:39.797Z" }
+wheels = [
+    { url = "https://files.pythonhosted.org/packages/a7/84/1e363301c06f509963d134f5479e82b3ade87fb1495ddacf9bf7ff24ac42/cbor2-6.1.4-cp311-cp311-macosx_11_0_arm64.whl", hash = "sha256:8156fdeb73c3ff6c8cf67ad414fb5c887cd708ff0af6d61f62629f41cb4c17b2", size = 414947, upload-time = "2026-08-01T20:40:37.405Z" },
+    { url = "https://files.pythonhosted.org/packages/8d/96/d8e1ed3e79ea20a3423a96b5c89ce794fa02cb428e4429e601f8ebcbac7c/cbor2-6.1.4-cp311-cp311-manylinux_2_28_aarch64.whl", hash = "sha256:e1fe2d62c50df290576280b18247ec63486f78be73e285bae269c2456c6ddff0", size = 457343, upload-time = "2026-08-01T20:40:38.868Z" },
+    { url = "https://files.pythonhosted.org/packages/d5/0c/5796c2ed2dcd0696fc4abedf0ea0dfd5361b3f022a311481f977fa51b2b8/cbor2-6.1.4-cp311-cp311-manylinux_2_28_x86_64.whl", hash = "sha256:c204a75f91f8cd9ed0881f6b88ec395c59aeac9fcf4d08155e7f899db2a1c46e", size = 464314, upload-time = "2026-08-01T20:40:40.63Z" },
+    { url = "https://files.pythonhosted.org/packages/b1/88/de524c6c2c91b740e5df6e6955a113fb616e979b26fd2e6a0693082d36e0/cbor2-6.1.4-cp311-cp311-musllinux_1_2_aarch64.whl", hash = "sha256:28fa5db05a7eae8fd80709959988d8a7f12838c6d4e5c58ec951414058641195", size = 523053, upload-time = "2026-08-01T20:40:42.602Z" },
+    { url = "https://files.pythonhosted.org/packages/84/07/cb5fd92834633508d680a5b5695aeaf99d33ca0bdc5b844550d538f335b0/cbor2-6.1.4-cp311-cp311-musllinux_1_2_x86_64.whl", hash = "sha256:316e217a496640418d3137483279d0e70053b000cdd4b52a4dbf20ea478bc40a", size = 532177, upload-time = "2026-08-01T20:40:44.058Z" },
+    { url = "https://files.pythonhosted.org/packages/c9/19/be98721365edfe6fc23e6bcd1385afa0e960b247c5f0b50bb67f5d05e2d9/cbor2-6.1.4-cp311-cp311-win32.whl", hash = "sha256:4903f24e0f9087275a0b6606c8b0aa586277001d51e4844fcdbc5b7211330aa8", size = 281660, upload-time = "2026-08-01T20:40:45.761Z" },
+    { url = "https://files.pythonhosted.org/packages/16/23/d54f679d4b155918f5a0879dab78203ce4fd514d311b7cfeba27dafe480b/cbor2-6.1.4-cp311-cp311-win_amd64.whl", hash = "sha256:5b99305d4013867e059f147752b95f728680682ab03d75a3f4dcfbb270d8dfe9", size = 303207, upload-time = "2026-08-01T20:40:47.293Z" },
+    { url = "https://files.pythonhosted.org/packages/53/3c/b3839d6213c88b249ba860525df05ff18b27bdc28ebc09cb1547790f001a/cbor2-6.1.4-cp311-cp311-win_arm64.whl", hash = "sha256:bd20ecc5c8ece24db952e48a91c8c47319eaa6358af707c85ac2bb388a79abc8", size = 296123, upload-time = "2026-08-01T20:40:48.808Z" },
+    { url = "https://files.pythonhosted.org/packages/2e/76/fb64293c19cafb860060310c57b768fd9cfb7cf592449660b756538cc116/cbor2-6.1.4-cp312-cp312-macosx_11_0_arm64.whl", hash = "sha256:1fc15061553e4494dc10883237501e3402c645fe509248dd698e1faf2460d68b", size = 404608, upload-time = "2026-08-01T20:40:50.219Z" },
+    { url = "https://files.pythonhosted.org/packages/96/ac/f58b3bafce7c86ada2ad8eaf189453136d2cf5bae526ea0540e1b9bc9d06/cbor2-6.1.4-cp312-cp312-manylinux_2_28_aarch64.whl", hash = "sha256:d9ada5a6ccfbb8ea7a3aa2aeb028421b52d8e0cd9323f0a2aeaa9c09d25fbce2", size = 449851, upload-time = "2026-08-01T20:40:51.725Z" },
+    { url = "https://files.pythonhosted.org/packages/f0/a5/10c6c126d59b07f2bd005094dd12a20afa46146f7e2673ed6f61a57641a7/cbor2-6.1.4-cp312-cp312-manylinux_2_28_x86_64.whl", hash = "sha256:310f3dfb296ba48fe9b63c5cf26e691e3548a1eae6901d2f0c18e941d151f220", size = 461193, upload-time = "2026-08-01T20:40:53.446Z" },
+    { url = "https://files.pythonhosted.org/packages/15/e4/4445e6237088d1cca3b8536daeb90d6b4e23776de5609c9fa46773874757/cbor2-6.1.4-cp312-cp312-musllinux_1_2_aarch64.whl", hash = "sha256:5e6c76004d674ad1c620660cb0bc5a8a0b72a5d8c7b70926d8e09e6d7e87332f", size = 516937, upload-time = "2026-08-01T20:40:54.952Z" },
+    { url = "https://files.pythonhosted.org/packages/8c/87/9c0959510f7a402e5995c81ccfd82cb9f314140dc0cce88c12836e5b93f1/cbor2-6.1.4-cp312-cp312-musllinux_1_2_x86_64.whl", hash = "sha256:32a4663425fbca4a4a7aa918eb5789d844c406439e58424cf34511f79f559242", size = 529229, upload-time = "2026-08-01T20:40:56.365Z" },
+    { url = "https://files.pythonhosted.org/packages/91/8e/6811e4ee84203ac657f6f461a37c7c9ba0287bde80eb83c7971e9b3fe156/cbor2-6.1.4-cp312-cp312-win32.whl", hash = "sha256:2310f07db3f9ba26f2a623774ff9f3dc7185af54f732ea119785a6b1bf7e1e7e", size = 278810, upload-time = "2026-08-01T20:40:57.76Z" },
+    { url = "https://files.pythonhosted.org/packages/da/27/87440788fc0d9513534c3c699238e2a9ca6010f8cb72e9c203b7af20a9f6/cbor2-6.1.4-cp312-cp312-win_amd64.whl", hash = "sha256:cc8cd300e236e9797b2e1ce306109dc481fcccf78bfa2682bf36d99e6eab1ec6", size = 299971, upload-time = "2026-08-01T20:40:59.256Z" },
+    { url = "https://files.pythonhosted.org/packages/23/f9/77981e6e63092de19d7306a09a12b0eb3fd2907dc22c10dd5d389eb27faf/cbor2-6.1.4-cp312-cp312-win_arm64.whl", hash = "sha256:553a46bda7d09552631a714e22b91e6ff2c867ecd91511596ce290d8879b8d5b", size = 290662, upload-time = "2026-08-01T20:41:00.89Z" },
+    { url = "https://files.pythonhosted.org/packages/0d/17/0b20c88e76942ede86c98cdce138681690f95908c540c264fff847729cd4/cbor2-6.1.4-cp313-cp313-macosx_11_0_arm64.whl", hash = "sha256:c48a7c938fc5fa5300ff82b5df09068dcb4838685ae8556b5ee8279d74f97ab4", size = 403677, upload-time = "2026-08-01T20:41:02.561Z" },
+    { url = "https://files.pythonhosted.org/packages/35/3d/93eed770864540c5c9ea0841008208e9db686b7335f42520705b7d6dc6b2/cbor2-6.1.4-cp313-cp313-manylinux_2_28_aarch64.whl", hash = "sha256:4bd29f21529e279d50fc14f1a811f7b05b4d8e66a7969163cce98983b6817245", size = 449762, upload-time = "2026-08-01T20:41:04.094Z" },
+    { url = "https://files.pythonhosted.org/packages/e3/21/69e4d37f00319b3d37322355aedc83154b4d8b75dc9e9789c06e1fbd8a92/cbor2-6.1.4-cp313-cp313-manylinux_2_28_x86_64.whl", hash = "sha256:36ae16d64b1f7b620c1af748e7b6947e20069ef80eee56871c5fbb84cc635905", size = 460420, upload-time = "2026-08-01T20:41:05.891Z" },
+    { url = "https://files.pythonhosted.org/packages/be/26/2cfdd5ee826205a88a826bb38b7a572c676ec3efa29574be5cdbd04b4859/cbor2-6.1.4-cp313-cp313-musllinux_1_2_aarch64.whl", hash = "sha256:69978901302ecbc8cda57b520487c5c5240ed217de783eb7728fceb258311d76", size = 516490, upload-time = "2026-08-01T20:41:07.52Z" },
+    { url = "https://files.pythonhosted.org/packages/82/86/d687cd1c2c9f9a986e8552ad1fdbd22411cc86389b5705dba6ec6f7e3226/cbor2-6.1.4-cp313-cp313-musllinux_1_2_x86_64.whl", hash = "sha256:ad4efa23fee6447e56a269191044e06eb39e809458bcd674e164fe9445feafd0", size = 528810, upload-time = "2026-08-01T20:41:09.144Z" },
+    { url = "https://files.pythonhosted.org/packages/40/08/88cecf20b8825bdd991c47b317415c08ef9e7d5f05a1def9acd346edabde/cbor2-6.1.4-cp313-cp313-win32.whl", hash = "sha256:d2560c2ba6a95904ba2a0ca257af878c4344409d9b46d8e646d8ebb617b1e0dd", size = 278058, upload-time = "2026-08-01T20:41:10.48Z" },
+    { url = "https://files.pythonhosted.org/packages/0e/67/ba140234a6415c16dcfbe0585ce12f905157b70e9cb1bb63a2b6d5721e70/cbor2-6.1.4-cp313-cp313-win_amd64.whl", hash = "sha256:c08b9c7d2ea013e24a0cb819b872b0119dde404f64a1182c0b24095b7bba781f", size = 299315, upload-time = "2026-08-01T20:41:12.067Z" },
+    { url = "https://files.pythonhosted.org/packages/5f/7f/35d53ff4252a5a85656480d3a81d5a5af823979ccd0c5cac95196a7548a6/cbor2-6.1.4-cp313-cp313-win_arm64.whl", hash = "sha256:598710183daae69cbdeb177a870ec64aa601de8138a61491fd256826d15a860f", size = 289976, upload-time = "2026-08-01T20:41:13.63Z" },
+    { url = "https://files.pythonhosted.org/packages/05/5d/c5374c76471ab41dff4420a276569a56352e83166374fba6f40fd0bde7ad/cbor2-6.1.4-cp314-cp314-macosx_11_0_arm64.whl", hash = "sha256:24da0a481294ac416e1e369e2d204b2b1d993cbd082d0d99fa3d6f5f27ae5e69", size = 407497, upload-time = "2026-08-01T20:41:15.189Z" },
+    { url = "https://files.pythonhosted.org/packages/46/f9/b9f12a5e24d5ae355e4c0f6d37330a2bbedad3331247a223a51c4cd39d5e/cbor2-6.1.4-cp314-cp314-manylinux_2_28_aarch64.whl", hash = "sha256:0859a0837e6e2d4fe5f5b849f6475797e4db545da98c19db4b1d3487bd47aa22", size = 452191, upload-time = "2026-08-01T20:41:16.705Z" },
+    { url = "https://files.pythonhosted.org/packages/67/22/8224b01f95a6fe07b1a64082aea34d9f49068392b3de93f5f3a10c73c62e/cbor2-6.1.4-cp314-cp314-manylinux_2_28_x86_64.whl", hash = "sha256:c0f5f2d6d3b58e44146860c049f3c082207a4005588b8926d51bf937ab66773c", size = 462383, upload-time = "2026-08-01T20:41:18.17Z" },
+    { url = "https://files.pythonhosted.org/packages/92/52/437e4aa4f5df1fb41020d64b3d99a8239f0f99a3a75eb6ffa5cb66004b7f/cbor2-6.1.4-cp314-cp314-musllinux_1_2_aarch64.whl", hash = "sha256:239db0f92d537fd29eaec4e40195fc3b2b48bc34a5887059658162489a9eb6ae", size = 518700, upload-time = "2026-08-01T20:41:19.592Z" },
+    { url = "https://files.pythonhosted.org/packages/7d/45/2f5ea5bfe0fd800b3739c7df8679bdffa9f7def6b2f2fee064ada1c63e85/cbor2-6.1.4-cp314-cp314-musllinux_1_2_x86_64.whl", hash = "sha256:3f4a434c36bb0d33aeb48ddae8e8b673ca7e1f14545ee7cf4a4c7c39380ea9a2", size = 531243, upload-time = "2026-08-01T20:41:21.21Z" },
+    { url = "https://files.pythonhosted.org/packages/bd/c6/0beac64cb74cd3217f295f9bb0d64675e1809c683a31ea2a49ac9d4d1504/cbor2-6.1.4-cp314-cp314-win32.whl", hash = "sha256:6abcf072b8c0fdc8ad7902ee26a906cafbf3427d026b662ff21166a253f85e18", size = 285248, upload-time = "2026-08-01T20:41:22.658Z" },
+    { url = "https://files.pythonhosted.org/packages/bb/7d/4afa096ddc94049f5a514690891b02a18319e146ceb14465ce30c8340a8b/cbor2-6.1.4-cp314-cp314-win_amd64.whl", hash = "sha256:855764e02dc60ab9413acd044e997c3170000fdea6155d6c43a923a1d966dbe6", size = 313044, upload-time = "2026-08-01T20:41:24.066Z" },
+    { url = "https://files.pythonhosted.org/packages/e5/b5/e614cee861772f6b5c4d926b066d2e7dbc11e220b50ba716ba91e430fb0f/cbor2-6.1.4-cp314-cp314-win_arm64.whl", hash = "sha256:c6b28b928c5f2dbf47dffa12dce9c8e36fe6ac1c1358bc326499c0736263b66f", size = 304088, upload-time = "2026-08-01T20:41:25.431Z" },
+    { url = "https://files.pythonhosted.org/packages/9e/41/3b28184154f6cbf7e47c1b7fb4a7a291c54f27a6f3a0a2f64b078c6a13e1/cbor2-6.1.4-cp314-cp314t-macosx_11_0_arm64.whl", hash = "sha256:7336ff4cb7d161ec43b65eef43bf3e9bcab44bd152efb54dd637b7afe711254f", size = 401042, upload-time = "2026-08-01T20:41:26.819Z" },
+    { url = "https://files.pythonhosted.org/packages/d5/1a/a8624023b84b41c43a150a89517c104aed0e467bd258866f13be4c3ac0c6/cbor2-6.1.4-cp314-cp314t-manylinux_2_28_aarch64.whl", hash = "sha256:8f1019494b0ec81a3df3ebb01b6acb446d5b946fe35845b1726379abd66a71da", size = 445301, upload-time = "2026-08-01T20:41:28.35Z" },
+    { url = "https://files.pythonhosted.org/packages/60/39/07dd0ea957c1f48673d3947f97ee36826efd4a824053dd0ec4df2f0c89d6/cbor2-6.1.4-cp314-cp314t-manylinux_2_28_x86_64.whl", hash = "sha256:179a794bf4be1d46ff190695929f65f0b42019c156919846ae539d2a7ec42e54", size = 459816, upload-time = "2026-08-01T20:41:29.839Z" },
+    { url = "https://files.pythonhosted.org/packages/23/8e/2015175132a27c1daed434f671ac6d9c1311461995df47f201307700e0da/cbor2-6.1.4-cp314-cp314t-musllinux_1_2_aarch64.whl", hash = "sha256:9b904b8d0f4ddac9259197d21d121fae4cb8b555700d65bc12c5d46a2e6c2025", size = 511565, upload-time = "2026-08-01T20:41:31.939Z" },
+    { url = "https://files.pythonhosted.org/packages/82/66/420991095d9473614b205d4c4e40b5d3b9f1ee4410eb3c48c1e902947837/cbor2-6.1.4-cp314-cp314t-musllinux_1_2_x86_64.whl", hash = "sha256:71fcf4f237d68bf4445bf45070f36f82b333f2e6a62612aa2c256683b51378a9", size = 527709, upload-time = "2026-08-01T20:41:33.413Z" },
+    { url = "https://files.pythonhosted.org/packages/cc/7c/73057e7a38488a816a0d40ff9e7cd9f418800894582e2e48fb2f47ce66a2/cbor2-6.1.4-cp314-cp314t-win32.whl", hash = "sha256:7deccc50fd0b55c4c7dd265b144c5358a645121e457c0ae3722b5ad59832b257", size = 281462, upload-time = "2026-08-01T20:41:35.127Z" },
+    { url = "https://files.pythonhosted.org/packages/99/5d/d5db22837cb566de733b9d1c418cdf1912ccb1efc7b179e295430b1d81a2/cbor2-6.1.4-cp314-cp314t-win_amd64.whl", hash = "sha256:f3fc7d15cba4174373df2496070faa4a927fe3ed772130d281808120aec7b61c", size = 309165, upload-time = "2026-08-01T20:41:36.716Z" },
+    { url = "https://files.pythonhosted.org/packages/29/5f/ff2c6da83553a692219a0a62a21b57a27ded4405200e50db758a17fbaf15/cbor2-6.1.4-cp314-cp314t-win_arm64.whl", hash = "sha256:164ca22b509408435b2d8236c80c964e4fc77c085ab034569cd04c40d5cc8883", size = 298386, upload-time = "2026-08-01T20:41:38.392Z" },
+]
+
 [[package]]
 name = "certifi"
 version = "2025.8.3"
@@ -575,7 +624,7 @@ requires-dist = [
     { name = "huggingface-hub", specifier = ">=0.20" },
     { name = "jinja2" },
     { name = "kernelguard", specifier = "==0.3.1" },
-    { name = "modal" },
+    { name = "modal", specifier = ">=1.5.0" },
     { name = "pre-commit", marker = "extra == 'dev'" },
     { name = "psycopg2-binary" },
     { name = "pyarrow", specifier = ">=14.0" },
@@ -804,15 +853,35 @@ wheels = [
 name = "grpclib"
 version = "0.4.8"
 source = { registry = "https://pypi.org/simple" }
+resolution-markers = [
+    "python_full_version == '3.13.*'",
+    "python_full_version < '3.13'",
+]
 dependencies = [
-    { name = "h2" },
-    { name = "multidict" },
+    { name = "h2", marker = "python_full_version < '3.14'" },
+    { name = "multidict", marker = "python_full_version < '3.14'" },
 ]
 sdist = { url = "https://files.pythonhosted.org/packages/19/75/0f0d3524b38b35e5cd07334b754aa9bd0570140ad982131b04ebfa3b0374/grpclib-0.4.8.tar.gz", hash = "sha256:d8823763780ef94fed8b2c562f7485cf0bbee15fc7d065a640673667f7719c9a", size = 62793, upload-time = "2025-05-04T16:27:30.051Z" }
 wheels = [
     { url = "https://files.pythonhosted.org/packages/03/8b/ad381ec1b8195fa4a9a693cb8087e031b99530c0d6b8ad036dcb99e144c4/grpclib-0.4.8-py3-none-any.whl", hash = "sha256:a5047733a7acc1c1cee6abf3c841c7c6fab67d2844a45a853b113fa2e6cd2654", size = 76311, upload-time = "2025-05-04T16:27:22.818Z" },
 ]
 
+[[package]]
+name = "grpclib"
+version = "0.4.9"
+source = { registry = "https://pypi.org/simple" }
+resolution-markers = [
+    "python_full_version >= '3.14'",
+]
+dependencies = [
+    { name = "h2", marker = "python_full_version >= '3.14'" },
+    { name = "multidict", marker = "python_full_version >= '3.14'" },
+]
+sdist = { url = "https://files.pythonhosted.org/packages/5b/28/5a2c299ec82a876a252c5919aa895a6f1d1d35c96417c5ce4a4660dc3a80/grpclib-0.4.9.tar.gz", hash = "sha256:cc589c330fa81004c6400a52a566407574498cb5b055fa927013361e21466c46", size = 84798, upload-time = "2025-12-14T22:23:14.349Z" }
+wheels = [
+    { url = "https://files.pythonhosted.org/packages/5c/90/b0cbbd9efcc82816c58f31a34963071aa19fb792a212a5d9caf8e0fc3097/grpclib-0.4.9-py3-none-any.whl", hash = "sha256:7762ec1c8ed94dfad597475152dd35cbd11aecaaca2f243e29702435ca24cf0e", size = 77063, upload-time = "2025-12-14T22:23:13.224Z" },
+]
+
 [[package]]
 name = "h11"
 version = "0.16.0"
@@ -1102,26 +1171,27 @@ wheels = [
 
 [[package]]
 name = "modal"
-version = "1.1.1"
+version = "1.5.5"
 source = { registry = "https://pypi.org/simple" }
 dependencies = [
     { name = "aiohttp" },
+    { name = "cbor2" },
     { name = "certifi" },
     { name = "click" },
-    { name = "grpclib" },
+    { name = "grpclib", version = "0.4.8", source = { registry = "https://pypi.org/simple" }, marker = "python_full_version < '3.14'" },
+    { name = "grpclib", version = "0.4.9", source = { registry = "https://pypi.org/simple" }, marker = "python_full_version >= '3.14'" },
     { name = "protobuf" },
     { name = "rich" },
     { name = "synchronicity" },
     { name = "toml" },
-    { name = "typer" },
     { name = "types-certifi" },
     { name = "types-toml" },
     { name = "typing-extensions" },
     { name = "watchfiles" },
 ]
-sdist = { url = "https://files.pythonhosted.org/packages/c2/26/68f143ce009b1f20c61edb6d8bbc18b505afd33c0f2b44830872269e46c9/modal-1.1.1.tar.gz", hash = "sha256:f353ca3b6151abec0b253c398198261859ceab79dbaa2238a58d068ce889b2c0", size = 569267, upload-time = "2025-08-01T13:32:50.337Z" }
+sdist = { url = "https://files.pythonhosted.org/packages/b9/a5/9e322043716e511b7f6c9120804f33b92e21a818bf4709b2d98b3c76d665/modal-1.5.5.tar.gz", hash = "sha256:30df363ed1898cc3d91a09ff3f95c38ab043f6b6294011b01085312c6a0ac777", size = 870356, upload-time = "2026-08-28T19:51:34.881Z" }
 wheels = [
-    { url = "https://files.pythonhosted.org/packages/e3/b1/26904e24dfde0420f12f963a7e35d50defb461edc6c9bbbbf3716789ffd4/modal-1.1.1-py3-none-any.whl", hash = "sha256:ec2e2f68848a29228b61a646d2af7ac239558c005b66ca996d7c6c5de587ae36", size = 657544, upload-time = "2025-08-01T13:32:48.037Z" },
+    { url = "https://files.pythonhosted.org/packages/04/19/b3dca8baec119126058b12ca620e44022e4d08a091956668bd04180c89a7/modal-1.5.5-py3-none-any.whl", hash = "sha256:8d10d3ee09818aaba1973b73ce2521ab8961b63a29b5b52e3ff0d25e7a74808e", size = 985163, upload-time = "2026-08-28T19:51:32.404Z" },
 ]
 
 [[package]]
@@ -1965,18 +2035,6 @@ wheels = [
     { url = "https://files.pythonhosted.org/packages/e0/f9/0595336914c5619e5f28a1fb793285925a8cd4b432c9da0a987836c7f822/shellingham-1.5.4-py2.py3-none-any.whl", hash = "sha256:7ecfff8f2fd72616f7481040475a65b2bf8af90a56c89140852d1120324e8686", size = 9755, upload-time = "2023-10-24T04:13:38.866Z" },
 ]
 
-[[package]]
-name = "sigtools"
-version = "4.0.1"
-source = { registry = "https://pypi.org/simple" }
-dependencies = [
-    { name = "attrs" },
-]
-sdist = { url = "https://files.pythonhosted.org/packages/5f/db/669ca14166814da187b3087b908ca924cf83f5b504fe23b3859a3ef67d4f/sigtools-4.0.1.tar.gz", hash = "sha256:4b8e135a9cd4d2ea00da670c093372d74e672ba3abb87f4c98d8e73dea54445c", size = 71910, upload-time = "2022-10-13T07:03:54.149Z" }
-wheels = [
-    { url = "https://files.pythonhosted.org/packages/1f/91/853dbf6ec096197dba9cd5fd0c836c5fc19142038b7db60ebe6332b1bab1/sigtools-4.0.1-py2.py3-none-any.whl", hash = "sha256:d216b4cf920bbab0fce636ddc429ed8463a5b533d9e1492acb45a2a1bc36ac6c", size = 76419, upload-time = "2022-10-13T07:03:52.658Z" },
-]
-
 [[package]]
 name = "sniffio"
 version = "1.3.1"
@@ -2010,15 +2068,14 @@ wheels = [
 
 [[package]]
 name = "synchronicity"
-version = "0.10.2"
+version = "0.12.5"
 source = { registry = "https://pypi.org/simple" }
 dependencies = [
-    { name = "sigtools" },
     { name = "typing-extensions" },
 ]
-sdist = { url = "https://files.pythonhosted.org/packages/77/b6/e977f03915cc02406bb52ac15398ea44dbde47805e5955b6bac9268acc12/synchronicity-0.10.2.tar.gz", hash = "sha256:e0dfd8a2ba4fb89c60ee53365c5fa2d2d69aabce60709055d38f736f6a592c86", size = 53891, upload-time = "2025-07-30T20:23:19.122Z" }
+sdist = { url = "https://files.pythonhosted.org/packages/5d/1c/f51dc54bbd302991026a53f9790735540e0e9e1184e9d5939f02446aa5bc/synchronicity-0.12.5.tar.gz", hash = "sha256:94d96b1d85698e3056b96a793b8c0949af6584e4a7d877fabdeb5385efe230aa", size = 60745, upload-time = "2026-06-18T21:06:23.545Z" }
 wheels = [
-    { url = "https://files.pythonhosted.org/packages/f8/f9/ce041b9531022a0b5999a47e6da14485239f7bce9c595d1bfb387fe60e89/synchronicity-0.10.2-py3-none-any.whl", hash = "sha256:4ba1f8c02ca582ef068033300201e3c403e08d81e42553554f4e67b27f0d9bb1", size = 38766, upload-time = "2025-07-30T20:23:18.04Z" },
+    { url = "https://files.pythonhosted.org/packages/06/74/ad9b99520f70c0bc3318e582e359d360cfc0f7afd7bf368a7f24013cece7/synchronicity-0.12.5-py3-none-any.whl", hash = "sha256:fdbbb10d437bc08a6b0f814fc66fddd1b58ffed314533d42f1ab555801e781af", size = 41107, upload-time = "2026-06-18T21:06:22.505Z" },
 ]
 
 [[package]]