Static variable in OpenCL C

Question

I'm writing a renderer from scratch using openCL and I have a little compilation problem on my kernel with the error : CL_BUILD_PROGRAM : error: program scope variable must reside in constant address space static float* objects;

The problem is that this program compiles on my desktop (with nvidia drivers) and doesn't work on my laptop (with nvidia drivers), also I have the exact same kernel file in another project that works fine on both computers... Does anyone have an idea what I could be doing wrong ?

As a clarification, I'm coding a raymarcher which's kernel takes a list of objects "encoded" in a float array that is needed a lot in the program and that's why I need it accessible to the hole kernel.

Here is the kernel code simplified :

float* objects;

float4 getDistCol(float3 position) {
    int arr_length = objects[0];

    float4 distCol = {INFINITY, 0, 0, 0};

    int index = 1;
    while (index < arr_length) {
        float objType = objects[index];

        if (compare(objType, SPHERE)) {
            // Treats the part of the buffer as a sphere
            index += SPHERE_ATR_LENGTH;

        } else if (compare(objType, PLANE)) {
            //Treats the part of the buffer as a plane
            index += PLANE_ATR_LENGTH;

        } else {
            float4 errCol = {500, 1, 0, 0};
            return errCol;
        }

    }
}

__kernel void mkernel(__global int *image, __constant int *dimension,
                      __constant float *position, __constant float *aimDir, __global float *objs) {


    objects = objs;

    // Gets ray direction and stuf
    // ...
    // ...

    float4 distCol = RayMarch(ro, rd);
    float3 impact = rd*distCol.x + ro;

    col = distCol.yzw * GetLight(impact);

    image[dimension[0]*dimension[1] - idx*dimension[1]+idy] = toInt(col);

Where getDistCol(float3 position) gets called a lot by a lot of functions and I would like to avoid having to pass my float buffer to every function that needs to call getDistCol()...

ProjectPhysX · Answer 1 · 2023-01-29T17:59:27.477

There is no "static" variables allowed in OpenCL C that you can declare outside of kernels and use across kernels. Some compilers might still tolerate this, others might not. Nvidia has recently changed their OpenCL compiler from LLVM 3.4 to NVVM 7 in a driver update, so you may have the 2 different compilers on your desktop/laptop GPUs.

In your case, the solution is to hand the global kernel parameter pointer over to the function:

float4 getDistCol(float3 position, __global float *objects) {
    int arr_length = objects[0]; // access objects normally, as you would in the kernel
    // ...
}
kernel void mkernel(__global int *image, __constant int *dimension, __constant float *position, __constant float *aimDir, __global float *objs) {
    // ...
    getDistCol(position, objs); // hand global objs pointer over to function
    // ...
}

Lonely variables out in the wild are only allowed as constant memory space, which is useful for large tables. They are cached in L2$, so read-only access is potentially faster. Example

constant float objects[1234] = {
    1.0f, 2.0f, ...
};

You might be right I have ver 526 on desktop and 527 on the laptop. I'm might have expressed myself wrongly, it's not really important for me if it's accessible by multiple work-items, I just need it accessible by the functions declared within the kernel file, I'll update my post to show the code and how I would like it to work — Matt, Jan 29 '23 at 10:41
I see. You can't have `float* objects;` standing there alone in the wild. Remove that, instead, add `__global float *objs` as a second function parameter to `getDistCol`, and in `mkernel`, remove `objects = objs;` and call `getDistCol(position, objs);` with the now two parameters. You can hand global pointers over to functions. — ProjectPhysX, Jan 29 '23 at 17:54

Static variable in OpenCL C

1 Answers1