Backend-neutral wrappers for device-wide scan/reduce/sort algorithms. More...

#include "backend/api.h"
#include "backend/types.h"
#include "cuda_check.h"
#include "mem/mempool.h"
#include <thrust/device_ptr.h>
#include <thrust/execution_policy.h>
#include <thrust/scan.h>
#include <cstddef>

Classes
struct	fz::backend::TempStorage

Namespaces
namespace	fz

Functions
auto	fz::backend::detail::parOn (fz::stream_t stream)

template<typename Fn >
TempStorage	fz::backend::withTempStorage (MemoryPool pool, fz::stream_t stream, const char tag, Fn &&fn)

void	fz::backend::freeTempStorage (MemoryPool *pool, const TempStorage &temp, fz::stream_t stream)

template<typename T >
void	fz::backend::exclusiveScan (fz::stream_t stream, T d_in, T d_out, size_t n)

Detailed Description

Backend-neutral wrappers for device-wide scan/reduce/sort algorithms.

FZGPUModules stages lean on cub's Device-level primitives (DeviceScan, DeviceReduce, DeviceRadixSort), all of which share the same two-call shape: query the required scratch bytes, allocate scratch (preferring the pipeline's MemoryPool, falling back to a raw device allocation), then run the real call. withTempStorage() below collapses that repeated query/allocate/fallback dance into one call; the actual algorithm invocation (still cub::Device* under the CUDA backend) stays at the call site via a lambda, so this header does not need to know about cub itself.

exclusiveScan() wraps the two ADM-stage call sites that use thrust::exclusive_scan directly (no scratch-buffer dance — thrust manages its own temporary storage internally).

CUDA and HIP are implemented. The two share one implementation: the allocation calls below are spelled with their CUDA names and re-pointed at HIP by backend/api.h, and hipCUB mirrors cub's API closely enough that the cub::Device* calls left at the call sites are a namespace swap. The only genuine divergence is thrust's execution policy (thrust::cuda::par vs rocThrust's thrust::hip::par), isolated in detail::parOn() below.

SYCL is added when that backend lands (oneDPL for exclusiveScan(), since oneDPL has no direct cub::Device* analogue).

Function Documentation

◆ parOn()

auto fz::backend::detail::parOn ( fz::stream_t stream )

inline

The active backend's stream-bound thrust execution policy.

◆ withTempStorage()

template<typename Fn >

TempStorage fz::backend::withTempStorage	(	MemoryPool *	pool,
		fz::stream_t	stream,
		const char *	tag,
		Fn &&	fn
	)

Runs a cub-shaped two-call device algorithm: fn(nullptr, bytes) sizes the scratch buffer, bytes of scratch is allocated (via pool if non-null and it succeeds, else a raw device allocation), then fn(d_temp, bytes) runs the real call.

Returns the scratch pointer plus its provenance — freeing it via freeTempStorage() is the caller's responsibility, since call sites differ on when it's safe to free (some free immediately after the call, some defer until after a downstream kernel that also reads scratch-adjacent output).

◆ freeTempStorage()

void fz::backend::freeTempStorage	(	MemoryPool *	pool,
		const TempStorage &	temp,
		fz::stream_t	stream
	)

inline

Frees storage returned by withTempStorage(), routing to the pool or a raw free as appropriate.

◆ exclusiveScan()

template<typename T >

void fz::backend::exclusiveScan	(	fz::stream_t	stream,
		T *	d_in,
		T *	d_out,
		size_t	n
	)

Exclusive prefix sum over n elements. No scratch-buffer management needed.

Classes