University:
Boston University
Course:
GRS CS 910 | Directed Study: Systems
Academic year:

2023
Views:

252

Pages:

17
Author:

Braydon Hess

Introduction to OpenACC Outline • Introduction to GPU and OpenACC • Basic syntax and the first OpenACC program: SAXPY • Kernels vs. parallel directives • An example: Laplace solver in OpenACC The first try Data transfer between GPU and CPU/host Data race and the reduction clause • GPU and OpenACC task granularities GPU and GPGPU • Originally, graphics processing unit (GPU) is dedicated for manipulating computer graphics and image processing. Traditionally GPU is known as “video card”. • GPU’s highly parallel structure makes it efficient for parallel programs. Nowadays GPUs are used for tasks that were formerly the domain of CPUs, such as scientific computation. This kind of GPU is called general-purpose GPU (GPGPU) . • In many cases, a parallel program runs faster on GPU than on CPU. Note that a serial program runs faster on CPU than on GPU. • The most popular type of GPU in the high-performance computing world is NVIDIA GPU. We will only focus on NVIDIA GPU here. GPU is an accelerator • GPU is a device on a CPU-based system. GPU is connected to CPU through PCI bus. • Computer program can be parallelized and accelerated on GPU. • CPU and GPU has separated memories. Data transfer between CPU and GPU is required for programming. Three ways to accelerate applications on GPU What is OpenACC • OpenACC (for Open Accelerators) is a programming standard for parallel computing on accelerators (mostly on NIVDIA GPU). • It is designed to simplify GPU programming. • The basic approach is to insert special comments (directives) into the code so as to offload computation onto GPUs and parallelize the code at the level of GPU (CUDA) cores. • It is possible for programmers to create an efficient parallel OpenACC code with only minor modifications to a serial CPU code. What are compiler directives?  The directives tell the compiler or runtime to ……  Generate parallel code for GPU GPU  Allocate GPU memory and copy input data  Execute parallel code on GPU  Copy output data to CPU and deallocate GPU memory  The first OpenACC directive: kernels  ask the compiler to generate a GPU code  let the compiler determine safe parallelism and data transfer . // ... serial code ... #pragma acc kernels for (int i= 0; i prevents parallelization Loop carried backward dependence of x-> prevents vectorization Accelerator scalar kernel generated Loop carried backward dependence of x-> prevents vectorization  The compiler creates a serial program, which runs slower on GPU than on CPU! Pointer aliasing in C (1)  An improper version of the SAXPY code (using pointers): int N=1000; float a = 3.0f; float * x = (float*)malloc(N * sizeof(float)); float * y = (float*)malloc(N * sizeof(float)); for (int i = 0; i < N; ++i) { x[i] = 2.0f; y[i] = 1.0f; } #pragma acc kernels for (int i = 0; i < N; ++i) { y[i] = a * x[i] + y[i]; }  Pointer aliasing: Different pointers are allowed to access the same object. This may induce implicit data dependency in a loop.  In this case, it is possible that the pointers x an y access to the same object. Potentially there is data dependency in the loop. Pointer aliasing in C (2)  The compiler refuses to parallelize the loop that is involved with pointer aliasing.  Compiling output of the improper SAXPY code: …… 20, Loop carried dependence of y-> prevents parallelization Complex loop carried dependence of x-> prevents parallelization Loop carried backward dependence of y-> prevents vectorization Accelerator scalar kernel generated Use restrict to avoid pointer aliasing  A proper version of the SAXPY code (using pointers): int N=1000; float a = 3.0f; float *x = (float*)malloc(N * sizeof(float)); float * restrict y = (float*)malloc(N * sizeof(float)); for (int i = 0; i < N; ++i) { x[i] = 2.0f; y[i] = 1.0f; } #pragma acc kernels for (int i = 0; i < N; ++i) { y[i] = a * x[i] + y[i]; }  To avoid pointer aliasing, use the keyword restrict.  restrict means: For the lifetime of the pointer ptr, only it or a value directly derived from it (such as ptr + 1) will be used to access the object to which it points. Parallel directive (1)  An improper version of SAXPY code (using parallel directive): C #pragma acc parallel for (int i = 0; i < N; ++i) { y[i] = a * x[i] + y[i]; } Fortran !$acc parallel do i=1,n y(i) = a*x(i)+y(i) enddo !$acc end parallel  The parallel directive tells the compiler to create a parallel region. But differently from the kernels region, the code in the parallel region (the loop in this case) is executed (by all gangs) redundantly. There is no work sharing (among gangs)! Parallel directive (2)  A proper version of SAXPY code (using parallel loop directive): C #pragma acc parallel loop for (int i = 0; i < N; ++i) { y[i] = a * x[i] + y[i]; } Fortran !$acc parallel loop do i=1,n y(i) = a*x(i)+y(i) enddo !$acc end parallel loop  It is necessary to add the keyword to loop share the works (among gangs).  In Fortran, the keyword loop can be replaced by do here.  In C, the keyword loop can be replaced by for.

Introduction to OpenACC

Related Documents

Recommended Documents

Description

Related Documents

Free up your schedule!

Take 5 seconds to unlock