Debugging with GDB (Linux & Windows)

GDB Plugin

cuda-edu provides a plugin that helps work around some of the problems GDB has with stepping through a kernel. The plugin will be loaded if you use the -g flag of the ./run script, e.g.:

./run -g 0

If you don't want to use the plugin, use the -p flag:

./run -gp 0

If the plugin has been successfully loaded, then the initialization of your GDB session will look like the following:

---
--- Loading CUDA-EDU Plugin
---
(cuda-edu) 

That is, you'll see a (cuda-edu) prompt instead of the normal (gdb).

The plugin may not play nicely with other tools. It relies on creating a number of breakpoints that are "invisible" and "silent." If other tools don't respect those properties of breakpoints, then things could get messy.

n command

The plugin's n command is used for stepping through a kernel -- it can also step through host code. The n command will step to the next line. If the line results in a cuda thread context switch (e.g. __syncthreads()__ or writing to shared memory), then your debugging session will switch to the next thread.

tn command

The plugin's tn command is similar to the n command, but it will not switch to another thread on a context switch. Instead, on a context switch, the tn command will wait for the other threads to execute until a context switch back to the thread you're debugging.

next command

You can still use GDB's standard next command by typing its full name. You probably shouldn't use it in a kernel, though, because if you use it to step over a line that causes a context switch, then your debugger will fail.

Breaking on a Specific Thread

cuda-edu provides a couple convenience functions, referred to as index predicates, that help in creating breakpoint conditions:

• ist(unsigned x, unsigned y, unsigned z) - Returns true if the current threadIdx is {x,y,z}
• isb(unsigned x, unsigned y, unsigned z) - Returns true if the current blockIdx is {x,y,z}

These two predicates can then be used in your debugger to break on a specific thread. For example, to break at mp.cu:20 for thread {0,1,2} of the current block issue the following:

break 20 if ist(0,1,2)

You can also combine the two, such as:

break 20 if ist(0,1,2) && isb(1,0,0)