Custom OpenCL kernel in opencv.

edit

I am trying to create a custom OpenCL kernel and I would like to use OpenCV to help me to initialize all the opencl stuff.

I downloaded and compiled Opencv 2.4.7 from source on a ubuntu. I enabled the OCL module, and OCL samples works well.

I did this very simple program which compute the element wise square of a matrix :

int main(int _argc, const char** _argv)
{
    cv::ocl::DevicesInfo devInfo;
    int res = cv::ocl::getOpenCLDevices(devInfo);
    if(res == 0)
    {
        std::cout << "There is no OPENCL Here !" << std::endl;
    }else
    {
        for(int i = 0 ; i < devInfo.size() ; ++i)
        {
            std::cout << "Device : " << devInfo[i]->deviceName << " is present" << std::endl;
        }
    }
    cv::ocl::setDevice(devInfo[0]);        // select device to use
    std::cout << CV_VERSION_EPOCH << "." << CV_VERSION_MAJOR << "." << CV_VERSION_MINOR << std::endl;
    const char *KernelSource = "\n" \
    "__kernel void square(                       \n" \
    "   __global uchar* input,                   \n" \
    "   __global uchar* output)                  \n" \
    "{                                           \n" \
    "   int i = get_global_id(0);                \n" \
    "   output[i] = input[i] * input[i];         \n" \
    "}\n";
    cv::ocl::ProgramSource src("square", KernelSource);
    std::size_t globalThreads[3]={1,0,0};
    std::size_t localThreads[3]={5,5,0};
    cv::ocl::oclMat source(cv::Size(500,500), CV_8UC1);
    cv::ocl::oclMat dest(cv::Size(500,500), CV_8UC1);

    std::vector<std::pair<size_t , const void *> > args;
    args.push_back( std::make_pair( 4, (void *) &source.data ));
    args.push_back( std::make_pair( 4, (void *) &dest.data ));

    cv::ocl::openCLExecuteKernelInterop(cv::ocl::Context::getContext(), src, "square",
                                        globalThreads, localThreads, args, 1, CV_8U, "");
    return 0;
}

The output of this program is an OpenCL error :

OpenCV is : 2.4.7
Platform : NVIDIA CUDA is present
Device : GeForce GT 425M is present, it's device number 0
0xa004f90
__kernel void square(                                                       
   __global float* input,                                              
   __global float* output,                                             
   const unsigned int count)                                           
{                                                                      
   int i = get_global_id(0);                                           
   if(i < count)                                                       
       output[i] = input[i] * input[i];                                
}                                                                      

OpenCV Error: Gpu API call (CL_INVALID_KERNEL_NAME) in openCLGetKernelFromSource, file /home/aba/ExternalSources/opencv-2.4.7/modules/ocl/src/cl_operations.cpp, line 328
terminate called after throwing an instance of 'cv::Exception'
  what():  /home/aba/ExternalSources/opencv-2.4.7/modules/ocl/src/cl_operations.cpp:328: error: (-217) CL_INVALID_KERNEL_NAME in function openCLGetKernelFromSource

I also replaced "square" as third argument of openCLExecuteKernelInterop, but the result is the same.

So, does any body can explain to me what happens ?

I also read in the header file which define openCLExecuteKernelInterop that this method is deprecated. Does someone can tell me which method will replace it ?

Regards,