如何在 OpenCL 内核中与多个用户定义的函数共享全局变量（数组）

Question

我的 OpenCL 内核有问题。我正在尝试进行 Runge-Kutta 4 集成。我已经在 OpenGL 计算着色器中实现了它并且它可以工作，现在我想在 OpenCL 中实现它。

我认为我的问题与不知道如何在所有函数调用中正确共享全局数组的单个实例有关，因为现在我必须将指向数组的指针作为每个函数的参数发送调用，在我看来，这实际上在这些函数中创建了一个本地副本，因为我当前的实现适用于小型数据集，但不适用于大型数据集（它们抛出 CL_OUT_OF_RESOURCES）。

在我的计算着色器中，我将所有全局数组声明为：

layout(std430, binding=0) buffer pblock { coherent volatile restrict vec4 mcPosition[]; };

layout(std430, binding=1) buffer vblock { coherent volatile restrict vec4 mcVelocity[]; };

而且我可以在我的函数中很好地使用它们：

vec4 calculateAcceleration(int numPoints, int step, ...) {...}

void rk4Step(int numPoints, int index, float timeStepToUse, ...) {...}

void calculateError(int index) {...}

但是在 OpenCL 实现上，我知道如何做的唯一方法是这样的（非常精简的示例）：

void rk4Step(
    const __constant int* numPoints,
    const int index,
    const float timeStepToUse,
    const bool calculateHalfTimeStep,
    const __constant float* squaredSofteningFactor,
    const __constant float* gravitationalConstant,
    __global float4* kvel,
    __global float4* dydx,
    __global float4* kpos,
    __global float4* mcPositionHalf,
    __global float4* mcVelocityHalf,
    __global float4* mcPositionFull,
    __global float4* mcVelocityFull
    )
{
    ...

    // Actual time step
    if(!calculateHalfTimeStep)
    {
        mcVelocityFull[index] += (kvel[index] + (2.0f*kvel[index+numPoints[0]]) + (2.0f*kvel[index+numPoints[0]*2]) + kvel[index+numPoints[0]*3]) * (1.0f/6.0f);
        mcPositionFull[index] += (kpos[index] + (2.0f*kpos[index+numPoints[0]]) + (2.0f*kpos[index+numPoints[0]*2]) + kpos[index+numPoints[0]*3]) * (timeStepToUse/6.0f);
    }
    else
    {
        mcVelocityHalf[index] += (kvel[index] + (2.0f*kvel[index+numPoints[0]]) + (2.0f*kvel[index+numPoints[0]*2]) + kvel[index+numPoints[0]*3]) * (1.0f/6.0f);
        mcPositionHalf[index] += (kpos[index] + (2.0f*kpos[index+numPoints[0]]) + (2.0f*kpos[index+numPoints[0]*2]) + kpos[index+numPoints[0]*3]) * (timeStepToUse/6.0f);
    }
}

void calculateError(const int index, __global float4* scale)
{
    float partialError = 0.0f;
    partialError = fmax(partialError, fabs(deltaPos[index].x / scale[index].x));
}

// Adaptive step 4th order Runge-Kutta
__kernel
void main( const __constant float* timeStep, const __constant float* accuracy, const __constant int* maxSteps,
    __global float4* mcPosition, __global float4* mcVelocity, __global float4* scale)
{
    // Scaling used to monitor accuracy
    scale[index] = calculateAcceleration(bi, index, numPoints, 1, false,
        squaredSofteningFactor, gravitationalConstant,
        mcPositionHalf, mcPositionFull, kvel);

    scale[index] = fabs(mcVelocity[index]) + fabs(scale[index] * timeStep[0]);

    for(int step=1; step<=maxSteps[0]; ++step)
    {
        // Take two half steps
        rk4Step(numPoints, index, timeStep[0], true,
            squaredSofteningFactor, gravitationalConstant,
            mcPosition, mcVelocity);
        rk4Step(numPoints, index, timeStep[0], true,
            squaredSofteningFactor, gravitationalConstant,
            mcPosition, mcVelocity);

        // Take one full step
        timeStep[0] *= 2.0f;
        rk4Step(numPoints, index, timeStep[0], false,
            squaredSofteningFactor, gravitationalConstant,
            mcPosition, mcVelocity);

        // Evaluate accuracy
        calculateError(index, accuracy, scale, deltaPos);
    }
}

如您所见，不同之处在于，在计算着色器版本中，我可以在文件顶部声明共享全局数组，并在我的任何一个函数中使用它们。

但在 OpenCL 内核版本中，我必须将这些数组作为参数传递给每个函数调用，对于大型数据集，这会给我一个 CL_OUT_OF_RESOURCES 错误。

我认为我的问题与这样一个事实有关，即即使我将数组声明为全局数组，每个函数调用都会尝试制作数组的本地副本，但也许我错了。我通过阅读文档来假设这一点，这个问题指出了同样的事情：

How many copies of a global variable declared inside an opencl kernel function is maintained in the global address space

所以我的问题是：我如何真正在用户定义的函数和我的 OpenCL 内核之间共享一个全局数组？

Answer 1

你提到的数组是作为指针传递的，没有理由期待整个数组的本地副本，还有 __constant 参数会停止写入并在读取 __constant 时复制 -只要。 no-local-copy 的主要原因可能是 gpu-opencl 实现没有堆栈。人们编写虚假堆栈来实现虚假递归，但即使那样也不能大于主机代码中定义的大小。

你什么时候得到"CL_OUT_OF_RESOURCES"？更改 __constant 缓冲区大小或 __global 大小后？通常 __constant 每个 GPU 只有 50-100 kB，而 __global 可以大到每个 gpu 每个缓冲区的视频内存的 1/4。 __constant 个参数的偶数是有限的。您可以将多个常量数组连接成单个常量数组以消除它。请查询常量和全局的常量内存限制。使用 CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE.

从 clGetDeviceInfo 开始

其他情况：

堆碎片 ---> 没有大数组。只有较小的可以分配为缓冲区。您是否发送使用所有 vram（或常量 vram）的并发内核？
本地工作组大小大于设备大小（例如：amd 在 gpu 上有 256，nvidia 有 1024）（它至少是全局大小的分频器）
标量寄存器太多，每个线程或每个向量寄存器太多 smx/cu。

测试：

每组有 1024 个线程。
函数中至少有 7 个 float4 变量。
每个 float4 是 16 个字节。
如果每个函数都使用这些单个变量（从任何来源读取），
每个 smx 需要 112 kB，这比它拥有的 (48 kB).
每个线程仅用于 float4s 就需要 112 个字节。您也使用标量变量。你可以用分析器检查一下。

帮助：

您可以 change/reorder 内核和函数中的东西，因此它在任何时候都需要更少的寄存器。仅在使用前声明一些内容。不是一开始。您也可以在完成工作后重新使用寄存器（例如使用 v1 而不是 v2、v3，以 v2、v3、v4 ... 的名义再次使用 v1）。
减小本地工作组大小，使每个 smx 的线程减少意味着每个 smx 的寄存器使用量减少。即使每个线程的使用也很重要，但只是为了性能。
有时低至 32（或 64）-local-work-group-size 可能是有利的，尽管有一半的内核空闲，以获得每个线程更多的内存 space。
内联函数也会增加注册压力。也许你应该降低展开和内联的级别，然后重试。
从 rk4Step 参数中删除 const(不是 __constant) 关键字。也许那些是预先分配在每个线程的 __constant 内存 space 中（并非不可能）

如何在 OpenCL 内核中与多个用户定义的函数共享全局变量（数组）

How to share global variables (arrays) in an OpenCL kernel with several user defined functions

arrays

kernel

global-variables

local-variables

opencl