Directive-Based GPU Programming
Jumanazarov Mardonbek
Plan
There was a real competition to define directive-oriented language standards for GPUs. OpenMP, a prominent directive-oriented language released in 1997, was a natural candidate to be seen as a simpler way to program GPUs. At the time, OpenMP was playing catch-up and was primarily focused on new CPU capabilities. To address the GPU affordability challenge, in 2011 a small group of compiler developers (Cray, PGI, and CAPS), along with NVIDIA as a GPU vendor, joined in releasing the OpenACC standard, providing a simpler path to GPU programming. Similar to what you saw in Chapter 7 on OpenMP, OpenACC also uses pragmas. In this case, OpenACC pragmas direct the compiler to generate GPU code. A couple of years later, the OpenMP Architecture Review Board (ARB) added its own support for GPU pragmas to the OpenMP standard. We'll walk through some basic examples in OpenACC and OpenMP to give you an idea of how they work. We encourage you to try the examples on your target system to see what compilers are available and their current state.
NOTE: As always, we recommend consulting the examples in this chapter at https://github.com/EssentialsofParallelComputing/Chapter11.
Many programmers are torn between using OpenACC or OpenMP, a directive-based language. Often, the choice becomes clear only after you understand what's available in your underlying framework. Keep in mind that the biggest barrier to overcome is simply getting started. If you later decide to switch to GPU languages, the up-front work will still be useful, as the core concepts transcend the language. We hope that after seeing how little effort it takes to generate GPU code using pragmas and directives, you'll try it on a few of your own code snippets. And with just a little effort, you might even achieve a slight speedup.
The History of OpenMP and OpenACC
The development of the OpenMP and OpenACC standards is largely a friendly rivalry; some members of the OpenACC committee also serve on the OpenMP committee. Implementations of these standards continue to emerge, led by efforts at Lawrence Livermore National Laboratory, IBM, and GCC. Attempts have been made to merge these two approaches, but they continue to coexist and will likely continue to do so for the foreseeable future.
OpenMP is rapidly gaining traction and is considered a more reliable long-term approach, but OpenACC currently has more mature implementations and broader developer support. The figure below shows the full release history of these standards. Note that version 4.0 of the OpenMP standard is the first to support GPUs and accelerators.
Release dates for pragma-oriented GPU languages
1. The process of applying directives and pragmas for GPU-based implementation
Directive-based annotations, or pragmas, for C, C++, or Fortran applications provide one of the most attractive ways to access the computing power of GPU processors. Similar to the OpenMP threading model described in Chapter 7, you can add just a few lines of code to your application, and the compiler will generate source code that can run on either the GPU or CPU. As first described in Chapters 6 and 7, pragmas are preprocessor instructions in C and C++ that provide special commands to the compiler. They take the form:
#pragma acc <directive> [expression]
#pragma omp <directive> [expression]
Corresponding capabilities for Fortran code are provided through directives in the form of special comments. Directives begin with a comment symbol followed by the acc or omp keyword to identify them as OpenACC and OpenMP directives, respectively. !$acc <directive> [expression]
!$omp <directive> [expression]
The same general steps are used to implement OpenACC and OpenMP in applications. These steps are shown in Figure 11.1 and will be discussed in detail in the following sections.
1. The process of applying directives and pragmas for GPU-based implementation
Рис. 11.1 Шаги по имплементированию переноса работы на GPU с помощью прагма-ориентированных языков.
Выгрузка работы на GPU приводит к передаче данных, которая замедляет работу приложения, пока не будет редуцировано перемещение данных
1. The process of applying directives and pragmas for GPU-based implementation
The three steps we'll use to convert source code for GPU execution using OpenACC or OpenMP are summarized as follows:
1. Offload computationally intensive work to the GPU. This results in data transfers between the CPU and GPU, which slows down the code, but the work must first be moved;
2. Reduce data transfers between the CPU and GPU. Move memory allocations to the GPU if the data is only used there;
3. Adjust the workgroup size, number of workgroups, and other kernel parameters to improve kernel performance.
At this point, you'll have an application that will run significantly faster on the GPU. Further optimizations are possible to improve performance, although they are typically more application-specific.
2. OpenACC: The Easiest Way to Run on Your GPU
We'll start by creating a simple application running with OpenACC. This is to demonstrate the basic details of how it works. Then, once we have the application up and running, we'll work on optimizing it. As you might expect, with a pragma-oriented approach, little effort yields big returns. But first, you'll need to work through the initial slowdown in your code. Don't despair! It's perfectly normal to encounter an initial slowdown on the way to faster GPU computing.
Often, the hardest step is getting the OpenACC compiler toolchain working. There are several reliable OpenACC compilers. The most well-known compilers available are listed below[1]:
[1] One of the original OpenACC compilers, CAPS, was decommissioned in 2016 and is no longer available.
2. OpenACC: The Easiest Way to Run on Your GPU
In the examples below, we will use the PGI compiler (version 19.7) and CUDA (version 10.1). The PGI compiler is the most mature of the available compilers. The GCC compiler is another option, but be sure to use the latest version. The Cray compiler is a great option if you have access to their system.
2. OpenACC: The Easiest Way to Run on Your GPU
NOTE: What if you don't have a suitable GPU? You can still try the examples by running the source code on your CPU using OpenACC-generated kernels. Performance will vary, but the underlying code should be the same.
Using the PGI compiler, you can first get information about your system using the pgaccelinfo command. This also lets you know whether your system and environment are operational. After running this command, the output should look similar to Figure 11.2.
2. OpenACC: The Easiest Way to Run on Your GPU
Fig. 11.2 The output from the pgaccelinfo command shows the GPU type and its characteristics.
2. OpenACC: The Easiest Way to Run on Your GPU
Compiling the OpenACC Source Code
Listing 11.1 shows several excerpts from OpenACC makefiles. CMake provides the FindOpenACC.cmake module, which is called on line 18 of the listing below. The complete CMakeLists.txt file is included in the source code accompanying this chapter in the OpenACC/StreamTriad directory at https://github.com/EssentialsofParallelComputing/Chapter11. We set several flags to elicit feedback from the compiler and to make the compiler less conservative about potential aliasing. The subdirectory contains the CMake file and a simple makefile.
2. OpenACC: The Easiest Way to Run on Your GPU
2. OpenACC: The Easiest Way to Run on Your GPU
Simple makefiles can also be used to build source code samples by copying or linking them to a Makefile using any of the commands shown below:
ln -s Makefile.simple.pgi Makefile
cp Makefile.simple.pgi Makefile
From the makefiles for the PGI and GCC compilers, we show the suggested flags for OpenACC:
Makefile.simple.pgi
6 CFLAGS:= -g -O3 -c99 -alias=ansi -Mpreprocess -acc -Mcuda -Minfo=accel
7
8 %.o: %.c
9 ${CC} ${CFLAGS} -c $^ 10
11 StreamTriad: StreamTriad.o timer.o
12 ${CC} ${CFLAGS} $^ -o StreamTriad
2. OpenACC: The Easiest Way to Run on Your GPU
Makefile.simple.gcc
6 CFLAGS:= -g -O3 -std=gnu99 -fstrict-aliasing -fopenacc \
-fopt-info-optimized-omp
7
8 %.o: %.c
9 ${CC} ${CFLAGS} -c $^ 10
11 StreamTriad: StreamTriad.o timer.o
12 ${CC} ${CFLAGS} $^ -o StreamTriad
2. OpenACC: The Easiest Way to Run on Your GPU
In the case of PGI, the flags for enabling OpenACC compilation for GCC are -acc -Mcuda. The Minfo=accel flag tells the compiler to provide feedback on acceleration directives. We also include the -alias=ansi flag, telling the compiler to be less concerned about pointer aliasing. This helps it generate parallel computing kernels more freely. It's still a good idea to include the restrict attribute on source code arguments, telling the compiler that variables don't point to overlapping memory locations. Both makefiles also include a flag to specify the C 1999 standard, allowing for defining index variables in a loop for clearer scoping. The -fopenacc flag enables OpenACC directive parsing for GCC. The -fopt-info-optimized-omp flag tells the compiler to provide feedback on accelerator code generation.
For the Cray compiler, OpenACC is enabled by default. You can use the -hnoacc compiler option if you need to disable it. OpenACC compilers must also define the _OPENACC macro. This macro is especially important because OpenACC is still being implemented in many compilers. You can use it to determine the OpenACC version your compiler supports and implement conditional compilations for newer features by comparing it with the _OPENACC == yyyymm compiler macro, where the version dates are as follows:
2. OpenACC: The Easiest Way to Run on Your GPU
Parallel Computation Sections in OpenACC for Computation Acceleration
There are two different options for declaring an accelerated block of code for computation. The first is the kernels pragma, which gives the compiler the freedom to automatically parallelize a block of code. This block of code can include larger sections of code with multiple loops. The second is the parallel loop pragma, which tells the compiler to generate code for the GPU or other accelerator. We will walk through examples of each approach.
2. OpenACC: The Easiest Way to Run on Your GPU
Using the kernels pragma to get automatic parallelization from the compiler
The kernels pragma allows the compiler to automatically parallelize a block of code. It is often used primarily to get feedback from the compiler about a section of code. We'll cover the formal syntax for the kernels pragma, including its optional expressions. Then, we'll look at the thread triad example we've used throughout the programming chapters and apply the kernels pragma. First, we'll list the kernels pragma specification from the OpenACC 2.6 standard:
#pragma acc kernels [ выражение данных | оптимизация ядра | асинхронное выражение | условный блок ]
Где
выражения данных - [ copy | copyin | copyout | create | no_create |present | deviceptr | attach | default(none|present) ]
оптимизация ядра - [ num_gangs | num_workers | vector_length |device_type | self ]
асинхронные выражения - [ async | wait ] условный блок - [ if ]
2. OpenACC: The Easiest Way to Run on Your GPU
We'll discuss data expressions in more detail in Section 11.2.3, although you can also use data expressions in the kernels pragma if they apply only to a single loop. We'll cover kernel optimizations in Section 11.2.4 and briefly mention asynchronous and conditional expressions in Section 11.2.5.
First, we start by specifying the location where we want to parallelize the work by adding #pragma acc kernels around the target code blocks. The kernels pragma applies to the block of code following the specified directive, or to the code in the for loop in the following listing.
2. OpenACC: The Easiest Way to Run on Your GPU
2. OpenACC: The Easiest Way to Run on Your GPU
The following output shows the output from the PGI compiler:
15, Generating implicit copyout(b[:20000000],a[:20000000])
[if not already present]
16, Loop is parallelizable
Generating Tesla code
16, #pragma acc loop gang, vector(128)
/* blockIdx.x threadIdx.x */
16, Complex loop carried dependence of a-> prevents parallelization
Loop carried dependence of b-> prevents parallelization
24, Generating implicit copyout(c[:20000000]) [if not already present]
Generating implicit copyin(b[:20000000],a[:20000000])
[if not already present]
2. OpenACC: The Easiest Way to Run on Your GPU
25, Complex loop carried dependence of a->,b-> prevents parallelization
Loop carried dependence of c-> prevents parallelization
Loop carried backward dependence of c-> prevents vectorization
Accelerator serial kernel generated
Generating Tesla code
25, #pragma acc loop seq
25, Complex loop carried dependence of b-> prevents parallelization
Loop carried backward dependence of c-> prevents vectorization
2. OpenACC: The Easiest Way to Run on Your GPU
What's unclear about this listing is that OpenACC treats each for loop as if it were preceded by an automatic #pragma acc loop. We've left it up to the compiler to decide whether the loop can be parallelized. The bolded output shows that the compiler doesn't think this is possible. The compiler is telling us it needs help. The simplest fix is to add the restrict attribute to lines 8–10 of Listing 11.2.
8. double* restrict a = malloc(nsize * sizeof(double));
9. double* restrict b = malloc(nsize * sizeof(double));
10. double* restrict c = malloc(nsize * sizeof(double));
Our second fix, to help the compiler, is to change the directive to tell the compiler that everything is OK and it's okay to generate parallel GPU code. The problem lies in the default loop (loop auto) directive we mentioned earlier. Here's its specification from the OpenACC 2.6 standard:
#pragma acc loop [ auto | independent | seq | collapse | gang | worker | vector | tile | device_type | private | reduction ]
2. OpenACC: The Easiest Way to Run on Your GPU
We'll touch on many of these expressions in subsequent sections. For now, let's focus on the first three: auto, independent, and seq.
Replacing the expression from auto to independent tells the compiler to parallelize the loop:
15 #pragma acc kernels loop independent
<skip unchanged code>
24 #pragma acc kernels loop independent
Note that in these directives, we've combined two constructs. You can optionally combine valid individual expressions into a single directive. Now the output shows that the loop is parallelized:
2. OpenACC: The Easiest Way to Run on Your GPU
main:
15, Generating implicit copyout(a[:20000000],b[:20000000]) [if not already present]
16, Loop is parallelizable Generating Tesla code
16, #pragma acc loop gang, vector(128)
/* blockIdx.x threadIdx.x */
24, Generating implicit copyout(c[:20000000]) [if not already present] Generating implicit copyin(b[:20000000],a[:20000000])
[if not already present]
25, Loop is parallelizable Generating Tesla code
25, #pragma acc loop gang, vector(128)
/* blockIdx.x threadIdx.x */
In this printout, it's important to note the data transfer response (highlighted in bold). We'll discuss how to respond to this response in Section 11.2.3.
2. OpenACC: The Easiest Way to Run on Your GPU
Trying the parallel loop pragma for greater control over parallelization
Next, we'll cover techniques for using the parallel loop pragma. We recommend using this technique in your application. It's more consistent with the form used in other parallel languages, such as OpenMP. It also provides more consistent and portable performance across compilers. Not all compilers can be relied upon to adequately perform the analysis required by the kernels directive.
The parallel loop pragma is actually two separate directives. The first is the parallel directive, which opens a parallel section. The second is the loop pragma, which distributes work among parallel work units. We'll look at the parallel pragma first. The parallel pragma accepts the same expressions as the kernels directive. In the example below, we've highlighted the additional expressions of the kernels directive in bold:
2. OpenACC: The Easiest Way to Run on Your GPU
#pragma acc parallel [ clause ]
выражения данных - [ reduction | private | firstprivate | copy |
copyin | copyout | create | no_create | present | deviceptr | attach | default(none|present) ]
�
оптимизация ядра - [ num_gangs | num_workers |
vector_length | device_type | self ] асинхронные выражения - [ async | wait ]
условный блок - [ if ]
Expressions for the loop construct were mentioned earlier in the kernels section. It's important to note that by default, the loop construct in a parallel section is independent, not automatic. Again, as with the kernels directive, the combined parallel loop construct can accept any expression that can be used by individual directives. With this explanation of the parallel loop construct, we'll move on to adding it to the stream triad example, as shown in the following listing.
2. OpenACC: The Easiest Way to Run on Your GPU
2. OpenACC: The Easiest Way to Run on Your GPU
The PGI compiler output looks like this:
main:
15, Generating Tesla code
16, #pragma acc loop gang, vector(128)
/* blockIdx.x threadIdx.x */
15, Generating implicit copyout(a[:20000000],b[:20000000]) [if not already present]
24, Generating Tesla code
25, #pragma acc loop gang, vector(128)
/* blockIdx.x threadIdx.x */
24, Generating implicit copyout(c[:20000000]) [if not already present] Generating implicit copyin(b[:20000000],a[:20000000])
[if not already present]
2. OpenACC: The Easiest Way to Run on Your GPU
The loop is parallelizable even without the restrict attribute, since the default for the loop directive is the independent clause. This differs from the default for kernels, which we saw earlier. However, we recommend using the restrict attribute in your code to help the compiler generate better code.
The result is similar to the output from the previous kernels directive. At this point, the code's performance will likely slow down due to the data movement, which we highlighted in bold in the compiler output above. Don't worry, we'll speed it up again in the next step.
Before we look at data movement, let's briefly review reductions and the serieal construct. Listing 11.4 shows an example of the mass sum first introduced in Section 6.3.3. The mass sum is a simple reduction operation. Instead of OpenMP's SIMD vectorization pragma, we placed OpenACC’s parallel loop pragma with a reduction clause before the loop. The syntax of the reduction is already familiar because it is the same as in the OpenMP streaming standard.
2. OpenACC: The Easiest Way to Run on Your GPU
2. OpenACC: The Easiest Way to Run on Your GPU
Other operators can also be used in the reduction clause. These include *, max, min, &, |, &&, and ||. For OpenACC versions prior to 2.6, the variable or comma-separated list of variables is limited to scalars, not arrays. However, OpenACC 2.7 allows arrays and composite variables in the reduction clause.
The last construct we'll cover in this section is for sequential execution. Some loops cannot be executed in parallel. Instead of breaking out of the parallel section, we stay within it and tell the compiler to simply execute that part sequentially. This is done with the serial directive:
#pragma acc serial
Blocks of this code with the serial directive are executed by a single gang, consisting of one worker with a vector length of one. Now let's turn our attention to the response to data movement.
2. OpenACC: The Easiest Way to Run on Your GPU
Using Directives to Reduce Data Movement Between CPU and GPU
In this section, we return to a theme we've explored throughout this book: data movement is more important than flops. While we sped up the computation by moving it to the GPU, the overall execution time slowed down due to the cost of data movement. Addressing the problem of excessive data movement will begin to yield an overall speedup. To achieve this, we add the data construct to our code. In the OpenACC standard, version 2.0, the specification of the data construct is as follows:
#pragma acc data [ copy | copyin | copyout | create | no_create | present | deviceptr | attach | default(none|present) ]
You'll also see references to expressions such as present_or_copy or the shortened form pcopy, which check for the presence of data before copying. These are no longer necessary, although they are retained for backward compatibility. This behavior has been included in the standard expressions since version 2.5 of the OpenACC standard.
2. OpenACC: The Easiest Way to Run on Your GPU
Many data expressions use an argument listing the data to be copied or otherwise manipulated. The compiler must obtain a range specification for the array. An example of this is the following fragment:
#pragma acc data copy(x[0:nsize])
The range specification differs slightly between C/C++ and Fortran. In C/C++, the first argument in the specification is the starting index, and the second is the length. In Fortran, the first argument is the starting index, and the second is the ending index.
There are two types of data ranges. The first is the structured data range from the original OpenACC standard, version 1.0. The second, the dynamic data range, was introduced in OpenACC version 2.0. We'll look at the structured data range first.
2. OpenACC: The Easiest Way to Run on Your GPU
Structured Data Region for Simple Code Blocks
A structured data region is delimited by a code block. This can be a natural code block formed by a loop or a code section contained within a set of curly braces. In Fortran, this region is marked with a start directive and ends with an end directive. Listing 11.5 shows an example of a structured data region that begins with a directive on line 16 and is delimited by an opening brace on line 17 and an ending brace on line 37. We've included a comment within the ending brace to help identify the code block that the brace ends.
2. OpenACC: The Easiest Way to Run on Your GPU
2. OpenACC: The Easiest Way to Run on Your GPU
A structured data region specifies that three arrays should be created at the beginning of the data region. These arrays will be destroyed at the end of the data region. The present expression is used in two parallel loops to avoid copying data for computational regions.
Dynamic data region for more flexible data scoping
The structured data region originally used by the OpenACC language, which allocates memory and then places multiple loops, does not work with more complex programs. In particular, memory allocation in object-oriented code occurs when an object is created. How can a data region be placed around something with such a program structure?
To address this issue, dynamic (also known as unstructured) data regions were added in OpenACC version 2.0. The dynamic data region construct was specifically designed for more complex data management scenarios, such as constructors and destructors in C++. Instead of using curly braces to define a data region, the pragma uses an enter and exit expression:
#pragma acc enter data
#pragma acc exit data
2. OpenACC: The Easiest Way to Run on Your GPU
The exit data directive has an additional delete statement that can be used. This use of the enter/exit data directive is best suited where memory allocations and deallocations occur. The enter data directive should be placed immediately after the allocation, and the exit data directive should be inserted immediately before the deallocation. This approach more naturally aligns with the existing variable scope in the application. As soon as you want to improve performance beyond what can be achieved with a loop-level strategy, the importance of these dynamic data regions increases. With larger dynamic data region scopes, an additional data update directive becomes necessary:
#pragma acc update [self(x) | device(x)]
The device argument specifies that the data on the device should be updated. The self argument specifies updating local data, which is typically the host version of the data.
Let's look at an example of using the data dynamic data pragma in Listing 11.6. The enter data directive is placed after the allocation on line 12. The exit data directive on line 35 is inserted before the deallocations. We suggest using dynamic data regions, as opposed to structured data regions, in all but the simplest code.
2. OpenACC: The Easiest Way to Run on Your GPU
2. OpenACC: The Easiest Way to Run on Your GPU
2. OpenACC: The Easiest Way to Run on Your GPU
If you examine the listing above carefully, you'll notice that the arrays a, b, and c are allocated on both the host and the device, but are only used on the device. In Listing 11.7, we demonstrate one approach to correcting this situation, using the acc_malloc procedure and then placing the deviceptr expression in the compute regions.
2. OpenACC: The Easiest Way to Run on Your GPU
2. OpenACC: The Easiest Way to Run on Your GPU
2. OpenACC: The Easiest Way to Run on Your GPU
As shown below, the output from the PGI compiler is now much shorter:
16 Generating Tesla code
17 #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
25 Generating Tesla code
26 #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
Data movement is eliminated, and host memory requirements are reduced. We still have some output data with feedback about the generated compute kernel, which we will discuss in Section 11.2.4. This example (Listing 11.7) works for one-dimensional arrays. For two-dimensional arrays, the deviceptr expression does not accept a descriptor argument, so the kernel must be modified to perform native two-dimensional indexing in a one-dimensional array. When accessing data regions, you have a rich set of data directives and data movement expressions that can be used to reduce unnecessary data movement. However, we haven't covered many more OpenACC expressions and functions that may be useful in specialized situations.
2. OpenACC: The Easiest Way to Run on Your GPU
Optimizing GPU Compute Cores
As a general rule, you'll have a greater impact by increasing the number of compute cores running on a GPU and reducing data movement than by optimizing the GPU compute cores themselves. The OpenACC compiler does a good job of generating cores, and the potential gains from further optimizations are small. Occasionally, you can help the compiler improve the performance of key cores enough to be worth the effort.
In this section, we'll cover the general strategies for these optimizations. First, let's discuss the terminology used in the OpenACC standard. As shown in Figure 11.3, OpenACC defines abstract levels of parallelism that apply to multiple hardware devices.
2. OpenACC: The Easiest Way to Run on Your GPU
Рис. 11.3 Иерархия уровней
в OpenACC: бригады, работники и векторы
2. OpenACC: The Easiest Way to Run on Your GPU
OpenACC defines the following levels of parallelism:
Below are some examples of specifying the level of a specific loop directive:
#pragma acc parallel loop vector
#pragma acc parallel loop gang
#pragma acc parallel loop gang vector
2. OpenACC: The Easiest Way to Run on Your GPU
The outer loop must be a gang loop, and the inner loop must be a vector loop. A worker loop can appear between them. A sequential loop (seq) can appear at any level.
For most modern GPUs, the vector length must be a multiple of 32, so it is an integer multiple of the warp size. It should be no greater than the maximum number of threads per block, which is typically around 1024 on current GPUs (see the output of the pgaccelinfo command in Figure 11.2). In the examples shown here, the PGI compiler sets the vector length to a reasonable value of 128. This value can be changed for the loop using the vector_length(x) directive.
In what scenario should the vector_length parameter be changed? If the inner loop of continuous data is less than 128, part of the vector will remain unused. In this case, it can be useful to reduce this value. Another option, as we'll discuss shortly, is to collapse a couple of the inner loops to produce a longer vector.
2. OpenACC: The Easiest Way to Run on Your GPU
The worker setting can be modified using the num_workers expression. However, it is not used in the examples in this chapter. However, it can be useful to increase it when the vector length is shortened or for an additional level of parallelization. If your code requires synchronization within a parallel workgroup, you should use a worker, but OpenACC does not provide a synchronization directive for the user. The worker layer also shares resources such as cache and local memory.
The remaining parallelization is performed using gangs, which are an asynchronous parallel layer. A large number of gangs on GPU processors is important because they hide latency and ensure high frequency. Typically, the compiler sets their number to a large number, so the user does not need to override it. If you might need to do this, the num_gangs expression is provided for this case.
2. OpenACC: The Easiest Way to Run on Your GPU
Many settings will only apply to a specific piece of hardware. The device_type(type) parameter before the expression limits the device to the specified type. This parameter remains active until the next device_type expression is found. For example:
}
2. OpenACC: The Easiest Way to Run on Your GPU
For a list of supported device types, see the openacc.h header file for PGI v19.7. Note that the lines in the openacc.h header shown earlier do not include the acc_device_radeon device, so the PGI compiler does not support the AMD Radeon™ device. This means we need a C preprocessor ifdef near line 3 in the code example above to prevent the PGI compiler from complaining.
Excerpt from openacc.h file for PGI 27 typedef enum{
2. OpenACC: The Easiest Way to Run on Your GPU
The kernels directive syntax is slightly different, with the parallel type applied to each loop directive separately and accepting an int argument directly:
#pragma acc kernels loop gang for (int j=0; j<jmax; j++){
#pragma acc loop vector(64) for (int i=0; i<imax; i++){
<work>
}
}
Loops can be combined with the collapse(n) expression. This is especially useful if there are two small inner loops that step through the data in a contiguous order. Combining these expressions allows for a longer vector length. Loops must be tightly nested.
DEFINITION: Two or more loops in which there are no additional statements between the for or do statements or between the ends of the loops are tightly nested.
2. OpenACC: The Easiest Way to Run on Your GPU
An example of combining two loops to use a long vector is the following code fragment:
#pragma acc parallel loop collapse(2) vector(32) for (int j=0; j<8; j++){
for (int i=0; i<4; i++){
<работа>
}
}
OpenACC version 2.0 added the tile expression, which can be used for optimization. You can specify the tile size or use asterisks to let the compiler choose automatically:
#pragma acc parallel loop tile(*,*) for (int j=0; j<jmax; j++){
for (int i=0; i<imax; i++){
<работа>
}
}
2. OpenACC: The Easiest Way to Run on Your GPU
Now is the time to try out various optimizations to the compute kernel. The streaming triad example didn't show any real benefit from our optimization attempts, so we'll work with the stencil example used in several previous chapters.
Related to the stencil example, the source code for this chapter goes through the same first two steps of moving computation cycles to the GPU, then reduces the data movement. The stencil code also requires one additional change. On the CPU, we exchange pointers at the end of the cycle. On the GPU, in lines 45–50, we must copy the new data back to the original array. The following listing shows the stencil code example, with these steps shown in full.
2. OpenACC: The Easiest Way to Run on Your GPU
2. OpenACC: The Easiest Way to Run on Your GPU
2. OpenACC: The Easiest Way to Run on Your GPU
First, note that we're using dynamic data region directives, so the data region isn't wrapped in curly braces like with a structured data region. A dynamic data region begins with the data region when it encounters the enter directive and ends when it reaches the exit directive, regardless of the path between the two directives. In this case, it's a straight line of execution from the enter directive to the exit directive. We'll add a collapse statement to the parallel loop to reduce the overhead of the two loops. This change is shown in the following listing.
2. OpenACC: The Easiest Way to Run on Your GPU
We can also try using the tile expression. We'll start by letting the compiler determine the tile size, as shown in lines 41 and 48 in the listing below.
2. OpenACC: The Easiest Way to Run on Your GPU
2. OpenACC: The Easiest Way to Run on Your GPU
The change in execution times resulting from these optimizations is small compared to the improvement observed with the initial implementation of OpenACC. Table 11.1 shows the results for an NVIDIA V100 GPU with PGI compiler version 19.7.
Table 11.1 Execution times for OpenACC stencil kernel optimizations
2. OpenACC: The Easiest Way to Run on Your GPU
We tried changing the vector length to 64 or 256 and different tile sizes, but saw no improvement in execution time. More complex code may benefit more from kernel optimizations, but note that any specialization of parameters such as vector length impacts compiler portability across different architectures.
Another optimization target is the implementation of a pointer swap at the end of the loop. Pointer swapping is used in the original CPU code as a fast way to return data to the original array. Copying data back to the original array doubles the execution time on the GPU. The difficulty of pragma-oriented languages is that swapping pointers in a parallel section requires simultaneous exchange of host and device pointers.
2. OpenACC: The Easiest Way to Run on Your GPU
Summary of Resulting Performance for the Streaming Triad
The runtime performance when converted to GPUs exhibits a typical pattern. Moving compute cores to the GPU results in a slowdown of approximately 3x, as shown in the implementations of core 2 and parallel case 1 in Table 11.2. In case 1, the computational cycle cannot be parallelized. When executed serially on the GPU, performance was even slower. After reducing data movement in core 3 and parallel cases 2–4, execution times increased by 67x. The specific data chunk type is not as critical for performance, but can be important for enabling additional cycle migration in more complex source codes.
2. OpenACC: The Easiest Way to Run on Your GPU
Table 11.2 Execution times of OpenACC stream triad computing kernel optimizations
2. OpenACC: The Easiest Way to Run on Your GPU
Advanced OpenACC Techniques
OpenACC has many other features for manipulating more complex source code. We'll cover them briefly to give you an idea of the possibilities.
Manipulating Functions with the openacc routine directive
OpenACC v1.0 requires functions used in computation kernels to be inlineable. In version 2.0, two different versions of the routine directive were added to simplify procedure calls. These two versions are:
#pragma acc routine [gang | worker | vector | seq | bind | no_host | device_type]
#pragma acc routine(name) [gang | worker | vector | seq | bind | no_host |
device_type]
In C and C++, the routine directive must appear immediately before the function prototype or definition. The named version can appear anywhere before the function is defined or used. The Fortran version must include the !#acc routine directive in the function body itself or in the interface body.
2. OpenACC: The Easiest Way to Run on Your GPU
Avoiding Race Conditions with OpenACC Atomic Directives
Many threaded routines have a shared variable that must be updated by multiple threads. This programming construct is a common performance bottleneck and potential race condition. To address this, OpenACC v2 provides atomics, which allow a storage location to be accessed by only one thread at a time. The syntax and valid expressions for the atomic directive are as follows:
#pragma acc atomic [read | write | update | capture]
If expression is omitted, update is used by default.
Below is an example of using the atomic expression:
#pragma acc atomic
cnt++;
2. OpenACC: The Easiest Way to Run on Your GPU
Asynchronous Operations in OpenACC
OpenACC's stacking of operations helps improve performance. The proper term for stacking operations is "asynchronous." OpenACC provides these asynchronous operations as expressions and as async and wait directives. The async expression is added to a work directive or data directive with an optional integer argument:
#pragma acc parallel loop async([<integer>])
The wait keyword can be either a directive or an expression added to a work directive or data directive. The following pseudocode in Listing 11.11 shows how it can be used to start computations on the x and y faces of a computational grid and then wait for the results to update cell values for the next iteration.
2. OpenACC: The Easiest Way to Run on Your GPU
Unified Memory Avoiding Data Movement Management
While unified memory is not currently part of the OpenACC standard, there are experimental developments that enable the system to manage memory movement. This experimental unified memory implementation is available in CUDA and the PGI OpenACC compiler. Using the -ta=tesla:managed flag with the PGI compiler and the latest NVIDIA GPUs, you can try out their unified memory implementation. Although the coding is simplified, the performance impact is currently unknown and will change as compilers mature.
2. OpenACC: The Easiest Way to Run on Your GPU
Interoperability with CUDA Libraries and Compute Kernels
OpenACC provides several directives and functions that enable interoperability with CUDA libraries. When calling the libraries, you must inform the compiler to use device pointers instead of host data. The host_data directive can be used for this purpose:
#pragma acc host_data use_device(x, y)
cublasDaxpy(n, 2.0, x, 1, y, 1);
We showed a similar example when allocating memory using acc_malloc in Listing 11.7. When using acc_malloc or cudaMalloc, the returned pointer is already on the device. In this case, we used the deviceptr expression to pass a pointer to the data region.
One of the most common errors in GPU programming in any language is confusing the device pointer with the host pointer. Try searching for "86 Pike Place, San Francisco" when it's actually "86 Pike Place, Seattle." The device pointer points to a different physical memory block on the GPU hardware.
Figure 11.4 shows the three different operations we've discussed to help you understand the differences. In the first case, the malloc procedure returns the host pointer.
2. OpenACC: The Easiest Way to Run on Your GPU
Рис. 11.4 Чем является этот указатель: указателем устройства либо указателем хоста? Один указывает соответственно на память GPU, а другой – на память CPU. OpenACC поддерживает соответствие между массивами в двух адресных пространствах и предоставляет процедуры для извлечения каждого из них
2. OpenACC: The Easiest Way to Run on Your GPU
The present expression converts this into a device pointer for the device's compute kernel. In the second case, when we allocate memory on the device using acc_malloc or cudaMalloc, we are given a device pointer. The deviceptr expression is used to send it to the GPU without any modifications. In the latter case, we don't have a host pointer at all. We must use the host_data directive use_device(var) to extract the device pointer to the host. This is done so that we have a pointer to send back to the device in the argument list of the device function.
It is recommended to append _h or _d to pointers to clarify their valid context. Our examples assume that all pointers and arrays are on the host, except for those ending in _d, which applies to any device pointer.
Managing Multiple Devices in OpenACC
Many modern HPC systems already have multiple GPUs. Furthermore, it is predictable that we will have nodes with different accelerators. The ability to manage the devices you use is becoming increasingly important. OpenACC provides this capability through the following functions:
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
OpenMP's acceleration capabilities are an exciting addition to the traditional threading model. In this section, we'll show you how to get started with its directives. We'll use the same examples as in Section 11.2 of OpenACC. By the end of this section, you should have some idea of how the two similar languages compare and which one might be a more suitable choice for your application.
Where do OpenMP's acceleration directives compare to OpenACC's? Currently, OpenMP implementations are noticeably less mature, although they are rapidly improving. Implementations are currently available for the following GPUs:
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
The two most mature implementations, from Cray and IBM, are available only on their respective systems. Unfortunately, not all developers have access to systems from these vendors, but more widely available compilers exist. Two of these compilers, Clang and GCC, are under development, and marginal versions are currently available. Follow the development of these compilers. The examples in this section use CUDA v10 and the IBM® XL 16 compiler.
Compiling OpenMP Source Code
We'll begin by setting up the build environment and compiling the OpenMP source code. CMake has an OpenMP module, but it doesn't explicitly support OpenMP acceleration directives. We include the OpenMPAccel module, which calls the regular OpenMP module and adds the necessary accelerator flags. It also checks the supported OpenMP version, and if it isn't 4.0 or later, it generates an error. The CMake module is included with the chapter's source code. Listing 11.12 shows excerpts from the main CMakeLists.txt file for this chapter. The output of most OpenMP compilers is currently poorly organized, so setting the
-DCMAKE_OPENMPACCEL flag to CMake will provide only minimal benefit. In these examples, we'll use other tools to fill this gap.
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
A simple makefile can also be used to generate source code examples by copying or linking them to the Makefile module in one of the following ways:
ln -s Makefile.simple.xl Makefile
cp Makefile.simple.xl Makefile
The following code fragment shows the suggested flags for OpenMP acceleration directives in simple makefiles for the IBM XL and GCC compilers:
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
Makefile.simple.xl
6 CFLAGS:=-qthreaded -g -O3 -std=gnu99 -qalias=ansi -qhot -qsmp=omp \
-qoffload -qreport
7
8 %.o: %.c
10
11 StreamTriad: StreamTriad.o timer.o
12 ${CC} ${CFLAGS} $^ -o StreamTriad
Makefile.simple.gcc
6 CFLAGS:= -g -O3 -std=gnu99 -fstrict-aliasing \
7 -fopenmp -foffload=nvptx-none -foffload=-lm -fopt-info-omp
8
9 %.o: %.c
11
12 StreamTriad: StreamTriad.o timer.o
13 ${CC} ${CFLAGS} $^ -o StreamTriad
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
Generating Parallel Jobs on the GPU with OpenMP
Now we need to generate parallel jobs on the GPU. OpenMP's parallel device abstractions are more complex than those we've seen in OpenACC. But this complexity also provides greater flexibility in scheduling future jobs. For now, you must preface each loop with the following directive:
#pragma omp target teams distribute parallel for simd
This is a long and confusing directive. Let's examine each part, as shown in Figure 11.5. The first three statements specify hardware resources:
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
Рис. 11.5 Директивы target, teams и distribute позволяют задействовать больше ресурсов оборудования. Директива parallel for simd распределяет работу внутри каждой рабочей группы
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
The remaining three are parallel execution expressions. All three expressions are required for portability. This is because compiler implementations distribute work differently:
For kernels with three nested loops, one way to distribute the work is as follows:
Loop k: #pragma omp target teams distribute
Loop j: #pragma omp parallel for
Loop i: #pragma omp simd
Each OpenMP compiler may distribute the work differently, requiring several variations of this scheme. The simd loop must be an inner loop in all contiguous memory locations. Some simplification of this complexity is introduced by the loop expression in OpenMP v5.0, as we will present in Section 11.3.5. You can also add expressions to this directive:
private, firstprivate, lastprivate, shared, reduction, collapse, dist_schedule
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
Many of these expressions are familiar to OpenACC and behave identically. One of the major differences from OpenACC is the default way of manipulating data when entering a parallel region. OpenACC compilers typically move all necessary arrays to the device. With OpenMP, there are two options:
Let's look at a simple example of adding a parallel directive in Listing 11.13. We use statically allocated arrays, which behave as if they were allocated on the stack; however, due to their large size, the actual memory could be allocated on the heap by the compiler.
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
The IBM XL compiler output shows that two cores are offloaded to the GPU, but offers no other information. GCC provides no feedback at all. The IBM XL output is:
"" 1586-672 (I) GPU OpenMP Runtime elided for offloaded kernel 'xl_main_l15_OL_1'
"" 1586-672 (I) GPU OpenMP Runtime elided for offloaded kernel 'xl_main_l23_OL_2'
To get some information about what the IBM XL compiler did, we'll use the NVIDIA profiler:
nvprof ./StreamTriad_par1
The first part of the output is:
==141409== Profiling application: ./StreamTriad_par1
==141409== Profiling result:
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
From this output, we now know that a memory copy is made from the host to the device (HtoD in the output), and then back from the device to the host (DtoH in the output). The nvprof output from GCC is similar, but without line numbers. More detailed information about the execution order can be obtained by running the following command:
nvprof --print-gpu-trace ./StreamTriad_par1
Most programs are not written with statically allocated arrays. Let's look at the more common case where arrays are dynamically allocated, as shown in the following listing.
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
Note that the map statement is added to lines 16 and 26. If you try the directive without this statement, although it compiles fine with the IBM XLC compiler, you will receive this message at runtime[1]:
1587-164 Encountered a zero-length array section that points to memory
starting at address 0x200020000010. Because this memory is not currently
mapped on the target device 0, a NULL pointer will be passed to the device.
1587-175 The underlying GPU runtime reported the following error: "an illegal
memory access was encountered."
1587-163 Error encountered while attempting to execute on the target device 0. The program will stop.
However, the GCC compiler compiles and executes the code normally without the map directive. Therefore, the GCC compiler moves heap-allocated memory to the device, while IBM XLC does not. For portability, it is necessary to include a map statement in the application source code.
OpenMP also has a reduction statement for directives designed for parallel processing regions. Its syntax is similar to that of OpenMP and OpenACC directives for threaded operation. An example directive looks like this:
#pragma omp commands distribute parallel for simd reduction(+:sum)
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
Creating Data Chunks to Manage Data Movement to the GPU with OpenMP
Now that we've moved the work to the GPU, we can add data chunks to manage data movement to and from the GPU. Data movement directives in OpenMP are similar to those in OpenACC, both structured and dynamic. The form of the directive is:
#pragma omp target data [ map() | use_device_ptr() ]
The work directives are enclosed in a structured data chunk, as shown in Listing 11.15. The data is copied to the GPU if it isn't already there. This data is then maintained until the block ends (line 35), when it is copied back. This significantly reduces data transfers for each parallel execution cycle and should result in a net speedup in the overall application execution time.
[1] Translation: 1587-164 A section with a zero-length array was encountered that points to memory starting at address 0x200020000010. Since this memory is not currently mapped to target device 0, a NULL pointer will be passed to the device. 1587-175 The GPU Reference Runtime reported the following error: "illegal memory access detected". 1587-163 An error was encountered while attempting to execute on target device 0. The program will stop. – Translator's note:
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
Structured data regions cannot manipulate more general programming patterns. OpenACC and OpenMP (version 4.5) added dynamic data regions, often referred to as unstructured data regions. The directive's form contains enter and exit statements with a map mapping modifier to specify the data transfer operation (such as the default to and from):
#pragma omp target enter data map([alloc | to]:array[[start]:[length]])
#pragma omp target exit data map([from | release | delete]:
array[[start]:[length]])
In Listing 11.16, we convert the omp target data directive to the omp target enter data directive (line 13). The data's scope on the GPU ends when it encounters the omp target exit data directive (line 36). These directives operate similarly to the structured data region in Listing 11.15. But the dynamic data section can be used in more complex data management scenarios, such as constructors and destructors in C++.
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
We can further optimize data transfers by allocating arrays on the device and deleting arrays when leaving a data region, thereby eliminating another data transfer. When a transfer is necessary to move data between the CPU and GPU, the omp target update directive can be used. The syntax for this directive is:
#pragma omp target update [to | from] (array[start:length])
We should also recognize that in this example, the CPU never uses the array memory. For memory that exists only on the GPU, we can allocate it there and then inform the parallel work regions that it already exists. This is done in several ways. One is to use OpenMP function calls to allocate and deallocate memory on the device. These calls look like this and require the inclusion of the OpenMP header file:
#include <omp.h>
double *a = omp_target_alloc(nsize*sizeof(double), omp_get_default_device()); omp_target_free(a, omp_get_default_device());
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
We could also use CUDA memory allocation routines. To use these routines, we need to include the cuda_runtime header file:
#include <cuda_runtime.h>
cudaMalloc((void *)&a,nsize*sizeof(double));
cudaFree(a);
In the parallel execution directives, we then need to add another statement related to passing device pointers to the cores on the device:
#pragma omp target commands distribute parallel for is_device_ptr(a)
Putting it all together, we get the code changes shown in the following listing.
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
OpenMP offers another way to allocate data on a device. This method uses the omp declare target directive, as shown in Listing 11.18. We first declare the array pointers in lines 10–12 and then allocate them on the device in the following block of code (lines 14–19). A similar block is used in lines 42–47 to free the data on the device.
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
As we've already seen, there are a wide range of options for managing GPU data. We've now covered the most common OpenMP directives and data region expressions. Recent additions to the OpenMP standard allow for the manipulation of more complex data structures and data transfers.
Optimizing OpenMP for GPUs
Let's move on to a stencil example of optimizing a compute kernel, as we did for OpenACC. Several things can be tried to speed up individual kernels, but for portability reasons, it's generally better to let the compiler handle the optimizations itself. The core of a stencil kernel with OpenMP data and work regions in the following listing provides a starting point for optimization.
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
Simply adding a single work directive for a two-dimensional loop and data construct is not enough to efficiently generate GPU work for version 16 of the IBM XL compiler. The execution time is almost twice that of the production version (see Table 11.4 at the end of this section). You can use nvprof to see where the time is spent. Here's the result:
==11376== Profiling application: ./Stencil_par2
==11376== Profiling result:
Time(%) Time Calls Avg Min Max Name
51.63% 9.73622s 1000 9.7362ms 9.6602ms 15.378ms xl_main_l42_OL_3
48.26% 9.10010s 1000 9.1001ms 9.0323ms 13.588ms xl_main_l41_OL_2
0.11% 20.439ms 1 20.439ms 20.439ms 20.439ms xl_main_l18_OL_1
0.00% 7.2960us 5 1.4590us 1.2160us 2.1440us [CUDA memcpy DtoH]
0.00% 5.3760us 2 2.6880us 2.5600us 2.8160us [CUDA memcpy HtoD]
<continued printout>
The first line shows that the third computational core takes up more than 50% of the execution time. Copying back to the original array takes an additional 48% of the execution time. The problem lies in the computational core code, not in the data transfer! To fix this, you should first try collapsing the two nested loops into a single parallel construct. The corresponding changes involve adding a collapse statement along with the number of loops to collapse in the execution directives. This is shown on lines 22, 30, 42, and 49 in the listing below.
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
The execution time is now faster than on the CPU (see Table 11.3), although not as fast as the version generated by the PGI OpenACC compiler (Table 11.1). We expect this situation to improve as the IBM XL compiler improves. Let's try another approach by splitting the parallel execution directives into two loops, as shown in the following listing.
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
The IBM XL compiler's output timing for parallel processing directives is similar to the collapse statement. Table 11.3 shows the results of our experiments with kernel optimizations.
Table 11.3. OpenMP Stencil Kernel Optimization Execution Times
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
In Table 11.4, we also look at the execution times of the streaming triad example output from the IBM XL v16 compiler on a Power 9 processor with an NVIDIA V100 GPU. The CPU performance differs because in one case we used an Intel Skylake processor, and in this case, a Power 9 processor. However, it is encouraging to note that the OpenMP streaming core performance on the V100 GPU is essentially the same as that of the PGI OpenACC compiler in Table 11.2.
Table 11.4: OpenMP Streaming Triad Computing Core Optimizations
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
OpenMP performance with the IBM XL compiler is good for a simple one-dimensional benchmark, but could be improved for the two-dimensional stencil case. So far, the focus has been on properly implementing the OpenMP standard for device offloading. We expect performance to improve with each compiler release and as more compiler vendors offer support for device offloading within OpenMP.
Advanced OpenMP for GPUs
OpenMP has a number of additional advanced features. OpenMP is also evolving based on experience with early implementations on GPUs and as hardware evolves. We will cover just a few additional directives and expressions that are important for:
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
Controlling OpenMP Compiler-Implemented GPU Computing Kernel Parameters
We'll begin by examining expressions that can be used to fine-tune the performance of the computing kernel. We can add these expressions to directives to modify the kernels the compiler generates for the GPU:
These expressions can be useful in special situations, but in general, it's best to leave optimization of these parameters to the compiler.
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
OpenMP Device Function Declaration
When we call a function in a parallel region on a device, we need to somehow inform the compiler that it should also be on the device. This is done by adding the declare target directive to the function. Its syntax is similar to that for variable declarations. Here's an example:
#pragma omp declare target int my_compute(<args>){
<job>
}
New scan reduction type
We discussed the importance of the scan algorithm in Section 5.6, where we also saw the difficulty of implementing this algorithm on a GPU. This operation is widely used in parallel computing and is difficult to write, so adding this type would be very welcome. The scan type will be available in OpenMP version 5.0.
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
int run_sum = 0;
#pragma omp parallel for simd reduction(inscan,+: run_sum) for (int i = 0; i < n; ++i) {
run_sum += ncells[i];
#pragma omp scan exclusive(run_sum) cell_start[i] = run_sum;
#pragma omp scan inclusive(run_sum) cell_end[i] = run_sum;
}
Preventing Race Conditions with the OpenMP atomic Statement
It's perfectly normal for multiple threads to access a shared variable in an algorithm. This often becomes a bottleneck in procedure execution. Various compilers and threading implementations provide this functionality through atomics. OpenMP also offers the atomic directive. An example of its use is provided below.
#pragma omp atomic
i++;
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
OpenMP's Version of Asynchronous Operations
In Section 10.5, we discussed the importance of overlapping data transfers and computations using asynchronous operations. OpenMP also offers its own version of these operations.
You create asynchronous device operations using the nowait clause in a data or work directive. You can then use the depend clause to specify that a new operation cannot begin until the previous operation completes. These operations can be chained together to form a sequence of operations. We can use a simple taskwait directive to wait for all tasks to complete:
#pragma omp taskwait
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
Accessing Special Memory Spaces
Memory bandwidth is often one of the most critical performance constraints. Pragma-oriented languages haven't always provided the ability to control memory allocation and its resulting bandwidth. The addition of features giving the programmer greater control over this was one of the most anticipated additions to OpenMP. With OpenMP 5.0, you can target special memory spaces, such as shared memory and high-bandwidth memory. This capability is enabled by the new allocator modifier. The allocate expression takes an optional modifier as follows:
allocate([allocator:] list)
The following pair of functions can be used to directly allocate and free memory:
omp_alloc(size_t size, omp_allocator_t *allocator)
omp_free(void *ptr, const omp_allocator_t *allocator)
The OpenMP 5.0 standard defines several predefined memory spaces for allocators, as shown in the table below.
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
There are a number of functions for defining new memory allocators. The two main procedures are:
omp_init_allocator
omp_destroy_allocator
These allocators accept one of a series of predefined space arguments and allocator characteristics, such as whether it should be pinned, aligned, private, nearby, and many others. Implementations of this capability are still under development. This functionality will be increasingly important in new architectures where specialized memory types exist with different performance, latency, and bandwidth characteristics.
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
Deep Copy Support for Transferring Complex Data Structures
OpenMP 5.0 also adds the declare mapper construct, which can make deep copies. Deep copies duplicate not only the data structure with pointers but also the data referenced by the pointers. Programs with complex data structures and classes have faced difficulties porting to GPU processors. The ability to make deep copies significantly simplifies these implementations.
Simplifying Work Distribution with the New Loop Directive
The OpenMP 5.0 standard introduces more flexible work directives. One of these is the loop directive, which is simpler and closer to the functionality in OpenACC. The loop directive replaces the distribute parallel directive for simd. With the loop directive, you inform the compiler that loop iterations can be executed concurrently, but the actual implementation is left to the compiler. The following listing shows an example of using this directive in a stencil kernel.
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
The loop expression is actually a loop-independent or concurrent expression, which tells the compiler that the loop iterations have no dependencies. The loop expression provides information or a declarative expression to the compiler, rather than telling the compiler what to do, being a prescriptive expression. Most compilers don't implement this new functionality, so we continue working with prescriptive expressions from the previous examples in this chapter. If you're unfamiliar with these concepts, a definition of each is provided below.
3. OpenMP: A Heavyweight Champion Enters the Accelerator World
OpenMP specifications traditionally use prescriptive statements. This reduces variation between implementations and improves portability. However, in the case of GPU processors, this has led to long, complex directives with subtle differences in interthread synchronization capabilities and other hardware-specific functionality.
A descriptive approach is closer to the OpenACC philosophy and is not so constrained by hardware details. This gives the compiler the freedom and responsibility to correctly and efficiently generate code for the target software. Note that for OpenMP, this is not only a significant shift but also a significant one. If OpenMP continues to pursue the path of prescriptive directives as hardware complexity increases, the OpenMP language will become too complex, and source code portability will decrease.
4. Exercises
1. Find compilers suitable for your local GPU system. Are both OpenACC and OpenMP compilers available in your case? If not, do you have access to any systems that would allow you to try out these pragma-oriented languages?
2. Run the stream triad examples from the OpenACC/StreamTriad and/or OpenMP/StreamTriad directories on a local GPU-based development system. You can find these directories at https://github.com/EssentialsofParallelComputing/Chapter11.
3. Compare your results from Exercise 2 with the BabelStream results at https://uob-hpc.github.io/BabelStream/results/. For the stream triad, the bytes moved are 3 * nsize * sizeof(datatype).
4. Modify the OpenMP data region mapping in Listing 11.16 to reflect actual array usage in the compute kernels.
5. Implement the mass sum example from Listing 11.4 in OpenMP.
6. For arrays x and y of size 20,000,000, find the maximum radius in the arrays using both OpenMP and OpenACC. Initialize the arrays to double-precision values that increase linearly from 1.0 to 2.0e7 for the x array and decrease linearly from 2.0e7 to 1.0 for the y array.