Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
30 changes: 30 additions & 0 deletions .github/workflows/ci.yml
Original file line numberDiff line numberDiff line change
Expand Up@@ -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.
Expand Down
26 changes: 21 additions & 5 deletions CLAUDE.md
Original file line numberDiff line numberDiff line change
Expand Up@@ -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
Expand DownExpand Up@@ -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
Expand Down
4 changes: 2 additions & 2 deletions include/fastfields/impl/cpu/distance_euclidean.h
Original file line numberDiff line numberDiff line change
Expand Up@@ -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;
Expand Down
12 changes: 7 additions & 5 deletions include/fastfields/impl/cuda/distance_euclidean.h
Original file line numberDiff line numberDiff line change
Expand Up@@ -3,6 +3,7 @@
#include <fastfields/impl/kernels/distance.h>
#include <fastfields/impl/kernels/batch.h>
#include "utils.h"
#include "launch.h" // FF_CUDA_LAUNCH -- the only checked kernel launch
#include <exception>
#include <cstdint>

Expand DownExpand Up@@ -76,14 +77,15 @@ FF_CUHOST void dt(
buffer = allocDevice<char>(buffer_size);
size_device = copyToDeviceAsync(size, ndim, s);
stride_device = copyToDeviceAsync(stride, ndim, s);
dt_kernel<scalar_t, offset_t>
<<<num_blocks, CUDA_NUM_THREADS, 0, s>>>
(ndim, f, buffer, w, size_device, stride_device);
FF_CUDA_LAUNCH(
(dt_kernel<scalar_t, offset_t>),
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);
}
Expand Down
12 changes: 7 additions & 5 deletions include/fastfields/impl/cuda/distance_l1.h
Original file line numberDiff line numberDiff line change
Expand Up@@ -3,6 +3,7 @@
#include <fastfields/impl/kernels/distance.h>
#include <fastfields/impl/kernels/batch.h>
#include "utils.h"
#include "launch.h" // FF_CUDA_LAUNCH -- the only checked kernel launch
#include <exception>
#include <cstdint>

Expand DownExpand Up@@ -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<scalar_t, offset_t>
<<<GET_BLOCKS(batch_size), CUDA_NUM_THREADS, 0, s>>>
(ndim, f, w, size_device, stride_device);
FF_CUDA_LAUNCH(
(dt_kernel<scalar_t, offset_t>),
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);
}
Expand Down
134 changes: 74 additions & 60 deletions include/fastfields/impl/cuda/distance_mesh.h
Original file line numberDiff line numberDiff line change
Expand Up@@ -4,6 +4,7 @@
#include <fastfields/impl/kernels/batch.h>
#include <fastfields/impl/kernels/utils.h>
#include "utils.h"
#include "launch.h" // FF_CUDA_LAUNCH -- the only checked kernel launch
#include <cstdint>
#include <memory> // std::unique_ptr
#include <type_traits> // std::is_trivially_copyable
Expand DownExpand Up@@ -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<scalar_t, offset_t>
<<<GET_BLOCKS(numel), CUDA_NUM_THREADS, 0, stream>>>
(ndim, out, inp, size_copy, stride_out_copy, stride_inp_copy);
FF_CUDA_LAUNCH(
(copy_tensor_kernel<scalar_t, offset_t>),
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);
Expand DownExpand Up@@ -305,9 +307,21 @@ index_t * copy_faces(
{
offset_t stride0 = stride[0], stride1 = stride[1];
index_t * faces_out = allocDevice<index_t>(nb_faces * ndim);
copy_faces_kernel<ndim, index_t, offset_t>
<<<GET_BLOCKS(nb_faces), CUDA_NUM_THREADS, 0, stream>>>
(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<ndim, index_t, offset_t>),
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;
}

Expand DownExpand Up@@ -1142,33 +1156,33 @@ sdt(
treetrace_device = allocDevice<char>(stride_buf * treesize);

// Compute SDT
sdt_kernel<ndim, scalar_t, index_t, offset_t>
<<<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
);
FF_CUDA_LAUNCH(
(sdt_kernel<ndim, scalar_t, index_t, offset_t>),
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,
Expand All@@ -1194,7 +1208,7 @@ sdt(
normverts_host,
normedges_host
);
throw e;
throw;
}

freeDevice(
Expand DownExpand Up@@ -1355,31 +1369,31 @@ sdt_naive(
int num_blocks = GET_BLOCKS(numel);

// Compute SDT
sdt_naive_kernel<ndim, scalar_t, index_t, offset_t>
<<<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
);
FF_CUDA_LAUNCH(
(sdt_naive_kernel<ndim, scalar_t, index_t, offset_t>),
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,
Expand All@@ -1401,7 +1415,7 @@ sdt_naive(
normverts_host,
normedges_host
);
throw e;
throw;
}

freeDevice(
Expand Down
Loading
Loading