OpenCL规约算法例子

本文给出一个规约算法求数组的和的例子。本例子求128000个整数的和。其实我想用浮点数做例子的但是我的电脑用atomic_float类型编译会报错。运算过程是每个工作组先把数据加载到局部内存中，工作组的大小是128，然后再求和，把结果累积到全局变量中。实际运行对比发现GPU的效率不如CPU直接求和。下述算法运行环境是VS2015、OpenCL3，CPU是AMD A4-9125，显卡就是CPU自带的核芯显卡。

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

string strKernel = R"(
    global volatile atomic_int fullsum = ATOMIC_VAR_INIT(0);
    global volatile atomic_int accumu = ATOMIC_VAR_INIT(0);
    kernel void addsum(global int* input, int count, global int* output)
    {
        local int temp[128];
        int i = get_global_id(0);
        int j = get_local_id(0);
        temp[j] = input[i];
        work_group_barrier(CLK_LOCAL_MEM_FENCE);
        if (j == 0)
        {
            int summary = 0;
            for (int k = 0; k < 128; k++)
            {
                summary += temp[k];
            }
            atomic_fetch_add(&fullsum, summary);
            atomic_fetch_add(&accumu, 128);
        }
        if (count == atomic_load(&accumu))
        {
            output[0] = atomic_load(&fullsum);
        }
    })";

void 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;
    }
    auto kernel = cl::KernelFunctor<cl::Buffer, int, cl::Buffer>(program, "addsum");

    int64 t1, t2;
    t1 = getTickCount();

    vector<int> number(128000, 2);
    vector<int> summary(1, 0);
    cl::Buffer input(number.begin(), number.end(), true);
    cl::Buffer output(summary.begin(), summary.end(), false);

    kernel(cl::EnqueueArgs(cl::NDRange(128000), cl::NDRange(128)), input, 128000, output);

    cl::copy(output, summary.begin(), summary.end());
    cout << summary[0] << endl;

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

    int c;
    cin >> c;
}

上述代码的核函数里变量accumu是用来判断求和是否已经完成，也即工作组间同步，如果accumu和count相等则完成求和可以向输出变量写数据。向输出output[0]写数据可能会执行很多次，不过不会导致结果出错。在Release版下测试，GPU速度远低于CPU，不太懂怎么能让GPU快于CPU。需要注意代码里捕捉异常的地方，需要开启OpenCL的异常，OpenCL默认是用返回错误代替抛异常的。getTickCount()是OpenCV的计时函数不需要可以删掉。

另外，此代码在用Qt调试过程中出现使Qt程序无法关闭的情况，数据正常也没有死循环，暂时不知道为什么。在控制台程序中调试程序可以正常退出。