OpenCL - 工作组轴可交换吗？

Question

我正在努力为一个问题寻找最佳的工作组规模，我想出了一些我无法为自己辩解的事情。

这是我的结果：

• GlobalWorkSize {6400 6400 1}，WorkGroupSize {64 4 1}，时间（毫秒）= 44.18
• GlobalWorkSize {6400 6400 1}，WorkGroupSize {4 64 1 }，时间（毫秒）= 24.39

交换轴引起两倍的执行速度。为什么！？

顺便说一下，我使用的是AMD GPU。

感谢:-)

编辑： 这是内核（一个简单的矩阵转置）：

__kernel void transpose(__global float *input, __global float *output, const int size){ 
    int i = get_global_id(0); 
    int j = get_global_id(1); 
    output[i*size + j] = input[j*size + i]; 
} 

Answer 1

我@Thomas同意，这很可能取决于你的内核。最有可能的是，在第二种情况下，您可以通过合并的方式访问内存和/或充分利用内存事务。

聚结：当线程需要访问元素在存储器中的硬件试图访问在这些元件尽可能少的交易，即如果线程0和线程1具有访问连续元素，将仅存在一个事务。

完全使用内存事务：假设您有一个在一次事务中获取32个字节的GPU。因此，如果您有4个线程需要获取一个int，则您只使用事务获取的一半数据;你浪费其余的（假设int是4个字节）。

为了说明这一点，假设您有一个n乘n的矩阵来访问。你的矩阵在主要行中，并且你使用在一个维度中组织的n个线程。您有两种可能性：

1. 每个工作项都负责一个列，每次循环一个列元素。
2. 每个工作项都需要一条线，每次循环一个线条元素。

这可能是违反直觉的，但第一个解决方案将能够使第二个解决方案不会成为聚结点。原因是当第一个工作项需要访问第一列中的第一个元素时，第二个工作项将访问第二个列中的第一个元素，依此类推。这些元素在内存中是连续的。第二种解决方案并非如此。

现在，如果你采用相同的例子，并应用解决方案1，但这次你有4个工作项目，而不是n和我刚刚讲过的同一个GPU，因为你很可能将时间增加了2倍你会浪费一半的内存交易。

编辑：现在你发布你的内核，我看到我忘了提到别的东西。

对于你的内核，似乎选择（1,256）或（256,1）的本地大小总是一个不好的选择。在第一种情况下，需要256个事务来读取一列（每次读取32个字节，其中只有4个将被使用 - 请记住前面例子中的相同GPU），而输出中需要32个事务写入：您可以在一个事务中写8个浮动，因此32个事务可以写入256个元素。

这与工作组大小为（256,1）的问题相同，但是这次使用32个事务读取和256个写入。

那么为什么第一个尺寸更好？这是因为有一个缓存系统，可以减轻读取部分的不良访问。因此，写入部分的大小（1,256）是有利的，并且高速缓存系统处理不太好的读取部分，减少了必需的读取事务的数量。

请注意事务处理数量总体上减少（考虑到NDRange中的所有工作组）。例如，第一个工作组发布256个事务，读取第一列的256个第一个元素。第二个工作组可能只是进入缓存以检索第二列的元素，因为它们是由第一个工作组发布的事务（32字节）提取的。

现在，我几乎可以肯定，你可以做得比（1,256）尝试（8,32）更好。