macos - OpenCL:如何优化缩减内核(列总和)，目前 CPU 速度更快

Question

我是第一次开始使用 OpenCL，我正在尝试优化缩减内核。内核采用大小为宽度像素的 float 正方形网格(数据表示灰度图像的亮度值)。内核对每一列求和并将每一列的总和返回到输出数组。

/* 
input  -- "2D" array of floats with width * height number of elements
output -- 1D array containing summation of column values with width number of elements
width  -- number of elements horizontally
height -- number of elements vertically

Both width and height must be a multiple of 64.
*/
kernel void sum_columns(global float* input, global float* output,  int width, int height)
{        
    size_t j = get_global_id(0);
    float sum = 0.0;
    int i;
    for(i=0; i<height; i++) {
        sum += input[i + width*j];
    }
    output[j] = sum;
}

OpenCL 应该同时执行每列求和，因为我将全局维度设置为数据中的列数。我在 MacOS 上使用了 Instruments.app，并计算了在 CPU 和 GPU 上执行内核 1000 次迭代所花费的时间。这可以通过将设备指定为 CL_DEVICE_TYPE_CPU 或 CL_DEVICE_TYPE_GPU 来完成。

性能不好!事实上，CPU 一直比 GPU 快，这让我觉得很奇怪。内核有问题吗？最多只能同时执行8个线程左右的CPU怎么可能更快呢？

此项目的代码在这里(Xcode 项目)，https://github.com/danieljfarrell/Xcode-4-OpenCL-Example .

这是我增加数据大小时的计时结果，

更新

通过查看使用 Instruments.app 运行内核所花费的时间，可以简单地隐含 CPU 和 GPU 的时序，

Answer 1

尝试的一个简单改进是使输入常量内存而不是全局内存。您需要在创建缓冲区时使用 CL_MEM_READ_ONLY 对其进行设置。当我只将该参数更改为 __constant 时，我使用的分析器似乎喜欢你的内核。

另一种选择是转置输入矩阵，这样您就不会尝试读取内存列。我制作了一个内核，它使用整个工作组对一行数据求和并在输出中产生一个条目。 __constant 参数也对这个内核有很大帮助，使它成为我运行的四次试验中唯一受 ALU 绑定(bind)而不是全局获取绑定(bind)的运行。

我没有循环遍历高度参数，但您可以设置它或为您的输出数据创建足够的工作组(每个元素 1 个)。

kernel void sum_rows(__constant  float* input, global float* output,  int width, int height)
{        
    int gid = get_local_id(0);
    int gsize = get_local_size(0);
    local float sum[64]; //assumes work group size of 64
    sum[gid] = 0;
    int i;
    int rowStart = width * get_group_id(0);
    for(i=gid; i<width; i+=gsize) {
        sum[gid] += input[rowStart + i];
    }
    barrier(CLK_LOCAL_MEM_FENCE);
    if(gid == 0){
        for(i = 0;i<64;i++){
            sum[0] += sum[i];
        }
        output[get_group_id(0)] = sum[0];
    }
}

除此之外，我建议查看主机级优化。有了足够大的数据集，gpu 在缩减内核中胜过 cpu 应该没有问题。