[Regression][CUDA] Multiple cts/math_builtin_* tests crashed on build with error: llvm-foreach: Segmentation fault (core dumped)

[skosnits]

Describe the bug
[Regression][CUDA] Multiple cts/math_builtin_* tests crashed on build with error: llvm-foreach: Segmentation fault (core dumped)

First failed
https://github.com/intel/llvm.git c89a914
Last passed
https://github.com/intel/llvm.git 9f2b7bd

To Reproduce

1. Get SYCL with CUDA support (GSG) and (SYCL-CTS)
2. Build CTS tests

cd SYCL-CTS
mkdir build && cd build
cmake -G "Ninja" -DSYCL_IMPLEMENTATION=Intel_SYCL -DINTEL_SYCL_ROOT=<path to built sycl> -DOpenCL_LIBRARY=<path to built sycl>/lib/libOpenCL.so -DOpenCL_INCLUDE_DIR=<path to built sycl>/include/sycl -DCMAKE_BUILD_TYPE=Release -DCMAKE_CXX_FLAGS_RELEASE=" " -DINTEL_SYCL_FLAGS="-fsycl-device-code-split=per_source;-Wno-deprecated-declarations;-fsycl;-fsycl-unnamed-lambda" -DCMAKE_EXE_LINKER_FLAGS=" -Wl,-no-relax " -DINTEL_SYCL_TRIPLE=nvptx64-nvidia-cuda -DSYCL_CTS_ENABLE_OPENCL_INTEROP_TESTS=Off -DSYCL_CTS_ENABLE_DOUBLE_TESTS=On -DSYCL_CTS_ENABLE_HALF_TESTS=On ..
ninja -j 7 -k 0 -v test_math_builtin_api

Error message:

FAILED: bin/test_math_builtin_api.
: && <path to built sycl>/bin/clang++ -DSYCL2020_CONFORMANT_APIS -DSYCL2020_DISABLE_DEPRECATION_WARNINGS -D__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ llvm-foreach: Segmentation fault (core dumped)
clang++: error: ptxas command failed with exit code 254 (use -v to see invocation)
clang version 14.0.0 (https://github.com/intel/llvm.git c89a9148355b2e30d68513b5bed0d73b09515561)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: <path to built sycl>/bin
clang++: note: diagnostic msg: Error generating preprocessed source(s).
ninja: build stopped: cannot make progress due to previous errors.

Environment

• OS: Linux
• Target device and vendor: Nvidia GPU
• DPC++ version: https://github.com/intel/llvm.git c89a914
• CUDA version:

nvidia-smi
+-----------------------------------------------------------------------------+
| NVIDIA-SMI 470.57.02    Driver Version: 470.57.02    CUDA Version: 11.4     |
|-------------------------------+----------------------+----------------------+

[JackAKirk]

Thanks for the detailed description of this bug.

It occurs for some test cases in the math_builtin_float_double.cpp test file only, such as test cases using sycl::cos, and only when using -O1, -O2 or -O3 opt levels. This bug has only become apparent now because the default opt level used in the CUDA backend has been set to -O3, whereas previously it was -O0.
The origin of the bug appears to be the same as that of a previous issue: #4436 : the origin can be traced to a CUDA driver bug that is described here: https://forums.developer.nvidia.com/t/libdevice-functions-causing-ptxas-segfault/193352/2
Since CUDA toolkit 11.6 this bug has been fixed for all opt levels. Therefore it is recommended that the CTS test documentation should be updated to require CUDA toolkit 11.6 if possible. It is also worth noting that the CTS bug will not occur if the tests are compiled with no Opt (-O0) for all CUDA toolkit versions.

[bader]

Therefore it is recommended that the CTS test documentation should be updated to require CUDA toolkit 11.6 if possible.

@JackAKirk, don't you think it's should be required in DPC++ documentation rather than CTS? If you agree, please, update CDUA toolkit requirements in DPC++ documentation.

[JackAKirk]

Therefore it is recommended that the CTS test documentation should be updated to require CUDA toolkit 11.6 if possible.

@JackAKirk, don't you think it's should be required in DPC++ documentation rather than CTS? If you agree, please, update CDUA toolkit requirements in DPC++ documentation.

You're right. I'll update the DPC++ doc.

[bader]

Another option: instead of requiring 11.6 version, we can continue require 11.0, but add a note that there is a know issue with -O1/O2/O3 optimization options for CUDA toolkits prior to 11.6.
Both options seems fine to me.

[AerialMantis]

With the update to the Getting Started Guide merged we believe this issue to be resolved now, so we are closing the ticket.