[PATCH] D79866: [CUDA][HIP] Do not emit debug info for stub function

[Yaxun Liu via Phabricator via cfe-commits]

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