yaxunl added a comment. In D79866#2034460 <https://reviews.llvm.org/D79866#2034460>, @tra wrote:
> 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. can you try set bp by using file name and line number on the kernel? 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