MODULE SIX:

LOOP OPTIMIZATIONS
Speaker/Date

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

 Majority of program runtime is spent in loops

 Every loop can execute in a very different way
 Using OpenACC loop optimization, we can speed-up our most time-consuming
portions of code

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
SAMPLE LOOP CODE
Matrix multiplication

 Our code is a 3-Dimensional Matrix

Multiplication code for( i = 0; i < size; i++ )
for( j = 0; j < size; j++ )
 The code allows for many different for( k = 0; k < size; k++ )
levels and types of parallelism, and c[i][j] += a[i][k] * b[k][j];
works well with all of our loop clauses

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

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

 The auto clause tells the compiler to

decide whether or not the loop is #pragma acc parallel loop auto
parallelizable for( i = 0; i < size; i++ )
for( j = 0; j < size; j++ )
 The auto clause can be very useful for( k = 0; k < size; k++ )
when you are unsure of whether or not c[i][j] += a[i][k] * b[k][j];
a loop is safe to parallelize

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

 When using the kernels directive, the

auto clause is implied
#pragma acc kernels loop auto
 This means that you do not need to for( i = 0; i < size; i++ )
include the auto clause when using the for( j = 0; j < size; j++ )
kernels directive for( k = 0; k < size; k++ )
c[i][j] += a[i][k] * b[k][j];
 However, the auto clause can be very
useful when using the parallel directive

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
INDEPENDENT CLAUSE
 The independent clause asserts to the
compiler that the loop is parallelizable #pragma acc kernels loop independent
for( i = 0; i < size; i++ )
 This will overwrite any decision that the for( j = 0; j < size; j++ )
compiler makes about the loop for( k = 0; k < size; k++ )
c[i][j] += a[i][k] * b[k][j];
 Adding the independent clause could
force the compiler to parallelize a non-
parallel loop
 Allows the programmer to force
parallelism when using the kernels
directive

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
INDEPENDENT CLAUSE
 When using the parallel directive, the
independent clause is implied
#pragma acc parallel loop independent
 With the parallel directive, the for( i = 0; i < size; i++ )
programmer is determining which loops for( j = 0; j < size; j++ )
are parallelizable and thus the for( k = 0; k < size; k++ )
independent clause is not needed c[i][j] += a[i][k] * b[k][j];

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

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
SEQ CLAUSE
 The seq clause (short for sequential)
will tell the compiler to run the loop
sequentially
#pragma acc parallel loop
 In the sample code, the compiler will for( i = 0; i < size; i++ )
parallelize the outer loops across the #pragma acc loop
parallel threads, but each thread will for( j = 0; j < size; j++ )
run the inner-most loop sequentially #pragma acc loop seq
for( k = 0; k < size; k++ )
 The compiler may automatically apply c[i][j] += a[i][k] * b[k][j];
the seq clause to loops that have too
many dimensions

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

 The private clause allows the double tmp[3];

programmer to define a list of variables
as “thread-private”. #pragma acc kernels loop private(tmp[0:3])
for( i = 0; i < size; i++ )
 Each thread will be given a private {
copy of every variable in the comma- tmp[0] = <value>;
tmp[1] = <value>;
separated list tmp[2] = <value>;
}
 firstprivate is like private except that
the private values are initialized to the // note that the host value of “tmp”
same value used on the host. private // remains unchanged.
variables are uninitialized.

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

double tmp[3];
 Variables in private or
firstprivate clause are #pragma acc kernels loop private(tmp[0:3])
private to the loop level for( i = 0; i < size; i++ ) {
on which the clause // the tmp array is private to each iteration
appears. // of the outer loop
tmp[0] = <value>;
 Private variables on an tmp[1] = <value>;
outer loop are shared tmp[2] = <value>;
within inner loops. #pragma acc loop
for ( j = 0; j < size2; j++) {
// but tmp is shared amongst the threads
// in the inner loop
array[i][j] = tmp[0]+tmp[1]+tmp[2];
}
}
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
SCALARS AND PRIVATE CLAUSE

 By default, scalars are firstprivate when used in a parallel region and private when
used in a kernels region.
 Except in some cases, scalars do not need to be added to a private clause. These
cases may include but are not limited to:
1. Scalars with global storage such as global variables in C/C++, Module variables in
Fortran
2. When the scalar is passed by reference to a device subroutine
3. When the scalar is used as an rvalue after the compute region, aka “live-out”
 Note that putting scalars in a private clause may actually hurt performance!

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

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

 collapse( N ) #pragma acc parallel loop collapse(2)

for( i = 0; i < size; i++ )
 Combine the next N tightly nested loops for( j = 0; j < size; j++ )
double tmp = 0.0f;
 Can turn a multidimensional loop nest #pragma acc loop reduction(+:tmp)
into a single-dimension loop for( k = 0; k < size; k++ )
tmp += a[i][k] * b[k][j];
 This can be extremely useful for c[i][j] = tmp;
increasing memory locality, as well as
creating larger loops to expose more
parallelism

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

collapse( 2 )

(0,0) (0,1) (0,2) (0,3)

(1,0) (1,1) (1,2) (1,3)

(2,0) (2,1) (2,2) (2,3)

(3,0) (3,1) (3,2) (3,3)

#pragma acc parallel loop collapse(2)

for( i = 0; i < 4; i++ )
for( j = 0; j < 4; j++ )
array[i][j] = 0.0f;
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
TILE CLAUSE

 tile ( x , y , z, ...)

 Breaks multidimensional loops into

#pragma acc kernels loop tile(32, 32)
“tiles” or “blocks” for( i = 0; i < size; i++ )
for( j = 0; j < size; j++ )
 Can increase data locality in some for( k = 0; k < size; k++ )
codes c[i][j] += a[i][k] * b[k][j];
 Will be able to execute multiple “tiles”
simultaneously

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

(0,0) (0,1) (0,2) (0,3)

(0,0) (0,1) (0,2) (0,3)
#pragma acc kernels loop tile(2,2)
for(int x = 0; x < 4; x++){ (1,0) (1,1) (1,2) (1,3)
for(int y = 0; y < 4; y++){
array[x][y]++; (1,0) (1,1) (1,2) (1,3)
}
}
(2,0) (2,1) (2,2) (2,3)
(2,0) (2,1) (2,2) (2,3)
(3,0) (3,1) (3,2) (3,3)
(3,0) (3,1) (3,2) (3,3)

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

 Gang / Worker / Vector defines the

various levels of parallelism we can
achieve with OpenACC
 This parallelism is most useful when
parallelizing multi-dimensional loop Vector Workers
nests
 OpenACC allows us to define a generic Gang
Gang / Worker / Vector model that will
be applicable to a variety of hardware,
but we fill focus a little bit on a GPU
specific implementation

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
GANG WORKER VECTOR
 When paralleling our loops, the highest
level of parallelism is gang level
parallelism
 When encountering either the kernels or
parallel directive, multiple gangs will be
generated, and loop iterations will be Gang
spread across the gangs
 These gangs are completely
independent of each other, and there is
no way to for the programmer to know
exactly how many gangs are running at
a given time
 In many architecures, the gangs have
completely separate (or private) memory

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
GANG WORKER VECTOR
 In our code example, we see that we are
applying the gang clause to an outer-
loop
 This means that the outer-loop iterations
will be split across some number of
gangs Gang

 These gangs will then execute in parallel

with each other
#pragma acc parallel loop gang
 Whenever a parallel compute region is for( i = 0; i < N; i++ )
encountered, some number of gangs will for( j = 0; j < M; j++ )
be created < loop code >

 The programmer is able to specify

exactly how many gangs to create

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
GANG WORKER VECTOR
 A vector is the lowest level of
parallelism
Vector
 Every gang will have at least 1 vector
 A vector has the ability to run a single
instruction on multiple data elements
 Many different architectures can
implement vectors in different ways,
however, OpenACC allows for us to
define them in a general, non-hardware-
specific way

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

 In our code example, the inner-loop Vector

iterations will be evenly divided across a
vector
 This means that those loop iterations will
be executing in parallel with one-another
 Any loop that is inside of our vector loop
cannot be parallelized further #pragma acc parallel loop gang
for( i = 0; i < N; i++ )
#pragma acc loop vector
for( j = 0; j < M; j++ )
< loop code >

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
GANG WORKER VECTOR
 The worker clause is a way for the
programmer to have multiple vectors
within a gang 3 Workers

 The primary use of the worker clause is

to split up one large vector into multiple
smaller vectors
 This can be useful when our inner
parallel loops are very small, and will not
benefit from having a large vector

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

 In our sample code, we apply both gang

and worker level parallelism to our outer- 3 Workers
loop
 The main difference this creates for our
code is that we can now have smaller
vectors running the inner loop
 This will most likely improve #pragma acc parallel loop gang worker
performance if the inner loop is relatively for( i = 0; i < N; i++ )
small #pragma acc loop vector
for( j = 0; j < M; j++ )
< loop code >

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
PARALLEL DIRECTIVE SYNTAX
 When using the parallel directive, you may
define the number of gangs/workers/vectors #pragma acc parallel num_gangs(2) \
with num_gangs(N), num_workers(M), num_workers(2) vector_length(32)
vector_length(Q) {
#pragma acc loop gang worker
 Then, you may define where they belong in for(int x = 0; x < 4; x++){
the loops using gang, worker, vector #pragma acc loop vector
for(int y = 0; y < 32; y++){
array[x][y]++;
}
}
}

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

 You may also apply gang/worker/vector

when using the parallel loop construct

#pragma acc parallel loop num_gangs(2) num_workers(2) \

vector_length(32) gang worker
for(int x = 0; x < 4; x++){
#pragma acc loop vector
for(int y = 0; y < 32; y++){
array[x][y]++;
}
}

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
 KERNELS DIRECTIVE SYNTAX
 When using the kernels directive, the
process is somewhat simplified
#pragma acc kernels loop gang(2) worker(2)
 You may define the location and for(int x = 0; x < 4; x++){
number by using gang(N), #pragma acc loop vector(32)
worker(M), vector(Q) for(int y = 0; y < 32; y++){
array[x][y]++;
 You may also define gang, worker, }
and vector using the same method }
as with the parallel directive
 If you do not specify a number, the
compiler will decide one

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
 KERNELS DIRECTIVE SYNTAX
 When using the kernels directive, the #pragma acc kernels
process is somewhat simplified {
#pragma acc loop gang(2) worker(2)
 You may define the location and for(int x = 0; x < 4; x++){
number by using gang(N), #pragma acc loop vector(32)
worker(M), vector(Q) for(int y = 0; y < 32; y++){
array[x][y]++;
 You may also define gang, worker, }
}
and vector using the same method
as with the parallel directive #pragma acc loop gang(4) worker(4)
for(int x = 0; x < 16; x++){
 If you do not specify a number, the #pragma acc loop vector(16)
compiler will decide one for(int y = 0; y < 16; y++){
array2[x][y]++;
 Each loop nest can have different }
values for gang, worker, and vector }
}
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
GANG WORKER VECTOR

 We have a simple 2-
dimensional loop nest
 We have specified that
there is 1 worker and a
#pragma acc kernels loop gang worker(1) vector length of 8
for(int x = 0; x < 4; x++){
#pragma acc loop vector(8)  We do not specify how
for(int y = 0; y < 8; y++){ many gangs to generate,
array[x][y]++;
}
so the compiler will create
} enough gangs to cover
the loop

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
GANG WORKER VECTOR
Vector
Worker  The diagram shows a single
gang, though the compiler will be
able to generate as many gangs
Gang
as it wants
 These gangs are completely
#pragma acc kernels loop gang worker(1)
separate from each other, and
for(int x = 0; x < 4; x++){ are indistinguishable
#pragma acc loop vector(8)
for(int y = 0; y < 8; y++){  We will show these gangs apply
array[x][y]++; to a physical loop diagram, but
} this representation may not be
} 100% accurate to what the
compiler might decide

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
GANG WORKER VECTOR
#pragma acc kernels loop gang worker(1)
Vector for(int x = 0; x < 4; x++){
#pragma acc loop vector(8)
Worker for(int y = 0; y < 8; y++){
array[x][y]++;
}
Gang }

Vector
(0,0) (0,1) (0,2) (0,3) (0,4) (0,5) (0,6) (0,7) 1 Worker Gang
 The vectors are colored, so
(1,0) (1,1) (1,2) (1,3) (1,4) (1,5) (1,6) (1,7) that we can observe which
loop iterations they are
(2,0) (2,1) (2,2) (2,3) (2,4) (2,5) (2,6) (2,7) being applied to
 Based on the size of this
(3,0) (3,1) (3,2) (3,3) (3,4) (3,5) (3,6) (3,7) loop nest, the compiler will
(theoretically) generate 4
gangs
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
GANG WORKER VECTOR
Vector
Worker

Gang  We have now reduced the vector

length to 4, but have kept
everything else the same
 The dimension of the outer-loop is
#pragma acc kernels loop gang worker(1) still the same, and is still being
for(int x = 0; x < 4; x++){ distributed across the gangs, so the
#pragma acc loop vector(4) numbers of gangs will not change
for(int y = 0; y < 8; y++){
array[x][y]++;  Let’s observe how our code will
} function with a smaller vector size
}

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
GANG WORKER VECTOR
#pragma acc kernels loop gang worker(1)
Vector for(int x = 0; x < 4; x++){
#pragma acc loop vector(4)
Worker for(int y = 0; y < 8; y++){
array[x][y]++;
}
Gang }

Gang 1 Worker (0,0) (0,1) (0,2) (0,3) (0,4) (0,5) (0,6) (0,7)  We are still generating 4
gangs, but now each
vector is computing two
(1,0) (1,1) (1,2) (1,3) (1,4) (1,5) (1,6) (1,7) loop iterations

(2,0) (2,1) (2,2) (2,3) (2,4) (2,5) (2,6) (2,7)  If we wanted to generate
more gangs, we would
need to increase the
(3,0) (3,1) (3,2) (3,3) (3,4) (3,5) (3,6) (3,7) size of the outer-loop

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

Workers

Gang
 For our last trivial example,
let’s increase the number of
workers to 2

#pragma acc kernels loop gang worker(2)  There are now two vectors
for(int x = 0; x < 4; x++){ per gang, and each vector
#pragma acc loop vector(4) is of length 4
for(int y = 0; y < 8; y++){
array[x][y]++;
}
}

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
GANG WORKER VECTOR
#pragma acc kernels loop gang worker(2)
Vector for(int x = 0; x < 4; x++){
#pragma acc loop vector(4)
2 Workers for(int y = 0; y < 8; y++){
array[x][y]++;
}
Gang }

(0,0) (0,1) (0,2) (0,3) (0,4) (0,5) (0,6) (0,7)  Since we have
Gang 2 Workers increased the number
of workers, we will now
(1,0) (1,1) (1,2) (1,3) (1,4) (1,5) (1,6) (1,7) only generate 2 gangs

(2,0) (2,1) (2,2) (2,3) (2,4) (2,5) (2,6) (2,7)

(3,0) (3,1) (3,2) (3,3) (3,4) (3,5) (3,6) (3,7)

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
GANG WORKER VECTOR
Vector
Worker

Gang
 Now let’s look at a situation
where the gang/worker/vector
model is very useful
 We have reduced the size of our
#pragma acc kernels loop gang worker(1) inner-loop to 4 iterations
for(int x = 0; x < 4; x++){
#pragma acc loop vector(8)  Let’s try to run this loop with a
for(int y = 0; y < 4; y++){
array[x][y]++;
vector length of 8
}
}

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
GANG WORKER VECTOR
#pragma acc kernels loop gang worker(1)
Vector for(int x = 0; x < 4; x++){
#pragma acc loop vector(8)
Worker for(int y = 0; y < 4; y++){
array[x][y]++;
}
}
Gang

Gang 1 Worker (0,0) (0,1) (0,2) (0,3)  We can see that our
vector length is much
larger than our inner-
(1,0) (1,1) (1,2) (1,3) loop

(2,0) (2,1) (2,2) (2,3)  We are wasting half of

our vector, meaning
our code is performing
(3,0) (3,1) (3,2) (3,3) half as well as it could

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

Workers

Gang  We can fix this by breaking

our vector up among 2
workers
 Now instead of having 1
#pragma acc kernels loop gang worker(2) long vector, we have 2
for(int x = 0; x < 4; x++){ shorter vectors
#pragma acc loop vector(4)
for(int y = 0; y < 4; y++){  This setup should fit the
array[x][y]++; organization of our loop
} better
}

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
GANG WORKER VECTOR
#pragma acc kernels loop gang worker(2)
Vector for(int x = 0; x < 4; x++){
#pragma acc loop vector(4)
for(int y = 0; y < 4; y++){
Workers array[x][y]++;
}
}
Gang

Vector
(0,0) (0,1) (0,2) (0,3)
 We are no longer wasting a portion
2 Workers Gang of our vectors, since the smaller
Vector
(1,0) (1,1) (1,2) (1,3) vector size now fits our loop
properly
(2,0) (2,1) (2,2) (2,3)
 We always need to consider the
size of the loop when choosing the
(3,0) (3,1) (3,2) (3,3) gang worker vector dimensions

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
GANG WORKER VECTOR
Vector
Worker

Gang
 Another way we could have
fixed this problem is by
using the collapse clause

#pragma acc kernels loop collapse(2) gang worker(1) vector(8)

for(int x = 0; x < 4; x++){
for(int y = 0; y < 4; y++){
array[x][y]++;
}
}

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
GANG WORKER VECTOR
#pragma acc kernels loop collapse(2) \
Vector gang worker(1) vector(8)
for(int x = 0; x < 4; x++){
Worker for(int y = 0; y < 4; y++){
array[x][y]++;
}
}
Gang

collapse( 2 )

(0,0) (0,1) (0,2) (0,3) (1,0) (1,1) (1,2) (1,3) (2,0) (2,1) (2,2) (2,3) (3,0) (3,1) (3,2) (3,3)

(1,0) (1,1) (1,2) (1,3)  The collapse clause allows

us to combine two small
(2,0) (2,1) (2,2) (2,3) loops into a larger one

(3,0) (3,1) (3,2) (3,3)  This exposes additional

parallelism, and allows us
to use a longer vector
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
WARPS
 So far we have been using a very small number
of gangs/worker/vectors, simply because they’re
easier to understand
 When actually programming, the number of
gangs/worker/vectors will be much larger
 When specifically programming for an NVIDIA
GPU, you will always want your vectors large
enough to fully utilize warps
 A warp, simply put, is an optimized group of 32
threads
 To utilize warps in OpenACC, always make sure
that your vector length is a multiple of 32

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

 device_type ( <type> )

 Clauses that follow only apply to the #pragma acc parallel loop collapse(3)\
device_type(nvidia) \
specified device type. vector_length(256)
for( i = 0; i < size; i++ )
 This allows you to optimize for one type for( j = 0; j < size; j++ )
(GPU) without hurting the performance for( k = 0; k < size; k++ )
of another (CPU) c[i][j] += a[i][k] * b[k][j];
 Multiple device types can be specified
on a single directive.

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
LOOP OPTIMIZATION RULES OF THUMB

 It is rarely a good idea to set the number of gangs in your code, let the compiler
decide.
 Most of the time you can effectively tune a loop nest by adjusting only the vector
length.
 It is rare to use a worker loop. When the vector length is very short, a worker loop
can increase the parallelism in your gang.
 When possible, the vector loop should step through your arrays
 Use the device_type clause to ensure that tuning for one architecture doesn’t
negatively affect other architectures.

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…
 The loop directive enables the programmer to give more information to
the compiler about specific loops
 This information may be used for correctness or to improve
performance.
 The device_type clause allows the programmer to optimize for one
device type without hurting others.

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 in attempt to improve the
performance
 Use PGProf to analyze the performance difference when changing
your loops
 Experiment with the device_type clause to ensure GPU optimizations
don’t slow down the multicore speed-up, or vice versa

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