From 0c3c45e653e912fab70eceb58b05c05ff437baba Mon Sep 17 00:00:00 2001 From: kferjaoui Date: Mon, 28 Sep 2026 11:47:55 +0200 Subject: [PATCH] CUDA: unify GPU code under aare::cuda - aare::device -> aare::cuda::detail; drivers, RAII helpers, StreamContext* moved into aare::cuda(::detail) - Algo contract: add compute_type, assert floating point - LaunchParams::nSigma uses compute_type - test: use FastPedestal in the manual route --- docs/conf.py.in | 4 +- docs/src/ClusterFinderCUDA.rst | 6 +-- include/aare/ClusterFinderCUDA.hpp | 19 +++++---- include/aare/ClusterFinderCUDA_graph.hpp | 45 ++++++++++++---------- include/aare/RAII_device.hpp | 4 +- include/aare/clusterfinder_algo.cuh | 18 ++++++--- include/aare/clusterfinder_kernel.cuh | 4 +- python/src/bind_ClusterFinderCUDA.hpp | 4 +- python/src/bind_ClusterFinderCUDAGraph.hpp | 2 +- src/ClusterFinderCUDA_testimpl.cu | 16 +++++--- 10 files changed, 73 insertions(+), 49 deletions(-) diff --git a/docs/conf.py.in b/docs/conf.py.in index f0051ef3..6c53243c 100644 --- a/docs/conf.py.in +++ b/docs/conf.py.in @@ -38,8 +38,8 @@ napoleon_use_ivar = True # CUDA execution-space and pointer qualifiers are not C++, so Sphinx's C++ # domain cannot parse declarations such as the __global__ kernels in -# aare::device. Declaring them as attributes lets the parser skip them instead -# of failing, which matters because the docs are built with -W. +# aare::cuda::detail. Declaring them as attributes lets the parser skip them +# instead of failing, which matters because the docs are built with -W. cpp_id_attributes = [ '__global__', '__device__', diff --git a/docs/src/ClusterFinderCUDA.rst b/docs/src/ClusterFinderCUDA.rst index b8a93097..55864fba 100644 --- a/docs/src/ClusterFinderCUDA.rst +++ b/docs/src/ClusterFinderCUDA.rst @@ -5,7 +5,7 @@ GPU cluster finder. Available only when aare is configured with ``-DAARE_CUDA=ON``; on a CPU-only build the class is not compiled and the Python factory raises ``RuntimeError``. -Single frames go through :cpp:func:`aare::ClusterFinderCUDA::find_clusters`, +Single frames go through :cpp:func:`aare::cuda::ClusterFinderCUDA::find_clusters`, which mirrors the CPU :cpp:class:`aare::ClusterFinder` interface. Throughput comes from ``find_clusters_batched``, which distributes a batch of frames round-robin over ``n_streams`` CUDA streams (4 by default) and overlaps H2D, @@ -14,13 +14,13 @@ returns one ``ClusterVector`` per frame, or with ``collect_view``, which hands back a ``BatchView`` reading clusters in place from the pinned host buffer with no per-frame allocation or copy. -.. doxygenclass:: aare::ClusterFinderCUDA +.. doxygenclass:: aare::cuda::ClusterFinderCUDA :members: :undoc-members: Device kernel ------------- -.. doxygennamespace:: aare::device +.. doxygennamespace:: aare::cuda::detail :members: :undoc-members: diff --git a/include/aare/ClusterFinderCUDA.hpp b/include/aare/ClusterFinderCUDA.hpp index 2fdb7c65..3ee75bf8 100644 --- a/include/aare/ClusterFinderCUDA.hpp +++ b/include/aare/ClusterFinderCUDA.hpp @@ -14,7 +14,9 @@ #include #include -namespace aare { +namespace aare::cuda { + +namespace detail { /** * @brief One stream's device resources, owned. @@ -61,6 +63,8 @@ template class StreamContext { typename Ped::View pedestal_view() const { return Ped::make_view(m_ped); } }; +} // namespace detail + /** * @brief GPU cluster-finding driver: streams, batching, host staging. * @@ -84,15 +88,16 @@ class ClusterFinderCUDA { using FRAME_TYPE = typename Algo::frame_type; using Ped = typename Algo::pedestal; using Output = typename Algo::Output; + using COMPUTE_TYPE = + typename Algo::compute_type; // match the kernel's internal precision + // i.e. detail::COMPUTE_TYPE private: static_assert(no_2x2_cluster::value, "ClusterFinderCUDA: both cluster dimensions must exceed 2"); - static_assert(cuda::contract::require_algo::value, + static_assert(contract::require_algo::value, "ClusterFinderCUDA: Algo does not satisfy the contract in " "clusterfinder_algo.cuh"); - using COMPUTE_TYPE = - device::COMPUTE_TYPE; // match the kernel's internal precision // Two ping-pong output slots. At most 2 batches can be in flight at once: // the caller submits batch B before collecting batch A, so A and B occupy @@ -138,7 +143,7 @@ class ClusterFinderCUDA { bool m_pedestal_dirty = true; // One per stream. unique_ptr so the context can stay non-movable. - std::vector>> v_sc; + std::vector>> v_sc; float m_total_kernel_ms = 0.0f; size_t m_frames_processed = 0; @@ -473,7 +478,7 @@ class ClusterFinderCUDA { // algorithm's pedestal arrays and the packed output block. v_sc.reserve(n_streams); for (int k = 0; k < n_streams; ++k) - v_sc.push_back(std::make_unique>( + v_sc.push_back(std::make_unique>( m_image_size, static_cast(m_max_clusters_per_frame))); for (int s = 0; s < NUM_SLOTS; ++s) { @@ -962,4 +967,4 @@ class ClusterFinderCUDA { } }; -} // namespace aare +} // namespace aare::cuda diff --git a/include/aare/ClusterFinderCUDA_graph.hpp b/include/aare/ClusterFinderCUDA_graph.hpp index a2757a35..5e750a80 100644 --- a/include/aare/ClusterFinderCUDA_graph.hpp +++ b/include/aare/ClusterFinderCUDA_graph.hpp @@ -12,7 +12,9 @@ #include #include -namespace aare { +namespace aare::cuda { + +namespace detail { // Per-stream device resources and CUDA Graph state. template @@ -21,10 +23,10 @@ struct StreamContextGraph { FRAME_TYPE *d_frame = nullptr; // Device pedestal precision is set by DEVICE_PED_TYPE in the kernel header. // Accumulators hold CENTERED moments of Y = X - d_pd_off (see kernel). - device::DEVICE_PED_TYPE *d_pd_mean = nullptr; - device::DEVICE_PED_TYPE *d_pd_sum = nullptr; - device::DEVICE_PED_TYPE *d_pd_sum2 = nullptr; - device::DEVICE_PED_TYPE *d_pd_off = nullptr; // frozen per-pixel baseline X0 + DEVICE_PED_TYPE *d_pd_mean = nullptr; + DEVICE_PED_TYPE *d_pd_sum = nullptr; + DEVICE_PED_TYPE *d_pd_sum2 = nullptr; + DEVICE_PED_TYPE *d_pd_off = nullptr; // frozen per-pixel baseline X0 uint8_t *d_output = nullptr; // [uint32_t count | ClusterType clusters[max]] // CUDA Graph handles — rebuilt on pedestal change or h_output_pinned resize @@ -36,12 +38,12 @@ struct StreamContextGraph { // Kernel argument storage for cudaKernelNodeParams. // kargs_ptrs holds addresses into these fields; members must not move. FRAME_TYPE *karg_d_frame = nullptr; - device::DEVICE_PED_TYPE *karg_d_pd_mean = nullptr; - device::DEVICE_PED_TYPE *karg_d_pd_sum = nullptr; - device::DEVICE_PED_TYPE *karg_d_pd_sum2 = nullptr; - device::DEVICE_PED_TYPE *karg_d_pd_off = nullptr; + DEVICE_PED_TYPE *karg_d_pd_mean = nullptr; + DEVICE_PED_TYPE *karg_d_pd_sum = nullptr; + DEVICE_PED_TYPE *karg_d_pd_sum2 = nullptr; + DEVICE_PED_TYPE *karg_d_pd_off = nullptr; uint32_t karg_n_pd_samples = 0; - device::COMPUTE_TYPE karg_nSigma = 0.0f; + COMPUTE_TYPE karg_nSigma = 0.0f; int32_t karg_nrows = 0; // must match the kernel signature exactly: int32_t karg_ncols = 0; // kernelParams is untyped ClusterType *karg_d_clusters = nullptr; @@ -54,12 +56,14 @@ struct StreamContextGraph { cudaMemcpy3DParms d2h_params = {}; }; +} // namespace detail + template , typename FRAME_TYPE = uint16_t, typename PEDESTAL_TYPE = double, typename = std::enable_if_t::value>> class ClusterFinderCUDAGraph { using COMPUTE_TYPE = - device::COMPUTE_TYPE; // match the kernel's internal precision + detail::COMPUTE_TYPE; // match the kernel's internal precision static constexpr int BLOCK_X = 16; static constexpr int BLOCK_Y = 16; @@ -99,9 +103,10 @@ class ClusterFinderCUDAGraph { bool m_pedestal_dirty = true; bool m_graphs_dirty = true; // set when pedestal or h_output_pinned changes // Frozen per-pixel baseline X0 (~mean at t=0), captured once on first sync. - std::vector m_offset; + std::vector m_offset; - using SC = StreamContextGraph; + using SC = + detail::StreamContextGraph; std::vector v_sc; size_t m_frames_processed = 0; @@ -190,13 +195,13 @@ class ClusterFinderCUDAGraph { cudaStreamCreateWithFlags(&sc.stream, cudaStreamNonBlocking)); CUDA_CHECK(cudaMalloc(&sc.d_frame, m_image_bytes)); CUDA_CHECK(cudaMalloc( - &sc.d_pd_mean, m_image_size * sizeof(device::DEVICE_PED_TYPE))); + &sc.d_pd_mean, m_image_size * sizeof(detail::DEVICE_PED_TYPE))); CUDA_CHECK(cudaMalloc( - &sc.d_pd_sum, m_image_size * sizeof(device::DEVICE_PED_TYPE))); + &sc.d_pd_sum, m_image_size * sizeof(detail::DEVICE_PED_TYPE))); CUDA_CHECK(cudaMalloc( - &sc.d_pd_sum2, m_image_size * sizeof(device::DEVICE_PED_TYPE))); + &sc.d_pd_sum2, m_image_size * sizeof(detail::DEVICE_PED_TYPE))); CUDA_CHECK(cudaMalloc( - &sc.d_pd_off, m_image_size * sizeof(device::DEVICE_PED_TYPE))); + &sc.d_pd_off, m_image_size * sizeof(detail::DEVICE_PED_TYPE))); CUDA_CHECK(cudaMalloc(&sc.d_output, m_output_bytes_per_frame)); } } @@ -438,7 +443,7 @@ class ClusterFinderCUDAGraph { auto h_sum = m_pedestal.get_sum(); auto h_sum2 = m_pedestal.get_sum2(); - using DPT = device::DEVICE_PED_TYPE; + using DPT = detail::DEVICE_PED_TYPE; const double n = static_cast(m_pedestal.n_samples()); // Capture the frozen per-pixel baseline X0 (≈ mean at t=0) ONCE, then @@ -563,7 +568,7 @@ class ClusterFinderCUDAGraph { cudaKernelNodeParams kp = {}; // NOLINTNEXTLINE(cppcoreguidelines-pro-type-reinterpret-cast) kp.func = reinterpret_cast( - device::find_clusters_in_single_frame); kp.gridDim = grid; kp.blockDim = block; @@ -595,4 +600,4 @@ class ClusterFinderCUDAGraph { } }; -} // namespace aare +} // namespace aare::cuda diff --git a/include/aare/RAII_device.hpp b/include/aare/RAII_device.hpp index 5696e783..4aec518e 100644 --- a/include/aare/RAII_device.hpp +++ b/include/aare/RAII_device.hpp @@ -3,7 +3,7 @@ #include #include -namespace aare { +namespace aare::cuda { template struct CudaFreeDeleter { void operator()(T *p) const noexcept { @@ -52,4 +52,4 @@ inline Stream make_stream() { return Stream{raw_ptr}; } -} // namespace aare \ No newline at end of file +} // namespace aare::cuda \ No newline at end of file diff --git a/include/aare/clusterfinder_algo.cuh b/include/aare/clusterfinder_algo.cuh index 95ad19d4..511de96a 100644 --- a/include/aare/clusterfinder_algo.cuh +++ b/include/aare/clusterfinder_algo.cuh @@ -43,7 +43,10 @@ namespace aare::cuda { * * Algo (the search algorithm, e.g. FixedWindow): * cluster_type, frame_type, pedestal pedestal must satisfy the above - * LaunchParams per-launch algorithm scalars + * compute_type floating-point type of nSigma; the driver stores it + * LaunchParams aggregate, brace-initialised by the driver as + * {int32_t nrows, int32_t ncols, compute_type nSigma, + * uint32_t max_clusters} * Output owning, move-only device output * block Output Output::allocate(uint32_t max_clusters_per_frame) static * constexpr size_t Output::clusters_offset static size_t @@ -83,6 +86,9 @@ template struct require_algo { "clusters_offset' and 'static size_t bytes_per_frame(uint32_t)': the " "driver sizes its pinned host slots from them, before any Output " "exists."); + static_assert(std::is_floating_point::value, + "Algo::compute_type is the nSigma threshold type: an " + "integral type compiles and truncates nSigma (3.5 -> 3)."); static constexpr bool value = true; }; @@ -105,7 +111,7 @@ template struct require_algo { * and the caller supplies it. */ struct CenteredRunningPedestal { - using DevicePedT = device::DEVICE_PED_TYPE; + using DevicePedT = detail::DEVICE_PED_TYPE; struct Buffers { DeviceBuffer mean; @@ -282,6 +288,8 @@ struct FixedWindow { using cluster_type = ClusterType; using frame_type = FrameType; using pedestal = Ped; + using compute_type = detail::COMPUTE_TYPE; // exposed here but fixed by the + // kernel (not a knob) /// The output block one frame writes: a count word followed by the /// clusters, in one allocation. Owns the device copy and knows how to read @@ -336,7 +344,7 @@ struct FixedWindow { struct LaunchParams { int32_t nrows, ncols; - float nSigma; + compute_type nSigma; uint32_t max_clusters; }; @@ -349,13 +357,13 @@ struct FixedWindow { constexpr int row_radius = ClusterType::cluster_size_y / 2; constexpr size_t shmem_bytes = (BLOCK_X + 2 * col_radius) * (BLOCK_Y + 2 * row_radius) * - sizeof(device::COMPUTE_TYPE); + sizeof(compute_type); const dim3 block(BLOCK_X, BLOCK_Y); const dim3 grid( (static_cast(lp.ncols) + BLOCK_X - 1) / BLOCK_X, (static_cast(lp.nrows) + BLOCK_Y - 1) / BLOCK_Y); - device::find_clusters_in_single_frame <<>>( d_frame, pv.d_pd_mean, pv.d_pd_sum, pv.d_pd_sum2, pv.d_pd_off, diff --git a/include/aare/clusterfinder_kernel.cuh b/include/aare/clusterfinder_kernel.cuh index c490c6d1..ef7b43e4 100644 --- a/include/aare/clusterfinder_kernel.cuh +++ b/include/aare/clusterfinder_kernel.cuh @@ -5,7 +5,7 @@ #include #include -namespace aare::device { +namespace aare::cuda::detail { /// Stencil arithmetic precision. Set both aliases to double for an f64 build. /// float is safe here because the device pedestal stores centered moments @@ -288,4 +288,4 @@ __global__ void find_clusters_in_single_frame( } } -} // namespace aare::device \ No newline at end of file +} // namespace aare::cuda::detail \ No newline at end of file diff --git a/python/src/bind_ClusterFinderCUDA.hpp b/python/src/bind_ClusterFinderCUDA.hpp index 4529e61f..4b81a857 100644 --- a/python/src/bind_ClusterFinderCUDA.hpp +++ b/python/src/bind_ClusterFinderCUDA.hpp @@ -29,8 +29,8 @@ void define_ClusterFinderCUDA(py::module &m, const std::string &typestr) { using ClusterType = Cluster; // The driver is templated on the algorithm; a second algorithm is a second // instantiation registered under its own suffix (see ClusterFinder.py). - using Algo = aare::cuda::FixedWindow; - using CF = ClusterFinderCUDA; + using Algo = cuda::FixedWindow; + using CF = cuda::ClusterFinderCUDA; using ContigArr = py::array_t; diff --git a/python/src/bind_ClusterFinderCUDAGraph.hpp b/python/src/bind_ClusterFinderCUDAGraph.hpp index a7af6d3b..972419b7 100644 --- a/python/src/bind_ClusterFinderCUDAGraph.hpp +++ b/python/src/bind_ClusterFinderCUDAGraph.hpp @@ -26,7 +26,7 @@ void define_ClusterFinderCUDAGraph(py::module &m, const std::string &typestr) { auto class_name = fmt::format("ClusterFinderCUDAGraph_{}", typestr); using ClusterType = Cluster; - using CF = ClusterFinderCUDAGraph; + using CF = cuda::ClusterFinderCUDAGraph; using ContigArr = py::array_t; diff --git a/src/ClusterFinderCUDA_testimpl.cu b/src/ClusterFinderCUDA_testimpl.cu index ed49622b..9e5e49e1 100644 --- a/src/ClusterFinderCUDA_testimpl.cu +++ b/src/ClusterFinderCUDA_testimpl.cu @@ -78,10 +78,16 @@ ManualVsDriver> run(int nrows, int ncols, int n_ped, // ---- Route 1: caller-owned buffers, nothing from the driver ---------- { - Pedestal ped(nrows, ncols); - for (int f = 0; f < n_ped; ++f) - ped.push(NDView( - const_cast(s.ped.data() + f * npx), shape)); + FastPedestal ped(nrows, ncols); + for (int f = 0; f < n_ped; ++f) { + NDView v( + const_cast(s.ped.data() + f * npx), shape); + // Same sequence as ClusterFinderCUDA::push_pedestal_frame. + if (!ped.ready()) + ped.add_init_frame(v); + else + ped.push_ema(v); + } using Ped = typename Algo::pedestal; const auto img = Ped::prepare_pedestal(ped); @@ -147,7 +153,7 @@ ManualVsDriver> run(int nrows, int ncols, int n_ped, // ---- Route 2: the driver. One stream so the per-frame pedestal update // follows the same sequence as route 1. -------------------------------- { - ClusterFinderCUDA cf(shape, nSigma, max_clusters, 1); + cuda::ClusterFinderCUDA cf(shape, nSigma, max_clusters, 1); for (int f = 0; f < n_ped; ++f) cf.push_pedestal_frame(NDView( const_cast(s.ped.data() + f * npx), shape));