Как запустить пользовательское ядро ​​OpenCL в OpenCV (3.0.0) OCL?

Question

Я, вероятно, неправильно использую OpenCV, используя его как оболочку для официальных ссылок OpenCL C++, чтобы я мог запускать свои собственные ядра.

Однако в OpenCV есть классы, такие как Program, ProgramSource, Kernel, Queue и т. Д., Которые, похоже, говорят мне, что я могу запускать свои собственные (даже не имиджные) ядра с OpenCV. У меня возникли проблемы с поиском документации для этих классов, не говоря уже о примерах. Итак, я до сих пор на него наносил удар:

#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; 
} 

Отметьте, что я еще не установил переменные хоста. Я застрял в kernel->args(...). Существует 15 перегрузок, и ни один из них не указывает, какой порядок я должен указать для аргумента:

1. Индекс параметров, поэтому я вручную сопоставляю параметр в порядке, указанном в ядре.
2. Сама переменная хоста.
3. Размер массива переменной хоста - обычно я говорю что-то вроде sizeof(int) * ARRAY_SIZE, хотя я использовал, чтобы указать это в функции clEnqueueWriteBuffer в простом OpenCL. Доступ к памяти буфера
4. устройств, например CL_MEM_READ_ONLY

Это не выглядит, как я называю enqueueWriteBufer (...), enqueueNDRangeKernel (...), или enqueueReadBuffer (...), потому что (я думаю,) kernel-> run() делает все это для меня под капотом. Я предполагаю, что kernel-> run() будет записывать новые значения в мой выходной параметр.

Я не указывал очередь команд, устройство или контекст. Я думаю, что есть только одна командная очередь и один контекст, а устройство по умолчанию - все создано под капотом и доступно из этих классов.

Итак, как я могу использовать функцию args ядра?

Answer 1

Хотя я не уверен на 100%, я выяснил, как это сделать. В этом примере содержатся советы о том, как передавать/извлекать данные в/из пользовательского ядра с использованием cv :: UMat, базовых типов (например, int/float/uchar) и 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(); 
} 

Ниже представлен файл "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; 
} 

Точка использования UMat. Мы получаем 5 параметров в ядре (* data_ptr, int step, int offset, int rows, int cols) с помощью KernelArg :: ReadOnly (umat); 3 (* data_ptr, int step, int offset) с KernelArg :: ReadOnlyNoSize (umat); и только 1 (* data_prt) с KernelArg :: PtrReadOnly (umat). Это правило одинаково для WriteOnly и ReadWrite.

Шаг и смещение требуются при доступе к массиву данных, поскольку UMat может быть не плотной матрицей из-за выравнивания адреса памяти.

cv :: ocl :: Image2D может быть сконструирован из экземпляра UMat и может быть напрямую передан kernel.args(). С помощью image2D_t и sampler_t мы можем использовать аппаратные текстурные единицы графического процессора для линейно-интерполяционной выборки (с реальными координатами пикселей).

Обратите внимание, что «-D xxx = yyy» build-option предлагает замену текста с xxx на yyy в коде ядра.

Вы можете найти больше кодов на своем посту: http://qiita.com/tackson5/items/8dac6b083071d31baf00