From 71dd943a012868ef04a0172508e6bb6f3cc4cb68 Mon Sep 17 00:00:00 2001 From: fastfields agent Date: Thu, 20 Aug 2026 16:19:09 +0000 Subject: [PATCH 1/4] feat(cuda): route every kernel launch through one checked helper 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 --- .../fastfields/impl/cpu/distance_euclidean.h | 4 +- .../fastfields/impl/cuda/distance_euclidean.h | 12 +- include/fastfields/impl/cuda/distance_l1.h | 12 +- include/fastfields/impl/cuda/distance_mesh.h | 134 ++++--- include/fastfields/impl/cuda/launch.h | 374 ++++++++++++++++++ include/fastfields/impl/cuda/posdef.h | 72 ++-- include/fastfields/impl/cuda/pushpull.h | 127 +++--- include/fastfields/impl/cuda/reg_field.h | 73 ++-- include/fastfields/impl/cuda/reg_flow.h | 49 +-- include/fastfields/impl/cuda/resize.h | 10 +- include/fastfields/impl/cuda/restrict.h | 10 +- include/fastfields/impl/cuda/splinc.h | 7 +- include/fastfields/impl/cuda/utils.h | 32 +- 13 files changed, 678 insertions(+), 238 deletions(-) create mode 100644 include/fastfields/impl/cuda/launch.h diff --git a/include/fastfields/impl/cpu/distance_euclidean.h b/include/fastfields/impl/cpu/distance_euclidean.h index e9cf7b5..2a40211 100755 --- a/include/fastfields/impl/cpu/distance_euclidean.h +++ b/include/fastfields/impl/cpu/distance_euclidean.h @@ -39,12 +39,12 @@ dt( kernel(f + offset, v, z, d, w, n, s); } } - catch (const std::exception &exc) + catch (const std::exception &) { if (v) delete[] v; if (z) delete[] z; if (d) delete[] d; - throw exc; + throw; } delete[] v; delete[] z; diff --git a/include/fastfields/impl/cuda/distance_euclidean.h b/include/fastfields/impl/cuda/distance_euclidean.h index eefea71..8abf389 100755 --- a/include/fastfields/impl/cuda/distance_euclidean.h +++ b/include/fastfields/impl/cuda/distance_euclidean.h @@ -3,6 +3,7 @@ #include #include #include "utils.h" +#include "launch.h" // FF_CUDA_LAUNCH -- the only checked kernel launch #include #include @@ -76,14 +77,15 @@ FF_CUHOST void dt( buffer = allocDevice(buffer_size); size_device = copyToDeviceAsync(size, ndim, s); stride_device = copyToDeviceAsync(stride, ndim, s); - dt_kernel - <<>> - (ndim, f, buffer, w, size_device, stride_device); + FF_CUDA_LAUNCH( + (dt_kernel), + num_blocks, CUDA_NUM_THREADS, 0, s, + ndim, f, buffer, w, size_device, stride_device); } - catch (const std::exception &exc) + catch (const std::exception &) { freeDevice(buffer, size_device, stride_device); - throw exc; + throw; } freeDevice(buffer, size_device, stride_device); } diff --git a/include/fastfields/impl/cuda/distance_l1.h b/include/fastfields/impl/cuda/distance_l1.h index f2d46c6..bce89a0 100755 --- a/include/fastfields/impl/cuda/distance_l1.h +++ b/include/fastfields/impl/cuda/distance_l1.h @@ -3,6 +3,7 @@ #include #include #include "utils.h" +#include "launch.h" // FF_CUDA_LAUNCH -- the only checked kernel launch #include #include @@ -55,14 +56,15 @@ FF_CUHOST void dt( cudaStream_t s = (cudaStream_t)(std::intptr_t)stream; size_device = copyToDeviceAsync(size, ndim, s); stride_device = copyToDeviceAsync(stride, ndim, s); - dt_kernel - <<>> - (ndim, f, w, size_device, stride_device); + FF_CUDA_LAUNCH( + (dt_kernel), + GET_BLOCKS(batch_size), CUDA_NUM_THREADS, 0, s, + ndim, f, w, size_device, stride_device); } - catch (const std::exception &exc) + catch (const std::exception &) { freeDevice(size_device, stride_device); - throw exc; + throw; } freeDevice(size_device, stride_device); } diff --git a/include/fastfields/impl/cuda/distance_mesh.h b/include/fastfields/impl/cuda/distance_mesh.h index 04075c9..66114bc 100755 --- a/include/fastfields/impl/cuda/distance_mesh.h +++ b/include/fastfields/impl/cuda/distance_mesh.h @@ -4,6 +4,7 @@ #include #include #include "utils.h" +#include "launch.h" // FF_CUDA_LAUNCH -- the only checked kernel launch #include #include // std::unique_ptr #include // std::is_trivially_copyable @@ -245,15 +246,16 @@ scalar_t * copyTensorToContiguous( stride_out_copy = copyToDeviceAsync(stride_out, ndim, stream); stride_inp_copy = copyToDeviceAsync(stride_inp, ndim, stream); // Copy data - copy_tensor_kernel - <<>> - (ndim, out, inp, size_copy, stride_out_copy, stride_inp_copy); + FF_CUDA_LAUNCH( + (copy_tensor_kernel), + GET_BLOCKS(numel), CUDA_NUM_THREADS, 0, stream, + ndim, out, inp, size_copy, stride_out_copy, stride_inp_copy); } - catch (const std::exception & e) + catch (const std::exception &) { freeHost(stride_out); freeDevice(out, size_copy, stride_out_copy, stride_inp_copy); - throw e; + throw; } freeHost(stride_out); freeDevice(size_copy, stride_out_copy, stride_inp_copy); @@ -305,9 +307,21 @@ index_t * copy_faces( { offset_t stride0 = stride[0], stride1 = stride[1]; index_t * faces_out = allocDevice(nb_faces * ndim); - copy_faces_kernel - <<>> - (nb_faces, faces_out, faces, stride0, stride1); + // The launch can now throw, and `faces_out` is this function's to own + // until it returns -- before the launch was checked, nothing between the + // allocation and the return could fail, so there was no handler here. + try + { + FF_CUDA_LAUNCH( + (copy_faces_kernel), + GET_BLOCKS(nb_faces), CUDA_NUM_THREADS, 0, stream, + nb_faces, faces_out, faces, stride0, stride1); + } + catch (const std::exception &) + { + freeDevice(faces_out); + throw; + } return faces_out; } @@ -1142,33 +1156,33 @@ sdt( treetrace_device = allocDevice(stride_buf * treesize); // Compute SDT - sdt_kernel - <<>> - ( - nbatch, - dist, - nearest_vertex, - coord, - verts_device, - faces_device, - tree_device, - treetrace_device, - treesize, - normfaces_device, - normverts_device, - normedges_device, - size_device, - stride_dist_device, - stride_nearest_device, - stride_coord_device, - stride_vec_device, - stride_vec_device, - stride_vec_device, - stride_vec_device, - stride_mat_device - ); + FF_CUDA_LAUNCH( + (sdt_kernel), + num_blocks, CUDA_NUM_THREADS, 0, s, + nbatch, + dist, + nearest_vertex, + coord, + verts_device, + faces_device, + tree_device, + treetrace_device, + treesize, + normfaces_device, + normverts_device, + normedges_device, + size_device, + stride_dist_device, + stride_nearest_device, + stride_coord_device, + stride_vec_device, + stride_vec_device, + stride_vec_device, + stride_vec_device, + stride_mat_device + ); } - catch (const std::exception & e) + catch (const std::exception &) { freeDevice( faces_device, @@ -1194,7 +1208,7 @@ sdt( normverts_host, normedges_host ); - throw e; + throw; } freeDevice( @@ -1355,31 +1369,31 @@ sdt_naive( int num_blocks = GET_BLOCKS(numel); // Compute SDT - sdt_naive_kernel - <<>> - ( - nbatch, - dist, - nearest_vertex, - coord, - verts_device, - faces_device, - normfaces_device, - normverts_device, - normedges_device, - size_device, - nb_faces, - stride_dist_device, - stride_nearest_device, - stride_coord_device, - stride_vec_device, - stride_vec_device, - stride_vec_device, - stride_vec_device, - stride_mat_device - ); + FF_CUDA_LAUNCH( + (sdt_naive_kernel), + num_blocks, CUDA_NUM_THREADS, 0, s, + nbatch, + dist, + nearest_vertex, + coord, + verts_device, + faces_device, + normfaces_device, + normverts_device, + normedges_device, + size_device, + nb_faces, + stride_dist_device, + stride_nearest_device, + stride_coord_device, + stride_vec_device, + stride_vec_device, + stride_vec_device, + stride_vec_device, + stride_mat_device + ); } - catch (const std::exception & e) + catch (const std::exception &) { freeDevice( faces_device, @@ -1401,7 +1415,7 @@ sdt_naive( normverts_host, normedges_host ); - throw e; + throw; } freeDevice( diff --git a/include/fastfields/impl/cuda/launch.h b/include/fastfields/impl/cuda/launch.h new file mode 100644 index 0000000..1e9a303 --- /dev/null +++ b/include/fastfields/impl/cuda/launch.h @@ -0,0 +1,374 @@ +#pragma once +#include +#include // std::snprintf +#include // std::getenv +#include // std::strcmp +#include // std::runtime_error +#include // std::string +#include // std::forward + +/*********************************************************************** + * * + * THE ONE CHECKED KERNEL LAUNCH * + * * + *********************************************************************** + * + * A CUDA kernel launch is asynchronous and does not throw. It reports + * failure by setting the runtime's error state, which somebody has to go + * and look at. Nothing in this tree ever did: there were 39 `<<<` sites + * and zero `cudaGetLastError` / `cudaPeekAtLastError` calls, so a kernel + * that never ran was indistinguishable from one that ran correctly -- + * the launcher returned normally and the caller read whatever was + * already in the output buffer (fastfields-lib#152). + * + * The fix is structural rather than conventional. `<<<` appears in + * exactly ONE place in this repository -- `launchKernel` below -- and + * every launcher goes through the `FF_CUDA_LAUNCH` macro that wraps it. + * `tools/check-cuda-launches.py --check` fails the build if a `<<<` + * shows up anywhere else, which is what makes this survive contact with + * future patches: a check nobody *can* bypass beats one people have to + * remember. (The same reasoning as `-Wl,--no-undefined` in + * make/common.mk and the `FF_MEM_BUDGET_KB` gate in CI; #152 exists + * precisely because "remember to check" was never applied even once.) + * + * WHAT THIS COSTS AT RUNTIME: nothing measurable, and in particular no + * device synchronisation. `cudaGetLastError` is a host-side read of a + * thread-local error word that the driver wrote while enqueueing the + * launch; it does not touch the device and does not wait for the kernel. + * The stream ordering of the calling code is completely unchanged. + * + * WHAT IT THEREFORE CANNOT SEE: errors raised by the kernel while it + * *executes* (illegal address, misaligned access, device-side assert). + * Those only become observable at a synchronisation point, and + * synchronising after every launch would serialise the pipeline and + * destroy exactly the asynchrony the design depends on. So that is an + * opt-in debug mode, off by default -- see `FF_CUDA_LAUNCH_SYNC` below, + * which is this project's `CUDA_LAUNCH_BLOCKING`. + */ + +FF_NAMESPACE_BEGIN(FF_NS) +FF_NAMESPACE_BEGIN(FF_DEVICE) + +/*********************************************************************** + * OPT-IN POST-LAUNCH SYNCHRONISATION * + ***********************************************************************/ + +// Compile-time default for the debug sync described above. It is 0, and it +// must stay 0: turning it on inserts a `cudaStreamSynchronize` after every +// single launch, which makes every launcher blocking and removes any overlap +// between the host and the device. It exists so that a debug build can be +// produced with `-DFF_CUDA_LAUNCH_SYNC=1` without editing sources. +#ifndef FF_CUDA_LAUNCH_SYNC +# define FF_CUDA_LAUNCH_SYNC 0 +#endif + +// Runtime override, read from the environment variable of the same name. +// "0", "false" and the empty value mean off; anything else means on. 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: +// +// FF_CUDA_LAUNCH_SYNC=1 python -c "import fastfields; ..." +// +// This is deliberately the same shape as the CUDA runtime's own +// CUDA_LAUNCH_BLOCKING, which is the knob a CUDA programmer already reaches +// for. The two compose: CUDA_LAUNCH_BLOCKING makes the launch itself +// synchronous, this makes the *error check* synchronous. +FF_CUHOST inline bool _launchSyncFromEnv() +{ + const char * v = std::getenv("FF_CUDA_LAUNCH_SYNC"); + if (!v || !*v) return FF_CUDA_LAUNCH_SYNC != 0; + if (std::strcmp(v, "0") == 0) return false; + if (std::strcmp(v, "false") == 0) return false; + return true; +} + +// C++11 guarantees the initialisation of a function-local static is run once +// and is thread-safe, so the getenv is paid once per process rather than once +// per launch. (Namespace-scope state in a header is what fastfields-kernels' +// threadpool bug was, hence the accessor rather than a variable.) +FF_CUHOST inline bool launchSyncEnabled() +{ + static const bool enabled = _launchSyncFromEnv(); + return enabled; +} + +/*********************************************************************** + * STICKY vs. NON-STICKY CLASSIFICATION * + ***********************************************************************/ + +// The difference matters enormously to whoever reads the report, and it is +// not visible from the error string: +// +// * a launch-CONFIGURATION error is raised before the kernel starts. The +// context is untouched, the next launch can succeed, and the fix is on +// the caller's side ("your block size is too big for this kernel"). +// * an EXECUTION error poisons the CUDA context. Every subsequent CUDA +// call in the process -- including in code that has nothing to do with +// fastfields -- returns the same error until the process exits. The +// honest report is "this process is dead", not "retry with less". +// +// Anything not on either list is reported as unclassified rather than +// guessed at; a wrong claim about stickiness is worse than no claim. +enum class LaunchErrorClass { Configuration, Execution, Unclassified }; + +FF_CUHOST inline LaunchErrorClass launchErrorClass(cudaError_t err) +{ + switch (err) + { + // --- rejected before the kernel ran: recoverable ---------------- + case cudaErrorInvalidConfiguration: // grid/block out of range + case cudaErrorLaunchOutOfResources: // registers or shared mem + case cudaErrorInvalidDeviceFunction: // no such kernel on device + case cudaErrorNoKernelImageForDevice: // no cubin for this arch + case cudaErrorInvalidValue: + case cudaErrorInvalidDevice: + case cudaErrorInvalidPtx: + case cudaErrorUnsupportedPtxVersion: + // Not launch *configuration* so much as "there was nothing to launch + // on", but they belong on this side of the line all the same: no + // kernel ran, so no kernel poisoned anything. + case cudaErrorNoDevice: + case cudaErrorInsufficientDriver: + return LaunchErrorClass::Configuration; + + // --- raised by the running kernel: sticky ------------------------ + case cudaErrorLaunchFailure: + case cudaErrorIllegalAddress: + case cudaErrorMisalignedAddress: + case cudaErrorIllegalInstruction: + case cudaErrorInvalidAddressSpace: + case cudaErrorInvalidPc: + case cudaErrorHardwareStackError: + case cudaErrorLaunchTimeout: + case cudaErrorAssert: + case cudaErrorECCUncorrectable: + return LaunchErrorClass::Execution; + + default: + return LaunchErrorClass::Unclassified; + } +} + +/*********************************************************************** + * THE ERROR REPORT * + ***********************************************************************/ + +// Deliberately NOT a template: it is called from the cold path of a function +// template that is instantiated once per (kernel signature, dtype, offset +// type, ndim, spline, bound) combination, and none of this text needs to be +// duplicated across those. Everything kernel-specific arrives as an argument. +// +// `max_threads` / `num_regs` are `cudaFuncGetAttributes` results, or -1 when +// the query itself failed. They are what turns "too many resources requested +// for launch" into an actionable sentence. +FF_CUHOST inline void _throwLaunchError( + const char * name, + dim3 grid, + dim3 block, + size_t shared_mem, + cudaStream_t stream, + cudaError_t err, + bool preexisting, + bool from_sync, + int max_threads, + int num_regs) +{ + // `#KERNEL` stringises a parenthesised template-id (see FF_CUDA_LAUNCH), + // so strip the wrapping parentheses for readability. + std::string kernel(name ? name : ""); + if (kernel.size() >= 2 && kernel.front() == '(' && kernel.back() == ')') + kernel = kernel.substr(1, kernel.size() - 2); + + const char * classification; + switch (launchErrorClass(err)) + { + case LaunchErrorClass::Configuration: + classification = + "the launch was REJECTED before the kernel started, so no " + "kernel has poisoned the CUDA context. The fix is in the " + "launch configuration, the inputs, or the environment"; + break; + case LaunchErrorClass::Execution: + classification = + "this is an EXECUTION error raised by a running kernel. It is " + "STICKY: the CUDA context is now unusable and every subsequent " + "CUDA call in this process will fail with the same error"; + break; + default: + classification = + "this error is not one that fastfields classifies, so whether " + "the CUDA context survived it is unknown"; + break; + } + + char buf[2048]; + int n = std::snprintf( + buf, sizeof(buf), + "ff::cuda: CUDA kernel launch failed.\n" + " kernel : %s\n" + " configuration : grid=(%u,%u,%u) block=(%u,%u,%u) " + "shared=%llu bytes stream=%p -- %llu threads total\n" + " error : %s (%d): %s\n" + " meaning : %s.\n" + " observed : %s", + kernel.c_str(), + grid.x, grid.y, grid.z, block.x, block.y, block.z, + static_cast(shared_mem), + static_cast(stream), + static_cast(grid.x) * grid.y * grid.z * + block.x * block.y * block.z, + cudaGetErrorName(err), static_cast(err), cudaGetErrorString(err), + classification, + from_sync + ? "on synchronising the stream after the launch " + "(FF_CUDA_LAUNCH_SYNC is on), so it was raised by the kernel " + "itself rather than by the launch configuration" + : "immediately after the launch, without synchronising"); + + // The register budget, when we could get it and it is the thing that went + // wrong. `CUDA_NUM_THREADS` is 1024 -- the *architectural maximum*, not a + // safe default -- so a register-heavy instantiation whose + // maxThreadsPerBlock is below that fails here, and this line says so in + // as many words rather than leaving it to be rediscovered with a + // debugger. See fastfields-lib#152 and `GET_BLOCKS`'s block-size + // parameter in utils.h. + if (n > 0 && n < static_cast(sizeof(buf)) && max_threads > 0) + { + const bool over = static_cast(block.x * block.y * block.z) + > max_threads; + std::snprintf( + buf + n, sizeof(buf) - static_cast(n), + "\n this kernel : maxThreadsPerBlock=%d, %d registers/thread" + "%s", + max_threads, num_regs, + over ? " -- THE LAUNCH ASKED FOR MORE THREADS PER BLOCK THAN THIS " + "KERNEL CAN TAKE. Its register usage is the limit. Read the " + "note above CUDA_NUM_THREADS in impl/cuda/utils.h before " + "changing the block size: two launchers size a device " + "scratch buffer from the launch geometry, so the grid may " + "not simply be scaled up to compensate." + : ""); + } + + if (preexisting) + { + // See the note at the peek in `launchKernel`. Say it out loud rather + // than quietly attributing somebody else's failure to this launch. + const size_t len = std::strlen(buf); + std::snprintf( + buf + len, sizeof(buf) - len, + "\n ATTRIBUTION : the CUDA error state was ALREADY set with " + "this same code before this launch was issued, so it may have " + "been raised by an earlier CUDA call outside fastfields. It is " + "reported here because this is the first place in the process " + "that inspects it."); + } + + throw std::runtime_error(buf); +} + +/*********************************************************************** + * THE LAUNCH * + ***********************************************************************/ + +// The only `<<<` in the tree. Launches `kernel` on the CALLER'S STREAM with +// the given configuration, then checks the error state and throws on failure. +// +// `kernel` is a pointer to a __global__ function; nvcc lowers `ptr<<<...>>>()` +// through the same host stub as a direct launch, so routing every site through +// this template costs nothing and gains a compile-time check that the argument +// list matches the kernel signature. +template +FF_CUHOST inline void launchKernel( + const char * name, + Kernel kernel, + dim3 grid, + dim3 block, + size_t shared_mem, + cudaStream_t stream, + Args &&... args) +{ + // MISATTRIBUTION, handled rather than ignored. + // + // `cudaGetLastError` returns *and clears* the last error from any + // preceding CUDA call in this thread -- not just from a launch. So if + // something entirely unrelated failed earlier and nobody looked, the + // check below would find that error sitting there and blame this launch + // for it. `cudaPeekAtLastError` reads the same word WITHOUT clearing it, + // so we can tell the two apart: if the state was already dirty with the + // same code, the report says so explicitly instead of asserting that this + // kernel is what failed. + // + // Why observe here rather than draining at the dispatch boundary + // (src/lib-cuda/*.cpp): draining would *discard* an error belonging to + // the host application -- PyTorch, CuPy -- which is entitled to see it on + // its own next check. fastfields is a library in someone else's process + // and does not own that state. Peeking takes nothing away. We do clear on + // the failure path below, but only because we immediately convert it into + // an exception: the error is escalated, never swallowed. + const cudaError_t prior = cudaPeekAtLastError(); + + kernel<<>>(std::forward(args)...); + + // Host-side only. No `cudaDeviceSynchronize`, no `cudaStreamSynchronize`, + // nothing that waits on the device -- see the header comment. + const cudaError_t err = cudaGetLastError(); + if (err != cudaSuccess) + { + int max_threads = -1, num_regs = -1; + cudaFuncAttributes attr; + if (cudaFuncGetAttributes(&attr, kernel) == cudaSuccess) + { + max_threads = attr.maxThreadsPerBlock; + num_regs = attr.numRegs; + } + // The query above sets the error state if it failed; drop that so the + // state we leave behind reflects the launch, not our own diagnosis. + cudaGetLastError(); + + _throwLaunchError(name, grid, block, shared_mem, stream, err, + /*preexisting=*/err == prior, /*from_sync=*/false, + max_threads, num_regs); + } + + // Opt-in, off by default, and never on by default: this serialises the + // caller against the device on every launch. It is the only way to + // observe a fault raised while the kernel runs, which is a different and + // strictly larger class of failure than the check above can reach. + if (launchSyncEnabled()) + { + const cudaError_t serr = cudaStreamSynchronize(stream); + if (serr != cudaSuccess) + { + cudaGetLastError(); // drain: we are about to report it + _throwLaunchError(name, grid, block, shared_mem, stream, serr, + /*preexisting=*/false, /*from_sync=*/true, + -1, -1); + } + } +} + +FF_NAMESPACE_END(FF_DEVICE) +FF_NAMESPACE_END(FF_NS) + +/*********************************************************************** + * THE CALL SITE FORM * + ***********************************************************************/ + +// Replaces +// kernel<<>>(x, y, z); +// with +// FF_CUDA_LAUNCH((kernel), grid, block, shmem, stream, x, y, z); +// +// The macro exists for exactly one reason: `#KERNEL` records the kernel's +// spelling so the failure report can name it. Everything else is the function +// template above, per this project's "prefer an inline function to a macro" +// rule -- but a function cannot stringise its own argument. +// +// PARENTHESES ARE REQUIRED around a template-id: the commas inside +// `kernel` would otherwise split it across macro parameters. A +// parenthesised function name still decays to a function pointer, and +// `_throwLaunchError` strips the parentheses back out of the printed name. +#define FF_CUDA_LAUNCH(KERNEL, GRID, BLOCK, SHMEM, STREAM, ...) \ + ::FF_NS::FF_DEVICE::launchKernel( \ + #KERNEL, KERNEL, (GRID), (BLOCK), (SHMEM), (STREAM), __VA_ARGS__) diff --git a/include/fastfields/impl/cuda/posdef.h b/include/fastfields/impl/cuda/posdef.h index 5ca2d48..d612360 100755 --- a/include/fastfields/impl/cuda/posdef.h +++ b/include/fastfields/impl/cuda/posdef.h @@ -3,6 +3,7 @@ #include #include #include "utils.h" +#include "launch.h" // FF_CUDA_LAUNCH -- the only checked kernel launch #include #include @@ -330,10 +331,11 @@ void sym_matvec( d_sh = copyToDevice(stride_hes, ndim); d_si = copyToDevice(stride_inp, ndim); FF_POSDEF_PROLOGUE; -# define FF_LAUNCH(NB, CC) \ - sym_matvec_k \ - <<>>( \ - out, hes, inp, d_size, d_so, d_sh, d_si) +# define FF_LAUNCH(NB, CC) \ + FF_CUDA_LAUNCH( \ + (sym_matvec_k), \ + blocks, threads, 0, s, \ + out, hes, inp, d_size, d_so, d_sh, d_si) FF_POSDEF_NC_SWITCH; # undef FF_LAUNCH } catch (...) { freeDevice(d_size, d_so, d_sh, d_si); throw; } @@ -362,10 +364,11 @@ void sym_matvec_backward( d_sg = copyToDevice(stride_grd, ndim); d_si = copyToDevice(stride_inp, ndim); FF_POSDEF_PROLOGUE; -# define FF_LAUNCH(NB, CC) \ - sym_matvec_backward_k \ - <<>>( \ - out, grd, inp, d_size, d_so, d_sg, d_si) +# define FF_LAUNCH(NB, CC) \ + FF_CUDA_LAUNCH( \ + (sym_matvec_backward_k), \ + blocks, threads, 0, s, \ + out, grd, inp, d_size, d_so, d_sg, d_si) FF_POSDEF_NC_SWITCH; # undef FF_LAUNCH } catch (...) { freeDevice(d_size, d_so, d_sg, d_si); throw; } @@ -394,10 +397,11 @@ void sym_addmatvec_( d_sh = copyToDevice(stride_hes, ndim); d_si = copyToDevice(stride_inp, ndim); FF_POSDEF_PROLOGUE; -# define FF_LAUNCH(NB, CC) \ - sym_addmatvec__k \ - <<>>( \ - out, hes, inp, d_size, d_so, d_sh, d_si) +# define FF_LAUNCH(NB, CC) \ + FF_CUDA_LAUNCH( \ + (sym_addmatvec__k), \ + blocks, threads, 0, s, \ + out, hes, inp, d_size, d_so, d_sh, d_si) FF_POSDEF_NC_SWITCH; # undef FF_LAUNCH } catch (...) { freeDevice(d_size, d_so, d_sh, d_si); throw; } @@ -426,10 +430,11 @@ void sym_submatvec_( d_sh = copyToDevice(stride_hes, ndim); d_si = copyToDevice(stride_inp, ndim); FF_POSDEF_PROLOGUE; -# define FF_LAUNCH(NB, CC) \ - sym_submatvec__k \ - <<>>( \ - out, hes, inp, d_size, d_so, d_sh, d_si) +# define FF_LAUNCH(NB, CC) \ + FF_CUDA_LAUNCH( \ + (sym_submatvec__k), \ + blocks, threads, 0, s, \ + out, hes, inp, d_size, d_so, d_sh, d_si) FF_POSDEF_NC_SWITCH; # undef FF_LAUNCH } catch (...) { freeDevice(d_size, d_so, d_sh, d_si); throw; } @@ -462,10 +467,11 @@ void sym_solve( d_sh = copyToDevice(stride_hes, ndim); d_sw = stride_wgt ? copyToDevice(stride_wgt, ndim) : nullptr; FF_POSDEF_PROLOGUE; -# define FF_LAUNCH(NB, CC) \ - sym_solve_k \ - <<>>( \ - out, inp, hes, wgt, d_size, d_so, d_si, d_sh, d_sw) +# define FF_LAUNCH(NB, CC) \ + FF_CUDA_LAUNCH( \ + (sym_solve_k), \ + blocks, threads, 0, s, \ + out, inp, hes, wgt, d_size, d_so, d_si, d_sh, d_sw) FF_POSDEF_NC_SWITCH; # undef FF_LAUNCH } catch (...) { freeDevice(d_size, d_so, d_si, d_sh, d_sw); throw; } @@ -494,10 +500,11 @@ void sym_solve_( d_sh = copyToDevice(stride_hes, ndim); d_sw = stride_wgt ? copyToDevice(stride_wgt, ndim) : nullptr; FF_POSDEF_PROLOGUE; -# define FF_LAUNCH(NB, CC) \ - sym_solve__k \ - <<>>( \ - out, hes, wgt, d_size, d_so, d_sh, d_sw) +# define FF_LAUNCH(NB, CC) \ + FF_CUDA_LAUNCH( \ + (sym_solve__k), \ + blocks, threads, 0, s, \ + out, hes, wgt, d_size, d_so, d_sh, d_sw) FF_POSDEF_NC_SWITCH; # undef FF_LAUNCH } catch (...) { freeDevice(d_size, d_so, d_sh, d_sw); throw; } @@ -523,10 +530,11 @@ void sym_invert( d_so = copyToDevice(stride_out, ndim); d_sh = copyToDevice(stride_hes, ndim); FF_POSDEF_PROLOGUE; -# define FF_LAUNCH(NB, CC) \ - sym_invert_k \ - <<>>( \ - out, hes, d_size, d_so, d_sh) +# define FF_LAUNCH(NB, CC) \ + FF_CUDA_LAUNCH( \ + (sym_invert_k), \ + blocks, threads, 0, s, \ + out, hes, d_size, d_so, d_sh) FF_POSDEF_NC_SWITCH; # undef FF_LAUNCH } catch (...) { freeDevice(d_size, d_so, d_sh); throw; } @@ -549,9 +557,11 @@ void sym_invert_( d_size = copyToDevice(size, ndim); d_st = copyToDevice(stride, ndim); FF_POSDEF_PROLOGUE; -# define FF_LAUNCH(NB, CC) \ - sym_invert__k \ - <<>>(hes, d_size, d_st) +# define FF_LAUNCH(NB, CC) \ + FF_CUDA_LAUNCH( \ + (sym_invert__k), \ + blocks, threads, 0, s, \ + hes, d_size, d_st) FF_POSDEF_NC_SWITCH; # undef FF_LAUNCH } catch (...) { freeDevice(d_size, d_st); throw; } diff --git a/include/fastfields/impl/cuda/pushpull.h b/include/fastfields/impl/cuda/pushpull.h index 9267001..3e10653 100755 --- a/include/fastfields/impl/cuda/pushpull.h +++ b/include/fastfields/impl/cuda/pushpull.h @@ -5,6 +5,7 @@ #include #include #include "utils.h" // allocDevice / copyToDevice / freeDevice / GET_BLOCKS +#include "launch.h" // FF_CUDA_LAUNCH -- the only checked kernel launch #include // std::intptr_t #include // std::logic_error @@ -754,18 +755,19 @@ FF_CUHOST void pull( d_si = copyToDevice(stride_inp, n1); d_sgr = copyToDevice(stride_grid, n1); -#define FF_PP_PULL(NB) \ - pull \ - <<>> \ - (bnd, spl, extrapolate, out, inp, grid, d_sg, d_ss, d_so, d_si, d_sgr) +#define FF_PP_PULL(NB) \ + FF_CUDA_LAUNCH( \ + (pull), \ + GET_BLOCKS(numel), CUDA_NUM_THREADS, 0, cstream, \ + bnd, spl, extrapolate, out, inp, grid, d_sg, d_ss, d_so, d_si, d_sgr) FF_PP_DISPATCH(FF_PP_PULL); #undef FF_PP_PULL } - catch (const std::exception &exc) + catch (const std::exception &) { freeDevice(d_sg, d_ss, d_so, d_si, d_sgr); - throw exc; + throw; } freeDevice(d_sg, d_ss, d_so, d_si, d_sgr); } @@ -806,18 +808,19 @@ FF_CUHOST void push( d_si = copyToDevice(stride_inp, n1); d_sgr = copyToDevice(stride_grid, n1); -#define FF_PP_PUSH(NB) \ - push \ - <<>> \ - (bnd, spl, extrapolate, out, inp, grid, d_sg, d_ss, d_so, d_si, d_sgr) +#define FF_PP_PUSH(NB) \ + FF_CUDA_LAUNCH( \ + (push), \ + GET_BLOCKS(numel), CUDA_NUM_THREADS, 0, cstream, \ + bnd, spl, extrapolate, out, inp, grid, d_sg, d_ss, d_so, d_si, d_sgr) FF_PP_DISPATCH(FF_PP_PUSH); #undef FF_PP_PUSH } - catch (const std::exception &exc) + catch (const std::exception &) { freeDevice(d_sg, d_ss, d_so, d_si, d_sgr); - throw exc; + throw; } freeDevice(d_sg, d_ss, d_so, d_si, d_sgr); } @@ -855,18 +858,19 @@ FF_CUHOST void count( d_so = copyToDevice(stride_out, n1); d_sgr = copyToDevice(stride_grid, n1); -#define FF_PP_COUNT(NB) \ - count \ - <<>> \ - (bnd, spl, extrapolate, out, grid, d_sg, d_ss, d_so, d_sgr) +#define FF_PP_COUNT(NB) \ + FF_CUDA_LAUNCH( \ + (count), \ + GET_BLOCKS(numel), CUDA_NUM_THREADS, 0, cstream, \ + bnd, spl, extrapolate, out, grid, d_sg, d_ss, d_so, d_sgr) FF_PP_DISPATCH(FF_PP_COUNT); #undef FF_PP_COUNT } - catch (const std::exception &exc) + catch (const std::exception &) { freeDevice(d_sg, d_ss, d_so, d_sgr); - throw exc; + throw; } freeDevice(d_sg, d_ss, d_so, d_sgr); } @@ -907,18 +911,19 @@ FF_CUHOST void grad( d_si = copyToDevice(stride_inp, n1); d_sgr = copyToDevice(stride_grid, n1); -#define FF_PP_GRAD(NB) \ - grad \ - <<>> \ - (bnd, spl, extrapolate, out, inp, grid, d_sg, d_ss, d_so, d_si, d_sgr) +#define FF_PP_GRAD(NB) \ + FF_CUDA_LAUNCH( \ + (grad), \ + GET_BLOCKS(numel), CUDA_NUM_THREADS, 0, cstream, \ + bnd, spl, extrapolate, out, inp, grid, d_sg, d_ss, d_so, d_si, d_sgr) FF_PP_DISPATCH(FF_PP_GRAD); #undef FF_PP_GRAD } - catch (const std::exception &exc) + catch (const std::exception &) { freeDevice(d_sg, d_ss, d_so, d_si, d_sgr); - throw exc; + throw; } freeDevice(d_sg, d_ss, d_so, d_si, d_sgr); } @@ -971,19 +976,20 @@ FF_CUHOST void pull_backward( d_sgi = copyToDevice(stride_ginp, n1); d_sgr = copyToDevice(stride_grid, n1); -#define FF_PP_PULLB(NB) \ - pull_backward \ - <<>> \ - (bnd, spl, extrapolate, out, gout, inp, ginp, grid, \ - d_sg, d_ss, d_so, d_sgo, d_si, d_sgi, d_sgr) +#define FF_PP_PULLB(NB) \ + FF_CUDA_LAUNCH( \ + (pull_backward), \ + GET_BLOCKS(numel), CUDA_NUM_THREADS, 0, cstream, \ + bnd, spl, extrapolate, out, gout, inp, ginp, grid, \ + d_sg, d_ss, d_so, d_sgo, d_si, d_sgi, d_sgr) FF_PP_DISPATCH(FF_PP_PULLB); #undef FF_PP_PULLB } - catch (const std::exception &exc) + catch (const std::exception &) { freeDevice(d_sg, d_ss, d_so, d_sgo, d_si, d_sgi, d_sgr); - throw exc; + throw; } freeDevice(d_sg, d_ss, d_so, d_sgo, d_si, d_sgi, d_sgr); } @@ -1031,19 +1037,20 @@ FF_CUHOST void push_backward( d_sgi = copyToDevice(stride_ginp, n1); d_sgr = copyToDevice(stride_grid, n1); -#define FF_PP_PUSHB(NB) \ - push_backward \ - <<>> \ - (bnd, spl, extrapolate, out, gout, inp, ginp, grid, \ - d_sg, d_ss, d_so, d_sgo, d_si, d_sgi, d_sgr) +#define FF_PP_PUSHB(NB) \ + FF_CUDA_LAUNCH( \ + (push_backward), \ + GET_BLOCKS(numel), CUDA_NUM_THREADS, 0, cstream, \ + bnd, spl, extrapolate, out, gout, inp, ginp, grid, \ + d_sg, d_ss, d_so, d_sgo, d_si, d_sgi, d_sgr) FF_PP_DISPATCH(FF_PP_PUSHB); #undef FF_PP_PUSHB } - catch (const std::exception &exc) + catch (const std::exception &) { freeDevice(d_sg, d_ss, d_so, d_sgo, d_si, d_sgi, d_sgr); - throw exc; + throw; } freeDevice(d_sg, d_ss, d_so, d_sgo, d_si, d_sgi, d_sgr); } @@ -1085,18 +1092,19 @@ FF_CUHOST void count_backward( d_sgr = copyToDevice(stride_grid, n1); #define FF_PP_COUNTB(NB) \ - count_backward \ - <<>> \ - (bnd, spl, extrapolate, gout, ginp, grid, \ - d_sg, d_ss, d_sgo, d_sgi, d_sgr) + FF_CUDA_LAUNCH( \ + (count_backward), \ + GET_BLOCKS(numel), CUDA_NUM_THREADS, 0, cstream, \ + bnd, spl, extrapolate, gout, ginp, grid, \ + d_sg, d_ss, d_sgo, d_sgi, d_sgr) FF_PP_DISPATCH(FF_PP_COUNTB); #undef FF_PP_COUNTB } - catch (const std::exception &exc) + catch (const std::exception &) { freeDevice(d_sg, d_ss, d_sgo, d_sgi, d_sgr); - throw exc; + throw; } freeDevice(d_sg, d_ss, d_sgo, d_sgi, d_sgr); } @@ -1144,19 +1152,20 @@ FF_CUHOST void grad_backward( d_sgi = copyToDevice(stride_ginp, n1 + 1); // extra (D) axis d_sgr = copyToDevice(stride_grid, n1); -#define FF_PP_GRADB(NB) \ - grad_backward \ - <<>> \ - (bnd, spl, extrapolate, out, gout, inp, ginp, grid, \ - d_sg, d_ss, d_so, d_sgo, d_si, d_sgi, d_sgr) +#define FF_PP_GRADB(NB) \ + FF_CUDA_LAUNCH( \ + (grad_backward), \ + GET_BLOCKS(numel), CUDA_NUM_THREADS, 0, cstream, \ + bnd, spl, extrapolate, out, gout, inp, ginp, grid, \ + d_sg, d_ss, d_so, d_sgo, d_si, d_sgi, d_sgr) FF_PP_DISPATCH(FF_PP_GRADB); #undef FF_PP_GRADB } - catch (const std::exception &exc) + catch (const std::exception &) { freeDevice(d_sg, d_ss, d_so, d_sgo, d_si, d_sgi, d_sgr); - throw exc; + throw; } freeDevice(d_sg, d_ss, d_so, d_sgo, d_si, d_sgi, d_sgr); } diff --git a/include/fastfields/impl/cuda/reg_field.h b/include/fastfields/impl/cuda/reg_field.h index d59f44e..8a0501d 100755 --- a/include/fastfields/impl/cuda/reg_field.h +++ b/include/fastfields/impl/cuda/reg_field.h @@ -6,6 +6,7 @@ #include #include #include "utils.h" // allocDevice / copyToDevice / freeDevice / GET_BLOCKS +#include "launch.h" // FF_CUDA_LAUNCH -- the only checked kernel launch #include // std::logic_error using namespace std; @@ -1719,27 +1720,27 @@ void relax_bending_jrls_( // explicit template-argument lists select unambiguously in both directions. // Inner dispatch: given a compile-time NB (nbatch), pick the channel count. -#define FF_REGFIELD_LAUNCH_C(KERN, NB, ...) \ - switch (nc) { \ - case 1: KERN \ - <<>>(bnd, __VA_ARGS__); break; \ - case 2: KERN \ - <<>>(bnd, __VA_ARGS__); break; \ - case 3: KERN \ - <<>>(bnd, __VA_ARGS__); break; \ - default: throw std::logic_error( \ - "ff::cuda::reg_field: channel count outside [1, 3] is not " \ - "supported by the CUDA launcher"); \ +#define FF_REGFIELD_LAUNCH_C(KERN, NB, ...) \ + switch (nc) { \ + case 1: FF_CUDA_LAUNCH((KERN), \ + blocks, CUDA_NUM_THREADS, 0, stream, bnd, __VA_ARGS__); break; \ + case 2: FF_CUDA_LAUNCH((KERN), \ + blocks, CUDA_NUM_THREADS, 0, stream, bnd, __VA_ARGS__); break; \ + case 3: FF_CUDA_LAUNCH((KERN), \ + blocks, CUDA_NUM_THREADS, 0, stream, bnd, __VA_ARGS__); break; \ + default: throw std::logic_error( \ + "ff::cuda::reg_field: channel count outside [1, 3] is not " \ + "supported by the CUDA launcher"); \ } // Outer dispatch: pick the (compile-time) number of batch dimensions. -#define FF_REGFIELD_LAUNCH(KERN, ...) \ - switch (nbatch) { \ - case 0: FF_REGFIELD_LAUNCH_C(KERN, 0, __VA_ARGS__); break; \ - case 1: FF_REGFIELD_LAUNCH_C(KERN, 1, __VA_ARGS__); break; \ - default: throw std::logic_error( \ - "ff::cuda::reg_field: nbatch > 1 is not supported by the CUDA " \ - "launcher"); \ +#define FF_REGFIELD_LAUNCH(KERN, ...) \ + switch (nbatch) { \ + case 0: FF_REGFIELD_LAUNCH_C(KERN, 0, __VA_ARGS__); break; \ + case 1: FF_REGFIELD_LAUNCH_C(KERN, 1, __VA_ARGS__); break; \ + default: throw std::logic_error( \ + "ff::cuda::reg_field: nbatch > 1 is not supported by the CUDA " \ + "launcher"); \ } // --- ABSOLUTE --------------------------------------------------------- @@ -2082,26 +2083,26 @@ FF_CUHOST void kernel_bending( // overwrite `sol` in place, so — unlike matvec/diag/kernel — they carry no // `op` template param; hence a dedicated dispatch macro. -#define FF_REGFIELD_LAUNCH_RELAX_C(KERN, NB, ...) \ - switch (nc) { \ - case 1: KERN \ - <<>>(bnd, __VA_ARGS__); break; \ - case 2: KERN \ - <<>>(bnd, __VA_ARGS__); break; \ - case 3: KERN \ - <<>>(bnd, __VA_ARGS__); break; \ - default: throw std::logic_error( \ - "ff::cuda::reg_field: channel count outside [1, 3] is not " \ - "supported by the CUDA relax launcher"); \ +#define FF_REGFIELD_LAUNCH_RELAX_C(KERN, NB, ...) \ + switch (nc) { \ + case 1: FF_CUDA_LAUNCH((KERN), \ + blocks, CUDA_NUM_THREADS, 0, stream, bnd, __VA_ARGS__); break; \ + case 2: FF_CUDA_LAUNCH((KERN), \ + blocks, CUDA_NUM_THREADS, 0, stream, bnd, __VA_ARGS__); break; \ + case 3: FF_CUDA_LAUNCH((KERN), \ + blocks, CUDA_NUM_THREADS, 0, stream, bnd, __VA_ARGS__); break; \ + default: throw std::logic_error( \ + "ff::cuda::reg_field: channel count outside [1, 3] is not " \ + "supported by the CUDA relax launcher"); \ } -#define FF_REGFIELD_LAUNCH_RELAX(KERN, ...) \ - switch (nbatch) { \ - case 0: FF_REGFIELD_LAUNCH_RELAX_C(KERN, 0, __VA_ARGS__); break; \ - case 1: FF_REGFIELD_LAUNCH_RELAX_C(KERN, 1, __VA_ARGS__); break; \ - default: throw std::logic_error( \ - "ff::cuda::reg_field: nbatch > 1 is not supported by the CUDA " \ - "relax launcher"); \ +#define FF_REGFIELD_LAUNCH_RELAX(KERN, ...) \ + switch (nbatch) { \ + case 0: FF_REGFIELD_LAUNCH_RELAX_C(KERN, 0, __VA_ARGS__); break; \ + case 1: FF_REGFIELD_LAUNCH_RELAX_C(KERN, 1, __VA_ARGS__); break; \ + default: throw std::logic_error( \ + "ff::cuda::reg_field: nbatch > 1 is not supported by the CUDA " \ + "relax launcher"); \ } template #include #include "utils.h" // allocDevice / copyToDevice / freeDevice / GET_BLOCKS +#include "launch.h" // FF_CUDA_LAUNCH -- the only checked kernel launch #include // std::logic_error using namespace std; @@ -1402,35 +1403,35 @@ void relax_lame_jrls_( // Dispatch the runtime `nbatch` to a compile-time device-kernel launch. // `KERN` is the (device) kernel name; the trailing args are the kernel args. -#define FF_REGFLOW_LAUNCH_NBATCH(KERN, ...) \ - switch (nbatch) { \ - case 0: KERN<0, ndim, op, reduce_t, scalar_t, offset_t, BOUND...> \ - <<>>(bnd, __VA_ARGS__); break; \ - case 1: KERN<1, ndim, op, reduce_t, scalar_t, offset_t, BOUND...> \ - <<>>(bnd, __VA_ARGS__); break; \ - case 2: KERN<2, ndim, op, reduce_t, scalar_t, offset_t, BOUND...> \ - <<>>(bnd, __VA_ARGS__); break; \ - case 3: KERN<3, ndim, op, reduce_t, scalar_t, offset_t, BOUND...> \ - <<>>(bnd, __VA_ARGS__); break; \ - default: throw std::logic_error( \ - "ff::cuda::reg_flow: nbatch > 3 is not supported by the CUDA launcher"); \ +#define FF_REGFLOW_LAUNCH_NBATCH(KERN, ...) \ + switch (nbatch) { \ + case 0: FF_CUDA_LAUNCH((KERN<0, ndim, op, reduce_t, scalar_t, offset_t, BOUND...>), \ + blocks, CUDA_NUM_THREADS, 0, stream, bnd, __VA_ARGS__); break; \ + case 1: FF_CUDA_LAUNCH((KERN<1, ndim, op, reduce_t, scalar_t, offset_t, BOUND...>), \ + blocks, CUDA_NUM_THREADS, 0, stream, bnd, __VA_ARGS__); break; \ + case 2: FF_CUDA_LAUNCH((KERN<2, ndim, op, reduce_t, scalar_t, offset_t, BOUND...>), \ + blocks, CUDA_NUM_THREADS, 0, stream, bnd, __VA_ARGS__); break; \ + case 3: FF_CUDA_LAUNCH((KERN<3, ndim, op, reduce_t, scalar_t, offset_t, BOUND...>), \ + blocks, CUDA_NUM_THREADS, 0, stream, bnd, __VA_ARGS__); break; \ + default: throw std::logic_error( \ + "ff::cuda::reg_flow: nbatch > 3 is not supported by the CUDA launcher"); \ } // Same as FF_REGFLOW_LAUNCH_NBATCH but for the relaxers, whose device kernels // take no `op` template parameter (they always accumulate in place). The last // kernel argument is the red-black colour index `col`. -#define FF_REGFLOW_LAUNCH_RELAX(KERN, ...) \ - switch (nbatch) { \ - case 0: KERN<0, ndim, reduce_t, scalar_t, offset_t, BOUND...> \ - <<>>(bnd, __VA_ARGS__); break; \ - case 1: KERN<1, ndim, reduce_t, scalar_t, offset_t, BOUND...> \ - <<>>(bnd, __VA_ARGS__); break; \ - case 2: KERN<2, ndim, reduce_t, scalar_t, offset_t, BOUND...> \ - <<>>(bnd, __VA_ARGS__); break; \ - case 3: KERN<3, ndim, reduce_t, scalar_t, offset_t, BOUND...> \ - <<>>(bnd, __VA_ARGS__); break; \ - default: throw std::logic_error( \ - "ff::cuda::reg_flow: nbatch > 3 is not supported by the CUDA launcher"); \ +#define FF_REGFLOW_LAUNCH_RELAX(KERN, ...) \ + switch (nbatch) { \ + case 0: FF_CUDA_LAUNCH((KERN<0, ndim, reduce_t, scalar_t, offset_t, BOUND...>), \ + blocks, CUDA_NUM_THREADS, 0, stream, bnd, __VA_ARGS__); break; \ + case 1: FF_CUDA_LAUNCH((KERN<1, ndim, reduce_t, scalar_t, offset_t, BOUND...>), \ + blocks, CUDA_NUM_THREADS, 0, stream, bnd, __VA_ARGS__); break; \ + case 2: FF_CUDA_LAUNCH((KERN<2, ndim, reduce_t, scalar_t, offset_t, BOUND...>), \ + blocks, CUDA_NUM_THREADS, 0, stream, bnd, __VA_ARGS__); break; \ + case 3: FF_CUDA_LAUNCH((KERN<3, ndim, reduce_t, scalar_t, offset_t, BOUND...>), \ + blocks, CUDA_NUM_THREADS, 0, stream, bnd, __VA_ARGS__); break; \ + default: throw std::logic_error( \ + "ff::cuda::reg_flow: nbatch > 3 is not supported by the CUDA launcher"); \ } // --- ABSOLUTE --------------------------------------------------------- diff --git a/include/fastfields/impl/cuda/resize.h b/include/fastfields/impl/cuda/resize.h index 5be012f..dcb13ef 100755 --- a/include/fastfields/impl/cuda/resize.h +++ b/include/fastfields/impl/cuda/resize.h @@ -10,6 +10,7 @@ #include #include #include "utils.h" +#include "launch.h" // FF_CUDA_LAUNCH -- the only checked kernel launch #include #include @@ -157,10 +158,11 @@ void loop( const int threads = CUDA_NUM_THREADS; # define FF_RESIZE_LAUNCH(NB) \ - kernel \ - <<>>( \ - out, inp, shift, d_scale, d_so, d_si, d_to, d_ti) + FF_CUDA_LAUNCH( \ + (kernel), \ + blocks, threads, 0, s, \ + out, inp, shift, d_scale, d_so, d_si, d_to, d_ti) switch (nbatch) { diff --git a/include/fastfields/impl/cuda/restrict.h b/include/fastfields/impl/cuda/restrict.h index d43d5d3..12bafd6 100755 --- a/include/fastfields/impl/cuda/restrict.h +++ b/include/fastfields/impl/cuda/restrict.h @@ -11,6 +11,7 @@ #include #include #include "utils.h" +#include "launch.h" // FF_CUDA_LAUNCH -- the only checked kernel launch #include #include @@ -233,10 +234,11 @@ void loop( const int threads = CUDA_NUM_THREADS; # define FF_RESTRICT_LAUNCH(NB) \ - kernel \ - <<>>( \ - out, inp, shift, d_scale, d_so, d_si, d_to, d_ti) + FF_CUDA_LAUNCH( \ + (kernel), \ + blocks, threads, 0, s, \ + out, inp, shift, d_scale, d_so, d_si, d_to, d_ti) switch (nbatch) { diff --git a/include/fastfields/impl/cuda/splinc.h b/include/fastfields/impl/cuda/splinc.h index 02a99ed..279aedf 100755 --- a/include/fastfields/impl/cuda/splinc.h +++ b/include/fastfields/impl/cuda/splinc.h @@ -4,6 +4,7 @@ #include #include #include "utils.h" +#include "launch.h" // FF_CUDA_LAUNCH -- the only checked kernel launch #include #include @@ -78,8 +79,10 @@ void loop( const int threads = CUDA_NUM_THREADS; # define FF_SPLINC_LAUNCH(NB) \ - kernel \ - <<>>(inp, d_size, d_stride, d_poles) + FF_CUDA_LAUNCH( \ + (kernel), \ + blocks, threads, 0, s, \ + inp, d_size, d_stride, d_poles) switch (nbatch) { diff --git a/include/fastfields/impl/cuda/utils.h b/include/fastfields/impl/cuda/utils.h index 61b20e7..60a3846 100644 --- a/include/fastfields/impl/cuda/utils.h +++ b/include/fastfields/impl/cuda/utils.h @@ -13,6 +13,26 @@ FF_NAMESPACE_BEGIN(FF_DEVICE) ***********************************************************************/ // Number of threads per block for CUDA kernels. (Copied from PyTorch) +// +// 1024 is the ARCHITECTURAL MAXIMUM, not a safe default. A kernel whose +// register usage is high enough has `cudaFuncGetAttributes(...) +// .maxThreadsPerBlock < 1024`, and launching it at 1024 fails immediately +// with `cudaErrorLaunchOutOfResources`. Until fastfields-lib#152 that failure +// was silently discarded; it is now reported, and the report names this +// kernel's `maxThreadsPerBlock` and register count so the number to clamp to +// is in the message (see `_throwLaunchError` in launch.h). +// +// Making the block size fit automatically -- `cudaOccupancyMaxPotentialBlock +// Size`, or a clamp to `maxThreadsPerBlock` -- is deliberately NOT done here +// yet, because two launchers couple the launch geometry to a host-side +// allocation: `distance_euclidean::dt` and `distance_mesh::sdt` size a +// per-lane scratch buffer as `num_blocks * CUDA_NUM_THREADS` while the kernel +// derives its lane stride from `blockDim.x * gridDim.x`. Shrinking the block +// alone is safe there (the grid-stride loops still cover every element and +// the buffer is merely over-allocated); raising the grid to compensate is an +// out-of-bounds device write. Untangling that is its own change, and one that +// no CI here can exercise -- there is no GPU, and `cudaFuncGetAttributes` +// needs a device. static constexpr int CUDA_NUM_THREADS = 1024; // Set the number of blocks for CUDA kernel launches. (Copied from PyTorch) @@ -168,11 +188,11 @@ FF_CUHOST inline O * copyToDevice(const I * inp, S size, O * out = nullptr) ); if (err) throw std::bad_alloc(); } - catch (const std::exception &exc) + catch (const std::exception &) { freeDevice(ownout); freeHost(owntmp); - throw exc; + throw; } freeHost(owntmp); return out; @@ -258,11 +278,11 @@ FF_CUHOST inline O * copyToDeviceAsync( if (err) throw std::bad_alloc(); } } - catch (const std::exception &exc) + catch (const std::exception &) { freeDevice(ownout); freeHost(owntmp); - throw exc; + throw; } freeHost(owntmp); return out; @@ -308,10 +328,10 @@ FF_CUHOST inline O * copyToHost(const I * inp, S size, O * out = nullptr) freeHost(stage); } } - catch (const std::exception &exc) + catch (const std::exception &) { freeHost(ownout); - throw exc; + throw; } return out; } From a6699db69b1a69872827a12e1776a4fec25b8634 Mon Sep 17 00:00:00 2001 From: fastfields agent Date: Thu, 20 Aug 2026 16:19:09 +0000 Subject: [PATCH 2/4] build(ci): fail the build when a kernel launch bypasses the helper 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 --- .github/workflows/ci.yml | 30 ++++ tools/check-cuda-launches.py | 297 +++++++++++++++++++++++++++++++++++ 2 files changed, 327 insertions(+) create mode 100755 tools/check-cuda-launches.py diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 81298d0..e7929a4 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -162,6 +162,36 @@ jobs: - run: pip install "codespell==2.4.3" - run: codespell + # Every CUDA kernel launch must go through ff::cuda::launchKernel, so that + # the error state a launch sets is actually inspected (fastfields-lib#152: + # 39 `<<<` sites, zero cudaGetLastError, so a kernel that never ran was + # indistinguishable from one that ran correctly). + # + # This job is why that survives. There is no GPU in CI and there is not + # going to be one, so nothing here can execute a launch and no test can + # notice the check being dropped -- but a launch written the old way is a + # purely textual property, and this catches it in seconds on any runner. + # Unconditional, not path-filtered: the rule is about the whole tree, and + # the cost is a Python startup. + # + # `--selftest` runs first and deliberately: a checker that has quietly + # stopped matching anything prints "clean" forever. It asserts the analyser + # still flags a bare launch, still ignores one inside a comment or a string + # literal, and still reports the right line number. + cuda-launches: + name: lint (cuda launch sites) + runs-on: ubuntu-latest + timeout-minutes: 10 + steps: + - uses: actions/checkout@v5 + - uses: actions/setup-python@v5 + with: + python-version: "3.x" + - name: Verify the checker itself + run: python3 tools/check-cuda-launches.py --selftest + - name: Every `<<<` is inside the launch helper + run: python3 tools/check-cuda-launches.py --check + clang-format: name: lint (clang-format, changed lines) # Needs a base ref to diff against; there is none on a push to main. diff --git a/tools/check-cuda-launches.py b/tools/check-cuda-launches.py new file mode 100755 index 0000000..fdccc6c --- /dev/null +++ b/tools/check-cuda-launches.py @@ -0,0 +1,297 @@ +#!/usr/bin/env python3 +""" +check-cuda-launches.py -- every CUDA kernel launch goes through one helper. + +THE RULE THIS ENFORCES +-------------------------------------------------------------------------- +A CUDA kernel launch is asynchronous and does not throw. It reports failure +by setting the runtime's error state, and if nobody inspects that state the +failure is simply discarded: the launcher returns normally and the caller +reads whatever was already in the output buffer. That was the state of this +tree in fastfields-lib#152 -- 39 `<<<` sites, zero `cudaGetLastError`. + +The fix was structural rather than conventional. `<<<` now appears in +exactly ONE place: + + include/fastfields/impl/cuda/launch.h ff::cuda::launchKernel + +and every launcher reaches it through the `FF_CUDA_LAUNCH` macro. This +script is what keeps that true. It fails if: + + 1. `<<<` appears in code anywhere outside the helper; + 2. the helper does not contain exactly one `<<<`; + 3. the helper stops calling `cudaGetLastError` -- i.e. the funnel is + still a funnel but no longer checks anything; + 4. `cudaLaunchKernel` / `cudaLaunchCooperativeKernel` is called outside + the helper, which is the way to launch a kernel without writing `<<<`. + +WHY A SCRIPT AND NOT A CONVENTION +-------------------------------------------------------------------------- +There is no GPU in CI, so nothing here can *execute* a launch; compile+link +is the whole CUDA gate. A regression is still perfectly preventable, but +only mechanically. This project has learned the same lesson twice already -- +`-Wl,--no-undefined` (#87) and the `FF_MEM_BUDGET_KB` budget (#95) both +exist because a rule people were supposed to remember was not remembered -- +and #152 exists precisely because "check the launch" was never applied even +once, at any of the 39 sites, over the life of the code. + +COMMENTS AND STRINGS ARE NOT CODE +-------------------------------------------------------------------------- +Comments and string literals are blanked before scanning, so prose may +discuss `<<<` freely (launch.h's own header comment does, six times). That +is not laxity: it is what lets the "exactly one" half of the rule be stated +about the helper itself rather than waived for it. + +USAGE +-------------------------------------------------------------------------- + python3 tools/check-cuda-launches.py # check (the default) + python3 tools/check-cuda-launches.py --check # same, explicit + python3 tools/check-cuda-launches.py --selftest # verify the checker + +`--selftest` runs the analyser over synthetic inputs and asserts that it +both flags the violations it is supposed to flag and passes clean code. A +lint that has quietly stopped matching anything reports success forever; +this is what makes that failure mode visible. CI runs both modes. + +Exit status + 0 the tree obeys the rule (or, with --selftest, the checker works) + 1 a violation was found (or the checker is broken) + 2 usage / environment error +""" + +import os +import re +import sys + +ROOT = os.path.dirname(os.path.dirname(os.path.abspath(__file__))) + +SOURCE_DIRS = ("include", "src", "tests") +SOURCE_EXTS = (".h", ".hpp", ".inl", ".cpp", ".cu", ".cuh") + +# The one file allowed to contain a launch. +HELPER = os.path.join("include", "fastfields", "impl", "cuda", "launch.h") + +# The check the helper must still be performing. `cudaPeekAtLastError` alone +# is not enough: it does not clear, so a helper that only peeked would report +# the same stale error on every subsequent launch. +REQUIRED_IN_HELPER = "cudaGetLastError" + +# Launching without writing `<<<`. +BACKDOORS = ("cudaLaunchKernel", "cudaLaunchCooperativeKernel") + + +# ---------------------------------------------------------------- stripping + +def blank_comments_and_strings(text): + """Replace comments and string/char literals with spaces. + + Newlines are preserved so reported line numbers stay meaningful, and the + result has the same length as the input so column offsets do too. + """ + out = [] + i, n = 0, len(text) + while i < n: + c = text[i] + nxt = text[i + 1] if i + 1 < n else "" + if c == "/" and nxt == "/": + while i < n and text[i] != "\n": + out.append(" ") + i += 1 + elif c == "/" and nxt == "*": + out.append(" ") + i += 2 + while i < n and not (text[i] == "*" and i + 1 < n and text[i + 1] == "/"): + out.append("\n" if text[i] == "\n" else " ") + i += 1 + out.append(" ") + i += 2 + elif c in ('"', "'"): + quote = c + out.append(" ") + i += 1 + while i < n and text[i] != quote: + if text[i] == "\\" and i + 1 < n: + out.append(" ") + i += 2 + continue + out.append("\n" if text[i] == "\n" else " ") + i += 1 + out.append(" ") + i += 1 + else: + out.append(c) + i += 1 + return "".join(out) + + +def code_hits(text, needle): + """[(line number, source line)] for `needle` outside comments/strings.""" + blanked = blank_comments_and_strings(text) + raw = text.split("\n") + hits = [] + for k, line in enumerate(blanked.split("\n")): + if needle in line: + hits.append((k + 1, raw[k].strip() if k < len(raw) else "")) + return hits + + +# ------------------------------------------------------------------- checks + +def iter_sources(): + for d in SOURCE_DIRS: + top = os.path.join(ROOT, d) + if not os.path.isdir(top): + continue + for dirpath, _dirnames, filenames in os.walk(top): + for fn in sorted(filenames): + if fn.endswith(SOURCE_EXTS): + path = os.path.join(dirpath, fn) + yield os.path.relpath(path, ROOT) + + +def check_tree(): + """Returns (list of problem strings, stats dict).""" + problems = [] + launch_sites = 0 + helper_seen = False + scanned = 0 + + for rel in iter_sources(): + with open(os.path.join(ROOT, rel), "r", encoding="utf-8", errors="replace") as fh: + text = fh.read() + scanned += 1 + blanked = blank_comments_and_strings(text) + is_helper = rel.replace("\\", "/") == HELPER.replace("\\", "/") + if not is_helper: + # The helper's own `#define` is not a call site. + launch_sites += len(re.findall(r"\bFF_CUDA_LAUNCH\s*\(", blanked)) + + if is_helper: + helper_seen = True + hits = code_hits(text, "<<<") + if len(hits) != 1: + problems.append( + "%s: expected exactly one `<<<` in code, found %d%s" + % (rel, len(hits), + "".join("\n line %d: %s" % h for h in hits))) + if REQUIRED_IN_HELPER not in blanked: + problems.append( + "%s: the launch helper no longer calls `%s` -- every launch " + "still funnels through it, but nothing checks the result " + "(this is exactly fastfields-lib#152)" + % (rel, REQUIRED_IN_HELPER)) + continue + + for lineno, src in code_hits(text, "<<<"): + problems.append( + "%s:%d: kernel launch outside the helper\n %s\n" + " use FF_CUDA_LAUNCH((kernel<...>), grid, block, shmem, " + "stream, args...) from " + % (rel, lineno, src)) + + for backdoor in BACKDOORS: + for lineno, src in code_hits(text, backdoor): + problems.append( + "%s:%d: `%s` outside the helper -- that launches a kernel " + "without the error check\n %s" + % (rel, lineno, backdoor, src)) + + if not helper_seen: + problems.append( + "%s is missing: the launch helper is the whole mechanism this " + "check exists to protect" % HELPER) + if launch_sites == 0: + problems.append( + "no FF_CUDA_LAUNCH call sites found anywhere -- either the funnel " + "was removed or this check is looking in the wrong place") + + return problems, {"scanned": scanned, "launch_sites": launch_sites} + + +# ----------------------------------------------------------------- selftest + +_SELFTEST_CASES = [ + # (name, source, needle, expected number of code hits) + ("bare launch", "void f() { k<<<1, 2, 0, s>>>(x); }", "<<<", 1), + ("launch in a // comment", "// k<<<1,2>>>(x) is what this replaces\n", "<<<", 0), + ("launch in a /* */ comment", "/* k<<<1,2>>>(x)\n more */\nint x;\n", "<<<", 0), + ("launch in a string", 'const char* s = "k<<<1,2>>>(x)";\n', "<<<", 0), + ("stream operators are not launches", "std::cout << x << y;\n", "<<<", 0), + ("shift then template", "a = b << c;\n", "<<<", 0), + ("escaped quote does not swallow code", + 'const char* s = "\\"";\nvoid f() { k<<<1,1,0,s>>>(x); }\n', "<<<", 1), + ("comment marker inside a string is not a comment", + 'const char* s = "// not a comment";\nvoid f() { k<<<1,1,0,s>>>(x); }\n', + "<<<", 1), + ("backdoor call", "cudaLaunchKernel(f, g, b, a, 0, s);\n", "cudaLaunchKernel", 1), + ("backdoor in a comment", "// cudaLaunchKernel is the other way in\n", + "cudaLaunchKernel", 0), +] + + +def selftest(): + failures = [] + for name, src, needle, expected in _SELFTEST_CASES: + got = len(code_hits(src, needle)) + if got != expected: + failures.append( + " %-45s expected %d hit(s) for %r, got %d" + % (name, expected, needle, got)) + + # Line numbers must survive the blanking, or a report points at the wrong + # place -- which is how a lint stops being actionable. + hits = code_hits("/* a\n b */\nint x;\nvoid f(){ k<<<1,1,0,s>>>(y); }\n", "<<<") + if hits != [] and hits[0][0] != 4: + failures.append(" %-45s expected the hit on line 4, got line %d" + % ("line numbers survive blanking", hits[0][0])) + + # And the whole-tree check has to actually be able to fail. + if not code_hits("k<<<1,1,0,s>>>(x);", "<<<"): + failures.append(" %-45s the detector matches nothing at all" + % "detector is live") + + if failures: + print("check-cuda-launches: SELFTEST FAILED -- the checker is broken, " + "so a green run below would mean nothing:") + print("\n".join(failures)) + return 1 + print("check-cuda-launches: selftest passed (%d cases)" % (len(_SELFTEST_CASES) + 2)) + return 0 + + +# --------------------------------------------------------------------- main + +def main(argv): + mode = "check" + for arg in argv[1:]: + if arg in ("--check", "--selftest"): + mode = arg[2:] + elif arg in ("-h", "--help"): + print(__doc__.strip()) + return 0 + else: + sys.stderr.write("check-cuda-launches: unknown argument: %s\n" % arg) + return 2 + + if mode == "selftest": + return selftest() + + problems, stats = check_tree() + if problems: + print("check-cuda-launches: %d problem(s):\n" % len(problems)) + for p in problems: + print(" " + p) + print("\nEvery CUDA kernel launch must go through " + "ff::cuda::launchKernel in %s, which checks cudaGetLastError and " + "throws with the kernel name and launch configuration. See " + "fastfields-lib#152." % HELPER) + return 1 + + print("check-cuda-launches: clean -- %d source file(s) scanned, " + "1 launch site in %s, %d FF_CUDA_LAUNCH call site(s)." + % (stats["scanned"], HELPER, stats["launch_sites"])) + return 0 + + +if __name__ == "__main__": + sys.exit(main(sys.argv)) From 25d6e1dbbd0f744ece8cd01b8a24e9ec62be7983 Mon Sep 17 00:00:00 2001 From: fastfields agent Date: Thu, 20 Aug 2026 16:19:09 +0000 Subject: [PATCH 3/4] docs: record the one-launch-site rule in CLAUDE.md 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 --- CLAUDE.md | 26 +++++++++++++++++++++----- 1 file changed, 21 insertions(+), 5 deletions(-) diff --git a/CLAUDE.md b/CLAUDE.md index 843d61d..9142b16 100644 --- a/CLAUDE.md +++ b/CLAUDE.md @@ -110,11 +110,11 @@ the CPU path is the tested source of truth and CUDA is **compile+link only**. ## CI -`.github/workflows/ci.yml`, path-filtered. `codespell` always; `test-cpu` (a -3-leg `BOUNDFLAGS`/`SPLINEFLAGS` matrix + an `INDEXFLAGS` leg + a g++ leg), -`sanitize` (ASan+UBSan) and `tsan` on kernels/cpu/hub changes; `test-hub` on -hub changes; `build-cuda` (two legs, one per `FF_INDEX32` position) and -`compile-probe-cuda` on kernels/cuda changes. +`.github/workflows/ci.yml`, path-filtered. `codespell` and `lint (cuda launch +sites)` always; `test-cpu` (a 3-leg `BOUNDFLAGS`/`SPLINEFLAGS` matrix + an +`INDEXFLAGS` leg + a g++ leg), `sanitize` (ASan+UBSan) and `tsan` on +kernels/cpu/hub changes; `test-hub` on hub changes; `build-cuda` (two legs, one +per `FF_INDEX32` position) and `compile-probe-cuda` on kernels/cuda changes. **The `tsan` leg is the only one that runs anything in parallel.** With the shipping `GRAIN_SIZE` (32768) every workload in `tests/lib-cpu/` is below the @@ -152,6 +152,22 @@ pushpull's fully-static order×bound compile is nightly `FF_CUDA::` entry points undefined with every build green. The hub link is now the gate on backend completeness — forget a `MODULES` entry and it fails there. +- **Every CUDA kernel launch goes through `FF_CUDA_LAUNCH`, and `<<<` appears + in exactly one file** — `include/fastfields/impl/cuda/launch.h`, whose + `ff::cuda::launchKernel` launches on the caller's stream, calls + `cudaGetLastError()`, and throws with the kernel name and the grid/block + configuration. A launch is asynchronous and does not throw; it only sets an + error state, and until fastfields-lib#152 there were 39 launches and zero + reads of that state, so a kernel that never ran was indistinguishable from + one that ran correctly. The post-launch check is host-side and does **not** + synchronise, so it costs nothing; the price of that is that it cannot see a + fault raised while the kernel *executes*. Observing those needs a + synchronisation point, which is `FF_CUDA_LAUNCH_SYNC` — build flag *and* + environment variable, **off by default and it stays off**, this project's + `CUDA_LAUNCH_BLOCKING`. `tools/check-cuda-launches.py --check` fails if a + `<<<` (or a raw `cudaLaunchKernel`) shows up outside the helper, and + `--selftest` checks the checker; both run in CI on every push. No GPU can + ever catch a regression here, so the lint is the whole guard. - Op renames from the impl layer: `resize -> resample`, `restrict -> restriction`, `splinc -> spline_coeff` (a namespace cannot share a name with a function inside `ff::cpu`). **`restriction` accumulates into From bb078fedcc8ac8b5b9e0fbdb5ff6f406996dbbe7 Mon Sep 17 00:00:00 2001 From: fastfields agent Date: Thu, 20 Aug 2026 17:06:37 +0000 Subject: [PATCH 4/4] tools: teach the launch check about raw string literals 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 #152 --- tools/check-cuda-launches.py | 34 +++++++++++++++++++++++++++++++++- 1 file changed, 33 insertions(+), 1 deletion(-) diff --git a/tools/check-cuda-launches.py b/tools/check-cuda-launches.py index fdccc6c..92cfe38 100755 --- a/tools/check-cuda-launches.py +++ b/tools/check-cuda-launches.py @@ -42,6 +42,11 @@ is not laxity: it is what lets the "exactly one" half of the rule be stated about the helper itself rather than waived for it. +Raw string literals are handled too. The tree has none today, but a raw +string's body may hold unescaped `"` and `\`, and a scanner that mishandled +one would desynchronise and stop seeing launches *after* it -- a silent +weakening, which is the one failure mode this whole file exists to prevent. + USAGE -------------------------------------------------------------------------- python3 tools/check-cuda-launches.py # check (the default) @@ -79,6 +84,10 @@ # Launching without writing `<<<`. BACKDOORS = ("cudaLaunchKernel", "cudaLaunchCooperativeKernel") +# `R"delim(` -- the opening of a raw string literal, with optional encoding +# prefix. Group 1 is the delimiter, so the matching `)delim"` can be found. +_RAW_STRING = re.compile(r'(?:u8|u|U|L)?R"([^ ()\\\t\n]{0,16})\(') + # ---------------------------------------------------------------- stripping @@ -93,7 +102,20 @@ def blank_comments_and_strings(text): while i < n: c = text[i] nxt = text[i + 1] if i + 1 < n else "" - if c == "/" and nxt == "/": + # Raw string literal: R"delim( ... )delim". Handled before the ordinary + # quote case because its body may contain unescaped `"` and `\`, which + # would otherwise desynchronise the scanner and let a later launch pass + # unseen. The tree has no raw strings today; this is here so that adding + # one cannot silently weaken the check. + prev = text[i - 1] if i else "" + m = (None if (prev.isalnum() or prev == "_") + else _RAW_STRING.match(text, i)) + if m: + end = text.find(')' + m.group(1) + '"', m.end()) + end = n if end < 0 else end + len(m.group(1)) + 2 + out.append("".join("\n" if ch == "\n" else " " for ch in text[i:end])) + i = end + elif c == "/" and nxt == "/": while i < n and text[i] != "\n": out.append(" ") i += 1 @@ -226,6 +248,16 @@ def check_tree(): ("backdoor call", "cudaLaunchKernel(f, g, b, a, 0, s);\n", "cudaLaunchKernel", 1), ("backdoor in a comment", "// cudaLaunchKernel is the other way in\n", "cudaLaunchKernel", 0), + # Raw strings: their body may hold unescaped quotes and backslashes, which + # would desynchronise a naive scanner and hide every launch after them. + ("launch inside a raw string", + 'const char* s = R"(k<<<1,1,0,s>>>(x) "quoted" \\ )";\n', "<<<", 0), + ("code after a raw string is still scanned", + 'const char* s = R"x(a "b" \\ )x";\nvoid f(){ k<<<1,1,0,s>>>(y); }\n', + "<<<", 1), + # `R` preceded by an identifier character is not a raw-string prefix. + ("identifier ending in R is not a raw string", + 'void f(){ MYR("a(b"); k<<<1,1,0,s>>>(y); }\n', "<<<", 1), ]