2 Comments An emerging standard uses pragmas to move parallel computations in C/C++ and Fortran to the gpu

By Rob Farber, June 11, 2012

This is the first in a series of articles by Rob Farber on OpenACC directives, which enable existing C/C++ and Fortran code to run with high performance on massively parallel devices such as GPUs. The magic in OpenACC lies in how it extends the familiar face of OpenMP pragma programming to encompass coprocessors. As a result, OpenACC opens the door to scalable, massively parallel GPU — accelerating millions of lines of legacy application code without requiring a new language such as CUDA or OpenCL, or fork application source tree to support multiple languages

OpenACC is a set of standardized, high-level pragmas that enables C/C++ and Fortran programmers to utilize massively parallel coprocessors with much of the convenience of OpenMP. A pragma is a form of code annotation that informs the compiler of something about the code. In this case, it identifies the succeeding block of code or structured loop as a good candidate for parallelization. OpenMP is a well-known and widely supported standard that defines pragmas programmers have used since 1997 to parallelize applications on shared memory multicore processors. The OpenACC standard has generated excitement because it preserves the familiarity of OpenMP code annotation while extending the execution model to encompass devices that reside in separate memory spaces. To support coprocessors, OpenACC pragmas annotate data placement and transfer as well as loop and block parallelism.

The success of GPU computing in recent years has motivated compiler vendors to extend the OpenMP shared memory pragma programming approach to coprocessors. Approved by the OpenACC standards committee in November 2011, the OpenACC version 1.0 standard creates a unified syntax and prevents a "tower of babel" proliferation of incompatible pragmas. Adoption has been rapid by companies such as NVIDIA, PGI (The Portland Group), CAPS Enterprise, and Cray.

Pragmas and high-level APIs are designed to provide software functionality. They hide many details of the underlying implementation to free a programmer's attention for other tasks.A colleague humorously refers to pragma-based programming as a negotiation that occurs between the developer and the compiler. Note that pragmas are informational statements provided by the programmer to the assist the compiler. This means that pragmas are not subject to the same level of syntax, type, and sanity checking as the rest of the source code. The compiler is free to ignore any pragma for any reason including: it does not support the pragma, syntax errors, code complexity, unresolved (or potentially unresolved) dependencies, edge cases where the compiler cannot guarantee that vectors or matrices do not overlap, use of pointers, and many others. Profiling tools and informational messages from the compiler about parallelization, or an inability to parallelize, are essential to a successful to achieving high performance.

An OpenACC pragma for C/C++ can be identified from the string "#pragma acc" just like an OpenMP pragma can be identified from "#pragma omp". Similarly, Fortran pragmas can be identified by "! $acc". Always ensure that these strings begin all OpenACC (or OpenMP) pragmas. Moreover, it is legal to mix OpenMP, OpenACC, and other pragmas in a single source file.

OpenACC Syntax

OpenACC provides a fairly rich pragma language to annotate data location, data transfer, and loop or code block parallelism. The syntax of OpenACC pragmas (sometimes referred to as OpenACC directives) is:

• 
  C/C++: "#pragma acc directive-name [clause [[,] clause]…] new-line"
  
• 
  Fortran: "!$acc directive-name [clause [[,] clause]…] new-line"
  

Table 1 shows a list of OpenACC version 1.0 pragmas and clauses.

Two OpenACC environment variables, ACC_DEVICE_TYPE and ACC_DEVICE_NUM can be set by the user:

• 
  ACC_DEVICE_TYPE: Controls the default device type to use when executing accelerator parallel and kernels regions, when the program has been compiled to use more than one different type of device.
  The allowed values of this environment variable are implementation-defined.
  Examples include ACC_DEVICE_TYPE=NVIDIA.
  
• 
  ACC_DEVICE_NUM: Specifies the default device number to use when executing accelerator regions. The value of this environment variable must be a nonnegative integer between zero and the number of devices of the desired type attached to the host.
  If the value is zero, the implementation-defined default is used.
  If the value is greater than the number of devices attached, the behavior is implementation-defined.
  On multi-GPU systems, this variable will avoid the TDR (Timeout Detection and Recovery) watchdog reset for long-running GPU applications by running on the GPU that is not used for the display. (Consult the vendor driver information to see how to modify the TDR time for your operation system.
  

Vendor specific information can be found on the Nvidia, PGI, CAPS, and Cray websites.

This tutorial uses The Portland Group (PGI) Accelerator C and Fortran compilers release 12.5 with OpenACC support. PGI has been deeply involved in developing pragma-based programming for coprocessors since 2008, plus they are a founding member of the OpenACC standards body. The PGI OpenACC compilers currently target NVIDIA GPUs, but it is important to note that OpenACC can support other coprocessors (such as AMD GPUs and Intel MIC) as well. More information about the PGI compilers is available on the company's website.

An extended 30-day trial license for the PGI software can be obtained by registering with NVIDIA. The Portland Group also provides a free 15 day OpenACC trial license, which can be obtained by following the following three steps:

1. Download any of the available software packages for your operating system.
2. Review the PGI Installation Guide [PDF] or the PGI Visual Fortran Installation Guide [PDF] and configure your environment.
3. Generate the trial license keys. Note the trial keys and all executable files compiled using them will cease operating at the end of the trial period.

The following set of examples multiply two matrices a and b and store the result in matrix c. They utilize a useful set of basic OpenACC data transfer, parallelization, and memory creation/access clauses. A C-language OpenMP matrix multiply is also provided to show the similarity between OpenACC and OpenMP and provide CPU and GPU performance comparisons. While the PGI matrix multiplication performance is good, please look to the highly optimized BLAS (Basic Linear Algebra Subroutines) packages such as CUBLAS and phiGEMM for production GPU and hybrid CPU + GPU implementations.

Following is our first OpenACC program, matix-acc-check.c. This simple code creates a static set of square matrices (a,b,c,seq), initializes them, and then performs a matrix multiplication on the OpenACC device. The test code then performs the matrix multiplication sequentially on the host processor and double-checks the OpenACC result.

?

1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55

/* matrix-acc-check.c */ #define SIZE 1000 float a[SIZE][SIZE]; float b[SIZE][SIZE]; float c[SIZE][SIZE]; float seq[SIZE][SIZE];    int main() {   int i,j,k;        // Initialize matrices.   for (i = 0; i < SIZE; ++i) {     for (j = 0; j < SIZE; ++j) {       a[i][j] = (float)i + j;       b[i][j] = (float)i - j;       c[i][j] = 0.0f;     }   }        // Compute matrix multiplication. #pragma acc kernels copyin(a,b) copy(c)   for (i = 0; i < SIZE; ++i) {     for (j = 0; j < SIZE; ++j) {       for (k = 0; k < SIZE; ++k) {     c[i][j] += a[i][k] * b[k][j];       }     }   }      // ****************   // double-check the OpenACC result sequentially on the host   // ****************   // Initialize the seq matrix   for(i = 0; i < SIZE; ++i)      for(j = 0; j < SIZE; ++j)        seq[i][j] = 0.f;        // Perform the multiplication   for (i = 0; i < SIZE; ++i)      for (j = 0; j < SIZE; ++j)        for (k = 0; k < SIZE; ++k)      seq[i][j] += a[i][k] * b[k][j];        // check all the OpenACC matrices   for (i = 0; i < SIZE; ++i)     for (j = 0; j < SIZE; ++j)       if(c[i][j] != seq[i][j]) {     printf("Error %d %d\n", i,j);     exit(1);       }   printf("OpenACC matrix multiplication test was successful!\n");        return 0; }

The OpenACC pragma tells the compiler the following:

• 
  #pragma acc: This is an OpenACC pragma.
  
• 
  kernels: A kernels region.
  No jumps are allowed into/out of the kernels region.
  Loops will be sent to the OpenACC device.
  The scope of the kernels region code block is denoted by the curly brackets in a C program.
  
• 
  copyin(): copy the contiguous region of memory from the host to the device.
  The variables, arrays or subarrays in the list have values in the host memory that need to be copied to the device memory.
  If a subarray is specified, then only that subarray of the array needs to be copied.
  
• 
  copy(): copy the contiguous memory region from the host to the device and back again.
  The variables, arrays or subarrays in the list have values in the host memory that need to be copied to the device memory.
  If a subarray is specified, then only that subarray of the array needs to be copied.
  The data is copied to the device memory before entry to the kernles region, and data copied back to the host memory when the code block is complete.
  

Note the similarity between matrix-acc.c and the following OpenMP implementation, matrix-omp.c. Only the pragmas are different as the OpenACC pragma includes copy operations that are not required in the OpenMP implementation.

Fortran programmers will find the corresponding source code in Example 4. Again, the OpenACC pragmas annotate data movement with the copy() and copyin() clauses. Note that the C-based pragmas know the extent of the code block due to the use of curly brackets while the Fortran version must explicitly specify the end of the scope of the pragma with "!$acc end …".

The following commands compile the source code for each application with the PGI C and Fortran compilers. These commands assume the source code has been saved to the file name provided in the comment at the beginning of each example.

pgcc -fast -mp -Minfo -Mconcur=allcores matrix-omp.c -o matrix-omp
pgcc -fast -acc -Minfo matrix-acc.c -o matrix-acc-gpu
pgfortran -fast -acc -Minfo matrix-acc.f -o matrix-acc-gpuf

The command line arguments to the PGI C compiler (pgcc) and Fortran compiler (pgfortran) are:

• 
  -fast: Chooses generally optimal flags for the target platform.
  
• 
  -mp: Interpret OpenMP pragmas to explicitly parallelize regions of code for execution by multiple threads on a multi-processor system.
  
• 
  -acc: Interpret OpenACC pragmas.
  
• 
  -Minfo: Emit useful information to stderr.
  
• 
  -Mconcur: Instructs the compiler to enable auto-concurrentization of loops.
  

This output from the PGI runtime profiling tells us that the application spent 7 milliseconds transferring data and 43 milliseconds computing the matrix multiply kernel.

Notice that there are:

• 
  Three host to device data transfers at the start of the computation. These transfers correspond to the copyin() clauses for matrices a and b plus the copy() clause for matrix c.
  
• 
  A GPU computation that requires 39.1% of the time for kernel main_24_gpu. A helpful feature of the PGI OpenACC compiler is that it intelligently labels the kernel with the routine name and line number to make these timelines intelligible.
  
• 
  A single data transfer back from the device to the host, which was required by the copy clause for matrix c at the end of the kernel.
  

For example, the following screenshot shows the initial analysis of the timeline shown in Figure 1: