[CUDA] CTS specialization_constants are failed for access violation on Windows #4896
Description
Describe the bug
CTS specialization_constants are failed for access violation with CUDA backend on Windows.
The following tests have the same issue:
specialization_constants_class_with_member_fun
specialization_constants_defined_various_ways_fp64
specialization_constants_multiple_fp16
specialization_constants_multiple_fp64
specialization_constants_same_name_inter_link_1st_tu_fp16
specialization_constants_same_name_inter_link_1st_tu_fp64
specialization_constants_same_name_inter_link_2nd_tu_fp16
specialization_constants_same_name_inter_link_2nd_tu_fp64
specialization_constants_same_name_stress_fp64
To Reproduce
- Get SYCL with CUDA support (GSG) and [SYCL-CTS](https://github.com/KhronosGroup/SYCL-CTS KhronosGroup/SYCL-CTS@14fac44)
- Build CTS tests
cd SYCL-CTS
mkdir build && cd build
cmake -GNinja -DSYCL_IMPLEMENTATION=Intel_SYCL -DINTEL_SYCL_ROOT=<path to built sycl> -Dopencl_platform_name=nvidia -Dopencl_device_name=opencl_gpu -DCMAKE_BUILD_TYPE=Release -DCMAKE_CXX_FLAGS_RELEASE="-fsycl -fsycl-unnamed-lambda /EHsc /EHsc /MD -Wno-deprecated-declarations" -DINTEL_SYCL_FLAGS="-Xsycl-target-backend;--cuda-gpu-arch=sm_50" -DCMAKE_EXE_LINKER_FLAGS=" -Wl,-no-relax " -DINTEL_SYCL_TRIPLE=nvptx64-nvidia-cuda-sycldevice -DSYCL_CTS_ENABLE_OPENCL_INTEROP_TESTS=Off -DSYCL_CTS_ENABLE_DOUBLE_TESTS=On -DSYCL_CTS_ENABLE_HALF_TESTS=On ..
ninja test_specialization_constants
- Run test
export SYCL_DEVICE_FILTER="cuda:gpu:0"
export SYCL_ENABLE_HOST_DEVICE="1"
bin/test_specialization_constants -p nvidia -d opencl_gpu --test specialization_constants_class_with_member_fun
Error message:
$ ./bin/test_specialization_constants.exe -p nvidia -d opencl_gpu --test specialization_constants_class_with_member_fun
--- specialization_constants_class_with_member_fun
Segmentation fault
Environment
- OS: Windows
- Target device and vendor: Nvidia GPU
- DPC++ version: https://github.com/intel/llvm.git 35ccdd5
- CUDA version:
Cuda compilation tools, release 11.4, V11.4.120
Build cuda_11.4.r11.4/compiler.30300941_0
Additional context
The exception is shown in queue.submit from function run in specialization_constants_class_with_member_fun.cpp:
$ sed -n "39,68p" tests/specialization_constants/specialization_constants_class_with_member_fun.cpp
void run(util::logger &log) override {
try {
auto queue = util::get_cts_object::queue();
sycl::range<1> range(1);
const int val_A = 3;
const int val_B = 4;
const int val_C = 2;
sc_class_with_memb result(0, 0);
sc_class_with_memb ref(val_A, val_B);
int kernel_result_val = 0;
{
sycl::buffer<sc_class_with_memb, 1> result_buffer(&result, range);
sycl::buffer<int, 1> kernel_result_val_buffer(&kernel_result_val,
range);
queue.submit([&](sycl::handler &cgh) { //<-- throw exception
auto res_acc =
result_buffer.template get_access<sycl::access_mode::write>(cgh);
auto kernel_res_val_acc =
kernel_result_val_buffer
.template get_access<sycl::access_mode::write>(cgh);
cgh.set_specialization_constant<sc_cl_w_mem_fn>(ref);
cgh.single_task<class sc_cl_w_mem_fn_kernel>(
[=](sycl::kernel_handler h) {
res_acc[0] = h.get_specialization_constant<sc_cl_w_mem_fn>();
kernel_res_val_acc[0] =
h.get_specialization_constant<sc_cl_w_mem_fn>().calculate(
val_C);
});
});
}