MODULE FIVE:

DATA MANAGEMENT
Speaker, Date

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
MODULE OVERVIEW
OpenACC Data Management

 Explicit Data Management

 OpenACC Data Regions and Clauses
 Unstructured Data Lifetimes
 Data Synchronization

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

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

 Data must be visible on the device when Host

we run our parallel code
Device
 Data must be visible on the host when we
run our sequential code
 When the host and device don’t share
memory, data movement must occur
 To maximize performance, the Host
programmer should avoid all unnecessary Memory
data transfers Device
Memory

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
EXPLICIT MEMORY MANAGEMENT
Key problems

 Many parallel accelerators (such as Host

devices) have a separate memory space
from the host Device
 These separate memories can become
out-of-sync and contain completely
different data
 Transferring between these two memories
can be a very time consuming process Host
Memory
Device
Memory

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
EXPLICIT MEMORY MANAGEMENT device
CPU
Key problems
$ $ $ $ $ $

 Many parallel accelerators (such as

$ $ $ $ $ $

Shared Cache
devices) have a separate memory pool
from the host
$ $ $ $ $ $ $ $

 These separate memories can become

out-of-sync and contain completely Shared Cache

different data
 Transferring between these two memories
can be a very time consuming process CPU
Memory IO Bus
device
Memory

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

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 #pragma acc data clauses
{
for data on the device
 During the region data should be < Sequential and/or
thought of as residing on the Parallel code >
accelerator
}
 Data clauses allow the programmer
to control the allocation and !$acc data clauses
movement of data
< Sequential and/or
Parallel code >

!$acc end data

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
DATA CLAUSES
copy( list ) Allocates memory on device and copies data from host to device
when entering region and copies data to the host when exiting region.

Principal use: For many important data structures in your code, this is a
logical default to input, modify and return the data.

copyin( list ) Allocates memory on device and copies data from host to device
when entering region.

Principal use: Think of this like an array that you would use as just an
input to a subroutine.

copyout( list ) Allocates memory on device 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 device but does not copy.

Principal use: Temporary arrays.

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 transferred
 In Fortran, the second number is the ending index

copy(array[starting_index:length]) C/C++

copy(array(starting_index:ending_inde
Fortran
x))
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)) Fortran

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
ARRAY SHAPING (CONT.)
Partial Arrays

copy(array[i*N/4:N/4]) C/C++

Both of these examples copy only ¼ of the full array

copy(array(i*N/4:i*N/4+N/4)) Fortran

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

This parallel loop will

execute on the #pragma acc parallel loop
accelerator, so a, b, for(int i = 0; i < N; i++){
and c must be visible c[i] = a[i] + b[i];
on the accelerator. }

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

#pragma acc data copyin(a[0:N],b[0:N])

Start of copyout(c[0:N])
Data Region {
#pragma acc parallel loop
for(int i = 0; i < N; i++){
c[i] = a[i] + b[i];
End of }
Data Region }

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
STRUCTURED DATA DIRECTIVE
Example
#pragma acc data copyin(a[0:N],b[0:N])
copyout(c[0:N]) Action
{
#pragma acc parallel loop Deallocate
Execute
Allocate
Copy CA
Bloop
from
CA
B
C
A
B on
from
on
for(int i = 0; i < N; i++){ CPUdevice
device
totodevice
CPU
c[i] = a[i] + b[i];
}
}

Host Memory Device memory

A B C’
C A B C’
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
IMPLIED DATA REGIONS

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
IMPLIED DATA REGIONS
Definition
 Every kernels and parallel region has
an implicit data region surrounding it
 This allows data to exist solely for the #pragma acc kernels
duration of the region copyin(a[0:100])
{
 All data clauses usable on a data for( int i = 0; i < 100; i+
+ )
directive can be used on a parallel and {
kernels as well a[i] = 0;
}
}

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
IMPLIED DATA REGIONS
Explicit vs Implicit Data Regions

Explicit Implicit
#pragma acc data copyin(a[0:100])
{
#pragma acc kernels
{ #pragma acc kernels copyin(a[0:100])
for( int i = 0; i < {
100; i++ ) for( int i = 0; i < 100; i++ )
{ {
a[i] = 0; a[i] = 0;
} }
} }
}

These two codes are functionally the same.

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
EXPLICIT VS. IMPLICIT DATA REGIONS
Limitation

Explicit 1 Data Copy Implicit 2 Data Copies

#pragma acc data copyout(a[0:100])
{
#pragma acc kernels
#pragma acc kernels copyout(a[0:100])
{ {
a[i] = i; a[i] = i;
} }

#pragma acc kernels #pragma acc kernels

{ copy(a[0:100])
a[i] = 2 * a[i]; {
} a[i] = 2 * a[i];
}
}

The code on the left will perform better than the code on the right.

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

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
 UNSTRUCTURED DATA DIRECTIVES
Enter Data Directive
 Data lifetimes aren’t always neatly #pragma acc enter data clauses
structured.
 The enter data directive handles < Sequential and/or
Parallel code >
device memory allocation
 You may use either the create or the #pragma acc exit data clauses
copyin clause for memory allocation
!$acc enter data clauses
 The enter data directive is not the start
of a data region, because you may < Sequential and/or
have multiple enter data directives Parallel code >

!$acc exit data clauses

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
 UNSTRUCTURED DATA DIRECTIVES
Exit Data Directive
 The exit data directive handles device
memory deallocation #pragma acc enter data clauses

 You may use either the delete or the < Sequential and/or
copyout clause for memory deallocation Parallel code >
 You should have as many exit data for a #pragma acc exit data clauses
given array as enter data
 These can exist in different functions !$acc enter data clauses

< Sequential and/or

Parallel code >

!$acc exit data clauses

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

copyin ( list ) Allocates memory on device and copies data from host to device
on enter data.
copyout ( list ) Allocates memory on device and copies data back to the host on
exit data.
create ( list ) Allocates memory on device without data transfer on enter data.
delete ( list ) Deallocates memory on device without data transfer on exit data

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
UNSTRUCTURED DATA DIRECTIVES
Basic Example

#pragma acc parallel loop

for(int i = 0; i < N; i++){
c[i] = a[i] + b[i];
}

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
UNSTRUCTURED DATA DIRECTIVES
Basic Example

#pragma acc enter data copyin(a[0:N],b[0:N])

create(c[0:N])

#pragma acc parallel loop

for(int i = 0; i < N; i++){
c[i] = a[i] + b[i];
}

#pragma acc exit data copyout(c[0:N])

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
UNSTRUCTURED DATA DIRECTIVES
Basic Example

#pragma acc enter data copyin(a[0:N],b[0:N])

create(c[0:N]) Action
#pragma acc parallel loop
Copy C A
B
for(int i = 0; i < N; i++){ Execute
Deallocate
Allocateloop
from
C
A
BC
c[i] = a[i] + b[i]; from
on
device
CPU toto
} device
device
CPU

#pragma acc exit data copyout(c[0:N])

CPU MEMORY device MEMORY

C A B C’
A B C’ This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
UNSTRUCTURED DATA DIRECTIVES
Basic Example – proper memory deallocation

#pragma acc enter data copyin(a[0:N],b[0:N])

create(c[0:N]) Action
#pragma acc parallel loop
for(int i = 0; i < N; i++){ Deallocate A
B
c[i] = a[i] + b[i]; from
} device

#pragma acc exit data copyout(c[0:N]) delete(a,b)

CPU MEMORY device MEMORY

C A B
A B C’ This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
UNSTRUCTURED VS STRUCTURED
With a simple code
Unstructured Structured
 Can have multiple starting/ending points  Must have explicit start/end points
 Can branch across multiple functions  Must be within a single function
 Memory exists until explicitly deallocated  Memory only exists within the data region
#pragma acc enter data #pragma acc data copyin(a[0:N],b[0:N]) \
copyin(a[0:N],b[0:N]) \ create(c[0:N]) copyout(c[0:N])
{
#pragma acc parallel loop #pragma acc parallel loop
for(int i = 0; i < N; i++){ for(int i = 0; i < N; i++){
c[i] = a[i] + b[i]; c[i] = a[i] + b[i];
} }

#pragma acc exit data copyout(c[0:N]) \ }

delete(a,b)

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
UNSTRUCTURED DATA DIRECTIVES
Branching across multiple functions
int* allocate_array(int N){
int* ptr = (int *) malloc(N *
sizeof(int));  In this example enter data and exit data are
#pragma acc enter data
create(ptr[0:N]) in different functions
return ptr;
}  This allows the programmer to put device
void deallocate_array(int* ptr){ allocation/deallocation with the matching
#pragma acc exit data delete(ptr) host versions
free(ptr);
}
 This pattern is particularly useful in C++,
int main(){ where structured scopes may not be
int* a = allocate_array(100);
#pragma acc kernels possible.
{
a[0] = 0;
}
deallocate_array(a);
}

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

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
OPENACC UPDATE DIRECTIVE
update: Explicitly transfers data between the host and the device
Useful when you want to synchronize data in the middle of a data region
Clauses:
self: makes host data agree with device data
device: makes device data agree with host data

#pragma acc update self(x[0:count])

#pragma acc update device(x[0:count])
C/C++
!$acc update self(x(1:end_index))
!$acc update device(x(1:end_index))
Fortran
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
 OPENACC UPDATE DIRECTIVE
#pragma acc update device(A[0:N])

The data must exist on

A A*
A
both the CPU and device
CPU Memory device Memory
for the update directive
to work.

B*
B B*
#pragma acc update self(A[0:N])
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
SYNCHRONIZE DATA WITH UPDATE
int* allocate_array(int N){
int* A=(int*)
malloc(N*sizeof(int));  Inside the initialize function we alter the
#pragma acc enter data
create(A[0:N]) host copy of ‘A’
return A;
}  This means that after calling initialize the
host and device copy of ‘A’ are out-of-sync
void deallocate_array(int* A){
#pragma acc exit data  We use the update directive with the
delete(A)
free(A); device clause to update the device copy of
} ‘A’
void initialize_array(int* A, int N){  Without the update directive later compute
for(int i = 0; i < N; i++){
A[i] = i;
regions will use incorrect data.
}
#pragma acc update
device(A[0:N])
}
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
C/C++ STRUCTS/CLASSES

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
C STRUCTS typedef struct {
Without dynamic data members float x, y, z;
} float3;
 Dynamic data members are anything
int main(int argc, char* argv[]){
contained within a struct that can have int N = 10;
a variable size, such as dynamically float3* f3 = malloc(N *
allocated arrays sizeof(float3));
 OpenACC is easily able to copy our #pragma acc enter data
struct to device memory because create(f3[0:N])
everything in our float3 struct has a
#pragma acc kernels
fixed size for(int i = 0; i < N; i++){
f3[i].x = 0.0f;
 But what if the struct had dynamically f3[i].y = 0.0f;
allocated members? f3[i].z = 0.0f;
}

#pragma acc exit data delete(f3)

free(f3);
}
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
C STRUCTS typedef struct {
float *arr;
With dynamic data members int n;
} vector;
 OpenACC does not have enough int main(int argc, char* argv[]){
information to copy the struct and its
dynamic members vector v;
v.n = 10;
 You must first copy the struct into [Link] = (float*)
malloc(v.n*sizeof(float));
device memory, then allocate/copy the
dynamic members into device memory #pragma acc enter data copyin(v)
#pragma acc enter data
 To deallocate, first deal with the create([Link][0:v.n])
dynamic members, then the struct
...
 OpenACC will automatically attach
your dynamic members to the struct #pragma acc exit data
delete([Link])
#pragma acc exit data delete(v)
free([Link]);
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
C++ STRUCTS/CLASSES
class vector {
With dynamic data members private:
float *arr;
int n;
 C++ Structs/Classes work the same public:
vector(int size){
exact way as they do in C n = size;
arr = new float[n];
 The main difference is that now we #pragma acc enter
have to account for the implicit “this” data copyin(this)
pointer #pragma acc enter
data create(arr[0:n])
}
~vector(){
#pragma acc exit
data delete(arr)
#pragma acc exit
data delete(this)
delete(arr);
}
};
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
 C++ CLASS DATA SYNCHRONIZATION
[Link]();

 Since data is encapsulated, the class

needs to be extended to include data
synchronization methods
 Including explicit methods for
host/device synchronization may ease
C++ data management
CPU Memory device Memory
 Allows the class to be able to naturally
handle synchronization, creating less
code clutter [Link]();

void accUpdateSelf() {
#pragma acc update self(arr[0:n])
}
void accUpdateDevice() {
#pragma acc update
device(arr[0:n])
}
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
 USING A OPENACC AWARE C++ CLASS
#include “vector.h"

int main() { A A
vector A(N), B(N);
for (int i=0; i < [Link](); ++i) {
B[i]=2.5;
}
[Link](); B B
#pragma acc parallel loop present(A,B)
for (int i=0; i < [Link](); ++i) {
A[i]=B[i]+i;
}
[Link]();
for(int i=0; i<10; ++i) {
cout << "A[" << i << "]: " << A[i] << endl;
}
exit(0);
}

Host Memory Device Memory

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

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
KEY CONCEPTS
In this module we discussed…
 Why explicit data management is necessary for best performance
 Structured and Unstructured Data Lifetimes
 Explicit and Implicit Data Regions
 The data, enter data, exit data, and update directives
 Data Clauses

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
LAB ASSIGNMENT
In this module’s lab you will…
 Update the code from the previous module to use explicit data
directives
 Analyze the different between using CUDA Managed Memory and
explicit data management in the lab code.

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
 ADDITIONAL RESOURCES
YouTube OpenACC Introduction Series by Michael Wolfe
Introduction to Parallel Programming with OpenACC – Part 5

Follow along by downloading the code here!

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