[Click image to view at full size]

The Visual Profiler shows that only one data transfer occurs at the end of the data region as required by the copyout() clause.

Removal of the data transfers speeds the OpenACC performance over the OpenMP version:

Speedup with copyin() and copy() clauses over OpenMP: 6.4x

Speedup with create() and copyout() clauses over OpenMP: 6.9x

The previous examples utilized static 2D globally accessible C arrays. Most applications utilize dynamic allocation of all data structures including multidimensional arrays that are frequently passed to functions and subroutines. A particular challenge for C/C++ programmers is that OpenACC transfers occur between contiguous regions of host and device memory. The use of non-contiguous multidimensional arrays (such as float ** arrays) is not recommended because they require individual transfers of each contiguous memory region.

The following example, matrix-acc-func.c, dynamically allocates the 2D matrices for the test and passes them to doTest(), which performs the matrix initializations and multiplication. This test utilizes the C-language (as of C99) restrict keyword that indicates the matrices do not overlap. For convenience, the 2D nature of the arrays was defined in the function to make array accesses straightforward:

?

1 2 3 4 5 6

int doTest(restrict float a[][SIZE], restrict float b[][SIZE],        restrict float c[][SIZE], int size) {   …     c[i][j] = 0.0f; }

• 
  Example 7: Code snippet for straightforward 2D array indexing.
  
• 
  Of course, the programmer can pass the pointer to the contiguous region of memory and manually calculate the offsets into the multidimensional array as will be demonstrated in the next example.
  
• 
  ?
  

The matrix-acc-func.c example also demonstrates the use of the OpenACC pragma "#pragma acc loop independent". The independent clause tells the compiler to ignore its own dependency analysis and trust that the programmer knows the loops have no dependencies. Incorrect and non-deterministic program behavior can happen if the programmer is mistaken. Conversely, the OpenACC pragma "#pragma acc loop seq" tells the compiler to generate code that will execute sequentially on the device.

OpenACC also provides the ability to use previously allocated device memory with the deviceptr() clause. The following example matrix-acc-alloc.c demonstrates how to allocate memory in main() with the OpenACC runtime method acc_malloc(). The pointer is then passed to doTest() where it is accessed via deviceptr(). The copyout() clause also includes the size of the contiguous region of memory. For timing purposes, this code utilizes a size specified by the user on the command-line.

• 
  Example 10: Source code for matrix-acc-alloc.c.
  

OpenACC has been designed to provide OpenMP-style programmers with an easy transition to GPU programming. Following the common sense adage, "Make your life easy and use the highest level API first," OpenACC provides a natural starting point to transition any C or Fortran code to massive parallelism. For legacy code, OpenACC can be the only viable route to massively parallel coprocessors because it eliminates the need for a total rewrite of the software and Fortran is supported. As a result, OpenACC opens the door to scalable, massively parallel GPU (or, more generically, coprocessor) acceleration of millions of lines of legacy application code. Currently, OpenACC is supported by compilers that must be purchased from either PGI or CAPS Enterprise. The PGI compiler used in this article is free for evaluation but continued use after the trial period expires requires a license that must be purchased. As with OpenMP, it is assumed that open-source compilers will eventually provide free OpenACC support.

Profiling and informational compiler messages play a key role in achieving high performance in pragma-based programming. Instead of having to blindly add pragmas and then guess at the impact of each might have on an application, free tools like the NVIDIA Visual Profiler let the developer actually see what is happening during runtime on Windows, Linux, and Mac computers. Being able to see what effect OpenACC pragmas have on runtime behavior greatly speeds the OpenACC learning process as well as application acceleration.

My next article in this series will discuss the OpenACC memory and execution model including the gang and worker clauses plus more sophisticated ways to handle data.