tra added a comment. I do not see the behavior the patch is supposed to fix in CUDA. If I compile a simple program, host-side debugger does not see the `kernel`, sees `__device_stub_kernel` and, if the breakpoint is set on `kernel`, it treats it as a yet-to-be-loaded one and does end up breaking on intry into the kernel on the GPU side.
E.g.: (cuda-gdb) info symbol kernel No symbol "kernel" in current context. (cuda-gdb) info symbol __device_stub__kernel __device_stub__kernel() in section .text (cuda-gdb) b kernel Function "kernel" not defined. Make breakpoint pending on future shared library load? (y or [n]) y Breakpoint 1 (kernel) pending. (cuda-gdb) r Starting program: /usr/local/google/home/tra/work/llvm/build/debug/print [Thread debugging using libthread_db enabled] Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1". Hello from host [New Thread 0x7fffdffff700 (LWP 227347)] [New Thread 0x7fffdf7fe700 (LWP 227348)] [New Thread 0x7fffdeffd700 (LWP 227349)] [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 "print" hit Breakpoint 1, kernel<<<(1,1,1),(1,1,1)>>> () at print.cu:3 3 printf("Hello\n"); Perhaps it's HIP-specific behavior that needs this tweak. For CUDA, I would rather that we continue to emit debug info for the stub (in it's __device_stub form). It is useful for debugging some issues. CHANGES SINCE LAST ACTION https://reviews.llvm.org/D79866/new/ https://reviews.llvm.org/D79866 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits