OpenCL中局部变量和全局变量间的异步复制内存

本文主要是async_work_group_copy(...)和wait_group_events(...)的使用范例。展示了从全局内存到局部内存加载数据和从局部内存到全局内存写入数据。这系列内置函数可以用来取代直接赋值式的代码。不过我没有仔细对比过使用异步复制和直接赋值那个效率更高。本系列函数中还有定步长跨元素异步复制的函数，这里没有展示。

函数wait_group_events(...)用来等待异步操作执行完成。你可以在调用此函数之前执行其他操作以节省时间。以下代码运行环境是VS2017、OpenCL3。显卡是Intel Core i5的核芯显卡，main.cpp的代码如下：

string kernelStr = R"(
    #define GROUP_SIZE 128
    kernel void copyData(global const float* input, global float* output)
    {
        local float temp[GROUP_SIZE];
        int index = get_group_id(0);
        int start = index * GROUP_SIZE;
        event_t copyEvent = async_work_group_copy(temp, &input[start], GROUP_SIZE, 0);
        // 等待完成，这里还可以执行其它操作
        int id = get_local_id(0);
        wait_group_events(1, &copyEvent);
        work_group_barrier(CLK_LOCAL_MEM_FENCE);
        temp[id] *= 2;
        copyEvent = async_work_group_copy(&output[start], temp, GROUP_SIZE, 0);
        // 等待完成，这里还可以执行其它操作
        wait_group_events(1, &copyEvent);
    })";

int 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 1;
    }
    auto kernel = cl::KernelFunctor<cl::Buffer, cl::Buffer>(program, "copyData");
    vector<float> a(12800, 4);
    vector<float> b(12800, 0);
    cl::Buffer inputa(a.begin(), a.end(), true);
    cl::Buffer outputb(b.begin(), b.end(), false);

    int64 t1, t2;
    t1 = getTickCount();

    kernel(cl::EnqueueArgs(cl::NDRange(12800), cl::NDRange(128)), inputa, outputb);
    cl::copy(outputb, b.begin(), b.end());
    cout << b[0] << endl;

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

    int c;
    cin >> c;
    return 0;
}

此测试的输入是12800个元素的内容为4的数组（inputa），执行后将输入乘以2后输出12800个元素的内容为8的数组（outputb）。此例重要描述代码实现，因此不比较代码执行效率。另外，此例中 temp[id] *= 2; 代码后面没调用work_group_barrier(...)函数同步工作组数据，不知道为什么这里不需要同步等待temp数组全部更新数据完成，教科书上也没讲。