1

我通过尝试优化 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 时会有非常沉重的开销。

4

1 回答 1

2

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

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

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

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

于 2013-11-25T04:07:50.767 回答