Kernel stack unwind stops after 5 frames

[RichardABunt]

Hi all,

I am observing stack truncation when debugging a HIP application with ROCgdb. This truncation is shown below:

Thread 6 "deep-stack" hit Breakpoint 1, base_case () at deep-stack.cpp:8
8       return;
(gdb) bt
#0  base_case () at deep.cpp:6
#1  0x00007f847d008454 in deep<0u> () at deep.cpp:17
#2  0x00007f847d0085bc in deep<1u> () at deep.cpp:12
#3  0x00007f847d008668 in deep<2u> () at deep.cpp:12
#4  0x00007f847d008720 in deep<3u> () at deep.cpp:12
#5  0xbebebebebebebebe in ?? ()

Environment

• GNU gdb (rocm-rel-4.0-23) 10.1
• 5.3.18-22-default x86_64 x86_64 x86_64 GNU/Linux
• gfx906
• SUSE Linux Enterprise Server 15 SP2

Reproducer

1. Compile the following:

   
   #include <cstdio>
   #include <cstdlib>
   #include "hip/hip_runtime.h"

device void base_case() { printf("Hello Device\n"); return; }

template device void deep() { deep(); }

template <> device void deep<0> () { base_case(); }

global void hip_deep() { deep<10>(); }

int main() { hipLaunchKernelGGL(HIP_KERNEL_NAME(hip_deep), dim3(1), dim3(1), 0, 0); hipDeviceSynchronize(); return EXIT_SUCCESS; }

with:

hipcc -ggdb -O0 -o deep deep.cpp

2. Start under rocgdb.
3. Place a break point on the return in base_case().
4. Run to the breakpoint.
5. bt.
6. Observe the unwind stop.

Many thanks,

Rich.