GPU上的OpenCL性能问题

Question

我正在使用opencl在Raspberry Pi GPU（Videocore IV）中优化一些代码。我正在使用VC4CL实现，最大工作组大小为12。

但是，使用简单的内核，例如求和两个阵列，GPU的性能比CPU最差。

例如，对于以下内核：

#define GLOBAL_SIZE 12
#define LOCAL_SIZE  1
#define WIDTH       12*12*12
#define NTIMES      1000
__attribute__ ((reqd_work_group_size(LOCAL_SIZE, LOCAL_SIZE, LOCAL_SIZE)))
__kernel void int_sum(global const uint* A, global const uint* B, global uint* C)
{
  int g_id0 = get_global_id(0);
  int g_id1 = get_global_id(1);
  int g_id2 = get_global_id(2);
  int index = g_id0 + g_id1 * GLOBAL_SIZE + g_id2 * GLOBAL_SIZE * GLOBAL_SIZE;
  for(int k = 0; k < NTIMES; k++)
    C[index + k * WIDTH] = A[index + k * WIDTH] + B[index + k * WIDTH];
}

概括了两个1E6位置的阵列，CPU的性能要好得多...我试图将工作组更改为一维，并且还使用其他组合（6x6x6->全局大小，2x2x2->本地大小）。有什么暗示我可能做错了什么？

预先感谢。

Answer 1

我不熟悉此特定的GPU，但是在您的代码中，一些可能脱颖而出的危险信号：

• 这是整数Alu重型代码，而不是使用浮点操作。许多GPU根本没有为此进行优化。
• 我不会依靠编译器优化数组偏移计算；一个特别愚蠢的编译器可能会在每个循环迭代中为C[index + k * WIDTH] = A[index + k * WIDTH] + B[index + k * WIDTH];发射3个整数乘法。我将偏移保持在变量中，并在每次迭代中添加到它，无需乘以。
• 环路的1000倍率通常是更好的并行性的潜在来源。许多GPU在长期运行的内核中表现不佳。
• 内存访问模式似乎是最佳的。尝试安排这些内容，以使相邻的工作项目在组中访问相邻的内存位置。2x2x2本地规模似乎是一个特别差的选择。您是否尝试过12x1x1？
• 为什么您甚至以这种方式安排工作项目？看来您实际上是在计算i = 0..1000*12*12*12的C[i] = A[i] + B[i]。如何编写您的内核并在一个维度上提交1728000个工作项目？这节省了所有复杂的索引计算。

如果您可以从驱动程序中获得任何形式的反馈，那么GPU受（ALU，内存加载，线程调度等）的限制，这将有很大的帮助，可以选择在哪里寻找可以加快速度的方法。

Answer 2

除了其他所有人在评论中所说的话，根据RPI的OPENCL实现作者，GPU"内存访问速度"（CPU-GPU内存复制？）比CPU慢得多。。因此，像阵列总和这样的"算术上的光"内核将受到内存带宽的限制，并且比在CPU上要慢得多。再加上理论上的GPU GFLOPS并不比GPU高得多（24 vs 6）。

除非您有一些非常的计算重核，也可以完全矢量化，否则您可能会发现使用GPU根本不值得。