OpenCL Graphics Driver 30.0.101.1069 Issue

[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