Kernels with MSVC STL constructs sometimes give errors on debug checks

[TomArnaez]

Hi,

When using std::array for example in kernels with MSVC in debug:

q.submit([&](sycl::handler& cgh) {
        cgh.single_task([=]() {
            std::array<size_t, 5> arr;
            arr[0] = 3;
        });
    });

error: SYCL kernel cannot call a variadic function
STL_VERIFY(_Pos < _Size, "array subscript out of range");

I assume it's because SYCL device compiler cannot parse the MSVC STL macros. Is there any way around it? We're using CUDA

[AlexeySachkov]

Hi @TomArnaez,

That's kind of a known issue: the problem is that we do not control headers provided by MSVC and therefore cannot guarantee that they don't use anything that is illegal in SYCL.

For portability reasons I would advise against using std:: functionality in SYCL kernels, because SYCL spec does not guarantee full (or any) support for standard C++ library within SYCL kernels. With respect to our specific implementation I'm also not sure: AFAIK we do use std::array ourselves under the hood of sycl::vec, but at the same time I can't find any official documentation saying that we do guarantee support for std::array within kernels.

As to whether it is possible to workaround that. There seems to be a macro that allows to control level of instrumentation of MSVC STL headers: _ITERATOR_DEBUG_LEVEL. I'm not sure if it helps, but I would definitely try setting that macro to 0 from its default (in debug build) value 2.

Note that doing so plainly will also affect host part of the app. So, if you want to preserve those assertions for host part of the app, then you should use -Xsycl-target-frontend=<triple> <arg> or -Xsycl-target-frontend <arg> to only set the macro for device compilation, whilst preserving its default on host.