GPU Compute Software
Ask questions about Intel® Graphics Compute software technologies, such as OpenCL* GPU driver and oneAPI Level Zero
140 Discussions

OpenCL CL_INVALID_KERNEL_ARGS with driver > 30.0.101.1340

Scott_S_2
Novice
‎10-06-2022 08:34 AM
2,083 Views

Hi,

I've been getting CL_INVALID_KERNEL_ARGS for all my kernels after updating to the latest intel driver version.

I ran into this when I was getting ".1" appended to my kernel names, where I updated to the latest driver to try and fix it.

I methodically tried installing all previous driver versions and the first version to not give me the CL_INVALID_KERNEL_ARGS was 30.0.101.1340.

Did something change in subsequent driver versions that would cause SPIR-V consumption to go wrong?

 

 

Here is the cpp code to duplicate the issue:

 

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

#define MAX_PLATFORM_SIZE 256
#define MAX_DEVICE_SIZE 256

namespace {
    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 check_error(int result,int line){
        if(result != CL_SUCCESS){
            throw std::runtime_error(std::to_string(line) + std::string(": error: ") + std::to_string(result));
        }
    }
}


int main()
{  
    std::vector<unsigned char> data = load_file("window_level.spv");

    // Get platform and device information
    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);
    check_error(ret,__LINE__);

    ret = clGetPlatformIDs(ret_num_platforms, platform_id, &ret_num_platforms);
    check_error(ret,__LINE__);

    for (unsigned int i=0; i<ret_num_platforms; i++) 
    {
        ret = clGetPlatformInfo(platform_id[i], CL_PLATFORM_NAME, sizeof(buf), buf, NULL);
        check_error(ret,__LINE__);

        ret = clGetDeviceIDs(platform_id[i], CL_DEVICE_TYPE_ALL, MAX_DEVICE_SIZE, device_id, &ret_num_devices);
        check_error(ret,__LINE__);

        for (unsigned int j=0; j<ret_num_devices; j++) {
            ret = clGetDeviceInfo(device_id[j], CL_DEVICE_NAME, sizeof(buf), buf, NULL);
            check_error(ret,__LINE__);

            ret = clGetDeviceInfo(device_id[j], CL_DEVICE_VENDOR, sizeof(buf), buf, NULL);
            check_error(ret,__LINE__);
        }
    }
 

    // Create an OpenCL context
    cl_context context = clCreateContext( NULL, 1, &device_id[0], NULL, NULL, &ret);
 
    // Create a command queue
    cl_command_queue command_queue = clCreateCommandQueueWithProperties(context, device_id[0], 0, &ret);

    cl_ulong max_constants_size = 0;
    ret = clGetDeviceInfo(device_id[0], CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, sizeof(max_constants_size), &max_constants_size, NULL);
    check_error(ret,__LINE__);
 
    // create some images
    cl_image_format image_format_gray16s;
    image_format_gray16s.image_channel_data_type = CL_SIGNED_INT16;
    image_format_gray16s.image_channel_order = CL_R;

    int width = 512;
    int height = 512;
    std::vector<short> slice(512*512, 0); 

    cl_image_desc image_desc_gray16s;
    image_desc_gray16s.image_type = CL_MEM_OBJECT_IMAGE2D;
    image_desc_gray16s.image_width = width;
    image_desc_gray16s.image_height = height;
    image_desc_gray16s.image_depth = 1;
    image_desc_gray16s.image_array_size = 1;
    image_desc_gray16s.image_row_pitch = 0;
    image_desc_gray16s.image_slice_pitch = 0;
    image_desc_gray16s.num_mip_levels = 0;
    image_desc_gray16s.num_samples = 0;
    image_desc_gray16s.buffer = NULL;

    cl_mem cl_image_gray16s = clCreateImage(
        context, 
        CL_MEM_READ_ONLY,
        &image_format_gray16s, 
        &image_desc_gray16s,
        NULL, 
        &ret
    );

    size_t origin[] = {0,0,0};
    size_t region[] = {(size_t)width,(size_t)height,1};
    ret = clEnqueueWriteImage(
        command_queue,
        cl_image_gray16s,
        CL_TRUE,
        origin, 
        region,
        width*sizeof(signed short),
        0,
        slice.data(), 
        0, 
        NULL, 
        NULL
    );

    check_error(ret,__LINE__);

    cl_image_format image_format_gray8;
    image_format_gray8.image_channel_data_type = CL_UNSIGNED_INT8;
    image_format_gray8.image_channel_order = CL_R;

    cl_image_desc image_desc_gray8;
    image_desc_gray8.image_type = CL_MEM_OBJECT_IMAGE2D;
    image_desc_gray8.image_width = width;
    image_desc_gray8.image_height = height;
    image_desc_gray8.image_depth = 1;
    image_desc_gray8.image_array_size = 1;
    image_desc_gray8.image_row_pitch = 0;
    image_desc_gray8.image_slice_pitch = 0;
    image_desc_gray8.num_mip_levels = 0;
    image_desc_gray8.num_samples = 0;
    image_desc_gray8.buffer = NULL;

    cl_mem cl_image_gray8 = clCreateImage(
        context, 
        CL_MEM_WRITE_ONLY,
        &image_format_gray8, 
        &image_desc_gray8,
        NULL, 
        &ret
    );

    cl_int error = 0;
    cl_program program = clCreateProgramWithIL(context,data.data(),data.size(),&error); // TODO
 
    // Build the program
    ret = clBuildProgram(program, 1, &device_id[0], NULL, NULL, NULL);
    check_error(ret,__LINE__);

    std::string kernel_name;
    {
        size_t ret_val_size = 0;
        ret = clGetProgramInfo(program, CL_PROGRAM_KERNEL_NAMES, 0, NULL, &ret_val_size);
        check_error(ret,__LINE__);

        kernel_name.resize(ret_val_size+1);
        ret = clGetProgramInfo(program, CL_PROGRAM_KERNEL_NAMES, ret_val_size, kernel_name.data(), NULL);
        check_error(ret,__LINE__);

        kernel_name[ret_val_size] = '\0';
    }
 
    // Create the OpenCL kernel
    cl_kernel kernel = clCreateKernel(program, kernel_name.c_str(), &ret);
    check_error(ret,__LINE__);

    // Set the arguments of the kernel
    ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&cl_image_gray16s);
    ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&cl_image_gray8);

    float window_width = 1000.0f;
    float window_level = 0.0f;
    float half_window_width = window_width/2;
    float window_bottom = window_level - half_window_width;
    float reciprocal_width = 255.0f / window_width;
    
    ret = clSetKernelArg(kernel, 2, sizeof(float), (void *)&window_bottom);
    ret = clSetKernelArg(kernel, 3, sizeof(float), (void *)&reciprocal_width);
     
    check_error(ret,__LINE__);
     
    // Execute the OpenCL kernel on the list
    size_t local_work[3]  = {8, 8, 1};
    size_t global_work[3]  = {(size_t)(local_work[0] * (size_t)ceilf((float)width/local_work[0])), (size_t)(local_work[1] * (size_t)ceilf((float)height/local_work[1])), 1};

    ret = clEnqueueNDRangeKernel(
        command_queue, 
        kernel,
        2, 
        NULL, 
        global_work, 
        local_work, 
        0, 
        NULL, 
        NULL
    );

    // error -52: CL_INVALID_KERNEL_ARGS
    check_error(ret,__LINE__);

    return 0;
}

 

I generate the spirv file via clang-15 and llvm-spriv-15:

 

clang -target spirv64 -emit-llvm -O1 -c -o window_level.bc window_level.clcpp
llvm-spirv -spirv-max-version=1.2 -o window_level.spv window_level.bc

 

the actual kernel:


kernel void window_level(
    __read_only image2d_t wl_image,
    __write_only image2d_t dst_image,
    float window_bottom,
    float width_rcp
) {
    int2 coords2d;
    coords2d.x = get_global_id(0);
    coords2d.y = get_global_id(1);

    int4 wl_sample = read_imagei(wl_image, coords2d);

    int4 dst_value;
    dst_value.x = clamp((int)(wl_sample.x - window_bottom) * width_rcp,0.0f,255.0f);

    write_imagei(dst_image, coords2d, dst_value);
}

I've attached the SPIR-V file too.

 

 

 

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

window_level.zip
0 Kudos

Link Copied

2 Replies
SantoshY_Intel
Moderator
‎10-06-2022 10:37 PM
2,060 Views

Hi,


Thank you for posting in the Intel forums.


Since this is a duplicate thread of https://community.intel.com/t5/GPU-Compute-Software/OpenCL-Kernel-Name-appends-quot-1-quot/m-p/1419720#M644, we will no longer monitor this thread. We will continue addressing this issue in the other thread.  


Thanks & Regards,

Santosh


0 Kudos
Copy link
Scott_S_2
Novice
‎10-07-2022 05:44 AM
2,052 Views

Hi,

 

This is not really a duplicate thread, its an additional issue discovered in the original thread.

The advice in that thread to upgrade the driver fixed the ".1" name issue, but exposed another CL_INVALID_KERNEL_ARGS issue. 

I'm not certain they are the same issue, so I posted another issue here to get advice on that.

 

Regards

Scott

0 Kudos
Copy link
Reply

Community support is provided during standard business hours (Monday to Friday 7AM - 5PM PST). Other contact methods are available here.

Intel does not verify all solutions, including but not limited to any file transfers that may appear in this community. Accordingly, Intel disclaims all express and implied warranties, including without limitation, the implied warranties of merchantability, fitness for a particular purpose, and non-infringement, as well as any warranty arising from course of performance, course of dealing, or usage in trade.

For more complete information about compiler optimizations, see our Optimization Notice.