OpenCL 内核的最大细分级别

Question

我有一个关于我的总体理解的问题。对于这个问题，我构建了一个场景以使其尽可能简单。

假设： 我有一个包含 2 个变量的结构（x 和 y）。而且我在一个数组中的一个缓冲区中有成千上万个这种结构的对象。这些结构的初始值不同。但是以后总是应该将相同的算术运算应用于这些结构中的每一个。（所以这对 GPU 来说非常好，因为每个 worker 都在做完全相同的操作，只是在没有分支的情况下使用不同的值。）此外，CPU 根本不需要这个结构。所以只有在程序的整个结束时，所有值才应该存储回 CPU.

对这些结构的操作也是有限的！比方说，我们有 8 个可以应用的操作：

x + y，将结果存储在 x
x + y，将结果存储在 y
x + x，将结果存储在 x
y + y，将结果存储在 y
x * y，将结果存储在 x
x * y，将结果存储在 y
x * x，将结果存入x
y * y，将结果存储在 y

当为一个操作创建一个内核程序时，操作 1 的内核程序如下所示：

__kernel void operation1(__global float *structArray) 
{

    // Get the index of the current element to be processed
    int i = get_global_id(0) * 2;

    // Do the operation
    structArray[i] = structArray[i] + structArray[i + 1]; //this line will change for different operations (+, *, store to x, y)
}

当以某种顺序多次执行这些内核时，例如：操作 1、2、2、3、1、7、3、5.... 然后每次执行至少有一个全局内存读取操作和一个全局内存写入操作。但理论上，如果每个工作人员将其结构（x 和 y 值）存储在私有内存中，则执行速度将加快 50 倍左右。

是否可以这样做？:

__private float x;
__private float y;

__kernel void operation1(void) 
{       
    // Do the operation
    x = x + y; //this line will change for different operations (+, *, store to x, y)
}

为此，您首先需要存储值...例如如下所示：

__private float x;
__private float y;

__kernel void operationStore(__global float *structArray) 
{       
    int i = get_global_id(0) * 2;
    //store the x and y value from global to private memory
    x = structArray[i];
    y = structArray[i + 1];
}

当然，在程序的整个结尾，您需要将它们存储回全局内存，以便稍后再次将其推送到 CPU：

__private float x;
__private float y;

__kernel void operationStoreToGlobal(__global float *structArray) 
{       
    int i = get_global_id(0) * 2;
    //store the x and y value from private to global memory
    structArray[i] = x;
    structArray[i + 1] = y;
}

所以我的问题是：

我能否以某种方式设法在不同的内核调用期间将值存储在私有内存或本地内存中？如果是这样，我只会通过程序队列降低性能。
程序队列从一个内核切换到另一个内核需要多少个时钟周期？
这个改变内核的时机，内核大小是特定的吗？如果是这样：是取决于内核中的操作数量还是取决于缓冲区绑定的数量（重新绑定）
是否有经验法则，内核至少应该具有怎样的混合操作（按时钟周期计算）？

Answer 1

这是不可能的。您不能在 private 或 local 内存 space 中的“全局变量”中跨内核传递数据。您需要使用 global 内核参数来临时存储结果，从而将值临时写入显存并在下一个内核中从显存中读取。 “全局变量”唯一允许的内存 space 是 constant：有了它，您可以创建大型查找表。这些是只读的。 constant 尽可能将变量缓存在 L2 中。
可能有几千个。当您完成一个内核并启动另一个内核时，您就有了一个全局同步点。内核 1 的所有实例需要在内核 2 启动之前完成。
是的。它取决于全局范围、本地（工作组）范围、操作数（尤其是 if-else 分支，因为一个工作组可能比另一个工作组花费的时间长得多），但不取决于内核参数/缓冲区绑定的数量.全局大小越大，内核耗时越长，工作组之间的相对时间变化越小，内核变化（同步点）的相对性能损失越小。
更好的问题：内核的性能应该有多大？答：非常大，大约是 CUDA 核心/流处理器数量的 100 倍。

有一些技巧可以减少所需的全局同步点数。例如：如果一个内核可以组合来自不同内核的多个不同任务，则将两个内核压缩为一个。这里的示例：格子玻尔兹曼方法，两步交换与一步交换。

另一个常见的技巧是在视频内存中分配两次缓冲区。在偶数步骤中，从 A 读取并写入 B，在奇数步骤中则相反。避免从 A 读取并同时写入 A 的其他元素（引入竞争条件）。

OpenCL 内核的最大细分级别

maximum level of subdivisions of OpenCL Kernels

opencl