Assertion failure in device_kernel_info.cpp:67 while executing on any SYCL device

[sycloid]

Describe the bug

This is related to IntelPython/dpctl#2156

Executing Python test suite for dpctl Python project compiled with SYCL nightly build, and using its SYCL runtime library, and running on AMD Epyc CPU for which CPU compiler identifies ARCH as core-avx2 an assertion assert(Info == *this); at https://github.com/intel/llvm/blob/sycl/sycl/source/detail/device_kernel_info.cpp#L71

To reproduce it on Intel CPU one needs to either set CL_CONFIG_CPU_VECTORIZER_MODE=8 or CL_CONFIG_CPU_TARGET_ARCH=core-avx2.

To reproduce

Do not have a simple reproducible.

1. Start with SYCL nightly bundle, and TBB 2022.2.0 and OCLCPU=2025.20.6.0.04_224945 dependencies, so that sycl-ls sees the CPU using this OCLCPU driver. Set environment variable DPCPP_HOME to that ${DPCPP_HOME}/bin/clang++ locates C++ clang compiler.
2. Download Miniforge Python package manager: wget https://github.com/conda-forge/miniforge/releases/latest/download/Miniforge3-Linux-x86_64.sh
3. Install it locally: chmod 755 ./Miniforge3-Linux-x86_64.sh && ./Miniforge3-Linux-x86_64.sh -b -p $(pwd)/miniforge
4. Activate it: source $(pwd)/miniforge/etc/profile.d/conda.sh
5. Create Python environment to build dpctl: conda create -n dev_dpctl --yes python=3.13 cython numpy scikit-build pytest versioneer cmake ninja.
6. Clone repository: git clone https://github.com/IntelPython/dpctl.git
7. Change to the cloned directory cd dpctl
8. Start building python scripts/build_locally.py --compiler-root=${DPCPP_HOME}--c-compiler=${DPCPP_HOME}/bin/clang --cxx-compiler=${DPCPP_HOME}/bin/clang++
9. Execute test that results in assertion failure: OCL_ICD_FILENAMES=~/sycl-builds/oclcpuexp/x64/libintelocl.so OCL_ICD_VENDORS=. CL_CONFIG_CPU_VECTORIZER_MODE=8 ONEAPI_DEVICE_SELECTOR=opencl:cpu gdb --args python -m pytest -s dpctl/tests

Environment

• OS: Linux, Ubuntu 24.04 for WSL
• Target device and vendor: CPU
• DPC++ version:

clang version 21.0.0git (https://github.com/intel/llvm 3906cea47a1e9e564ea3a10ef21d9b7e65c560c0)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /home/sycloid/sycl-builds/16-09-2025/bin
Build config: +assertions

• Dependencies version: [e.g. the output of sycl-ls --verbose]

INFO: Output filtered by ONEAPI_DEVICE_SELECTOR environment variable, which is set to opencl:cpu.
To see device ids, use the --ignore-device-selectors CLI option.

<LOADER>[INFO]: loaded adapter 0x0x55b25b34cd00 (libur_adapter_opencl.so.0) from /home/sycloid/sycl-builds/07-09-2025/lib/libur_adapter_opencl.so.0
[opencl:cpu] Intel(R) OpenCL, Intel(R) Core(TM) Ultra 7 155U OpenCL 3.0 (Build 0) [2025.20.6.0.04_224945]

Platforms: 1
Platform [#1]:
    Version  : OpenCL 3.0 LINUX
    Name     : Intel(R) OpenCL
    Vendor   : Intel(R) Corporation
    Devices  : 1
        Type              : cpu
        Version           : OpenCL 3.0 (Build 0)
        Name              : Intel(R) Core(TM) Ultra 7 155U
        Vendor            : Intel(R) Corporation
        Driver            : 2025.20.6.0.04_224945
        UUID              : 8680a406-0a00-2332-faff-fffb8b1f0000
        DeviceID          : 657060
        Num SubDevices    : 0
        Num SubSubDevices : 0
        Aspects           : cpu fp16 fp64 online_compiler online_linker queue_profiling usm_device_allocations usm_host_allocations usm_shared_allocations usm_system_allocations ext_intel_gpu_slices ext_intel_gpu_subslices_per_slice ext_intel_gpu_eu_count_per_subslice usm_atomic_host_allocations usm_atomic_shared_allocations atomic64 ext_intel_device_info_uuid ext_oneapi_srgb ext_oneapi_native_assert ext_intel_gpu_hw_threads_per_eu ext_intel_device_id ext_intel_legacy_image ext_oneapi_fragment ext_oneapi_chunk ext_oneapi_tangle ext_oneapi_limited_graph ext_oneapi_private_alloca ext_oneapi_atomic16 ext_oneapi_virtual_functions ext_intel_device_info_node_mask
        info::device::sub_group_sizes: 4 8 16 32 64
        Architecture: x86_64
default_selector()      : cpu, Intel(R) OpenCL, Intel(R) Core(TM) Ultra 7 155U OpenCL 3.0 (Build 0) [2025.20.6.0.04_224945]
accelerator_selector()  : No device of requested type available.
cpu_selector()          : cpu, Intel(R) OpenCL, Intel(R) Core(TM) Ultra 7 155U OpenCL 3.0 (Build 0) [2025.20.6.0.04_224945]
gpu_selector()          : No device of requested type available.
custom_selector(gpu)    : No device of requested type available.
custom_selector(cpu)    : cpu, Intel(R) OpenCL, Intel(R) Core(TM) Ultra 7 155U OpenCL 3.0 (Build 0) [2025.20.6.0.04_224945]
custom_selector(acc)    : No device of requested type available.

Additional context

The backtrace of crash from GDB:

.............................................................................................python: /__w/llvm/llvm/src/sycl/source/detail/device_kernel_info.cpp:67: void sycl::_V1::detail::DeviceKernelInfo::setCompileTimeInfoIfNeeded(const sycl::_V1::detail::compile_time_kernel_info_v1::CompileTimeKernelInfoTy&): Assertion `Info == *this' failed.

Thread 1 "python" received signal SIGABRT, Aborted.
Download failed: Invalid argument.  Continuing without source file ./nptl/./nptl/pthread_kill.c.
__pthread_kill_implementation (no_tid=0, signo=6, threadid=<optimized out>) at ./nptl/pthread_kill.c:44
warning: 44     ./nptl/pthread_kill.c: No such file or directory
(gdb) bt
#0  __pthread_kill_implementation (no_tid=0, signo=6, threadid=<optimized out>) at ./nptl/pthread_kill.c:44
#1  __pthread_kill_internal (signo=6, threadid=<optimized out>) at ./nptl/pthread_kill.c:78
#2  __GI___pthread_kill (threadid=<optimized out>, signo=signo@entry=6) at ./nptl/pthread_kill.c:89
#3  0x00007ffff7cef27e in __GI_raise (sig=sig@entry=6) at ../sysdeps/posix/raise.c:26
#4  0x00007ffff7cd28ff in __GI_abort () at ./stdlib/abort.c:79
#5  0x00007ffff7cd281b in __assert_fail_base (fmt=0x7ffff7e7a1e8 "%s%s%s:%u: %s%sAssertion `%s' failed.\n%n",
    assertion=assertion@entry=0x7ffff67d22fc "Info == *this",
    file=file@entry=0x7ffff67d21c8 "/__w/llvm/llvm/src/sycl/source/detail/device_kernel_info.cpp", line=line@entry=67,
    function=function@entry=0x7ffff67d2250 "void sycl::_V1::detail::DeviceKernelInfo::setCompileTimeInfoIfNeeded(const sycl::_V1::detail::compile_time_kernel_info_v1::CompileTimeKernelInfoTy&)") at ./assert/assert.c:96
#6  0x00007ffff7ce5517 in __assert_fail (assertion=0x7ffff67d22fc "Info == *this",
    file=0x7ffff67d21c8 "/__w/llvm/llvm/src/sycl/source/detail/device_kernel_info.cpp", line=67,
    function=0x7ffff67d2250 "void sycl::_V1::detail::DeviceKernelInfo::setCompileTimeInfoIfNeeded(const sycl::_V1::detail::compile_time_kernel_info_v1::CompileTimeKernelInfoTy&)") at ./assert/assert.c:105
#7  0x00007ffff668b270 in sycl::_V1::detail::DeviceKernelInfo::setCompileTimeInfoIfNeeded(sycl::_V1::detail::compile_time_kernel_info_v1::CompileTimeKernelInfoTy const&) () from /home/sycloid/sycl-builds/16-09-2025/lib/libsycl.so.8
#8  0x00007ffff66c499e in sycl::_V1::detail::ProgramManager::getOrCreateDeviceKernelInfo(sycl::_V1::detail::compile_time_kernel_info_v1::CompileTimeKernelInfoTy const&) () from /home/sycloid/sycl-builds/16-09-2025/lib/libsycl.so.8
#9  0x00007fff811050c4 in ?? () from /home/sycloid/repos/dpctl/dpctl/tensor/_tensor_impl.cpython-313-x86_64-linux-gnu.so
#10 0x00007fff81104be0 in ?? () from /home/sycloid/repos/dpctl/dpctl/tensor/_tensor_impl.cpython-313-x86_64-linux-gnu.so
#11 0x00007ffff66e72f1 in sycl::_V1::detail::queue_impl::submit_impl(sycl::_V1::detail::type_erased_cgfo_ty const&, bool, sycl::_V1::detail::code_location const&, bool, sycl::_V1::detail::v1::SubmissionInfo const&) () from /home/sycloid/sycl-builds/16-09-2025/lib/libsycl.so.8
#12 0x00007ffff67aaab6 in sycl::_V1::queue::submit_with_event_impl(sycl::_V1::detail::type_erased_cgfo_ty const&, sycl::_V1::detail::v1::SubmissionInfo const&, sycl::_V1::detail::code_location const&, bool) const () from /home/sycloid/sycl-builds/16-09-2025/lib/libsycl.so.8
#13 0x00007fff810b69f9 in ?? () from /home/sycloid/repos/dpctl/dpctl/tensor/_tensor_impl.cpython-313-x86_64-linux-gnu.so
#14 0x00007fff81101ddd in ?? () from /home/sycloid/repos/dpctl/dpctl/tensor/_tensor_impl.cpython-313-x86_64-linux-gnu.so
#15 0x00007fff810ffef5 in ?? () from /home/sycloid/repos/dpctl/dpctl/tensor/_tensor_impl.cpython-313-x86_64-linux-gnu.so
#16 0x00007fff810ff361 in ?? () from /home/sycloid/repos/dpctl/dpctl/tensor/_tensor_impl.cpython-313-x86_64-linux-gnu.so
#17 0x00007fff810af129 in ?? () from /home/sycloid/repos/dpctl/dpctl/tensor/_tensor_impl.cpython-313-x86_64-linux-gnu.so
#18 0x00007fff810ab20b in ?? () from /home/sycloid/repos/dpctl/dpctl/tensor/_tensor_impl.cpython-313-x86_64-linux-gnu.so
#19 0x00007fff810ab18f in ?? () from /home/sycloid/repos/dpctl/dpctl/tensor/_tensor_impl.cpython-313-x86_64-linux-gnu.so
#20 0x00007fff8107e5b3 in ?? () from /home/sycloid/repos/dpctl/dpctl/tensor/_tensor_impl.cpython-313-x86_64-linux-gnu.so
#21 0x000055555572f644 in cfunction_call (func=0x7fff81f903b0, args=0x7fff19630bc0, kwargs=<optimized out>)
    at /usr/local/src/conda/python-3.13.7/Objects/methodobject.c:539
#22 0x00005555556fdddc in _PyObject_MakeTpCall (tstate=0x555555ac45b0 <_PyRuntime+283024>, callable=0x7fff81f903b0, args=0x7ffff7ae7820, nargs=2,
    keywords=<optimized out>) at /usr/local/src/conda/python-3.13.7/Objects/call.c:242
#23 0x00005555557176a9 in _PyEval_EvalFrameDefault (tstate=<optimized out>, frame=<optimized out>, throwflag=<optimized out>)

[sycloid]

Actually, the assertion is tripped even without setting any CL_CONFIG_* variables:

$ OCL_ICD_FILENAMES=~/sycl-builds/oclcpuexp/x64/libintelocl.so OCL_ICD_VENDORS=. ONEAPI_DEVICE_SELECTOR=opencl:cpu python -m pytest -s dpctl/tests/
<snipped>
dpctl/tests/test_usm_ndarray_indexing.py ......................................................................................................................................................................................................................................................................................................................................................................................................................................................................................python: /__w/llvm/llvm/src/sycl/source/detail/device_kernel_info.cpp:67: void sycl::_V1::detail::DeviceKernelInfo::setCompileTimeInfoIfNeeded(const sycl::_V1::detail::compile_time_kernel_info_v1::CompileTimeKernelInfoTy&): Assertion `Info == *this' failed.
Fatal Python error: Aborted

<snipped>

[sycloid]

I can reliably reproduce the issue running just two test files out of the entire test suite:

$ OCL_ICD_FILENAMES=~/sycl-builds/oclcpuexp/x64/libintelocl.so OCL_ICD_VENDORS=. ONEAPI_DEVICE_SELECTOR=opencl:cpu python -m pytest -s dpctl/tests/test_tensor_accumulation.py dpctl/tests/test_usm_ndarray_indexing.py

[sycloid]

Smaller yet, just two tests:

OCL_ICD_FILENAMES=~/sycl-builds/oclcpuexp/x64/libintelocl.so \
    OCL_ICD_VENDORS=. \
    ONEAPI_DEVICE_SELECTOR=opencl:cpu \
     python -m pytest -sv \
         dpctl/tests/test_tensor_accumulation.py::test_contig_cumsum_sint[dt2] \
         dpctl/tests/test_usm_ndarray_indexing.py::test_nonzero_large

[sycloid]

It appears this issue is not specific to CPU at all. It happens with Level-Zero too:

$ ONEAPI_DEVICE_SELECTOR=level_zero:gpu python -m pytest -sv dpctl/tests/test_tensor_accumulation.py::test_contig_cumsum_sint[dt2] dpctl/tests/test_usm_ndarray_indexi
ng.py::test_nonzero_large
============================================================================================== test session starts ==============================================================================================
platform linux -- Python 3.13.7, pytest-8.4.2, pluggy-1.6.0 -- /home/sycloid/miniforge/envs/dev_dpctl/bin/python
cachedir: .pytest_cache
rootdir: /home/sycloid/repos/dpctl
configfile: pyproject.toml
collected 2 items

dpctl/tests/test_tensor_accumulation.py::test_contig_cumsum_sint[dt2] PASSED
dpctl/tests/test_usm_ndarray_indexing.py::test_nonzero_large python: /__w/llvm/llvm/src/sycl/source/detail/device_kernel_info.cpp:67: void sycl::_V1::detail::DeviceKernelInfo::setCompileTimeInfoIfNeeded(const sycl::_V1::detail::compile_time_kernel_info_v1::CompileTimeKernelInfoTy&): Assertion `Info == *this' failed.
Fatal Python error: Aborted

Current thread 0x00007f02e14f0740 (most recent call first):
  File "/home/sycloid/repos/dpctl/dpctl/tensor/_copy_utils.py", line 795 in _extract_impl
  File "/home/sycloid/repos/dpctl/dpctl/tests/test_usm_ndarray_indexing.py", line 1473 in test_nonzero_large
  File "/home/sycloid/miniforge/envs/dev_dpctl/lib/python3.13/site-packages/_pytest/python.py", line 157 in pytest_pyfunc_call
<snipped>

[sycloid]

Having built the dpctl with Debug build of DPC++ this is info from GDB:

(gdb) list
62      void DeviceKernelInfo::setCompileTimeInfoIfNeeded(
63          const CompileTimeKernelInfoTy &Info) {
64        if (!isCompileTimeInfoSet())
65          CompileTimeKernelInfoTy::operator=(Info);
66        assert(isCompileTimeInfoSet());
67        assert(Info == *this);
68      }
69
70      FastKernelSubcacheT &DeviceKernelInfo::getKernelSubcache() {
71        assertInitialized();
(gdb) p Info
$5 = (const sycl::_V1::detail::compile_time_kernel_info_v1::CompileTimeKernelInfoTy &) @0x7fff814ffcd8: {Name = {
    str = 0x7fff7ee066dc "_ZTSN5dpctl6tensor7kernels12accumulators42inclusive_scan_iter_local_scan_blocked_krnIiiLj4ENS0_12offset_utils22TwoZeroOffsets_IndexerENS4_11NoOpIndexerES6_NS2_15NoOpTransformerIiEESt4plusIiELb0EEE"}, NumParams = 13, IsESIMD = false, FileName = {str = 0x7fff7ee00350 ""}, FunctionName = {str = 0x7fff7ee00350 ""}, LineNumber = 0, ColumnNumber = 0,
  KernelSize = 120, ParamDescGetter = 0x7fff7e44c520, HasSpecialCaptures = true}
(gdb) p *this
$6 = {<sycl::_V1::detail::compile_time_kernel_info_v1::CompileTimeKernelInfoTy> = {Name = {
      str = 0x7fff721955b2 "_ZTSN5dpctl6tensor7kernels12accumulators42inclusive_scan_iter_local_scan_blocked_krnIiiLj4ENS0_12offset_utils22TwoZeroOffsets_IndexerENS4_11NoOpIndexerES6_NS2_15NoOpTransformerIiEESt4plusIiELb0EEE"}, NumParams = 13, IsESIMD = false, FileName = {str = 0x7fff72155efd ""}, FunctionName = {str = 0x7fff72155efd ""}, LineNumber = 0, ColumnNumber = 0,
    KernelSize = 120, ParamDescGetter = 0x7fff71e83500, HasSpecialCaptures = true},
  Name = "_ZTSN5dpctl6tensor7kernels12accumulators42inclusive_scan_iter_local_scan_blocked_krnIiiLj4ENS0_12offset_utils22TwoZeroOffsets_IndexerENS4_11NoOpIndexerES6_NS2_15NoOpTransformerIiEESt4plusIiELb0EEE", MInitialized = std::atomic<bool> = { true }, MFastKernelSubcache = {Entries = std::vector of length 1, capacity 1 = {{Key = {first = 0x555559e732d0,
          second = 0x555559e76500}, Value = std::shared_ptr<sycl::_V1::detail::FastKernelCacheVal> (use count 1, weak count 0) = {get() = 0x55555b38f2d0}}}, Mutex = {
      MLock = {<std::__atomic_flag_base> = {_M_i = false}, <No data fields>}}}, MUsesAssert = false, MImplicitLocalArgPos = std::optional [no contained value]}

[sycloid]

@bratpiorka This is not specific to OpenCL. The assertion is tripped for Level-Zero devices as well as for OpenCL devices (both GPU and CPU)

[sycloid]

Judging from the dpctl logs for workflow building and testing the project with nightly builds of SYCL bundle, the problem was introduced between Sept. 10 and Sept. 11:

I can reproduce the problem using installed SYCL nightly from Sept 16, but it is not reproducible using build from Sept. 7

[bratpiorka]

@sycloid thanks for the investigation, this would really help to identify the root cause of this failure

[sycloid]

@bratpiorka The commit that introduced the issue is cb03a1b that merged #20003

The SYCL built from the previous merge commit e26848d does not exhibit the assertion failure.