Kernel verification failure during execution of axpby with complex float inputs

[maleadt]

Summary

When executing oneapi::mkl::blas::column_major::axpby with std::complex<float> inputs under ZE_ENABLE_VALIDATION_LAYER=1, the kernel generated by MKL fails verification.

Version

MKL from oneAPI 2024.1.0 as downloaded from Conda:

_libgcc_mutex             0.1                 conda_forge    intel
_openmp_mutex             4.5                       2_gnu    intel
_sysroot_linux-64_curr_repodata_hack 3                   h69a702a_14    conda-forge
dpcpp-cpp-rt              2024.1.0              intel_963    intel
dpcpp_impl_linux-64       2024.1.0              intel_963    intel
dpcpp_linux-64            2024.1.0              intel_963    intel
icc_rt                    2024.1.0              intel_963    intel
intel-cmplr-lib-rt        2024.1.0              intel_963    intel
intel-cmplr-lic-rt        2024.1.0              intel_963    intel
intel-opencl-rt           2024.1.0              intel_963    intel
intel-openmp              2024.1.0              intel_963    intel
kernel-headers_linux-64   4.18.0               he073ed8_2    conda-forge
libgcc-ng                 13.2.0               h807b86a_5    intel
libgomp                   13.2.0               h807b86a_5    intel
libstdcxx-ng              13.2.0               h7e041cc_5    intel
mkl                       2024.1.0              intel_691    intel
mkl-devel                 2024.1.0              intel_691    intel
mkl-devel-dpcpp           2024.1.0              intel_691    intel
mkl-dpcpp                 2024.1.0              intel_691    intel
mkl-include               2024.1.0              intel_691    intel
onemkl-sycl-blas          2024.1.0              intel_691    intel
onemkl-sycl-datafitting   2024.1.0              intel_691    intel
onemkl-sycl-dft           2024.1.0              intel_691    intel
onemkl-sycl-lapack        2024.1.0              intel_691    intel
onemkl-sycl-rng           2024.1.0              intel_691    intel
onemkl-sycl-sparse        2024.1.0              intel_691    intel
onemkl-sycl-stats         2024.1.0              intel_691    intel
onemkl-sycl-vm            2024.1.0              intel_691    intel
sysroot_linux-64          2.28                 he073ed8_2    conda-forge
tbb                       2021.12.0             intel_495    intel

This is a regression, and AFAICT not present on 2024.0.0

Environment

❯ /home/tim/Julia/depot/scratchspaces/8f75cd03-7ff8-4ecb-9b8f-daf728133b1b/conda/bin/icpx --version
Intel(R) oneAPI DPC++/C++ Compiler 2024.1.0 (2024.1.0.20240308)

NEO v24.9.28717 with IGC v1.0.16238

`clinfo`

Number of platforms                               1
  Platform Name                                   Intel(R) OpenCL Graphics
  Platform Vendor                                 Intel(R) Corporation
  Platform Version                                OpenCL 3.0
  Platform Profile                                FULL_PROFILE
  Platform Extensions                             cl_khr_byte_addressable_store cl_khr_device_uuid cl_khr_fp16 cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_icd cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_intel_command_queue_families cl_intel_subgroups cl_intel_required_subgroup_size cl_intel_subgroups_short cl_khr_spir cl_intel_accelerator cl_intel_driver_diagnostics cl_khr_priority_hints cl_khr_throttle_hints cl_khr_create_command_queue cl_intel_subgroups_char cl_intel_subgroups_long cl_khr_il_program cl_intel_mem_force_host_memory cl_khr_subgroup_extended_types cl_khr_subgroup_non_uniform_vote cl_khr_subgroup_ballot cl_khr_subgroup_non_uniform_arithmetic cl_khr_subgroup_shuffle cl_khr_subgroup_shuffle_relative cl_khr_subgroup_clustered_reduce cl_intel_device_attribute_query cl_khr_suggested_local_work_size cl_intel_split_work_group_barrier cl_intel_spirv_media_block_io cl_intel_spirv_subgroups cl_khr_spirv_linkonce_odr cl_khr_spirv_no_integer_wrap_decoration cl_intel_unified_shared_memory cl_khr_mipmap_image cl_khr_mipmap_image_writes cl_ext_float_atomics cl_khr_external_memory cl_intel_planar_yuv cl_intel_packed_yuv cl_khr_int64_base_atomics cl_khr_int64_extended_atomics cl_khr_image2d_from_buffer cl_khr_depth_images cl_khr_3d_image_writes cl_intel_media_block_io cl_intel_subgroup_local_block_io cl_khr_integer_dot_product cl_khr_gl_sharing cl_khr_gl_depth_images cl_khr_gl_event cl_khr_gl_msaa_sharing cl_intel_sharing_format_query cl_khr_pci_bus_info
  Platform Extensions with Version                cl_khr_byte_addressable_store                                    0x400000 (1.0.0)
                                                  cl_khr_device_uuid                                               0x400000 (1.0.0)
                                                  cl_khr_fp16                                                      0x400000 (1.0.0)
                                                  cl_khr_global_int32_base_atomics                                 0x400000 (1.0.0)
                                                  cl_khr_global_int32_extended_atomics                             0x400000 (1.0.0)
                                                  cl_khr_icd                                                       0x400000 (1.0.0)
                                                  cl_khr_local_int32_base_atomics                                  0x400000 (1.0.0)
                                                  cl_khr_local_int32_extended_atomics                              0x400000 (1.0.0)
                                                  cl_intel_command_queue_families                                  0x400000 (1.0.0)
                                                  cl_intel_subgroups                                               0x400000 (1.0.0)
                                                  cl_intel_required_subgroup_size                                  0x400000 (1.0.0)
                                                  cl_intel_subgroups_short                                         0x400000 (1.0.0)
                                                  cl_khr_spir                                                      0x400000 (1.0.0)
                                                  cl_intel_accelerator                                             0x400000 (1.0.0)
                                                  cl_intel_driver_diagnostics                                      0x400000 (1.0.0)
                                                  cl_khr_priority_hints                                            0x400000 (1.0.0)
                                                  cl_khr_throttle_hints                                            0x400000 (1.0.0)
                                                  cl_khr_create_command_queue                                      0x400000 (1.0.0)
                                                  cl_intel_subgroups_char                                          0x400000 (1.0.0)
                                                  cl_intel_subgroups_long                                          0x400000 (1.0.0)
                                                  cl_khr_il_program                                                0x400000 (1.0.0)
                                                  cl_intel_mem_force_host_memory                                   0x400000 (1.0.0)
                                                  cl_khr_subgroup_extended_types                                   0x400000 (1.0.0)
                                                  cl_khr_subgroup_non_uniform_vote                                 0x400000 (1.0.0)
                                                  cl_khr_subgroup_ballot                                           0x400000 (1.0.0)
                                                  cl_khr_subgroup_non_uniform_arithmetic                           0x400000 (1.0.0)
                                                  cl_khr_subgroup_shuffle                                          0x400000 (1.0.0)
                                                  cl_khr_subgroup_shuffle_relative                                 0x400000 (1.0.0)
                                                  cl_khr_subgroup_clustered_reduce                                 0x400000 (1.0.0)
                                                  cl_intel_device_attribute_query                                  0x400000 (1.0.0)
                                                  cl_khr_suggested_local_work_size                                 0x400000 (1.0.0)
                                                  cl_intel_split_work_group_barrier                                0x400000 (1.0.0)
                                                  cl_intel_spirv_media_block_io                                    0x400000 (1.0.0)
                                                  cl_intel_spirv_subgroups                                         0x400000 (1.0.0)
                                                  cl_khr_spirv_linkonce_odr                                        0x400000 (1.0.0)
                                                  cl_khr_spirv_no_integer_wrap_decoration                          0x400000 (1.0.0)
                                                  cl_intel_unified_shared_memory                                   0x400000 (1.0.0)
                                                  cl_khr_mipmap_image                                              0x400000 (1.0.0)
                                                  cl_khr_mipmap_image_writes                                       0x400000 (1.0.0)
                                                  cl_ext_float_atomics                                             0x400000 (1.0.0)
                                                  cl_khr_external_memory                                             0x9001 (0.9.1)
                                                  cl_intel_planar_yuv                                              0x400000 (1.0.0)
                                                  cl_intel_packed_yuv                                              0x400000 (1.0.0)
                                                  cl_khr_int64_base_atomics                                        0x400000 (1.0.0)
                                                  cl_khr_int64_extended_atomics                                    0x400000 (1.0.0)
                                                  cl_khr_image2d_from_buffer                                       0x400000 (1.0.0)
                                                  cl_khr_depth_images                                              0x400000 (1.0.0)
                                                  cl_khr_3d_image_writes                                           0x400000 (1.0.0)
                                                  cl_intel_media_block_io                                          0x400000 (1.0.0)
                                                  cl_intel_subgroup_local_block_io                                 0x400000 (1.0.0)
                                                  cl_khr_integer_dot_product                                       0x800000 (2.0.0)
                                                  cl_khr_gl_sharing                                                0x400000 (1.0.0)
                                                  cl_khr_gl_depth_images                                           0x400000 (1.0.0)
                                                  cl_khr_gl_event                                                  0x400000 (1.0.0)
                                                  cl_khr_gl_msaa_sharing                                           0x400000 (1.0.0)
                                                  cl_intel_sharing_format_query                                    0x400000 (1.0.0)
                                                  cl_khr_pci_bus_info                                              0x400000 (1.0.0)
  Platform Numeric Version                        0xc00000 (3.0.0)
  Platform Extensions function suffix             INTEL
  Platform Host timer resolution                  1ns
  Platform External memory handle types           DMA buffer

  Platform Name                                   Intel(R) OpenCL Graphics
Number of devices                                 1
  Device Name                                     Intel(R) Iris(R) Xe Graphics
  Device Vendor                                   Intel(R) Corporation
  Device Vendor ID                                0x8086
  Device Version                                  OpenCL 3.0 NEO
  Device UUID                                     8680499a-0100-0000-0002-000000000000
  Driver UUID                                     32342e30-392e-3032-3837-313700000000
  Valid Device LUID                               No
  Device LUID                                     e0c9-b644fc7f0000
  Device Node Mask                                0
  Device Numeric Version                          0xc00000 (3.0.0)
  Driver Version                                  24.09.028717
  Device OpenCL C Version                         OpenCL C 1.2
  Device OpenCL C all versions                    OpenCL C                                                         0x400000 (1.0.0)
                                                  OpenCL C                                                         0x401000 (1.1.0)
                                                  OpenCL C                                                         0x402000 (1.2.0)
                                                  OpenCL C                                                         0xc00000 (3.0.0)
  Device OpenCL C features                        __opencl_c_int64                                                 0xc00000 (3.0.0)
                                                  __opencl_c_3d_image_writes                                       0xc00000 (3.0.0)
                                                  __opencl_c_images                                                0xc00000 (3.0.0)
                                                  __opencl_c_read_write_images                                     0xc00000 (3.0.0)
                                                  __opencl_c_atomic_order_acq_rel                                  0xc00000 (3.0.0)
                                                  __opencl_c_atomic_order_seq_cst                                  0xc00000 (3.0.0)
                                                  __opencl_c_atomic_scope_all_devices                              0xc00000 (3.0.0)
                                                  __opencl_c_atomic_scope_device                                   0xc00000 (3.0.0)
                                                  __opencl_c_generic_address_space                                 0xc00000 (3.0.0)
                                                  __opencl_c_program_scope_global_variables                        0xc00000 (3.0.0)
                                                  __opencl_c_work_group_collective_functions                       0xc00000 (3.0.0)
                                                  __opencl_c_subgroups                                             0xc00000 (3.0.0)
                                                  __opencl_c_ext_fp32_global_atomic_add                            0xc00000 (3.0.0)
                                                  __opencl_c_ext_fp32_local_atomic_add                             0xc00000 (3.0.0)
                                                  __opencl_c_ext_fp32_global_atomic_min_max                        0xc00000 (3.0.0)
                                                  __opencl_c_ext_fp32_local_atomic_min_max                         0xc00000 (3.0.0)
                                                  __opencl_c_ext_fp16_global_atomic_load_store                     0xc00000 (3.0.0)
                                                  __opencl_c_ext_fp16_local_atomic_load_store                      0xc00000 (3.0.0)
                                                  __opencl_c_ext_fp16_global_atomic_min_max                        0xc00000 (3.0.0)
                                                  __opencl_c_ext_fp16_local_atomic_min_max                         0xc00000 (3.0.0)
                                                  __opencl_c_integer_dot_product_input_4x8bit                      0xc00000 (3.0.0)
                                                  __opencl_c_integer_dot_product_input_4x8bit_packed               0xc00000 (3.0.0)
  Latest conformance test passed                  v2023-05-16-00
  Device Type                                     GPU
  Device PCI bus info (KHR)                       PCI-E, 0000:00:02.0
  Device Profile                                  FULL_PROFILE
  Device Available                                Yes
  Compiler Available                              Yes
  Linker Available                                Yes
  Max compute units                               96
  Max clock frequency                             1300MHz
  Device IP (Intel)                               0x3000000 (12.0.0)
  Device ID (Intel)                               39497
  Slices (Intel)                                  1
  Sub-slices per slice (Intel)                    6
  EUs per sub-slice (Intel)                       16
  Threads per EU (Intel)                          7
  Feature capabilities (Intel)                    DP4A
  Device Partition                                (core)
    Max number of sub-devices                     0
    Supported partition types                     None
    Supported affinity domains                    (n/a)
  Max work item dimensions                        3
  Max work item sizes                             512x512x512
  Max work group size                             512
  Preferred work group size multiple (device)     64
  Preferred work group size multiple (kernel)     64
  Max sub-groups per work group                   64
  Sub-group sizes (Intel)                         8, 16, 32
  Preferred / native vector sizes
    char                                                16 / 16
    short                                                8 / 8
    int                                                  4 / 4
    long                                                 1 / 1
    half                                                 8 / 8        (cl_khr_fp16)
    float                                                1 / 1
    double                                               0 / 0        (n/a)
  Half-precision Floating-point support           (cl_khr_fp16)
    Denormals                                     Yes
    Infinity and NANs                             Yes
    Round to nearest                              Yes
    Round to zero                                 Yes
    Round to infinity                             Yes
    IEEE754-2008 fused multiply-add               Yes
    Support is emulated in software               No
  Single-precision Floating-point support         (core)
    Denormals                                     Yes
    Infinity and NANs                             Yes
    Round to nearest                              Yes
    Round to zero                                 Yes
    Round to infinity                             Yes
    IEEE754-2008 fused multiply-add               Yes
    Support is emulated in software               No
    Correctly-rounded divide and sqrt operations  No
  Double-precision Floating-point support         (n/a)
  Address bits                                    64, Little-Endian
  External memory handle types                    DMA buffer
  Global memory size                              30668070912 (28.56GiB)
  Error Correction support                        No
  Max memory allocation                           4294959104 (4GiB)
  Unified memory for Host and Device              Yes
  Shared Virtual Memory (SVM) capabilities        (core)
    Coarse-grained buffer sharing                 Yes
    Fine-grained buffer sharing                   No
    Fine-grained system sharing                   No
    Atomics                                       No
  Unified Shared Memory (USM)                     (cl_intel_unified_shared_memory)
  Host USM capabilities (Intel)                   USM access, USM atomic access
  Device USM capabilities (Intel)                 USM access, USM atomic access
  Single-Device USM caps (Intel)                  USM access, USM atomic access
  Cross-Device USM caps (Intel)                   (n/a)
  Shared System USM caps (Intel)                  (n/a)
  Minimum alignment for any data type             128 bytes
  Alignment of base address                       1024 bits (128 bytes)
  Preferred alignment for atomics
    SVM                                           64 bytes
    Global                                        64 bytes
    Local                                         64 bytes
  Atomic memory capabilities                      relaxed, acquire/release, sequentially-consistent, work-group scope, device scope, all-devices scope
  Atomic fence capabilities                       relaxed, acquire/release, sequentially-consistent, work-item scope, work-group scope, device scope, all-devices scope
  Max size for global variable                    65536 (64KiB)
  Preferred total size of global vars             4294959104 (4GiB)
  Global Memory cache type                        Read/Write
  Global Memory cache size                        3932160 (3.75MiB)
  Global Memory cache line size                   64 bytes
  Image support                                   Yes
    Max number of samplers per kernel             16
    Max size for 1D images from buffer            268434944 pixels
    Max 1D or 2D image array size                 2048 images
    Base address alignment for 2D image buffers   4 bytes
    Pitch alignment for 2D image buffers          4 pixels
    Max 2D image size                             16384x16384 pixels
    Max planar YUV image size                     16384x16352 pixels
    Max 3D image size                             2048x2048x2048 pixels
    Max number of read image args                 128
    Max number of write image args                128
    Max number of read/write image args           128
  Pipe support                                    No
  Max number of pipe args                         0
  Max active pipe reservations                    0
  Max pipe packet size                            0
  Local memory type                               Local
  Local memory size                               65536 (64KiB)
  Max number of constant args                     8
  Max constant buffer size                        4294959104 (4GiB)
  Generic address space support                   Yes
  Max size of kernel argument                     2048 (2KiB)
  Queue properties (on host)
    Out-of-order execution                        Yes
    Profiling                                     Yes
  Device enqueue capabilities                     (n/a)
  Queue properties (on device)
    Out-of-order execution                        No
    Profiling                                     No
    Preferred size                                0
    Max size                                      0
  Max queues on device                            0
  Max events on device                            0
  Device queue families                           rcs                                                              (1)
                                                                                                 Queue properties  Out-of-order execution, Profiling
                                                                                                     Capabilities  create single-queue events, create cross-queue events
  Prefer user sync for interop                    Yes
  Profiling timer resolution                      52ns
  Execution capabilities
    Run OpenCL kernels                            Yes
    Run native kernels                            No
    Non-uniform work-groups                       Yes
    Work-group collective functions               Yes
    Sub-group independent forward progress        No
    IL version                                    SPIR-V_1.2
    ILs with version                              SPIR-V                                                           0x402000 (1.2.0)
    SPIR versions                                 1.2
  printf() buffer size                            4194304 (4MiB)
  Built-in kernels                                (n/a)
  Built-in kernels with version                   (n/a)
  Device Extensions                               cl_khr_byte_addressable_store cl_khr_device_uuid cl_khr_fp16 cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_icd cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_intel_command_queue_families cl_intel_subgroups cl_intel_required_subgroup_size cl_intel_subgroups_short cl_khr_spir cl_intel_accelerator cl_intel_driver_diagnostics cl_khr_priority_hints cl_khr_throttle_hints cl_khr_create_command_queue cl_intel_subgroups_char cl_intel_subgroups_long cl_khr_il_program cl_intel_mem_force_host_memory cl_khr_subgroup_extended_types cl_khr_subgroup_non_uniform_vote cl_khr_subgroup_ballot cl_khr_subgroup_non_uniform_arithmetic cl_khr_subgroup_shuffle cl_khr_subgroup_shuffle_relative cl_khr_subgroup_clustered_reduce cl_intel_device_attribute_query cl_khr_suggested_local_work_size cl_intel_split_work_group_barrier cl_intel_spirv_media_block_io cl_intel_spirv_subgroups cl_khr_spirv_linkonce_odr cl_khr_spirv_no_integer_wrap_decoration cl_intel_unified_shared_memory cl_khr_mipmap_image cl_khr_mipmap_image_writes cl_ext_float_atomics cl_khr_external_memory cl_intel_planar_yuv cl_intel_packed_yuv cl_khr_int64_base_atomics cl_khr_int64_extended_atomics cl_khr_image2d_from_buffer cl_khr_depth_images cl_khr_3d_image_writes cl_intel_media_block_io cl_intel_subgroup_local_block_io cl_khr_integer_dot_product cl_khr_gl_sharing cl_khr_gl_depth_images cl_khr_gl_event cl_khr_gl_msaa_sharing cl_intel_sharing_format_query cl_khr_pci_bus_info
  Device Extensions with Version                  cl_khr_byte_addressable_store                                    0x400000 (1.0.0)
                                                  cl_khr_device_uuid                                               0x400000 (1.0.0)
                                                  cl_khr_fp16                                                      0x400000 (1.0.0)
                                                  cl_khr_global_int32_base_atomics                                 0x400000 (1.0.0)
                                                  cl_khr_global_int32_extended_atomics                             0x400000 (1.0.0)
                                                  cl_khr_icd                                                       0x400000 (1.0.0)
                                                  cl_khr_local_int32_base_atomics                                  0x400000 (1.0.0)
                                                  cl_khr_local_int32_extended_atomics                              0x400000 (1.0.0)
                                                  cl_intel_command_queue_families                                  0x400000 (1.0.0)
                                                  cl_intel_subgroups                                               0x400000 (1.0.0)
                                                  cl_intel_required_subgroup_size                                  0x400000 (1.0.0)
                                                  cl_intel_subgroups_short                                         0x400000 (1.0.0)
                                                  cl_khr_spir                                                      0x400000 (1.0.0)
                                                  cl_intel_accelerator                                             0x400000 (1.0.0)
                                                  cl_intel_driver_diagnostics                                      0x400000 (1.0.0)
                                                  cl_khr_priority_hints                                            0x400000 (1.0.0)
                                                  cl_khr_throttle_hints                                            0x400000 (1.0.0)
                                                  cl_khr_create_command_queue                                      0x400000 (1.0.0)
                                                  cl_intel_subgroups_char                                          0x400000 (1.0.0)
                                                  cl_intel_subgroups_long                                          0x400000 (1.0.0)
                                                  cl_khr_il_program                                                0x400000 (1.0.0)
                                                  cl_intel_mem_force_host_memory                                   0x400000 (1.0.0)
                                                  cl_khr_subgroup_extended_types                                   0x400000 (1.0.0)
                                                  cl_khr_subgroup_non_uniform_vote                                 0x400000 (1.0.0)
                                                  cl_khr_subgroup_ballot                                           0x400000 (1.0.0)
                                                  cl_khr_subgroup_non_uniform_arithmetic                           0x400000 (1.0.0)
                                                  cl_khr_subgroup_shuffle                                          0x400000 (1.0.0)
                                                  cl_khr_subgroup_shuffle_relative                                 0x400000 (1.0.0)
                                                  cl_khr_subgroup_clustered_reduce                                 0x400000 (1.0.0)
                                                  cl_intel_device_attribute_query                                  0x400000 (1.0.0)
                                                  cl_khr_suggested_local_work_size                                 0x400000 (1.0.0)
                                                  cl_intel_split_work_group_barrier                                0x400000 (1.0.0)
                                                  cl_intel_spirv_media_block_io                                    0x400000 (1.0.0)
                                                  cl_intel_spirv_subgroups                                         0x400000 (1.0.0)
                                                  cl_khr_spirv_linkonce_odr                                        0x400000 (1.0.0)
                                                  cl_khr_spirv_no_integer_wrap_decoration                          0x400000 (1.0.0)
                                                  cl_intel_unified_shared_memory                                   0x400000 (1.0.0)
                                                  cl_khr_mipmap_image                                              0x400000 (1.0.0)
                                                  cl_khr_mipmap_image_writes                                       0x400000 (1.0.0)
                                                  cl_ext_float_atomics                                             0x400000 (1.0.0)
                                                  cl_khr_external_memory                                             0x9001 (0.9.1)
                                                  cl_intel_planar_yuv                                              0x400000 (1.0.0)
                                                  cl_intel_packed_yuv                                              0x400000 (1.0.0)
                                                  cl_khr_int64_base_atomics                                        0x400000 (1.0.0)
                                                  cl_khr_int64_extended_atomics                                    0x400000 (1.0.0)
                                                  cl_khr_image2d_from_buffer                                       0x400000 (1.0.0)
                                                  cl_khr_depth_images                                              0x400000 (1.0.0)
                                                  cl_khr_3d_image_writes                                           0x400000 (1.0.0)
                                                  cl_intel_media_block_io                                          0x400000 (1.0.0)
                                                  cl_intel_subgroup_local_block_io                                 0x400000 (1.0.0)
                                                  cl_khr_integer_dot_product                                       0x800000 (2.0.0)
                                                  cl_khr_gl_sharing                                                0x400000 (1.0.0)
                                                  cl_khr_gl_depth_images                                           0x400000 (1.0.0)
                                                  cl_khr_gl_event                                                  0x400000 (1.0.0)
                                                  cl_khr_gl_msaa_sharing                                           0x400000 (1.0.0)
                                                  cl_intel_sharing_format_query                                    0x400000 (1.0.0)
                                                  cl_khr_pci_bus_info                                              0x400000 (1.0.0)

NULL platform behavior
  clGetPlatformInfo(NULL, CL_PLATFORM_NAME, ...)  No platform
  clGetDeviceIDs(NULL, CL_DEVICE_TYPE_ALL, ...)   No platform
  clCreateContext(NULL, ...) [default]            No platform
  clCreateContext(NULL, ...) [other]              Success [INTEL]
  clCreateContextFromType(NULL, CL_DEVICE_TYPE_DEFAULT)  Success (1)
    Platform Name                                 Intel(R) OpenCL Graphics
    Device Name                                   Intel(R) Iris(R) Xe Graphics
  clCreateContextFromType(NULL, CL_DEVICE_TYPE_CPU)  No devices found in platform
  clCreateContextFromType(NULL, CL_DEVICE_TYPE_GPU)  Success (1)
    Platform Name                                 Intel(R) OpenCL Graphics
    Device Name                                   Intel(R) Iris(R) Xe Graphics
  clCreateContextFromType(NULL, CL_DEVICE_TYPE_ACCELERATOR)  No devices found in platform
  clCreateContextFromType(NULL, CL_DEVICE_TYPE_CUSTOM)  No devices found in platform
  clCreateContextFromType(NULL, CL_DEVICE_TYPE_ALL)  Success (1)
    Platform Name                                 Intel(R) OpenCL Graphics
    Device Name                                   Intel(R) Iris(R) Xe Graphics

ICD loader properties
  ICD loader Name                                 Khronos OpenCL ICD Loader
  ICD loader Vendor                               Khronos Group
  ICD loader Version                              3.0.6
  ICD loader Profile                              OpenCL 3.0

Steps to reproduce

C++ reproducer:

#include <CL/sycl.hpp>
#include <oneapi/mkl.hpp>
#include <iostream>
#include <cstdlib>
#include <ctime>

int main() {
    // Initialize SYCL queue on default device
    sycl::device device((sycl::gpu_selector()));
    sycl::queue queue(device);

    std::complex<float> alpha;
    std::complex<float> beta;

    std::complex<float> *x = sycl::malloc_device<std::complex<float>>(10 * sizeof(std::complex<float>), queue);
    std::complex<float> *y = sycl::malloc_device<std::complex<float>>(10 * sizeof(std::complex<float>), queue);

    // Execute axpby
    oneapi::mkl::blas::column_major::axpby(queue, 10, alpha, x, 1, beta, y, 1);

    return 0;
}

Observed behavior

❯ /conda/bin/icpx wip2.cpp -isystem /conda/include -L /conda/lib -lmkl_sycl -lmkl_intel_ilp64 -lmkl_sequential -lmkl_core -lsycl -lOpenCL -lpi_level_zero -lze_loader -o mwe

❯ ZE_ENABLE_VALIDATION_LAYER=1 LD_LIBRARY_PATH=/conda/lib ./mwe
internal error, assertion failed: G4Verification failure
file: /workspace/srcdir/intel-graphics-compiler/visa/G4_Verifier.cpp
function name: assertIfEnable
line: 197

Full symbolized backtrace:

abort at /lib/x86_64-linux-gnu/libc.so.6 (unknown line)
assert_and_exit at /workspace/srcdir/intel-graphics-compiler/visa/Assertions.cpp:42
assertIfEnable at /workspace/srcdir/intel-graphics-compiler/visa/G4_Verifier.cpp:197
verifyDefUseChain at /workspace/srcdir/intel-graphics-compiler/visa/G4_Verifier.cpp:151
verifyInst at /workspace/srcdir/intel-graphics-compiler/visa/G4_Verifier.cpp:99
verify at /workspace/srcdir/intel-graphics-compiler/visa/G4_Verifier.cpp:62
verifyG4Kernel at /workspace/srcdir/intel-graphics-compiler/visa/G4_Verifier.cpp:19
runPass at /workspace/srcdir/intel-graphics-compiler/visa/Optimizer.cpp:630
optimization at /workspace/srcdir/intel-graphics-compiler/visa/Optimizer.cpp:887
compileTillOptimize at /workspace/srcdir/intel-graphics-compiler/visa/VISAKernelImpl.cpp:230
compileFastPath at /workspace/srcdir/intel-graphics-compiler/visa/VISAKernelImpl.cpp:150
Compile at /workspace/srcdir/intel-graphics-compiler/visa/BuildCISAIRImpl.cpp:1752
Compile at /workspace/srcdir/intel-graphics-compiler/IGC/Compiler/CISACodeGen/CISABuilder.cpp:6300
runOnFunction at /workspace/srcdir/intel-graphics-compiler/IGC/Compiler/CISACodeGen/EmitVISAPass.cpp:1308
runOnFunction at /workspace/srcdir/intel-graphics-compiler/build/IGC/llvm-deps/src/llvm/lib/IR/LegacyPassManager.cpp:1434
runOnModule at /workspace/srcdir/intel-graphics-compiler/build/IGC/llvm-deps/src/llvm/lib/IR/LegacyPassManager.cpp:1480
runOnModule at /workspace/srcdir/intel-graphics-compiler/build/IGC/llvm-deps/src/llvm/lib/IR/LegacyPassManager.cpp:1549
run at /workspace/srcdir/intel-graphics-compiler/build/IGC/llvm-deps/src/llvm/lib/IR/LegacyPassManager.cpp:539
run at /workspace/srcdir/intel-graphics-compiler/build/IGC/llvm-deps/src/llvm/lib/IR/LegacyPassManager.cpp:1676
CodeGen at /workspace/srcdir/intel-graphics-compiler/IGC/Compiler/CISACodeGen/OpenCLKernelCodeGen.cpp:3369
CodeGen at /workspace/srcdir/intel-graphics-compiler/IGC/Compiler/CISACodeGen/OpenCLKernelCodeGen.cpp:3395
TranslateBuildSPMD at /workspace/srcdir/intel-graphics-compiler/IGC/AdaptorOCL/dllInterfaceCompute.cpp:1508
TranslateBuildSPMDAndESIMD at /workspace/srcdir/intel-graphics-compiler/IGC/VISALinkerDriver/VLD.cpp:367
TranslateBuild at /workspace/srcdir/intel-graphics-compiler/IGC/AdaptorOCL/dllInterfaceCompute.cpp:1761
Translate at /workspace/srcdir/intel-graphics-compiler/IGC/AdaptorOCL/ocl_igc_interface/impl/igc_ocl_translation_ctx_impl.h:359
TranslateImpl at /workspace/srcdir/intel-graphics-compiler/IGC/AdaptorOCL/ocl_igc_interface/impl/igc_ocl_translation_ctx_impl.cpp:68
Translate<> at /workspace/destdir/include/igc/ocl_igc_interface/igc_ocl_translation_ctx.h:93
translate<IGC::IgcOclTranslationCtx<3> > at /workspace/srcdir/compute-runtime/shared/source/compiler_interface/compiler_interface.inl:75
build at /workspace/srcdir/compute-runtime/shared/source/compiler_interface/compiler_interface.cpp:153
compileGenBinary at /workspace/srcdir/compute-runtime/level_zero/core/source/module/module_imp.cpp:207
buildFromSpirV at /workspace/srcdir/compute-runtime/level_zero/core/source/module/module_imp.cpp:290
initializeTranslationUnit at /workspace/srcdir/compute-runtime/level_zero/core/source/module/module_imp.cpp:734
initialize at /workspace/srcdir/compute-runtime/level_zero/core/source/module/module_imp.cpp:536
create at /workspace/srcdir/compute-runtime/level_zero/core/source/module/module_imp.cpp:1241
createModule at /workspace/srcdir/compute-runtime/level_zero/core/source/device/device_imp.cpp:487
createModule at /workspace/srcdir/compute-runtime/level_zero/core/source/context/context_imp.cpp:935
zeModuleCreate at /workspace/srcdir/compute-runtime/level_zero/api/core/ze_module_api_entrypoints.h:23
validation_layer::zeModuleCreate(_ze_context_handle_t*, _ze_device_handle_t*, _ze_module_desc_t const*, _ze_module_handle_t**, _ze_module_build_log_handle_t**) at /home/tim/.julia/artifacts/07d2a0b1b466f4d6fab3f80843bd68cb0036c027/lib/libze_validation_layer.so (unknown line)
urProgramBuildExp at /home/tim/.julia/scratchspaces/8f75cd03-7ff8-4ecb-9b8f-daf728133b1b/conda/lib/libpi_level_zero.so (unknown line)
piProgramBuild at /home/tim/.julia/scratchspaces/8f75cd03-7ff8-4ecb-9b8f-daf728133b1b/conda/lib/libpi_level_zero.so (unknown line)
_pi_result sycl::_V1::detail::plugin::call_nocheck<(sycl::_V1::detail::PiApiKind)43, _pi_program*, int, _pi_device* const*, char const*, decltype(nullptr), decltype(nullptr)>(_pi_program*, int, _pi_device* const*, char const*, decltype(nullptr), decltype(nullptr)) const at /home/tim/.julia/scratchspaces/8f75cd03-7ff8-4ecb-9b8f-daf728133b1b/conda/lib/libsycl.so.7 (unknown line)
sycl::_V1::detail::ProgramManager::build(std::unique_ptr<_pi_program, _pi_result (*)(_pi_program*)>, std::shared_ptr<sycl::_V1::detail::context_impl>, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, _pi_device* const&, unsigned int) at /home/tim/.julia/scratchspaces/8f75cd03-7ff8-4ecb-9b8f-daf728133b1b/conda/lib/libsycl.so.7 (unknown line)
_ZZN4sycl3_V16detail14ProgramManager17getBuiltPIProgramERKSt10shared_ptrINS1_12context_implEERKS3_INS1_11device_implEERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEbENK3 at /home/tim/.julia/scratchspaces/8f75cd03-7ff8-4ecb-9b8f-daf728133b1b/conda/lib/libsycl.so.7 (unknown line)
sycl::_V1::detail::ProgramManager::getBuiltPIProgram(std::shared_ptr<sycl::_V1::detail::context_impl> const&, std::shared_ptr<sycl::_V1::detail::device_impl> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, bool) at /home/tim/.julia/scratchspaces/8f75cd03-7ff8-4ecb-9b8f-daf728133b1b/conda/lib/libsycl.so.7 (unknown line)
sycl::_V1::detail::ProgramManager::getOrCreateKernel(std::shared_ptr<sycl::_V1::detail::context_impl> const&, std::shared_ptr<sycl::_V1::detail::device_impl> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&) at /home/tim/.julia/scratchspaces/8f75cd03-7ff8-4ecb-9b8f-daf728133b1b/conda/lib/libsycl.so.7 (unknown line)
sycl::_V1::detail::enqueueImpKernel(std::shared_ptr<sycl::_V1::detail::queue_impl> const&, sycl::_V1::detail::NDRDescT&, std::vector<sycl::_V1::detail::ArgDesc, std::allocator<sycl::_V1::detail::ArgDesc> >&, std::shared_ptr<sycl::_V1::detail::kernel_bundle_impl> const&, std::shared_ptr<sycl::_V1::detail::kernel_impl> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::vector<_pi_event*, std::allocator<_pi_event*> >&, std::shared_ptr<sycl::_V1::detail::event_impl> const&, std::function<void* (sycl::_V1::detail::AccessorImplHost*)> const&, _pi_kernel_cache_config) at /home/tim/.julia/scratchspaces/8f75cd03-7ff8-4ecb-9b8f-daf728133b1b/conda/lib/libsycl.so.7 (unknown line)
_ZZN4sycl3_V17handler8finalizeEvENK3 at /home/tim/.julia/scratchspaces/8f75cd03-7ff8-4ecb-9b8f-daf728133b1b/conda/lib/libsycl.so.7 (unknown line)
sycl::_V1::handler::finalize() at /home/tim/.julia/scratchspaces/8f75cd03-7ff8-4ecb-9b8f-daf728133b1b/conda/lib/libsycl.so.7 (unknown line)
void sycl::_V1::detail::queue_impl::finalizeHandler<sycl::_V1::handler>(sycl::_V1::handler&, sycl::_V1::detail::CG::CGTYPE const&, sycl::_V1::event&) at /home/tim/.julia/scratchspaces/8f75cd03-7ff8-4ecb-9b8f-daf728133b1b/conda/lib/libsycl.so.7 (unknown line)
sycl::_V1::detail::queue_impl::submit_impl(std::function<void (sycl::_V1::handler&)> const&, std::shared_ptr<sycl::_V1::detail::queue_impl> const&, std::shared_ptr<sycl::_V1::detail::queue_impl> const&, std::shared_ptr<sycl::_V1::detail::queue_impl> const&, sycl::_V1::detail::code_location const&, std::function<void (bool, bool, sycl::_V1::event&)> const*) at /home/tim/.julia/scratchspaces/8f75cd03-7ff8-4ecb-9b8f-daf728133b1b/conda/lib/libsycl.so.7 (unknown line)
sycl::_V1::detail::queue_impl::submit(std::function<void (sycl::_V1::handler&)> const&, std::shared_ptr<sycl::_V1::detail::queue_impl> const&, sycl::_V1::detail::code_location const&, std::function<void (bool, bool, sycl::_V1::event&)> const*) at /home/tim/.julia/scratchspaces/8f75cd03-7ff8-4ecb-9b8f-daf728133b1b/conda/lib/libsycl.so.7 (unknown line)
sycl::_V1::queue::submit_impl(std::function<void (sycl::_V1::handler&)>, sycl::_V1::detail::code_location const&) at /home/tim/.julia/scratchspaces/8f75cd03-7ff8-4ecb-9b8f-daf728133b1b/conda/lib/libsycl.so.7 (unknown line)
oneapi::mkl::gpu::caxpby_sycl_internal(sycl::_V1::queue*, long, oneapi::mkl::value_or_pointer<std::complex<float> >, std::complex<float> const*, long, oneapi::mkl::value_or_pointer<std::complex<float> >, std::complex<float>*, long, std::vector<sycl::_V1::event, std::allocator<sycl::_V1::event> > const&) at /home/tim/.julia/scratchspaces/8f75cd03-7ff8-4ecb-9b8f-daf728133b1b/conda/lib/libmkl_sycl_blas.so.4 (unknown line)
oneapi::mkl::gpu::caxpby_sycl(sycl::_V1::queue*, long, oneapi::mkl::value_or_pointer<std::complex<float> >, std::complex<float> const*, long, oneapi::mkl::value_or_pointer<std::complex<float> >, std::complex<float>*, long, std::vector<sycl::_V1::event, std::allocator<sycl::_V1::event> > const&) at /home/tim/.julia/scratchspaces/8f75cd03-7ff8-4ecb-9b8f-daf728133b1b/conda/lib/libmkl_sycl_blas.so.4 (unknown line)
oneapi::mkl::blas::caxpby(sycl::_V1::queue&, long, oneapi::mkl::value_or_pointer<std::complex<float> >, std::complex<float> const*, long, oneapi::mkl::value_or_pointer<std::complex<float> >, std::complex<float>*, long, std::vector<sycl::_V1::event, std::allocator<sycl::_V1::event> > const&) at /home/tim/.julia/scratchspaces/8f75cd03-7ff8-4ecb-9b8f-daf728133b1b/conda/lib/libmkl_sycl_blas.so.4 (unknown line)
oneapi::mkl::blas::column_major::axpby(sycl::_V1::queue&, long, oneapi::mkl::value_or_pointer<std::complex<float> >, std::complex<float> const*, long, oneapi::mkl::value_or_pointer<std::complex<float> >, std::complex<float>*, long, std::vector<sycl::_V1::event, std::allocator<sycl::_V1::event> > const&) at /home/tim/.julia/scratchspaces/8f75cd03-7ff8-4ecb-9b8f-daf728133b1b/conda/lib/libmkl_sycl_blas.so.4 (unknown line)

cc @pengtu

[mmeterel]

@maleadt This looks like a driver issue or driver mismatch with the base tool kit to me. Are you using same driver version for 2024.0 and 2024.1?

I see this line: Is this the version used for both base tool kits?
Driver Version 24.09.028717

Can you try a newer driver with 2024.1?
Can you also try using Level-zero backend?

[maleadt]

Are you using same driver version for 2024.0 and 2024.1?

Yes, we build the driver separately from the MKL bits we use from Conda. Are there driver requirements on MLK from 2024.1? If so, are these documented anywhere?

I see this line: Is this the version used for both base tool kits?
Driver Version 24.09.028717

Correct; that's what I meant with NEO v24.9.28717 with IGC v1.0.16238.
This is also the latest driver available on https://github.com/intel/compute-runtime/releases, so which one exactly do you want me to try?

Can you also try using Level-zero backend?

How do I do that?

[eero-t]

Can you also try using Level-zero backend?

How do I do that?

NEO/IGC/SyCL is level-zero backend, but as you're selecting GPU, maybe that question was about whether it happens also when selecting CPU backend (with same NEO/IGC version)?

[maleadt]

whether it happens also when selecting CPU backend (with same NEO/IGC version)

I'm confused; does NEO/IGC have a cpu back-end? FWIW, using sycl::cpu_selector simply throws a PI_ERROR_DEVICE_NOT_FOUND.

In any case, upgrading to NEO 24.13.29138.7 seems to have fixed this.

[maleadt]

Scratch that, I'm still seeing the error with axpby and trmv: https://buildkite.com/julialang/oneapi-dot-jl/builds/1062#018f80f4-a375-44b0-9a4b-8a45ce9fb9e7

It doesn't reproduce consistently though; is it possible the validation isn't deterministic, or that compiled kernels are cached somewhere?

[eero-t]

I'm confused; does NEO/IGC have a cpu back-end? FWIW, using sycl::cpu_selector simply throws a PI_ERROR_DEVICE_NOT_FOUND.

Sorry, I was confused myself...

It doesn't reproduce consistently though; is it possible the validation isn't deterministic, or that compiled kernels are cached somewhere?

One way to check that it's always the same kernel code, could be dumping it on disk:
https://github.com/intel/intel-graphics-compiler/blob/master/documentation/shader_dumps_instruction.md