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/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 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; } diff --git a/tools/check-cuda-launches.py b/tools/check-cuda-launches.py new file mode 100755 index 0000000..92cfe38 --- /dev/null +++ b/tools/check-cuda-launches.py @@ -0,0 +1,329 @@ +#!/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. + +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) + 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") + +# `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 + +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 "" + # 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 + 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), + # 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), +] + + +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))