N-WAYS GPU BOOTCAMPOPENACCOPENACC
What to expect?
•Basic introduction to OpenACC directives•

HPC SDK Usage•

Portability across Multicore and GPU

2This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

OpenACC is…

Add Simple Compiler Directive

a directives-based

parallel programming model
designed formain(){<serial code>#pragma acc kernels{<parallel code>}}performance and portability.

3This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

4This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

OpenACC Directives
ManageData
Movement
Initiate
Parallel
Execution
Optimize
LoopMappings#pragma acc data copyin(a,b) copyout(c){...
#pragma acc parallel{#pragma acc loop gang vectorfor (i = 0; i < n; ++i) {c[i] = a[i] + b[i];...
}}
...
}

• Incremental• Single source• Interoperable• Performance portable• CPU, GPU, Manycore

5This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

OPENACC SYNTAXSyntax for using OpenACC directives in code
C/C++
#pragma acc directive clauses<code>

Fortran!$acc directive clauses<code>

A pragma in C/C++ gives instructions to the compiler on how to compile the code. Compilers that do not understand aparticular pragma can freely ignore it.
A directive in Fortran is a specially formatted comment that likewise instructions the compiler in it compilation of thecode and can be freely ignored.
“acc” informs the compiler that what will come is an OpenACC directive

Directives are commands in OpenACC for altering our code.
Clauses are specifiers or additions to directives.

6
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

OPENACC PARALLEL DIRECTIVE
Expressing parallelism#pragma acc parallel{When encountering theparallel directive, thecompiler will generate
1 or more parallelgangs, which executeredundantly.
}

gang

gang

gang

gang

gang

gang
7

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

OPENACC PARALLEL DIRECTIVE
loop
gang

This loop will be
executed redundantly
on each gang

gang

loopgangloop

for(int i = 0; i < N; i++){// Do Something}looploop

gang

loop

#pragma acc parallel{}loop

Expressing parallelism

gang

gang

8This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

OPENACC PARALLEL DIRECTIVE
Expressing parallelism#pragma acc parallel{#pragma acc loopfor(int i = 0; i < N; i++){// Do Something}}

The loop directive
informs the compiler
which loops toparallelize.
9This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

OPENACC PARALLEL DIRECTIVE
Parallelizing a single loop
C/C++
#pragma acc parallel{#pragma acc loopfor(int i = 0; j < N; i++)a[i] = 0;}

Use a parallel directive to mark a region of code where you wantparallel execution to occur
This parallel region is marked by curly braces in C/C++ or a start andend directive in FortranThe loop directive is used to instruct the compiler to parallelize theiterations of the next loop to run across the parallel gangs

Fortran!$acc parallel!$acc loop
do i = 1, Na(i) = 0
end do!$acc end parallel
10
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

OPENACC PARALLEL DIRECTIVE
Parallelizing a single loop
C/C++
#pragma acc parallel loopfor(int i = 0; j < N; i++)a[i] = 0;

Fortran!$acc parallel loop
do i = 1, Na(i) = 0
end doThis pattern is so common that you can do all of this in a single line ofcode
In this example, the parallel loop directive applies to the next loop
This directive both marks the region for parallel execution anddistributes the iterations of the loop.
When applied to a loop with a data dependency, parallel loop may
produce incorrect results

11This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

BUILD AND RUN THE CODE
12This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

NVIDIA HPC SDK• Comprehensive suite of compilers, libraries, and tools used to GPU accelerate HPC
modeling and simulation application• The NVIDIA HPC SDK includes the new NVIDIA HPC compiler supporting OpenACC C andFortran•The command to compile C code is ‘nvc’•

The command to compile C++ code is ‘nvc++’•

The command to compile Fortran code is ‘nvfortran’nvc –fast –Minfo=accel –ta=tesla:managed main.cornvc -fast -Minfo=accel -acc=gpu -gpu=managed main.c

nvfortran –fast –Minfo=accel –ta=tesla:managed main.f90
ornvfortran -fast -Minfo=accel -acc=gpu -gpu=managed main.f90
13This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)BUILDING THE CODE-Minfo shows more details$ nvc -fast -ta=multicore -Minfo=accel laplace2d_uvm.cmain:
63, Generating Multicore code
64, #pragma acc loop gang
64, Accelerator restriction: size of the GPU copy of Anew,A is unknown
Generating reduction(max:error)
66, Loop is parallelizable$ nvc -fast -ta=tesla:managed -Minfo=accel rdf.c
main:
63, Accelerator kernel generated
Generating Tesla code
64, #pragma acc loop gang /* blockIdx.x */Generating reduction(max:error)
66, #pragma acc loop vector(128) /* threadIdx.x */63, Generating implicit copyin(A[:])
Generating implicit copy(error)
66, Loop is parallelizable1414This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

RDF

Pseudo Codefor (int frame=0;frame<nconf;frame++){for(int id1=0;id1<numatm;id1++){for(int id2=0;id2<numatm;id2++){
dx=d_x[]-d_x[];dy=d_y[]-d_y[];dz=d_z[]-d_z[];r=sqrtf(dx*dx+dy*dy+dz*dz);}

if (r<cut) {ig2=(int)(r/del);d_g2[ig2] = d_g2[ig2] +1 ;}• Across Frames• Find Distance• Reduction}}

15
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

RDF

Pseudo Code -Cfor (int frame=0;frame<nconf;frame++) {• Parallel Loop construct#pragma acc parallel loopfor(int id1=0;id1<numatm;id1++) {for(int id2=0;id2<numatm;id2++) {dx=d_x[]-d_x[];dy=d_y[]-d_y[];dz=d_z[]-d_z[];r=sqrtf(dx*dx+dy*dy+dz*dz);}

if (r<cut) {ig2=(int)(r/del);#pragma acc atomic
d_g2[ig2] = d_g2[ig2] +1 ;}• Atomic Construct}}

16
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

RDF

Pseudo Code - Fortrando iconf=1,nframes
if (mod(iconf,1).eq.0) print*,iconf!$acc parallel loop
do i=1,natoms
do j=1,natoms
dx=x(iconf,i)-x(iconf,j)dy=y(iconf,i)-y(iconf,j)dz=z(iconf,i)-z(iconf,j)...
if(r<cut)then!$acc atomicg(ind)=g(ind)+1.0d0
endif
enddo
enddo
enddo

• Parallel Loop construct• Atomic Construct

17This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

OPENACC SPEEDUP
Speed-up
450.00X
393.78X

400.00X
350.00X

Speed-Up

300.00X
250.00X
200.00X
150.00X
100.00X
50.00X
0.00X

1.00X
SERIAL

17.04X
MULTICORE

NVIDIA TESLA V100
HPC SDK 20.11, NVIDIA Tesla V100, DGX1
18This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

REFERENCES

https://www.openacc.org/get-started
https://developer.nvidia.com/hpc-sdk

19This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

ADDITIONAL EXERCISECONTENT
22
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

PROCESSING FLOW – STEP 1

PCI Bus

1. Copy input data from CPU memory to GPUmemory

23
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

PROCESSING FLOW – STEP 2

PCI Bus

1. Copy input data from CPU memory to GPU
memory2. Load GPU program and execute,caching data on chip for performance24
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

PROCESSING FLOW – STEP 3

PCI Bus

1. Copy input data from CPU memory to GPU
memory2. Load GPU program and execute,caching data on chip for performance
3. Copy results from GPU memory to CPU
memory4. Unified Memory changes the nature of flow
• Some of the basics remains same

25
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

CUDA UNIFIED MEMORY
Simplified Developer Effort

Without Managed Memory

Commonly referred to as“managed memory.”

With Managed Memory

CPU and GPU memories are
combined into a single, shared pool

SystemMemory

GPU Memory

Managed Memory

26
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

MANAGED MEMORY
Limitations

With Managed Memory•The programmer will almost always be able to get better
performance by manually handling data transfers•Memory allocation/deallocation takes longer with managedmemory•Cannot transfer data asynchronously•Currently only available from NVIDIA Compiler on NVIDIA GPUs.

Managed Memory

27
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

CUDA MANAGED MEMORYUsefulness• OpenACC:•

Enabled using –ta=tesla:managed or –gpu=managed• std::par•nvc++ –fast –ta=tesla:managed –Minfo=accel main.c
or
nvc++ -fast –acc=gpu –gpu=managed –Minfo=accel main.c

nvc++ -stdpar=gpu program.cpp -o program

All allocations uses managed memory• OpenMP:•Current Beta release does not support Unified memory. Need explicitly map target directive tocopy data
28
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

OPENACC DATA DIRECTIVE
Definition
The data directive defines a lifetime for data on thedevice beyond individual loops
During the region data is essentially “owned by” theaccelerator
Data clauses express shape and data movement for
the region#pragma acc data clauses{< Sequential and/or Parallel code >}!$acc data clauses< Sequential and/or Parallel code >!$acc end data

29
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

DATA CLAUSEScopy( list )Allocates memory on GPU and copies data from host to GPU whenentering region and copies data to the host when exiting region.
Principal use: For many important data structures in your code, this is alogical default to input, modify and return the data.

copyin( list )Allocates memory on GPU and copies data from host to GPU whenentering region.
Principal use: Think of this like an array that you would use as just aninput to a subroutine.

copyout( list )Allocates memory on GPU and copies data to the host when exiting
region.
Principal use: A result that isn’t overwriting the input data structure.

create( list )Allocates memory on GPU but does not copy.
Principal use: Temporary arrays.

30

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

ARRAY SHAPING

Sometimes the compiler needs help understanding the shape of an array
The first number is the start index of the array
In C/C++, the second number is how much data is to be transferredIn Fortran, the second number is the ending index

copy(array[starting_index:length])

C/C++

copy(array(starting_index:ending_index))

Fortran
31

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

ARRAY SHAPING (CONT.)Multi-dimensional Array shaping

copy(array[0:N][0:M])

C/C++

Both of these examples copy a 2D array to the device

copy(array(1:N, 1:M))

Fortran32
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

STRUCTURED DATA DIRECTIVEExample#pragma acc data copyin(a[0:N],b[0:N]) copyout(c[0:N]){#pragma acc parallel loopfor(int i = 0; i < N; i++){c[i] = a[i] + b[i];}}Host Memory

Device memory

A B C
C’

A B C’

ActionAllocate A
Deallocate
ExecuteAllocate
Copy CABloop
from
CB
C
B on
A
fromononCPUdevicedevicetoto
deviceCPUdevice

33
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)