# RAJA::Launch Basics¶

There are no exercise files to work through for this section. Instead, there is an example source file RAJA/examples/tut_launch_basic.cpp which contains complete code examples of the concepts described here.

Key RAJA features shown in the following examples are:

• RAJA::launch method to create a run-time selectable host/device execution space.
• RAJA::loop methods to express algorithms in terms of nested for loops.

In this example, we introduce the RAJA Launch framework and discuss hierarchical loop-based parallelism. Kernel execution details with RAJA Launch occur inside the lambda expression passed to the RAJA::launch method, which defines an execution space:

RAJA::launch<launch_policy>(RAJA::ExecPlace ,
RAJA::LaunchParams(RAJA::Teams(Nteams,Nteams),
[=] RAJA_HOST_DEVICE (RAJA::LaunchContext ctx) {

/* Kernel code goes here */

});


The RAJA::launch method accepts a RAJA::LaunchPolicy template parameter that can be defined using up to two policies (host and device). For example, the following creates an execution space for a sequential and CUDA kernel dispatch:

using launch_policy = RAJA::LaunchPolicy
<RAJA::seq_launch_t, RAJA::cuda_launch_t<false>>;


Whether a kernel executes on the host or device is determined by the first argument passed to the RAJA::launch method, which is a RAJA::ExecPlace enum value, either HOST or DEVICE. Similar to GPU thread and block programming models, RAJA Launch carries out computation in a predefined compute grid made up of threads which are then grouped into teams when executing on the device. The execution space is then enclosed by a host/device lambda which takes a RAJA::LaunchContext object, which may be used to control the flow within the kernel, for example by creating thread-team synchronization points.

Inside the execution space, developers write a kernel using nested RAJA::loop methods. The manner in which each loop is executed is determined by a template parameter type, which indicates how the corresponding iterates are mapped to the Teams/Threads configuration defined by the RAJA::LaunchParams type passed as the second argument to the RAJA::launch method. Following the CUDA and HIP programming models, this defines an hierarchical structure in which outer loops are executed by thread-teams and inner loops are executed by threads in a team.

     RAJA::loop<teams_y>(ctx, RAJA::TypedRangeSegment<int>(0, Nteams), [&] (int by) {
RAJA::loop<teams_x>(ctx, RAJA::TypedRangeSegment<int>(0, Nteams), [&] (int bx) {

tx, ty, bx, by);

});
});

});
});


The mapping between teams and threads to the underlying programming model depends on how the RAJA::loop template parameter types are defined. For example, we may define host and device mapping strategies as:

using teams_x = RAJA::LoopPolicy<RAJA::loop_exec,
RAJA::cuda_block_x_direct>;
RAJA::cuda_block_x_direct>;


Here, the RAJA::LoopPolicy type holds both the host (CPU) and device (CUDA GPU) loop mapping strategies. On the host, both the team/thread strategies expand out to standard C-style loops for execution:

    for (int by=0; by<Nteams; ++by) {
for (int bx=0; bx<Nteams; ++bx) {

for (int ty=0; ty<Nthreads; ++ty) {
for (int tx=0; tx<Nthreads; ++tx) {

printf("c-iter: iter_tx %d iter_ty %d iter_bx %d iter_by %d \n",
tx, ty, bx, by);
}
}

}
}


On the device the teams_x/y policies will map loop iterations directly to CUDA (or HIP) thread blocks, while the thread_x/y policies will map loop iterations directly to threads in a CUDA (or HIP) thread block. The direct CUDA equivalent of the kernel body using the policy shown above is:

  {int by = blockIdx.y;
{int bx = blockIdx.x;