GPU Basics: OpenACC Tutorial

OpenAcc is programming tool for parallel computing designed to simplify parallel programming of heterogeneous CPU/GPU systems. It utilizes sets of compiler directives to enable the compiler to generate cuda code for the GPU. Most of the compiler directives are entered in the user’s code as comments. This example shows a minimal conversion of a Vector Addition CPU code to an OpenACC accelerator directives version. Consider this an OpenACC ‘Hello World’. Modifications from the CPU version will be briefly discussed. This tutorial will also show you how to compile run and submit jobs for CADES gpu nodes.

Prepare Your Programming Environment

Cades uses modules to control the paths for your software. For example, if no modules are loaded, the LD_LIBRARY_PATH, which is the default library path accessed to check for available dynamic and shared libraries, is set to this default:


However, if the pgi/19.4 module is loaded, LD_LIBRARY_PATH is reset to hold the pgi/19.4 libraries.

$ module load pgi/19.4

The ‘module show pgi/19.4’ will show you what other paths the module sets.

As a best practice on CADES SHPC condo, always clear your modules before running or compiling code. This is done with ‘module purge’. To use OpenACC on CADES Open SHPC condo, you will need to use the pgi/19.4 compiler module.

Prepare your environment:

$ module purge
$ module load pgi/19.4


VecAdd.c adds two vectors A and B to produce C, where Ci = Ai + Bi. This version is by Adam Simpson and can be downloaded from OLCF Git Hub.

#include <stdio.h>
#include <stdlib.h>
#include <math.h>

int main( int argc, char* argv[] )

    // Size of vectors
    int n = 10000;

    // Input vectors
    double *restrict a;
    double *restrict b;
    // Output vector
    double *restrict c;

    // Size, in bytes, of each vector
    size_t bytes = n*sizeof(double);

    // Allocate memory for each vector
    a = (double*)malloc(bytes);
    b = (double*)malloc(bytes);
    c = (double*)malloc(bytes);

    // Initialize content of input vectors, vector a[i] = sin(i)^2 vector b[i] = cos(i)^2
    int i;
    for(i=0; i<n; i++) {
        a[i] = sin(i)*sin(i);
        b[i] = cos(i)*cos(i);

    // sum component wise and save result into vector c
    #pragma acc kernels copyin(a[0:n],b[0:n]), copyout(c[0:n])
    for(i=0; i<n; i++) {
        c[i] = a[i] + b[i];

    // Sum up vector c and print result divided by n, this should equal 1 within error
    double sum = 0.0;
    for(i=0; i<n; i++) {
        sum += c[i];
    sum = sum/n;
    printf("final result: %f\n", sum);

    // Release memory

    return 0;

Changes to VecAdd.c

There are two main differences in this code that make it vectorize well and be able to run on the GPU. 1. The restrict key word:

// Input vectors
double *restrict a;
double *restrict b;
// Output vector
double *restrict c;

The restrict keyword is a non-enforced guarantee to the compiler that the pointers are not aliased. This is not required by the OpenACC standard but in some cases allows the compiler to better vectorize the code.

  1. The pragma:
// sum component wise and save result into vector c
#pragma acc kernels copyin(a[0:n],b[0:n]), copyout(c[0:n])
for(i=0; i<n; i++) {
    c[i] = a[i] + b[i];

The compiler will analyze code following a pragma to see if it can be run on the GPU. Areas of code where many operations can be done at once, without depending on each other are good candidates for the GPU.

In this case, the acc kernels directive will make the sperate iterations of the loop run in parallel on the GPU and memory is copied from the CPU to the GPU at the start of the loop and back from the GPU to the CPU at the end of the loop.

There are many different OpenAcc directives that can be used to move code to the GPU. This is just one example. To see more go to

Compiling VecAdd.c with OpenAcc

To compile this code with the pgcc compiler and OpenAcc use the -acc flag:

$ module purge
$ module load pgi/19.4
$ pgcc -acc vecAdd.c -o VecAdd.c


The pgi compiler -Minfo flag allow you to see what the compiler is doing with your code:

Recompile with the -Minfo flag:

$ pgcc -acc -Minfo vecAdd.c -o vecAdd.o
     32, Generating copyin(a[:n])
         Generating copyout(c[:n])
         Generating copyin(b[:n])
     33, Loop is parallelizable
         Generating Tesla code
         33, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */


Your CADES home directory sits on the Network File System (NFS). HPC code needs to run in the Lustre Parallel file system.

Copy vecAdd.c and vecAdd.o to your Lustre directory. Below the example of the CADE Open Lustre and Birthright condo are used.

cp vecAdd.c /lustre/or-scratch/cades-birthright/user_id
cp vecAdd.o /lustre/or-scratch/cades-birthright/user_id

One way to run this is to start an interactive job on the compute nodes with srun:

$srun -A birthright -p gpu -N 1 -n 1 -c 2 --gres=gpu:2 --mem=8G -t 1:00:00 --pty /bin/bash
srun: job 7034 queued and waiting for resources

Note that the, partition for using the gpu must be specified with -p gpu and --gres==gpu:2 must be specified.

When the job starts:

$ module purge
$ module load pgi/19.4
$ mpirun ./vecAdd.o

final result: 1.000000

Did the code really run on the GPU?

There are a few ways to check if your job really ran on the GPU.


Nvprof is nvida's built-in profiler. It will show you that your code is running on the GPU and also give you performance information about the code.

To use nvprof issue:

mpirun nvprof ./vecAdd.o

If the code ran on the GPU you will get a result like this:

==119959== NVPROF is profiling process 119959, command: ./vecAddC.o
final result: 1.000000
==119959== Profiling application: ./vecAddC.o
==119959== Profiling result:
Time(%)      Time     Calls       Avg       Min       Max  Name
 61.77%  27.040us         2  13.520us  12.320us  14.720us  [CUDA memcpy HtoD]
 29.17%  12.768us         1  12.768us  12.768us  12.768us  [CUDA memcpy DtoH]
  9.06%  3.9680us         1  3.9680us  3.9680us  3.9680us  main_33_gpu

==119959== API calls:
Time(%)      Time     Calls       Avg       Min       Max  Name
 59.74%  139.67ms         1  139.67ms  139.67ms  139.67ms  cuDevicePrimaryCtxRetain
 21.15%  49.448ms         1  49.448ms  49.448ms  49.448ms  cuDevicePrimaryCtxRelease
 15.39%  35.987ms         1  35.987ms  35.987ms  35.987ms  cuMemHostAlloc
  3.19%  7.4473ms         1  7.4473ms  7.4473ms  7.4473ms  cuMemFreeHost
  0.21%  487.30us         1  487.30us  487.30us  487.30us  cuMemAllocHost
  0.20%  458.15us         4  114.54us  5.1670us  219.86us  cuMemAlloc
  0.04%  97.782us         1  97.782us  97.782us  97.782us  cuModuleLoadData
  0.03%  62.684us         1  62.684us  62.684us  62.684us  cuLaunchKernel
  0.02%  54.570us         2  27.285us  7.5190us  47.051us  cuMemcpyHtoDAsync
  0.01%  23.821us         3  7.9400us  2.7650us  12.937us  cuStreamSynchronize
  0.01%  12.297us         1  12.297us  12.297us  12.297us  cuMemcpyDtoHAsync
  0.00%  9.2830us         1  9.2830us  9.2830us  9.2830us  cuStreamCreate
  0.00%  7.6830us         3  2.5610us     675ns  3.5170us  cuPointerGetAttributes
  0.00%  6.6520us         2  3.3260us  1.1020us  5.5500us  cuEventCreate
  0.00%  5.8830us         1  5.8830us  5.8830us  5.8830us  cuEventRecord
  0.00%  4.2680us         3  1.4220us     177ns  3.6380us  cuDeviceGetCount
  0.00%  4.2390us        20     211ns     103ns     777ns  cuDeviceGetAttribute
  0.00%  4.1070us        12     342ns     108ns     863ns  cuDeviceGet
  0.00%  3.8610us         6     643ns     124ns  1.7480us  cuCtxSetCurrent
  0.00%  3.2310us         1  3.2310us  3.2310us  3.2310us  cuEventSynchronize
  0.00%  3.1580us         1  3.1580us  3.1580us  3.1580us  cuModuleGetFunction
  0.00%     869ns         4     217ns     106ns     507ns  cuDeviceComputeCapability
  0.00%     331ns         1     331ns     331ns     331ns  cuDriverGetVersion
  0.00%     261ns         1     261ns     261ns     261ns  cuCtxGetCurrent

==119959== OpenACC (excl):
Time(%)      Time     Calls       Avg       Min       Max  Name
 98.76%  36.448ms         1  36.448ms  36.448ms  36.448ms  acc_enter_data@vecAdd.c:32
  0.34%  124.50us         1  124.50us  124.50us  124.50us  acc_device_init@vecAdd.c:32
  0.19%  69.908us         1  69.908us  69.908us  69.908us  acc_enqueue_launch@vecAdd.c:33 (main_33_gpu)
  0.19%  68.985us         2  34.492us  9.3040us  59.681us  acc_enqueue_upload@vecAdd.c:32
  0.16%  58.884us         1  58.884us  58.884us  58.884us  acc_wait@vecAdd.c:37
  0.14%  52.340us         1  52.340us  52.340us  52.340us  acc_exit_data@vecAdd.c:32
  0.11%  39.090us         1  39.090us  39.090us  39.090us  acc_enqueue_download@vecAdd.c:37
  0.05%  17.071us         1  17.071us  17.071us  17.071us  acc_compute_construct@vecAdd.c:32
  0.04%  16.552us         1  16.552us  16.552us  16.552us  acc_wait@vecAdd.c:32
  0.03%  9.5830us         1  9.5830us  9.5830us  9.5830us  acc_wait@vecAdd.c:33
  0.00%       0ns         3       0ns       0ns       0ns  acc_delete@vecAdd.c:37
  0.00%       0ns         3       0ns       0ns       0ns  acc_alloc@vecAdd.c:32
  0.00%       0ns         3       0ns       0ns       0ns  acc_create@vecAdd.c:32

Note that in the first line of the timing information, you can see that about 62% of the code’s running time was spent coping memory from the "CPU host" to the "GPU device". To really make code run faster using the GPU, memory should be move to and from the GPU as little as is strictly needed.

Running OpenACC code with a batch script

A batch script is the most convenient way to submit code to the compute nodes. The following is an example for submitting vecAdd.o to the GPU using the birthright condo.



#SBATCH -A birthright
#SBATCH -p -gpu
#SBATCH -n 1
#SBATCH -c 1
#SBATCH --gres=gpu:2
#SBATCH -J gpu-test-job
#SBATCH --mem=1G
#SBATCH -t 10:00
#SBATCH -o ./%j-gpu-output.txt
#SBATCH -e ./%j-gpu-error.txt
#SBATCH --mail-type=FAIL
#SBATCH --mail-user=<your_email>

module purge
module load pgi/19.4
mpirun ./vecAdd.o

The --gpus option above is used to specify the number of GPU resources your job needs. An additional specifier can be added to indicate the type of GPU that is needed if there are multiple models available. To specify K80 GPUs, the above option would change to --gpus=k80:2.

New Arguments

 --gres=gpu:2 : Reserve nodes with the specified gpu resources

Submit the job script

To submit your job script to the compute nodes issue:

sbatch run_vecadd.sbatch