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 必须访问连续元素，则只有一个事务。

完全使用内存事务 :假设您有一个 GPU，可以在一个事务中获取 32 个字节。因此，如果您有 4 个线程需要每个线程获取一个 int，那么您只使用事务获取的数据的一半；你浪费了剩下的(假设一个 int 是 4 个字节)。

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

每个工作项处理一列，一次一个地循环遍历每个列元素。

每个工作项处理一行，一次一个地循环遍历每个行元素。

这可能违反直觉，但第一个解决方案将能够进行合并访问，而第二个则不能。原因是当第一个工作项需要访问第一列中的第一个元素时，第二个工作项将访问第二列中的第一个元素，依此类推。这些元素在内存中是连续的。第二种解决方案不是这种情况。

现在，如果您采用相同的示例，并应用解决方案 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) try (8, 32) 更好。