Policies¶

RAJA kernel execution methods take an execution policy type template parameter to specialize execution behavior. Typically, the policy indicates which programming model back-end to use and other information about the execution pattern, such as number of CUDA threads per thread block, whether execution is synchronous or asynchronous, etc. This section describes RAJA policies for loop kernel execution, scans, sorts, reductions, atomics, etc. Please see detailed examples in RAJA Tutorial and Examples for a variety of use cases.

As RAJA functionality evolves, new policies are added and some may be redefined and to work in new ways.

Note

All RAJA policies are in the namespace RAJA.
All RAJA policies have a prefix indicating the back-end implementation that they use; e.g., omp_ for OpenMP, cuda_ for CUDA, etc.

RAJA Loop/Kernel Execution Policies¶

The following tables summarize RAJA policies for executing kernels. Please see notes below policy descriptions for additional usage details and caveats.

Sequential CPU Policies¶

For the sequential CPU back-end, RAJA provides policies that allow developers to have some control over the optimizations that compilers are allowed to apply.

Sequential/SIMD Execution Policies

Works with

Brief description

seq_launch_t

launch

Creates a sequential execution space.

seq_exec

forall, kernel (For), scan, sort

Sequential execution, where the compiler is allowed to apply any optimizations that its heuristics deem beneficial; i.e., no loop decorations (pragmas or intrinsics) are used in the RAJA implementation.

simd_exec

forall, kernel (For), scan

Try to force generation of SIMD instructions via compiler hints in RAJA’s internal implementation.

OpenMP Parallel CPU Policies¶

For the OpenMP CPU multithreading back-end, RAJA has policies that create an OpenMP parallel region and execute a kernel within it. We refer to these as full policies. They are provided to support common OpenMP use cases.

RAJA also provides other OpenMP policies, which we refer to as partial policies, since they need to be used in combination with other policies. Typically, they work by providing an outer policy and an inner policy as a template parameter to the outer policy. These give users flexibility to create more complex execution patterns.

Note

To control the number of threads used by OpenMP, users may set the value of the environment variable ‘OMP_NUM_THREADS’ (which is fixed for duration of run), or call the OpenMP routine ‘omp_set_num_threads(nthreads)’ in their applications, which allows changing the number of threads at run time.

The full policies are described in the following table. Partial policies are described in other tables below.

OpenMP CPU Full Policies

Works with

Brief description

omp_parallel_for_exec

forall, kernel (For), launch (loop), scan, sort

Same as applying the OpenMP pragma ‘omp parallel for schedule(auto)’

omp_parallel_for_static_exec<ChunkSize>

forall, kernel (For)

Same as applying ‘omp parallel for schedule(static, ChunkSize)’

omp_parallel_for_dynamic_exec<ChunkSize>

forall, kernel (For)

Same as applying ‘omp parallel for schedule(dynamic, ChunkSize)’

omp_parallel_for_guided_exec<ChunkSize>

forall, kernel (For)

Same as applying ‘omp parallel for schedule(guided, ChunkSize)’

omp_parallel_for_runtime_exec

forall, kernel (For)

Same as applying ‘omp parallel for schedule(runtime)’

Note

For the OpenMP scheduling policies above that take a ChunkSize parameter, the chunk size is optional. If not provided, the default chunk size that the OpenMP implementation applies is used. For this case, the RAJA policy syntax is omp_parallel_for_{static|dynamic|guided}_exec< >, which will result in the OpenMP pragma omp parallel for schedule({static|dynamic|guided}) being applied.

RAJA provides outer OpenMP CPU policies to create a parallel region in which to execute a kernel. The outer policies require an inner policy that defines how a kernel will execute in parallel inside the region.

OpenMP CPU Outer Policies

Works with

Brief description

omp_launch_t

launch

Creates an OpenMP parallel region. Same as applying ‘omp parallel’ pragma.

omp_parallel_exec<InnerPolicy>

forall, kernel (For), scan

Creates OpenMP parallel region and requires an InnerPolicy. Same as applying ‘omp parallel’ pragma.

The table below summarizes the inner policies that RAJA provides for OpenMP. These policies are passed to the RAJA omp_parallel_exec outer policy as a template argument as described above.

OpenMP CPU Inner Policies

Works with

Brief description

omp_for_exec

forall, kernel (For), launch (loop) scan

Parallel execution within existing parallel region, specifically apply the OpenMP pragma ‘omp for schedule (auto)’ pragma.

omp_for_static_exec<ChunkSize>

forall, kernel (For)

Same as applying ‘omp for schedule(static, ChunkSize)’

omp_for_nowait_static_exec<ChunkSize>

forall, kernel (For)

Same as applying ‘omp for schedule(static, ChunkSize) nowait’

omp_for_dynamic_exec<ChunkSize>

forall, kernel (For)

Same as applying ‘omp for schedule(dynamic, ChunkSize)’

omp_for_guided_exec<ChunkSize>

forall, kernel (For)

Same as applying ‘omp for schedule(guided, ChunkSize)’

omp_for_runtime_exec

forall, kernel (For)

Same as applying ‘omp for schedule(runtime)’

omp_parallel_collapse_exec

kernel (Collapse + ArgList)

Use in Collapse statement to parallelize multiple loop levels in loop nest indicated using ArgList

Important

RAJA only provides a nowait policy option for static scheduling since that is the only case in which nowait can be used and be correct in general when executing multiple loops in a single parallel region. Paraphrasing the OpenMP standard: programs that depend on which thread executes a particular loop iteration under any circumstance other than static schedule are non-conforming.

Note

As in the RAJA full policies for OpenMP scheduling, the ChunkSize is optional. If not provided, the default chunk size that the OpenMP implementation applies is used.

Note

As noted above, RAJA inner OpenMP policies must be used within an existing parallel region to work properly. Embedding an inner policy inside the RAJA outer omp_parallel_exec will allow you to apply the OpenMP execution prescription specified by the policies to a single kernel. To support use cases with multiple kernels inside an OpenMP parallel region, RAJA provides a region construct that takes a template argument to specify the execution back-end. For example:

RAJA::region<RAJA::omp_parallel_region>([=]() {

  RAJA::forall<RAJA::omp_for_nowait_static_exec< > >(segment,
    [=] (int idx) {
      // do something at iterate 'idx'
    }
  );

  RAJA::forall<RAJA::omp_for_static_exec< > >(segment,
    [=] (int idx) {
      // do something else at iterate 'idx'
    }
  );

});

Here, the RAJA::region<RAJA::omp_parallel_region> method call creates an OpenMP parallel region, which contains two RAJA::forall kernels. The first uses the RAJA::omp_for_nowait_static_exec< > policy, meaning that no thread synchronization is needed after the kernel. Thus, threads can start working on the second kernel while others are still working on the first kernel. In general, this will be correct when the iteration segments used in the two kernels are the same and their are no loop carried dependences in either kernel. Static scheduling is applied to both kernels. The second kernel uses the RAJA::omp_for_static_exec policy (without ‘no wait’ clause), which means that all threads will complete before the kernel exits. In this example, this is not really needed since there is no more code to execute in the parallel region and the RAJA::omp_parallel_region construct applies a barrier at the end of it.

GPU Policies for CUDA and HIP¶

RAJA policies for GPU execution using CUDA or HIP are essentially identical. The only difference is that CUDA policies have the prefix cuda_ and HIP policies have the prefix hip_. The policies in the following table that contain angle brackets indicate template parameters used to specialize execution behavior.

When a CUDA or HIP policy leaves parameters like the block size and/or grid size unspecified, such as the cuda/hip_exec_occ_custom in the table above, a concretizer object is used to decide those parameters. RAJA provides the following concretizers to use with the cuda/hip_exec_occ_custom policies:

Execution Policy	Brief description
Cuda/HipDefaultConcretizer	The default concretizer, expected to provide good performance in general. Note that it may not use max occupancy.
Cuda/HipRecForReduceConcretizer	Expected to provide good performance in loops with reducers. Note that it may not use max occupancy.
Cuda/HipMaxOccupancyConcretizer	Uses max occupancy.
Cuda/HipAvoidDeviceMaxThreadOccupancyConcretizer	Avoids using the max occupancy of the device in terms of threads. Note that it may use the max occupancy of the kernel if that is below the max occupancy of the device.
Cuda/HipFractionOffsetOccupancyConcretizer< Fraction<size_t, numerator, denomenator>, BLOCKS_PER_SM_OFFSET>	Uses a fraction and offset to choose an occupancy based on the max occupancy using the following formula: (Fraction * kernel_max_blocks_per_sm + BLOCKS_PER_SM_OFFSET) * sm_per_device

Several notable constraints apply to RAJA CUDA/HIP direct_unchecked policies.

Note

Direct unchecked policies do not mask out threads that are out-of-range. So they should only be used when the size of the range matches the size of the block or grid.
Repeating direct_unchecked policies with the same dimension in perfectly nested loops is not recommended. Your code may do something, but likely will not do what you expect and/or be correct.
If multiple direct_unchecked policies are used in a kernel (using different dimensions), the product of sizes of the corresponding iteration spaces cannot be greater than the maximum allowable threads per block or blocks per grid. Typically, this is 1024 threads per block. Attempting to execute a kernel with more than the maximum allowed threads per block or blocks per grid will cause the CUDA/HIP runtime to complain about illegal launch parameters.
Block-direct-unchecked policies are recommended for most tiled loop patterns. In these cases the CUDA/HIP kernel is launched with the exact number of blocks needed so no iteration space size checking is needed.

Several notable constraints apply to RAJA CUDA/HIP direct policies.

Note

Direct policies mask out threads that are out-of-range. So they should only be used when the size of the range is less than or equal to the size of the block or grid.
Repeating direct policies with the same dimension in perfectly nested loops is not recommended. Your code may do something, but likely will not do what you expect and/or be correct.
If multiple direct policies are used in a kernel (using different dimensions), the product of sizes of the corresponding iteration spaces cannot be greater than the maximum allowable threads per block or blocks per grid. Typically, this is 1024 threads per block. Attempting to execute a kernel with more than the maximum allowed causes the CUDA/HIP runtime to complain about illegal launch parameters.
Global-direct-sized policies are recommended for most loop patterns, but may be inappropriate for kernels using block level synchronization.
Thread-direct policies are recommended only for certain loop patterns, such as block tiling, that produce small fixed size iteration spaces within each block.

Several notes regarding CUDA/HIP loop policies are also good to know.

Note

Loop policies perform a block or grid stride loop. Thus, they can be used when the size of the loop iteration space exceeds the size of the block or grid.
There is no constraint on the product of sizes of the associated loop iteration space.
These polices allow having a larger number of iterates than threads/blocks in the x, y, or z dimension.
The cuda/hip_thread_loop policies are not safe to use with Cuda/HipSyncThreads, use the cuda/hip_thread_syncable_loop<dims…> policies instead. For example cuda_thread_x_loop -> cuda_thread_syncable_loop<named_dim::x>.
CUDA/HIP loop policies are recommended for some loop patterns where a large or unknown sized iteration space is mapped to a small or fixed number of threads.

Finally

Note

CUDA/HIP block-direct-unchecked or block-direct policies may be preferable to block-loop policies in situations where block load balancing may be an issue as the block-direct-unchecked or block-direct policies may yield better performance.

Several notes regarding the CUDA/HIP policy implementation allow you to write more explicit policies.

Note

Policies are a class template like cuda/hip_exec_explicit or cuda/hip_indexer. The various template parameters specify the behavior of the policy.
Policies have a mapping from loop iterations to iterates in the index set via a iteration_mapping enum template parameter. The possible values are DirectUnchecked, Direct, and StridedLoop.
Policies can be safely used with some synchronization constructs via a kernel_sync_requirement enum template parameter. The possible values are none and sync.
Policies get their indices via an iteration getter class template like cuda/hip::IndexGlobal.
Iteration getters can be used with different dimensions via the named_dim enum. The possible values are x, y and z.
Iteration getters know the number of threads per block (block_size) and number of blocks per grid (grid_size) via integer template parameters. These can be positive integers, in which case they must match the number used in the kernel launch. These can also be values of the named_usage enum. The possible values are unspecified and ignored. For example, in cuda_thread_x_direct block_size is unspecified so a runtime number of threads is used, but grid_size is ignored so blocks are ignored when getting indices.

GPU Policies for SYCL¶

Note

SYCL uses C++-style ordering for its work group and global thread dimension/indexing types. This is due, in part, to SYCL’s closer alignment with C++ multi-dimensional indexing, which is “row-major”. This is the reverse of the thread indexing used in CUDA or HIP, which is “column-major”. For example, suppose we have a thread-block or work-group where we specify the shape as (nx, ny, nz). Consider an element in the thread-block or work-group with id (x, y, z). In CUDA or HIP, the element index is x + y * nx + z * nx * ny. In SYCL, the element index is z + y * nz + x * nz * ny.

In terms of the CUDA or HIP built-in variables to support threads, we have:

Thread ID: threadIdx.x/y/z
Block ID: blockIdx.x/y/z
Block dimension: blockDim.x/y/z
Grid dimension: gridDim.x/y/z

The analogues in SYCL are:

Thread ID: sycl::nd_item.get_local_id(2/1/0)
Work-group ID: sycl::nd_item.get_group(2/1/0)
Work-group dimensions: sycl::nd_item.get_local_range().get(2/1/0)
ND-range dimensions: sycl::nd_item.get_group_range(2/1/0)

When using RAJA::launch, thread and block configuration follows CUDA and HIP programming models and is always configured in three-dimensions. This means that SYCL dimension 2 always exists and should be used as one would use the x dimension for CUDA and HIP.

Similarly, RAJA::kernel uses a three-dimensional work-group configuration. SYCL dimension 2 always exists and should be used as one would use the x dimension in CUDA and HIP.

Device policy aliases¶

To simplify transitions between GPU back-ends (CUDA/HIP/SYCL) and reduce downstream preprocessor conditionals, RAJA provides a set of device_* policy aliases that resolve to the active GPU back-end.

In particular, the following aliases are available when building with a GPU device back-end (i.e., when ENABLE_CUDA, ENABLE_HIP, or RAJA_ENABLE_SYCL is enabled):

`device_*` alias coverage¶
Alias family	CUDA	HIP	SYCL	Notes
device_exec*	yes	yes	partial	Block-size templated exec aliases, for example `device_exec<256>`. `device_exec` and `device_exec_async` exist for SYCL; the CUDA/HIP occupancy and reduction variants do not.
device_atomic and device_atomic_explicit<host_policy>	yes	yes	yes	Back-end atomic policy aliases are available on all three back-ends, for example `device_atomic_explicit<RAJA::seq_atomic>`.
device_reduce	yes	yes	yes	Back-end default reduce policy alias.
device_reduce_atomic and device_reduce_base<with_atomic>	yes	yes	no	CUDA/HIP expose tuning and base reduce aliases, for example `device_reduce_base<true>`; SYCL currently only exposes `device_reduce`.
device_multi_reduce_atomic and device_multi_reduce_atomic_low_performance_low_overhead	yes	yes	no	CUDA/HIP expose multi_reduce aliases; SYCL does not currently provide a back-end-equivalent device_* mapping.
device_launch_t	yes	yes	yes	Back-end launch policy alias, for example `device_launch_t<false>`.
device_global_size_{x,y,z}_{direct,direct_unchecked,loop}<N>	yes	yes	partial	Size-templated aliases, for example `device_global_size_x_direct<64>`. SYCL currently exposes only the direct x/y/z aliases; the unchecked and loop variants are CUDA/HIP only.
device_thread_{x,y,z}_{direct,loop}	yes	yes	yes	Single-dimension thread mapping.
device_thread_size_{x,y,z}_{direct,direct_unchecked,loop}<N>	yes	yes	no	Size-templated aliases, for example `device_thread_size_x_direct<128>`. These aliases are declared as compile-time errors under SYCL.
device_block_{x,y,z}_{direct,loop}	yes	yes	yes	Single-dimension block mapping.
device_block_size_{x,y,z}_{direct,direct_unchecked,loop}<N>	yes	yes	no	Size-templated aliases, for example `device_block_size_x_direct<128>`. These aliases are declared as compile-time errors under SYCL.
device_flatten_thread_size_, device_flatten_block_size_, device_flatten_global_size_*	yes	yes	no	Size-templated aliases, for example `device_flatten_thread_size_x_direct<128>`, `device_flatten_block_size_x_direct<128>`, and `device_flatten_global_size_x_direct<128>`. These flattened size aliases exist for CUDA/HIP but are declared as compile-time errors under SYCL.

Important

For SYCL, these aliases use CUDA-like (x,y,z) naming with the standard RAJA mapping described above: x corresponds to SYCL dimension 2, y to SYCL dimension 1, and z to SYCL dimension 0. These build options enable the corresponding internal RAJA_*_ACTIVE compile-time macros used by the implementation. Device aliases that have no SYCL equivalent are intentionally not defined under SYCL as usable policies. Attempting to use them will cause compile time failure so unsupported code paths are caught immediately.

See also the example code examples/device-policy-aliases.cpp.

SYCL Execution Policies	Works with	Brief description
sycl_exec<WORK_GROUP_SIZE>	forall,	Execute loop iterations in a GPU kernel launched with given work group size.
sycl_launch_t	launch	Launches a SYCL kernel, any code express within the lambda is executed on the device.
sycl_global_0<WORK_GROUP_SIZE>	kernel (For)	Map loop iterates directly to GPU global ids in first dimension, one iterate per work item. Group execution into work groups of given size.
sycl_global_1<WORK_GROUP_SIZE>	kernel (For)	Same as above, but map to global ids in second dim
sycl_global_2<WORK_GROUP_SIZE>	kernel (For)	Same as above, but map to global ids in third dim
sycl_global_item_0	launch (loop)	Creates a unique thread id for each thread for dimension 0 of the grid. Same as computing itm.get_group(0) * itm.get_local_range(0) + itm.get_local_id(0).
sycl_global_item_1	launch (loop)	Same as above, but uses threads in dimension 1 Same as computing itm.get_group(1) + itm.get_local_range(1) * itm.get_local_id(1).
sycl_global_item_2	launch (loop)	Same as above, but uses threads in dimension 2 Same as computing itm.get_group(2) + itm.get_local_range(2) * itm.get_local_id(2).
sycl_local_0_direct	kernel (For) launch (loop)	Map loop iterates directly to GPU work items in first dimension, one iterate per work item (see note below about limitations)
sycl_local_1_direct	kernel (For) launch (loop)	Same as above, but map to work items in second dim
sycl_local_2_direct	kernel (For) launch (loop)	Same as above, but map to work items in third dim
sycl_local_0_loop	kernel (For) launch (loop)	Similar to local-1-direct policy, but use a work group-stride loop which doesn’t limit number of loop iterates
sycl_local_1_loop	kernel (For) launch (loop)	Same as above, but for work items in second dimension
sycl_local_2_loop	kernel (For) launch (loop)	Same as above, but for work items in third dimension
sycl_group_0_direct	kernel (For) launch (loop)	Map loop iterates directly to GPU group ids in first dimension, one iterate per group
sycl_group_1_direct	kernel (For) launch (loop)	Same as above, but map to groups in second dimension
sycl_group_2_direct	kernel (For) launch (loop)	Same as above, but map to groups in third dimension
sycl_group_0_loop	kernel (For) launch (loop)	Similar to group-1-direct policy, but use a group-stride loop.
sycl_group_1_loop	kernel (For) launch (loop)	Same as above, but use groups in second dimension
sycl_group_2_loop	kernel (For) launch (loop)	Same as above, but use groups in third dimension

OpenMP Target Offload Policies¶

RAJA provides policies to use OpenMP to offload kernel execution to a GPU device, for example. They are summarized in the following table.

OpenMP Target Execution Policies

Works with

Brief description

omp_target_parallel_for_exec<#>

forall, kernel(For)

Create parallel target region and execute with given number of threads per team inside it. Number of teams is calculated internally; i.e., apply omp teams distribute parallel for num_teams(iteration space size/#) thread_limit(#) pragma

omp_target_parallel_collapse_exec

kernel (Collapse)

Similar to above, but collapse perfectly-nested loops, indicated in arguments to RAJA Collapse statement. Note: compiler determines number of thread teams and threads per team

RAJA IndexSet Execution Policies¶

When an IndexSet iteration space is used in RAJA by passing a RAJA::IndexSet to a RAJA::forall method, an index set execution policy is required. An index set execution policy is a two-level policy: an ‘outer’ policy for iterating over segments in the index set, and an ‘inner’ policy used to execute the iterations defined by each segment. An index set execution policy type has the form:

RAJA::ExecPolicy< segment_iteration_policy, segment_execution_policy >

In general, any policy that can be used with a RAJA::forall method can be used as an (inner) segment execution policy. The following policies are available to use for the outer segment iteration policy:

Execution Policy	Brief description
Serial
seq_segit	Iterate over index set segments sequentially.
OpenMP CPU multithreading
omp_parallel_segit	Create OpenMP parallel region and iterate over segments in parallel inside it; i.e., apply `omp parallel for` pragma on loop over segments.
omp_parallel_for_segit	Same as above.

Parallel Region Policies¶

Earlier, we discussed using the RAJA::region construct to execute multiple kernels in an OpenMP parallel region. To support source code portability, RAJA provides a sequential region concept that can be used to surround code that uses execution back-ends other than OpenMP. This simplifies switching user code between sequential and parallel execution and does not require changing the code structure. For example:

RAJA::region<RAJA::seq_region>([=]() {

   RAJA::forall<RAJA::seq_exec>(segment, [=] (int idx) {
       // do something at iterate 'idx'
   } );

   RAJA::forall<RAJA::seq_exec>(segment, [=] (int idx) {
       // do something else at iterate 'idx'
   } );

 });

Note

The sequential region specialization is essentially a pass through operation. It is provided so that if you want to turn off OpenMP in your code, for example, you can simply replace the region policy type and you do not have to change your algorithm source code.

Reduction Policies¶

Each RAJA reduction object must be defined with a ‘reduction policy’ type. Reduction policy types are distinct from loop execution policy types. It is important to note the following constraint about RAJA reduction usage:

Important

To guarantee correctness, a reduction policy must support the loop execution policy used. For example, a CUDA reduction policy must be used when the execution policy is a CUDA policy. However an OpenMP reduction policy or a CUDA reduction policy may be used when the execution policy is an OpenMP policy, and so on.

Note

It is undefined behavior to use a reducer object with a loop execution policy that does not match the RAJA::Policy enum argument used to most recently setup the reducer object. RAJA::Policy::undefined may be used with any of the loop policies supported by the reduction policy. For example, if a reducer object with a CUDA reduction policy is setup with RAJA::Policy::cuda, but used in a sequential loop, then that is undefined behavior. Using either RAJA::Policy::undefined or RAJA::Policy::sequential is correct.

The following table summarizes RAJA reduction policy types:

Reduction Policy	Loop Policies Supported	Brief description
seq_reduce	seq_exec	Non-parallel (sequential) reduction.
omp_reduce	any OpenMP policy, seq_exec	OpenMP parallel reduction.
omp_reduce_ordered	any OpenMP policy, seq_exec	OpenMP parallel reduction with result guaranteed to be reproducible.
omp_target_reduce	any OpenMP Target policy, seq_exec	OpenMP parallel target offload reduction.
cuda/hip_reduce	any CUDA/HIP policy, any OpenMP policy, seq_exec	Parallel reduction in a CUDA/HIP kernel (device synchronization will occur when reduction value is finalized).
cuda/hip_reduce_atomic	any CUDA/HIP policy, any OpenMP policy, seq_exec	Same as above, but reduction may use atomic operations leading to run to run variability in the results.
cuda/hip_reduce_base<with_atomic>	any CUDA/HIP policy, any OpenMP policy, seq_exec	Choose between cuda/hip_reduce and cuda/hip_reduce_atomic policies based on the with_atomic boolean.
cuda/hip_reduce_device_fence	any CUDA/HIP policy, any OpenMP policy, seq_exec	Same as above, and reduction uses normal memory accesses that are not visible across the whole device and device scope fences to ensure visibility and ordering. This works on all architectures but incurs higher overheads on some architectures.
cuda/hip_reduce_block_fence	any CUDA/HIP policy, any OpenMP policy, seq_exec	Same as above, and reduction uses special memory accesses to a level of cache visible to the whole device and block scope fences to ensure ordering. This improves performance on some architectures.
cuda/hip_reduce_atomic_host_init_device_fence	any CUDA/HIP policy, any OpenMP policy, seq_exec	Same as above with device fence, but initializes the memory used for atomics on the host. This works well on recent architectures and incurs lower overheads.
cuda/hip_reduce_atomic_host_init_block_fence	any CUDA/HIP policy, any OpenMP policy, seq_exec	Same as above with block fence, but initializes the memory used for atomics on the host. This works well on recent architectures and incurs lower overheads.
cuda/hip_reduce_atomic_device_init_device_fence	any CUDA/HIP policy, any OpenMP policy, seq_exec	Same as above with device fence, but initializes the memory used for atomics on the device. This works on all architectures but incurs higher overheads.
cuda/hip_reduce_atomic_device_init_block_fence	any CUDA/HIP policy, any OpenMP policy, seq_exec	Same as above with block fence, but initializes the memory used for atomics on the device. This works on all architectures but incurs higher overheads.
sycl_reduce	any SYCL policy, seq_exec	Reduction in a SYCL kernel (device synchronization will occur when the reduction value is finalized).

Important

RAJA reductions used with SIMD execution policies are not guaranteed to generate correct results. So they should not be used for kernels containing reductions.

MultiReduction Policies¶

Each RAJA multi-reduction object must be defined with a ‘multi-reduction policy’ type. Multi-reduction policy types are distinct from loop execution policy types. It is important to note the following constraints about RAJA multi-reduction usage:

Important

To guarantee correctness, a multi-reduction policy must support the loop execution policy used. For example, a CUDA multi-reduction policy must be used when the execution policy is a CUDA policy. However an OpenMP multi-reduction policy or a CUDA multi-reduction policy may be used when the execution policy is an OpenMP policy, and so on.

Note

It is undefined behavior to use a multi-reducer object with a loop execution policy that does not match the RAJA::Policy enum argument used to most recently setup the multi-reducer object. RAJA::Policy::undefined may be used with any of the loop policies supported by the multi-reduction policy. For example, if a multi-reducer object with a CUDA multi-reduction policy is setup with RAJA::Policy::cuda, but used in a sequential loop, then that is undefined behavior. Using either RAJA::Policy::undefined or RAJA::Policy::sequential is correct.

The following table summarizes RAJA multi-reduction policy types:

MultiReduction Policy	Loop Policies Supported	Brief description
seq_multi_reduce	seq_exec	Non-parallel (sequential) multi-reduction.
omp_multi_reduce	any OpenMP policy, seq_exec	OpenMP parallel multi-reduction.
omp_multi_reduce_ordered	any OpenMP policy, seq_exec	OpenMP parallel multi-reduction with result guaranteed to be reproducible.
cuda/hip_multi_reduce_atomic	any CUDA/HIP policy, any OpenMP policy, seq_exec	Parallel multi-reduction in a CUDA/HIP kernel. Multi-reduction may use atomic operations leading to run to run variability in the results. (device synchronization will occur when reduction value is finalized)
cuda/hip_multi_reduce_atomic_low_performance_low_overhead	any CUDA/HIP policy, any OpenMP policy, seq_exec	Same as above, but multi-reduction uses a low overhead algorithm with a minimal set of resources. This minimally effects the performance of loops containing the multi-reducer though it may cause the multi-reducer itself to perform poorly if it is used.
cuda/hip_multi_reduce_atomic_block_then_atomic_grid_host_init	any CUDA/HIP policy, any OpenMP policy, seq_exec	The multi-reduction uses atomics into shared memory and global memory. Atomics into shared memory are used each time a value is combined into the multi-reducer and at the end of the life of the block the shared values are combined into global memory with atomics. If there is not enough shared memory available this will fall back to using atomics into global memory only, which may have a performance penalty. The memory for global atomics is initialized on the host.
cuda/hip_multi_reduce_atomic_global_host_init	any CUDA/HIP policy, any OpenMP policy, seq_exec	The multi-reduction uses atomics into global global memory only. Atomics into global memory are used each time a value is combined into the multi-reducer. The memory for global atomics is initialized on the host.
cuda/hip_multi_reduce_atomic_global_no_replication_host_init	any CUDA/HIP policy, any OpenMP policy, seq_exec	Same as above, but uses minimal memory by not replicating global atomics.

Important

RAJA multi-reductions used with SIMD execution policies are not guaranteed to generate correct results. So they should not be used for kernels containing multi-reductions.

Atomic Policies¶

Each RAJA atomic operation must be defined with an ‘atomic policy’ type. Atomic policy types are distinct from loop execution policy types.

Important

An atomic policy type must be consistent with the loop execution policy for the kernel in which the atomic operation is used.

The following table summarizes RAJA atomic policies and usage.

Atomic Policy	Loop Policies to Use With	Brief description
seq_atomic	seq_exec,	Atomic operation performed in a non-parallel (sequential) kernel.
omp_atomic	any OpenMP policy	Atomic operation in OpenMP multithreading or target kernel; i.e., apply `omp atomic` pragma.
cuda/hip/sycl_atomic	any CUDA/HIP/SYCL policy	Atomic operation performed in a CUDA/HIP/SYCL kernel.
cuda/hip_atomic_explicit	any CUDA/HIP policy	Atomic operation performed in a CUDA/HIP kernel that may also be used in a host execution context. The atomic policy takes a host atomic policy template argument. See additional explanation and example below.
builtin_atomic	seq_exec, any OpenMP policy	Compiler builtin atomic operation.
auto_atomic	seq_exec, any OpenMP policy, any CUDA/HIP/SYCL policy	Atomic operation compatible with loop execution policy. See example below. Cannot be used inside CUDA or HIP explicit atomic policies.

Note

The cuda_atomic_explicit and hip_atomic_explicit policies take a host atomic policy template parameter. They are intended to be used with kernels that are host-device decorated to be used in either a host or device execution context, possibly decided at run time.

Here is an example illustrating use of the cuda_atomic_explicit policy with an OpenMP host policy:

auto kernel = [=] RAJA_HOST_DEVICE (RAJA::Index_type i) {
  RAJA::atomicAdd< RAJA::cuda_atomic_explicit<omp_atomic> >(&sum, 1);
};

RAJA::forall< RAJA::cuda_exec<BLOCK_SIZE> >(RAJA::TypedRangeSegment<int> seg(0, N), kernel);

RAJA::forall< RAJA::omp_parallel_for_exec >(RAJA::TypedRangeSegment<int> seg(0, N), kernel);

In this case, the atomic operation knows when it is compiled for the device in a CUDA kernel context and the CUDA atomic operation is applied. Similarly when it is compiled for the host in an OpenMP kernel the omp_atomic policy is used and the OpenMP version of the atomic operation is applied.

Here is an example illustrating use of the auto_atomic policy:

RAJA::forall< RAJA::cuda_exec<BLOCK_SIZE> >(RAJA::TypedRangeSegment<int> seg(0, N),
  [=] RAJA_DEVICE (RAJA::Index_type i) {

  RAJA::atomicAdd< RAJA::auto_atomic >(&sum, 1);

});

In this case, the atomic operation knows that it is used in a CUDA kernel context and the CUDA atomic operation is applied. Similarly, if an OpenMP execution policy was used, the OpenMP version of the atomic operation would be used.

Note

The builtin_atomic policy may be preferable to the omp_atomic policy in terms of performance.

Local Array Memory Policies¶

RAJA::LocalArray types must use a memory policy indicating where the memory for the local array will live. These policies are described in Local Array.

The following memory policies are available to specify memory allocation for RAJA::LocalArray objects:

RAJA::cpu_tile_mem - Allocate CPU memory on the stack

RAJA::cuda/hip_shared_mem - Allocate CUDA or HIP shared memory

RAJA::cuda/hip_thread_mem - Allocate CUDA or HIP thread private memory

RAJA Kernel Execution Policies¶

RAJA kernel execution policy constructs form a simple domain specific language for composing and transforming complex loops that relies solely on standard C++20 template support. RAJA kernel policies are constructed using a combination of Statements and Statement Lists. A RAJA Statement is an action, such as execute a loop, invoke a lambda, set a thread barrier, etc. A StatementList is an ordered list of Statements that are composed in the order that they appear in the kernel policy to construct a kernel. A Statement may contain an enclosed StatementList. Thus, a RAJA::KernelPolicy type is really just a StatementList.

The main Statement types provided by RAJA are RAJA::statement::For and RAJA::statement::Lambda, that we discussed in Complex Loops (RAJA::kernel). A RAJA::statement::For<ArgID, ExecPolicy, Enclosed Statements> type indicates a for-loop structure. The ArgID parameter is an integral constant that identifies the position of the iteration space in the iteration space tuple passed to the RAJA::kernel method to be used for the loop. The ExecPolicy is the RAJA execution policy to use on the loop, which is similar to RAJA::forall usage. The EnclosedStatements type is a nested template parameter that contains whatever is needed to execute the kernel and which forms a valid StatementList. The RAJA::statement::Lambda<LambdaID> type invokes the lambda expression corresponding to its position ‘LambdaID’ in the sequence of lambda expressions in the RAJA::kernel argument list. For example, a simple sequential for-loop:

for (int i = 0; i < N; ++i) {
  // loop body
}

can be represented using the RAJA kernel interface as:

using KERNEL_POLICY =
  RAJA::KernelPolicy<
    RAJA::statement::For<0, RAJA::seq_exec,
      RAJA::statement::Lambda<0>
    >
  >;

RAJA::kernel<KERNEL_POLICY>(
  RAJA::make_tuple(range),
  [=](int i) {
    // loop body
  }
);

Note

All RAJA::forall functionality can be done using the RAJA::kernel interface. We maintain the RAJA::forall interface since it is less verbose and thus more convenient for users.

RAJA::kernel Statement Types¶

The list below summarizes the current collection of statement types that can be used with RAJA::kernel and RAJA::kernel_param. More detailed explanation along with examples of how they are used can be found in the RAJA::kernel examples in RAJA Tutorial and Examples.

Note

All of the statement types described below are in the namespace RAJA::statement. For brevity, we omit the namespaces in the discussion in this section.

Note

RAJA::kernel_param functions similarly to RAJA::kernel except that the second argument is a tuple of parameters used in a kernel for local arrays, thread local variables, tiling information, etc.

Several RAJA statements can be specialized with auxiliary types, which are described in Auxilliary Types.

The following list contains the most commonly used statement types.

For< ArgId, ExecPolicy, EnclosedStatements > abstracts a for-loop associated with kernel iteration space at tuple index ArgId, to be run with ExecPolicy execution policy, and containing the EnclosedStatements which are executed for each loop iteration.
Lambda< LambdaId > invokes the lambda expression that appears at position ‘LambdaId’ in the sequence of lambda arguments. With this statement, the lambda expression must accept all arguments associated with the tuple of iteration space segments and tuple of parameters (if kernel_param is used).
Lambda< LambdaId, Args...> extends the Lambda statement. The second template parameter indicates which arguments (e.g., which segment iteration variables) are passed to the lambda expression.
Collapse< ExecPolicy, ArgList<...>, EnclosedStatements > collapses multiple perfectly nested loops specified by tuple iteration space indices in ArgList, using the ExecPolicy execution policy, and places EnclosedStatements inside the collapsed loops which are executed for each iteration. Note that this only works for CPU execution policies (e.g., sequential, OpenMP). It may be available for CUDA in the future if such use cases arise.

There is one statement specific to OpenMP kernels.

OmpSyncThreads applies the OpenMP #pragma omp barrier directive.

Statement types that launch CUDA or HIP GPU kernels are listed next. They work similarly for each back-end and their names are distinguished by the prefix Cuda or Hip. For example, CudaKernel or HipKernel.

Cuda/HipKernel< EnclosedStatements> launches EnclosedStatements as a GPU kernel; e.g., a loop nest where the iteration spaces of each loop level are associated with threads and/or thread blocks as described by the execution policies applied to them. This kernel launch is synchronous.
Cuda/HipKernelAsync< EnclosedStatements> asynchronous version of Cuda/HipKernel.
Cuda/HipKernelFixed<num_threads, EnclosedStatements> similar to Cuda/HipKernel but enables a fixed number of threads (specified by num_threads). This kernel launch is synchronous.
Cuda/HipKernelFixedAsync<num_threads, EnclosedStatements> asynchronous version of Cuda/HipKernelFixed.
CudaKernelFixedSM<num_threads, min_blocks_per_sm, EnclosedStatements> similar to CudaKernelFixed but enables a minimum number of blocks per sm (specified by min_blocks_per_sm), this can help increase occupancy. This kernel launch is synchronous. Note: there is no HIP variant of this statement.
CudaKernelFixedSMAsync<num_threads, min_blocks_per_sm, EnclosedStatements> asynchronous version of CudaKernelFixedSM. Note: there is no HIP variant of this statement.
Cuda/HipKernelOcc<EnclosedStatements> similar to CudaKernel but uses the CUDA occupancy calculator to determine the optimal number of threads/blocks. Statement is intended for use with RAJA::cuda/hip_block_{xyz}_loop policies. This kernel launch is synchronous.
Cuda/HipKernelOccAsync<EnclosedStatements> asynchronous version of Cuda/HipKernelOcc.
Cuda/HipKernelExp<num_blocks, num_threads, EnclosedStatements> similar to CudaKernelOcc but with the flexibility to fix the number of threads and/or blocks and let the CUDA occupancy calculator determine the unspecified values. This kernel launch is synchronous.
Cuda/HipKernelExpAsync<num_blocks, num_threads, EnclosedStatements> asynchronous version of Cuda/HipKernelExp.
Cuda/HipSyncThreads invokes CUDA or HIP __syncthreads() barrier.
Cuda/HipSyncWarp invokes CUDA __syncwarp() barrier. Warp sync is not supported in HIP, so the HIP variant is a no-op.

Statement types that launch SYCL kernels are listed next.

SyclKernel<EnclosedStatements> launches EnclosedStatements as a SYCL kernel. This kernel launch is synchronous.
SyclKernelAsync<EnclosedStatements> asynchronous version of SyclKernel.

RAJA provides statements to define loop tiling which can improve performance; e.g., by allowing CPU cache blocking or use of GPU shared memory.

Tile< ArgId, TilePolicy, ExecPolicy, EnclosedStatements > abstracts an outer tiling loop containing an inner for-loop over each tile. The ArgId indicates which entry in the iteration space tuple to which the tiling loop applies and the TilePolicy specifies the tiling pattern to use, including its dimension. The ExecPolicy and EnclosedStatements are similar to what they represent in a statement::For type.
TileTCount< ArgId, ParamId, TilePolicy, ExecPolicy, EnclosedStatements > abstracts an outer tiling loop containing an inner for-loop over each tile, where it is necessary to obtain the tile number in each tile. The ArgId indicates which entry in the iteration space tuple to which the loop applies and the ParamId indicates the position of the tile number in the parameter tuple. The TilePolicy specifies the tiling pattern to use, including its dimension. The ExecPolicy and EnclosedStatements are similar to what they represent in a statement::For type.
ForICount< ArgId, ParamId, ExecPolicy, EnclosedStatements > abstracts an inner for-loop within an outer tiling loop where it is necessary to obtain the local iteration index in each tile. The ArgId indicates which entry in the iteration space tuple to which the loop applies and the ParamId indicates the position of the tile index parameter in the parameter tuple. The ExecPolicy and EnclosedStatements are similar to what they represent in a statement::For type.

It is often advantageous to use local arrays for data accessed in tiled loops. RAJA provides a statement for allocating data in a Local Array object according to a memory policy. See Local Array Memory Policies for more information about such policies.

InitLocalMem< MemPolicy, ParamList<...>, EnclosedStatements > allocates memory for a RAJA::LocalArray object used in kernel. The ParamList entries indicate which local array objects in a tuple will be initialized. The EnclosedStatements contain the code in which the local array will be accessed; e.g., initialization operations.

RAJA provides some statement types that apply in specific kernel scenarios.

Reduce< ReducePolicy, Operator, ParamId, EnclosedStatements > reduces a value across threads in a multithreaded code region to a single thread. The ReducePolicy is similar to what it represents for RAJA reduction types. ParamId specifies the position of the reduction value in the parameter tuple passed to the RAJA::kernel_param method. Operator is the binary operator used in the reduction; typically, this will be one of the operators that can be used with RAJA scans (see RAJA Scan Operators). After the reduction is complete, the EnclosedStatements execute on the thread that received the final reduced value.
If< Conditional > chooses which portions of a policy to run based on run-time evaluation of conditional statement; e.g., true or false, equal to some value, etc.
Hyperplane< ArgId, HpExecPolicy, ArgList<...>, ExecPolicy, EnclosedStatements > provides a hyperplane (or wavefront) iteration pattern over multiple indices. A hyperplane is a set of multi-dimensional index values: i0, i1, … such that h = i0 + i1 + … for a given h. Here, ArgId is the position of the loop argument we will iterate on (defines the order of hyperplanes), HpExecPolicy is the execution policy used to iterate over the iteration space specified by ArgId (often sequential), ArgList is a list of other indices that along with ArgId define a hyperplane, and ExecPolicy is the execution policy that applies to the loops in ArgList. Then, for each iteration, everything in the EnclosedStatements is executed.

Auxilliary Types¶

The following list summarizes auxiliary types used in the above statements. These types live in the RAJA namespace.

tile_fixed<TileSize> tile policy argument to a Tile or TileTCount statement; partitions loop iterations into tiles of a fixed size specified by TileSize. This statement type can be used as the TilePolicy template parameter in the Tile statements above.

tile_dynamic<ParamIdx> TilePolicy argument to a Tile or TileTCount statement; partitions loop iterations into tiles of a size specified by a TileSize{} positional parameter argument. This statement type can be used as the TilePolicy template parameter in the Tile statements above.

Segs<...> argument to a Lambda statement; used to specify which segments in a tuple will be used as lambda arguments.

Offsets<...> argument to a Lambda statement; used to specify which segment offsets in a tuple will be used as lambda arguments.

Params<...> argument to a Lambda statement; used to specify which params in a tuple will be used as lambda arguments.

ValuesT<T, ...> argument to a Lambda statement; used to specify compile time constants, of type T, that will be used as lambda arguments.

Examples that show how to use a variety of these statement types can be found in Complex Loops (RAJA::kernel).