So starten Sie benutzerdefinierte OpenCL-kernel in OpenCV (3.0.0) OCL?

Ich bin wahrscheinlich missbrauchen, OpenCV, indem Sie Sie als wrapper, um die offizielle OpenCL C++ - Bindungen, so dass ich starten kann meinen eigenen Kernel.

Jedoch OpenCV hat Klassen, wie Programm, ProgramSource, Kernel, Queue, etc. das scheint mir zu sagen, dass ich starten kann, meine eigene (auch nicht-Bild-Basis) Kernel mit OpenCV. Ich habe Probleme bei der Suche nach Dokumentation gibt es für diese Klassen, geschweige denn Beispiele. Also, ich nahm einen Stich an es so weit:

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

Hinweis, dass ich nicht auf meinen host-Variablen noch. Ich bin am stecken kernel->args(...). Es gibt 15 überlastungen und keiner von Ihnen angeben, in welcher Reihenfolge sollte ich geben Sie den folgenden, pro-argument:

Den parameter index, so dass ich manuell zuordnen, die parameter in der angegebenen Reihenfolge in den kernel.
Die host-variable selbst.
Die host-Variablen-array-Größe - normalerweise sage ich so etwas wie sizeof(int) * ARRAY_SIZE, obwohl ich verwendet, um anzugeben, dass auf die clEnqueueWriteBuffer Funktion in plain OpenCL.
Gerät buffer memory access, zum Beispiel CL_MEM_READ_ONLY

Sieht es nicht aus, ich nenne enqueueWriteBufer(...), enqueueNDRangeKernel(...), oder enqueueReadBuffer (...), weil (vermute ich mal) die kernel->run() macht alles für mich unter die Haube. Ich gehe davon aus, dass die kernel->run() schreiben Sie die neuen Werte meiner output-parameter.

Ich nicht geben Sie einen Befehl queue, ein Gerät oder ein Kontext. Ich denke, dass es nur eine command-queue und ein Kontext, und das Standard - Gerät erstellt unter der Haube und zugänglich sind, aus diesen Klassen.

So wieder, wie ich die args-Funktion des Kernels?

InformationsquelleAutor Mickael Caruso | 2015-02-15

c++gpgpu opencl opencv

Obwohl ich bin mir nicht 100% sicher, dass ich herausgefunden, einen Weg, dies zu tun.
Dieses Beispiel enthält Tipps, wie Sie pass/empfangen von Daten zu/von einem custom-kernel mittels cv::UMat, grundlegende Typen (z.B. int/float/uchar), und 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();
}

Unten ist ein "shift.cl" - Datei.

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

Dem Punkt ist die Verwendung von UMat. Wir erhalten 5 Parameter im kernel (*data_ptr, int step, int offset, int rows, int cols) mit KernelArg::ReadOnly(umat); 3 (*data_ptr, int step, int offset) mit KernelArg::ReadOnlyNoSize(umat); und nur 1 (*data_prt) mit KernelArg::PtrReadOnly(umat). Von dieser Regel ist das gleiche für WriteOnly und ReadWrite.

Den Schritt und offset sind erforderlich, wenn der Zugriff auf das Daten-array, da UMat kann nicht dichten matrix aufgrund der Speicher-Adresse-Ausrichtung.

cv::ocl::Image2D konstruiert werden können, die aus einer UMat-Instanz und kann direkt übergeben kernel.args(). Mit image2D_t und sampler_t, wir können profitieren von der GPU-hardware textur-Einheiten für linear-interpolation Abtastung (mit reellwertige pixel-Koordinaten).

Beachten Sie, dass die "-D xxx=yyy " build-option bietet text-Ersatz von xxx zu yyy in den kernel-code.

Finden Sie weitere codes in meinem post: http://qiita.com/tackson5/items/8dac6b083071d31baf00

Ich kann nicht kompilieren von OpenCL-kernel. error: parameter may not be qualified with an address space: __global const image2d_t src Meine OCL-Gerät ist eine Intel-Iris-GPU. Irgendwelche Vorschläge?
In meinem Fall habe ich das Problem gelöst, indem Sie ersetzen: __global const image2d_t src durch read_only image2d_t src. Nicht sicher, ob es der richtige Weg, da bin ich angefangen zu lernen OpenCL.
Wie wissen Sie, welche input-argument sollten wir verwenden? Keinem offiziellen Dokument hier? thx im Voraus
Hier finden Sie einige Verweise auf den OpenCL-Website. Es ist ein Spickzettel und der Referenz-Anleitung für die OpenCL 2.2 hier.
Thx für Ihre freundliche Antwort!

InformationsquelleAutor user4588286

Schreibe einen Kommentar

Du musst angemeldet sein, um einen Kommentar abzugeben.