OpenACC Labs

General Setup

For the hands-on sessions open a terminal on the local PC and log in to HAWK via

To work interactively on a compute node submit a job using

Several nodes on HAWK have been reserved specially for the course. R_dl is the reservation name which is only available during the course.

Further details can be found under:

https://kb.hlrs.de/platforms/index.php/HPE_Hawk

https://kb.hlrs.de/platforms/index.php/Big_Data,_AI_Aplications_and_Frameworks#Hawk

Copy the exercises to your own home directory via

This folder contains both C and Fortran versions of the code.

Initialise the environment on the compute node via

This runs the following commands in the current shell

In the labs you will implement 12 different version of the code for multicore CPUs and GPUs. Note the execution time for each of those versions in the results sheet handed out during the first hands-on session (also available online under https://doku.lrz.de/files/29609547/35881452/1/1688912940323/OpenACC-Labs-Resultsheet-v1.pdf).

Lab 1: Profiling with Nsight Systems

Check the configuration of the GPUs via the following commands which are provided with the NVIDIA HPC Software Development Kit (SDK) on the compute node:

We are using the NVIDIA HPC SDK C and Fortran compilers for the hands-on.

Compile the code using the following commands:

For C:

For Fortran:

Alternatively you can adopt the Makefile located in the C/Fortran directory and just run "make" to compile and execute the program.

To profile the code run the Nsight Systems command line interface on the compute node via:

This will produce the trace files laplace.nsys-rep and laplace.sqlite. Use "nsys --help" and "nsys --help profile" to get more information about the available profiler options.

For performance reasons, visualisation of the trace files should be always done on the local PCs.

Copy the trace files to your local PC in a terminal on the local PC via e.g.

Run "Nsight Systems 2021.2.1" on the local PC and open the trace file via File → Open. Zoom in the timeline and also test different views like "Top-Down view". Identify the 2 most time consuming routines of the code.

Further information is available under https://developer.nvidia.com/nsight-systems

Using NVIDIA Tools Extension (NVTX)

For C:

Add #include "nvtx3/nvToolsExt.h" in a copy jacobi-nvtx.c of the source code jacobi.c and wrap parts of the code which you want to capture events with calls to the NVTX API functions. For example, try adding nvtxRangePush("calc") before calling calcNext() and nvtxRangePop() just after calling it. We recommend to wrap calls to the functions initialize(), calcNext() and swap() and also wrap the while-loop in which the last 2 functions are called.

For C compile the NVTX annotated code as before:

For Fortran:

For Fortran codes an additional Fortran module nvtx in the file nvtx.f90 is needed to instrument the code. Add use nvtx in a copy jacobi-nvtx.f90 of the source code jacobi.f90 and wrap parts of the code which you want to capture events with calls to the NVTX API functions. For example, try adding call nvtxStartRange("calc") before calling calcNext() and call nvtxEndRange() just after calling it. We recommend to wrap calls to the functions initialize(), calcNext() and swap() and also wrap the while-loop in which the last 2 functions are called.

For Fortran compile the NVTX annotated code including the special nvtx module and link the library libnvToolsExt as follows:

Profile again on the compute node using

and visualise the profile data on the local PC.

Further information can be found in the CUDA Profiler User’s Guide https://docs.nvidia.com/cuda/profiler-users-guide/index.html

You can use the NVTX annotated code also for profiling in the next Labs if you want.

Lab 2: OpenACC Directives

Compiling for multicore CPUs

Compile the code for a muliticore CPU using the following commands

For C:

For Fortran:

Run and profile the code on the compute node and compare with the profiling results of the previous lab.

Using OpenACC for multicore CPUs with the parallel directive

For C:

Make a copy laplace2d-multicore-parallel.c of the file laplace2d.c  and include the following directives at the right place to parallelise the calcNext() and swap() functions:

For Fortran:

Make a copy laplace2d-multicore-parallel.f90 of the file laplace2d.f90  and include the following directives at the right place to parallelise the calcNext() and swap() functions:

Compile the code for multicore CPUs using the following commands:

For C:

For Fortran:

Run the code on the compute node.

Using OpenACC for multicore CPUs with the kernels directive

For C:

Make a copy laplace2d-multicore-kernels.c of the file laplace2d.c  and include the following directives at the right place to parallelise the calcNext() and swap() functions:

If you run into a situation where the compiler refuses to parallelise a loop, you may override the compilers decision (on your own risk!) via

For Fortran:

Make a copy laplace2d-multicore-kernels.f90 of the file laplace2d.f90  and include the following directives at the right place to parallelise the calcNext() and swap() functions:

If you run into a situation where the compiler refuses to parallelise a loop, you may override the compilers decision (on your own risk!) via

Compile the code for a muliticore CPU using the following commands

For C:

For Fortran:

Run the code on the compute node.

Profiling OpenACC code

Profile both versions of the code on the compute node via

Compare with the previous profiling results.

Lab 3: GPU Programming

Programming GPUs using OpenACC with CUDA managed memory

Compile the code using the following command including the option -gpu=managed to use CUDA managed memory:

For C:

Reuse the files laplace2d-multicore-parallel.c and laplace2d-multicore-kernels.c from the previous lab.

For Fortran:

Reuse the files laplace2d-multicore-parallel.f90 and laplace2d-multicore-kernels.f90 from the previous lab.

Run and profile both versions of the code on the compute node and compare with the previous profiling results.

Programming GPUs using OpenACC without CUDA managed memory

For C:

Make copies laplace2d-gpu-parallel.c and laplace2d-gpu-kernels.c of laplace2d-multicore-parallel.c and laplace2d-multicore-kernels.c, respectively.

Compile the code using the following commands

For Fortran:

Make copies laplace2d-gpu-parallel.f90 and laplace2d-gpu-kernels.f90 of laplace2d-multicore-parallel.f90 and laplace2d-multicore-kernels.f90, respectively.

Compile the code using the following commands

Check if the code compiles and runs correctly.

Modify the code to specify the sizes of the arrays to be copied from/to the GPU correctly. The sizes of the arrays can be specified with the copy, copyin, create etc. clauses of the parallel/kernels directive, e.g.:

In C:

In Fortran:

Run and profile both versions of the code on the compute node and compare with the previous profiling results.

Lab 4: Data Management

OpenACC Structured Data Directive

For C:

Make a copy jacobi-gpu-data-structured.c of the file jacobi.c  and include the following directive at the right place with the right data clauses (copy, copyin, create etc.) to create a data region:

#pragma acc data <data clauses>
{
...
}

Compile the code using the following commands

Run and profile both versions of the code on the compute node and compare with the previous profiling results.

For Fortran:

Make a copy jacobi-gpu-data-structured.f90 of the file jacobi.f90  and include the following directive at the right place with the right data clauses (copy, copyin, create etc.) to create a data region:

!$acc data <data clauses>
...
!$acc end data

Compile the code using the following commands

Run and profile both versions of the code on the compute node and compare with the previous profiling results.

OpenACC Unstructured Data Directives

For C:

Make copies laplace2d-gpu-data-unstructured-parallel.c and laplace2d-gpu-data-unstructured-kernels.c of laplace2d-gpu-parallel.c and laplace2d-gpu-kernels.c, respectively.

Include the following directive at the end of the  initialize() function with the right data clauses (copy, copyin, create, delete etc.):

#pragma acc enter data <data clauses>

Include the following directive at the beginning of the  deallocate() function with the right data clauses (copy, copyin, create, delete etc.):

#pragma acc exit data <data clauses>

Compile the code using the following commands

Run and profile both versions of the code on the compute node and compare with the previous profiling results.

For Fortran:

Make copies laplace2d-gpu-data-unstructured-parallel.f90 and laplace2d-gpu-data-unstructured-kernels.f90 of laplace2d-gpu-parallel.f90 and laplace2d-gpu-kernels.f90, respectively.

Include the following directive at the end of the  initialize() function with the right data clauses (copy, copyin, create, delete etc.):

!$acc enter data <data clauses>

Include the following directive at the beginning of the  deallocate() function with the right data clauses (copy, copyin, create, delete etc.):

!$acc exit data <data clauses>

Compile the code using the following commands

Run and profile both versions of the code on the compute node and compare with the previous profiling results.