[SYCL][CUDA][HIP] Segmentation fault from emitFunctionWithArgsBeginTrace

[zjin-lcf]

Running the updated example
https://github.com/zjin-lcf/HeCBench/tree/master/reverse-sycl

on an Nvidia/AMD GPU causes segfault. This is not reproducible on an Intel GPU.

To reproduce:
make CUDA=yes CUDA_ARCH=sm_60 run

gdb shows the source of the error:

Thread 1 "main" received signal SIGSEGV, Segmentation fault.
0x00007ffff71437a0 in cl::sycl::detail::pi::emitFunctionWithArgsBeginTrace(unsigned int, char const*, unsigned char*, _pi_plugin)

Thanks

[pgorlani]

I had a look at the segfault, and I'm able to reproduce it for CUDA and Intel backends. This segfault happens in the deconstruction of command group execution graph. In its actual form, reverse-sycl benchmark spawns a long chain of dependent events which need to be deleted on exit. Actually, this deletion occurs in a recursive manner. In our case, this recursion fills the stack memory causing the segfault.

You have two options to avoid this, the first is to set as unlimited the stack size with $ ulimit -s unlimited.
The second is moving buffer d_test definition (line 40) within the outer loop (line 45). This shorts the event chain allowing its deconstruction to fit in the default stack size.

[zjin-lcf]

I ran the program with the first option. I assume that buffer may be created/destroyed for every outer loop iteration in the second option. Thank you for your comments and options!

[zjin-lcf]

Please advise if I should close the issue.

[pgorlani]

Thanks for your feedback. If you are happy of the proposed solutions, feel free to close this issue.

[zjin-lcf]

I will add a readme with your explanation in the example.

[zjin-lcf]

There is no segfault when running the CUDA program. After porting CUDA programs to SYCL, people may encounter the same issue without any clue. I think this is serious. May your solution be added to the SYCL user guide ? Thanks.

[zjin-lcf]

Some examples observe the segfaults too. For example,

https://github.com/zjin-lcf/HeCBench/tree/master/mnist-sycl

[alexbatashev]

There is no segfault when running the CUDA program. After porting CUDA programs to SYCL, people may encounter the same issue without any clue. I think this is serious. May your solution be added to the SYCL user guide ? Thanks.

This is not something specific to SYCL. You can run into the same issue by writing poorly designed merge-sort, for example. There's no good way to describe the symptoms user may see.

[zjin-lcf]

Users will continue to report the issue when they port CUDA programs to SYCL, and then run the SYCL programs on an Nvidia GPU. All I ask is whether you may add pgorlani's suggestion to your user guide.

It is not about writing a poorly designed merge-sort; it is about letting developers and researchers know that they need to adjust stack size sometimes to avoid the segfault when running a ported program.

@bader
Is this a good suggestion ?