This post is basically a duplicate yet condensed version of all that happened on this post on the codeplay forum.
I have an issue where my SYCL kernels hangs and never finishes after it’s been computing for some time (~10s). This only happens when sycl::gpu_selector_v is selected for the sycl::queue command queue. Using sycl::cpu_selector_v shows no issues.
The kernel only hangs when there is a lot of computations to be done (see the details about LOOP_ITERATION and N below). One detail that might be relevant is that at first, just after launching the program, my laptop is barely usable due to the integrated GPU being used at its maximum. After a few seconds though, my laptop becomes usable again (as if it weren’t computing anything anymore) but the program still is running (and will never stop). At that point, my CPU will show a usage of ~50% (when using sycl::gpu_selector_v) until I decide to manually stop the program.
I managed to reproduce this issue on a simple example:
#include <sycl/sycl.hpp>
#include <vector>
#define LOOP_ITERATION 10000000
#define N 1000000
int main()
{
std::vector<float> v(N);
sycl::queue q{sycl::gpu_selector_v};
sycl::buffer buf{v};
q.submit([&](sycl::handler& cgh)
{
auto acc {buf.get_access(cgh,sycl::read_write)};
cgh.parallel_for(N, [=](sycl::id<1> id)
{
float x = 0.0f;
for (int i = 0; i < LOOP_ITERATION; i++)
x += i / 2;
acc[id] = x;
});
}).wait();
std::cout << "Done!" << std::endl;
return 0;
}
With the code posted above, I can only get the kernel to hang when N * LOOP_ITERATION is > 1 000 000 * 10 000 000. However, the kernel can still hang with lower LOOP_ITERATION (or N) values if we increase the complexity of the code inside the for (int i = 0; i < LOOP_ITERATION; i++) loop:
#define LOOP_ITERATION (10000000 / 10) //10 times less iterations
#define N 1000000
cgh.parallel_for(N, [=](sycl::id<1> id)
{
float x = 0.0f;
for (int i = 0; i < LOOP_ITERATION; i++)
{
x += i / 2;
float cosine = sycl::cos(sycl::sqrt(x));
float sine = sycl::sin(x);
float length = sycl::sqrt(cosine * cosine + sine * sine);
x /= sycl::cos(length) * sycl::sin(length);
}
acc[id] = x;
});
With LOOP_ITERATION divided by 10, the kernel never hangs unless the LOOP_ITERATION loop becomes more computationally demanding.
Sometimes but not always, this runtime error is thrown:
terminate called after throwing an instance of 'sycl::_V1::runtime_error'
what(): Native API failed. Native API returns: -14 (PI_ERROR_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST) -14 (PI_ERROR_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST)
I tried disabling the GPU hangcheck using this guide but to no avail. I also tried writing N to /sys/module/i915/parameters/enable_hangcheck and the solutions proposed here but none of this changed anything either…
For my immediate use case (ray tracing application where rendering an image with too many samples takes too long and hangs the kernel), there is a way for me to get around this issue by calling multiple smaller kernels (smaller amount of samples) but this isn't really a satisfactory solution, more like a workaround.
If relevant, I attached the output of the execution of the `clinfo` command on my system.
