Unknown CFA rule. A problem internal to GDB has been detected

[zjin-lcf]

While working with DPCPP (https://github.com/intel/llvm) compiling SYCL kernels targeting ROCm (gfx908) I came across the gdb error. It is true that there are some issues with compiling the SYCL kernels. I am not familiar with gdb internals. Does the gdb error provide some hint/clue to the compiler issue ?

https://github.com/zjin-lcf/oneAPI-DirectProgramming/tree/master/babelstream-sycl

rocgdb ./main
Precision: float
Array size: 134.2 MB (=0.1 GB)
Total size: 402.7 MB (=0.4 GB)
Memory access fault by GPU node-2 (Agent handle: 0x532080) on address 0x7ffeb0019000. Reason: Page not present or supervisor privilege.

Thread 6 "main" received signal SIGSEGV, Segmentation fault.
[Switching to AMDGPU Thread 2:3:1:1 (130944,0,0)/0]
copy<float>(cl::sycl::queue&, cl::sycl::buffer<float, 1, cl::sycl::detail::aligned_allocator<char>, void>&, cl::sycl::buffer<float, 1, cl::sycl::detail::aligned_allocator<char>, void>&)::{lambda(cl::sycl::handler&)#1}::operator()(cl::sycl::handler&) const::{lambda(cl::sycl::nd_item<1>)#1}::operator()(cl::sycl::nd_item<1>) const (
/home/release/git/aomp13/ROCgdb/gdb/dwarf2/frame.c:1029: internal-error: Unknown CFA rule.
A problem internal to GDB has been detected,
further debugging may prove unreliable.

[t-tye]

Thank you for the report. We actually recently encountered the same issue and have a fix prepared in the compiler to address the issue of the CFI being incorrect. In addition, I believe gdb was not gracefully handling this and that issue was also being fixed.

Will report back once those fixes are upstreamed.

[t-tye]

ROCgdb requires correct CFI DWARF information to be present in the GPU code objects. The upstream LLVM compiler does not contain this code.

Gdb does not handle unwinding well in the absence of CFI information which results in the internal error. That needs to be fixed in upstream gdb but is non-trivial due to the current design so a fix is not expected soon.

The ROCm LLVM based AMD GPU compiler is open source in github so it may be possible to use that to generate correct CFI information. There were some recent fixes which may not be available in the github until the next release is published.

[zjin-lcf]

Thank you for explaining the technical issue and updates.