Debugging with LLDB (Mac)

Stepping through your kernel

Currently, stepping through your kernel is a potentially confusing experience, since you may switch between cuda threads just by stepping a single line. You'll see such a thing happen, for example, on a call to __syncthreads() or a write to shared memory. An lldb plugin is in the works that will make this a little easier, but for now, you can use thread-specific breakpoints to get around this issue. If there's a spot in your kernel where a context switch happens, but you would prefer to keep following the current thread, just create a thread-specific breakpoint (see the following "Breaking on a Specific Thread" section) past the line causing the context switch and run to it via the "cont" command.

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:

b -l 20 -c ist(0,1,2)

You can also combine the two, such as:

b -l 20 -c "ist(0,1,2) && isb(1,0,0)"