2

我第一次开始使用 OpenCL，我正在尝试优化缩减内核。内核采用大小为宽乘长像素的浮点数方格（数据表示灰度图像的亮度值）。内核沿每一列求和，并将每一列的总和返回到输出数组。

/* 
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 快，这让我觉得很奇怪。内核有问题吗？CPU最多只能同时执行大约8个线程，怎么能更快呢？

这个项目的代码在这里（一个 Xcode 项目），https://github.com/danieljfarrell/Xcode-4-OpenCL-Example。

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

执行时间作为图像大小的函数。

更新

通过查看使用 Instruments.app 运行内核所花费的时间，简单地暗示了 CPU 和 GPU 的时序， 显示 Instruments.app 使用界面的屏幕截图，其中突出显示了内核的执行时间。

4

4 回答 4

4

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

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

我没有遍历 height 参数，但您可以设置它或为输出数据创建足够的工作组（每个元素 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 在缩减内核中的表现应该没有问题。

于 2013-03-27T02:45:00.597 回答
3

实际上，在使用 OPenCL 时需要权衡取舍，这实际上可能会导致在将 OpenCL 用于某些任务时性能变慢。

也许您可以尝试增加在 GPU 上实际执行的工作，从而超过设置上下文的成本。

请注意，并非所有工作都可以从 OpenCL 中受益，而且大多数时候您必须进行一些测试才能确定您的应用程序是否可以从中受益。
请参阅：OpenCL 对哪些工作有利

  • 设置上下文和传输数据的开销

由于设置上下文和通过 PCI 总线传输数据的开销，您必须处理相当大的数据集才能看到使用 OpenCL 的任何好处。您开始看到好处的确切时间取决于 OpenCL 实现和所使用的硬件，因此您必须进行试验以了解您的算法执行速度有多快。一般来说，计算与数据访问的高比率和大量的数学计算对于 OpenCL 程序是有益的。

资料来源：MAC 的 OpenCL 编程指南

  • 小型作业/过于分散的作业在 CPU 上的性能更好

例如，如果您使用 OpenCL 执行的作业太小或太小，那么与使用 OpenCL 执行实际作业相比，您将花费更多时间来设置所有内容。

我添加的分段代码越多，OpenCL 代码就越慢。
[...]
3 件事会杀了你。
调用 OpenCL 的延迟。这意味着，调用 OpenCL 函数比调用“真正的 Java/C# 函数”需要更多时间。
其次，GPU 访问主计算机内存并将内容复制到其中需要相当长的时间。[...]

资料来源：OpenCL 中的当前问题 (2010)

  • 还有这个页面，基准测试人员在其中指出：

显然，您看到 GPU 在更大的大小值上优于 CPU，因为程序能够使用 GPU 提供的多个线程。在较小的大小值下，与 GPU 相关的访问时间可观，因此 CPU 执行速度更快。

资料来源：CPU 与 GPU 性能对比与 OpenCL（2011 年 10 月）

于 2013-03-24T14:07:37.327 回答
3

您需要将问题分解为多个内核运行。在 OpenCL 中拥有大循环（在这种情况下“高度”可能很大）是不好的，而且您需要展开任何循环。您可以使用“#pragma unroll X”自动执行此操作，其中“X”是循环运行的次数。

现在您的问题变得稍微困难​​了，因为您需要将其分解为多个内核运行。以下是步骤：

  1. 确定可以展开的最大数量，然后让每个线程只对那么多元素进行操作。
  2. 使用讨论的循环展开
  3. 这将为每列生成部分总和。
  4. 一次又一次地运行内核，直到你只有少量的总和（通过反复试验确定这个数字）
  5. 将 CPU 上的最终部分总和相加。注意：在最后的 CPU 求和步骤之前，不要将部分总和复制回 CPU。

还有其他策略，但这将是真正让 GPU 在此计算中获胜的第一步。

于 2013-03-26T16:24:01.553 回答
1

首先，我同意mfa，如果您将数据转置会更好。这样，您将按顺序从全局内存中读取数据（搜索银行冲突）。但这只是一件事。

另一件事是改变你的算法。您的方法当前的缺点 1）您的工作项目数量很少 - 等于图像的高度。2）长周期。

我建议你重写你的算法，让它真正平行下一个方式：例如，你必须对 512 个项目求和。然后运行具有 256 个工作项的工作组。每个工作项加起来 2 个值。例如，第一个将添加 v[1]=v[0]+v[1]，第二个 v[3] = v[2]+v[3] 等等。因此，在第一次操作之后，您将在奇数索引中获得对和。下一个周期您进行类似的程序，但只有 128 个工作项完成这项工作，因为您已经只有 256 个元素要处理。现在第一个工作项的唯一区别是 v[3] = v[1]+v[3]，第二个 v[7] = v[5] + v[7] 等等。这样你就有 1) O(logN) 复杂性而不是 O(N) 2) 你产生更多的项目做更少的工作。-> 从并行化中受益。

当然，您需要在每次写入后调用 barrier(...) 指令来同步工作组中工作项之间的计算。

为了进一步加快速度，首先每个工作组将其值从全局内存复制到本地内存，并使用本地内存执行计算。

一个问题，您可能想问：“如果我需要汇总太多值（比如 100000）并且我无法创建具有如此大量工作项的工作组”。在这种情况下，您进行部分求和，然后第二次运行内核以对这些部分求和。因为您在执行内核时无法在工作组之间进行同步。

这里更清楚的是代码。它将 2 x blockSize 值相加。希望我没有犯任何错误（实际上并没有编译这个）

// run Work group with local size = (BLK_SIZ, 1), global size = ( width, height )
__kernel void calc_sum(__global float* d_in, __global float* d_sums, const int rowLen)
{

    int our_row = get_global_id(1);
    int lx = get_local_id(0);
    int gr = get_group_id(0);

    __local float our_mem[(2*BLK_SIZ)];

    // copy glob -> loc mem
    our_mem[2*lx + 0] = d_in[gr*2*BLK_SIZ + 2*lx + 0];
    if(gr*2*BLK_SIZ + 2*lx + 0 >= rowLen)
        our_mem[2*lx + 0] = 0;

    our_mem[2*lx + 1] = d_in[gr*2*BLK_SIZ + 2*lx + 1];
    if(gr*2*BLK_SIZ + 2*lx + 1 >= rowLen)
        our_mem[2*lx + 1] = 0;

    // do the calculations

    int width = 2;
    int num_el = 2*BLK_SIZ / width;
    int wby2 = width>>1;
    for(int i = 0;i<7;++i) 
    {
        barrier(CLK_LOCAL_MEM_FENCE);
        if(lx < num_el)
        {
            int idx = width*(lx + 1) - 1;
            our_mem[idx] = min(our_mem[idx], our_mem[idx-wby2]);
        }
        width<<=1;
        wby2 = width>>1;
        num_el>>=1;
    }
    barrier(CLK_LOCAL_MEM_FENCE);
    // store res
    if(lx == 0) // choose some element from work group to actualy write the sum
    {
        d_sums[our_row] = our_mem[2*lx-1]; // sum is in last element
    }
}

还可以在互联网上搜索 blelloch / hillis Steele 并行前缀和算法。nVidia 也有很好的示例，其中包含有关并行前缀和算法的良好文档。它做的比你需要的要多，但与我描述的方法相同。

希望这可以帮助。

于 2013-04-04T21:44:04.357 回答