CODEWITHSUNDEEP
crenderingopenclopencl-c
The A
The A

Reputation: 300

OpenCL: clBuildProgram Fails without log

What I am trying to accomplish:
I am trying to render some stuff in OpenCL and write it to the OpenGL Framebuffer (As it is the only Framebuffer I can get via Renderbuffers etc., but I will gladly accept any others I could use - You will not help me tho by telling to use glsl shaders)

The Problem:
As the Title says, the OpenCL function clBuildProgram Fails with error -11 (CL_​BUILD_​PROGRAM_​FAILURE). This wouldn't be an issue, but the Log from the CL compiler is empty. I doublechecked my log code, but it should be fine. I posted it down below none the less, just so you can see yourself.

What I tried to fix:

What I found out:
From the last point of what I tried, I noticed, that the write_imageui and read_imageui opencl functions make the compiler fail to compile my code (this is why I checked for the "cl_khr_gl_sharing" extension)

Furthermore:
My operating System is Windows 10, the C Compiler I am using is GCC (I do not know, how that could help, since the host program compiles fine, but here it is none the less)

Some Code:
The Shader/Kernel (as minified as possible to reproduce the problem, I hope also for you; the last two lines of calls are what I think causes the opencl compiler to not work; the other stuff is in there, to make it a shader, that can actually process something, once it is working):

#define ScreenWidth 1000
#define ScreenHight 1000
const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;
__kernel void rainbow(__read_write image2d_t asd) {
    int i = get_global_id(0);
    unsigned int x = i%ScreenWidth;
    unsigned int y = i/ScreenHight;
    uint4 pixel;
    pixel = read_imageui(asd, sampler, (int2)(x, y));
    write_imageui(asd, (int2)(x, y), pixel);
}

Minified Call Code (C) which does all the initialization stuff necessary (note: the log buffer is dynamically changeable):

cl_program program = clCreateProgramWithSource(contextZ, 1, (const char **)&source_str, (const size_t *)&source_size, &ret);

size_t retSize = 0;
clGetDeviceInfo(device_id, CL_DEVICE_EXTENSIONS, 0, NULL, &retSize);

char extensions[retSize];
clGetDeviceInfo(device_id, CL_DEVICE_EXTENSIONS, retSize, extensions, &retSize);
printf("%s\n", extensions);

// Build the program
ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
if (ret == CL_BUILD_PROGRAM_FAILURE) {
    l_logError("Could not build Kernel!");
    // Determine the size of the log
    size_t log_size;
    printf(" reta: %i\n", clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size));

    // Allocate memory for the log
    char *log = (char *) malloc(log_size);

    // Get the log
    printf(" retb: %i\n", clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, log_size, log, NULL));

    // Print the log
    printf(" ret-val: %i\n", ret);
    printf("%s\n", log);
}

You might be interested in the output (Last 2 lines are caused by the Kernel not beeing built. The program could be made from the source though - look at the code):

E: Could not build Kernel!
 reta: 0
 retb: 0
 ret-val: -11

E: Could not create Kernel!
 kernel error: -45

Did anybody else have a similar problem? Any Ideas, what I should do about it? Might there be a Header for the cl Kernel/Shader, which I need to include in it? It could be possible, that my clBuildProgram call is incorrect? (I read somebody did not pass a device, so maybe something else could be missing in my code)

Make sure to tell me, if you need further details, so I can provide them (I cannot think of any other you might need right now)

Thanks in Advance for your time!

EDIT:
According to the specification, a device needs to support the CL_​DEVICE_​IMAGE_​SUPPORT extension, which it does

I checked it using this:

cl_bool image_support = CL_FALSE;
clGetDeviceInfo(device_id, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool), &image_support, NULL);
printf("image_support: %i\n", image_support);

Which outputs:

image_support: 1

aka. CL_TRUE

Edit 2:
It turns out, OpenCL extensions need to be enabled in the kernel: https://www.khronos.org/registry/OpenCL/sdk/2.2/docs/man/html/EXTENSION.html
Adding #pragma OPENCL EXTENSION all : enable in the first line of the kernel/shader does results in the same issue tho

EDIT 3:
Removing the __read_write flag from the kernel image paramters or replace it with something else (like __read_only) causes the OpenCL compiler to crash or loop infinetly, as clBuildProgram never returns (or will return in like a very long time)

Upvotes: 1

Views: 1023

Answers (1)

The A
The A

Reputation: 300

What I found out in the last few days may have been slightly incorrect. My Edit 3 (from the original post/question) states, that replacing the __read_write with an __read_only causes the compiler to completly crash. This is incorrect, as I can now confirm, I just didn't add any more debug-output code after the compilation call. Adding some more debug lines after the clBuildProgram call shows, that it actually works.

I do not know, why this causes the OpenCL compiler to output literally nothing as an error, and the vendors of the driver should definetly fix this/output something (Device info in comment), to make development somewhat easier. (Even just a warning will be helpful)

I found this stackoverflow post, discussing a similar problem: OpenCL - Pass image2d_t twice to get both read and write from kernel?. This is how I even figured out, that this can cause devastating problems.
To be fair, the official docs state this, but I understood it more like they can be combined to __read_write (read_only | write_only == __read_write):

aQual in the following table refers to one of the access qualifiers. For write functions this may be write_only or read_write.

I replaced my kernel with only a write_image call (write_imageui) and set the image2d_t to be __write_only, to get minimal debug possibilities. This caused the shader/kernel to compile succesfully, but the screen is still empty. But latter is a matter for another question.

Upvotes: 2

Related Questions