From b16acc350fb8185286019e7535d0c8fb7fc9f9fc Mon Sep 17 00:00:00 2001 From: Mark Saroufim Date: Mon, 7 Sep 2026 12:04:17 -0700 Subject: [PATCH 1/8] Cache PyTorch inline headers on Modal volumes --- docs/benchmarks/modal-pch-2026-09-07.json | 311 +++++++++++++++++++ docs/modal-pch-cache.md | 140 +++++++++ examples/vectoradd_py/submission_cuda_pch.py | 62 ++++ pyproject.toml | 2 +- scripts/modal_pch_benchmark.py | 127 ++++++++ src/runners/modal_runner.py | 48 ++- src/runners/modal_runner_archs.py | 8 +- src/runners/pch/compiler.py | 145 +++++++++ src/runners/pch/warm.py | 36 +++ tests/test_pch_compiler.py | 111 +++++++ tests/test_pch_warmup.py | 57 ++++ uv.lock | 107 +++++-- 12 files changed, 1126 insertions(+), 28 deletions(-) create mode 100644 docs/benchmarks/modal-pch-2026-09-07.json create mode 100644 docs/modal-pch-cache.md create mode 100644 examples/vectoradd_py/submission_cuda_pch.py create mode 100644 scripts/modal_pch_benchmark.py create mode 100755 src/runners/pch/compiler.py create mode 100644 src/runners/pch/warm.py create mode 100644 tests/test_pch_compiler.py create mode 100644 tests/test_pch_warmup.py diff --git a/docs/benchmarks/modal-pch-2026-09-07.json b/docs/benchmarks/modal-pch-2026-09-07.json new file mode 100644 index 000000000..20a9169f6 --- /dev/null +++ b/docs/benchmarks/modal-pch-2026-09-07.json @@ -0,0 +1,311 @@ +{ + "date": "2026-09-07", + "base_commit": "ef6c2ae3", + "gpu": "Tesla T4", + "torch": "2.11.0+cu129", + "cuda_toolkit": "12.9.1", + "cpu": 4, + "memory_mib": 16384, + "ninja_parallelism": "default; MAX_JOBS unset", + "volume": "kernelbot-pch-test-20260907", + "app_urls": { + "minimal": "https://modal.com/apps/coreauto/main/ap-wVKHKwT4kzMrR2LFKFyacU", + "implicit": "https://modal.com/apps/coreauto/main/ap-AaFYKFJ026F8fv2YygaQQ9" + }, + "summary": { + "minimal": { + "cpp_seconds": { + "cold": 15.665, + "warm": 7.1195, + "speedup": 2.200294964534026 + }, + "load_inline_seconds": { + "cold": 16.425, + "warm": 7.930499999999999, + "speedup": 2.071117836202005 + } + }, + "implicit": { + "cpp_seconds": { + "cold": 16.018, + "warm": 6.2940000000000005, + "speedup": 2.5449634572608835 + }, + "load_inline_seconds": { + "cold": 33.899, + "warm": 32.831500000000005, + "speedup": 1.0325145058861154 + } + } + }, + "runs": { + "minimal": [ + { + "api_write_denied": true, + "enabled": false, + "gpu": "Tesla T4", + "torch": "2.11.0+cu129", + "container_id": "ta-01M1YKQ2B391GV5ZTXCAJ1QM3R", + "image_id": "im-YLQND06EX1zfw9SViNe2hr", + "load_inline_seconds": 16.812, + "cpp_seconds": [ + 15.997 + ], + "pch_consumed": false, + "keys": [], + "tests": { + "test-count": "5", + "test.0.spec": "size: 127; seed: 4242", + "test.0.status": "pass", + "test.1.spec": "size: 128; seed: 5236", + "test.1.status": "pass", + "test.2.spec": "size: 129; seed: 1001", + "test.2.status": "pass", + "test.3.spec": "size: 256; seed: 5531", + "test.3.status": "pass", + "test.4.spec": "size: 512; seed: 9173", + "test.4.status": "pass", + "check": "pass" + }, + "remote_seconds": 28.43327539, + "headers": "minimal", + "variant": 0 + }, + { + "api_write_denied": true, + "enabled": true, + "gpu": "Tesla T4", + "torch": "2.11.0+cu129", + "container_id": "ta-01M1YKR6RMMM4F8YD2VJGDDNVR", + "image_id": "im-YLQND06EX1zfw9SViNe2hr", + "load_inline_seconds": 7.629, + "cpp_seconds": [ + 6.878 + ], + "pch_consumed": true, + "keys": [ + "be683069d70b9c37dade41ea011df5640c61b0b65021da582efb3a834a420975" + ], + "tests": { + "test-count": "5", + "test.0.spec": "size: 127; seed: 4242", + "test.0.status": "pass", + "test.1.spec": "size: 128; seed: 5236", + "test.1.status": "pass", + "test.2.spec": "size: 129; seed: 1001", + "test.2.status": "pass", + "test.3.spec": "size: 256; seed: 5531", + "test.3.status": "pass", + "test.4.spec": "size: 512; seed: 9173", + "test.4.status": "pass", + "check": "pass" + }, + "remote_seconds": 18.014319609, + "headers": "minimal", + "variant": 0 + }, + { + "api_write_denied": true, + "enabled": false, + "gpu": "Tesla T4", + "torch": "2.11.0+cu129", + "container_id": "ta-01M1YKRX5H153JV4FANKNG13QR", + "image_id": "im-YLQND06EX1zfw9SViNe2hr", + "load_inline_seconds": 16.038, + "cpp_seconds": [ + 15.333 + ], + "pch_consumed": false, + "keys": [], + "tests": { + "test-count": "5", + "test.0.spec": "size: 127; seed: 4242", + "test.0.status": "pass", + "test.1.spec": "size: 128; seed: 5236", + "test.1.status": "pass", + "test.2.spec": "size: 129; seed: 1001", + "test.2.status": "pass", + "test.3.spec": "size: 256; seed: 5531", + "test.3.status": "pass", + "test.4.spec": "size: 512; seed: 9173", + "test.4.status": "pass", + "check": "pass" + }, + "remote_seconds": 26.422016425, + "headers": "minimal", + "variant": 1 + }, + { + "api_write_denied": true, + "enabled": true, + "gpu": "Tesla T4", + "torch": "2.11.0+cu129", + "container_id": "ta-01M1YKSVHQ6P6ZQ52KM71PMJ5R", + "image_id": "im-YLQND06EX1zfw9SViNe2hr", + "load_inline_seconds": 8.232, + "cpp_seconds": [ + 7.361 + ], + "pch_consumed": true, + "keys": [ + "be683069d70b9c37dade41ea011df5640c61b0b65021da582efb3a834a420975" + ], + "tests": { + "test-count": "5", + "test.0.spec": "size: 127; seed: 4242", + "test.0.status": "pass", + "test.1.spec": "size: 128; seed: 5236", + "test.1.status": "pass", + "test.2.spec": "size: 129; seed: 1001", + "test.2.status": "pass", + "test.3.spec": "size: 256; seed: 5531", + "test.3.status": "pass", + "test.4.spec": "size: 512; seed: 9173", + "test.4.status": "pass", + "check": "pass" + }, + "remote_seconds": 19.2062712, + "headers": "minimal", + "variant": 1 + } + ], + "implicit": [ + { + "api_write_denied": true, + "enabled": false, + "gpu": "Tesla T4", + "torch": "2.11.0+cu129", + "container_id": "ta-01M1YKQ9TWG5EYW7QVKN7EVSPR", + "image_id": "im-YLQND06EX1zfw9SViNe2hr", + "load_inline_seconds": 35.325, + "cpp_seconds": [ + 16.695 + ], + "pch_consumed": false, + "keys": [], + "tests": { + "test-count": "5", + "test.0.spec": "size: 127; seed: 4242", + "test.0.status": "pass", + "test.1.spec": "size: 128; seed: 5236", + "test.1.status": "pass", + "test.2.spec": "size: 129; seed: 1001", + "test.2.status": "pass", + "test.3.spec": "size: 256; seed: 5531", + "test.3.status": "pass", + "test.4.spec": "size: 512; seed: 9173", + "test.4.status": "pass", + "check": "pass" + }, + "remote_seconds": 48.656836319, + "headers": "implicit", + "variant": 0 + }, + { + "api_write_denied": true, + "enabled": true, + "gpu": "Tesla T4", + "torch": "2.11.0+cu129", + "container_id": "ta-01M1YKS008QV74K5X6KQFPJBBR", + "image_id": "im-YLQND06EX1zfw9SViNe2hr", + "load_inline_seconds": 32.884, + "cpp_seconds": [ + 6.462 + ], + "pch_consumed": true, + "keys": [ + "be683069d70b9c37dade41ea011df5640c61b0b65021da582efb3a834a420975" + ], + "tests": { + "test-count": "5", + "test.0.spec": "size: 127; seed: 4242", + "test.0.status": "pass", + "test.1.spec": "size: 128; seed: 5236", + "test.1.status": "pass", + "test.2.spec": "size: 129; seed: 1001", + "test.2.status": "pass", + "test.3.spec": "size: 256; seed: 5531", + "test.3.status": "pass", + "test.4.spec": "size: 512; seed: 9173", + "test.4.status": "pass", + "check": "pass" + }, + "remote_seconds": 43.767019222, + "headers": "implicit", + "variant": 0 + }, + { + "api_write_denied": true, + "enabled": false, + "gpu": "Tesla T4", + "torch": "2.11.0+cu129", + "container_id": "ta-01M1YKTG7PN1TYDKHMCKZ36DFR", + "image_id": "im-YLQND06EX1zfw9SViNe2hr", + "load_inline_seconds": 32.473, + "cpp_seconds": [ + 15.341 + ], + "pch_consumed": false, + "keys": [], + "tests": { + "test-count": "5", + "test.0.spec": "size: 127; seed: 4242", + "test.0.status": "pass", + "test.1.spec": "size: 128; seed: 5236", + "test.1.status": "pass", + "test.2.spec": "size: 129; seed: 1001", + "test.2.status": "pass", + "test.3.spec": "size: 256; seed: 5531", + "test.3.status": "pass", + "test.4.spec": "size: 512; seed: 9173", + "test.4.status": "pass", + "check": "pass" + }, + "remote_seconds": 42.962847054, + "headers": "implicit", + "variant": 1 + }, + { + "api_write_denied": true, + "enabled": true, + "gpu": "Tesla T4", + "torch": "2.11.0+cu129", + "container_id": "ta-01M1YKVZN7DTGCKA5V4DGMA0VR", + "image_id": "im-YLQND06EX1zfw9SViNe2hr", + "load_inline_seconds": 32.779, + "cpp_seconds": [ + 6.126 + ], + "pch_consumed": true, + "keys": [ + "be683069d70b9c37dade41ea011df5640c61b0b65021da582efb3a834a420975" + ], + "tests": { + "test-count": "5", + "test.0.spec": "size: 127; seed: 4242", + "test.0.status": "pass", + "test.1.spec": "size: 128; seed: 5236", + "test.1.status": "pass", + "test.2.spec": "size: 129; seed: 1001", + "test.2.status": "pass", + "test.3.spec": "size: 256; seed: 5531", + "test.3.status": "pass", + "test.4.spec": "size: 512; seed: 9173", + "test.4.status": "pass", + "check": "pass" + }, + "remote_seconds": 43.532378496, + "headers": "implicit", + "variant": 1 + } + ] + }, + "gpu_compute_estimate_usd": 0.04444317404926, + "gpu_estimate_note": "Eight final GPU calls only, excluding startup, exploratory runs, CPU warmup, memory and image builds. T4 rate from https://modal.com/pricing on 2026-09-07.", + "warmup": { + "profiles": 8, + "pch_bytes": 5762619418, + "app_url": "https://modal.com/apps/coreauto/main/ap-prmRQYEi7xLWq6qdgjtIuB", + "completed": true + } +} diff --git a/docs/modal-pch-cache.md b/docs/modal-pch-cache.md new file mode 100644 index 000000000..732b9b520 --- /dev/null +++ b/docs/modal-pch-cache.md @@ -0,0 +1,140 @@ +# Modal `load_inline` header cache + +Kernelbot's Modal image sets `CXX` to a small GCC wrapper. For a C++ translation +unit starting with `#include ` (the default generated by +PyTorch `load_inline`), it looks up a precompiled header in the +`kernelbot-torch-pch` Modal Volume. Existing submissions need no source changes. +Leave PyTorch's separate `use_pch` option at its default, `False`. + +The Volume contains shared headers only. Kernel objects and shared libraries +remain in each container's normal local extension cache. +This accelerates the **host C++ binding**; it does not cache NVCC's CUDA header +parsing, kernel compilation, or arbitrary user-provided headers. + +## Warm and deploy + +From the repository root, using the same Modal profile/environment for both: + +```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 +``` + +Run warmup again after changing the image, compiler, installed headers, Python, +or torch. Warmup uses a CPU container, not a GPU. It builds both CPU-only and +CUDA-linked `load_inline` profiles for these exact `extra_cflags` lists: + +- omitted / `[]` +- `["-O2"]` +- `["-O3"]` +- `["-O3", "-ffast-math"]` + +The warm function also accepts a comma-separated selection of profile names, such as +`--profiles cuda-default,cuda-O3`. Extend the trusted profiles +in `pch/warm.py` and the allowed names in `warm_pch` to support additional common +flag combinations. Arbitrary submission code and compiler arguments are never +sent to the shared-cache writer. + +`KERNELBOT_PCH_VOLUME` selects a different Volume at launch/deploy time. The +benchmark below uses this to avoid touching the production cache. A missing +profile compiles normally without writing to the Volume. No special handling is +needed for a cold deployment, though it will not receive a PCH speedup until +warmup completes and new containers start. Warm before deploying; after adding +profiles to an existing deployment, replace existing runner containers to refresh +their snapshots. The restricted GPU functions cannot call the Volume reload API. + +## Isolation and compatibility + +- GPU runners mount the Volume read-only and disable Modal API access with + `restrict_modal_access=True`. The API restriction prevents bypassing a read-only + filesystem mount with Volume upload APIs. Only the trusted CPU warmup writes it. +- Keys include the image's installed-header/compiler fingerprint, compiler + flags, include paths and header-search environment. Different module names, + kernel source filenames, object paths and dependency filenames share a key. +- ABI macros, optimization settings, C++ standards and include-path changes + select different keys. Unsupported profiles fall back to normal compilation. +- The wrapper leaves linking, compiler probes and sources without the initial + torch header alone. An explicit alternative `CXX` also bypasses it. +- GCC validates the PCH when loading it; `-Winvalid-pch` reports rejection and + GCC can parse the header normally. A cache-hit log alone is not proof of use. +- Warmup publishes a completed PCH with an atomic rename, then commits the + Volume. It does not rely on a distributed filesystem lock. + +PyTorch's native `use_pch=True` maintains a single `extension.h.gch` beside its +installed headers, and ordinary `load_inline` calls remove that native cache. +Our separate per-configuration cache is unaffected by those removals. Explicit +`use_pch=True` may still spend time building PyTorch's native PCH, so it is +unnecessary here. + +References: [PyTorch `load_inline`](https://docs.pytorch.org/docs/main/cpp_extension.html), +[Modal Volume consistency](https://modal.com/docs/guide/volumes), +[GCC precompiled headers](https://gcc.gnu.org/onlinedocs/gcc/Precompiled-Headers.html). + +## Keep CUDA source headers small + +`load_inline` normally injects `torch/types.h` into CUDA sources too. That can +make NVCC the bottleneck even after the C++ binding becomes faster. Where CUDA +code only needs device pointers, keep Tensor handling in C++ and use +`no_implicit_headers=True`. Explicitly include `torch/extension.h` first in the +C++ source so the PCH still applies, and include only the CUDA headers needed +by the device code. See the working FP16 +[vector-add submission](../examples/vectoradd_py/submission_cuda_pch.py). + +## Real submission benchmark + +```sh +PYTHONPATH=src:src/runners KERNELBOT_PCH_VOLUME=kernelbot-pch-test \ + uv run modal run scripts/modal_pch_benchmark.py + +# Also check the original example with implicit Torch headers in CUDA: +PYTHONPATH=src:src/runners KERNELBOT_PCH_VOLUME=kernelbot-pch-test \ + uv run modal run scripts/modal_pch_benchmark.py --headers implicit \ + --output /tmp/kernelbot-pch-implicit-results.json +``` + +This uses the runner image and `modal_run_config` evaluation path, with real CUDA +vector-add submissions on a T4 (4 requested CPU cores, 16 GiB memory). It warms +one header profile on CPU, then runs two different kernel variants with the PCH +disabled and enabled, each in a fresh GPU container. It uses normal Ninja +parallelism (no `MAX_JOBS=1` override). The default fixture uses minimal CUDA +headers; `--headers implicit` uses the original vector-add example. It checks +all five vector-add test cases, confirms filesystem and API writes are denied, +checks GCC's `-H` output for actual PCH consumption, and records both C++ +compilation and `load_inline` wall times. +It asserts that the warmed C++ median is faster, and additionally checks the +end-to-end median for the minimal-header fixture. NVCC can hide that improvement +in the implicit-header fixture, so its total-time speedup is reported without an +assertion. The script writes raw JSON results to +`/tmp/kernelbot-pch-results.json`. It does not deploy the production app. + +To diagnose a submission, `KERNELBOT_PCH_TRACE=1` prints the host compile time and +GCC's header trace on PCH-enabled compilations. `KERNELBOT_PCH_DISABLE=1` provides +a baseline without the cache. + +## Measured results (2026-09-07) + +T4, CUDA 12.9.1, PyTorch 2.11.0+cu129, normal Ninja parallelism. Each row below +is the median of two different kernel variants, each tested in separate cold +and PCH-enabled containers. Both fixtures share the same warmed header entry. + +| Submission | C++ compilation, cold → PCH | Total `load_inline`, cold → PCH | +| --- | --- | --- | +| Minimal CUDA headers | 15.67s → 7.12s (2.20×) | 16.43s → 7.93s (2.07×) | +| Original implicit headers | 16.02s → 6.29s (2.54×) | 33.90s → 32.83s (1.03×) | + +All 40 correctness checks passed across eight fresh Modal container IDs. GCC +confirmed actual PCH consumption on every cache-enabled build; filesystem and +Modal API writes were denied. All eight warmup profiles were built and committed +to `kernelbot-pch-test-20260907`, including successful reuse after CPU preemption. +The implementation also passed 36 local tests and the repository Ruff check. + +The original example's CUDA compilation dominates its total time, so speeding +up the parallel C++ stage alone has little effect on that workload. The +minimal-header example demonstrates the PCH benefit without that bottleneck. +The initial exploratory single-worker timings are not included in this table. + +[Raw measurements and run links](benchmarks/modal-pch-2026-09-07.json). +The eight final GPU calls used approximately $0.044 of T4 GPU compute at +[Modal's published rate](https://modal.com/pricing), excluding startup, +exploratory runs, CPU warmup, memory and image builds. The production runner +has not been deployed by this experiment. diff --git a/examples/vectoradd_py/submission_cuda_pch.py b/examples/vectoradd_py/submission_cuda_pch.py new file mode 100644 index 000000000..f18720e93 --- /dev/null +++ b/examples/vectoradd_py/submission_cuda_pch.py @@ -0,0 +1,62 @@ +#!POPCORN leaderboard vectoradd_py-dev +"""FP16 vector add: keep Torch headers in C++, and only CUDA headers in CUDA. + +Kernelbot's Modal CXX wrapper reuses the C++ torch header PCH automatically. +no_implicit_headers avoids unnecessarily parsing torch/types.h again with nvcc. +""" + +from task import input_t, output_t +from torch.utils.cpp_extension import load_inline + +add_cpp_source = r""" +#include +#include +#include + +void add_cuda_impl(const void*, const void*, void*, int, cudaStream_t); + +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 C = torch::empty_like(A); + add_cuda_impl(A.data_ptr(), B.data_ptr(), C.data_ptr(), A.numel(), + c10::cuda::getCurrentCUDAStream().stream()); + TORCH_CHECK(cudaGetLastError() == cudaSuccess, "CUDA launch failed"); + return C; +} +""" + +add_cuda_source = r""" +#include +#include + +__global__ void add_kernel(const __half* A, const __half* B, __half* C, int N) { + const int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < N) C[idx] = __hadd(A[idx], B[idx]); +} + +void add_cuda_impl(const void* A, const void* B, void* C, int N, cudaStream_t stream) { + const int threads = 256; + add_kernel<<<(N + threads - 1) / threads, threads, 0, stream>>>( + static_cast(A), static_cast(B), static_cast<__half*>(C), N); +} +""" + +add_module = load_inline( + name='add_cuda', + cpp_sources=add_cpp_source, + cuda_sources=add_cuda_source, + functions=['add_cuda'], + no_implicit_headers=True, + verbose=True, +) + + +def add(A, B): + return add_module.add_cuda(A, B) + + +def custom_kernel(data: input_t) -> output_t: + return add(*data) 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..9cce1d261 --- /dev/null +++ b/scripts/modal_pch_benchmark.py @@ -0,0 +1,127 @@ +"""Real submissions in fresh Modal containers; no production deployment required. + +PYTHONPATH=src:src/runners KERNELBOT_PCH_VOLUME=kernelbot-pch-test \\ + uv run modal run scripts/modal_pch_benchmark.py +""" + +import io +import json +import os +import re +import statistics +import time +from pathlib import Path + +import modal +from modal_runner import PCH_MOUNT, cuda_image, modal_run_config, pch_volume, warm_pch + +app = modal.App("kernelbot-pch-benchmark") +app.include(warm_pch.app) + + +@app.function( + image=cuda_image, + gpu="T4", + cpu=4, + memory=16384, + timeout=900, + single_use_containers=True, + restrict_modal_access=True, + volumes={PCH_MOUNT: pch_volume.with_mount_options(read_only=True)}, +) +def submit(config: dict, enabled: bool): + os.environ["KERNELBOT_PCH_DISABLE"] = "0" if enabled else "1" + os.environ["KERNELBOT_PCH_TRACE"] = "1" + os.environ.pop("MAX_JOBS", None) # Match normal load_inline/Ninja parallelism. + started = time.perf_counter() + result = modal_run_config(config) + assert result.success, result.error + test = result.runs["test"] + assert test.run.success and test.run.passed, test.run + assert test.compilation and test.compilation.exit_code == 0, test.compilation + logs = test.compilation.stdout + test.compilation.stderr + durations = re.findall(r"PCH_LOAD_INLINE_SECONDS=([0-9.]+)", logs) + assert durations, logs[-5000:] + consumed = "! /kernelbot-pch/" in logs + assert consumed == enabled, logs[-5000:] + # A read-only cache must remain read-only from submission containers. + try: + Path(PCH_MOUNT, "submission-write-probe").write_text("must fail") + except OSError: + pass + else: + raise AssertionError("Submission can write the shared PCH Volume") + try: + with pch_volume.batch_upload() as batch: + batch.put_file(io.BytesIO(b"must fail"), "/submission-api-write-probe") + except modal.exception.AuthError: + pass + else: + raise AssertionError("Submission can write the shared PCH Volume through the API") + return { + "api_write_denied": True, + "enabled": enabled, + "gpu": result.system.gpu, + "torch": result.system.torch, + "container_id": os.environ["MODAL_TASK_ID"], + "image_id": os.environ["MODAL_IMAGE_ID"], + "load_inline_seconds": float(durations[0]), + "cpp_seconds": [float(v) for v in re.findall(r"compile_seconds=([0-9.]+)", logs)], + "pch_consumed": consumed, + "keys": re.findall(r"\[kernelbot-pch\] hit (\w+)", logs), + "tests": test.run.result, + "remote_seconds": time.perf_counter() - started, + } + + +@app.local_entrypoint() +def main(output: str = "/tmp/kernelbot-pch-results.json", repeats: int = 2, headers: str = "minimal"): + if headers not in {"minimal", "implicit"}: + raise ValueError("headers must be minimal or implicit") + from libkernelbot.consts import SubmissionMode + from libkernelbot.task import build_task_config, make_task_definition + + print(warm_pch.remote(profiles="cuda-default")) + root = Path(__file__).resolve().parents[1] / "examples/vectoradd_py" + task = make_task_definition(root).task + filename = "submission_cuda_pch.py" if headers == "minimal" else "submission_cuda_inline.py" + source = (root / filename).read_text() + rows = [] + for index in range(repeats): + for enabled in (False, True): + variant = source.replace( + "const int threads = 256", f"const int threads = {128 if index else 256}" + ) + variant = variant.replace( + "name='add_cuda'", f"name='pch_variant_{index}_{int(enabled)}'" + ) + variant = variant.replace( + "add_module = load_inline(", + "import time\n_pch_started = time.perf_counter()\nadd_module = load_inline(", + ) + variant = variant.replace( + "\ndef add(A, B):", + '\nprint(f"PCH_LOAD_INLINE_SECONDS={time.perf_counter() - _pch_started:.3f}", flush=True)\n' + "\ndef add(A, B):", + ) + config = build_task_config( + task=task, submission_content=variant, arch="75", mode=SubmissionMode.TEST + ) + row = submit.remote(config, enabled) + row["headers"] = headers + row["variant"] = index + rows.append(row) + Path(output).write_text(json.dumps(rows, indent=2)) + print(json.dumps(row), flush=True) + assert len({row["container_id"] for row in rows}) == len(rows), "Expected fresh containers" + warm_keys = {key for row in rows if row["enabled"] for key in row["keys"]} + assert len(warm_keys) == 1, "Different kernel variants must reuse the same header entry" + cold = statistics.median(row["load_inline_seconds"] for row in rows if not row["enabled"]) + warm = statistics.median(row["load_inline_seconds"] for row in rows if row["enabled"]) + print(f"load_inline median: cold={cold:.3f}s warm={warm:.3f}s speedup={cold / warm:.2f}x") + cpp_cold = statistics.median(row["cpp_seconds"][0] for row in rows if not row["enabled"]) + cpp_warm = statistics.median(row["cpp_seconds"][0] for row in rows if row["enabled"]) + print(f"C++ median: cold={cpp_cold:.3f}s warm={cpp_warm:.3f}s speedup={cpp_cold / cpp_warm:.2f}x") + assert cpp_warm < cpp_cold, "PCH did not improve host compilation" + if headers == "minimal": + assert warm < cold, "PCH did not improve end-to-end compilation" 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*", 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 \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..1dfc6639e --- /dev/null +++ b/tests/test_pch_compiler.py @@ -0,0 +1,111 @@ +"""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 \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 = compiler.compile_flags(command(tmp_path, "second")) + 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 "] +) +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]] From 38b98928a72e24229be9e146f4897919984aadcb Mon Sep 17 00:00:00 2001 From: Mark Saroufim Date: Mon, 7 Sep 2026 13:29:17 -0700 Subject: [PATCH 2/8] Warm header profiles before the existing Modal rollout --- .github/workflows/modal-deploy.yml | 9 ++++++--- docs/modal-pch-cache.md | 4 +++- 2 files changed, 9 insertions(+), 4 deletions(-) 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/modal-pch-cache.md b/docs/modal-pch-cache.md index 732b9b520..7bd6724fd 100644 --- a/docs/modal-pch-cache.md +++ b/docs/modal-pch-cache.md @@ -13,7 +13,9 @@ parsing, kernel compilation, or arbitrary user-provided headers. ## Warm and deploy -From the repository root, using the same Modal profile/environment for both: +The existing main/dev deployment workflow warms all profiles before deploying +the GPU functions. A warmup failure prevents that rollout. For a manual rollout, +run these commands from the repository root in the same Modal profile/environment: ```sh PYTHONPATH=src:src/runners uv run modal run src/runners/modal_runner.py::warm_pch From 6c502bd1ce0d080b193e887bad37776bd99f738a Mon Sep 17 00:00:00 2001 From: Mark Saroufim Date: Mon, 7 Sep 2026 13:35:55 -0700 Subject: [PATCH 3/8] Record fresh T4 PCH benchmarks on current main --- .../modal-pch-current-main-2026-09-07.json | 302 ++++++++++++++++++ docs/modal-pch-cache.md | 29 +- 2 files changed, 330 insertions(+), 1 deletion(-) create mode 100644 docs/benchmarks/modal-pch-current-main-2026-09-07.json diff --git a/docs/benchmarks/modal-pch-current-main-2026-09-07.json b/docs/benchmarks/modal-pch-current-main-2026-09-07.json new file mode 100644 index 000000000..1d6c35373 --- /dev/null +++ b/docs/benchmarks/modal-pch-current-main-2026-09-07.json @@ -0,0 +1,302 @@ +{ + "date": "2026-09-07", + "base_commit": "727212cdcf1b9b4d587c12f6d1484b3fd54549d0", + "implementation_commit": "38b98928", + "cuda_image_version": "13.3.0", + "torch": "2.12.0+cu130", + "cpu": 4, + "memory_mib": 16384, + "normal_ninja_parallelism": true, + "volume": "kernelbot-pch-test-20260907", + "warmed_profiles": [ + "cuda-default" + ], + "run_links": { + "minimal": "https://modal.com/apps/coreauto/main/ap-iNSFXRa4TIP80OqKP6Evxs", + "implicit": "https://modal.com/apps/coreauto/main/ap-A5cGWJRnQYzTttK1ZeEbTV" + }, + "summary": { + "minimal": { + "cpp_seconds": { + "cold": 19.36, + "warm": 8.906500000000001, + "speedup": 2.173693370010666 + }, + "load_inline_seconds": { + "cold": 20.447000000000003, + "warm": 9.91, + "speedup": 2.063269424823411 + } + }, + "implicit": { + "cpp_seconds": { + "cold": 18.441, + "warm": 8.7895, + "speedup": 2.098071562659992 + }, + "load_inline_seconds": { + "cold": 38.096000000000004, + "warm": 43.245999999999995, + "speedup": 0.8809138417425891 + } + } + }, + "rows": [ + { + "api_write_denied": true, + "enabled": false, + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "container_id": "ta-01M1YS0N16K76Q8V4C3VF51N3R", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "load_inline_seconds": 19.445, + "cpp_seconds": [ + 18.61 + ], + "pch_consumed": false, + "keys": [], + "tests": { + "test-count": "5", + "test.0.spec": "size: 127; seed: 4242", + "test.0.status": "pass", + "test.1.spec": "size: 128; seed: 5236", + "test.1.status": "pass", + "test.2.spec": "size: 129; seed: 1001", + "test.2.status": "pass", + "test.3.spec": "size: 256; seed: 5531", + "test.3.status": "pass", + "test.4.spec": "size: 512; seed: 9173", + "test.4.status": "pass", + "check": "pass" + }, + "remote_seconds": 31.217458912999998, + "headers": "minimal", + "variant": 0 + }, + { + "api_write_denied": true, + "enabled": true, + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "container_id": "ta-01M1YS1TYTV9BKMDNRHENGDD9R", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "load_inline_seconds": 9.23, + "cpp_seconds": [ + 8.278 + ], + "pch_consumed": true, + "keys": [ + "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" + ], + "tests": { + "test-count": "5", + "test.0.spec": "size: 127; seed: 4242", + "test.0.status": "pass", + "test.1.spec": "size: 128; seed: 5236", + "test.1.status": "pass", + "test.2.spec": "size: 129; seed: 1001", + "test.2.status": "pass", + "test.3.spec": "size: 256; seed: 5531", + "test.3.status": "pass", + "test.4.spec": "size: 512; seed: 9173", + "test.4.status": "pass", + "check": "pass" + }, + "remote_seconds": 21.60613694, + "headers": "minimal", + "variant": 0 + }, + { + "api_write_denied": true, + "enabled": false, + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "container_id": "ta-01M1YS2KZ7R00G1ANY0ZQMS1TR", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "load_inline_seconds": 21.449, + "cpp_seconds": [ + 20.11 + ], + "pch_consumed": false, + "keys": [], + "tests": { + "test-count": "5", + "test.0.spec": "size: 127; seed: 4242", + "test.0.status": "pass", + "test.1.spec": "size: 128; seed: 5236", + "test.1.status": "pass", + "test.2.spec": "size: 129; seed: 1001", + "test.2.status": "pass", + "test.3.spec": "size: 256; seed: 5531", + "test.3.status": "pass", + "test.4.spec": "size: 512; seed: 9173", + "test.4.status": "pass", + "check": "pass" + }, + "remote_seconds": 33.334485347, + "headers": "minimal", + "variant": 1 + }, + { + "api_write_denied": true, + "enabled": true, + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "container_id": "ta-01M1YS3T71BGNZHNW70VDM0ZTR", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "load_inline_seconds": 10.59, + "cpp_seconds": [ + 9.535 + ], + "pch_consumed": true, + "keys": [ + "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" + ], + "tests": { + "test-count": "5", + "test.0.spec": "size: 127; seed: 4242", + "test.0.status": "pass", + "test.1.spec": "size: 128; seed: 5236", + "test.1.status": "pass", + "test.2.spec": "size: 129; seed: 1001", + "test.2.status": "pass", + "test.3.spec": "size: 256; seed: 5531", + "test.3.status": "pass", + "test.4.spec": "size: 512; seed: 9173", + "test.4.status": "pass", + "check": "pass" + }, + "remote_seconds": 26.458702556, + "headers": "minimal", + "variant": 1 + }, + { + "api_write_denied": true, + "enabled": false, + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "container_id": "ta-01M1YS2KSKB7VPRNKQXTVN7S2R", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "load_inline_seconds": 38.816, + "cpp_seconds": [ + 19.127 + ], + "pch_consumed": false, + "keys": [], + "tests": { + "test-count": "5", + "test.0.spec": "size: 127; seed: 4242", + "test.0.status": "pass", + "test.1.spec": "size: 128; seed: 5236", + "test.1.status": "pass", + "test.2.spec": "size: 129; seed: 1001", + "test.2.status": "pass", + "test.3.spec": "size: 256; seed: 5531", + "test.3.status": "pass", + "test.4.spec": "size: 512; seed: 9173", + "test.4.status": "pass", + "check": "pass" + }, + "remote_seconds": 51.25590044, + "headers": "implicit", + "variant": 0 + }, + { + "api_write_denied": true, + "enabled": true, + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "container_id": "ta-01M1YS4C0R2T6HGV7RYZYNKZVR", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "load_inline_seconds": 49.988, + "cpp_seconds": [ + 10.423 + ], + "pch_consumed": true, + "keys": [ + "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" + ], + "tests": { + "test-count": "5", + "test.0.spec": "size: 127; seed: 4242", + "test.0.status": "pass", + "test.1.spec": "size: 128; seed: 5236", + "test.1.status": "pass", + "test.2.spec": "size: 129; seed: 1001", + "test.2.status": "pass", + "test.3.spec": "size: 256; seed: 5531", + "test.3.status": "pass", + "test.4.spec": "size: 512; seed: 9173", + "test.4.status": "pass", + "check": "pass" + }, + "remote_seconds": 67.888768213, + "headers": "implicit", + "variant": 0 + }, + { + "api_write_denied": true, + "enabled": false, + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "container_id": "ta-01M1YS6W4ESSF84B2K39QJEMZR", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "load_inline_seconds": 37.376, + "cpp_seconds": [ + 17.755 + ], + "pch_consumed": false, + "keys": [], + "tests": { + "test-count": "5", + "test.0.spec": "size: 127; seed: 4242", + "test.0.status": "pass", + "test.1.spec": "size: 128; seed: 5236", + "test.1.status": "pass", + "test.2.spec": "size: 129; seed: 1001", + "test.2.status": "pass", + "test.3.spec": "size: 256; seed: 5531", + "test.3.status": "pass", + "test.4.spec": "size: 512; seed: 9173", + "test.4.status": "pass", + "check": "pass" + }, + "remote_seconds": 48.367798831, + "headers": "implicit", + "variant": 1 + }, + { + "api_write_denied": true, + "enabled": true, + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "container_id": "ta-01M1YS8GAY01Q5Y1XPHQ0E9DRR", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "load_inline_seconds": 36.504, + "cpp_seconds": [ + 7.156 + ], + "pch_consumed": true, + "keys": [ + "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" + ], + "tests": { + "test-count": "5", + "test.0.spec": "size: 127; seed: 4242", + "test.0.status": "pass", + "test.1.spec": "size: 128; seed: 5236", + "test.1.status": "pass", + "test.2.spec": "size: 129; seed: 1001", + "test.2.status": "pass", + "test.3.spec": "size: 256; seed: 5531", + "test.3.status": "pass", + "test.4.spec": "size: 512; seed: 9173", + "test.4.status": "pass", + "check": "pass" + }, + "remote_seconds": 47.203748833999995, + "headers": "implicit", + "variant": 1 + } + ] +} diff --git a/docs/modal-pch-cache.md b/docs/modal-pch-cache.md index 7bd6724fd..12c3c9a34 100644 --- a/docs/modal-pch-cache.md +++ b/docs/modal-pch-cache.md @@ -113,7 +113,34 @@ To diagnose a submission, `KERNELBOT_PCH_TRACE=1` prints the host compile time a GCC's header trace on PCH-enabled compilations. `KERNELBOT_PCH_DISABLE=1` provides a baseline without the cache. -## Measured results (2026-09-07) +## Current-main measurements (2026-09-07) + +Revalidated on main `727212cd`, preserving CUDA 13.3.0, PyTorch 2.12.0+cu130, +and MathDx. Both fixtures use the same image and freshly warmed `cuda-default` +profile, normal Ninja parallelism, and eight distinct T4 containers. + +| Submission | C++ compilation, cold → PCH | Total `load_inline`, cold → PCH | +| --- | --- | --- | +| Minimal CUDA headers | 19.36s → 8.91s (2.17×) | 20.45s → 9.91s (2.06×) | +| Original implicit headers | 18.44s → 8.79s (2.10×) | 38.10s → 43.25s (0.88×) | + +All 40 correctness checks passed. Every cache-enabled build consumed the same +PCH, and every container rejected filesystem and Modal API writes. The +Torch-heavy CUDA fixture did **not** improve end-to-end in this run: warm totals +were 49.99s and 36.50s, despite faster host compilation. These two-sample medians +show the host-stage benefit and the minimal-header fixture's total-time benefit; +they do not establish an end-to-end speedup for NVCC-heavy submissions. + +41 relevant local tests and Ruff passed on this base. GitHub unit, lint, Modal +integration, and GitHub integration checks also passed on `38b98928`. +Only `cuda-default` was rewarmed on the current image; the eight-profile warmup +below applies to the earlier image. The deployment workflow warms all profiles +for the actual image before deploying it. Production has not been deployed by +this experiment. + +[Current-main raw measurements and run links](benchmarks/modal-pch-current-main-2026-09-07.json). + +## Earlier-image measurements (2026-09-07) T4, CUDA 12.9.1, PyTorch 2.11.0+cu129, normal Ninja parallelism. Each row below is the median of two different kernel variants, each tested in separate cold From 3f9cd13115bc3447c537a25c984f73c565ae2406 Mon Sep 17 00:00:00 2001 From: Mark Saroufim Date: Mon, 7 Sep 2026 13:52:38 -0700 Subject: [PATCH 4/8] Test PCH hits, misses, changed headers and non-inline builds on Modal --- .../modal-pch-matrix-2026-09-07.json | 678 ++++++++++++++++++ docs/modal-pch-cache.md | 89 +++ scripts/modal_pch_matrix.py | 260 +++++++ tests/test_pch_compiler.py | 8 +- 4 files changed, 1034 insertions(+), 1 deletion(-) create mode 100644 docs/benchmarks/modal-pch-matrix-2026-09-07.json create mode 100644 scripts/modal_pch_matrix.py diff --git a/docs/benchmarks/modal-pch-matrix-2026-09-07.json b/docs/benchmarks/modal-pch-matrix-2026-09-07.json new file mode 100644 index 000000000..3a65d36c0 --- /dev/null +++ b/docs/benchmarks/modal-pch-matrix-2026-09-07.json @@ -0,0 +1,678 @@ +{ + "date": "2026-09-07", + "implementation_commit": "6c502bd1ce0d080b193e887bad37776bd99f738a", + "harness_sha256": "cac105d7621941eacd6fa0885f0acff0bddec067cfb29a1197331dde2f16c659", + "cuda": "13.3.0", + "torch": "2.12.0+cu130", + "gpu": "T4", + "requested_cpu": 4, + "memory_mib": 16384, + "volume": "kernelbot-pch-test-20260907", + "warmed_profile": "cuda-default", + "normal_ninja_parallelism": true, + "baseline_compiler": "/usr/bin/g++", + "diagnostic_header_tracing_excluded_from_timing": true, + "timing_notes": "Both modes use the same external timing shim. Total load/load_inline is timed for CUDA cases; plain_cpp reports the median of ten object compilations, excluding its link. Per-case summaries are medians of two different source/header variants in separate containers. No zero-slowdown guarantee is asserted.", + "run_links": { + "inline_simple": "https://modal.com/apps/coreauto/main/ap-ZBKWqPp5sSkMJCFjKmWwPI", + "remaining": "https://modal.com/apps/coreauto/main/ap-5aIt3temjLqpwFSfQhW6LI" + }, + "correctness_checks": 100, + "fresh_containers": 20, + "pch_keys": [ + "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" + ], + "estimated_gpu_compute_usd": 0.09262219526251199, + "cost_notes": "T4 GPU time for these 20 calls at $0.000164/s, including untimed diagnostics; excludes startup, setup/failing pilot runs, CPU warmup and memory.", + "summary": { + "inline_simple": { + "baseline_seconds": 18.8471031105, + "cache_seconds": 9.997744876999999, + "ratio": 1.8851354322771445, + "baseline_host_seconds": 18.042343731499997, + "cache_host_seconds": 8.803183696000001 + }, + "inline_complex": { + "baseline_seconds": 19.938248450499998, + "cache_seconds": 14.2401688755, + "ratio": 1.400141292200787, + "baseline_host_seconds": 19.111699750499998, + "cache_host_seconds": 7.875871039 + }, + "load_files": { + "baseline_seconds": 18.3017816695, + "cache_seconds": 8.774257883, + "ratio": 2.085849528648963, + "baseline_host_seconds": 17.528128097, + "cache_host_seconds": 7.7537195005 + }, + "plain_cpp": { + "baseline_seconds": 0.2017184177500002, + "cache_seconds": 0.28704759275, + "ratio": 0.7027350963562476, + "baseline_host_seconds": 0.1250452337500002, + "cache_host_seconds": 0.20968030274999983 + }, + "flag_miss": { + "baseline_seconds": 18.0761359535, + "cache_seconds": 20.301430385, + "ratio": 0.8903873082192184, + "baseline_host_seconds": 17.252728550999997, + "cache_host_seconds": 19.076173736499996 + } + }, + "rows": [ + { + "case": "inline_simple", + "variant": 0, + "enabled": false, + "extension_name": "pch_matrix_inline_simple_0_0", + "seconds": 18.800584544, + "timing_samples": [ + 18.800584544 + ], + "host_compile_seconds": 17.997382098, + "pch_consumed": false, + "hit_keys": [], + "cache_miss": false, + "correctness_checks": 6, + "outputs": [ + 2.0, + 2.0, + 2.0, + 2.0, + 2.0 + ], + "source_sha256": "4ec65ca0b08feea11ce48b7a0ebd322676aab622d35eba1da63b605643e46435", + "container_id": "ta-01M1YT654VFX478EV36GSNAGCR", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 1, + "remote_seconds": 41.135301596999994 + }, + { + "case": "inline_simple", + "variant": 0, + "enabled": true, + "extension_name": "pch_matrix_inline_simple_0_1", + "seconds": 9.449116172, + "timing_samples": [ + 9.449116172 + ], + "host_compile_seconds": 8.415279449, + "pch_consumed": true, + "hit_keys": [ + "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" + ], + "cache_miss": false, + "correctness_checks": 6, + "outputs": [ + 2.0, + 2.0, + 2.0, + 2.0, + 2.0 + ], + "source_sha256": "4ec65ca0b08feea11ce48b7a0ebd322676aab622d35eba1da63b605643e46435", + "container_id": "ta-01M1YT65B8E9149HBHXBTA0NAR", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 1, + "remote_seconds": 21.075640688 + }, + { + "case": "inline_simple", + "variant": 1, + "enabled": true, + "extension_name": "pch_matrix_inline_simple_1_1", + "seconds": 10.546373582, + "timing_samples": [ + 10.546373582 + ], + "host_compile_seconds": 9.191087943000001, + "pch_consumed": true, + "hit_keys": [ + "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" + ], + "cache_miss": false, + "correctness_checks": 6, + "outputs": [ + 3.0, + 3.0, + 3.0, + 3.0, + 3.0 + ], + "source_sha256": "7e064d69277762a33d2a65a4d6258e5422c7d2b3535694ce9afd6a6304f6993a", + "container_id": "ta-01M1YT654V66S3D05ZP1KSGBXR", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 1, + "remote_seconds": 22.933361906 + }, + { + "case": "inline_simple", + "variant": 1, + "enabled": false, + "extension_name": "pch_matrix_inline_simple_1_0", + "seconds": 18.893621677, + "timing_samples": [ + 18.893621677 + ], + "host_compile_seconds": 18.087305365, + "pch_consumed": false, + "hit_keys": [], + "cache_miss": false, + "correctness_checks": 6, + "outputs": [ + 3.0, + 3.0, + 3.0, + 3.0, + 3.0 + ], + "source_sha256": "7e064d69277762a33d2a65a4d6258e5422c7d2b3535694ce9afd6a6304f6993a", + "container_id": "ta-01M1YT65B8TXRZGSM820Z3Q16R", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 1, + "remote_seconds": 40.762046516 + }, + { + "case": "inline_complex", + "variant": 0, + "enabled": false, + "extension_name": "pch_matrix_inline_complex_0_0", + "seconds": 18.759751232, + "timing_samples": [ + 18.759751232 + ], + "host_compile_seconds": 18.012592808999997, + "pch_consumed": false, + "hit_keys": [], + "cache_miss": false, + "correctness_checks": 6, + "outputs": [ + 2.0, + 2.0, + 2.0, + 2.0, + 2.0 + ], + "source_sha256": "bb3eedfa1df83397af6a8f7edb81db46b0db7cda4a388c79a270329aa8ae85ee", + "container_id": "ta-01M1YT0HRZP289WPV9VMSSPF8R", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 1, + "remote_seconds": 40.230402710999996 + }, + { + "case": "inline_complex", + "variant": 0, + "enabled": true, + "extension_name": "pch_matrix_inline_complex_0_1", + "seconds": 14.449278815, + "timing_samples": [ + 14.449278815 + ], + "host_compile_seconds": 7.923061731999999, + "pch_consumed": true, + "hit_keys": [ + "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" + ], + "cache_miss": false, + "correctness_checks": 6, + "outputs": [ + 2.0, + 2.0, + 2.0, + 2.0, + 2.0 + ], + "source_sha256": "bb3eedfa1df83397af6a8f7edb81db46b0db7cda4a388c79a270329aa8ae85ee", + "container_id": "ta-01M1YT0J0EQEC8X2QHC6S0EAPR", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 1, + "remote_seconds": 26.335685509 + }, + { + "case": "inline_complex", + "variant": 1, + "enabled": true, + "extension_name": "pch_matrix_inline_complex_1_1", + "seconds": 14.031058936, + "timing_samples": [ + 14.031058936 + ], + "host_compile_seconds": 7.8286803460000005, + "pch_consumed": true, + "hit_keys": [ + "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" + ], + "cache_miss": false, + "correctness_checks": 6, + "outputs": [ + 3.0, + 3.0, + 3.0, + 3.0, + 3.0 + ], + "source_sha256": "8be4b20baeb5c3ea8fbb6ec714175f9d77dd02662be31f23bddffc7807074d66", + "container_id": "ta-01M1YT0HRZYPV863GAZA3K6T1R", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 1, + "remote_seconds": 26.304385472 + }, + { + "case": "inline_complex", + "variant": 1, + "enabled": false, + "extension_name": "pch_matrix_inline_complex_1_0", + "seconds": 21.116745669, + "timing_samples": [ + 21.116745669 + ], + "host_compile_seconds": 20.210806692, + "pch_consumed": false, + "hit_keys": [], + "cache_miss": false, + "correctness_checks": 6, + "outputs": [ + 3.0, + 3.0, + 3.0, + 3.0, + 3.0 + ], + "source_sha256": "8be4b20baeb5c3ea8fbb6ec714175f9d77dd02662be31f23bddffc7807074d66", + "container_id": "ta-01M1YT0HM8DWCT713GAJ6ZQRVR", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 1, + "remote_seconds": 45.54159261 + }, + { + "case": "load_files", + "variant": 0, + "enabled": false, + "extension_name": "pch_matrix_shared", + "seconds": 18.484804069, + "timing_samples": [ + 18.484804069 + ], + "host_compile_seconds": 17.716654833999996, + "pch_consumed": false, + "hit_keys": [], + "cache_miss": false, + "correctness_checks": 6, + "outputs": [ + 2.0, + 2.0, + 2.0, + 2.0, + 2.0 + ], + "source_sha256": "4ec65ca0b08feea11ce48b7a0ebd322676aab622d35eba1da63b605643e46435", + "container_id": "ta-01M1YT1MEHZQJTJ91K03EYFKJR", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 1, + "remote_seconds": 39.96083478 + }, + { + "case": "load_files", + "variant": 0, + "enabled": true, + "extension_name": "pch_matrix_shared", + "seconds": 8.844153404, + "timing_samples": [ + 8.844153404 + ], + "host_compile_seconds": 7.813371175, + "pch_consumed": true, + "hit_keys": [ + "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" + ], + "cache_miss": false, + "correctness_checks": 6, + "outputs": [ + 2.0, + 2.0, + 2.0, + 2.0, + 2.0 + ], + "source_sha256": "4ec65ca0b08feea11ce48b7a0ebd322676aab622d35eba1da63b605643e46435", + "container_id": "ta-01M1YT1MM8C9EBBGEZCWXZAW5R", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 1, + "remote_seconds": 20.400708645 + }, + { + "case": "load_files", + "variant": 1, + "enabled": true, + "extension_name": "pch_matrix_shared", + "seconds": 8.704362362000001, + "timing_samples": [ + 8.704362362000001 + ], + "host_compile_seconds": 7.6940678259999995, + "pch_consumed": true, + "hit_keys": [ + "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" + ], + "cache_miss": false, + "correctness_checks": 6, + "outputs": [ + 3.0, + 3.0, + 3.0, + 3.0, + 3.0 + ], + "source_sha256": "7e064d69277762a33d2a65a4d6258e5422c7d2b3535694ce9afd6a6304f6993a", + "container_id": "ta-01M1YT20M1E9G09T9PZBWNQ65R", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 1, + "remote_seconds": 20.25399071 + }, + { + "case": "load_files", + "variant": 1, + "enabled": false, + "extension_name": "pch_matrix_shared", + "seconds": 18.118759269999998, + "timing_samples": [ + 18.118759269999998 + ], + "host_compile_seconds": 17.33960136, + "pch_consumed": false, + "hit_keys": [], + "cache_miss": false, + "correctness_checks": 6, + "outputs": [ + 3.0, + 3.0, + 3.0, + 3.0, + 3.0 + ], + "source_sha256": "7e064d69277762a33d2a65a4d6258e5422c7d2b3535694ce9afd6a6304f6993a", + "container_id": "ta-01M1YT276G72J91NSPNG9VWAWR", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 1, + "remote_seconds": 39.779365328000004 + }, + { + "case": "plain_cpp", + "variant": 0, + "enabled": false, + "extension_name": "pch_matrix_plain_cpp_0_0", + "seconds": 0.19980638800000028, + "timing_samples": [ + 0.21276275, + 0.19661020400000018, + 0.20195385499999974, + 0.20159454799999965, + 0.20105536800000046, + 0.1982173860000005, + 0.1969059470000003, + 0.1985574080000001, + 0.19711016499999978, + 0.2051067179999997 + ], + "host_compile_seconds": 0.12402935050000008, + "pch_consumed": false, + "hit_keys": [], + "cache_miss": false, + "correctness_checks": 1, + "outputs": [ + 3 + ], + "source_sha256": "f5fbb559817c22a648c4be4067de69dd806e6e9f2c9be71aa84f2371ecf8d781", + "container_id": "ta-01M1YT2FP2X908DV2VFHCJPT6R", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 0, + "remote_seconds": 2.3064990610000002 + }, + { + "case": "plain_cpp", + "variant": 0, + "enabled": true, + "extension_name": "pch_matrix_plain_cpp_0_1", + "seconds": 0.2823680209999999, + "timing_samples": [ + 0.35445732099999994, + 0.3628376659999999, + 0.30029828899999966, + 0.28355820099999995, + 0.28399713199999965, + 0.27632356199999997, + 0.2811778409999999, + 0.2746912210000003, + 0.2782820089999998, + 0.2755435680000007 + ], + "host_compile_seconds": 0.20570012650000002, + "pch_consumed": false, + "hit_keys": [], + "cache_miss": false, + "correctness_checks": 1, + "outputs": [ + 3 + ], + "source_sha256": "f5fbb559817c22a648c4be4067de69dd806e6e9f2c9be71aa84f2371ecf8d781", + "container_id": "ta-01M1YT2QSWZR9SN6VGQJC16VSR", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 0, + "remote_seconds": 3.291071738 + }, + { + "case": "plain_cpp", + "variant": 1, + "enabled": true, + "extension_name": "pch_matrix_plain_cpp_1_1", + "seconds": 0.29172716450000014, + "timing_samples": [ + 0.49523322599999986, + 0.36589996000000014, + 0.28990115899999935, + 0.291525472, + 0.28335249200000057, + 0.29192885700000026, + 0.2972645429999998, + 0.2952913079999995, + 0.2844164109999996, + 0.2789291899999995 + ], + "host_compile_seconds": 0.21366047899999963, + "pch_consumed": false, + "hit_keys": [], + "cache_miss": false, + "correctness_checks": 1, + "outputs": [ + 4 + ], + "source_sha256": "4f5b6c5b34f3f71e1e94ad89fb640e348e9bd0232d5cdfc75075c5e827674378", + "container_id": "ta-01M1YT2TY6A6XZNT3DJDM9AW6R", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 0, + "remote_seconds": 3.5053143249999996 + }, + { + "case": "plain_cpp", + "variant": 1, + "enabled": false, + "extension_name": "pch_matrix_plain_cpp_1_0", + "seconds": 0.20363044750000014, + "timing_samples": [ + 0.2254486870000001, + 0.20830646500000016, + 0.20089933599999998, + 0.207017467, + 0.2131324689999996, + 0.20014710299999994, + 0.20089258399999999, + 0.2038948989999998, + 0.19941997500000053, + 0.2033659960000005 + ], + "host_compile_seconds": 0.12606111700000033, + "pch_consumed": false, + "hit_keys": [], + "cache_miss": false, + "correctness_checks": 1, + "outputs": [ + 4 + ], + "source_sha256": "4f5b6c5b34f3f71e1e94ad89fb640e348e9bd0232d5cdfc75075c5e827674378", + "container_id": "ta-01M1YT3294D1MHQ774CC40NP6R", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 0, + "remote_seconds": 2.300308057 + }, + { + "case": "flag_miss", + "variant": 0, + "enabled": false, + "extension_name": "pch_matrix_flag_miss_0_0", + "seconds": 18.314668328000003, + "timing_samples": [ + 18.314668328000003 + ], + "host_compile_seconds": 17.468810722999997, + "pch_consumed": false, + "hit_keys": [], + "cache_miss": false, + "correctness_checks": 6, + "outputs": [ + 2.0, + 2.0, + 2.0, + 2.0, + 2.0 + ], + "source_sha256": "4ec65ca0b08feea11ce48b7a0ebd322676aab622d35eba1da63b605643e46435", + "container_id": "ta-01M1YT32ZPHKD6NAA1RQHDFZNR", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 1, + "remote_seconds": 39.548441550999996 + }, + { + "case": "flag_miss", + "variant": 0, + "enabled": true, + "extension_name": "pch_matrix_flag_miss_0_1", + "seconds": 20.427356331, + "timing_samples": [ + 20.427356331 + ], + "host_compile_seconds": 19.146695692999998, + "pch_consumed": false, + "hit_keys": [], + "cache_miss": true, + "correctness_checks": 6, + "outputs": [ + 2.0, + 2.0, + 2.0, + 2.0, + 2.0 + ], + "source_sha256": "4ec65ca0b08feea11ce48b7a0ebd322676aab622d35eba1da63b605643e46435", + "container_id": "ta-01M1YT36X83KGN6Y6BSXPHJB5R", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 1, + "remote_seconds": 44.605277338 + }, + { + "case": "flag_miss", + "variant": 1, + "enabled": true, + "extension_name": "pch_matrix_flag_miss_1_1", + "seconds": 20.175504439, + "timing_samples": [ + 20.175504439 + ], + "host_compile_seconds": 19.005651779999997, + "pch_consumed": false, + "hit_keys": [], + "cache_miss": true, + "correctness_checks": 6, + "outputs": [ + 3.0, + 3.0, + 3.0, + 3.0, + 3.0 + ], + "source_sha256": "7e064d69277762a33d2a65a4d6258e5422c7d2b3535694ce9afd6a6304f6993a", + "container_id": "ta-01M1YT3AQERWZZTC7SBEKE1EFR", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 1, + "remote_seconds": 43.945379948 + }, + { + "case": "flag_miss", + "variant": 1, + "enabled": false, + "extension_name": "pch_matrix_flag_miss_1_0", + "seconds": 17.837603579, + "timing_samples": [ + 17.837603579 + ], + "host_compile_seconds": 17.036646379, + "pch_consumed": false, + "hit_keys": [], + "cache_miss": false, + "correctness_checks": 6, + "outputs": [ + 3.0, + 3.0, + 3.0, + 3.0, + 3.0 + ], + "source_sha256": "7e064d69277762a33d2a65a4d6258e5422c7d2b3535694ce9afd6a6304f6993a", + "container_id": "ta-01M1YT3MBVRB90BKSZW7DR6FHR", + "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", + "gpu": "Tesla T4", + "torch": "2.12.0+cu130", + "nvcc_translation_units": 1, + "remote_seconds": 40.553874818 + } + ] +} diff --git a/docs/modal-pch-cache.md b/docs/modal-pch-cache.md index 12c3c9a34..9ad4e5e24 100644 --- a/docs/modal-pch-cache.md +++ b/docs/modal-pch-cache.md @@ -113,6 +113,95 @@ To diagnose a submission, `KERNELBOT_PCH_TRACE=1` prints the host compile time a GCC's header trace on PCH-enabled compilations. `KERNELBOT_PCH_DISABLE=1` provides a baseline without the cache. +## Similar sources, other build APIs, and performance limits + +Reuse is based on the Torch header's compilation configuration, not source-text +similarity. Different kernel bodies, filenames, extension names, or user headers +included **after** `torch/extension.h` can share the same Torch PCH. Their source +and user headers are still compiled separately into new objects/shared libraries; +none of that submission code is persisted in the Volume. Changing a local header +therefore changes the resulting code, even when the Torch PCH key stays the same. + +Changing host compiler flags, ABI macros, include paths, or header-search +environment selects another key. If that exact profile was not warmed, the +wrapper compiles normally. A macro or other header before the initial Torch +include makes the wrapper bypass PCH entirely, preserving preprocessing order. +Installed header changes require rebuilding the image/fingerprint and rewarming. + +The wrapper operates at `CXX`, so it can also accelerate +`torch.utils.cpp_extension.load()` with file sources beginning with the same Torch +include and matching flags. Plain C++ without that prefix bypasses PCH. Builds +that explicitly choose another compiler bypass the wrapper. Python/Triton paths +that do not invoke `CXX` never invoke this cache. + +There is **no guarantee of zero slowdown**. Even a miss/bypass adds a Python +compiler-wrapper invocation; hits must load a large PCH from the Volume, and +NVCC-heavy builds may receive little or no end-to-end benefit. Warmup itself costs +CPU time and storage, although it runs before deployment rather than on a +submission's GPU. `KERNELBOT_PCH_DISABLE=1` skips PCH lookup, but still starts the +wrapper; `CXX=/usr/bin/g++` bypasses the wrapper entirely. + +The broader regression matrix runs simple CUDA headers, CuTe/MathDx headers, +file-based `load()`, a plain `` C++ program, and unwarmed compiler flags: + +```sh +PYTHONPATH=src:src/runners KERNELBOT_PCH_VOLUME=kernelbot-pch-test-20260907 \ + uv run modal run scripts/modal_pch_matrix.py +``` + +Each case changes a local header (`kVariant = 1` versus `2`); CUDA cases also +change the kernel body (256 versus 128 threads), producing different expected +outputs. The file-based `load()` case even keeps the extension name identical +across variants. Each variant +runs with direct GCC and with the cache wrapper in a fresh T4 container. The reference +compiler is real `/usr/bin/g++`, rather than the cache wrapper in disabled mode. +Both modes have identical external timing instrumentation. GCC `-H` verification +runs in a separate compile **after** timing and correctness checks. Plain C++ +uses ten compiles per container to make wrapper overhead visible. No timing +threshold forces a favorable result; the matrix reports all timings and asserts +correctness, the intended hit/miss behavior, actual PCH consumption, unique +container IDs, changed source hashes, and distinct outputs for the variants. +The complex fixture enables CUDA half operators/conversions via NVCC `-U` flags +needed by MathDx; these device-only flags do not change the host PCH key. + +## Broader matrix results (2026-09-07) + +CUDA 13.3 / PyTorch 2.12.0+cu130, T4, normal Ninja parallelism. Each table value +is the median across two changed source/header variants, each tested in separate +direct-GCC and cache-enabled containers. Plain C++ first takes the median of ten +object compilations per container; its linking is excluded. + +| Case | Direct GCC | Cache wrapper | Ratio | +| --- | ---: | ---: | ---: | +| Minimal CUDA headers, `load_inline` | 18.847s | 9.998s | 1.89× | +| CuTe/MathDx headers, `load_inline` | 19.938s | 14.240s | 1.40× | +| File-based `load()`, same extension name | 18.302s | 8.774s | 2.09× | +| Plain C++ `` object compilation | 0.202s | 0.287s | 0.70× | +| Unwarmed host flags, `load_inline` | 18.076s | 20.301s | 0.89× | + +All **100 correctness checks passed across 20 fresh containers**. The six warmed +CUDA builds (simple, complex, and file-based, two variants each) consumed the +same Torch PCH. The CUDA outputs changed from 2 to 3 as expected when the local +header changed; the file-based test used the same extension name for both. +Both changed-flag builds logged cache misses, and plain C++ bypassed PCH. + +The plain C++ compiler stage, excluding the common timing shim's startup, went +from about 0.125s to 0.210s: roughly **85ms wrapper overhead**. The cache-miss +case was also slower in this sample (18.08s to 20.30s); with two source variants +on different machines, this does not isolate wrapper/filesystem cost from +machine variability. These results explicitly do not establish zero regressions. +For a tiny build that must avoid wrapper cost, set `CXX=/usr/bin/g++`. + +CuTe/MathDx headers still compile through NVCC; the PCH only saves Torch host-header +work. The complex case initially failed because PyTorch's NVCC half-type +suppression macros conflict with MathDx. The final fixture undoes those macros +only for NVCC, then passes in both baseline and cached modes. The successful +matrix runs below use that corrected fixture, with no changes to production +compiler-cache code. + +The harness and updated key regression test pass Ruff and the 41 relevant local +tests. [Raw results, per-container timings, and successful Modal run links](benchmarks/modal-pch-matrix-2026-09-07.json). + ## Current-main measurements (2026-09-07) Revalidated on main `727212cd`, preserving CUDA 13.3.0, PyTorch 2.12.0+cu130, diff --git a/scripts/modal_pch_matrix.py b/scripts/modal_pch_matrix.py new file mode 100644 index 000000000..c5ae72342 --- /dev/null +++ b/scripts/modal_pch_matrix.py @@ -0,0 +1,260 @@ +"""Broader PCH regression matrix: real compilation and distinct GPU outputs. + +PYTHONPATH=src:src/runners KERNELBOT_PCH_VOLUME=kernelbot-pch-test-20260907 \ + uv run modal run scripts/modal_pch_matrix.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", "inline_complex", "load_files", "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() +code = subprocess.call([os.environ["PCH_TEST_COMPILER"], *sys.argv[1:]]) +if "-c" in sys.argv: + with open(os.environ["PCH_TEST_COMMANDS"], "a") as f: + f.write(json.dumps({"args": sys.argv[1:], "seconds": time.perf_counter()-started}) + "\n") +sys.exit(code) +''' + +CPP = r''' +#include +#include +#include +#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 +#include +#include "variant.h" +// EXTRA_HEADERS +__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(a), static_cast(b), static_cast<__half*>(out), n); +} +''' + + +@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): + import hashlib + import os + import re + import shlex + 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) + 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 == "inline_complex": + cpp = cpp.replace( + '#include "variant.h"', + '#include "variant.h"\n#include \n#include ', + ) + cuda = cuda.replace( + "// EXTRA_HEADERS", + "#include \n#include \n" + "static_assert(cute::Int<128>::value == 128);", + ) + name = "pch_matrix_shared" if case == "load_files" else f"pch_matrix_{case}_{variant}_{int(enabled)}" + if case == "plain_cpp": + cpp = '#include \n#include \n#include "variant.h"\nint main() {\n' + cpp += 'std::array 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 == "inline_complex": + # MathDx needs half operators/conversions disabled by PyTorch's NVCC defaults. + # Only device flags change; the host C++ PCH remains compatible. + 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", + ] +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=True, + extra_cflags=[f"-DPCH_CASE={variant+1}"] if case == "flag_miss" else [], **options) +seconds = time.perf_counter() - started +assert module.version() == variant + 1 +outputs = [] +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))) +''' + 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 = 6 + times = [wall] + observed = [json.loads(line) for line in commands.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 {"inline_simple", "inline_complex", "load_files"} + 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:] + ninja = root / "build.ninja" + nvcc_commands = [] + if ninja.exists(): + rendered = subprocess.check_output(["ninja", "-C", str(root), "-t", "commands"], text=True) + nvcc_commands = [shlex.split(line) for line in rendered.splitlines() if "/nvcc " in line] + return { + "case": case, "variant": variant, "enabled": enabled, "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_translation_units": len(nvcc_commands), + "remote_seconds": time.perf_counter() - started, + } + + +@app.local_entrypoint() +def main(output: str = "/tmp/kernelbot-pch-matrix.json", cases: str = ""): + 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) + 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) & {"inline_simple", "inline_complex", "load_files"} 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/tests/test_pch_compiler.py b/tests/test_pch_compiler.py index 1dfc6639e..09ab44227 100644 --- a/tests/test_pch_compiler.py +++ b/tests/test_pch_compiler.py @@ -35,7 +35,13 @@ def command(tmp_path, name="first", extra=()): def test_different_kernels_share_header_key(tmp_path): first = compiler.compile_flags(command(tmp_path, "first")) - second = compiler.compile_flags(command(tmp_path, "second")) + 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 \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") From 3b8fcb844bcb0737201725a0d2148340ad1e7850 Mon Sep 17 00:00:00 2001 From: Mark Saroufim Date: Mon, 7 Sep 2026 13:53:15 -0700 Subject: [PATCH 5/8] Explain automatic PCH usage for submission authors --- docs/modal-pch-cache.md | 33 +++++++++++++++++++++++++++++++++ 1 file changed, 33 insertions(+) diff --git a/docs/modal-pch-cache.md b/docs/modal-pch-cache.md index 9ad4e5e24..c631bffb8 100644 --- a/docs/modal-pch-cache.md +++ b/docs/modal-pch-cache.md @@ -11,6 +11,39 @@ remain in each container's normal local extension cache. This accelerates the **host C++ binding**; it does not cache NVCC's CUDA header parsing, kernel compilation, or arbitrary user-provided headers. +## For submission authors: automatic after deployment + +Keep using `load_inline` normally; no cache flag is required: + +```python +module = load_inline( + name="my_kernel", + cpp_sources=cpp_source, + cuda_sources=cuda_source, + functions=["my_kernel"], +) +``` + +Kernelbot configures `CXX` in the runner image. The normal implicit +`torch/extension.h` prefix and default host compiler flags match the warmed cache. +Do **not** add `use_pch=True` to enable this feature: that invokes PyTorch's +separate native PCH mechanism and can add redundant work. Leave it at its default. + +The standard warmup also covers exact `extra_cflags=["-O2"]`, `["-O3"]`, and +`["-O3", "-ffast-math"]` profiles. Other host flags or include environments may +miss, but still compile normally. `extra_cuda_cflags` apply to NVCC, so changing +only those does not change the host C++ cache key. No submission author needs to +create, mount, reload, or write a Modal Volume; the operator owns warmup/deployment. + +For `no_implicit_headers=True`, explicitly put `#include ` +first in the C++ source if you want this cache to apply. The same condition +allows file-based `load()` to benefit. Pure Python/Triton submissions do not use +this C++ cache. To bypass all wrapper overhead for a tiny or unusual C++ build, +set `os.environ["CXX"] = "/usr/bin/g++"` before calling the extension loader. + +This feature remains a proposed runner change until PR #527 is deployed; the +reported tests ran in separate Modal test apps. + ## Warm and deploy The existing main/dev deployment workflow warms all profiles before deploying From 8232215024c7f90786433c99ad907159abf0dfc4 Mon Sep 17 00:00:00 2001 From: Mark Saroufim Date: Mon, 7 Sep 2026 14:04:50 -0700 Subject: [PATCH 6/8] Benchmark instantiated CUTLASS GEMMs and trim PCH experiment artifacts --- docs/benchmarks/modal-pch-2026-09-07.json | 311 -------- .../modal-pch-current-main-2026-09-07.json | 302 -------- .../modal-pch-matrix-2026-09-07.json | 678 ------------------ docs/benchmarks/modal-pch.json | 52 ++ docs/modal-pch-cache.md | 339 ++------- examples/vectoradd_py/submission_cuda_pch.py | 62 -- scripts/modal_pch_benchmark.py | 413 ++++++++--- scripts/modal_pch_matrix.py | 260 ------- 8 files changed, 438 insertions(+), 1979 deletions(-) delete mode 100644 docs/benchmarks/modal-pch-2026-09-07.json delete mode 100644 docs/benchmarks/modal-pch-current-main-2026-09-07.json delete mode 100644 docs/benchmarks/modal-pch-matrix-2026-09-07.json create mode 100644 docs/benchmarks/modal-pch.json delete mode 100644 examples/vectoradd_py/submission_cuda_pch.py delete mode 100644 scripts/modal_pch_matrix.py diff --git a/docs/benchmarks/modal-pch-2026-09-07.json b/docs/benchmarks/modal-pch-2026-09-07.json deleted file mode 100644 index 20a9169f6..000000000 --- a/docs/benchmarks/modal-pch-2026-09-07.json +++ /dev/null @@ -1,311 +0,0 @@ -{ - "date": "2026-09-07", - "base_commit": "ef6c2ae3", - "gpu": "Tesla T4", - "torch": "2.11.0+cu129", - "cuda_toolkit": "12.9.1", - "cpu": 4, - "memory_mib": 16384, - "ninja_parallelism": "default; MAX_JOBS unset", - "volume": "kernelbot-pch-test-20260907", - "app_urls": { - "minimal": "https://modal.com/apps/coreauto/main/ap-wVKHKwT4kzMrR2LFKFyacU", - "implicit": "https://modal.com/apps/coreauto/main/ap-AaFYKFJ026F8fv2YygaQQ9" - }, - "summary": { - "minimal": { - "cpp_seconds": { - "cold": 15.665, - "warm": 7.1195, - "speedup": 2.200294964534026 - }, - "load_inline_seconds": { - "cold": 16.425, - "warm": 7.930499999999999, - "speedup": 2.071117836202005 - } - }, - "implicit": { - "cpp_seconds": { - "cold": 16.018, - "warm": 6.2940000000000005, - "speedup": 2.5449634572608835 - }, - "load_inline_seconds": { - "cold": 33.899, - "warm": 32.831500000000005, - "speedup": 1.0325145058861154 - } - } - }, - "runs": { - "minimal": [ - { - "api_write_denied": true, - "enabled": false, - "gpu": "Tesla T4", - "torch": "2.11.0+cu129", - "container_id": "ta-01M1YKQ2B391GV5ZTXCAJ1QM3R", - "image_id": "im-YLQND06EX1zfw9SViNe2hr", - "load_inline_seconds": 16.812, - "cpp_seconds": [ - 15.997 - ], - "pch_consumed": false, - "keys": [], - "tests": { - "test-count": "5", - "test.0.spec": "size: 127; seed: 4242", - "test.0.status": "pass", - "test.1.spec": "size: 128; seed: 5236", - "test.1.status": "pass", - "test.2.spec": "size: 129; seed: 1001", - "test.2.status": "pass", - "test.3.spec": "size: 256; seed: 5531", - "test.3.status": "pass", - "test.4.spec": "size: 512; seed: 9173", - "test.4.status": "pass", - "check": "pass" - }, - "remote_seconds": 28.43327539, - "headers": "minimal", - "variant": 0 - }, - { - "api_write_denied": true, - "enabled": true, - "gpu": "Tesla T4", - "torch": "2.11.0+cu129", - "container_id": "ta-01M1YKR6RMMM4F8YD2VJGDDNVR", - "image_id": "im-YLQND06EX1zfw9SViNe2hr", - "load_inline_seconds": 7.629, - "cpp_seconds": [ - 6.878 - ], - "pch_consumed": true, - "keys": [ - "be683069d70b9c37dade41ea011df5640c61b0b65021da582efb3a834a420975" - ], - "tests": { - "test-count": "5", - "test.0.spec": "size: 127; seed: 4242", - "test.0.status": "pass", - "test.1.spec": "size: 128; seed: 5236", - "test.1.status": "pass", - "test.2.spec": "size: 129; seed: 1001", - "test.2.status": "pass", - "test.3.spec": "size: 256; seed: 5531", - "test.3.status": "pass", - "test.4.spec": "size: 512; seed: 9173", - "test.4.status": "pass", - "check": "pass" - }, - "remote_seconds": 18.014319609, - "headers": "minimal", - "variant": 0 - }, - { - "api_write_denied": true, - "enabled": false, - "gpu": "Tesla T4", - "torch": "2.11.0+cu129", - "container_id": "ta-01M1YKRX5H153JV4FANKNG13QR", - "image_id": "im-YLQND06EX1zfw9SViNe2hr", - "load_inline_seconds": 16.038, - "cpp_seconds": [ - 15.333 - ], - "pch_consumed": false, - "keys": [], - "tests": { - "test-count": "5", - "test.0.spec": "size: 127; seed: 4242", - "test.0.status": "pass", - "test.1.spec": "size: 128; seed: 5236", - "test.1.status": "pass", - "test.2.spec": "size: 129; seed: 1001", - "test.2.status": "pass", - "test.3.spec": "size: 256; seed: 5531", - "test.3.status": "pass", - "test.4.spec": "size: 512; seed: 9173", - "test.4.status": "pass", - "check": "pass" - }, - "remote_seconds": 26.422016425, - "headers": "minimal", - "variant": 1 - }, - { - "api_write_denied": true, - "enabled": true, - "gpu": "Tesla T4", - "torch": "2.11.0+cu129", - "container_id": "ta-01M1YKSVHQ6P6ZQ52KM71PMJ5R", - "image_id": "im-YLQND06EX1zfw9SViNe2hr", - "load_inline_seconds": 8.232, - "cpp_seconds": [ - 7.361 - ], - "pch_consumed": true, - "keys": [ - "be683069d70b9c37dade41ea011df5640c61b0b65021da582efb3a834a420975" - ], - "tests": { - "test-count": "5", - "test.0.spec": "size: 127; seed: 4242", - "test.0.status": "pass", - "test.1.spec": "size: 128; seed: 5236", - "test.1.status": "pass", - "test.2.spec": "size: 129; seed: 1001", - "test.2.status": "pass", - "test.3.spec": "size: 256; seed: 5531", - "test.3.status": "pass", - "test.4.spec": "size: 512; seed: 9173", - "test.4.status": "pass", - "check": "pass" - }, - "remote_seconds": 19.2062712, - "headers": "minimal", - "variant": 1 - } - ], - "implicit": [ - { - "api_write_denied": true, - "enabled": false, - "gpu": "Tesla T4", - "torch": "2.11.0+cu129", - "container_id": "ta-01M1YKQ9TWG5EYW7QVKN7EVSPR", - "image_id": "im-YLQND06EX1zfw9SViNe2hr", - "load_inline_seconds": 35.325, - "cpp_seconds": [ - 16.695 - ], - "pch_consumed": false, - "keys": [], - "tests": { - "test-count": "5", - "test.0.spec": "size: 127; seed: 4242", - "test.0.status": "pass", - "test.1.spec": "size: 128; seed: 5236", - "test.1.status": "pass", - "test.2.spec": "size: 129; seed: 1001", - "test.2.status": "pass", - "test.3.spec": "size: 256; seed: 5531", - "test.3.status": "pass", - "test.4.spec": "size: 512; seed: 9173", - "test.4.status": "pass", - "check": "pass" - }, - "remote_seconds": 48.656836319, - "headers": "implicit", - "variant": 0 - }, - { - "api_write_denied": true, - "enabled": true, - "gpu": "Tesla T4", - "torch": "2.11.0+cu129", - "container_id": "ta-01M1YKS008QV74K5X6KQFPJBBR", - "image_id": "im-YLQND06EX1zfw9SViNe2hr", - "load_inline_seconds": 32.884, - "cpp_seconds": [ - 6.462 - ], - "pch_consumed": true, - "keys": [ - "be683069d70b9c37dade41ea011df5640c61b0b65021da582efb3a834a420975" - ], - "tests": { - "test-count": "5", - "test.0.spec": "size: 127; seed: 4242", - "test.0.status": "pass", - "test.1.spec": "size: 128; seed: 5236", - "test.1.status": "pass", - "test.2.spec": "size: 129; seed: 1001", - "test.2.status": "pass", - "test.3.spec": "size: 256; seed: 5531", - "test.3.status": "pass", - "test.4.spec": "size: 512; seed: 9173", - "test.4.status": "pass", - "check": "pass" - }, - "remote_seconds": 43.767019222, - "headers": "implicit", - "variant": 0 - }, - { - "api_write_denied": true, - "enabled": false, - "gpu": "Tesla T4", - "torch": "2.11.0+cu129", - "container_id": "ta-01M1YKTG7PN1TYDKHMCKZ36DFR", - "image_id": "im-YLQND06EX1zfw9SViNe2hr", - "load_inline_seconds": 32.473, - "cpp_seconds": [ - 15.341 - ], - "pch_consumed": false, - "keys": [], - "tests": { - "test-count": "5", - "test.0.spec": "size: 127; seed: 4242", - "test.0.status": "pass", - "test.1.spec": "size: 128; seed: 5236", - "test.1.status": "pass", - "test.2.spec": "size: 129; seed: 1001", - "test.2.status": "pass", - "test.3.spec": "size: 256; seed: 5531", - "test.3.status": "pass", - "test.4.spec": "size: 512; seed: 9173", - "test.4.status": "pass", - "check": "pass" - }, - "remote_seconds": 42.962847054, - "headers": "implicit", - "variant": 1 - }, - { - "api_write_denied": true, - "enabled": true, - "gpu": "Tesla T4", - "torch": "2.11.0+cu129", - "container_id": "ta-01M1YKVZN7DTGCKA5V4DGMA0VR", - "image_id": "im-YLQND06EX1zfw9SViNe2hr", - "load_inline_seconds": 32.779, - "cpp_seconds": [ - 6.126 - ], - "pch_consumed": true, - "keys": [ - "be683069d70b9c37dade41ea011df5640c61b0b65021da582efb3a834a420975" - ], - "tests": { - "test-count": "5", - "test.0.spec": "size: 127; seed: 4242", - "test.0.status": "pass", - "test.1.spec": "size: 128; seed: 5236", - "test.1.status": "pass", - "test.2.spec": "size: 129; seed: 1001", - "test.2.status": "pass", - "test.3.spec": "size: 256; seed: 5531", - "test.3.status": "pass", - "test.4.spec": "size: 512; seed: 9173", - "test.4.status": "pass", - "check": "pass" - }, - "remote_seconds": 43.532378496, - "headers": "implicit", - "variant": 1 - } - ] - }, - "gpu_compute_estimate_usd": 0.04444317404926, - "gpu_estimate_note": "Eight final GPU calls only, excluding startup, exploratory runs, CPU warmup, memory and image builds. T4 rate from https://modal.com/pricing on 2026-09-07.", - "warmup": { - "profiles": 8, - "pch_bytes": 5762619418, - "app_url": "https://modal.com/apps/coreauto/main/ap-prmRQYEi7xLWq6qdgjtIuB", - "completed": true - } -} diff --git a/docs/benchmarks/modal-pch-current-main-2026-09-07.json b/docs/benchmarks/modal-pch-current-main-2026-09-07.json deleted file mode 100644 index 1d6c35373..000000000 --- a/docs/benchmarks/modal-pch-current-main-2026-09-07.json +++ /dev/null @@ -1,302 +0,0 @@ -{ - "date": "2026-09-07", - "base_commit": "727212cdcf1b9b4d587c12f6d1484b3fd54549d0", - "implementation_commit": "38b98928", - "cuda_image_version": "13.3.0", - "torch": "2.12.0+cu130", - "cpu": 4, - "memory_mib": 16384, - "normal_ninja_parallelism": true, - "volume": "kernelbot-pch-test-20260907", - "warmed_profiles": [ - "cuda-default" - ], - "run_links": { - "minimal": "https://modal.com/apps/coreauto/main/ap-iNSFXRa4TIP80OqKP6Evxs", - "implicit": "https://modal.com/apps/coreauto/main/ap-A5cGWJRnQYzTttK1ZeEbTV" - }, - "summary": { - "minimal": { - "cpp_seconds": { - "cold": 19.36, - "warm": 8.906500000000001, - "speedup": 2.173693370010666 - }, - "load_inline_seconds": { - "cold": 20.447000000000003, - "warm": 9.91, - "speedup": 2.063269424823411 - } - }, - "implicit": { - "cpp_seconds": { - "cold": 18.441, - "warm": 8.7895, - "speedup": 2.098071562659992 - }, - "load_inline_seconds": { - "cold": 38.096000000000004, - "warm": 43.245999999999995, - "speedup": 0.8809138417425891 - } - } - }, - "rows": [ - { - "api_write_denied": true, - "enabled": false, - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "container_id": "ta-01M1YS0N16K76Q8V4C3VF51N3R", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "load_inline_seconds": 19.445, - "cpp_seconds": [ - 18.61 - ], - "pch_consumed": false, - "keys": [], - "tests": { - "test-count": "5", - "test.0.spec": "size: 127; seed: 4242", - "test.0.status": "pass", - "test.1.spec": "size: 128; seed: 5236", - "test.1.status": "pass", - "test.2.spec": "size: 129; seed: 1001", - "test.2.status": "pass", - "test.3.spec": "size: 256; seed: 5531", - "test.3.status": "pass", - "test.4.spec": "size: 512; seed: 9173", - "test.4.status": "pass", - "check": "pass" - }, - "remote_seconds": 31.217458912999998, - "headers": "minimal", - "variant": 0 - }, - { - "api_write_denied": true, - "enabled": true, - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "container_id": "ta-01M1YS1TYTV9BKMDNRHENGDD9R", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "load_inline_seconds": 9.23, - "cpp_seconds": [ - 8.278 - ], - "pch_consumed": true, - "keys": [ - "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" - ], - "tests": { - "test-count": "5", - "test.0.spec": "size: 127; seed: 4242", - "test.0.status": "pass", - "test.1.spec": "size: 128; seed: 5236", - "test.1.status": "pass", - "test.2.spec": "size: 129; seed: 1001", - "test.2.status": "pass", - "test.3.spec": "size: 256; seed: 5531", - "test.3.status": "pass", - "test.4.spec": "size: 512; seed: 9173", - "test.4.status": "pass", - "check": "pass" - }, - "remote_seconds": 21.60613694, - "headers": "minimal", - "variant": 0 - }, - { - "api_write_denied": true, - "enabled": false, - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "container_id": "ta-01M1YS2KZ7R00G1ANY0ZQMS1TR", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "load_inline_seconds": 21.449, - "cpp_seconds": [ - 20.11 - ], - "pch_consumed": false, - "keys": [], - "tests": { - "test-count": "5", - "test.0.spec": "size: 127; seed: 4242", - "test.0.status": "pass", - "test.1.spec": "size: 128; seed: 5236", - "test.1.status": "pass", - "test.2.spec": "size: 129; seed: 1001", - "test.2.status": "pass", - "test.3.spec": "size: 256; seed: 5531", - "test.3.status": "pass", - "test.4.spec": "size: 512; seed: 9173", - "test.4.status": "pass", - "check": "pass" - }, - "remote_seconds": 33.334485347, - "headers": "minimal", - "variant": 1 - }, - { - "api_write_denied": true, - "enabled": true, - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "container_id": "ta-01M1YS3T71BGNZHNW70VDM0ZTR", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "load_inline_seconds": 10.59, - "cpp_seconds": [ - 9.535 - ], - "pch_consumed": true, - "keys": [ - "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" - ], - "tests": { - "test-count": "5", - "test.0.spec": "size: 127; seed: 4242", - "test.0.status": "pass", - "test.1.spec": "size: 128; seed: 5236", - "test.1.status": "pass", - "test.2.spec": "size: 129; seed: 1001", - "test.2.status": "pass", - "test.3.spec": "size: 256; seed: 5531", - "test.3.status": "pass", - "test.4.spec": "size: 512; seed: 9173", - "test.4.status": "pass", - "check": "pass" - }, - "remote_seconds": 26.458702556, - "headers": "minimal", - "variant": 1 - }, - { - "api_write_denied": true, - "enabled": false, - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "container_id": "ta-01M1YS2KSKB7VPRNKQXTVN7S2R", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "load_inline_seconds": 38.816, - "cpp_seconds": [ - 19.127 - ], - "pch_consumed": false, - "keys": [], - "tests": { - "test-count": "5", - "test.0.spec": "size: 127; seed: 4242", - "test.0.status": "pass", - "test.1.spec": "size: 128; seed: 5236", - "test.1.status": "pass", - "test.2.spec": "size: 129; seed: 1001", - "test.2.status": "pass", - "test.3.spec": "size: 256; seed: 5531", - "test.3.status": "pass", - "test.4.spec": "size: 512; seed: 9173", - "test.4.status": "pass", - "check": "pass" - }, - "remote_seconds": 51.25590044, - "headers": "implicit", - "variant": 0 - }, - { - "api_write_denied": true, - "enabled": true, - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "container_id": "ta-01M1YS4C0R2T6HGV7RYZYNKZVR", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "load_inline_seconds": 49.988, - "cpp_seconds": [ - 10.423 - ], - "pch_consumed": true, - "keys": [ - "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" - ], - "tests": { - "test-count": "5", - "test.0.spec": "size: 127; seed: 4242", - "test.0.status": "pass", - "test.1.spec": "size: 128; seed: 5236", - "test.1.status": "pass", - "test.2.spec": "size: 129; seed: 1001", - "test.2.status": "pass", - "test.3.spec": "size: 256; seed: 5531", - "test.3.status": "pass", - "test.4.spec": "size: 512; seed: 9173", - "test.4.status": "pass", - "check": "pass" - }, - "remote_seconds": 67.888768213, - "headers": "implicit", - "variant": 0 - }, - { - "api_write_denied": true, - "enabled": false, - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "container_id": "ta-01M1YS6W4ESSF84B2K39QJEMZR", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "load_inline_seconds": 37.376, - "cpp_seconds": [ - 17.755 - ], - "pch_consumed": false, - "keys": [], - "tests": { - "test-count": "5", - "test.0.spec": "size: 127; seed: 4242", - "test.0.status": "pass", - "test.1.spec": "size: 128; seed: 5236", - "test.1.status": "pass", - "test.2.spec": "size: 129; seed: 1001", - "test.2.status": "pass", - "test.3.spec": "size: 256; seed: 5531", - "test.3.status": "pass", - "test.4.spec": "size: 512; seed: 9173", - "test.4.status": "pass", - "check": "pass" - }, - "remote_seconds": 48.367798831, - "headers": "implicit", - "variant": 1 - }, - { - "api_write_denied": true, - "enabled": true, - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "container_id": "ta-01M1YS8GAY01Q5Y1XPHQ0E9DRR", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "load_inline_seconds": 36.504, - "cpp_seconds": [ - 7.156 - ], - "pch_consumed": true, - "keys": [ - "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" - ], - "tests": { - "test-count": "5", - "test.0.spec": "size: 127; seed: 4242", - "test.0.status": "pass", - "test.1.spec": "size: 128; seed: 5236", - "test.1.status": "pass", - "test.2.spec": "size: 129; seed: 1001", - "test.2.status": "pass", - "test.3.spec": "size: 256; seed: 5531", - "test.3.status": "pass", - "test.4.spec": "size: 512; seed: 9173", - "test.4.status": "pass", - "check": "pass" - }, - "remote_seconds": 47.203748833999995, - "headers": "implicit", - "variant": 1 - } - ] -} diff --git a/docs/benchmarks/modal-pch-matrix-2026-09-07.json b/docs/benchmarks/modal-pch-matrix-2026-09-07.json deleted file mode 100644 index 3a65d36c0..000000000 --- a/docs/benchmarks/modal-pch-matrix-2026-09-07.json +++ /dev/null @@ -1,678 +0,0 @@ -{ - "date": "2026-09-07", - "implementation_commit": "6c502bd1ce0d080b193e887bad37776bd99f738a", - "harness_sha256": "cac105d7621941eacd6fa0885f0acff0bddec067cfb29a1197331dde2f16c659", - "cuda": "13.3.0", - "torch": "2.12.0+cu130", - "gpu": "T4", - "requested_cpu": 4, - "memory_mib": 16384, - "volume": "kernelbot-pch-test-20260907", - "warmed_profile": "cuda-default", - "normal_ninja_parallelism": true, - "baseline_compiler": "/usr/bin/g++", - "diagnostic_header_tracing_excluded_from_timing": true, - "timing_notes": "Both modes use the same external timing shim. Total load/load_inline is timed for CUDA cases; plain_cpp reports the median of ten object compilations, excluding its link. Per-case summaries are medians of two different source/header variants in separate containers. No zero-slowdown guarantee is asserted.", - "run_links": { - "inline_simple": "https://modal.com/apps/coreauto/main/ap-ZBKWqPp5sSkMJCFjKmWwPI", - "remaining": "https://modal.com/apps/coreauto/main/ap-5aIt3temjLqpwFSfQhW6LI" - }, - "correctness_checks": 100, - "fresh_containers": 20, - "pch_keys": [ - "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" - ], - "estimated_gpu_compute_usd": 0.09262219526251199, - "cost_notes": "T4 GPU time for these 20 calls at $0.000164/s, including untimed diagnostics; excludes startup, setup/failing pilot runs, CPU warmup and memory.", - "summary": { - "inline_simple": { - "baseline_seconds": 18.8471031105, - "cache_seconds": 9.997744876999999, - "ratio": 1.8851354322771445, - "baseline_host_seconds": 18.042343731499997, - "cache_host_seconds": 8.803183696000001 - }, - "inline_complex": { - "baseline_seconds": 19.938248450499998, - "cache_seconds": 14.2401688755, - "ratio": 1.400141292200787, - "baseline_host_seconds": 19.111699750499998, - "cache_host_seconds": 7.875871039 - }, - "load_files": { - "baseline_seconds": 18.3017816695, - "cache_seconds": 8.774257883, - "ratio": 2.085849528648963, - "baseline_host_seconds": 17.528128097, - "cache_host_seconds": 7.7537195005 - }, - "plain_cpp": { - "baseline_seconds": 0.2017184177500002, - "cache_seconds": 0.28704759275, - "ratio": 0.7027350963562476, - "baseline_host_seconds": 0.1250452337500002, - "cache_host_seconds": 0.20968030274999983 - }, - "flag_miss": { - "baseline_seconds": 18.0761359535, - "cache_seconds": 20.301430385, - "ratio": 0.8903873082192184, - "baseline_host_seconds": 17.252728550999997, - "cache_host_seconds": 19.076173736499996 - } - }, - "rows": [ - { - "case": "inline_simple", - "variant": 0, - "enabled": false, - "extension_name": "pch_matrix_inline_simple_0_0", - "seconds": 18.800584544, - "timing_samples": [ - 18.800584544 - ], - "host_compile_seconds": 17.997382098, - "pch_consumed": false, - "hit_keys": [], - "cache_miss": false, - "correctness_checks": 6, - "outputs": [ - 2.0, - 2.0, - 2.0, - 2.0, - 2.0 - ], - "source_sha256": "4ec65ca0b08feea11ce48b7a0ebd322676aab622d35eba1da63b605643e46435", - "container_id": "ta-01M1YT654VFX478EV36GSNAGCR", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 1, - "remote_seconds": 41.135301596999994 - }, - { - "case": "inline_simple", - "variant": 0, - "enabled": true, - "extension_name": "pch_matrix_inline_simple_0_1", - "seconds": 9.449116172, - "timing_samples": [ - 9.449116172 - ], - "host_compile_seconds": 8.415279449, - "pch_consumed": true, - "hit_keys": [ - "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" - ], - "cache_miss": false, - "correctness_checks": 6, - "outputs": [ - 2.0, - 2.0, - 2.0, - 2.0, - 2.0 - ], - "source_sha256": "4ec65ca0b08feea11ce48b7a0ebd322676aab622d35eba1da63b605643e46435", - "container_id": "ta-01M1YT65B8E9149HBHXBTA0NAR", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 1, - "remote_seconds": 21.075640688 - }, - { - "case": "inline_simple", - "variant": 1, - "enabled": true, - "extension_name": "pch_matrix_inline_simple_1_1", - "seconds": 10.546373582, - "timing_samples": [ - 10.546373582 - ], - "host_compile_seconds": 9.191087943000001, - "pch_consumed": true, - "hit_keys": [ - "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" - ], - "cache_miss": false, - "correctness_checks": 6, - "outputs": [ - 3.0, - 3.0, - 3.0, - 3.0, - 3.0 - ], - "source_sha256": "7e064d69277762a33d2a65a4d6258e5422c7d2b3535694ce9afd6a6304f6993a", - "container_id": "ta-01M1YT654V66S3D05ZP1KSGBXR", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 1, - "remote_seconds": 22.933361906 - }, - { - "case": "inline_simple", - "variant": 1, - "enabled": false, - "extension_name": "pch_matrix_inline_simple_1_0", - "seconds": 18.893621677, - "timing_samples": [ - 18.893621677 - ], - "host_compile_seconds": 18.087305365, - "pch_consumed": false, - "hit_keys": [], - "cache_miss": false, - "correctness_checks": 6, - "outputs": [ - 3.0, - 3.0, - 3.0, - 3.0, - 3.0 - ], - "source_sha256": "7e064d69277762a33d2a65a4d6258e5422c7d2b3535694ce9afd6a6304f6993a", - "container_id": "ta-01M1YT65B8TXRZGSM820Z3Q16R", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 1, - "remote_seconds": 40.762046516 - }, - { - "case": "inline_complex", - "variant": 0, - "enabled": false, - "extension_name": "pch_matrix_inline_complex_0_0", - "seconds": 18.759751232, - "timing_samples": [ - 18.759751232 - ], - "host_compile_seconds": 18.012592808999997, - "pch_consumed": false, - "hit_keys": [], - "cache_miss": false, - "correctness_checks": 6, - "outputs": [ - 2.0, - 2.0, - 2.0, - 2.0, - 2.0 - ], - "source_sha256": "bb3eedfa1df83397af6a8f7edb81db46b0db7cda4a388c79a270329aa8ae85ee", - "container_id": "ta-01M1YT0HRZP289WPV9VMSSPF8R", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 1, - "remote_seconds": 40.230402710999996 - }, - { - "case": "inline_complex", - "variant": 0, - "enabled": true, - "extension_name": "pch_matrix_inline_complex_0_1", - "seconds": 14.449278815, - "timing_samples": [ - 14.449278815 - ], - "host_compile_seconds": 7.923061731999999, - "pch_consumed": true, - "hit_keys": [ - "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" - ], - "cache_miss": false, - "correctness_checks": 6, - "outputs": [ - 2.0, - 2.0, - 2.0, - 2.0, - 2.0 - ], - "source_sha256": "bb3eedfa1df83397af6a8f7edb81db46b0db7cda4a388c79a270329aa8ae85ee", - "container_id": "ta-01M1YT0J0EQEC8X2QHC6S0EAPR", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 1, - "remote_seconds": 26.335685509 - }, - { - "case": "inline_complex", - "variant": 1, - "enabled": true, - "extension_name": "pch_matrix_inline_complex_1_1", - "seconds": 14.031058936, - "timing_samples": [ - 14.031058936 - ], - "host_compile_seconds": 7.8286803460000005, - "pch_consumed": true, - "hit_keys": [ - "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" - ], - "cache_miss": false, - "correctness_checks": 6, - "outputs": [ - 3.0, - 3.0, - 3.0, - 3.0, - 3.0 - ], - "source_sha256": "8be4b20baeb5c3ea8fbb6ec714175f9d77dd02662be31f23bddffc7807074d66", - "container_id": "ta-01M1YT0HRZYPV863GAZA3K6T1R", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 1, - "remote_seconds": 26.304385472 - }, - { - "case": "inline_complex", - "variant": 1, - "enabled": false, - "extension_name": "pch_matrix_inline_complex_1_0", - "seconds": 21.116745669, - "timing_samples": [ - 21.116745669 - ], - "host_compile_seconds": 20.210806692, - "pch_consumed": false, - "hit_keys": [], - "cache_miss": false, - "correctness_checks": 6, - "outputs": [ - 3.0, - 3.0, - 3.0, - 3.0, - 3.0 - ], - "source_sha256": "8be4b20baeb5c3ea8fbb6ec714175f9d77dd02662be31f23bddffc7807074d66", - "container_id": "ta-01M1YT0HM8DWCT713GAJ6ZQRVR", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 1, - "remote_seconds": 45.54159261 - }, - { - "case": "load_files", - "variant": 0, - "enabled": false, - "extension_name": "pch_matrix_shared", - "seconds": 18.484804069, - "timing_samples": [ - 18.484804069 - ], - "host_compile_seconds": 17.716654833999996, - "pch_consumed": false, - "hit_keys": [], - "cache_miss": false, - "correctness_checks": 6, - "outputs": [ - 2.0, - 2.0, - 2.0, - 2.0, - 2.0 - ], - "source_sha256": "4ec65ca0b08feea11ce48b7a0ebd322676aab622d35eba1da63b605643e46435", - "container_id": "ta-01M1YT1MEHZQJTJ91K03EYFKJR", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 1, - "remote_seconds": 39.96083478 - }, - { - "case": "load_files", - "variant": 0, - "enabled": true, - "extension_name": "pch_matrix_shared", - "seconds": 8.844153404, - "timing_samples": [ - 8.844153404 - ], - "host_compile_seconds": 7.813371175, - "pch_consumed": true, - "hit_keys": [ - "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" - ], - "cache_miss": false, - "correctness_checks": 6, - "outputs": [ - 2.0, - 2.0, - 2.0, - 2.0, - 2.0 - ], - "source_sha256": "4ec65ca0b08feea11ce48b7a0ebd322676aab622d35eba1da63b605643e46435", - "container_id": "ta-01M1YT1MM8C9EBBGEZCWXZAW5R", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 1, - "remote_seconds": 20.400708645 - }, - { - "case": "load_files", - "variant": 1, - "enabled": true, - "extension_name": "pch_matrix_shared", - "seconds": 8.704362362000001, - "timing_samples": [ - 8.704362362000001 - ], - "host_compile_seconds": 7.6940678259999995, - "pch_consumed": true, - "hit_keys": [ - "cd04740b9ab36ac9a2a8d533e50fc30250ab89f341186829da169aefb0865f38" - ], - "cache_miss": false, - "correctness_checks": 6, - "outputs": [ - 3.0, - 3.0, - 3.0, - 3.0, - 3.0 - ], - "source_sha256": "7e064d69277762a33d2a65a4d6258e5422c7d2b3535694ce9afd6a6304f6993a", - "container_id": "ta-01M1YT20M1E9G09T9PZBWNQ65R", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 1, - "remote_seconds": 20.25399071 - }, - { - "case": "load_files", - "variant": 1, - "enabled": false, - "extension_name": "pch_matrix_shared", - "seconds": 18.118759269999998, - "timing_samples": [ - 18.118759269999998 - ], - "host_compile_seconds": 17.33960136, - "pch_consumed": false, - "hit_keys": [], - "cache_miss": false, - "correctness_checks": 6, - "outputs": [ - 3.0, - 3.0, - 3.0, - 3.0, - 3.0 - ], - "source_sha256": "7e064d69277762a33d2a65a4d6258e5422c7d2b3535694ce9afd6a6304f6993a", - "container_id": "ta-01M1YT276G72J91NSPNG9VWAWR", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 1, - "remote_seconds": 39.779365328000004 - }, - { - "case": "plain_cpp", - "variant": 0, - "enabled": false, - "extension_name": "pch_matrix_plain_cpp_0_0", - "seconds": 0.19980638800000028, - "timing_samples": [ - 0.21276275, - 0.19661020400000018, - 0.20195385499999974, - 0.20159454799999965, - 0.20105536800000046, - 0.1982173860000005, - 0.1969059470000003, - 0.1985574080000001, - 0.19711016499999978, - 0.2051067179999997 - ], - "host_compile_seconds": 0.12402935050000008, - "pch_consumed": false, - "hit_keys": [], - "cache_miss": false, - "correctness_checks": 1, - "outputs": [ - 3 - ], - "source_sha256": "f5fbb559817c22a648c4be4067de69dd806e6e9f2c9be71aa84f2371ecf8d781", - "container_id": "ta-01M1YT2FP2X908DV2VFHCJPT6R", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 0, - "remote_seconds": 2.3064990610000002 - }, - { - "case": "plain_cpp", - "variant": 0, - "enabled": true, - "extension_name": "pch_matrix_plain_cpp_0_1", - "seconds": 0.2823680209999999, - "timing_samples": [ - 0.35445732099999994, - 0.3628376659999999, - 0.30029828899999966, - 0.28355820099999995, - 0.28399713199999965, - 0.27632356199999997, - 0.2811778409999999, - 0.2746912210000003, - 0.2782820089999998, - 0.2755435680000007 - ], - "host_compile_seconds": 0.20570012650000002, - "pch_consumed": false, - "hit_keys": [], - "cache_miss": false, - "correctness_checks": 1, - "outputs": [ - 3 - ], - "source_sha256": "f5fbb559817c22a648c4be4067de69dd806e6e9f2c9be71aa84f2371ecf8d781", - "container_id": "ta-01M1YT2QSWZR9SN6VGQJC16VSR", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 0, - "remote_seconds": 3.291071738 - }, - { - "case": "plain_cpp", - "variant": 1, - "enabled": true, - "extension_name": "pch_matrix_plain_cpp_1_1", - "seconds": 0.29172716450000014, - "timing_samples": [ - 0.49523322599999986, - 0.36589996000000014, - 0.28990115899999935, - 0.291525472, - 0.28335249200000057, - 0.29192885700000026, - 0.2972645429999998, - 0.2952913079999995, - 0.2844164109999996, - 0.2789291899999995 - ], - "host_compile_seconds": 0.21366047899999963, - "pch_consumed": false, - "hit_keys": [], - "cache_miss": false, - "correctness_checks": 1, - "outputs": [ - 4 - ], - "source_sha256": "4f5b6c5b34f3f71e1e94ad89fb640e348e9bd0232d5cdfc75075c5e827674378", - "container_id": "ta-01M1YT2TY6A6XZNT3DJDM9AW6R", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 0, - "remote_seconds": 3.5053143249999996 - }, - { - "case": "plain_cpp", - "variant": 1, - "enabled": false, - "extension_name": "pch_matrix_plain_cpp_1_0", - "seconds": 0.20363044750000014, - "timing_samples": [ - 0.2254486870000001, - 0.20830646500000016, - 0.20089933599999998, - 0.207017467, - 0.2131324689999996, - 0.20014710299999994, - 0.20089258399999999, - 0.2038948989999998, - 0.19941997500000053, - 0.2033659960000005 - ], - "host_compile_seconds": 0.12606111700000033, - "pch_consumed": false, - "hit_keys": [], - "cache_miss": false, - "correctness_checks": 1, - "outputs": [ - 4 - ], - "source_sha256": "4f5b6c5b34f3f71e1e94ad89fb640e348e9bd0232d5cdfc75075c5e827674378", - "container_id": "ta-01M1YT3294D1MHQ774CC40NP6R", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 0, - "remote_seconds": 2.300308057 - }, - { - "case": "flag_miss", - "variant": 0, - "enabled": false, - "extension_name": "pch_matrix_flag_miss_0_0", - "seconds": 18.314668328000003, - "timing_samples": [ - 18.314668328000003 - ], - "host_compile_seconds": 17.468810722999997, - "pch_consumed": false, - "hit_keys": [], - "cache_miss": false, - "correctness_checks": 6, - "outputs": [ - 2.0, - 2.0, - 2.0, - 2.0, - 2.0 - ], - "source_sha256": "4ec65ca0b08feea11ce48b7a0ebd322676aab622d35eba1da63b605643e46435", - "container_id": "ta-01M1YT32ZPHKD6NAA1RQHDFZNR", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 1, - "remote_seconds": 39.548441550999996 - }, - { - "case": "flag_miss", - "variant": 0, - "enabled": true, - "extension_name": "pch_matrix_flag_miss_0_1", - "seconds": 20.427356331, - "timing_samples": [ - 20.427356331 - ], - "host_compile_seconds": 19.146695692999998, - "pch_consumed": false, - "hit_keys": [], - "cache_miss": true, - "correctness_checks": 6, - "outputs": [ - 2.0, - 2.0, - 2.0, - 2.0, - 2.0 - ], - "source_sha256": "4ec65ca0b08feea11ce48b7a0ebd322676aab622d35eba1da63b605643e46435", - "container_id": "ta-01M1YT36X83KGN6Y6BSXPHJB5R", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 1, - "remote_seconds": 44.605277338 - }, - { - "case": "flag_miss", - "variant": 1, - "enabled": true, - "extension_name": "pch_matrix_flag_miss_1_1", - "seconds": 20.175504439, - "timing_samples": [ - 20.175504439 - ], - "host_compile_seconds": 19.005651779999997, - "pch_consumed": false, - "hit_keys": [], - "cache_miss": true, - "correctness_checks": 6, - "outputs": [ - 3.0, - 3.0, - 3.0, - 3.0, - 3.0 - ], - "source_sha256": "7e064d69277762a33d2a65a4d6258e5422c7d2b3535694ce9afd6a6304f6993a", - "container_id": "ta-01M1YT3AQERWZZTC7SBEKE1EFR", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 1, - "remote_seconds": 43.945379948 - }, - { - "case": "flag_miss", - "variant": 1, - "enabled": false, - "extension_name": "pch_matrix_flag_miss_1_0", - "seconds": 17.837603579, - "timing_samples": [ - 17.837603579 - ], - "host_compile_seconds": 17.036646379, - "pch_consumed": false, - "hit_keys": [], - "cache_miss": false, - "correctness_checks": 6, - "outputs": [ - 3.0, - 3.0, - 3.0, - 3.0, - 3.0 - ], - "source_sha256": "7e064d69277762a33d2a65a4d6258e5422c7d2b3535694ce9afd6a6304f6993a", - "container_id": "ta-01M1YT3MBVRB90BKSZW7DR6FHR", - "image_id": "im-VGOID7nd5iWWNRTr1Lw2O7", - "gpu": "Tesla T4", - "torch": "2.12.0+cu130", - "nvcc_translation_units": 1, - "remote_seconds": 40.553874818 - } - ] -} diff --git a/docs/benchmarks/modal-pch.json b/docs/benchmarks/modal-pch.json new file mode 100644 index 000000000..56ba72916 --- /dev/null +++ b/docs/benchmarks/modal-pch.json @@ -0,0 +1,52 @@ +{ + "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 only for the CUTLASS GEMM run. The earlier matrix instruments CXX only.", + "benchmark_sha256": "3bf02d121088b5e3d564c538846615cb5637271afd60b3e12b15e48b36dd63fe", + "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" + }, + "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":"inline_complex","variant":0,"repeat":0,"cached":false,"build_s":18.759751232,"cpp_s":18.012592808999997,"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":"bb3eedfa1df83397af6a8f7edb81db46b0db7cda4a388c79a270329aa8ae85ee","container_id":"ta-01M1YT0HRZP289WPV9VMSSPF8R"}, + {"case":"inline_complex","variant":0,"repeat":0,"cached":true,"build_s":14.449278815,"cpp_s":7.923061731999999,"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":"bb3eedfa1df83397af6a8f7edb81db46b0db7cda4a388c79a270329aa8ae85ee","container_id":"ta-01M1YT0J0EQEC8X2QHC6S0EAPR"}, + {"case":"inline_complex","variant":1,"repeat":0,"cached":true,"build_s":14.031058936,"cpp_s":7.8286803460000005,"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":"8be4b20baeb5c3ea8fbb6ec714175f9d77dd02662be31f23bddffc7807074d66","container_id":"ta-01M1YT0HRZYPV863GAZA3K6T1R"}, + {"case":"inline_complex","variant":1,"repeat":0,"cached":false,"build_s":21.116745669,"cpp_s":20.210806692,"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":"8be4b20baeb5c3ea8fbb6ec714175f9d77dd02662be31f23bddffc7807074d66","container_id":"ta-01M1YT0HM8DWCT713GAJ6ZQRVR"}, + {"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"} + ] +} diff --git a/docs/modal-pch-cache.md b/docs/modal-pch-cache.md index c631bffb8..a76e068e0 100644 --- a/docs/modal-pch-cache.md +++ b/docs/modal-pch-cache.md @@ -1,291 +1,94 @@ -# Modal `load_inline` header cache +# Modal Torch header cache -Kernelbot's Modal image sets `CXX` to a small GCC wrapper. For a C++ translation -unit starting with `#include ` (the default generated by -PyTorch `load_inline`), it looks up a precompiled header in the -`kernelbot-torch-pch` Modal Volume. Existing submissions need no source changes. -Leave PyTorch's separate `use_pch` option at its default, `False`. +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. +File-based `load()` also benefits when its C++ source starts with the same include. -The Volume contains shared headers only. Kernel objects and shared libraries -remain in each container's normal local extension cache. -This accelerates the **host C++ binding**; it does not cache NVCC's CUDA header -parsing, kernel compilation, or arbitrary user-provided headers. +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. -## For submission authors: automatic after deployment +## Profiles and fallback -Keep using `load_inline` normally; no cache flag is required: +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. -```python -module = load_inline( - name="my_kernel", - cpp_sources=cpp_source, - cuda_sources=cuda_source, - functions=["my_kernel"], -) -``` - -Kernelbot configures `CXX` in the runner image. The normal implicit -`torch/extension.h` prefix and default host compiler flags match the warmed cache. -Do **not** add `use_pch=True` to enable this feature: that invokes PyTorch's -separate native PCH mechanism and can add redundant work. Leave it at its default. - -The standard warmup also covers exact `extra_cflags=["-O2"]`, `["-O3"]`, and -`["-O3", "-ffast-math"]` profiles. Other host flags or include environments may -miss, but still compile normally. `extra_cuda_cflags` apply to NVCC, so changing -only those does not change the host C++ cache key. No submission author needs to -create, mount, reload, or write a Modal Volume; the operator owns warmup/deployment. +Unknown profiles compile normally; sources without the initial Torch include +bypass PCH. With `no_implicit_headers=True`, explicitly include +`` first in C++ to use the cache. Python/Triton code that does +not invoke `CXX` is unaffected. -For `no_implicit_headers=True`, explicitly put `#include ` -first in the C++ source if you want this cache to apply. The same condition -allows file-based `load()` to benefit. Pure Python/Triton submissions do not use -this C++ cache. To bypass all wrapper overhead for a tiny or unusual C++ build, -set `os.environ["CXX"] = "/usr/bin/g++"` before calling the extension loader. +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. -This feature remains a proposed runner change until PR #527 is deployed; the -reported tests ran in separate Modal test apps. +## Deployment -## Warm and deploy - -The existing main/dev deployment workflow warms all profiles before deploying -the GPU functions. A warmup failure prevents that rollout. For a manual rollout, -run these commands from the repository root in the same Modal profile/environment: +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 ``` -Run warmup again after changing the image, compiler, installed headers, Python, -or torch. Warmup uses a CPU container, not a GPU. It builds both CPU-only and -CUDA-linked `load_inline` profiles for these exact `extra_cflags` lists: - -- omitted / `[]` -- `["-O2"]` -- `["-O3"]` -- `["-O3", "-ffast-math"]` - -The warm function also accepts a comma-separated selection of profile names, such as -`--profiles cuda-default,cuda-O3`. Extend the trusted profiles -in `pch/warm.py` and the allowed names in `warm_pch` to support additional common -flag combinations. Arbitrary submission code and compiler arguments are never -sent to the shared-cache writer. - -`KERNELBOT_PCH_VOLUME` selects a different Volume at launch/deploy time. The -benchmark below uses this to avoid touching the production cache. A missing -profile compiles normally without writing to the Volume. No special handling is -needed for a cold deployment, though it will not receive a PCH speedup until -warmup completes and new containers start. Warm before deploying; after adding -profiles to an existing deployment, replace existing runner containers to refresh -their snapshots. The restricted GPU functions cannot call the Volume reload API. - -## Isolation and compatibility - -- GPU runners mount the Volume read-only and disable Modal API access with - `restrict_modal_access=True`. The API restriction prevents bypassing a read-only - filesystem mount with Volume upload APIs. Only the trusted CPU warmup writes it. -- Keys include the image's installed-header/compiler fingerprint, compiler - flags, include paths and header-search environment. Different module names, - kernel source filenames, object paths and dependency filenames share a key. -- ABI macros, optimization settings, C++ standards and include-path changes - select different keys. Unsupported profiles fall back to normal compilation. -- The wrapper leaves linking, compiler probes and sources without the initial - torch header alone. An explicit alternative `CXX` also bypasses it. -- GCC validates the PCH when loading it; `-Winvalid-pch` reports rejection and - GCC can parse the header normally. A cache-hit log alone is not proof of use. -- Warmup publishes a completed PCH with an atomic rename, then commits the - Volume. It does not rely on a distributed filesystem lock. - -PyTorch's native `use_pch=True` maintains a single `extension.h.gch` beside its -installed headers, and ordinary `load_inline` calls remove that native cache. -Our separate per-configuration cache is unaffected by those removals. Explicit -`use_pch=True` may still spend time building PyTorch's native PCH, so it is -unnecessary here. - -References: [PyTorch `load_inline`](https://docs.pytorch.org/docs/main/cpp_extension.html), -[Modal Volume consistency](https://modal.com/docs/guide/volumes), -[GCC precompiled headers](https://gcc.gnu.org/onlinedocs/gcc/Precompiled-Headers.html). +`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. -## Keep CUDA source headers small +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. -`load_inline` normally injects `torch/types.h` into CUDA sources too. That can -make NVCC the bottleneck even after the C++ binding becomes faster. Where CUDA -code only needs device pointers, keep Tensor handling in C++ and use -`no_implicit_headers=True`. Explicitly include `torch/extension.h` first in the -C++ source so the PCH still applies, and include only the CUDA headers needed -by the device code. See the working FP16 -[vector-add submission](../examples/vectoradd_py/submission_cuda_pch.py). - -## Real submission benchmark +## Benchmark ```sh PYTHONPATH=src:src/runners KERNELBOT_PCH_VOLUME=kernelbot-pch-test \ uv run modal run scripts/modal_pch_benchmark.py - -# Also check the original example with implicit Torch headers in CUDA: +# 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 --headers implicit \ - --output /tmp/kernelbot-pch-implicit-results.json + uv run modal run scripts/modal_pch_benchmark.py --cases cutlass_gemm --repeats 2 ``` -This uses the runner image and `modal_run_config` evaluation path, with real CUDA -vector-add submissions on a T4 (4 requested CPU cores, 16 GiB memory). It warms -one header profile on CPU, then runs two different kernel variants with the PCH -disabled and enabled, each in a fresh GPU container. It uses normal Ninja -parallelism (no `MAX_JOBS=1` override). The default fixture uses minimal CUDA -headers; `--headers implicit` uses the original vector-add example. It checks -all five vector-add test cases, confirms filesystem and API writes are denied, -checks GCC's `-H` output for actual PCH consumption, and records both C++ -compilation and `load_inline` wall times. -It asserts that the warmed C++ median is faster, and additionally checks the -end-to-end median for the minimal-header fixture. NVCC can hide that improvement -in the implicit-header fixture, so its total-time speedup is reported without an -assertion. The script writes raw JSON results to -`/tmp/kernelbot-pch-results.json`. It does not deploy the production app. - -To diagnose a submission, `KERNELBOT_PCH_TRACE=1` prints the host compile time and -GCC's header trace on PCH-enabled compilations. `KERNELBOT_PCH_DISABLE=1` provides -a baseline without the cache. - -## Similar sources, other build APIs, and performance limits - -Reuse is based on the Torch header's compilation configuration, not source-text -similarity. Different kernel bodies, filenames, extension names, or user headers -included **after** `torch/extension.h` can share the same Torch PCH. Their source -and user headers are still compiled separately into new objects/shared libraries; -none of that submission code is persisted in the Volume. Changing a local header -therefore changes the resulting code, even when the Torch PCH key stays the same. - -Changing host compiler flags, ABI macros, include paths, or header-search -environment selects another key. If that exact profile was not warmed, the -wrapper compiles normally. A macro or other header before the initial Torch -include makes the wrapper bypass PCH entirely, preserving preprocessing order. -Installed header changes require rebuilding the image/fingerprint and rewarming. - -The wrapper operates at `CXX`, so it can also accelerate -`torch.utils.cpp_extension.load()` with file sources beginning with the same Torch -include and matching flags. Plain C++ without that prefix bypasses PCH. Builds -that explicitly choose another compiler bypass the wrapper. Python/Triton paths -that do not invoke `CXX` never invoke this cache. - -There is **no guarantee of zero slowdown**. Even a miss/bypass adds a Python -compiler-wrapper invocation; hits must load a large PCH from the Volume, and -NVCC-heavy builds may receive little or no end-to-end benefit. Warmup itself costs -CPU time and storage, although it runs before deployment rather than on a -submission's GPU. `KERNELBOT_PCH_DISABLE=1` skips PCH lookup, but still starts the -wrapper; `CXX=/usr/bin/g++` bypasses the wrapper entirely. - -The broader regression matrix runs simple CUDA headers, CuTe/MathDx headers, -file-based `load()`, a plain `` C++ program, and unwarmed compiler flags: - -```sh -PYTHONPATH=src:src/runners KERNELBOT_PCH_VOLUME=kernelbot-pch-test-20260907 \ - uv run modal run scripts/modal_pch_matrix.py -``` - -Each case changes a local header (`kVariant = 1` versus `2`); CUDA cases also -change the kernel body (256 versus 128 threads), producing different expected -outputs. The file-based `load()` case even keeps the extension name identical -across variants. Each variant -runs with direct GCC and with the cache wrapper in a fresh T4 container. The reference -compiler is real `/usr/bin/g++`, rather than the cache wrapper in disabled mode. -Both modes have identical external timing instrumentation. GCC `-H` verification -runs in a separate compile **after** timing and correctness checks. Plain C++ -uses ten compiles per container to make wrapper overhead visible. No timing -threshold forces a favorable result; the matrix reports all timings and asserts -correctness, the intended hit/miss behavior, actual PCH consumption, unique -container IDs, changed source hashes, and distinct outputs for the variants. -The complex fixture enables CUDA half operators/conversions via NVCC `-U` flags -needed by MathDx; these device-only flags do not change the host PCH key. - -## Broader matrix results (2026-09-07) - -CUDA 13.3 / PyTorch 2.12.0+cu130, T4, normal Ninja parallelism. Each table value -is the median across two changed source/header variants, each tested in separate -direct-GCC and cache-enabled containers. Plain C++ first takes the median of ten -object compilations per container; its linking is excluded. - -| Case | Direct GCC | Cache wrapper | Ratio | -| --- | ---: | ---: | ---: | -| Minimal CUDA headers, `load_inline` | 18.847s | 9.998s | 1.89× | -| CuTe/MathDx headers, `load_inline` | 19.938s | 14.240s | 1.40× | -| File-based `load()`, same extension name | 18.302s | 8.774s | 2.09× | -| Plain C++ `` object compilation | 0.202s | 0.287s | 0.70× | -| Unwarmed host flags, `load_inline` | 18.076s | 20.301s | 0.89× | - -All **100 correctness checks passed across 20 fresh containers**. The six warmed -CUDA builds (simple, complex, and file-based, two variants each) consumed the -same Torch PCH. The CUDA outputs changed from 2 to 3 as expected when the local -header changed; the file-based test used the same extension name for both. -Both changed-flag builds logged cache misses, and plain C++ bypassed PCH. - -The plain C++ compiler stage, excluding the common timing shim's startup, went -from about 0.125s to 0.210s: roughly **85ms wrapper overhead**. The cache-miss -case was also slower in this sample (18.08s to 20.30s); with two source variants -on different machines, this does not isolate wrapper/filesystem cost from -machine variability. These results explicitly do not establish zero regressions. -For a tiny build that must avoid wrapper cost, set `CXX=/usr/bin/g++`. - -CuTe/MathDx headers still compile through NVCC; the PCH only saves Torch host-header -work. The complex case initially failed because PyTorch's NVCC half-type -suppression macros conflict with MathDx. The final fixture undoes those macros -only for NVCC, then passes in both baseline and cached modes. The successful -matrix runs below use that corrected fixture, with no changes to production -compiler-cache code. - -The harness and updated key regression test pass Ruff and the 41 relevant local -tests. [Raw results, per-container timings, and successful Modal run links](benchmarks/modal-pch-matrix-2026-09-07.json). - -## Current-main measurements (2026-09-07) - -Revalidated on main `727212cd`, preserving CUDA 13.3.0, PyTorch 2.12.0+cu130, -and MathDx. Both fixtures use the same image and freshly warmed `cuda-default` -profile, normal Ninja parallelism, and eight distinct T4 containers. - -| Submission | C++ compilation, cold → PCH | Total `load_inline`, cold → PCH | -| --- | --- | --- | -| Minimal CUDA headers | 19.36s → 8.91s (2.17×) | 20.45s → 9.91s (2.06×) | -| Original implicit headers | 18.44s → 8.79s (2.10×) | 38.10s → 43.25s (0.88×) | - -All 40 correctness checks passed. Every cache-enabled build consumed the same -PCH, and every container rejected filesystem and Modal API writes. The -Torch-heavy CUDA fixture did **not** improve end-to-end in this run: warm totals -were 49.99s and 36.50s, despite faster host compilation. These two-sample medians -show the host-stage benefit and the minimal-header fixture's total-time benefit; -they do not establish an end-to-end speedup for NVCC-heavy submissions. - -41 relevant local tests and Ruff passed on this base. GitHub unit, lint, Modal -integration, and GitHub integration checks also passed on `38b98928`. -Only `cuda-default` was rewarmed on the current image; the eight-profile warmup -below applies to the earlier image. The deployment workflow warms all profiles -for the actual image before deploying it. Production has not been deployed by -this experiment. - -[Current-main raw measurements and run links](benchmarks/modal-pch-current-main-2026-09-07.json). - -## Earlier-image measurements (2026-09-07) - -T4, CUDA 12.9.1, PyTorch 2.11.0+cu129, normal Ninja parallelism. Each row below -is the median of two different kernel variants, each tested in separate cold -and PCH-enabled containers. Both fixtures share the same warmed header entry. - -| Submission | C++ compilation, cold → PCH | Total `load_inline`, cold → PCH | -| --- | --- | --- | -| Minimal CUDA headers | 15.67s → 7.12s (2.20×) | 16.43s → 7.93s (2.07×) | -| Original implicit headers | 16.02s → 6.29s (2.54×) | 33.90s → 32.83s (1.03×) | - -All 40 correctness checks passed across eight fresh Modal container IDs. GCC -confirmed actual PCH consumption on every cache-enabled build; filesystem and -Modal API writes were denied. All eight warmup profiles were built and committed -to `kernelbot-pch-test-20260907`, including successful reuse after CPU preemption. -The implementation also passed 36 local tests and the repository Ruff check. - -The original example's CUDA compilation dominates its total time, so speeding -up the parallel C++ stage alone has little effect on that workload. The -minimal-header example demonstrates the PCH benefit without that bottleneck. -The initial exploratory single-worker timings are not included in this table. - -[Raw measurements and run links](benchmarks/modal-pch-2026-09-07.json). -The eight final GPU calls used approximately $0.044 of T4 GPU compute at -[Modal's published rate](https://modal.com/pricing), excluding startup, -exploratory runs, CPU warmup, memory and image builds. The production runner -has not been deployed by this experiment. +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, including +file-based `load()` with the same extension name, against distinct expected outputs. + +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 plain C++ row measures object compilation only. + +| Build | Baseline | Cached | +| --- | ---: | ---: | +| Minimal CUDA headers | 18.85s | 10.00s | +| CuTe/MathDx includes (vector add) | 19.94s | 14.24s | +| CUTLASS tensor-core GEMM | 19.95s | 12.58s | +| File-based `load()` | 18.30s | 8.77s | +| Plain C++ | 0.202s | 0.287s | +| Unwarmed host flags | 18.08s | 20.30s | + +The CUTLASS GEMM build was **1.59× faster** (37% less time). Host compilation went +from 19.10s to 7.88s; NVCC took 12.54s and 11.53s respectively and became the +bottleneck. The cache does not accelerate NVCC; those separate-container times +include machine variability. All 80 GEMM checks passed, with maximum absolute +error 8.53e-6. The broader cases passed another 100 checks. + +The plain C++ wrapper added about 85ms; the miss case was also slower in this +sample. Larger CUDA template builds may hide the host-side savings entirely. +[Per-container results and Modal runs](benchmarks/modal-pch.json). diff --git a/examples/vectoradd_py/submission_cuda_pch.py b/examples/vectoradd_py/submission_cuda_pch.py deleted file mode 100644 index f18720e93..000000000 --- a/examples/vectoradd_py/submission_cuda_pch.py +++ /dev/null @@ -1,62 +0,0 @@ -#!POPCORN leaderboard vectoradd_py-dev -"""FP16 vector add: keep Torch headers in C++, and only CUDA headers in CUDA. - -Kernelbot's Modal CXX wrapper reuses the C++ torch header PCH automatically. -no_implicit_headers avoids unnecessarily parsing torch/types.h again with nvcc. -""" - -from task import input_t, output_t -from torch.utils.cpp_extension import load_inline - -add_cpp_source = r""" -#include -#include -#include - -void add_cuda_impl(const void*, const void*, void*, int, cudaStream_t); - -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 C = torch::empty_like(A); - add_cuda_impl(A.data_ptr(), B.data_ptr(), C.data_ptr(), A.numel(), - c10::cuda::getCurrentCUDAStream().stream()); - TORCH_CHECK(cudaGetLastError() == cudaSuccess, "CUDA launch failed"); - return C; -} -""" - -add_cuda_source = r""" -#include -#include - -__global__ void add_kernel(const __half* A, const __half* B, __half* C, int N) { - const int idx = blockIdx.x * blockDim.x + threadIdx.x; - if (idx < N) C[idx] = __hadd(A[idx], B[idx]); -} - -void add_cuda_impl(const void* A, const void* B, void* C, int N, cudaStream_t stream) { - const int threads = 256; - add_kernel<<<(N + threads - 1) / threads, threads, 0, stream>>>( - static_cast(A), static_cast(B), static_cast<__half*>(C), N); -} -""" - -add_module = load_inline( - name='add_cuda', - cpp_sources=add_cpp_source, - cuda_sources=add_cuda_source, - functions=['add_cuda'], - no_implicit_headers=True, - verbose=True, -) - - -def add(A, B): - return add_module.add_cuda(A, B) - - -def custom_kernel(data: input_t) -> output_t: - return add(*data) diff --git a/scripts/modal_pch_benchmark.py b/scripts/modal_pch_benchmark.py index 9cce1d261..06187799e 100644 --- a/scripts/modal_pch_benchmark.py +++ b/scripts/modal_pch_benchmark.py @@ -1,22 +1,124 @@ -"""Real submissions in fresh Modal containers; no production deployment required. +"""Benchmark Torch PCH reuse with CUDA, CUTLASS, file-based builds, and cache misses. -PYTHONPATH=src:src/runners KERNELBOT_PCH_VOLUME=kernelbot-pch-test \\ +PYTHONPATH=src:src/runners KERNELBOT_PCH_VOLUME=kernelbot-pch-test \ uv run modal run scripts/modal_pch_benchmark.py """ -import io import json -import os -import re import statistics -import time from pathlib import Path import modal -from modal_runner import PCH_MOUNT, cuda_image, modal_run_config, pch_volume, warm_pch +from modal_runner import PCH_MOUNT, cuda_image, pch_volume, warm_pch -app = modal.App("kernelbot-pch-benchmark") +app = modal.App("kernelbot-pch-matrix") app.include(warm_pch.app) +CASES = ("inline_simple", "inline_complex", "load_files", "plain_cpp", "flag_miss", "cutlass_gemm") + +# 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 +#include +#include +#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 +#include +#include "variant.h" +// EXTRA_HEADERS +__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(a), static_cast(b), static_cast<__half*>(out), n); +} +''' + +GEMM_CPP = r''' +#include +#include +#include +#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(), + 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 +#include +#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, + 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(a), k}, + {static_cast(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( @@ -25,103 +127,218 @@ 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 submit(config: dict, enabled: bool): - os.environ["KERNELBOT_PCH_DISABLE"] = "0" if enabled else "1" - os.environ["KERNELBOT_PCH_TRACE"] = "1" - os.environ.pop("MAX_JOBS", None) # Match normal load_inline/Ninja parallelism. +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() - result = modal_run_config(config) - assert result.success, result.error - test = result.runs["test"] - assert test.run.success and test.run.passed, test.run - assert test.compilation and test.compilation.exit_code == 0, test.compilation - logs = test.compilation.stdout + test.compilation.stderr - durations = re.findall(r"PCH_LOAD_INLINE_SECONDS=([0-9.]+)", logs) - assert durations, logs[-5000:] - consumed = "! /kernelbot-pch/" in logs - assert consumed == enabled, logs[-5000:] - # A read-only cache must remain read-only from submission containers. - try: - Path(PCH_MOUNT, "submission-write-probe").write_text("must fail") - except OSError: - pass - else: - raise AssertionError("Submission can write the shared PCH Volume") - try: - with pch_volume.batch_upload() as batch: - batch.put_file(io.BytesIO(b"must fail"), "/submission-api-write-probe") - except modal.exception.AuthError: - pass - else: - raise AssertionError("Submission can write the shared PCH Volume through the API") - return { - "api_write_denied": True, - "enabled": enabled, - "gpu": result.system.gpu, - "torch": result.system.torch, - "container_id": os.environ["MODAL_TASK_ID"], - "image_id": os.environ["MODAL_IMAGE_ID"], - "load_inline_seconds": float(durations[0]), - "cpp_seconds": [float(v) for v in re.findall(r"compile_seconds=([0-9.]+)", logs)], - "pch_consumed": consumed, - "keys": re.findall(r"\[kernelbot-pch\] hit (\w+)", logs), - "tests": test.run.result, - "remote_seconds": time.perf_counter() - started, - } + 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 == "cutlass_gemm": + cpp = GEMM_CPP + cuda = GEMM_CUDA.replace("TILE_N", str(128 if variant == 0 else 64)) + if case == "inline_complex": + cpp = cpp.replace( + '#include "variant.h"', + '#include "variant.h"\n#include \n#include ', + ) + cuda = cuda.replace( + "// EXTRA_HEADERS", + "#include \n#include \n" + "static_assert(cute::Int<128>::value == 128);", + ) + name = "pch_matrix_shared" if case == "load_files" else f"pch_matrix_{case}_{variant}_{int(enabled)}" + if case == "plain_cpp": + cpp = '#include \n#include \n#include "variant.h"\nint main() {\n' + cpp += 'std::array 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 in {"inline_complex", "cutlass_gemm"}: + # MathDx needs half operators/conversions disabled by PyTorch's NVCC defaults. + # Only device flags change; the host C++ PCH remains compatible. + 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", + ] +if case == "cutlass_gemm": + options["extra_cuda_cflags"] += ["-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=True, + 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 == "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 == "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 {"inline_simple", "inline_complex", "load_files", "cutlass_gemm"} + 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 == "cutlass_gemm" else 0, + "remote_seconds": time.perf_counter() - started, + } @app.local_entrypoint() -def main(output: str = "/tmp/kernelbot-pch-results.json", repeats: int = 2, headers: str = "minimal"): - if headers not in {"minimal", "implicit"}: - raise ValueError("headers must be minimal or implicit") - from libkernelbot.consts import SubmissionMode - from libkernelbot.task import build_task_config, make_task_definition - +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")) - root = Path(__file__).resolve().parents[1] / "examples/vectoradd_py" - task = make_task_definition(root).task - filename = "submission_cuda_pch.py" if headers == "minimal" else "submission_cuda_inline.py" - source = (root / filename).read_text() + 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 index in range(repeats): - for enabled in (False, True): - variant = source.replace( - "const int threads = 256", f"const int threads = {128 if index else 256}" - ) - variant = variant.replace( - "name='add_cuda'", f"name='pch_variant_{index}_{int(enabled)}'" - ) - variant = variant.replace( - "add_module = load_inline(", - "import time\n_pch_started = time.perf_counter()\nadd_module = load_inline(", - ) - variant = variant.replace( - "\ndef add(A, B):", - '\nprint(f"PCH_LOAD_INLINE_SECONDS={time.perf_counter() - _pch_started:.3f}", flush=True)\n' - "\ndef add(A, B):", - ) - config = build_task_config( - task=task, submission_content=variant, arch="75", mode=SubmissionMode.TEST - ) - row = submit.remote(config, enabled) - row["headers"] = headers - row["variant"] = index - rows.append(row) - Path(output).write_text(json.dumps(rows, indent=2)) - print(json.dumps(row), flush=True) - assert len({row["container_id"] for row in rows}) == len(rows), "Expected fresh containers" - warm_keys = {key for row in rows if row["enabled"] for key in row["keys"]} - assert len(warm_keys) == 1, "Different kernel variants must reuse the same header entry" - cold = statistics.median(row["load_inline_seconds"] for row in rows if not row["enabled"]) - warm = statistics.median(row["load_inline_seconds"] for row in rows if row["enabled"]) - print(f"load_inline median: cold={cold:.3f}s warm={warm:.3f}s speedup={cold / warm:.2f}x") - cpp_cold = statistics.median(row["cpp_seconds"][0] for row in rows if not row["enabled"]) - cpp_warm = statistics.median(row["cpp_seconds"][0] for row in rows if row["enabled"]) - print(f"C++ median: cold={cpp_cold:.3f}s warm={cpp_warm:.3f}s speedup={cpp_cold / cpp_warm:.2f}x") - assert cpp_warm < cpp_cold, "PCH did not improve host compilation" - if headers == "minimal": - assert warm < cold, "PCH did not improve end-to-end compilation" + 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"]} + cached_cases = {"inline_simple", "inline_complex", "load_files", "cutlass_gemm"} + 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/scripts/modal_pch_matrix.py b/scripts/modal_pch_matrix.py deleted file mode 100644 index c5ae72342..000000000 --- a/scripts/modal_pch_matrix.py +++ /dev/null @@ -1,260 +0,0 @@ -"""Broader PCH regression matrix: real compilation and distinct GPU outputs. - -PYTHONPATH=src:src/runners KERNELBOT_PCH_VOLUME=kernelbot-pch-test-20260907 \ - uv run modal run scripts/modal_pch_matrix.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", "inline_complex", "load_files", "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() -code = subprocess.call([os.environ["PCH_TEST_COMPILER"], *sys.argv[1:]]) -if "-c" in sys.argv: - with open(os.environ["PCH_TEST_COMMANDS"], "a") as f: - f.write(json.dumps({"args": sys.argv[1:], "seconds": time.perf_counter()-started}) + "\n") -sys.exit(code) -''' - -CPP = r''' -#include -#include -#include -#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 -#include -#include "variant.h" -// EXTRA_HEADERS -__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(a), static_cast(b), static_cast<__half*>(out), n); -} -''' - - -@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): - import hashlib - import os - import re - import shlex - 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) - 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 == "inline_complex": - cpp = cpp.replace( - '#include "variant.h"', - '#include "variant.h"\n#include \n#include ', - ) - cuda = cuda.replace( - "// EXTRA_HEADERS", - "#include \n#include \n" - "static_assert(cute::Int<128>::value == 128);", - ) - name = "pch_matrix_shared" if case == "load_files" else f"pch_matrix_{case}_{variant}_{int(enabled)}" - if case == "plain_cpp": - cpp = '#include \n#include \n#include "variant.h"\nint main() {\n' - cpp += 'std::array 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 == "inline_complex": - # MathDx needs half operators/conversions disabled by PyTorch's NVCC defaults. - # Only device flags change; the host C++ PCH remains compatible. - 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", - ] -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=True, - extra_cflags=[f"-DPCH_CASE={variant+1}"] if case == "flag_miss" else [], **options) -seconds = time.perf_counter() - started -assert module.version() == variant + 1 -outputs = [] -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))) -''' - 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 = 6 - times = [wall] - observed = [json.loads(line) for line in commands.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 {"inline_simple", "inline_complex", "load_files"} - 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:] - ninja = root / "build.ninja" - nvcc_commands = [] - if ninja.exists(): - rendered = subprocess.check_output(["ninja", "-C", str(root), "-t", "commands"], text=True) - nvcc_commands = [shlex.split(line) for line in rendered.splitlines() if "/nvcc " in line] - return { - "case": case, "variant": variant, "enabled": enabled, "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_translation_units": len(nvcc_commands), - "remote_seconds": time.perf_counter() - started, - } - - -@app.local_entrypoint() -def main(output: str = "/tmp/kernelbot-pch-matrix.json", cases: str = ""): - 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) - 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) & {"inline_simple", "inline_complex", "load_files"} 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") From e226b8627886189a589dae31c7b329aebbd4fc47 Mon Sep 17 00:00:00 2001 From: Mark Saroufim Date: Mon, 7 Sep 2026 14:18:37 -0700 Subject: [PATCH 7/8] Measure CUTLASS with default and minimal CUDA headers --- docs/benchmarks/modal-pch.json | 22 +++++++++++------- docs/modal-pch-cache.md | 25 +++++++++++++-------- scripts/modal_pch_benchmark.py | 41 +++++++++++++--------------------- 3 files changed, 46 insertions(+), 42 deletions(-) diff --git a/docs/benchmarks/modal-pch.json b/docs/benchmarks/modal-pch.json index 56ba72916..0d4a08f92 100644 --- a/docs/benchmarks/modal-pch.json +++ b/docs/benchmarks/modal-pch.json @@ -11,23 +11,21 @@ "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 only for the CUTLASS GEMM run. The earlier matrix instruments CXX only.", - "benchmark_sha256": "3bf02d121088b5e3d564c538846615cb5637271afd60b3e12b15e48b36dd63fe", + "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": "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":"inline_complex","variant":0,"repeat":0,"cached":false,"build_s":18.759751232,"cpp_s":18.012592808999997,"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":"bb3eedfa1df83397af6a8f7edb81db46b0db7cda4a388c79a270329aa8ae85ee","container_id":"ta-01M1YT0HRZP289WPV9VMSSPF8R"}, - {"case":"inline_complex","variant":0,"repeat":0,"cached":true,"build_s":14.449278815,"cpp_s":7.923061731999999,"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":"bb3eedfa1df83397af6a8f7edb81db46b0db7cda4a388c79a270329aa8ae85ee","container_id":"ta-01M1YT0J0EQEC8X2QHC6S0EAPR"}, - {"case":"inline_complex","variant":1,"repeat":0,"cached":true,"build_s":14.031058936,"cpp_s":7.8286803460000005,"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":"8be4b20baeb5c3ea8fbb6ec714175f9d77dd02662be31f23bddffc7807074d66","container_id":"ta-01M1YT0HRZYPV863GAZA3K6T1R"}, - {"case":"inline_complex","variant":1,"repeat":0,"cached":false,"build_s":21.116745669,"cpp_s":20.210806692,"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":"8be4b20baeb5c3ea8fbb6ec714175f9d77dd02662be31f23bddffc7807074d66","container_id":"ta-01M1YT0HM8DWCT713GAJ6ZQRVR"}, {"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"}, @@ -47,6 +45,14 @@ {"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","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 index a76e068e0..fb2b90012 100644 --- a/docs/modal-pch-cache.md +++ b/docs/modal-pch-cache.md @@ -53,7 +53,7 @@ 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 --repeats 2 + 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 @@ -77,18 +77,25 @@ CUTLASS. The plain C++ row measures object compilation only. | Build | Baseline | Cached | | --- | ---: | ---: | | Minimal CUDA headers | 18.85s | 10.00s | -| CuTe/MathDx includes (vector add) | 19.94s | 14.24s | -| CUTLASS tensor-core GEMM | 19.95s | 12.58s | +| CUTLASS GEMM, default implicit CUDA headers | 48.08s | 45.61s | +| CUTLASS GEMM, `no_implicit_headers=True` | 19.95s | 12.58s | | File-based `load()` | 18.30s | 8.77s | | Plain C++ | 0.202s | 0.287s | | Unwarmed host flags | 18.08s | 20.30s | -The CUTLASS GEMM build was **1.59× faster** (37% less time). Host compilation went -from 19.10s to 7.88s; NVCC took 12.54s and 11.53s respectively and became the -bottleneck. The cache does not accelerate NVCC; those separate-container times -include machine variability. All 80 GEMM checks passed, with maximum absolute -error 8.53e-6. The broader cases passed another 100 checks. +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 passed another 76 checks. All eight warmup profiles also completed and +committed on this image. The plain C++ wrapper added about 85ms; the miss case was also slower in this -sample. Larger CUDA template builds may hide the host-side savings entirely. +sample. Separate containers introduce timing variability, and larger CUDA +builds may hide the host-side savings entirely. [Per-container results and Modal runs](benchmarks/modal-pch.json). diff --git a/scripts/modal_pch_benchmark.py b/scripts/modal_pch_benchmark.py index 06187799e..b4191ff16 100644 --- a/scripts/modal_pch_benchmark.py +++ b/scripts/modal_pch_benchmark.py @@ -13,7 +13,12 @@ app = modal.App("kernelbot-pch-matrix") app.include(warm_pch.app) -CASES = ("inline_simple", "inline_complex", "load_files", "plain_cpp", "flag_miss", "cutlass_gemm") +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. @@ -54,7 +59,6 @@ #include #include #include "variant.h" -// EXTRA_HEADERS __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))); @@ -165,19 +169,9 @@ def trial(case: str, variant: int, enabled: bool, repeat: int = 0): (root / "variant.h").write_text(header) cpp = CPP cuda = CUDA.replace("THREADS", str(256 if variant == 0 else 128)) - if case == "cutlass_gemm": + if case.startswith("cutlass_gemm"): cpp = GEMM_CPP cuda = GEMM_CUDA.replace("TILE_N", str(128 if variant == 0 else 64)) - if case == "inline_complex": - cpp = cpp.replace( - '#include "variant.h"', - '#include "variant.h"\n#include \n#include ', - ) - cuda = cuda.replace( - "// EXTRA_HEADERS", - "#include \n#include \n" - "static_assert(cute::Int<128>::value == 128);", - ) name = "pch_matrix_shared" if case == "load_files" else f"pch_matrix_{case}_{variant}_{int(enabled)}" if case == "plain_cpp": cpp = '#include \n#include \n#include "variant.h"\nint main() {\n' @@ -219,28 +213,26 @@ def trial(case: str, variant: int, enabled: bool, repeat: int = 0): 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 in {"inline_complex", "cutlass_gemm"}: - # MathDx needs half operators/conversions disabled by PyTorch's NVCC defaults. - # Only device flags change; the host C++ PCH remains compatible. +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", ] -if case == "cutlass_gemm": - options["extra_cuda_cflags"] += ["-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=True, + 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 == "cutlass_gemm": +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)): @@ -277,12 +269,12 @@ def trial(case: str, variant: int, enabled: bool, repeat: int = 0): measurement = json.loads((root / "result.json").read_text()) wall = measurement["seconds"] device_outputs = measurement["outputs"] - correctness_checks = 10 if case == "cutlass_gemm" else 6 + 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 {"inline_simple", "inline_complex", "load_files", "cutlass_gemm"} + 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. @@ -308,7 +300,7 @@ def trial(case: str, variant: int, enabled: bool, repeat: int = 0): "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 == "cutlass_gemm" else 0, + "max_absolute_error": max(measurement["errors"], default=0) if case.startswith("cutlass_gemm") else 0, "remote_seconds": time.perf_counter() - started, } @@ -333,8 +325,7 @@ def main(output: str = "/tmp/kernelbot-pch-matrix.json", cases: str = "", repeat 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"]} - cached_cases = {"inline_simple", "inline_complex", "load_files", "cutlass_gemm"} - assert len(hit_keys) == (1 if set(selected) & cached_cases else 0) + 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 From 034fda3782cb57e0c5691678c0466255fbcf8cf0 Mon Sep 17 00:00:00 2001 From: Mark Saroufim Date: Mon, 7 Sep 2026 14:28:37 -0700 Subject: [PATCH 8/8] Focus PCH results on load_inline submissions --- docs/modal-pch-cache.md | 21 +++++++++++---------- 1 file changed, 11 insertions(+), 10 deletions(-) diff --git a/docs/modal-pch-cache.md b/docs/modal-pch-cache.md index fb2b90012..845900cba 100644 --- a/docs/modal-pch-cache.md +++ b/docs/modal-pch-cache.md @@ -3,7 +3,7 @@ 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. -File-based `load()` also benefits when its C++ source starts with the same include. +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. @@ -59,8 +59,9 @@ PYTHONPATH=src:src/runners KERNELBOT_PCH_VOLUME=kernelbot-pch-test \ 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, including -file-based `load()` with the same extension name, against distinct expected outputs. +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 @@ -72,15 +73,13 @@ time, not GEMM execution speed. 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 plain C++ row measures object compilation only. +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 | -| File-based `load()` | 18.30s | 8.77s | -| Plain C++ | 0.202s | 0.287s | | Unwarmed host flags | 18.08s | 20.30s | With default implicit CUDA headers, the CUTLASS GEMM improved only **1.05×**, @@ -92,10 +91,12 @@ 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 passed another 76 checks. All eight warmup profiles also completed and +cases, including the supplemental diagnostics, passed another 76 checks. +All eight warmup profiles also completed and committed on this image. -The plain C++ wrapper added about 85ms; the miss case was also slower in this -sample. Separate containers introduce timing variability, and larger CUDA -builds may hide the host-side savings entirely. +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).