Uh oh!
There was an error while loading. Please reload this page.
fix(cuda): route all 39 kernel launches through one checked helper - #154
Merged
Conversation
A CUDA kernel launch is asynchronous and does not throw. It reports failure by setting the runtime's error state, which somebody then has to inspect. Nothing here did: 39 `<<<` sites, zero `cudaGetLastError` and zero `cudaPeekAtLastError` in the whole repository (#152). A kernel that never ran was therefore indistinguishable from one that ran correctly -- the launcher returned normally and the caller read whatever was already in the output buffer. The fix is structural rather than conventional. `<<<` now appears in exactly one place in the tree, `impl/cuda/launch.h`, and all 39 sites go through the `FF_CUDA_LAUNCH` macro that wraps it. A convention applied at 39 call sites is a convention that will be forgotten at the 40th; #152 exists precisely because this one was never applied at any of them. The lint that makes the funnel unbypassable is the following commit. `ff::cuda::launchKernel` launches on the CALLER'S STREAM -- unchanged, no regression to the default stream -- then calls `cudaGetLastError()` and throws with the kernel name and the grid/block configuration. * No synchronisation. The post-launch check is a host-side read of a thread-local error word; it does not touch the device and does not wait for the kernel, so stream ordering and host/device overlap are exactly as before. * It therefore cannot see a fault raised while the kernel *executes*. Observing those needs a synchronisation point, which would serialise every launcher, so that is opt-in and off by default: FF_CUDA_LAUNCH_SYNC, a build flag and an environment variable of the same name (environment wins in both directions) -- this project's CUDA_LAUNCH_BLOCKING. * Misattribution is handled, not ignored. `cudaGetLastError` returns AND CLEARS the last error from any preceding CUDA call, so a stale unrelated failure would otherwise be blamed on this launch. The helper peeks first (`cudaPeekAtLastError`, which does not clear); if the state was already dirty with the same code the report says so explicitly. Draining at the dispatch boundary instead was rejected: that would discard an error belonging to the host application, which is entitled to see it. Peeking takes nothing away, and the failure path clears only because it immediately throws. * Sticky and non-sticky errors are distinguished in the message. A rejected launch leaves the context usable; an execution fault poisons it for the rest of the process. Codes on neither list are reported as unclassified rather than guessed at. * On failure the report also carries the kernel's `cudaFuncGetAttributes().maxThreadsPerBlock` and register count, and says in as many words when the block size exceeded it. That is the trigger #152 names -- CUDA_NUM_THREADS is 1024, the architectural maximum rather than a safe default -- turned from a silent wrong answer into a sentence naming the number to clamp to. Making the block size fit automatically is deliberately not done here; the reasons are recorded above CUDA_NUM_THREADS in utils.h. Two consequences of the launch now being able to throw, both required for the report to survive: * `distance_mesh::copy_faces` had no cleanup handler, because before this nothing between its `allocDevice` and its `return` could fail. It has one now, or the device buffer leaks on the new throw path. * Twelve `catch (const std::exception &exc) { ...; throw exc; }` sites copy-initialised a *std::exception* from the reference and threw that, discarding the dynamic type and the message. Four of them sit directly on a launcher's unwind path, where they would have replaced the whole diagnostic above with "std::exception". All twelve are now `throw;`, including the one in impl/cpu that has nothing to do with CUDA but is the same defect. No behaviour change on the success path: same kernels, same configurations, same stream, two host-side error-word reads per launch. Refs #152
The previous commit funnelled all 39 launches through one helper. On its own that is still a convention: nothing stops the fortieth launch being written the old way, and nothing here could notice if it were. There is no GPU in CI and there is not going to be one, so no test can execute a launch, and an unchecked one compiles, links and passes every gate this project has -- which is exactly how #152 lasted as long as it did. But it is a textual property, so it is checkable without hardware. `tools/check-cuda-launches.py` fails if: 1. `<<<` appears in code anywhere outside impl/cuda/launch.h; 2. launch.h does not contain exactly one `<<<`; 3. launch.h stops calling `cudaGetLastError` -- the funnel is still a funnel but no longer checks anything, which is #152 again with more steps; 4. `cudaLaunchKernel` / `cudaLaunchCooperativeKernel` is called outside the helper, which is how you launch a kernel without writing `<<<`. Comments and string literals are blanked before scanning, so prose may discuss `<<<` freely -- launch.h's own header comment does, six times. That is what lets rule 2 be stated about the helper rather than waived for it. `--selftest` is the other half and runs first in CI. A lint that has quietly stopped matching anything prints "clean" forever, so it asserts over synthetic inputs that the analyser still flags a bare launch, still ignores one inside a `//` comment, a `/* */` comment and a string literal, is not fooled by an escaped quote or by `//` inside a string, does not mistake `a << b` or `cout << x << y` for a launch, and still reports the right line number. Wired in as `lint (cuda launch sites)`, unconditional rather than path-filtered: the rule is about the whole tree and the cost is a Python startup. Follows tools/rename-macros.py and tools/test-baseline.sh -- a committed script with a self-verifying check mode. Refs #152
Where `<<<` is allowed to appear, what the helper does and does not check, why the post-launch check costs nothing, and that FF_CUDA_LAUNCH_SYNC is off by default and stays off. Also adds the new lint job to the CI summary, which previously listed codespell as the only unconditional gate. Refs #152
A blind spot in the blanker, found by asking what would make it stop
matching. A raw string's body may contain unescaped `"` and `\`, so the
ordinary quote scanner would take the first `"` inside `R"(...)"` as the
literal's end and then be one quote out of phase for the rest of the file
-- which does not produce a false alarm, it produces a false *pass*, and
every launch after that point goes unseen. That is the one failure mode
this checker exists to prevent, so it should not have it.
The tree contains no raw strings today; this is about the day somebody
adds one. `R` preceded by an identifier character is not a prefix, so
`MYR("a(b")` is still an ordinary string.
Three more --selftest cases cover it: a launch inside a raw string (not
counted), code after a raw string whose body holds a quote and a
backslash (still counted), and the identifier-ending-in-R case. 15 cases
now.
Refs #152Uh oh!
There was an error while loading. Please reload this page.
This was referenced Aug 20, 2026
Merged
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Fixes#152.
The defect
A CUDA kernel launch is asynchronous and does not throw. It reports failure by
setting the runtime's error state, and that state has to be read by somebody.
Nothing here read it: 39
<<<sites, zerocudaGetLastError, zerocudaPeekAtLastErrorin the whole repository. A kernel that never ran wastherefore indistinguishable from one that ran correctly — the launcher returned
normally and the caller read whatever was already in the output buffer.
The shape of the fix
Not a check macro at 39 call sites.
<<<now appears in exactly one place inthe tree —
include/fastfields/impl/cuda/launch.h— and all 39 sites reach itthrough
FF_CUDA_LAUNCH:ff::cuda::launchKernellaunches on the caller's stream — unchanged, noregression to the default stream — then calls
cudaGetLastError()and throwswith the kernel name and the grid/block configuration.
The macro exists only because
#KERNELrecords the kernel's spelling for thereport; the work is a function template, per this tree's "prefer an
inlinefunction to a macro" rule. nvcc lowers
ptr<<<…>>>(args)through the same hoststub as a direct launch, so the indirection is free and the argument list is
still type-checked against the kernel signature at the call.
This costs no synchronisation — please read this before worrying that it does
cudaGetLastError()is a host-side read of a thread-local error word thatthe driver wrote while enqueueing the launch. It does not touch the device, does
not wait for the kernel, and does not change stream ordering or host/device
overlap. Two such reads per launch (one peek, one get) against a launch that
costs microseconds.
The price of that is what it cannot see: a fault raised while the kernel
executes (illegal address, misaligned access, device-side assert). Observing
those requires a synchronisation point, which would serialise every launcher and
destroy the asynchrony the design depends on. So that is opt-in and off by
default:
The environment wins over the compile-time default in both directions, so a
release build can be asked for the diagnosis once without a rebuild. It is this
project's
CUDA_LAUNCH_BLOCKING, and it is never on by default.Misattribution: peek first, and say so
cudaGetLastError()returns and clears the last error from any precedingCUDA call — not just from a launch. So a stale, unrelated failure would
otherwise be blamed on whichever launch happened to look next.
What I chose:
cudaPeekAtLastError()before the launch (it does not clear).If the state was already dirty with the same code, the report says so in as many
words rather than asserting this kernel failed.
What I rejected, and why: draining at the dispatch boundary. fastfields is a
library inside somebody else's process — PyTorch, CuPy — and that error belongs
to them; they are entitled to see it on their own next check. Draining would
discard it. Peeking takes nothing away. The failure path does clear, but only
because it immediately converts the error into an exception: escalated, never
swallowed.
This is not theoretical, and it is already exercised — see the sample output
below, where an earlier
cudaMallocleft the state dirty and the helpercorrectly declined to blame the kernel.
Sticky vs non-sticky, in the message
The distinction is invisible in the error string and matters enormously:
untouched, the fix is on the caller's side;
process fails with the same code. That is "this process is dead", not "retry
with less".
launchErrorClass()classifies both lists explicitly and reports anything onneither as unclassified rather than guessing. A wrong claim about stickiness
is worse than no claim.
The trigger (
CUDA_NUM_THREADS = 1024) — reported now, fixed separately1024 is the architectural maximum, not a safe default. On failure the report
carries the kernel's
cudaFuncGetAttributes().maxThreadsPerBlockand registercount and says outright when the block size exceeded it, so the first person to
hit
cudaErrorLaunchOutOfResourcesgets the number to clamp to instead of amystery.
Making the block size fit automatically is deliberately not in this PR, and
the reasons are recorded above
CUDA_NUM_THREADSinutils.h:distance_euclidean::dtanddistance_mesh::sdtsize a per-lane scratchbuffer as
num_blocks * CUDA_NUM_THREADSon the host, while the kernelderives its lane stride from
blockDim.x * gridDim.xon the device. Shrinkingthe block alone is safe (the loops are grid-stride, and the buffer is merely
over-allocated); raising the grid to compensate — which is what
cudaOccupancyMaxPotentialBlockSizehands you — is an out-of-bounds devicewrite. Untangling that coupling is its own change with its own review.
cudaFuncGetAttributesneeds a device, soa clamp would ship completely unvalidated, whereas the reporting path is
validated by compile+link and by the message being exercisable on a
device-less host (below).
Per the issue's own ordering: reporting the failure is worth more than avoiding
it, and it is the half that can be done without hardware.
Exceptions at the call sites
The launchers can now throw where they previously could not, so I checked every
one.
noexceptfunction. Theonly
noexceptin the tree is onAbstractVector's accessors and the onlydestructors are
mesh_utils.h's virtual ones; none launches anything.distance_mesh::copy_faceshad no cleanup handler at all, because beforethis nothing between its
allocDeviceand itsreturncould fail. It has onenow, or the device buffer leaks on the new throw path.
catch (const std::exception &exc) { …; throw exc; }sites werecopy-initialising a
std::exceptionfrom the reference and throwing that —dynamic type and message both discarded. Four sit directly on a launcher's
unwind path, where they would have replaced the entire diagnostic below with
"std::exception". All twelve are nowthrow;, including the one inimpl/cpu/distance_euclidean.hthat has nothing to do with CUDA but is thesame defect.
The lint — the part that makes this stick
There is no GPU in CI, so nothing can execute a launch and no test can notice
the check being dropped: an unchecked launch compiles, links and passes every
gate this project has. That is exactly how #152 survived. But it is a textual
property, so
tools/check-cuda-launches.pyfails if:<<<appears in code anywhere outsidelaunch.h;launch.hdoes not contain exactly one<<<;launch.hstops callingcudaGetLastError— the funnel is still a funnelbut no longer checks anything, which is No CUDA launch error checking anywhere: 39 launches, zero cudaGetLastError — failures are silently discarded #152 again with more steps;
cudaLaunchKernel/cudaLaunchCooperativeKernelis called outside thehelper, which is how you launch without writing
<<<.Comments and string literals are blanked before scanning, so prose may discuss
<<<freely (launch.h's own header comment does, six times) — that is whatlets rule 2 be stated about the helper instead of waived for it.
--selftestruns first in CI and is the other half: a lint that has quietlystopped matching anything prints "clean" forever. It asserts the analyser still
flags a bare launch, still ignores one in a
//comment, a/* */comment anda string literal, is not fooled by an escaped quote or by
//inside a string,does not mistake
a << borcout << x << yfor a launch, and still reports theright line number.
Wired in as
lint (cuda launch sites)— unconditional rather than path-filtered(the rule is tree-wide and the cost is a Python startup), following
tools/rename-macros.py --checkandtools/test-baseline.sh --check.It also understands raw string literals, whose bodies may hold unescaped
"and
\: a scanner that mishandled one would go out of phase and stop seeinglaunches after it — a false pass, which is the only failure mode that
matters here. The tree has none today; that is the point of covering it now.
And with a launch reintroduced the old way:
What a failure looks like
Real output, from a binary built against this branch and run on a machine with
no GPU (so the launch fails with
cudaErrorNoDevice). It also happens todemonstrate the attribution logic: the preceding
cudaMallochad already failedand left the error state dirty, and the helper says so instead of blaming the
kernel.
Validation
CI on this branch (run 32391904003, at
25d6e1d; the latertools/-only commit re-runs the lint jobs only, by the path filter)lint (cuda launch sites)lint (codespell)test-cpu× 5 (clang-static,clang-dynamic,clang-cuda-default,clang-index64,gcc-static)test-cpu (sanitize asan+ubsan),test-cpu (tsan, grain=1)test-hubbuild-cuda (index64),build-cuda (index32)--no-undefinedandldd -rcompile-probe-cudalint (clang-format, changed lines)continue-on-error, and it cannot print its own findings (#89). Not chased.The gate, row-for-row
tools/test-baseline.sh --tree . --legs default,lib --checkagainst thedefaultandlibrows oftools/test-baseline.expected(the 13defaultrows sum to exactly 59,886 checks):
All 13 suites, 0 failures, every per-suite count unchanged. (
--checkcompareswhole reports and refuses a leg subset outright — "BASELINE NOT COMPARABLE" —
so the recorded file was filtered to the two legs first rather than the counts
being summed by hand.)
nvcc peak RSS is unchanged
From this run's own
build-cudaFFMEM tables, against the two columns recordedabove
MODULESinsrc/lib-cuda/Makefile(#143).reg_flow— the module the14 GiB budget exists for — lands on the recorded figure to 0.02%:
index32recordedindex32hereindex64recordedindex64herereg_flowreg_fieldreg_field_rlsresizereg_flow_rlspushpull_backwardpushpullrestrictdistancesplincposdefpeak nvcc RSS within budget (14680064 kB). The helper is host codeinstantiated once per launch-site signature, so it adds nothing to the
device-side instantiation matrix — which is what the 14 GiB
FF_MEM_BUDGET_KBgate protects. Reproduced locally on nvcc 12.0.140 (the version CI installs):
reg_flow5,969,432 kB local vs 5,970,612 kB in CI, a 0.02% difference.Explicitly not claimed: anything about runtime
Compile+link is the CUDA bar here. This change makes failures reportable; it
does not prove any launch succeeds. A GPU run would still be needed to
establish:
maxThreadsPerBlockat1024 threads. That is unknown, and it is the question No CUDA launch error checking anywhere: 39 launches, zero cudaGetLastError — failures are silently discarded #152 exists to make
answerable rather than the one it answers;
maxThreadsPerBlock/ register line of the report is produced for areal
cudaErrorLaunchOutOfResources—cudaFuncGetAttributesneeds a device,so that line is the one part of the message a device-less host cannot
exercise;
FF_CUDA_LAUNCH_SYNC=1does surface an execution fault (illegal address,device assert) that the default path is documented as unable to see;
in flight — i.e. that the
freeDevicecleanup paths, including the new one incopy_faces, actually free what they claim.One further gap, for the record: the fully-static pushpull order × bound matrix
(
.github/workflows/nightly-pushpull.yml, which has a CUDA leg) is scheduledand therefore never runs on a PR. It will first see
pushpull.h's convertedlaunch sites after merge, at 03:17 UTC.
Coordination
mainatb0463bc(after build: make the 32-bit index axis a per-backend option (FF_INDEX32), default unchanged #143, refactor: #pragma once in every header #145, refactor: spell the public interface <fastfields/...> #146, lib-cuda: record the measured cost of the 32-bit index axis on this backend #151). Followsrefactor: #pragma once in every header #145's
#pragma once-on-line-1 rule and refactor: spell the public interface <fastfields/...> #146's<fastfields/…>vs"…"rule;
tools/normalise-header-guards.py --check,tools/normalise-include-delimiters.py --checkandtools/rename-macros.py --checkare all clean.reg_flowper-(family, ndim) TUs): read before touchinganything. It moves
reg_flow.cpp's dispatch intoreg_flow_slice.inlbutdoes not touch
impl/cuda/reg_flow.h, which is where this PR's change toFF_REGFLOW_LAUNCH_*lives. The two should not collide; if perf(cuda): split reg_flow into per-(family, ndim) translation units #147 lands firstthe resolution is a re-apply of the same macro rewrite.
impl/kernels/vector/deletion) does not overlap.Reported, not fixed
CLAUDE.mdsaystools/normalise-include-delimiters.py --checkandtools/normalise-header-guards.py --checkenforce their conventions, andtools/rename-macros.py --checkis quoted as a gate in PR descriptions — butno workflow runs any of them. They all pass on
maintoday, so wiring theother three into the new
lintjob would be three lines; I left it out to keepthis PR to one subject. Worth doing.
Generated by Claude Code