intel / intel/intel-graphics-compiler

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

Open
#345 2 comments 0 reactions 0 assignees View on GitHub
Dominant language
C++
Stars
718
Forks
191
PR merge metrics
No merged PRs in 30d

Description

# 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:

```c
__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](https://github.com/KhronosGroup/SPIRV-LLVM-Translator/blob/065a94408d01bacf2ee86443f56cfaa64fda1534/lib/SPIRV/libSPIRV/SPIRVOpCodeEnum.h#L333), 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](https://github.com/KhronosGroup/SPIRV-LLVM-Translator/issues/1391) / [PR #2511](https://github.com/KhronosGroup/SPIRV-LLVM-Translator/pull/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`](https://github.com/intel/intel-graphics-compiler/blob/36a6b4e5190a181df613ab4cc6db5115ca8d56f8/IGC/AdaptorOCL/SPIRV/libSPIRV/SPIRVOpCodeEnum.h#L344).

Which version of SPIRVOpCodeEnum.h is used depends on the compile-time settings [`IGC_OPTION__LINK_KHRONOS_SPIRV_TRANSLATOR`](https://github.com/intel/intel-graphics-compiler/blob/36a6b4e5190a181df613ab4cc6db5115ca8d56f8/external/llvm/llvm_spirv_source_hook.cmake#L27) and [`IGC_OPTION__USE_KHRONOS_SPIRV_TRANSLATOR_IN_SC`](https://github.com/intel/intel-graphics-compiler/blob/36a6b4e5190a181df613ab4cc6db5115ca8d56f8/IGC/CMakeLists.txt#L272), and indeed, [Arch's package is built with both of them disabled](https://gitlab.archlinux.org/archlinux/packaging/packages/intel-graphics-compiler/-/blob/e9d3b3274ed24feb33ac17c9280f044179cbf55b/PKGBUILD#L72-73), 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](https://github.com/KhronosGroup/SPIRV-LLVM-Translator/pull/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`):

```c
#include
#include
#define CL_TARGET_OPENCL_VERSION 120
#include

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;
}
```

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.