Làm thế nào để khởi chạy hạt nhân OpenCL tùy chỉnh trong OpenCV (3.0.0) OCL?

Question

Tôi có thể lạm dụng OpenCV bằng cách sử dụng nó như là trình bao bọc cho các ràng buộc OpenCL C++ chính thức để tôi có thể khởi chạy hạt nhân của riêng mình. Tuy nhiên, OpenCV không có các lớp như Program, ProgramSource, Kernel, Queue, vv mà dường như cho tôi biết rằng tôi có thể khởi chạy các hạt nhân của riêng mình (ngay cả không dựa trên hình ảnh) với OpenCV. Tôi gặp khó khăn khi tìm tài liệu ở đó cho các lớp này, hãy để một mình ví dụ. Vì vậy, tôi đã đâm vào nó cho đến nay:

#include <fstream> 
#include <iostream> 

#include "opencv2/opencv.hpp" 
#include "opencv2/core/ocl.hpp" 

#define ARRAY_SIZE 128 

using namespace std; 
using namespace cv; 

int main(int, char) 
{ 
    std::ifstream file("kernels.cl"); 
    std::string kcode(std::istreambuf_iterator<char>(file), 
     (std::istreambuf_iterator<char>())); 

    cv::ocl::ProgramSource * programSource; 
    programSource = new cv::ocl::ProgramSource(kcode.c_str()); 

    cv::String errorMessage; 
    cv::ocl::Program * program; 
    program = new cv::ocl::Program(*programSource, NULL, errorMessage); 

    cv::ocl::Kernel * kernel; 
    kernel = new cv::ocl::Kernel("simple_add", *program); 
    /* I'm stuck here at the args. */ 

    size_t globalSize[2] = { ARRAY_SIZE, 1 }; 
    size_t localSize[2] = { ARRAY_SIZE, 1 };  
    kernel->run(ARRAY_SIZE, globalSize, localSize, true); 

    return 0; 
} 

Lưu ý rằng tôi chưa thiết lập biến chủ của mình. Tôi bị kẹt ở số kernel->args(...). Có 15 tình trạng quá tải và không có quá tải nào xác định thứ tự tôi nên chỉ định những thứ sau, cho mỗi đối số:

1. Chỉ số tham số, vì vậy tôi khớp thủ công tham số theo thứ tự được đưa ra trong hạt nhân.
2. Bản thân biến chủ.
3. Kích thước mảng của biến chủ nhà - thường tôi nói điều gì đó như sizeof(int) * ARRAY_SIZE, mặc dù tôi đã sử dụng để chỉ định rằng trên hàm clEnqueueWriteBuffer trong đồng bằng OpenCL.
4. Device truy cập bộ nhớ đệm, ví dụ CL_MEM_READ_ONLY

Nó không giống như tôi gọi enqueueWriteBufer (...), enqueueNDRangeKernel (...), hoặc enqueueReadBuffer (...) vì (tôi đoán) kernel-> run() thực hiện tất cả điều đó cho tôi dưới mui xe. Tôi giả định rằng kernel-> run() sẽ ghi các giá trị mới vào tham số đầu ra của tôi.

Tôi không chỉ định hàng đợi lệnh, thiết bị hoặc ngữ cảnh. Tôi nghĩ rằng chỉ có một hàng đợi lệnh và một bối cảnh, và thiết bị mặc định - tất cả đều được tạo dưới nắp và có thể truy cập được từ các lớp này.

Vì vậy, một lần nữa, làm cách nào để sử dụng hàm args của hạt nhân?

Answer 1

Mặc dù tôi không chắc chắn 100%, tôi đã tìm ra cách để thực hiện việc này. Ví dụ này chứa các mẹo về cách chuyển/truy xuất dữ liệu đến/từ hạt nhân tùy chỉnh bằng cách sử dụng cv :: UMat, các loại cơ bản (ví dụ: int/float/uchar) và Image2D.

#include <iostream> 
#include <fstream> 
#include <string> 
#include <iterator> 
#include <opencv2/opencv.hpp> 
#include <opencv2/core/ocl.hpp> 

using namespace std; 

void main() 
{ 
    if (!cv::ocl::haveOpenCL()) 
    { 
     cout << "OpenCL is not avaiable..." << endl; 
     return; 
    } 
    cv::ocl::Context context; 
    if (!context.create(cv::ocl::Device::TYPE_GPU)) 
    { 
     cout << "Failed creating the context..." << endl; 
     return; 
    } 

    // In OpenCV 3.0.0 beta, only a single device is detected. 
    cout << context.ndevices() << " GPU devices are detected." << endl; 
    for (int i = 0; i < context.ndevices(); i++) 
    { 
     cv::ocl::Device device = context.device(i); 
     cout << "name     : " << device.name() << endl; 
     cout << "available   : " << device.available() << endl; 
     cout << "imageSupport   : " << device.imageSupport() << endl; 
     cout << "OpenCL_C_Version  : " << device.OpenCL_C_Version() << endl; 
     cout << endl; 
    } 

    // Select the first device 
    cv::ocl::Device(context.device(0)); 

    // Transfer Mat data to the device 
    cv::Mat mat_src = cv::imread("Lena.png", cv::IMREAD_GRAYSCALE); 
    mat_src.convertTo(mat_src, CV_32F, 1.0/255); 
    cv::UMat umat_src = mat_src.getUMat(cv::ACCESS_READ, cv::USAGE_ALLOCATE_DEVICE_MEMORY); 
    cv::UMat umat_dst(mat_src.size(), CV_32F, cv::ACCESS_WRITE, cv::USAGE_ALLOCATE_DEVICE_MEMORY); 

    std::ifstream ifs("shift.cl"); 
    if (ifs.fail()) return; 
    std::string kernelSource((std::istreambuf_iterator<char>(ifs)), std::istreambuf_iterator<char>()); 
    cv::ocl::ProgramSource programSource(kernelSource); 

    // Compile the kernel code 
    cv::String errmsg; 
    cv::String buildopt = cv::format("-D dstT=%s", cv::ocl::typeToStr(umat_dst.depth())); // "-D dstT=float" 
    cv::ocl::Program program = context.getProg(programSource, buildopt, errmsg); 

    cv::ocl::Image2D image(umat_src); 
    float shift_x = 100.5; 
    float shift_y = -50.0; 
    cv::ocl::Kernel kernel("shift", program); 
    kernel.args(image, shift_x, shift_y, cv::ocl::KernelArg::ReadWrite(umat_dst)); 

    size_t globalThreads[3] = { mat_src.cols, mat_src.rows, 1 }; 
    //size_t localThreads[3] = { 16, 16, 1 }; 
    bool success = kernel.run(3, globalThreads, NULL, true); 
    if (!success){ 
     cout << "Failed running the kernel..." << endl; 
     return; 
    } 

    // Download the dst data from the device (?) 
    cv::Mat mat_dst = umat_dst.getMat(cv::ACCESS_READ); 

    cv::imshow("src", mat_src); 
    cv::imshow("dst", mat_dst); 
    cv::waitKey(); 
} 

Dưới đây là tệp "shift.cl".

__constant sampler_t samplerLN = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR; 
__kernel void shift(
    __global const image2d_t src, 
    float shift_x, 
    float shift_y, 
    __global uchar* dst, 
    int dst_step, int dst_offset, int dst_rows, int dst_cols) 
{ 
    int x = get_global_id(0); 
    int y = get_global_id(1); 
    if (x >= dst_cols) return; 
    int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT), dst_offset)); 
    __global dstT *dstf = (__global dstT *)(dst + dst_index); 
    float2 coord = (float2)((float)x+0.5f+shift_x, (float)y+0.5f+shift_y); 
    dstf[0] = (dstT)read_imagef(src, samplerLN, coord).x; 
} 

Vấn đề là sử dụng UMat. Chúng ta nhận 5 tham số trong kernel (* data_ptr, int step, int offset, int rows, int cols) với KernelArg :: ReadOnly (umat); 3 (* data_ptr, int step, int offset) với KernelArg :: ReadOnlyNoSize (umat); và chỉ 1 (* data_prt) với KernelArg :: PtrReadOnly (umat). Quy tắc này giống nhau đối với WriteOnly và ReadWrite.

Bước và bù được yêu cầu khi truy cập mảng dữ liệu, vì UMat có thể không phải là ma trận dày đặc do căn chỉnh địa chỉ bộ nhớ.

cv :: ocl :: Image2D có thể được tạo từ một cá thể UMat và có thể được chuyển trực tiếp đến kernel.args(). Với image2D_t và sampler_t, chúng ta có thể hưởng lợi từ các đơn vị kết cấu phần cứng của GPU để lấy mẫu nội suy tuyến tính (với tọa độ pixel có giá trị thực).

Lưu ý rằng tùy chọn xây dựng "-D xxx = yyy" cung cấp thay thế văn bản từ xxx thành yyy trong mã hạt nhân.

Bạn có thể tìm thêm mã tại bài đăng của tôi: http://qiita.com/tackson5/items/8dac6b083071d31baf00