OpenCL compiler backend issue with Drivers >= 30.0.101.1069

[Scott_S_2]

Hi all,

I have just upgraded to the latest Intel Graphics Driver and I can no longer build opencl programs from spirv. It seems image3d write is broken/missing as I get the following error for a simple opencl kernel:

failed to build opencl program:

error: undefined reference to `__builtin_spirv_OpImageWrite_img3d_wo_v3i32_v4f32'
in function: '__builtin_spirv_OpImageWrite_img3d_wo_v3i32_v4f32' called by kernel: 'driver_test2'

error: backend compiler failed build.

here is a simple kernel that triggers this:

#include <opencl_memory>
#include <opencl_work_item>
#include <opencl_image>
#include <opencl_integer>

using namespace cl;

kernel void driver_test(
    image3d<float4, image_access::write> dst_vol
) {
    const int3 dst_coords(get_global_id(0), get_global_id(1), get_global_id(2));

    float4 source_sample(0.0f,0.0f,0.0f,0.0f);
    dst_vol.write(dst_coords, source_sample);  
}

and the cpp part:

#include "CL/cl.h"
#include <iostream>
#include <fstream>

#define MAX_PLATFORM_SIZE 256
#define MAX_DEVICE_SIZE 256

std::vector<unsigned char> load_file(const std::string& file)
{
    std::fstream input(file, std::ios::in | std::ios::binary | std::ios::ate);
    auto size = input.tellg();
    input.seekg(0, std::ios::beg);
    std::vector<unsigned char> binary(size);
    input.read((char*)binary.data(),size);
    input.close();

    return binary;
}

void driver_test()
{  
    std::vector<unsigned char> data = load_file("driver_test.spir");

    cl_device_id device_id[256];
    cl_platform_id platform_id[256];
    cl_uint ret_num_devices;
    cl_uint ret_num_platforms;
    char buf[4096];

    cl_int ret = clGetPlatformIDs(0, 0, &ret_num_platforms);
    ret = clGetPlatformIDs(ret_num_platforms, platform_id, &ret_num_platforms);
    if (ret != CL_SUCCESS) {
        std::cout << "something went wrong in clGetPlatformIDs" << std::endl;
        return;
    }

    for (unsigned int i=0; i<ret_num_platforms; i++) 
    {
        ret = clGetPlatformInfo(platform_id[i], CL_PLATFORM_NAME, sizeof(buf), buf, NULL);
        if (ret != CL_SUCCESS) {
            std::cout << "something went wrong in clGetPlatformInfo" << std::endl;
            return;
        }
        std::cout << "Platform[" << i << "]: " << buf << std::endl;

        ret = clGetDeviceIDs(platform_id[i], CL_DEVICE_TYPE_ALL, MAX_DEVICE_SIZE, device_id, &ret_num_devices);
        if (ret != CL_SUCCESS) {
            std::cout << "something went wrong in clGetDeviceIDs" << std::endl;
            return;
        }

        for (unsigned int j=0; j<ret_num_devices; j++) {
            ret = clGetDeviceInfo(device_id[j], CL_DEVICE_NAME, sizeof(buf), buf, NULL);
            if (ret != CL_SUCCESS) {
                std::cout << "something went wrong in clGetDeviceInfo" << std::endl;
                return;
            }
            std::cout << "Device Name: " << buf << std::endl;

            ret = clGetDeviceInfo(device_id[j], CL_DEVICE_VENDOR, sizeof(buf), buf, NULL);
            if (ret != CL_SUCCESS) {
                std::cout << "something went wrong in clGetDeviceInfo" << std::endl;
                return;
            }
            std::cout << "Device Vendor: " << buf << std::endl;
        }
        std::cout << std::endl;
    }
 
    cl_context context = clCreateContext( NULL, 1, &device_id[0], NULL, NULL, &ret);
 
    cl_int err = 0;
    cl_program program = clCreateProgramWithIL((cl_context)context, data.data(), data.size(), &err); 
    if(err != CL_SUCCESS){
        std::cout << "failed to create opencl program with IL " << std::endl;
        return;
    }

    // sometimes crashes here
    err = clBuildProgram(program, 1, &device_id[0], NULL, NULL, NULL);

    cl_int build_status;
    err = clGetProgramBuildInfo(program, device_id[0], CL_PROGRAM_BUILD_STATUS, sizeof(cl_build_status), &build_status, NULL);

    if(build_status != CL_SUCCESS){
        size_t ret_val_size;
        err = clGetProgramBuildInfo(program, device_id[0], CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
                 
        std::vector<char> build_log(ret_val_size+1,0);
        err = clGetProgramBuildInfo(program, device_id[0], CL_PROGRAM_BUILD_LOG, ret_val_size, build_log.data(), NULL);

        build_log[ret_val_size] = '\0';
        std::cout << "failed to build opencl program: \n" << build_log.data() << std::endl;
    }
}

I've been rolling back the driver versions and the earliest driver where this happens is: 30.0.101.1069

Earlier driver version work fine

Sometimes it crashes at clBuildProgram too.

I have attached the .cl, cpp and .spir files. We use:
SPIRV-LLVM
SPIR
libclcxx

to generate the spirv file. I've attached the generated spirv for easier testing.

Any help would be appreciated.

Regards
Scott

System Setup Information:
-----------------------------------------

System Used: Inspiron 3670
CPU SKU: Intel(R) Core(TM) i3-8100 CPU @ 3.60GHz 3.60 GHz
GPU SKU: Intel(R) UHD Graphics 630
System BIOS Version: Dell Inc. 1.3.4, 24/05/2018
Graphics Driver Version: 30.0.101.1069
Operating System: Windows 10 Pro
OS Version: 10.0.19042 Build 19042
API: OpenCL
Occurs on non-Intel GPUs?: no tested

[AlHill]

@Scott_S_2 One post is enough.   Stay with your original post.

Doc (not an Intel employee or contractor)
[Maybe Windows 12 will be better]

[Scott_S_2]

@AlHill  the system continuously marked my posts as SPAM, without any extra information, hence the weird state of multiple posts appearing

[AlHill]

Well, the system did not mark as spam the multiple post that I saw.  

Doc (not an Intel employee or contractor)
[Maybe Windows 12 will be better]

[Ben_A_Intel]

Thanks for the thorough description and the reproducer!

As best I can tell, the issue is caused by using a three-component vector as the image coordinate for the 3D image write:

    const int3 dst_coords(get_global_id(0), get_global_id(1), get_global_id(2));
    ...
    dst_vol.write(dst_coords, source_sample);  

The SPIR-V spec - and specifically the OpenCL SPIR-V environment spec - isn't clear whether this should be supported or not.  But, since the corrosponding OpenCL C function takes a four-component vector as the image coordinate for a 3D image write, it looks like this is the only SPIR-V version we currently have implemented in our compiler:

https://github.com/intel/intel-graphics-compiler/blob/master/IGC/BiFModule/Implementation/images.cl#L997

I'll try to think if there is a workaround for this issue.

In the meantime, I'd suggest filing an OpenCL spec issue to clarify whether three-component coordinates are valid for 3D image writes, or if four-component coordinates are required:

https://github.com/KhronosGroup/OpenCL-Docs/issues

PS: For whatever it's worth, our CPU OpenCL device does not seem to be accepting the three-component vector either.

[Scott_S_2]

Hi,

This is interesting, although confusing as the previous driver versions seem to handle this fine.  I'm guessing something has changed between versions with the SPIRV handling. 

I also get errors with regular C kernels when using LLVM 15 and LLVM-SPIRV generated spirv e.g.

__kernel void window_level(
    __read_only image2d_t srcImg,
    __write_only image2d_t dstImg,
    float window_bottom,
    float width_rcp
)           
{                                                                                                                                    
   int2 coord = (int2) ((int)get_global_id(0), (int)get_global_id(1));                                                                         
                                                                                                                                     
   int2 size = get_image_dim(srcImg);                                                                                                
                                                                                                                                     
   if(coord.x >= size.x || coord.y >= size.y){                                                                                       
       return;                                                                                                                       
   }                                                                                                                                 
                                                                                                                                     
   sampler_t sampler = CLK_FILTER_NEAREST | CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP;                                         
                                                                                                                                     
   int4 src_value = read_imagei(srcImg, sampler, coord);                                                                             
                                                                                                                                     
   int4 dst_value;                                                                                                                   
   dst_value.x = (int)((src_value.x - window_bottom) * width_rcp);                                                                   
                                                                                                                                     
   if(dst_value.x > 255){ dst_value.x = 255; }                                                                                       
   if(dst_value.x < 0){ dst_value.x = 0; }                                                                                           
                                                                                                                                     
   write_imagei(dstImg, coord, dst_value);                                                                                           
} 

gives me the following error

error: undefined reference to `__builtin_spirv_OpImageWrite_img2d_wo_v2i32_v4i32_i32'
in function: '__builtin_spirv_OpImageWrite_img2d_wo_v2i32_v4i32_i32' called by kernel: 'window_level.1'

error: backend compiler failed build.

So this seems to be a common issue with newer driver versions. 

I can't just change the coords to int4 as it complains that write_imagei cannot convert the input from int4 to int2, and dst_vol.write(dst_coords, source_sample) cannot convert from int4 to int3 etc.

In the meantime I'll post in the OpenCL-Docs issue tracker.

Scott

[Scott_S_2]

for this particular C kernel case: passing -spirv-max-version=1.2 solves the 2D undefined reference issue