API reference — reduce

API reference — reduce#

GPU (CUDA) and CPU (OpenMP) parallel reductions. The two engines live in the device namespaces as eagle::cuda::Reduction and eagle::cpu::Reduction; eagle::reduce keeps the shared kernels and the native graph node.

eagle::cuda::Reduction#

struct Reduction

GPU parallel reduction over aether scalar arrays.

All overloads of reduceBlocking synchronise the CUDA stream before returning the result on the host. For graph-based (non-blocking) variants, use the overloads that accept a cuda::Graph&.

Note

Available only when EAGLE_CPU_ONLY is not defined.

Public Static Functions

template<typename T>
static inline aether::Array<T> makeBuffer(const idx_t &size)

Generate a work buffer for use with reduceBlocking, of the appropriate size to perform a reduction over arr.

Note: the returned array is only allocated on the host. You will need to call oarr.upload before invoking reduceBlocking.

template<typename T>
static inline aether::Array<T> makeBuffer(const aether::Array<T> &arr)

Generate a work buffer for use with reduceBlocking, of the appropriate size to perform a reduction over arr.

Note: the returned array is only allocated on the host. You will need to call oarr.upload before invoking reduceBlocking.

template<typename T, typename OP>
static inline T reduceBlocking(aether::Array<T> &arr, aether::Array<T> &oarr, const T &init = 0, const cudaStream_t &stream = 0)

Perform a reduction over a GPU-allocated array using a pre-allocated work buffer.

Note: causes a blocking synchronization of the given stream.

Template Parameters:

  • T – Data type of the elements

  • OP – Reduction operator to use. Must be callable as OP{}(a, b) (aether’s functor shape). Example: aether::SumOp<T>

Parameters:

  • arr – Input scalar array. The reduction will be performed over the device copy, but it will not be modified. Must be already allocated on the device.

  • oarr – Output array. Used for intermediate copies. Must be allocated on host and device. Its values on both host and device will be overwritten. It is recommended to generate it via makeBuffer to ensure it has the appropriate size. Must be distinct from arr.

  • init – Initial value for the accumulation.

  • stream – CUDA stream to use.

template<typename T, typename OP>
static inline T reduceBlocking(aether::Array<T> &arr, const T &init = 0, const cudaStream_t &stream = 0)

Perform a reduction over a GPU-allocated array, allocating a work buffer ad-hoc.

Note: causes a blocking synchronization of the given stream.

Note that this will allocate a new, temporary work buffer on host and device. If performing multiple reductions with same-sized arrays, it may be more efficient to generate the buffer once with makeBuffer and reuse it with the overloaded version of this function.

Template Parameters:

  • T – Data type of the elements

  • OP – Reduction operator to use. Must be callable as OP{}(a, b) (aether’s functor shape). Example: aether::SumOp<T>

Parameters:

  • arr – Input scalar array. The reduction will be performed over the device copy, but it will not be modified. Must be already allocated on the device.

  • init – Initial value for the accumulation.

  • stream – CUDA stream to use.

template<typename T, typename OP>
static inline T reduceBlocking(const CRefArrT<T> &arr, aether::Array<T> &oarr, const T &init = 0, const cudaStream_t &stream = 0)

Perform a reduction over a GPU-allocated array using a pre-allocated work buffer.

Note: causes a blocking synchronization of the given stream.

Template Parameters:

  • T – Data type of the elements

  • OP – Reduction operator to use. Must be callable as OP{}(a, b) (aether’s functor shape). Example: aether::SumOp<T>

Parameters:

  • arr – Input scalar array. The reduction will be performed over the device copy, but it will not be modified. Must be already allocated on the device.

  • oarr – Output array. Used for intermediate copies. Must be allocated on host and device. Its values on both host and device will be overwritten. It is recommended to generate it via makeBuffer to ensure it has the appropriate size. Must be distinct from arr.

  • init – Initial value for the accumulation.

  • stream – CUDA stream to use.

template<typename T, typename OP, bool DirectCopy = false>
static inline void reduceBlocking(cuda::Graph &graph, T *result, const CRefArrT<T> &arr, aether::Array<T> &oarr, const idx_t N, const T &init = 0, const cudaStream_t &stream = 0)

Build a graph to perform a reduction over a GPU-allocated array using a pre-allocated work buffer.

Note: causes a blocking synchronization of the given stream.

Template Parameters:

  • T – Data type of the elements

  • OP – Reduction operator to use. Must be callable as OP{}(a, b) (aether’s functor shape). Example: aether::SumOp<T>

Parameters:

  • result – Pointer to the reduction result

  • arr – Input scalar array. The reduction will be performed over the device copy, but it will not be modified. Must be already allocated on the device.

  • oarr – Output array. Used for intermediate copies. Must be allocated on host and device. Its values on both host and device will be overwritten. It is recommended to generate it via makeBuffer to ensure it has the appropriate size. Must be distinct from arr.

  • init – Initial value for the accumulation.

  • stream – CUDA stream to use.

template<typename T, typename OP, bool DirectCopy = false>
static inline void reduceBlocking(cuda::Graph &graph, T *result, const CRefArrT<T> &arr, aether::Array<T> &oarr, const T &init = 0, const cudaStream_t &stream = 0)

Build a graph to perform a reduction over a GPU-allocated array using a pre-allocated work buffer.

Note: causes a blocking synchronization of the given stream.

Template Parameters:

  • T – Data type of the elements

  • OP – Reduction operator to use. Must be callable as OP{}(a, b) (aether’s functor shape). Example: aether::SumOp<T>

Parameters:

  • result – Pointer to the reduction result

  • arr – Input scalar array. The reduction will be performed over the device copy, but it will not be modified. Must be already allocated on the device.

  • oarr – Output array. Used for intermediate copies. Must be allocated on host and device. Its values on both host and device will be overwritten. It is recommended to generate it via makeBuffer to ensure it has the appropriate size. Must be distinct from arr.

  • init – Initial value for the accumulation.

  • stream – CUDA stream to use.

template<typename T, typename OP>
static inline cuda::Graph reduceBlocking(T *result, const CRefArrT<T> &arr, aether::Array<T> &oarr, const T &init = 0, const cudaStream_t &stream = 0)

Build a graph to perform a reduction over a GPU-allocated array using a pre-allocated work buffer.

Note: causes a blocking synchronization of the given stream.

Template Parameters:

  • T – Data type of the elements

  • OP – Reduction operator to use. Must be callable as OP{}(a, b) (aether’s functor shape). Example: aether::SumOp<T>

Parameters:

  • result – Pointer to the reduction result

  • arr – Input scalar array. The reduction will be performed over the device copy, but it will not be modified. Must be already allocated on the device.

  • oarr – Output array. Used for intermediate copies. Must be allocated on host and device. Its values on both host and device will be overwritten. It is recommended to generate it via makeBuffer to ensure it has the appropriate size. Must be distinct from arr.

  • init – Initial value for the accumulation.

  • stream – CUDA stream to use.

eagle::cpu::Reduction#

template<typename T, typename OP>
struct Reduction#

CPU parallel reduction over aether scalar arrays using OpenMP.

Cache-line padding (alignas(64) on PaddedT) prevents false sharing across threads in the per-thread accumulator buffer.

Template Parameters:

  • T – Element type of the array being reduced.

  • OP – Reduction operator; must be CALLABLE as OP{}(a, b) (aether’s functor shape — there is no static OP::eval), or be one of aether::SumOp<T>, aether::MaxOp<T> or aether::LogicalAndOp (which use native OpenMP reduction clauses for better performance).

Public Static Functions

static inline T reduce(const CRefArrT<T> &in, std::vector<PaddedT> &buffer, const T &init = 0)#

Run the OpenMP based reduction.

static inline T reduce(const CRefArrT<T> &in, const T &init = 0)#

Allocate buffer on the fly and run the OpenMP based reduction.

struct PaddedT#

Namespace eagle::reduce#

namespace reduce#
template<typename T, typename OP>
class ReductionNode#
#include <ReductionNode.h>

A reduction as a native node — the arena-backed twin of cuda::Reduction::reduceBlocking, with both a CUDA and a host face.

Reduces the input array into a single result. The CUDA face (buildInto, EAGLE_CPU_ONLY undefined) captures the same multi-level detail::reduceOnce chain the free builder uses; the host face (runHost, always compiled) runs the OpenMP twin cpu::Reduction. Both share reserveScratch.

Contributed transparently with graph.addNative(ReductionNode{...}, deps) (CUDA Graph) or hostGraph.addNative(ReductionNode{...}, deps) (cpu::Graph). The node holds only PODs (result pointer, input handle, size, init, stream, scratch handle), so it is copyable — as addNative requires.

Template Parameters:

  • T – element type.

  • OP – reduction operator (as reduceOnce / cpu::Reduction).

Public Types

using HandleT = CRefArrT<T>#

The one non-owning reference tier (aether’s View).

using StreamT = nativeStream_t#

Mode-agnostic stream type (cudaStream_t under CUDA, int in EAGLE_CPU_ONLY).

Public Functions

inline ReductionNode(T *result, const HandleT &input, idx_t n, T init = 0, const StreamT &stream = 0)#
Parameters:

  • result – Host pointer the single reduced value is copied into.

  • input – Read-only device view over the array to reduce (e.g. arr.deviceRef().as_const()); must outlive the launch.

  • n – Element count of input.

  • init – Initial accumulator value.

  • stream – Stream the reduce chain is captured on.

inline void reserveScratch(ScratchArena &arena)#

Phase 1: reserve the reduction work buffer (sized like cuda::Reduction::makeBuffer — at least nBlocks elements).

inline idx_t buildInto(cuda::Graph &g, const std::vector<idx_t> &deps)#

Phase 2 (CUDA): capture the multi-level reduce chain (first node depends on deps, the rest auto-chain) plus the D2H result copy, using the committed arena’s pointer as the work buffer. Returns the last node.

inline void runHost(cpu::Graph&)#

Host face: run the OpenMP reduction twin over input_ and store the result. cpu::Reduction owns its own per-thread padded accumulators (outer thread loop over samples), so no arena scratch is touched — the reserved slot simply stays unused on the host path.

namespace detail#

Functions

template<typename T, typename OP, unsigned int blockSize_>
void reduceOnce(CRefArrT<T> g_idata, GRefArrT<T> g_odata, unsigned int n, const T init)#

Run a single reduction.

Block size is FIXED by the warp-shuffle reduction layout: the recursive sdata[threadIdx.x + offset] pattern is only valid when blockDim equals the templated blockSize_ exactly. NEVER feed this kernel through Launcher::setLogicalSize (which mutates blockDim). SASS register footprint: ~8 regs (sm_61).

struct Core#
#include <detail.h>

Reductions.

Public Static Functions

template<typename T, typename OP>
static inline T warpShuffleReduce(T val, unsigned int mask = 0xffffffffu)#

Warp reduction using shuffle intrinsics.