Comment opencl fonctionne avec opencv

Question

J'essaye de créer un algorithme de région groiwing en utilisant OpenCL, pour cela je vais ouvrir une image en utilisant OpenCV. Le problème est de savoir comment transformer les données en OpenCL.

J'utilise la version OpenCV: 2.4.9 et OpenCL: AMD APP SDK \ 2.9-1 en studio visuel

Est-ce que quelqu'un me dire après openning l'image en utilisant OpenCV que dois-je faire

Answer 1

En général , il existe deux façons de transférer des images (ou d'autres données) du programme hôte au programme de périphérique dans les applications OpenCL: 1-Utilisation des tampons 2- Utilisation de Image2d. Les deux utilisent le type cl_mem. Parce que l'utilisation de buffer est plus simple que l'utilisation de image2d (en particulier dans les images en niveaux de gris), j'explique comment transférer des images d'un programme hôte vers un périphérique en utilisant des tampons dans OpenCL. Après avoir lu l'image d'entrée par l'objet openCV Mat, convertissez-la en une image à échelle de gris. Ensuite, nous utilisons la méthode clCreateBuffer qui renvoie un tampon cl_mem. Nous pouvons simplement passer data (une propriété de Mat obeject) au clCreateBuffer pour initialiser notre tampon de noyau d'entrée par des données d'image d'entrée. Ensuite, nous pouvons transférer le tampon créé au noyau en utilisant la méthode clSetKernelArg. Enfin, lorsque le noyau a terminé son travail, nous pouvons lire les résultats par clEnqueueReadBuffer.

Lisez les commentaires pour comprendre ce code et n'hésitez pas à poser des questions.

code hôte:

// Make Contex, Kerenl and other requirements for OpenCL before this section.... 

Mat image = imread("logo.bmp", CV_LOAD_IMAGE_COLOR); // reading input image by opencv to Mat type 
Mat input_; 

cvtColor(image, input_, CV_BGR2GRAY); // convert input image to gray scale 

cl_mem inputSignalBuffer = clCreateBuffer(
     context, 
     CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 
     input_.rows * input_.cols * input_.elemSize(), 
     static_cast<void *>(input_.data),    // inputSignalBuffers will be initialized by input_.data which contains input image data 
     &errNum); 

cl_mem outputSignalBuffer = clCreateBuffer(  // make and preparing an empty output buffer to use after opencl kernel call back 
     context, 
     CL_MEM_WRITE_ONLY, 
     input_.rows * input_.cols * input_.elemSize(), 
     NULL, 
     &errNum); 
    checkErr(errNum, "clCreateBuffer(outputSignal)"); 


    errNum = clSetKernelArg(kernel, 0, sizeof(cl_mem), &inputSignalBuffer); // passing input buffer and output buffer to kernel in order to be used on device 
    errNum |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &maskBuffer); 
    errNum |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &outputSignalBuffer); 
    errNum |= clSetKernelArg(kernel, 3, sizeof(cl_uint), &input_.rows); 
    errNum |= clSetKernelArg(kernel, 4, sizeof(cl_uint), &input_.cols); 
    errNum |= clSetKernelArg(kernel, 5, sizeof(cl_uint), &maskWidth); 

    size_t localWorkSize[2] = { 16, 16 }; // Using 2 dimensional range with size of work group 16 
    size_t globalWorkSize[2] = { input_.rows, // Note: Global work size (input image rows and cols) should be multiple of size of work group. 
            input_.cols }; 

    // Queue the kernel up for execution across the array 
    errNum = clEnqueueNDRangeKernel(// enqueue kernel with enabling host blocking until finishing kernel execution 
     queue, 
     kernel, 
     2, 
     NULL, 
     globalWorkSize, 
     localWorkSize, 
     0, 
     NULL, 
     NULL); 
    checkErr(errNum, "clEnqueueNDRangeKernel"); 

    Mat output_ = cv::Mat(input_.rows, input_.cols, CV_8UC1); 


    errNum = clEnqueueReadBuffer(// reading from ourput parameter of kernel 
     queue, 
     outputSignalBuffer, 
     CL_TRUE, 
     0, 
     input_.rows * input_.cols * input_.elemSize(), 
     output_.data, //initialize OpenCV Mat by output_.data which contains output results of kernel 
     0, 
     NULL, 
     NULL); 
    checkErr(errNum, "clEnqueueReadBuffer"); 

    // cut the extra border spaces which has been added in the first part of the code in order to adjust image size with Work Group Size; 

    cv::imwrite("output.bmp",output_); // saving output in image file 

Code noyau:

__kernel void convolve(
    const __global uchar * const input, 
    __constant uint * const mask, 
    __global uchar * const output, 
    const int inputHeight, 
    const int inputWidth, 
    const int maskWidth) 
{  
    uint sum = 0; 

    const int curr_x = get_global_id(0); // current curr_x (row) 
    const int curr_y = get_global_id(1); // current curr_y (col) 
    int d = maskWidth/2;  

    if(curr_x>d-1 && curr_y>d-1 && curr_x<inputHeight-d && curr_y<inputWidth-d) // checking mask borders not to be out of input matrix 
     for(int i=-d; i<=d; i++) 
      for(int j=-d; j<=d; j++) { 
       int mask_ptr = maskWidth*(i+d)+(j+d); //you can also use mad24(maskWidth, i+d, j+d) which is faster. 
       sum += input[(curr_x+i)*inputWidth+curr_y+j]*mask[mask_ptr];    
      } 

    sum /= (maskWidth*maskWidth); // miangin gereftan 

    sum = clamp(sum, (uint)0, (uint)255);// clamp == min(max(x, minval), maxval) 

    output[curr_x*inputWidth+curr_y] = sum; 

}