Can I use conditional breakpoint in cuda-gdb?

[202476410arsmart]

Cuda-gdb 3.0 doesn’t support conditional breakpoint.

https://forums.developer.nvidia.com/t/cuda-gdb-and-conditional-breakpoints/16060/3

I want to see when “threadIdx.x == 0 && threadIdx.y == 0 && blockIdx.x == 0 && blockIdx.y ==0” Not sure whether can I do it?

[agontarek]

Hi - thanks for reaching out!

We can use conditional breakpoints to accomplish this. I believe in the referenced post one of the issues was that the application was not compiled with the -G flag. For a full debugging experience, we need to pass both -g and -G flags on the compliation line.

We can set a conditional breakpoint by adding “if” followed by the conditional expression when using the breakpoint command. We can also use the condition command to transform an existing breakpoint into a conditional breakpoint.

For example:

NVIDIA (R) CUDA Debugger
CUDA Toolkit 12.3 release
Portions Copyright (C) 2007-2023 NVIDIA Corporation
GNU gdb (GDB) 12.1
Copyright (C) 2022 Free Software Foundation, Inc.
License GPLv3+: GNU GPL version 3 or later <http://gnu.org/licenses/gpl.html>
This is free software: you are free to change and redistribute it.
There is NO WARRANTY, to the extent permitted by law.
Type "show copying" and "show warranty" for details.
This GDB was configured as "x86_64-pc-linux-gnu".
Type "show configuration" for configuration details.
For bug reporting instructions, please see:
<https://www.gnu.org/software/gdb/bugs/>.
Find the GDB manual and other documentation resources online at:
    <http://www.gnu.org/software/gdb/documentation/>.

For help, type "help".
Type "apropos word" to search for commands related to "word"...
Using python library libpython3.9.so.1
Reading symbols from a.out...
(cuda-gdb) break C if threadIdx.x == 0
No symbol "threadIdx" in current context.
Make conditional breakpoint pending on future shared library load? (y or [n]) y
Breakpoint 1 (C if threadIdx.x == 0) pending.
(cuda-gdb) run
Starting program: /home/agontarek/Sandbox/a.out 
Error in re-setting breakpoint 1: No symbol "threadIdx" in current context.
Error in re-setting breakpoint 1: No symbol "threadIdx" in current context.
Error in re-setting breakpoint 1: No symbol "threadIdx" in current context.
Error in re-setting breakpoint 1: No symbol "threadIdx" in current context.
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
Error in re-setting breakpoint 1: No symbol "threadIdx" in current context.
Error in re-setting breakpoint 1: No symbol "threadIdx" in current context.
Error in re-setting breakpoint 1: No symbol "threadIdx" in current context.
[New Thread 0x7ffff1854000 (LWP 384203)]
Error in re-setting breakpoint 1: No symbol "threadIdx" in current context.
Error in re-setting breakpoint 1: No symbol "threadIdx" in current context.
[New Thread 0x7fffebfff000 (LWP 384204)]
[Detaching after fork from child process 384205]
[New Thread 0x7fffea7dd000 (LWP 384217)]
[New Thread 0x7fffe8df4000 (LWP 384218)]
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]

Thread 1 "a.out" hit Breakpoint 1, C () at test_step.cu:5
5	   printf("C()\n");
(cuda-gdb) print threadIdx
$1 = {x = 0, y = 0, z = 0}
(cuda-gdb)

In the above, I set a pending conditional breakpoint on a device function. The error messages here can be safely ignored. The messages are informing us that the pending breakpoint was unable to be reset for the host side function corresponding to C. There is no threadIdx variable exposed in host code. We only have access to that convenience variable when executing on the device. Once the kernel has been loaded and execution proceeds into the function, the conditional breakpoint will be resolved and we hit the breakpoint when our expression evaluates to true.

One caveat to be aware of today is that conditional breakpoints are evaluated on the host. This can be an expensive operation. If you expect a breakpoint to hit hundreds of thousands of times before the expression evaluates to true it may take some time.