"Undefined temporary symbol" for symbols whose mangled names are embedded in SYCL kernel entry-point names

[mctylr-gh]

Pre-submission Checklist

• I am using the latest GPU driver version (releases)
• I have searched for similar issues and found none

GPU Hardware

Intel Arc Pro B70

DRI Devices Information

$ ls -l /dev/dri/*
crw-rw----+ root video    0 B Tue Sep 15 02:24:36 2026  /dev/dri/card0
crw-rw----+ root video    0 B Mon Sep 21 00:07:43 2026  /dev/dri/card1
crw-rw-rw-+ root render   0 B Tue Sep 15 02:24:35 2026  /dev/dri/renderD128
crw-rw-rw-+ root render   0 B Tue Sep 15 02:24:36 2026  /dev/dri/renderD129

/dev/dri/by-path:
lrwxrwxrwx root root  8 B Tue Sep 15 02:24:35 2026  pci-0000:06:00.0-card ⇒ ../card1
lrwxrwxrwx root root 13 B Tue Sep 15 02:24:35 2026  pci-0000:06:00.0-render ⇒ ../renderD128
lrwxrwxrwx root root  8 B Tue Sep 15 02:24:36 2026  pci-0000:0e:00.0-card ⇒ ../card0
lrwxrwxrwx root root 13 B Tue Sep 15 02:24:36 2026  pci-0000:0e:00.0-render ⇒ ../renderD129

Other video device present is an AMD Radeon 6600 (Navi 23) used for video display, not compute.

GPU Detailed Information (lspci output)

$ sudo lspci -vvv -k -s 0000:0e:00.0
0e:00.0 VGA compatible controller: Intel Corporation Battlemage G31 [Intel Graphics] (prog-if 00 [VGA controller])
        Subsystem: Intel Corporation Device 1701
        Control: I/O- Mem+ BusMaster+ SpecCycle- MemWINV- VGASnoop- ParErr- Stepping- SERR- FastB2B- DisINTx-
        Status: Cap+ 66MHz- UDF- FastB2B- ParErr- DEVSEL=fast >TAbort- <TAbort- <MAbort+ >SERR- <PERR- INTx-
        Latency: 0, Cache Line Size: 64 bytes
        Interrupt: pin ? routed to IRQ 140
        IOMMU group: 31
        Region 0: Memory at 7804000000 (64-bit, prefetchable) [size=16M]
        Region 2: Memory at 6800000000 (64-bit, prefetchable) [size=32G]
        Expansion ROM at fbe00000 [disabled] [size=2M]
        Capabilities: [40] Vendor Specific Information: Len=0c <?>
        Capabilities: [70] Express (v2) Endpoint, MSI 00
                DevCap: MaxPayload 256 bytes, PhantFunc 0, Latency L0s unlimited, L1 unlimited
                        ExtTag+ AttnBtn- AttnInd- PwrInd- RBE+ FLReset+ SlotPowerLimit 0W
                DevCtl: CorrErr- NonFatalErr- FatalErr- UnsupReq-
                        RlxdOrd+ ExtTag+ PhantFunc- AuxPwr- NoSnoop+ FLReset-
                        MaxPayload 256 bytes, MaxReadReq 512 bytes
                DevSta: CorrErr- NonFatalErr- FatalErr- UnsupReq- AuxPwr- TransPend-
                LnkCap: Port #0, Speed 2.5GT/s, Width x1, ASPM L0s L1, Exit Latency L0s <64ns, L1 <1us
                        ClockPM- Surprise- LLActRep- BwNot- ASPMOptComp+
                LnkCtl: ASPM Disabled; RCB 64 bytes, Disabled- CommClk-
                        ExtSynch- ClockPM- AutWidDis- BWInt- AutBWInt-
                LnkSta: Speed 2.5GT/s, Width x1
                        TrErr- Train- SlotClk- DLActive- BWMgmt- ABWMgmt-
                DevCap2: Completion Timeout: Range B, TimeoutDis+ NROPrPrP- LTR+
                         10BitTagComp+ 10BitTagReq+ OBFF Not Supported, ExtFmt+ EETLPPrefix-
                         EmergencyPowerReduction Not Supported, EmergencyPowerReductionInit-
                         FRS- TPHComp- ExtTPHComp-
                         AtomicOpsCap: 32bit- 64bit- 128bitCAS-
                DevCtl2: Completion Timeout: 50us to 50ms, TimeoutDis- LTR+ 10BitTagReq- OBFF Disabled,
                         AtomicOpsCtl: ReqEn-
                LnkCap2: Supported Link Speeds: 2.5GT/s, Crosslink- Retimer- 2Retimers- DRS-
                LnkCtl2: Target Link Speed: 2.5GT/s, EnterCompliance- SpeedDis-
                         Transmit Margin: Normal Operating Range, EnterModifiedCompliance- ComplianceSOS-
                         Compliance Preset/De-emphasis: -6dB de-emphasis, 0dB preshoot
                LnkSta2: Current De-emphasis Level: -6dB, EqualizationComplete- EqualizationPhase1-
                         EqualizationPhase2- EqualizationPhase3- LinkEqualizationRequest-
                         Retimer- 2Retimers- CrosslinkRes: unsupported
        Capabilities: [ac] MSI: Enable+ Count=1/1 Maskable+ 64bit+
                Address: 00000000fee00000  Data: 0000
                Masking: 00000000  Pending: 00000000
        Capabilities: [d0] Power Management version 3
                Flags: PMEClk- DSI- D1- D2- AuxCurrent=0mA PME(D0+,D1-,D2-,D3hot+,D3cold-)
                Status: D3 NoSoftRst+ PME-Enable- DSel=0 DScale=0 PME-
        Capabilities: [100 v1] Alternative Routing-ID Interpretation (ARI)
                ARICap: MFVC- ACS-, Next Function: 0
                ARICtl: MFVC- ACS-, Function Group: 0
        Capabilities: [110 v1] Null
        Capabilities: [200 v1] Address Translation Service (ATS)
                ATSCap: Invalidate Queue Depth: 00
                ATSCtl: Enable+, Smallest Translation Unit: 00
        Capabilities: [420 v1] Physical Resizable BAR
                BAR 2: current size: 32GB, supported: 256MB 512MB 1GB 2GB 4GB 8GB 16GB 32GB
        Capabilities: [220 v1] Virtual Resizable BAR
                BAR 2: current size: 8GB, supported: 256MB 512MB 1GB 2GB 4GB 8GB 16GB 32GB
        Capabilities: [320 v1] Single Root I/O Virtualization (SR-IOV)
                IOVCap: Migration- 10BitTagReq+ Interrupt Message Number: 000
                IOVCtl: Enable- Migration- Interrupt- MSE- ARIHierarchy+ 10BitTagReq-
                IOVSta: Migration-
                Initial VFs: 4, Total VFs: 4, Number of VFs: 0, Function Dependency Link: 00
                VF offset: 1, stride: 1, Device ID: e223
                Supported Page Size: 00000553, System Page Size: 00000001
                Region 0: Memory at 0000007800000000 (64-bit, prefetchable)
                Region 2: Memory at 0000007000000000 (64-bit, prefetchable)
                VF Migration: offset: 00000000, BIR: 0
        Capabilities: [400 v1] Latency Tolerance Reporting
                Max snoop latency: 1048576ns
                Max no snoop latency: 1048576ns
        Kernel driver in use: xe
        Kernel modules: xe

Driver Version

26.35.39758.10-0

Installed GPU Driver Packages

$ sudo dpkg --list | grep -iE "igc|gmm|opencl|level-zero|level_zero|ocloc|libze"
ii  amdrocm-opencl                                        10.0.0~pre4-32424980240                                         amd64        AMD ROCm opencl libraries
ii  amdrocm-opencl-dev10.0                                10.0.0-4                                                        amd64        AMD ROCm opencl development files.
ii  amdrocm-opencl10.0                                    10.0.0-4                                                        amd64        AMD ROCm opencl libraries
ii  clinfo                                                3.0.23.01.25-1build1                                            amd64        Query OpenCL system information
ii  cuda-opencl-13-3                                      13.3.27-1                                                       amd64        CUDA OpenCL native Libraries
ii  intel-igc-core-2                                      2.41.5                                                          amd64        Intel(R) Graphics Compiler for OpenCL(TM)
ii  intel-igc-core-devel                                  2.41.5                                                          amd64        Intel(R) Graphics Compiler for OpenCL(TM)
ii  intel-igc-opencl-2                                    2.41.5                                                          amd64        Intel(R) Graphics Compiler for OpenCL(TM)
ii  intel-igc-opencl-devel                                2.41.5                                                          amd64        Intel(R) Graphics Compiler for OpenCL(TM)
ii  intel-ocloc                                           26.35.39758.10-0                                                amd64        Tool for managing Intel Compute GPU device binary format
ii  intel-oneapi-runtime-dpcpp-sycl-opencl-cpu            2026.1.1-325                                                    amd64        Intel® CPU Runtime for OpenCL(TM) Applications runtime
ii  intel-oneapi-runtime-opencl                           2026.1.1-325                                                    amd64        Intel® CPU Runtime for OpenCL(TM) Applications runtime
ii  intel-opencl-icd                                      26.35.39758.10-0                                                amd64        Intel graphics compute runtime for OpenCL
ii  libclc-21-dev                                         1:21.1.8~++20251221032922+2078da43e25a-1~exp1~20251221153059.70 all          OpenCL C language implementation - development files
ii  libigc-dev                                            2.40.13+ds1-1~24.04                                             amd64        Core development files for Intel(R) Graphics Compiler for OpenCL(TM)
ii  libigc2                                               2.40.13+ds1-1~24.04                                             amd64        Core libraries for Intel(R) Graphics Compiler for OpenCL(TM)
ii  libigdfcl-dev                                         2.40.13+ds1-1~24.04                                             amd64        OpenCL development files for Intel(R) Graphics Compiler for OpenCL(TM)
ii  libigdfcl2                                            2.40.13+ds1-1~24.04                                             amd64        OpenCL library for Intel(R) Graphics Compiler for OpenCL(TM)
ii  libigdgmm12:amd64                                     22.10.0                                                         amd64        Intel Graphics Memory Management Library -- shared library
ii  libkf5configcore5:amd64                               5.115.0-0ubuntu5                                                amd64        configuration settings framework for Qt
ii  libsigc++-2.0-0v5:amd64                               2.12.1-2                                                        amd64        type-safe Signal Framework for C++ - runtime
ii  libsigc++-3.0-0:amd64                                 3.6.0-2                                                         amd64        type-safe Signal Framework for C++ - runtime
ii  libze-dev:amd64                                       1.32.0-1~24.04~ppa1                                             amd64        oneAPI Level Zero -- development files
ii  libze-intel-gpu-dev                                   26.31.39395.14-1~24.04~ppa1                                     amd64        Intel oneAPI L0 support implementation for Intel GPUs -- devel files
ii  libze-intel-gpu-raytracing                            1.2.4-2~24.04~ppa3                                              amd64        Level Zero Ray Tracing Support library
ii  libze-intel-gpu1                                      26.35.39758.10-0                                                amd64        Intel(R) Graphics Compute Runtime for oneAPI Level Zero.
ii  libze1:amd64                                          1.32.0-1~24.04~ppa1                                             amd64        oneAPI Level Zero -- share libraries
rc  mesa-opencl-icd:amd64                                 25.2.8-0ubuntu0.24.04.2                                         amd64        free implementation of the OpenCL API -- ICD runtime
ii  ocl-icd-libopencl1:amd64                              2.3.2-1build1                                                   amd64        Generic OpenCL ICD Loader
ii  ocl-icd-libopencl1:i386                               2.3.2-1build1                                                   i386         Generic OpenCL ICD Loader
ii  ocl-icd-opencl-dev:amd64                              2.3.2-1build1                                                   amd64        OpenCL development files
ii  opencl-c-headers                                      3.0~2023.12.14-1                                                all          OpenCL (Open Computing Language) C header files
ii  opencl-clhpp-headers                                  3.0~2023.12.14-1ubuntu1                                         all          C++ headers for OpenCL development

Driver Installation Details

Normally use intel-graphics PPA for Ubuntu, but as per your guidelines, I have installed the latest release from github.com/intel/compute-runtime/releases which is 26.35.39758.10
26.35.39758.10 commit: dfa6faf.

Linux Distribution

Ubuntu 24.04 LTS

Other Linux Distribution

No response

Kernel Version & Boot Parameters

$ uname -r
6.17.0-42-generic

$ cat /proc/cmdline
BOOT_IMAGE=/boot/vmlinuz-6.17.0-42-generic root=/dev/mapper/vgubuntu-root ro quiet splash vt.handoff=7

$ lsmod | grep -E 'i915|xe'
xe                   3874816  3
intel_vsec             24576  2 pmt_telemetry,xe
gpu_sched              65536  1 xe
drm_gpuvm              49152  1 xe
drm_gpusvm_helper      36864  1 xe
drm_buddy              28672  1 xe
drm_exec               12288  2 drm_gpuvm,xe
drm_suballoc_helper    24576  1 xe
amddrm_exec            12288  1 amdgpu
amdkcl                 32768  4 amd_sched,amdttm,amddrm_exec,amdgpu
drm_display_helper    290816  2 amdgpu,xe
cec                    98304  3 drm_display_helper,amdgpu,xe
video                  77824  3 asus_wmi,amdgpu,xe
i2c_algo_bit           16384  2 amdgpu,xe
drm_ttm_helper         16384  2 amdgpu,xe
ttm                   126976  2 drm_ttm_helper,xe

Actual Behavior

This is my best-guess / understanding of the situation. It may contain errors. I think the IGC libigc.so is the root cause.

Component: Intel Graphics Compiler (IGC), libigc.so.2, invoked from the Level Zero
GPU driver (libze_intel_gpu.so.1) during offline compilation of SPIR-V kernel
images submitted by the SYCL runtime.

When a SYCL kernel image is compiled, IGC reports

`<unknown>:0: error: Undefined temporary symbol .L_<mangled helper name>`

The reported names are exactly the mangled names of helper functions that appear embedded in the kernel entry-point name string: when a kernel lambda is instantiated from a template function that takes a function
pointer as a non-type template parameter (NTTP), the auto-generated kernel
name (the mangled lambda closure type) contains the helper's mangled name as
the XadL_<mangled name>EE NTTP encoding, e.g.

`_ZTSZL11submit_copyIXadL_ZL13cpy_1_t1PKcPcEEEvRN4sycl3_V15queueES2_S2_iEUlNS4_7nd_itemILi3EEEE_`

IGC then creates an undefined .L_-prefixed local symbol from that embedded
name (_ZL13cpy_1_t1PKcPc -> .L_ZL13cpy_1_t1PKcPc) and fails the image with
LLVM's ELF object-writer error "Undefined temporary symbol".

The error is spurious: the helper exists in the module, is referenced by a
normal direct call, and is not imported from anywhere. A/B tests show the
failure follows the name string, not the code: the identical kernel body
compiles cleanly when given an explicit clean kernel name, and compiles
cleanly when the function-pointer NTTP is replaced by an integer tag.

In production (llama.cpp SYCL backend) the same failure makes the cpy kernels
in a multi-kernel image silently never execute (no exception is raised),
corrupting data: attention copies are skipped, logits become garbage, and the
application eventually aborts in an unrelated code path.

• Level Zero driver: libze-intel-gpu1 (26.35.39758.10-0)
• oneAPI: 2026.1, icpx (Intel oneAPI DPC++/C++ Compiler 2026.1.1 (2026.1.1.20260724))
• SYCL runtime: intel-oneapi-runtime-dpcpp-sycl-rt (2026.1.1-325)

$ icpx --version
Intel(R) oneAPI DPC++/C++ Compiler 2026.1.1 (2026.1.1.20260724)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/intel/oneapi/compiler/2026.1/bin/compiler
Configuration file: /opt/intel/oneapi/compiler/2026.1/bin/compiler/../icpx.cfg

Expected Behavior

The two example programs should run without runtime errors.

 $ ONEAPI_DEVICE_SELECTOR=level_zero:gpu SYCL_UR_USE_LEVEL_ZERO_V2=1 SYCL_CACHE_PERSISTENT=0 SYCL_CACHE_IN_MEM=0 ./igc-min-repro
device: Intel(R) Arc(TM) Pro B70 Graphics
RESULT: OK

$ ./cpy-repro
device: Intel(R) Arc(TM) Pro B70 Graphics
dst nonzero bytes: 938 / 16384
RESULT: OK

Reproduction Rate

Always reproduces - 100%

Steps to Reproduce

I've included a minimal example, that unfortunately only seems to crash on the first invocation reliably. After that it tends to work.

Compile & link with

$ icpx -g -Wall -Rno-debug-disables-optimization -std=gnu++17 -fsycl \
   igc-min-repro.cpp   -o igc-min-repro

igc-min-repro.cpp

$ ONEAPI_DEVICE_SELECTOR=level_zero:gpu SYCL_UR_USE_LEVEL_ZERO_V2=1 SYCL_CACHE_PERSISTENT=0 SYCL_CACHE_IN_MEM=0 ./igc-min-repro
device: Intel(R) Arc(TM) Pro B70 Graphics
<unknown>:0: error: Undefined temporary symbol .L_ZL8cpy_1_t1PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL8cpy_1_t2PKcPc
RESULT: OK

Afterwords it will appear to run fine.

$ ONEAPI_DEVICE_SELECTOR=level_zero:gpu SYCL_UR_USE_LEVEL_ZERO_V2=1 SYCL_CACHE_PERSISTENT=0 SYCL_CACHE_IN_MEM=0 ./igc-min-repro
device: Intel(R) Arc(TM) Pro B70 Graphics
RESULT: OK

I've attached a not-so-minimal example that is fairly reliable at exhibiting the erroneous behaviour.

Compile & link with

$ icpx -g -Wall -Rno-debug-disables-optimization -std=gnu++17 -fsycl \
     cpy-repro.cpp -o cpy-repro

cpy-repro.cpp

$ ONEAPI_DEVICE_SELECTOR=level_zero:gpu  SYCL_CACHE_PERSISTENT=0 SYCL_CACHE_IN_MEM=0 ./cpy-repro
device: Intel(R) Arc(TM) Pro B70 Graphics
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f16PKcPc
dst nonzero bytes: 938 / 16384
RESULT: OK

Without modifying caching

$ ONEAPI_DEVICE_SELECTOR=level_zero:gpu ./cpy-repro
device: Intel(R) Arc(TM) Pro B70 Graphics
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f16PKcPc
dst nonzero bytes: 938 / 16384
RESULT: OK

Is this a regression?

• Yes, this is a regression - functionality that previously worked is now broken

Last Known Working Driver Version

No response

First Known Failing Driver Version

No response

API Call Logs

$ ONEAPI_DEVICE_SELECTOR=level_zero:gpu SYCL_CACHE_PERSISTENT=0 SYCL_CACHE_IN_MEM=0 ~/dev/intel-pti/pti-gpu/tools/unitrace/build/unitrace --device-timing --ccl-summary-report --kernel-submission ./cpy-repro >cpy-repro-unittrace-log.txt 2>&1

cpy-repro-unittrace-log.txt

strace Logs

ONEAPI_DEVICE_SELECTOR=level_zero:gpu SYCL_CACHE_PERSISTENT=0 SYCL_CACHE_IN_MEM=0 strace -k -f -o strace-cpy-repro-log.txt ./cpy-repro

stdout:

device: Intel(R) Arc(TM) Pro B70 Graphics
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f16_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f16PKcPc
dst nonzero bytes: 938 / 16384
RESULT: OK

strace-cpy-repro.zip

System Logs / dmesg Output

$ sudo dmesg | grep -i -E 'i915|xe|drm|gpu' | grep -v 'amdgpu' | tail -n 100
[523109.892998] xe 0000:0e:00.0: [drm] GT0: Engine reset: engine_class=bcs, logical_mask: 0x1, guc_id=6
[523237.594197] xe 0000:0e:00.0: [drm] GT0: Engine reset: engine_class=bcs, logical_mask: 0x1, guc_id=6
[523862.552880] xe 0000:0e:00.0: [drm] GT0: Engine reset: engine_class=ccs, logical_mask: 0x1, guc_id=2
[523862.553210] xe 0000:0e:00.0: [drm] GT0: Engine reset: engine_class=bcs, logical_mask: 0x1, guc_id=6
[524108.684622] xe 0000:0e:00.0: [drm] GT0: Engine reset: engine_class=bcs, logical_mask: 0x1, guc_id=6

Backtrace (if crash or hang occurred)

No response

Source Code / Reproducer

igc-min-repro.cpp

cpy-repro.cpp

Command Line / Application Details

No response

oneAPI Version (if applicable)

$ icpx --version
Intel(R) oneAPI DPC++/C++ Compiler 2026.1.1 (2026.1.1.20260724)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/intel/oneapi/compiler/2026.1/bin/compiler
Configuration file: /opt/intel/oneapi/compiler/2026.1/bin/compiler/../icpx.cfg

Screenshots / Video

No response

Additional Notes

No response

[mctylr-gh]

$ ONEAPI_DEVICE_SELECTOR=level_zero:0 SYCL_CACHE_PERSISTENT=0 SYCL_CACHE_IN_MEM=0 SYCL_RT_WARNING_LEVEL=2 SYCL_UR_USE_LEVEL_ZERO_V2=1 SYCL_UR_TRACE=2 SYCL_CACHE_TRACE=7 ./cpy-repro >cpy-repro-env-debug.txt 2>&1

excerpt from logs:

   <--- urDeviceGetInfo(.hDevice = 0xa054520, .propName = UR_DEVICE_INFO_COMPILER_AVAILABLE, .propSize = 1, .pPropValue = 0x7ffc97207650 (true), .pPropSizeRet = nullptr) -> UR_RESULT_SUCCESS;
   ---> urProgramCreateWithIL
   <--- urProgramCreateWithIL(.hContext = 0xa058810, .pIL = 0x9ea3a00, .length = 4048, .pProperties = nullptr, .phProgram = 0x7ffc972078d0 (0xa0ba7b0)) -> UR_RESULT_SUCCESS;
   ---> urProgramCompileExp
   <--- urProgramCompileExp(.hProgram = 0xa0ba290, .numDevices = 1, .phDevices = 0xa0b9be0 {0xa054520}, .flags = 0, .pOptions = 0xa0b9f20 (-g  -ze-opt-disable)) -> UR_RESULT_SUCCESS;
   ---> urProgramCompileExp
   <--- urProgramCompileExp(.hProgram = 0xa0ba7b0, .numDevices = 1, .phDevices = 0xa0b9be0 {0xa054520}, .flags = 0, .pOptions = 0xa0b9f20 (-g  -ze-opt-disable)) -> UR_RESULT_SUCCESS;
   ---> urProgramLinkExp
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f16_f16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i16_i16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_f32_i32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL13cpy_1_i32_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_f32_bf16PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL14cpy_1_bf16_f32PKcPc
<unknown>:0: error: Undefined temporary symbol .L_ZL15cpy_1_bf16_bf16PKcPc
   <--- urProgramLinkExp(.hContext = 0xa058810, .numDevices = 1, .phDevices = 0xa0b9be0 {0xa054520}, .flags = 0, .count = 2, .phPrograms = 0xa0bb9c0 {0xa0ba290, 0xa0ba7b0}, .pOptions = 0x7ffc97207ab0 (), .phProgram = 0x7ffc97207ce0 (0xa0bba70)) -> UR_RESULT_SUCCESS;
   ---> urProgramRelease
   <--- urProgramRelease(.hProgram = 0xa0ba290) -> UR_RESULT_SUCCESS;
   ---> urProgramRelease
   <--- urProgramRelease(.hProgram = 0xa0ba7b0) -> UR_RESULT_SUCCESS;
   ---> urProgramGetInfo
   <--- urProgramGetInfo(.hProgram = 0xa0bba70, .propName = UR_PROGRAM_INFO_DEVICES, .propSize = 0, .pPropValue = nullptr, .pPropSizeRet = 0x7ffc97207628 (8)) -> UR_RESULT_SUCCESS;
   ---> urProgramGetInfo
   <--- urProgramGetInfo(.hProgram = 0xa0bba70, .propName = UR_PROGRAM_INFO_DEVICES, .propSize = 8, .pPropValue = {0xa054520}, .pPropSizeRet = nullptr) -> UR_RESULT_SUCCESS;
   ---> urProgramGetBuildInfo
   <--- urProgramGetBuildInfo(.hProgram = 0xa0bba70, .hDevice = 0xa054520, .propName = UR_PROGRAM_BUILD_INFO_LOG, .propSize = 0, .pPropValue = nullptr, .pPropSizeRet = 0x7ffc97207620 (1402)) -> UR_RESULT_SUCCESS;
   ---> urProgramGetBuildInfo
   <--- urProgramGetBuildInfo(.hProgram = 0xa0bba70, .hDevice = 0xa054520, .propName = UR_PROGRAM_BUILD_INFO_LOG, .propSize = 1402, .pPropValue = 0xa9226c0 (
warning: in kernel 'typeinfo name for void submit_copy<&cpy_1_bf16_bf16(char const*, char*)>(sycl::_V1::queue&, char*, char*, int)::'lambda'(sycl::_V1::nd_item<3>)': Stack call has been detected, the private memory size is set to 1112B. You can change this size by setting environmental variable IGC_ForcePerThreadPrivateMemorySize to a value in range [1024:20480]. Greater values can affect performance, and lower ones may lead to incorrect results of your program.
To make sure your program runs correctly you can use StackOverflowDetection feature. See documentation:
https://github.com/intel/intel-graphics-compiler/tree/master/documentation/igc/StackOverflowDetection/StackOverflowDetection.md


warning: in kernel 'typeinfo name for void submit_copy<&cpy_1_bf16_bf16(char const*, char*)>(sycl::_V1::queue&, char*, char*, int)::'lambda'(sycl::_V1::nd_item<3>)': Stack call has been detected, the private memory size is set to 1112B. You can change this size by setting environmental variable IGC_ForcePerThreadPrivateMemorySize to a value in range [1024:20480]. Greater values can affect performance, and lower ones may lead to incorrect results of your program.
To make sure your program runs correctly you can use StackOverflowDetection feature. See documentation:
https://github.com/intel/intel-graphics-compiler/tree/master/documentation/igc/StackOverflowDetection/StackOverflowDetection.md
), .pPropSizeRet = nullptr) -> UR_RESULT_SUCCESS;

Complete output:

cpy-repro-env-debug.txt