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:
__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;
}