Namespace to contain device resident reduction functions. Includes functions and kernels for array reduction, warp level, block level, and grid level reductions.
More...
|
| __global__ void | 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.
|
| |
| __device__ Real | warpReduceMax (Real val) |
| | Perform a reduction within the warp/wavefront to find the maximum value of val
|
| |
| __device__ Real | blockReduceMax (Real val) |
| | Perform a reduction within the block to find the maximum value of val
|
| |
| __device__ int | encode (float val) |
| | Encode a float as an int.
|
| |
| __device__ long long | encode (double val) |
| | Encode a double as a long long int.
|
| |
| __device__ float | decode (int val) |
| | Decodes an int as a float.
|
| |
| __device__ double | decode (long long val) |
| | Decodes a long long int as a double.
|
| |
| __device__ float | atomicMaxBits (float *address, float val) |
| | Perform an atomic reduction to find the maximum value of val
|
| |
| __device__ double | atomicMaxBits (double *address, double val) |
| | Perform an atomic reduction to find the maximum value of val
|
| |
| __device__ float | atomicMinBits (float *address, float val) |
| | Perform an atomic reduction to find the minimum value of val
|
| |
| __device__ double | atomicMinBits (double *address, double val) |
| | Perform an atomic reduction to find the minimum value of val
|
| |
| __device__ void | 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.
|
| |
|
template<std::size_t N, std::size_t Blocksize> |
| __device__ void | blockAccumulateIntoNReals (Real *__restrict__ dest, Real *__restrict__ src_shared) |
| |
Namespace to contain device resident reduction functions. Includes functions and kernels for array reduction, warp level, block level, and grid level reductions.
◆ atomicMaxBits() [1/2]
| __device__ double reduction_utilities::atomicMaxBits |
( |
double * |
address, |
|
|
double |
val |
|
) |
| |
|
inline |
Perform an atomic reduction to find the maximum 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 maximum of across the grid. Typically this should be a partial reduction that has already been reduced to the block level |
◆ atomicMaxBits() [2/2]
| __device__ float reduction_utilities::atomicMaxBits |
( |
float * |
address, |
|
|
float |
val |
|
) |
| |
|
inline |
Perform an atomic reduction to find the maximum 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 maximum of across the grid. Typically this should be a partial reduction that has already been reduced to the block level |
◆ atomicMinBits() [1/2]
| __device__ double reduction_utilities::atomicMinBits |
( |
double * |
address, |
|
|
double |
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 across the grid. Typically this should be a partial reduction that has already been reduced to the block level |
◆ atomicMinBits() [2/2]
| __device__ float reduction_utilities::atomicMinBits |
( |
float * |
address, |
|
|
float |
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 across the grid. Typically this should be a partial reduction that has already been reduced to the block level |
◆ blockReduceMax()
| __device__ Real reduction_utilities::blockReduceMax |
( |
Real |
val | ) |
|
|
inline |
Perform a reduction within the block to find the maximum value of val
- Parameters
-
| [in] | val | The thread local variable to find the maximum of across the block |
- Returns
- Real The maximum value of
val within the block
◆ decode() [1/2]
| __device__ float reduction_utilities::decode |
( |
int |
val | ) |
|
|
inline |
Decodes an int as a float.
- Parameters
-
- Returns
- float The decoded float
◆ decode() [2/2]
| __device__ double reduction_utilities::decode |
( |
long long |
val | ) |
|
|
inline |
Decodes a long long int as a double.
- Parameters
-
| val | The long long to decode |
- Returns
- double The decoded double
◆ encode() [1/2]
| __device__ long long reduction_utilities::encode |
( |
double |
val | ) |
|
|
inline |
Encode a double as a long long int.
- Parameters
-
- Returns
- long long The encoded long long int
◆ encode() [2/2]
| __device__ int reduction_utilities::encode |
( |
float |
val | ) |
|
|
inline |
Encode a float as an int.
- Parameters
-
- Returns
- int The encoded int
◆ gridReduceMax()
| __device__ void reduction_utilities::gridReduceMax |
( |
Real |
val, |
|
|
Real * |
out |
|
) |
| |
|
inline |
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.
This function can perform a reduction to find the maximum of the thread local variable val across the entire grid. It relies on a warp-wise reduction using registers followed by a block-wise reduction using shared memory, and finally a grid-wise reduction using atomics. As a result the performance of this function is substantally improved by using as many threads per block as possible and as few blocks as possible since each block has to perform an atomic operation. To accomplish this it is reccommened that you use the AutomaticLaunchParams functions to get the optimal number of blocks and threads per block to launch rather than relying on Cholla defaults and then within the kernel using a grid-stride loop to make sure the kernel works with any combination of threads and blocks. Note that after this function call you cannot use the reduced value in global memory since there is no grid wide sync. You can get around this by either launching a second kernel to do the next steps or by using cooperative groups to perform a grid wide sync. During it's execution it also calls multiple __synchThreads and so cannot be called from within any kind of thread guard.
- Parameters
-
| [in] | val | The thread local variable to find the maximum of across the grid |
| [out] | out | The pointer to where to store the reduced scalar value in device memory |
◆ kernelReduceMax()
| __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.
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.
- Parameters
-
| [in] | in | The pointer to the array to reduce in device memory |
| [out] | out | The pointer to where to store the reduced scalar value in device memory |
| [in] | N | The size of the in array |
◆ warpReduceMax()
| __device__ Real reduction_utilities::warpReduceMax |
( |
Real |
val | ) |
|
|
inline |
Perform a reduction within the warp/wavefront to find the maximum value of val
- Parameters
-
| [in] | val | The thread local variable to find the maximum of across the warp |
- Returns
- Real The maximum value of
val within the warp
