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

兜尼完發表於2024-06-06

本例只是其中一個簡單的使用，在OpenCL中操作管道的函式有很多，這裡只用了其中2個讀寫函式。此例功能是計算12800個整數的積的和。例子不是很好但我實在想不到實用的例子了，因為我做影像處理方面的應用用不到管道來傳遞資料。程式碼如下，這裡不和CPU處理時間進行比較，這種簡單的運算一般是CPU更快。程式執行環境是VS2015、OpenCL306。下面的程式碼還使用了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類建立。原因是在我的電腦上用 cl::Pipe(sizeof(int), 12800) 構造程式會報錯，兩種建立管道的方法的區別是C++版執行庫設定了CL_MEM_READ_WRITE屬性。下面是程式執行結果的截圖，可知執行結果是對的（12800*3*3=115200）：

相關文章