Environment variables

# Inserts a read before every write, making page faults detectable while
# the thread is still alive.  Without this, the fault may go undetected or
# be reported at the wrong location.  Enabled by default for applications
# compiled without optimization.
# Note that it may impact the observable performance.
export IGC_VISAOptions="-enableBarrierWA"
# Disable caching to ensure environment variables take effect.
export NEO_CACHE_PERSISTENT=0
export SYCL_CACHE_PERSISTENT=0

Examples

Memory access requests of a GPU thread are asynchronous. While a request is processed, the thread that triggered the request may continue execution, up to the point where the requested memory would be used or the thread would exit, which ever comes first. That means, when the exception is triggered and the thread is stopped, thread IP may have proceeded further from the instruction that triggered the faulting request.

In the following example, the kernel attempts to read through a nullptr:

(gdb) list 37
32
33            cgh.single_task<> ([=] ()
34              {
35                int *src = nullptr;
36                int num = *src;
37                numbers[0] = num;
38              });
39          });
40        deviceQueue.wait ();
41
(gdb) run

Thread 2.97 received signal SIGSEGV, Segmentation fault
Warning: The location reported for the signal may be inaccurate.
[Switching to thread 2.97:0 (ZE 0.1.4.0 lane 0)]
main::{lambda(sycl::_V1::handler&)#1}::operator()(sycl::_V1::handler&) const::{lambda()#1}:: ...
    at kernel-pagefault-read.cpp:37
37                numbers[0] = num;
(gdb)

Note the warning “The location reported for the signal may be inaccurate”. The faulting read is requested at sourceline 36, but the value is used only at sourceline 37. Thus, it depends on timing whether the thread reaches the line 37 before the thread is stopped and the SIGSEGV signal is triggered.

In the next example, kernel attempts to write through a nullptr, and then exits without further accessing that variable:

(gdb) list 37
32
33            cgh.single_task<> ([=] ()
34              {
35                int *p = nullptr;
36                int num = numbers[0];
37                *p = num;
38              });
39          });
40        deviceQueue.wait ();
41
(gdb) run

Thread 2.97 received signal SIGSEGV, Segmentation fault
Warning: The location reported for the signal may be inaccurate.
[Switching to thread 2.97:0 (ZE 0.1.4.0 lane 0)]
0x00008000ffe80980 in _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_ (_arg_numbers=sycl: ...
    at /opt/intel/oneapi/compiler/2024.2/bin/compiler/../../include/sycl/handler.hpp:1535
1535      }
(gdb)

This time the thread has already returned from its main function and is preparing to exit, before the thread is stopped.

Identify the instruction triggering the faulting request

In case of a failing read request, the instruction triggering the faulting request can be identifyed by setting the following environment variable:

# Add a sync after every read request
export IGC_EnableForceDebugSWSB=1

With this setting, the thread stops immediatelly after executing the failing request.

(gdb) run

Thread 2.97 received signal SIGSEGV, Segmentation fault
Warning: The location reported for the signal may be inaccurate.
[Switching to thread 2.97:0 (ZE 0.1.4.0 lane 0)]
0x00008000ffe58c00 in main::{lambda(sycl::_V1::handler&)#1}::operator()(sycl::_V1::handler&) ...
    at kernel-pagefault-read.cpp:36
36                int num = *src;
(gdb)

In this example, thread stops at 0x8000ffe58c00. This is the address of the instruction following immediately the one that triggers the exception.

However, write faults may be detected much later. In some cases, the fault is detected after the thread returned from the kernel function and is about to exit, and finding the instruction that caused the page fault may require stepping from an earlier breakpoint.

All pending read and write requested of a thread are completed before the thread is stopped. Thus, we can use stepping to find the sourceline and the exact instruction that triggered the faulting request. If we have multiple threads, we should also set ‘scheduler-locking’ to avoid switching threads.

(gdb) break 37
Breakpoint 1 at 0x40625b: file kernel-pagefault-write.cpp, line 37.
(gdb) run

Thread 2.97 hit Breakpoint 1.2, with SIMD lane 0, main::{lambda(sycl::_V1::handler&)#1}::ope ...
    at kernel-pagefault-write.cpp:37
37                *p = num;             /* pagefault-line */
(gdb) set scheduler-locking step
(gdb) stepi
0x00008000ffe18e90      37                *p = num;
(gdb) stepi
0x00008000ffe18ea0      37                *p = num;
(gdb) stepi
0x00008000ffe18eb0      37                *p = num;
(gdb) stepi

Thread 2.97 received signal SIGSEGV, Segmentation fault
Warning: The location reported for the signal may be inaccurate.
main::{lambda(sycl::_V1::handler&)#1}::operator()(sycl::_V1::handler&) const::{lambda()#1}:: ...
    at kernel-pagefault-write.cpp:38
38              });
(gdb)

In the above example, we hit a breakpoint at sourceline 37, continue stepping by instruction until the exception occurs, and find that it was triggered by the instruction at 0x8000ffe18eb0.

Exit signal without a stopped thread

What should you do if the inferior exited with signal SIGSEGV? GDB removes the inferior because the underlying process that the inferior represents is already gone.

Program terminated with signal SIGSEGV, Segmentation fault.
The program no longer exists.
Detached from device [0000:3a:00.0]
[Inferior 2 (device [0000:3a:00.0]) detached]
intelgt: inferior 2 (gdbserver-ze) has been removed.
(gdb)

There are two alternative scenarios that could result in this behavior.

In the first scenario, if Barrier WA was not enabled when a thread requested a write, this would result in a page fault. Due to asynchronous memory accesses, the thread may have exited before the exception was triggered. In such a case, the debugger can only show the signal, but there is no crashing thread to stop. With the Barrier WA, the thread waits at a memory fence inserted by the compiler until all its memory requests are completed. This enables the debugger to stop the thread that triggered the exception. Use environment variables to enable the Barrier WA.

In the second scenario, if the Barrier WA was enabled when above described behavior is detected, the page fault likely occurred inside the driver.