我试图为一个问题找到最佳的工作组规模，但我发现了一些我无法为自己辩护的东西。
这些是我的结果：
交换轴导致执行速度提高一倍。为什么 ！？
顺便说一句，我使用的是 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];
}
我同意@Thomas，这很可能取决于您的内核。最有可能的是，在第二种情况下，您以合并方式访问内存和/或充分利用内存事务。
合并：当线程需要访问内存中的元素时，硬件会尝试在尽可能少的事务中访问这些元素，即如果线程 0 和线程 1 必须访问连续的元素，则只有一个事务。
充分利用内存事务：假设您有一个 GPU 可以在一次事务中获取 32 个字节。因此，如果您有 4 个线程需要每个线程获取一个 int，那么您只使用事务获取的数据的一半；你浪费了其余的（假设一个 int 是 4 个字节）。
为了说明这一点，假设您有一个 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) 尝试 (8, 32) 做得更好。