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

[Artem Belevich via Phabricator via cfe-commits]

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