OpenCL规约求和方法

本文给出一个规约算法求数组的和的例子。本例子求128128个整数的和。运算过程是每个工作组先把数据加载到局部内存中，工作组的大小是64，如果你的电脑支持更多也可以设置的更大一些，然后再求和，把结果放到到全局内存中。由于OpenCL没有全局内存的内存屏障，所以只能用CPU端同步全局内存。即每计算一次就用CPU同步一次全局内存。每次求和都能把输入缩小1/64，反复几次就可以把结果缩小到1，此时就是结果。实际运行对比发现GPU的效率不如CPU直接求和。下述算法运行环境是VS2019、OpenCL3，CPU是Intel i5，显卡是英伟达的。

本例不需要头文件，下面是CPP文件：

string strKernel = R"(
    kernel void addsum(global float* input, int count)
    {
        local float localMemory[64];
        int gi = get_global_id(0);
        int wi = get_group_id(0);
        int li = get_local_id(0);

        localMemory[li] = (gi < count ? input[gi] : 0);
        work_group_barrier(CLK_LOCAL_MEM_FENCE);

        for (int space = 32; space >= 1; space /= 2)
        {
            if (li < space) /* 分成两份，前面一份加后面一份 */
            {
                localMemory[li] += localMemory[li + space];
            }
            work_group_barrier(CLK_LOCAL_MEM_FENCE);
        }
        input[wi] = localMemory[0];
    })";

int main()
{

    cl::Program program(strKernel);
    try
    {
        program.build("-cl-std=CL2.0");
    }
    catch (...)
    {
        cl_int buildErr = CL_SUCCESS;
        auto buildInfo = program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(&buildErr);
        for (auto& pair : buildInfo)
        {
            std::cerr << pair.second << std::endl << std::endl;
        }
        return 0;
    }
    auto kernel = cl::KernelFunctor<cl::Buffer, int>(program, "addsum");

    /* 准备初始数据 */
    vector<float> number(128128, 1);
    number[2000] += 10000;
    number[4000] += 10000;

    int64 t1, t2;
    t1 = getTickCount();

    float sum = 0;
    for (int i = 0; i < 128128; i++)
    {
        sum += number[i];
    }
    cout << "sum1=" << sum << endl;

    t2 = getTickCount();
    cout << "CPU1(ms):" << (t2 - t1) / getTickFrequency() * 1000 << endl;

    cl::Buffer input(number.begin(), number.end(), false);

    t1 = getTickCount();

    vector<float> summary(1);
    /* 结果数量是输入的1/64倍，要考虑余数 */
    for (int i = 128128; i >= 2; i = i / 64 + bool(i % 64))
    {
        /* 输入global-size必须能被64整除，向上取64的整倍数 */
        int inputCount = (i / 64 + bool(i % 64)) * 64;
        kernel(cl::EnqueueArgs(cl::NDRange(inputCount), cl::NDRange(64)), input, i);
    }
    cl::copy(input, summary.begin(), summary.end());
    cout << "sum2=" << summary[0] << endl;

    t2 = getTickCount();
    cout << "GPU1(ms):" << (t2 - t1) / getTickFrequency() * 1000 << endl;

    return 0;
}

程序的输出截图如下。这个截图过时了，上面的代码是新改的。经过测试，CPU求和时间在0.3ms左右，GPU求和在1ms左右，GPU比CPU更慢一些。