MODULE TWO:
PROFILING
Dr. Volker Weinberg | LRZ
MODULE OVERVIEW
Topics to be covered
Compiling and profiling sequential code
Explanation of multicore programming
Compiling and profiling multicore code
COMPILING SEQUENTIAL CODE
NVIDIA’S HPC COMPILERS (AKA PGI)
NVIDIA Compiler Names (PGI names still work)
nvc - The command to compile C code (formerly known as ‘pgcc’)
nvc++ - The command to compile C++ code (formerly known as ‘pgc++’)
nvfortran - The command to compile Fortran code (formerly known As
pgfortran/pgf90/pgf95/pgf77)
The -fast flag instructs the compiler to optimize the code to the best of its abilities
$ nvc –fast main.c $ pgcc –fast main.c
$ nvc++ -fast main.cpp $ pgc++ -fast main.cpp
$ nvfortran –fast main.F90 $ pgfortran –fast main.F90
NVIDIA’S HPC COMPILERS (AKA PGI)
-Minfo flag
The Minfo flag will instruct the compiler to print feedback about the compiled code
-Minfo=accel will give us information about what parts of the code were accelerated
via OpenACC
-Minfo=opt will give information about all code optimizations
-Minfo=all will give all code feedback, whether positive or negative
$ pgcc –fast –Minfo=all main.c
$ pgc++ -fast -Minfo=all main.cpp
$ pgfortran –fast –Minfo=all main.f90
NVIDIA NSIGHT FAMILY
Nsight Product Family
Workflow
Nsight Systems -
Analyze application
algorithm system-wide
Nsight Compute -
Debug/optimize CUDA
kernel
Nsight Graphics -
Debug/optimize graphics
workloads
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
Thread/core
migration
Processes
and
threads Thread state
CUDA and
OpenGL API trace
cuDNN and
cuBLAS trace
Kernel and memory
transfer activities
Multi-GPU
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
PROFILING SEQUENTIAL CODE
OPENACC DEVELOPMENT CYCLE
Analyze your code to determine
most likely places needing Analyze
parallelization or optimization.
Parallelize your code by starting
with the most time consuming parts,
check for correctness and then
analyze it again.
Optimize your code to improve
observed speed-up from
parallelization.
Optimize Parallelize
PROFILING SEQUENTIAL CODE
Step 1: Run Your Code Terminal Window
$ pgcc –fast jacobi.c laplace2d.c
Record the time it takes for your $ ./a.out
sequential program to run. 0, 0.250000
100, 0.002397
200, 0.001204
300, 0.000804
Note the final results to verify
400, 0.000603
correctness later.
500, 0.000483
600, 0.000403
700, 0.000345
Always run a problem that is 800, 0.000302
representative of your real jobs. 900, 0.000269
total: 39.432648 s
PROFILING SEQUENTIAL CODE
Step 2: Profile Your Code Lab Code: Laplace Heat Transfer
Obtain detailed information about how
the code ran. Total Runtime: 39.43 seconds
This can include information such as:
Total runtime
Runtime of individual routines swap
calcNext
19.04s
Hardware counters 21.49s
Identify the portions of code that took
the longest to run. We want to focus on
these “hotspots” when parallelizing.
PROFILING WITH NSIGHT SYSTEM
AND NVTX
PROFILING SEQUENTIAL CODE
Using Command Line Interface (CLI)
NVIDIA Nsight Systems CLI provides
Simple interface to collect data
Can be copied to any system and analysed later
Profiles both serial and parallel code
For more info enter nsys --help on the terminal
To profile a serial application with NVIDIA Nsight Systems, we use NVIDIA Tools Extension
(NVTX) API functions in addition to collecting backtraces while sampling.
PROFILING SEQUENTIAL CODE
NVIDIA Tools Extension API (NVTX) library
What is it?
A C-based Application Programming Interface (API) for annotating events
Can be easily integrated to the application
Can be used with NVIDIA Nsight Systems
Why?
Allows manual instrumentation of the application
Allows additional information for profiling (e.g: tracing of CPU events and time ranges)
How?
Import the header only C library nvToolsExt.h
Wrap the code region or a specific function with nvtxRangePush() and nvtxRangPop()
#include <string.h> -t Selects the APIs to be traced (nvtx in this example)
#include <stdio.h>
#include <stdlib.h>
#include <omp.h> --status if true, generates summary of statistics after the collection
#include "laplace2d.h"
#include <nvtx3/nvToolsExt.h> -b Selects the backtrace method to use while sampling. The option dwarf
int main(int argc, char** argv) uses DWARF's CFI (Call Frame Information).
{
const int n = 4096; --force-overwrite if true, overwrites the existing results
const int m = 4096;
const int iter_max = 1000;
-o sets the output (qdrep) filename
const double tol = 1.0e-6;
double error = 1.0;
double *restrict A = (double*)malloc(sizeof(double)*n*m);
double *restrict Anew = (double*)malloc(sizeof(double)*n*m);
nvtxRangePushA("init");
initialize(A, Anew, m, n);
nvtxRangePop();
printf("Jacobi relaxation Calculation: %d x %d mesh\n", n, m);
double st = omp_get_wtime();
int iter = 0;
nvtxRangePushA("while");
while ( error > tol && iter < iter_max )
{
nvtxRangePushA("calc");
error = calcNext(A, Anew, m, n);
nvtxRangePop();
nvtxRangePushA("swap");
swap(A, Anew, m, n);
nvtxRangePop();
if(iter % 100 == 0) printf("%5d, %0.6f\n", iter, error);
iter++; NVTX range
} statistics
nvtxRangePop();
double runtime = omp_get_wtime() - st;
printf(" total: %f s\n", runtime);
“calc” region (calcNext function) takes 26.6%
deallocate(A, Anew); “swap” region (swap function) takes 23.4% of
return 0; total execution time
}
Open laplace-seq.qdrep with
jacobi.c Nsight System GUI to view the
(starting and ending of ranges are timeline
highlighted with the same color)
PROFILING SEQUENTIAL
CODE
Using Nsight Systems
Open the generated report files
(*.qdrep) from command line in the
Nsight Systems profiler.
File > Open
PROFILING SEQUENTIAL
CODE
Using Nsight Systems
Navigate through the “view selector”.
“Analysis summary” shows a summary of the profiling
session. To review the project configuration used to
generate this report, see next slide.
“Timeline View” contains the timeline at the top, and a
bottom pane that contains the events view and the
function table.
Read more: https://docs.nvidia.com/nsight-systems
PROFILING SEQUENTIAL
CODE
Using Nsight Systems
Timeline view
(charts and the hierarchy on the top pane)
Analysis Summary
Timeline view
(event view and function table on the bottom pane)
PROFILING SEQUENTIAL
CODE Enlarge view!
Using Nsight Systems
right click in
selected region
and Zoom into
selection!
PROFILING SEQUENTIAL
CODE
Using Nsight Systems
PROFILING SEQUENTIAL
CODE
Viewing captured NVTX events and time
ranges via Nsight Systems GUI
From the Timeline view, right click on the “NVTX” from
the top pane and choose “Show in Events View”.
From the bottom pane, you can now see name of the
events captured with the duration.
PLEASE START LAB NOW!
CSC_OPENACC_AMBASSADOR_MAY22
TRAINING SETUP
To get started, follow these steps:
Create an NVIDIA Developer account at http://courses.nvidia.com/join Select "Log in
with my NVIDIA Account" and then '"Create Account“
Visit http://courses.nvidia.com/dli-event and enter the event code
LRZ_OPENACC_AMBASSADOR_MY22
TRAINING SETUP
TRAINING SETUP
TRAINING SETUP
TRAINING SETUP
TRAINING SETUP
TRAINING SETUP
TRAINING SETUP
To be able to visualise Nsight System profiler output during the course, please install
Nsight System latest version on your local system before the course. The software
can be downloaded from https://developer.nvidia.com/nsight-systems.
PROFILING MULTICORE CODE
PROFILING MULTICORE CODE
What is multicore?
Multicore refers to using a CPU with multiple CPU
computational cores as our parallel device
These cores can run independently of each
other, but have shared access to memory
Loop iterations can be spread across CPU
threads and can utilize SIMD/vector instructions
(SSE, AVX, etc.)
Parallelizing on a multicore CPU is a good
starting place, since data management is
unnecessary
PROFILING MULTICORE CODE
Using a multicore CPU with OpenACC
OpenACC’s generic model involves a Host
combination of a host and a device =
Device
Host generally means a CPU, and the device
is some parallel hardware
When running with a multicore CPU as our
device, typically this means that our
host/device will be the same
Host
This also means that their memories will be Memory =
the same Device
Memory
PROFILING MULTICORE CODE
Compiling code for a specific parallel hardware
The ‘-ta’ flag will allow us to compile our code for a specific, target parallel hardware
‘ta’ stands for “Target Accelerator,” an accelerator being another way to refer to a
parallel hardware
Our OpenACC code can be compiled for many different kinds of parallel hardware
without having to change the code
$ pgcc –fast –Minfo=accel –ta=multicore laplace2d.c
calcNext:
35, Generating Multicore code
36, #pragma acc loop gang
PROFILING MULTICORE CODE
Compiling code for a specific parallel hardware
nsys profile -t nvtx --stats=true --force-overwrite true -o
laplace_parallel ./laplace_parallel
PROFILING OPENACC CODE
PARALLEL VS
SEQUENTIAL
Compiler feedback
Have a close look at the PGI compiler
Sequential
feedback for both sequential and parallel
implementation of the application.
It provides information about how your program
was optimized or why a particular optimization
was not made.
Note: Adding –Minfo flag or -Minfo=accel
or -Minfo=all when compiling, will enable
compiler feedback messages, giving details
about the parallel code generated.
Parallel
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)
#include <math.h>
#include <stdlib.h>
#define OFFSET(x, y, m) (((x)*(m)) + (y))
void initialize(double *restrict A, double *restrict Anew, int m, int n)
{
memset(A, 0, n * m * sizeof(double));
memset(Anew, 0, n * m * sizeof(double));
for(int i = 0; i < m; i++){
A[i] = 1.0;
Anew[i] = 1.0; CUDA API
}
}
statistics
double calcNext(double *restrict A, double *restrict Anew, int m, int n)
{
double error = 0.0;
#pragma acc parallel loop reduction(max:err)
for( int j = 1; j < n-1; j++)
{
#pragma acc loop
for( int i = 1; i < m-1; i++ )
{ CUDA Kernel
Anew[OFFSET(j, i, m)] = 0.25 * ( A[OFFSET(j, i+1, m)] + A[OFFSET(j, i-1, m)] statistics
+ A[OFFSET(j-1, i, m)] + A[OFFSET(j+1, i, m)]);
error = max( error, fabs(Anew[OFFSET(j, i, m)] - A[OFFSET(j, i , m)]));
}
}
return error;
} CUDA
Memory
void swap(double *restrict A, double *restrict Anew, int m, int n) Operation
{
#pragma acc parallel loop statistics
for( int j = 1; j < n-1; j++)
{
#pragma acc loop
for( int i = 1; i < m-1; i++ )
{
A[OFFSET(j, i, m)] = Anew[OFFSET(j, i, m)];
}
}
}
void deallocate(double *restrict A, double *restrict Anew)
{
free(A);
NVTX range
free(Anew); statistics
}
laplace2d.c “calc” region (calcNext function) takes 29.2% Open laplace-par.qdrep
(Parallelised using OpenACC parallel “swap” region (swap function) takes 18.3% of with Nsight System GUI to
directives (pragmas highlighted) view the timeline
total execution time
PARALLEL VS
SEQUENTIAL SPEEDUP
Viewing captured NVTX events
Have a close look at the captured
NVTX events for both serial and
parallel implementations. Parallel
Time spent in “while” loop has
significantly decreased.
Sequential
Achieved speedup: ~47
Parallel
Sequential
PROFILING PARALLEL
CODE
Viewing timeline via Nsight Systems
Contents of the tree-like hierarchy on
the left depend on the project settings
used to collect this report.
If a certain feature has not been
enabled, corresponding rows will not
be shown on the timeline.
In this example, we chose to trace
NVTX and OpenACC while sampling.
Note: Kernel launches are
represented by blue and memory
transfers are displayed in green.
LAB CODE
LAPLACE HEAT TRANSFER
Introduction to lab code - visual Very Hot Room Temp
We will observe a simple simulation
of heat distributing across a metal
plate.
We will apply a consistent heat to
the top of the plate.
Then, we will simulate the heat
distributing across the plate.
LAPLACE HEAT TRANSFER
Introduction to lab code - technical
The lab simulates a very basic
2-dimensional heat transfer problem. A Anew
We have two 2-dimensional arrays,
A and Anew. 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0
The arrays represent a 2-
dimensional, metal plate. Each 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0
element in the array is a double
value that represents temperature. 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0
We will simulate the distribution of 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0
heat until a minimum change value
is achieved, or until we exceed a
maximum number of iterations.
LAPLACE HEAT TRANSFER
Introduction to lab code - technical
We initialize the top row to be a
temperature of 1.0 A Anew
The calcNext function will iterate
through all of the inner elements of 0.0
1.0 0.0
1.0 0.0
1.0 0.0
1.0 0.0
1.0 0.0
1.0 0.0
1.0 0.0
1.0
array A, and update the
corresponding elements in Anew 0.0 0.0 0.0 0.0 0.0 0.0
0.25 0.0
0.25 0.0
We will take the average of the 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0
neighboring cells, and record it in
Anew.
0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0
The swap function will copy the
contents of Anew to A
LAPLACE HEAT TRANSFER
Introduction to lab code
A Anew
1.0 1.0 1.0 1.0 1.0 1.0 1.0 1.0
0.0 0.25 0.25 0.0 0.0 0.25 0.25 0.0
0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0
0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0
The swap function will copy the
contents of Anew to A
KEY CONCEPTS
In this module we discussed…
Compiling sequential and parallel code
CPU profiling for sequential and parallel execution
Specifics of our Laplace Heat Transfer lab code
LAB GOALS
In this lab you will do the following…
Build and run the example code using the NVIDIA’s HPC compiler
Use Nsight Systems to understand where the program spends its time
THANK YOU
