OpenCL 内核的最大细分级别

如何解决OpenCL 内核的最大细分级别

我有一个总体上的理解问题。对于这个问题，我建立了一个尽可能简单的场景。

让我们说： 我有一个由 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,3,1,7,5.... 然后我每次执行至少一个全局内存读取操作和一个全局内存写入操作。但理论上，如果每个 worker 都将其结构（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 (+,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;
}

所以我的问题是：

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

解决方法

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

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

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

OpenCL 内核的最大细分级别

如何解决OpenCL 内核的最大细分级别

解决方法

相关推荐