OpenCL C 中的静态变量

Question

我正在使用 openCL 从头开始​​编写一个渲染器，我的内核出现了一些编译问题，并出现以下错误： CL_BUILD_PROGRAM : error: program scope variable must reside in constant address space static float* objects;

问题是这个程序在我的桌面上编译（使用 nvidia 驱动程序）并且不能在我的笔记本电脑上运行（使用 nvidia 驱动程序），而且我在另一个项目中有完全相同的内核文件，在两台计算机上都可以正常工作...... 有谁知道我可能做错了什么？

作为澄清，我正在编写一个 raymarcher，它的内核采用一个浮点数组中“编码”的对象列表，该数组在程序中需要很多，这就是为什么我需要 hole 内核可以访问它。

这是简化的内核代码：

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);

getDistCol(float3 position) 被很多函数调用的地方，我想避免将我的浮点缓冲区传递给需要调用 getDistCol() 的每个函数...

Answer 1

OpenCL C 不允许您在内核外部声明并跨内核使用的“静态”变量。一些编译器可能仍然容忍这一点，其他人可能不会。 Nvidia 最近在驱动程序更新中将他们的 OpenCL 编译器从 LLVM 3.4 更改为 NVVM 7，因此您的台式机/笔记本电脑 GPU 上可能有 2 种不同的编译器。

在您的情况下，解决方案是将 global 内核参数指针移交给函数：

float4 getDistCol(float3 position, __global float *objects) {
    // ...
}
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
    // ...
}

野外的孤独变量只允许作为constant内存空间，这对大表很有用。它们缓存在 L2$ 中，因此只读访问可能更快。例子

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