OpenCL中的管道cl::Pipe的使用

本例只是其中一个简单的使用，在OpenCL中操作管道的函数有很多，这里只用了其中2个读写函数。此例功能是计算12800个整数的积的和。例子不是很好但我实在想不到实用的例子了，因为我做图像处理方面的应用用不到管道来传递数据。代码如下，这里不和CPU处理时间进行比较，这种简单的运算一般是CPU更快。程序运行环境是VS2015、OpenCL306，CPU是AMD A4-9125处理器，GPU是核芯显卡。下面的代码还使用了OpenCV中的getTickCount()函数用于计时。

string kernelStr = R"(
    kernel void add(global const int* arr1, global const int* arr2, write_only pipe int myPipe)
    {
        int id = get_global_id(0);
        int z = arr1[id] * arr2[id];
        write_pipe(myPipe, &z);
    }

    global atomic_int volatile summary = ATOMIC_VAR_INIT(0);
    global atomic_int volatile counter = ATOMIC_VAR_INIT(0);
    kernel void mul(read_only pipe int myPipe, global int* output)
    {
        int size = get_global_size(0);
        int z = 0;
        int status = read_pipe(myPipe, &z);
        if (status == 0)
        {
            atomic_fetch_add(&summary, z);
            atomic_fetch_add(&counter, 1);
        }
        if (size == atomic_load(&counter))
        {
            *output = atomic_load(&summary);
        }
    })";

void main()
{
    cl::Program program(kernelStr);
    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;
    }
    cl::KernelFunctor<cl::Buffer&, cl::Buffer&, cl::Pipe&> myAddKernel(program, "add");
    cl::KernelFunctor<cl::Pipe&, cl::Buffer&> myMulKernel(program, "mul");

    vector<int> inta(12800, 3);
    vector<int> intb(12800, 3);
    vector<int> ints(1, 0);
    cl::Buffer a(inta.begin(), inta.end(), true);
    cl::Buffer b(intb.begin(), intb.end(), true);
    cl::Buffer s(ints.begin(), ints.end(), false);
    cl_mem cpipe = clCreatePipe(cl::Context::getDefault()(), CL_MEM_HOST_NO_ACCESS, sizeof(int), 12800, 0, 0); /* 1 */
    cl::Pipe myPipe(cpipe, true);

    int64 t1, t2;
    t1 = getTickCount();

    myAddKernel(cl::EnqueueArgs(cl::NDRange(12800), cl::NDRange(128)), a, b, myPipe);
    myMulKernel(cl::EnqueueArgs(cl::NDRange(12800), cl::NDRange(128)), myPipe, s);
    cl::copy(s, ints.begin(), ints.end());
    cout << "和：" << ints[0] << endl;

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

    int c;
    cin >> c;
}

在核函数mul(...)中我们使用了atomic_int类型累加求和。这是低效的写法，因为原子操作是顺序操作的无法发挥GPU的并行计算的优势。另外，在注释/* 1 */处我使用的是C语言版的clCreatePipe(...)创建的管道对象，没有直接使用C++的cl::Pipe类创建。原因是在我的AMD电脑上用 cl::Pipe(sizeof(int), 12800) 构造程序会报错（不过在Intel的CPU核芯显卡上运行不会报错），两种创建管道的方法的区别是C++版运行库设置了CL_MEM_READ_WRITE属性。另外还要注意要先执行myAddKernel再执行myMulKernel，顺序不能反，否则会导致read_pipe(...)返回错误（请忽略这里变量命名的错误，写代码的时候没注意add和mul搞反了）。下面是程序运行结果的截图，可知执行结果是对的（12800*3*3=115200）：