OpenCV-OpenCL-kernel-arguments.cpp

// a 2-d mat
cv::Mat mat1(2,2,CV_8UC1);

int sz[] = {2,3,4};

// a 3-d mat
cv::Mat mat2(3, sz, CV_8UC1);

cv::UMat m1 = mat1.getUMat(cv::ACCESS_RW);
cv::UMat m2 = mat2.getUMat(cv::ACCESS_RW);

ocl::Kernel k; 
//....
// Initialize the kernel k

//============================
//    Example1 (2-d and 3-d) PTR_ONLY
//----------------------------
k.set(0, ocl::KernelArg::PtrReadWrite(m1));
// k.set(0, ocl::KernelArg::PtrReadWrite(m2));

/*
//----------------------------
// the corresponding kernel should be: 
//----------------------------
__kernel void myKernelName(__global uchar *data)
{

}
*/

//============================
//    Example2 (2-d) NO_SIZE
//----------------------------
k.set(0, ocl::KernelArg::ReadWriteNoSize(m1));
/*
//----------------------------
// the corresponding kernel should be: 
//----------------------------
__kernel void myKernelName(__global uchar *data, int step, int offset)
{

}
// int step;   ---> equals to m1.step[0], i.e, the number of bytes of each row
// int offset; ---> equals to 0 in our case, because it is not a submatrix
*/

//============================
//    Example3 (3-d) NO_SIZE
//----------------------------
k.set(0, ocl::KernelArg::ReadWriteNoSize(m2));
/*
//----------------------------
// the corresponding kernel should be: 
//----------------------------
__kernel void myKernelName(__global uchar *data, int slicestep, int step, int offset)
{

}
// int slicestep; ---> equals to m2.step[0], i.e, the number of bytes of each plane
// int step;      ---> equals to m2.step[1], i.e, the number of bytes of each row
// int offset;    ---> equals to 0 in our case, because it is not a submatrix
*/


//============================
//    Example4 (2-d) other flags without PTR_ONLY and NO_SIZE
//----------------------------
k.set(0, ocl::KernelArg::ReadWrite(m1));
/*
//----------------------------
// the corresponding kernel should be: 
//----------------------------
__kernel void myKernelName(__global uchar *data, int step, int offset, int rows, int cols)
{

}
// int step;   ---> equals to m1.step[0], i.e, the number of bytes of each row
// int offset; ---> equals to 0 in our case, because it is not a submatrix
// int rows;   ---> equals to m1.rows, number of rows of the matrix
// int cols;   ---> equals to m1.cols, number of columns of the matrix
*/

//============================
//    Example5 (3-d) other flags without PTR_ONLY and NO_SIZE
//----------------------------
k.set(0, ocl::KernelArg::ReadWrite(m2));
/*
//----------------------------
// the corresponding kernel should be: 
//----------------------------
__kernel void myKernelName(__global uchar *data, int slicestep, int step, int offset, int slices, int rows, int cols)
{

}
// int slicestep; ---> equals to m2.step[0], i.e, the number of bytes of each plane
// int step;   ---> equals to m2.step[0], i.e, the number of bytes of each row
// int offset; ---> equals to 0 in our case, because it is not a submatrix
// int slices; ---> equals to m2.size[0], i.e, number of planes of the matrix
// int rows;   ---> equals to m1.rows, number of rows of the matrix
// int cols;   ---> equals to m1.cols, number of columns of the matrix
*/