Debugging ESIMD

[DDEle]

Is your feature request related to a problem? Please describe
I'm trying to gdb-oneapi an esimd-based project on A770 GPU. I notice that I cannot get value for most of variables in the kernel.

Taking the vadd_usm as an example:

• Installing OneAPI Base Toolkit 2024.1from internal source

  $ icpx --version
  Intel(R) oneAPI DPC++/C++ Compiler 2024.1.0 (2024.1.0.20240117)
  Target: x86_64-unknown-linux-gnu
  Thread model: posix
  InstalledDir: /opt/intel/oneapi/compiler/2024.1/bin/compiler
  Configuration file: /opt/intel/oneapi/compiler/2024.1/bin/compiler/../icpx.cfg

• Setup environment via source /opt/intel/oneapi/setvars.sh
• Get the vadd_usm example

  mkdir test_gdb && cd test_gdb
  wget https://raw.githubusercontent.com/intel/llvm/sycl/sycl/test-e2e/ESIMD/vadd_usm.cpp
  wget https://raw.githubusercontent.com/intel/llvm/sycl/sycl/test-e2e/ESIMD/esimd_test_utils.hpp
  icpx -fsycl vadd_usm.cpp -O0 -g

• Run gdb-oneapi (and get almost nothing informative)

  ONEAPI_DEVICE_SELECTOR=level_zero1:* ZET_ENABLE_PROGRAM_DEBUGGING=1 gdb-oneapi --args a.out
  ...
  (gdb) b 54
  Breakpoint 1 at 0x406c86: file vadd_usm.cpp, line 55.
  (gdb) r
  ...
  [Switching to thread 2.449:0 (ZE 0.3.8.0 lane 0)]
  Thread 2.449 hit Breakpoint 1.2, with SIMD lanes [0-31], main::{lambda(sycl::_V1::handler&)#1}::operator()(sycl::_V1::handler&) const::{lambda(sycl::_V1::nd_item<1>)#1}::operator()(sycl::_V1::nd_item<1>) const (this=<optimized out>, ndi=...) at vadd_usm.cpp:55
  55                                           vb.copy_from(B + i * VL);
  (gdb) info local
  i = <optimized out>
  va = <optimized out>
  vb = <optimized out>
  vc = <optimized out>
  (gdb) info args
  this = <optimized out>
  ndi = <optimized out>
  

Describe the solution you would like
Enabling gdb-oneapi for esimd programs

Describe alternatives you have considered
N/A

Additional context
I installed the gpu driver based on instructions here, while manually specified the kernel version to be 5.19. After that, I set up the GPU debugger. So that,

• ONEAPI_DEVICE_SELECTOR=level_zero1:* ZET_ENABLE_PROGRAM_DEBUGGING=1 sycl-ls:
  [ext_oneapi_level_zero:gpu:0] Intel(R) Level-Zero, Intel(R) Arc(TM) A770 Graphics 1.3 [1.3.27191]
• uname -r:
  5.19.0-50-generic
• apt list --installed | grep i915:
  intel-i915-dkms/unknown,unknown,now 1.23.9.11.231003.15+i19-1 all [installed]

[sarnex]

Unfortunately IGC's vector backend doesn't support debugging, so we can't do much better at the moment.

FYI: @intel/dpcpp-esimd-reviewers