"IGC SPIRV Consumer does not support the following opcode: 402" building OpenCL programs since igc-1.0.17384.11

[joanbm]

Summary

On my Arch Linux system with an Intel(R) HD Graphics 620 iGPU (integrated into i5-7200U), I have found that since igc-1.0.17384.11, trying to build a simple OpenCL program like the following:

__kernel void dummy( __global int *ptr) { int x = ptr != NULL; }

Results in a program build failure (from clBuildProgram) with the following log from clGetProgramBuildInfo:

error: IGC SPIRV Consumer does not support the following opcode: 402

Investigation

From what I have been able to gather the opcode 402 corresponds to PtrNotEqual as defined in SPIRV-LLVM-Translator's SPIRVOpCodeEnum.h, which makes sense given the program that fails does a pointer comparison.

From what I have been able to gather, in SPIRV-LLVM-Translator issue #1391 / PR #2511 the PtrNotEqual is generated in more scenarios than before, which seems to explain why this is happening since igc-1.0.17384.11 (as that's the first version to include this change).

However, I noticed that intel-graphics-compiler also includes a copy of the SPIRVOpCodeEnum.h, and that one does not have a definition for PtrNotEqual.

Which version of SPIRVOpCodeEnum.h is used depends on the compile-time settings IGC_OPTION__LINK_KHRONOS_SPIRV_TRANSLATOR and IGC_OPTION__USE_KHRONOS_SPIRV_TRANSLATOR_IN_SC, and indeed, Arch's package is built with both of them disabled, which leads to it using the 'legacy' version (i.e. the one without the PtrNotEqual definition).

I have rebuilt the Arch package with both options enabled and it fixes the OpenCL program build error and OpenCL appears to work correctly.

If this is indeed the cause of the issue: Is this a bug in intel-graphics-compiler, or should it be considered a bug in Arch Linux's package? The problem as I see is that the old version is 'legacy' but the new one is 'experimental', so it's not clear what versions are actually supported.

Additional information

Package versions

This reproduces on a fully up-to-date Arch Linux install, with:

linux                   6.11.3.arch1-1
intel-compute-runtime   24.35.30872.22-2
intel-graphics-compiler 1:1.0.17537.20-1

As mentioned the hardware is a Intel(R) HD Graphics 620 iGPU (integrated into i5-7200U). Something else that I have noticed is that this SPIRV-LLVM-Translator PR #2511 checks for SPIR-V 1.4 support (the BM->isAllowedToUseVersion(VersionNumber::SPIRV_1_4) check), and based on clinfo's output, my iGPU does not support SPIR-V 1.4, but I'm unsure if that's actually related to the issue:

    IL version                                    SPIR-V_1.3 SPIR-V_1.2 SPIR-V_1.1 SPIR-V_1.0 
    ILs with version                              SPIR-V                                                           0x403000 (1.3.0)
                                                  SPIR-V                                                           0x402000 (1.2.0)
                                                  SPIR-V                                                           0x401000 (1.1.0)
                                                  SPIR-V                                                           0x400000 (1.0.0)
    SPIR versions                                 1.2 

Sample code

Full reproducer test program (build with gcc cltest.c -lOpenCL && ./cltest):

#include <stdio.h>
#include <stdlib.h>
#define CL_TARGET_OPENCL_VERSION 120
#include <CL/cl.h>

int main() {
    cl_platform_id platform;
    cl_device_id device;
    clGetPlatformIDs(1, &platform, NULL);
    clGetDeviceIDs(platform, CL_DEVICE_TYPE_DEFAULT, 1, &device, NULL);

    char device_name[128];
    clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(device_name), device_name, NULL);
    printf("Using OpenCL device: %s\n", device_name);

    cl_context context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL);

    const char* kernelSource = "__kernel void dummy( __global int *ptr) { int x = ptr != NULL; }";

    cl_program program = clCreateProgramWithSource(context, 1, &kernelSource, NULL, NULL);
    cl_int err = clBuildProgram(program, 1, &device, NULL, NULL, NULL);

    if (err != CL_SUCCESS) {
        size_t log_size;
        clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
        char *log = (char *)malloc(log_size);
        clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, log_size, log, NULL);
        fprintf(stderr, "Error in kernel build:\n%s\n", log);
        return 1;
    }

    return 0;
}

[pszymich]

Hi @joanbm, the build configuration with built-in translator is not supported anymore and linking with upstream is production, not experimental. The option cleanup removing the legacy path and enabling linking with upstream translator by default is in the backlog.

As a workaround for now please enable both options in your build environment.

@mateuszchudyk, please update this issue when the cleanup is done, thanks.

[joanbm]

Hi @pszymich, thanks for confirming what the current state is. I have opened an issue on the Arch Linux package to switch to the Khronos SPIR-V translator. This issue can be closed now if you wish so. Thanks!