Come avviare il kernel OpenCL personalizzato in OpenCV (3.0.0) OCL?

Question

Probabilmente sto utilizzando in modo errato OpenCV utilizzandolo come wrapper per i binding OpenCL C++ ufficiali in modo che possa avviare i miei kernel personali.

Tuttavia, OpenCV ha classi come Program, ProgramSource, Kernel, Queue, ecc. Che sembrano dirmi che posso lanciare i miei kernel (anche non basati su immagini) con OpenCV. Ho difficoltà a trovare la documentazione là fuori per queste classi, per non parlare di esempi. Quindi, ho preso una pugnalata fino ad ora:

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

Nota che non ho ancora impostato le mie variabili host. Sono bloccato a kernel->args(...). Ci sono 15 sovraccarichi e nessuno di loro specificano che ordine dovrei specificare quanto segue, per ogni argomento:

1. L'indice parametro, quindi abbino manualmente i parametri nell'ordine indicato nel kernel.
2. La variabile host stessa.
3. La dimensione dell'array della variabile host, in genere dico qualcosa come sizeof(int) * ARRAY_SIZE, anche se ero solito specificarlo nella funzione clEnqueueWriteBuffer in OpenCL. accesso alla memoria buffer di
4. dispositivo, ad esempio CL_MEM_READ_ONLY

E non sembra che io chiamo enqueueWriteBufer (...), enqueueNDRangeKernel (...), o enqueueReadBuffer (...) perché (suppongo) il kernel-> run() fa tutto questo per me sotto il cofano. Presumo che kernel-> run() scriverà i nuovi valori nel mio parametro di output.

Non ho specificato una coda comandi, dispositivo o contesto. Penso che ci sia solo una coda di comandi e un contesto, e il dispositivo predefinito - tutti creati sotto il cofano e accessibili da queste classi.

Quindi, ancora una volta, come utilizzare la funzione args del kernel?

Answer 1

Anche se non ne sono sicuro al 100%, ho trovato un modo per farlo. Questo esempio contiene suggerimenti su come passare/recuperare i dati su/da un kernel personalizzato usando cv :: UMat, i tipi di base (ad es. Int/float/uchar) e 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(); 
} 

Di seguito è riportato un file "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; 
} 

Il punto è l'uso di UMat. Riceviamo 5 parametri nel kernel (* data_ptr, int step, int offset, int rows, int cols) con KernelArg :: ReadOnly (umat); 3 (* data_ptr, int step, int offset) con KernelArg :: ReadOnlyNoSize (umat); e solo 1 (* data_prt) con KernelArg :: PtrReadOnly (umat). Questa regola è la stessa per WriteOnly e ReadWrite.

Il passaggio e l'offset sono necessari quando si accede all'array di dati, poiché UMat potrebbe non essere una matrice densa a causa dell'allineamento dell'indirizzo di memoria.

cv :: ocl :: Image2D può essere creato da un'istanza di UMat e può essere passato direttamente a kernel.args(). Con image2D_t e sampler_t, possiamo beneficiare delle unità di trama hardware della GPU per il campionamento dell'interpolazione lineare (con coordinate di pixel a valori reali).

Si noti che l'opzione di build "-D xxx = yyy" offre la sostituzione del testo da xxx a yyy nel codice del kernel.

Puoi trovare altri codici al mio posto: http://qiita.com/tackson5/items/8dac6b083071d31baf00