OpenCL 慢——不知道为什么答案

【问题标题】：OpenCL slow -- not sure whyOpenCL 慢——不知道为什么
【发布时间】：2013-11-24 02:39:22
【问题描述】：

我正在通过尝试优化 mpeg4dst 参考音频编码器来自学 OpenCL。我通过在 CPU 上使用矢量指令实现了 3 倍的加速，但我认为 GPU 可能会做得更好。

我的首要改进领域是计算 OpenCL 中的自相关向量。 CPU代码是：

for (int i = 0; i < NrOfChannels; i++) {
    for (int shift = 0; shift <= PredOrder[ChannelFilter[i]]; shift++)
        vDSP_dotpr(Signal[i] + shift, 1, Signal[i], 1, &out, NrOfChannelBits - shift);
}
NrOfChannels = 6
PredOrder = 129
NrOfChannelBits = 150528.

在我的测试文件中，这个函数大约需要 188 毫秒才能完成。

这是我的 OpenCL 方法：

kernel void calculateAutocorrelation(size_t offset,
                                 global const float *input,
                                 global float *output,
                                 size_t size) {
size_t index = get_global_id(0);
size_t end = size - index;
float sum = 0.0;

for (size_t i = 0; i < end; i++)
    sum += input[i + offset] * input[i + offset + index];

output[index] = sum;
}

它是这样称呼的：

gcl_memcpy(gpu_signal_in, Signal, sizeof(float) * NrOfChannels * MAXCHBITS);

for (int i = 0; i < NrOfChannels; i++) {
    size_t sz = PredOrder[ChannelFilter[i]] + 1;
    cl_ndrange range = { 1, { 0, 0, 0 }, { sz, 0, 0}, { 0, 0, 0 } };

    calculateAutocorrelation_kernel(&range, i * MAXCHBITS, (cl_float *)gpu_signal_in, (cl_float *)gpu_out, NrOfChannelBits);
    gcl_memcpy(out, gpu_out, sizeof(float) * sz);
}

根据 Instruments 的说法，我的 OpenCL 实现似乎需要大约 13 毫秒，内存复制开销大约为 54 毫秒 (gcl_memcpy)。

当我使用更大的测试文件时，2 声道音乐 1 分钟 vs 6 声道 1 秒，而 OpenCL 代码的测量性能似乎相同，CPU 使用率下降到 50% 左右整个程序需要大约 2 倍的时间来运行。

我在 Instruments 中找不到导致这种情况的原因，而且我还没有阅读任何内容表明我应该预计在进出 OpenCL 时会有非常沉重的开销。

【问题讨论】：

您的系统规格是什么？显卡？操作系统？一般来说，内核启动开销在离散 GPU 上非常大。在 Windows Vista 或更高版本上，您可以获得的最糟糕的启动开销是独立 GPU（由于驱动程序模型）。此外，您的 OpenCL 内核似乎运行的线程相对较少（GPU 的 128 个线程确实很低）。理想情况下，您的问题应该有数千个不同的线程（ndrange）。最重要的是，内存复制在 PciE 总线上非常昂贵。尝试编写它，以便在所有内核调用之后复制整个结果，而不是每次循环迭代的 memcpy。
我在 OSX 10.9 上使用带有 NVidia GT 650M 显卡的 Macbook Pro。我已经删除了第二个 gcl_memcpy 和循环，并且我正在学习 cl_ndrange 以尝试在 GPU 上获得更多线程。但是，从我的基准测试中，我看到执行内核需要 10 毫秒，而在 gcl_memcpy 中执行需要 50 毫秒（这应该是双向的 PCIe 传输。）这应该会给我一个很好的加速，~60 毫秒 vs~190 毫秒。相反，我看到时间“消失”了，CPU 仅使用了 40%，整个过程花费的时间几乎是以前的两倍。

标签： performance macos opencl grand-central-dispatch vdsp

【解决方案1】：

如果我正确读取了您的内核代码，则每个工作项都在迭代从其位置到最后的所有数据。这不会是有效的。一方面（也是主要的性能问题），内存访问不会被合并，因此不会处于完整的内存带宽。其次，由于每个工作项的工作量不同，工作组内会出现分支分歧，这会使一些线程空闲等待其他线程。

这似乎与归约问题有很多共同点，我建议阅读“并行归约”以获取有关并行执行此类操作的一些提示。

要查看内存是如何被读取的，请计算出 16 个工作项（例如 global_id 0 到 15）将如何读取每个步骤的数据。

请注意，如果工作组中的每个工作项都访问相同的内存，则硬件可以进行“广播”优化。因此，只需颠倒循环的顺序就可以改善事情。

【讨论】：