Intel CPU runtime: kernel compiler segfaults in FixedVectorType::get on vectorized kernels

Start by reproducing the supplied Julia kernels with intel_opencl_rt_jll and inspect the dumped .spv files using JULIA_OPENCL_DUMP_DIR. The report points to vectorized kernels containing aggregate SPIR-V values and identifies affected GPUArrays test groups: gpuarrays/linalg/core, gpuarrays/reductions/base, gpuarrays/reductions/contract, and gpuarrays/accumulate. Done means these kernels build and run correctly on the Intel CPU runtime, with the affected tests no longer needing to be skipped.

Building some kernels on Intel's OpenCL CPU runtime (intel_opencl_rt_jll) segfaults the whole process. The runtime's embedded LLVM crashes in llvm::FixedVectorType::get while handling the vectorized copy of the kernel (the _ZGVeN32... function the runtime's vectorizer creates). Depending on the kernel, the crash happens in instruction selection, loop-unroll or simplifycfg, but the top frames are the same: something asks for a vector of a type that can't be vectorized.

MWE

using intel_opencl_rt_jll, OpenCL

function kernel(a)
    i = get_global_id()
    @inbounds a[i] = a[i] / (a[i] + 1)
    return
end

a = CLArray(ComplexF32[1, 2, 3])
@opencl global_size=3 kernel(a)
println(Array(a))

PLEASE submit a bug report to https://software.intel.com/en-us/support/priority-support and include the crash backtrace and instructions to reproduce the bug.
Stack dump:
0.	Running pass 'Function Pass Manager' on module 'main'.
1.	Running pass 'X86 DAG->DAG Instruction Selection' on function ''

[1377712] signal 11 (1): Segmentation fault
_ZN4llvm15FixedVectorType3getEPNS_4TypeEj at libintelocl.so
_ZN4llvm3EVT19getExtendedVectorVTERNS_11LLVMContextES0_NS_12ElementCountE at libintelocl.so
_ZN4llvm19SelectionDAGBuilder12getValueImplEPKNS_5ValueE at libintelocl.so
_ZN4llvm19SelectionDAGBuilder8getValueEPKNS_5ValueE at libintelocl.so
_ZN4llvm19SelectionDAGBuilder18visitInsertElementERKNS_4UserE at libintelocl.so
...

Putting the division in a loop (for j in 1:n; a[i] = a[i] / (a[i] + 1); end) moves the crash to loop-unroll, as in gpu_trimatdiv (the triangular solve in GPUArrays):

1.	Running pass "loop-unroll<O3>" on function "_ZGVeN32uuu__Z1k13CLDeviceArrayI7ComplexI7Float32ELi1ELi1EE5Int64"
_ZN4llvm15FixedVectorType3getEPNS_4TypeEj at libintelocl.so
_ZN4llvm3EVT19getExtendedVectorVTERNS_11LLVMContextES0_NS_12ElementCountE at libintelocl.so
_ZNK4llvm16BasicTTIImplBaseINS_10X86TTIImplEE23getTypeLegalizationCostEPNS_4TypeE at libintelocl.so
_ZNK4llvm10X86TTIImpl18getVectorInstrCostEjPNS_4TypeENS_19TargetTransformInfo14TargetCostKindEjPKNS_5ValueES7_ at libintelocl.so
_ZNK4llvm31TargetTransformInfoImplCRTPBaseINS_10X86TTIImplEE18getInstructionCostEPKNS_4UserENS_8ArrayRefIPKNS_5ValueEEENS_19TargetTransformInfo14TargetCostKindE at libintelocl.so
_ZN4llvm11CodeMetrics17analyzeBasicBlockEPKNS_10BasicBlockERKNS_19TargetTransformInfoERKNS_15SmallPtrSetImplIPKNS_5ValueEEEbPKNS_4LoopE at libintelocl.so
_ZN4llvm19UnrollCostEstimatorC1EPKNS_4LoopERKNS_19TargetTransformInfoERKNS_15SmallPtrSetImplIPKNS_5ValueEEEj at libintelocl.so
...

The same crash, this time in simplifycfg, happens with GPUArrays reductions and scans whose operator has no known neutral element:

using intel_opencl_rt_jll, OpenCL
reduce((a, b) -> a + b, CLArray(Float32[1, 2, 3]))

1.	Running pass "simplifycfg<...>" on function "_ZGVeN32uuuuuu__Z21gpu__mapreduce_block_16CompilerMetadataI11DynamicSize12DynamicCheckv16CartesianIndicesILi1E5TupleI5OneToI5Int64EEE7NDRangeILi1ES0_10StaticSizeI6_256__ES8_vEE13CLDeviceArrayI7Float32Li1ELi1EESE_I5_LaneISF_ELi1ELi1EE8identity2_2SI_9_NoFinishS5_3ValILi256EESN_ILi16EE"
_ZN4llvm15FixedVectorType3getEPNS_4TypeEj at libintelocl.so
_ZN4llvm3EVT19getExtendedVectorVTERNS_11LLVMContextES0_NS_12ElementCountE at libintelocl.so
_ZNK4llvm16BasicTTIImplBaseINS_10X86TTIImplEE23getTypeLegalizationCostEPNS_4TypeE at libintelocl.so
_ZNK4llvm10X86TTIImpl18getVectorInstrCostEjPNS_4TypeENS_19TargetTransformInfo14TargetCostKindEjPKNS_5ValueES7_ at libintelocl.so
_ZNK4llvm31TargetTransformInfoImplCRTPBaseINS_10X86TTIImplEE18getInstructionCostEPKNS_4UserENS_8ArrayRefIPKNS_5ValueEEENS_19TargetTransformInfo14TargetCostKindE at libintelocl.so
_ZL19dominatesMergePointPN4llvm5ValueEPNS_10BasicBlockEPNS_11InstructionERNS_15SmallPtrSetImplIS5_EERNS_15InstructionCostES9_RKNS_19TargetTransformInfoEPNS_15AssumptionCacheES8_j at libintelocl.so
_ZN4llvm11simplifyCFGEPNS_10BasicBlockERKNS_19TargetTransformInfoEPNS_14DomTreeUpdaterERKNS_18SimplifyCFGOptionsENS_8ArrayRefINS_6WeakVHEEEPNS_17AndersensAAResultE at libintelocl.so
...

Other one-liners that crash the same way: accumulate((a, b) -> a + b, CLArray(rand(1:10, 100))), mapreduce(abs, (a, b) -> max(a, b), CLArray(rand(Float32, 100))), and UpperTriangular(CLArray(A)) \ CLArray(b) with ComplexF32 inputs. reduce(+, ...) and sum are fine, as are the Float32 and ComplexF64 versions of the kernel above.

What the kernels have in common

All of them handle aggregate-typed SPIR-V values inside the kernel: the ComplexF32 division is a call to a (non-inlined) function returning [2 x float], and the reductions without a neutral element load and store {float, i8} structs (AcceleratedKernels' _Lane) in local memory. My guess is that the vectorizer widens one of these aggregates into a vector type, which LLVM can't represent. To look at the SPIR-V we pass to the runtime, run the MWE with JULIA_OPENCL_DUMP_DIR=<dir>: OpenCL.jl writes the program to a .spv file there before handing it to the driver, so it's written even though the process crashes.

Workarounds

• CL_CONFIG_CPU_TARGET_ARCH=core-avx2 doesn't help (the crash moves to the _ZGVdN32... function).
• Requesting a different sub-group size (8 or 16, or none) doesn't help either: the kernel is still vectorized and still crashes.
• CL_CONFIG_USE_VECTORIZER=False on its own makes every kernel fail to build with Conflicting CL_CONFIG_CPU_VECTORIZER_MODE, intel_vec_len_hint and intel_reqd_sub_group_size are specified, because OpenCL.jl requests a sub-group size by default (intel_reqd_sub_group_size). Combined with not requesting one (@eval OpenCL _sub_group_size(dev) = nothing), all of the MWEs above build and give the right results.

PoCL (pocl_jll) runs all of these correctly.

Tests

Because of this, the tb/intel-cpu-driver branch (adding the Intel CPU runtime to CI, upcoming PR) skips these GPUArrays tests on it: gpuarrays/linalg/core (gpu_trimatdiv), gpuarrays/reductions/base, gpuarrays/reductions/contract and gpuarrays/accumulate. Each of them crashed with this signature at the first such kernel, so other problems later in those files would be hidden.

Environment

• OpenCL.jl main (7b8d2c4), GPUArrays 12.1.0, AcceleratedKernels 0.5.4, KernelAbstractions 0.9.44
• intel_opencl_rt_jll 2026.1.2 (driver version 2026.21.7.0.24_160000)
• AMD Ryzen 9 9950X (AVX-512), Linux
• Julia 1.13.1