reduction_utilities.h File Reference

Contains the declaration of the GPU resident reduction utilities. More...

#include <climits>
#include <cstdint>
#include "../global/global.h"
#include "../global/global_cuda.h"
#include "../utils/gpu.hpp"

Include dependency graph for reduction_utilities.h:

This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Namespaces
namespace	reduction_utilities
	Namespace to contain device resident reduction functions. Includes functions and kernels for array reduction, warp level, block level, and grid level reductions.

Functions
template<class To , class From >
__device__ constexpr To	reduction_utilities::backport::bit_cast (const From &from) noexcept
	Do a device side bit cast.
long long __device__	reduction_utilities::backport::atomicMin (long long *address, long long val)
	Perform an atomic reduction to find the minimum value of val
__device__ Real	reduction_utilities::warpReduceMax (Real val)
	Perform a reduction within the warp/wavefront to find the maximum value of val
__device__ Real	reduction_utilities::blockReduceMax (Real val)
	Perform a reduction within the block to find the maximum value of val
__device__ int	reduction_utilities::encode (float val)
	Encode a float as an int.
__device__ long long	reduction_utilities::encode (double val)
	Encode a double as a long long int.
__device__ float	reduction_utilities::decode (int val)
	Decodes an int as a float.
__device__ double	reduction_utilities::decode (long long val)
	Decodes a long long int as a double.
__device__ float	reduction_utilities::atomicMaxBits (float *address, float val)
	Perform an atomic reduction to find the maximum value of val
__device__ double	reduction_utilities::atomicMaxBits (double *address, double val)
	Perform an atomic reduction to find the maximum value of val
__device__ float	reduction_utilities::atomicMinBits (float *address, float val)
	Perform an atomic reduction to find the minimum value of val
__device__ double	reduction_utilities::atomicMinBits (double *address, double val)
	Perform an atomic reduction to find the minimum value of val
__device__ void	reduction_utilities::gridReduceMax (Real val, Real *out)
	Perform a reduction within the grid to find the maximum value of val. Note that the value of out should be set appropriately before the kernel launch that uses this function to avoid any potential race condition; the cuda_utilities::setScalarDeviceMemory function exists for this purpose. of val. Note that the value of out should be set appropriately before the kernel launch that uses this function to avoid any potential race condition; the cuda_utilities::setScalarDeviceMemory function exists for this purpose.
__global__ void	reduction_utilities::kernelReduceMax (Real *in, Real *out, size_t N)
	Find the maximum value in the array. Make sure to initialize out correctly before using this kernel; the cuda_utilities::setScalarDeviceMemory function exists for this purpose. If in and out are the same array that's ok, all the loads are completed before the overwrite occurs.
template<std::size_t N, std::size_t Blocksize>
__device__ void	reduction_utilities::blockAccumulateIntoNReals (Real *__restrict__ dest, Real *__restrict__ src_shared)

Detailed Description

Contains the declaration of the GPU resident reduction utilities.

Author

Robert 'Bob' Caddy (rvc@p.nosp@m.itt..nosp@m.edu)

Function Documentation

atomicMin()

long long __device__ reduction_utilities::backport::atomicMin ( long long *  address, long long  val  )

inline

Perform an atomic reduction to find the minimum value of val

Parameters

[out]	address	The pointer to where to store the reduced scalar value in device memory
[in]	val	The thread local variable to find the minimum of

bit_cast()

template<class To , class From >

__device__ constexpr To reduction_utilities::backport::bit_cast ( const From &  from )

constexprnoexcept

Do a device side bit cast.

Template Parameters

To	The output type
From	The input type

Parameters

from	The input value

Returns

To The bit cast version of From as type To