Tools for GPU Debugging and Profiling
June 02, 2022
Rory Kelly
Consulting Services
Overview
Debugging Tools
Profiling Tools
GPU Debugging and Profiling
Types of Bugs
The type of bug you have will have an impact on how you approach debugging it.
Think about what the nature of your bug may be telling you to guide your debugging approach.
GPU Debugging and Profiling
Before using a Debugger
There are a few easy to use debugging tools you may want to try before resorting to a debugger
GPU Debugging and Profiling
Hail printf(), Long May It Reign
__global__ void unsafe_inc(int *a_d){
...
if (blockIdx.x == 99 && threadIdx.x >= 16 && threadIdx.x < 32){
printf("Block x %d, Thread x: %d Value of a: %d\n",
blockIdx.x, threadIdx.x, a_d[idx]);
}
...
}
GPU Debugging and Profiling
Environment Variables for OpenACC
NV_ACC_NOTIFY=<1 for kernel launches, 2 for data xfer, 3 for both>
...
upload CUDA data file=/glade/work/rory/GPU-tut/c-openacc-prof/miniWeather_mpi_openacc.cpp function=_Z10reductionsRdS_ line=869 device=0 threadid=1 variable=te_loc bytes=8
launch CUDA kernel file=/glade/work/rory/GPU-tut/c-openacc-prof/miniWeather_mpi_openacc.cpp function=_Z10reductionsRdS_ line=869 device=0 threadid=1 num_gangs=625 num_workers=1 vector_length=128 grid=625 block=128 shared memory=2048
...
NV_ACC_DEBUG=1
GPU Debugging and Profiling
An Example CUDA Bug - Array Bounds
// buggy kernel will write one element off the end of c_d
__global__ void boundsBugAdd (int *a_d, int *b_d, int *c_d)
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
c_d[x+1] = a_d[x] + b_d[x];
}
called as:
arraySize=64;
boundsBugAdd <<<ceil((float) arraySize/32),32>>> (a_d, b_d, c_d);
Expect that last thread in last block will write out of bounds
GPU Debugging and Profiling
An Example CUDA Bug - Array Bounds
With arrays initialized as:
for (i=0; i < 64; i++){
a[i] = i+1;
b[i] = -(i+1);
c[i] = -1;
}
> ./boundsBug.exe
Result:
-1 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 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
As expected: first array element not updated, presumably an out of bounds write on the GPU, and no crash.
GPU Debugging and Profiling
An Example CUDA Bug - Array Bounds
Use compute sanitizer
> compute-sanitizer --tool memcheck ./boundsBug.exe
========= COMPUTE-SANITIZER
========= Invalid __global__ write of size 4 bytes
========= at 0x560 in boundsBug.cu:10:boundsBugAdd(int *, int *, int *)
========= by thread (31,0,0) in block (1,0,0)
========= Address 0x2b6beda00500 is out of bounds
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame: [0x21740c]
========= in /lib64/libcuda.so
========= Host Frame: [0x87eb]
========= in /glade/work/rory/GPU-tut/boundsBug.cuda/./boundsBug.exe
...
GPU Debugging and Profiling
An Example CUDA Bug - Race Condition
// buggy kernel has a shared memory race condition
__global__ void unsafe_inc(int *a_d){
__shared__ int s;
s = *a_d;
s += 1;
*a_d = s;
}
called as:
unsafe_inc<<<1000,1000>>>(a_d);
Expect that last a_d will end up with value < 1e6, due to the race condition
GPU Debugging and Profiling
An Example CUDA Bug - Race Condition
> ./race-cond.exe
GPU Time elapsed: 0.000082 seconds
a = 16
> ./race-cond.exe
GPU Time elapsed: 0.000077 seconds
a = 12
Use compute sanitizer
> compute-sanitizer --tool=racecheck ./race-cond.exe
========= COMPUTE-SANITIZER
========= ERROR: Race reported between Write access at 0x270 in race-cond.cu:18:unsafe_inc(int *)
========= and Write access at 0x1c0 in race-cond.cu:17:unsafe_inc(int *) [1 hazards]
========= and Read access at 0x210 in race-cond.cu:18:unsafe_inc(int *) [6732 hazards]
========= and Write access at 0x270 in race-cond.cu:18:unsafe_inc(int *) [987 hazards]
========= and Read access at 0x2c0 in race-cond.cu:19:unsafe_inc(int *) [17032 hazards]
...
GPU Debugging and Profiling
Debuggers - when you must
If you haven’t been able to find a bug with simpler methods, it may by time to use a debugger.
There are a few options available, and we’ll talk about two today:
Neither is perfect, but can provide additional insight into your code
Both seem to work better with CUDA codes than with OpenACC generated kernels.
GPU Debugging and Profiling
Debuggers - compiling for debugging
CUDA Flags:
-O3 -g -G or
-O0 -g -G
OpenACC
-O3 -g -acc=gpu -gpu=cc70,debug,nordc or
-O0 -g -acc=gpu -gpu=cc70,debug,nordc
may be interesting to keep generated kernels with
-gpu=...,keepgpu
GPU Debugging and Profiling
CUDA-GDB
The same gdb you are familiar with, including all the same CPU-side capabilities, but extended to work on NVIDIA GPUs and CUDA code.
Fairly feature rich, but usefulness of the tools depends on the nature of your bug. More useful for CUDA, has some limitations for OpenACC.
Can work through the CLI or from within an IDE.
Uses /tmp by default, but respects $TMPDIR environment variable, which you can point to /glade/scratch/$USER/tmp or similar
https://docs.nvidia.com/cuda/cuda-gdb/index.html
GPU Debugging and Profiling
CUDA-GDB
Can’t cover all the features today, but a quick way to get info on CUDA specific functionality is to start cuda-gdb, and then type
cuda <tab> ← commands for showing and selecting current focus
block device grid kernel lane sm thread warp
These are the software/hardware views of the currently executing focus. You can also use these commands to switch the current focus.
help <command> (e.g. help cuda thread) for more info.
GPU Debugging and Profiling
CUDA-GDB
set cuda <tab> ← commands to control debug behavior
api_failures disassemble_per notify
break_on_launch gpu_busy_check ptx_cache
coalescing hide_internal_frame single_stepping_optimizations
collect_stats kernel_events software_preemption
context_events kernel_events_depth stop_signal
launch_blocking thread_selection disassemble_from memcheck value_extrapolation
device_resume_on_cpu_dynamic_function_call
help <command> (e.g. help set cuda break_on_launch) for more info.
GPU Debugging and Profiling
CUDA-GDB - memcheck
> cuda-gdb ./boundsBug.exe
(cuda-gdb) set cuda memcheck on
(cuda-gdb) run
Thread 1 "boundsBug.exe" received signal CUDA_EXCEPTION_1, Lane Illegal Address.
[Switching focus to CUDA kernel 0, grid 1, block (1,0,0), thread (31,0,0), device 0, sm 2, warp 0, lane 31]
0x0000000001132e60 in boundsBugAdd<<<(2,1,1),(32,1,1)>>> (a_d=0x2aab04400000, b_d=0x2aab04400200, c_d=0x2aab04400400) at boundsBug.cu:10
10 c_d[x+1] = a_d[x] + b_d[x];
(cuda-gdb) list
5
6 // buggy kernel will write one element off the end of c_d
7 __global__ void boundsBugAdd (int *a_d, int *b_d, int *c_d)
8 {
9 int x = blockIdx.x * blockDim.x + threadIdx.x;
10 c_d[x+1] = a_d[x] + b_d[x];
11 }
GPU Debugging and Profiling
CUDA-GDB - breakpoints
(cuda-gdb) list 6,11
6 // buggy kernel will write one element off the end of c_d
7 __global__ void boundsBugAdd (int *a_d, int *b_d, int *c_d)
8 {
9 int x = blockIdx.x * blockDim.x + threadIdx.x;
10 c_d[x+1] = a_d[x] + b_d[x];
11 }
So, in this specific case, these commands would be equivalent:
(cuda-gdb) break 9
(cuda-gdb) break boundsBug:9
(cuda-gdb) break boundsBugAdd
(cuda-gdb) set cuda break_on_launch application
GPU Debugging and Profiling
CUDA-GDB - focus
Reading symbols from boundsBug.exe...
(cuda-gdb) set cuda break_on_launch application
(cuda-gdb) run
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]
boundsBugAdd<<<(2,1,1),(32,1,1)>>> (a_d=0x2aab03a00000, b_d=0x2aab03a00200, c_d=0x2aab03a00400) at boundsBug.cu:9
9 int x = blockIdx.x * blockDim.x + threadIdx.x;
(cuda-gdb) cuda block thread
block (0,0,0), thread (0,0,0)
GPU Debugging and Profiling
CUDA-GDB - focus
(cuda-gdb) step
10 c_d[x+1] = a_d[x] + b_d[x];
(cuda-gdb) p x
$2 = 0
(cuda-gdb) cuda block(1,0,0) thread(31,0,0)
[Switching focus to CUDA kernel 0, grid 1, block (1,0,0), thread (31,0,0), device 0, sm 2, warp 0, lane 31]
9 int x = blockIdx.x * blockDim.x + threadIdx.x;
(cuda-gdb) step
10 c_d[x+1] = a_d[x] + b_d[x];
(cuda-gdb) p x
$3 = 63
GPU Debugging and Profiling
CUDA-GDB - examining local state
> cuda-gdb ./test_voigt
(cuda-gdb) list voigt.cu:66,94
66 Z1_real = A6 * damping + A5;
67 Z1_imag = A6 * -V;
68 Z2_real = Z1_real * damping - Z1_imag * -V + A4;
69 Z2_imag = Z1_real * -V + Z1_imag * damping;
70 Z3_real = Z2_real * damping - Z2_imag * -V + A3;
71 Z3_imag = Z2_real * -V + Z2_imag * damping;
72 Z4_real = Z3_real * damping - Z3_imag * -V + A2;
73 Z4_imag = Z3_real * -V + Z3_imag * damping;
74 Z5_real = Z4_real * damping - Z4_imag * -V + A1;
75 Z5_imag = Z4_real * -V + Z4_imag * damping;
76 Z6_real = Z5_real * damping - Z5_imag * -V + A0;
77 Z6_imag = Z5_real * -V + Z5_imag * damping;
78 ZZ1_real = damping + B6;
79 ZZ1_imag = -V;
(list continued)
80 ZZ2_real = ZZ1_real * damping - ZZ1_imag * -V + B5;
81 ZZ2_imag = ZZ1_real * -V + ZZ1_imag * damping;
82 ZZ3_real = ZZ2_real * damping - ZZ2_imag * -V + B4;
83 ZZ3_imag = ZZ2_real * -V + ZZ2_imag * damping;
84 ZZ4_real = ZZ3_real * damping - ZZ3_imag * -V + B3;
85 ZZ4_imag = ZZ3_real * -V + ZZ3_imag * damping;
86 ZZ5_real = ZZ4_real * damping - ZZ4_imag * -V + B2;
87 ZZ5_imag = ZZ4_real * -V + ZZ4_imag * damping;
88 ZZ6_real = ZZ5_real * damping - ZZ5_imag * -V + B1;
89 ZZ6_imag = ZZ5_real * -V + ZZ5_imag * damping;
90 ZZ7_real = ZZ6_real * damping - ZZ6_imag * -V + B0;
91 ZZ7_imag = ZZ6_real * -V + ZZ6_imag * damping;
92 division_factor = 1.0f / (ZZ7_real * ZZ7_real + ZZ7_imag
* ZZ7_imag);
93 ZZZ_real = (Z6_real * ZZ7_real + Z6_imag * ZZ7_imag)
* division_factor;
94 voigt_value[idx] = ZZZ_real;
GPU Debugging and Profiling
CUDA-GDB - examining local state
(cuda-gdb) break voigt.cu:94
Breakpoint 1 at 0x405c0a: file voigt.cu, line 95.
(cuda-gdb) run
Thread 1 "test_voigt" hit Breakpoint 1, my_voigt<<<(8192,256,1),(32,1,1)>>> (damp_arr=0x2aab2e000000, offs_arr=0x2aab3e000000,
voigt_value=0x2aab4e000000) at voigt.cu:94
94 voigt_value[idx] = ZZZ_real;
(cuda-gdb) step
95 }
At this point there are a number of commands to get full info on available local state
GPU Debugging and Profiling
CUDA-GDB - examining local state
(cuda-gdb) backtrace full
#0 my_voigt<<<(8192,256,1),(32,1,1)>>> (damp_arr=0x2aab2e000000, offs_arr=0x2aab3e000000, voigt_value=0x2aab4e000000) at voigt.cu:95
Z1_real = 11.5567265
ZZ1_real = 20.4837646
ZZ3_real = -1291.85046
Z6_imag = 6668691.5
ZZ6_imag = 11803426
division_factor = 2.71227659e-17
damping = 10.0039062
Z1_imag = -5.64410019
Z2_real = 89.3294983
Z4_real = -33328.1367
...
Z5_imag = 145220.266
Z6_real = -3763029
ZZ5_imag = 255083.594
ZZ6_real = -6699409.5
ZZ7_imag = 185100640
idx = 0
ivsigno = 1
GPU Debugging and Profiling
CUDA-GDB - examining local state
(cuda-gdb) print V
$4 = 10.0039062
(cuda-gdb) print offset
$5 = 10.0039062
(cuda-gdb) print V - offset
$6 = 0
(cuda-gdb) print V / offset
$7 = 1
GPU Debugging and Profiling
CUDA-GDB - watchpoints
> cuda-gdb ./test_voigt
(cuda-gdb) run
Thread 1 "test_voigt" hit Breakpoint 1, my_voigt<<<(8192,256,1),(32,1,1)>>> (damp_arr=0x2aab2e000000, offs_arr=0x2aab3e000000,
voigt_value=0x2aab4e000000) at voigt.cu:66
(cuda-gdb) watch Z1_imag if Z1_imag < 0.0
Watchpoint 2: Z1_imag
(cuda-gdb) continue
Continuing.
Thread 1 "test_voigt" hit Watchpoint 2: Z1_imag
Old value = 0
New value = -5.64410019
my_voigt<<<(8192,256,1),(32,1,1)>>> (damp_arr=0x2aab2e000000, offs_arr=0x2aab3e000000, voigt_value=0x2aab4e000000) at voigt.cu:68
68 Z2_real = Z1_real * damping - Z1_imag * -V + A4;
GPU Debugging and Profiling
CUDA-GDB - Sometimes problematic with OpenACC
> cuda-gdb matrix_mult.exe
(cuda-gdb) set cuda break_on_launch application
(cuda-gdb) run
Starting program: /glade/work/rory/GPU-tut/f90-mmul/matrix_mult.exe[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]
cuda-gdb/10.1/gdb/cuda/cuda-regmap.c:703: internal-error: regmap_st* regmap_table_search(objfile*, const char*, const char*, uint64_t): Assertion `func_name' failed.
A problem internal to GDB has been detected,
further debugging may prove unreliable.
Quit this debugging session? (y or n)
y
Create a core file of GDB? (y or n)
n
GPU Debugging and Profiling
CUDA-GDB - Sometime problematic with OpenACC
GPU Debugging and Profiling
ARM Forge/DDT
In addition to being a scalable MPI and OpenMP debug tool for CPU codes, DDT is also able to debug on NVIDIA GPUs, including both CUDA and OpenACC codes.
Works a bit better for OpenACC codes (esp Fortran) vs cuda-gdb.
Many similar capabilities to cuda-gdb. The primary interface is a GUI, which you either like or don’t, but as GUI tools in HPC go, it’s pretty good.
We have full documentation on getting it set up at NCAR
GPU Debugging and Profiling
ARM Forge/DDT - Revisiting Fortran OpenACC
> forge --connect ./matrix_mult.exe
GPU Debugging and Profiling
GPU Debugging and Profiling
GPU Debugging and Profiling
GPU Debugging and Profiling
Also has similar memory debugging capabilities
Debugger Debrief - General Advice
GPU Debugging and Profiling
A Brief Word on Profiling
Both NVIDIA and ARM Forge toolchains include profiling and optimization tools as well.
Future GPU Workshop session will dive deeper on the NVIDIA NSight tool, particularly the GUI version, but the CLI version is also quite useful.
Using ARM Forge/MAP shares launch method and interface with DDT, so it should be a small leap for existing DDT users.
GPU Debugging and Profiling
NSight CLI
> nv-nsight-cu-cli --launch-count=100 --launch-skip=1 ./mw_openacc
==PROF== Connected to process 136370 (/glade/work/rory/GPU-tut/c-openacc-prof/build/mw_openacc)
==PROF== Profiling "_Z23reductions_869_gpu__redRdS_" - 1 of 100: 0%....50%....100% - 19 passes
==PROF== Profiling "_Z25set_halo_values_x_408_gpuPd" - 2 of 100: 0%....50%....100% - 19 passes
==PROF== Profiling "_Z25set_halo_values_x_431_gpuPd" - 3 of 100: 0%....50%....100% - 19 passes
...
GPU Debugging and Profiling
NSight CLI -- Example Report
_Z28compute_tendencies_z_379_gpuPdS_S_d, 2022-Jun-01 23:52:50, Context 1, Stream 14
Section: GPU Speed Of Light Throughput
----------------------------------------------------------------- --------------- -------------------------
DRAM Frequency cycle/usecond 608.42
SM Frequency cycle/usecond 880.03
Elapsed Cycles cycle 10,633
Memory [%] % 44.35
DRAM Throughput % 42.92
Duration usecond 12.06
L1/TEX Cache Throughput % 38.40
L2 Cache Throughput % 44.35
SM Active Cycles cycle 8,132.10
Compute (SM) [%] % 36.98
----------------------------------------------------------------- --------------- -------------------------
WRN This kernel exhibits low compute throughput and memory bandwidth utilization relative to the peak
performance of this device. Achieved compute throughput and/or memory bandwidth below 60.0% of peak typically
indicate latency issues. Look at Scheduler Statistics and Warp State Statistics for potential reasons.
GPU Debugging and Profiling
NSight CLI -- Example Report
Section: Launch Statistics
----------------------------------------------------------------- --------------- -------------------------
Block Size 128
Function Cache Configuration cudaFuncCachePreferNone
Grid Size 2,500
Registers Per Thread register/thread 44
Shared Memory Configuration Size byte 0
Driver Shared Memory Per Block byte/block 0
Dynamic Shared Memory Per Block byte/block 0
Static Shared Memory Per Block byte/block 0
Threads thread 320,000
Waves Per SM 3.12
----------------------------------------------------------------- --------------- -------------------------
WRN A wave of thread blocks is defined as the maximum number of blocks that can be executed in parallel on
the target GPU. The number of blocks in a wave depends on the number of multiprocessors and the theoretical
occupancy of the kernel. This kernel launch results in 3 full waves and a partial wave of 100 thread blocks.
Under the assumption of a uniform execution duration of all thread blocks, the partial wave may account for
up to 25.0% of the total kernel runtime with a lower occupancy of 27.8%. Try launching a grid with no partial
wave. The overall impact of this tail effect also lessens with the number of full waves executed for a grid.
GPU Debugging and Profiling
NSight CLI -- Example Report
Section: Occupancy
----------------------------------------------------------------- --------------- -------------------------
Block Limit SM block 32
Block Limit Registers block 10
Block Limit Shared Mem block 32
Block Limit Warps block 16
Theoretical Active Warps per SM warp 40
Theoretical Occupancy % 62.50
Achieved Occupancy % 45.15
Achieved Active Warps Per SM warp 28.90
----------------------------------------------------------------- --------------- -------------------------
WRN This kernel's theoretical occupancy (62.5%) is limited by the number of required registers The difference
between calculated theoretical (62.5%) and measured achieved occupancy (45.2%) can be the result of warp
scheduling overheads or workload imbalances during the kernel execution. Load imbalances can occur
between warps within a block as well as across blocks of the same kernel.
GPU Debugging and Profiling
ARM Forge/MAP
Uses reverse connect to launch the same way as DDT
> module load arm-forge/22.0.2
> map --connect ./mw_openacc
GPU Debugging and Profiling
Click to add footer
Profiling Summary
Just a quick overview of profiling tools. As mentioned, future session will focus on the NVIDIA profiling tools in more depth.
If you want more help with either of these tools, feel free to reach out to CSG and we can assist you. We will likely have future vendor provided trainings on these and other tools, and such trainings will be announced in many channels, including to this group.
Questions?
Now? ←→ At the end? ←→ Offline?
Thanks
GPU Debugging and Profiling