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, global float* output)
    {
        local float localMemory[64];
        int gs = get_global_size(0);
        int gi = get_global_id(0);
        int wi = get_group_id(0);
        int li = get_local_id(0);

        localMemory[li] = (gi < gs ? 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);
        }
        output[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, cl::Buffer>(program, "addsum");

    int64 t1, t2;
    t1 = getTickCount();

    vector<float> number(128128, 1); /* 输入要能整除64，因为下面的本地组大小是64 */
    vector<float> summary;
    for (int i = 128128; i >= 2; i = i / 64 + bool(i % 64))
    {
        /* 结果数量是输入的1/64倍，要考虑余数 */
        int resultCount = i / 64 + bool(i % 64);
        /* 结果数量也要能被64整除，要考虑余数 */
        resultCount = (resultCount / 64 + bool(resultCount % 64)) * 64;
        summary.resize(resultCount, 0);

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

        kernel(cl::EnqueueArgs(cl::NDRange(number.size()), cl::NDRange(64)), input, output);
        cl::copy(output, summary.begin(), summary.end());
        number = std::move(summary);
    }
    cout << "sum=" << number[0] << endl;

    return 0;
}

程序的输出截图如下：