title: Covariance matrix with opencl and opencv
body: <p>I recently start to learn opencl and how to create kernel for OpenCV as well.
  I am still working with the basics.</p>
  
  <p>I need to implement an opencl kernel for calculate the covariance matrix.
  This function don't have any opencl kernel already implemented by opencv's fundation.</p>
  
  <p>Technically what I want to process is that  :</p>
  
  <pre><code>cv::RNG tutu(std::time(nullptr)); // This is only for check the improvement of the method over the executions
  
  int sz = tutu.uniform(1,20);
  
  cv::Mat_&lt;float&gt; a1(1,sz);
  cv::Mat_&lt;float&gt; a2(a1.size());
  cv::Mat_&lt;float&gt; c2;
  
  for(std::size_t i=0;i&lt;sz;i++)
  {
      a1(i) = tutu.uniform(0,300);
      a2(i) = tutu.uniform(300,600);
  }
  
  cv::Mat_&lt;float&gt; f;
  
  cv::vconcat(a1,a2,f);
  
  // process the Covariance Matrix :
  
  cv::gemm(one,f,-0.5f,f,1.f,c2,0);
  cv::gemm(c2.t(),c2,1.f,cv::noArray(),0.f,c2);
  </code></pre>
  
  <p>I found OpenCV's fundation have implemented an OpenCL kernel for both gemm and the function transpose.</p>
  
  <p>So I tried to derivate an implementation from an exemple directly from OpenCV's source.</p>
  
  <p>I wrote this :</p>
  
  <p>ocl.h :</p>
  
  <pre><code>void CovarMatrix( cv::Mat_&lt;float&gt;&amp; src,cv::Mat_&lt;float&gt;&amp; covar);
  </code></pre>
  
  <p>ocl.cpp :</p>
  
  <pre><code>    #include &lt;memory&gt;
      #include &lt;fstream&gt;
      #include &lt;sstream&gt;
      #include &lt;iterator&gt;
  
      #include &lt;opencv2/core.hpp&gt;
      #include &lt;opencv2/core/ocl.hpp&gt;
  
  
      namespace test
      {
  
      namespace ocl
      {
  
      namespace
      {
  
      std::unique_ptr&lt;cv::ocl::ProgramSource&gt; cov_src;
  
      void init_cov()
      {
  
          std::ifstream stream("../mahalanobis/covarianceMatrix.cl");
          std::ostringstream sstream;
  
          sstream &lt;&lt; stream.rdbuf();
  
          cv::String norm_file_content = sstream.str();
  
          stream.close();
  
          cov_src.reset(new cv::ocl::ProgramSource(norm_file_content));
      }
  
  
      }
  
  static bool ocl_gemm( cv::Mat_&lt;float&gt;&amp; matA, cv::Mat_&lt;float&gt;&amp; matB, cv::Mat_&lt;float&gt;&amp; CV_OUT matD)
  {
  
      cv::Mat_&lt;float&gt; tmp;
      cv::Mat_&lt;float&gt; tmp2;
  
      cv::Size sizeA = matA.size(), sizeB = matB.size();
  
      cv::Size sizeD(sizeB.width, sizeA.height);
  
      const cv::ocl::Device &amp; dev = cv::ocl::Device::getDefault();
      int max_wg_size = (int)dev.maxWorkGroupSize();
      int block_size = (max_wg_size / 32 &lt; 32) ? (max_wg_size / 16 &lt; 16) ? (max_wg_size / 8 &lt; 8) ? 1 : 8 : 16 : 32;
  
  //    matD.create(sizeD);
  //    tmp2.create(matD.t().size());
  
      tmp.create(sizeD);
      tmp2.create(tmp.t().size());
      matD.create(sizeD.width,sizeD.width);
  
      cv::UMat A = matA.getUMat(cv::ACCESS_READ,cv::USAGE_ALLOCATE_DEVICE_MEMORY);
      cv::UMat B = matB.getUMat(cv::ACCESS_READ,cv::USAGE_ALLOCATE_DEVICE_MEMORY);
  //    cv::UMat D = matD.getUMat(cv::ACCESS_WRITE,cv::USAGE_ALLOCATE_DEVICE_MEMORY);
      cv::UMat D = tmp.getUMat(cv::ACCESS_WRITE,cv::USAGE_ALLOCATE_DEVICE_MEMORY);
  //    cv::UMat E(sizeD.width,sizeD.height,CV_32FC1,cv::Scalar::all(0.),cv::USAGE_ALLOCATE_DEVICE_MEMORY);
  
      cv::UMat E = tmp2.getUMat(cv::ACCESS_WRITE,cv::USAGE_ALLOCATE_DEVICE_MEMORY);
      cv::UMat F = matD.getUMat(cv::ACCESS_WRITE,cv::USAGE_ALLOCATE_DEVICE_MEMORY);
  
      matB.copyTo(D);
  
  
  
      int vectorWidths[] = { 4, 4, 2, 2, 1, 4, 1, -1 };
      int kercn = cv::ocl::checkOptimalVectorWidth(vectorWidths, B, D);
  
      cv::String opts = cv::format(
                  "-I /home/administrateur/lib_dir/opencv_dir/opencv_304/opencv/modules/core/src/opencl/ -D T=float -D T1=float -D WT=%s -D cn=1 -D kercn=%d -D LOCAL_SIZE=%d %s -D HAVE_C -D TILE_DIM=32 -D BLOCK_ROWS=8 -D rowsPerWI=1  ",
  
                             cv::ocl::typeToStr(CV_32FC(kercn)),
                             kercn, block_size,
                            (sizeA.width % block_size !=0) ? "-D NO_MULT" : "");
  
      cv::ocl::Kernel k("covarianceMatrix", *cov_src, opts);
  
  
      k.args(cv::ocl::KernelArg::ReadOnlyNoSize(A),
             cv::ocl::KernelArg::ReadOnlyNoSize(B, 1, kercn),
             cv::ocl::KernelArg::ReadWrite(D, 1, kercn),
             sizeA.width,
             cv::ocl::KernelArg::ReadWrite(E,kercn,1),
             cv::ocl::KernelArg::ReadWrite(F,kercn,kercn)
             );
  
      std::size_t globalsize[2] = { static_cast&lt;std::size_t&gt;(sizeD.width / kercn), static_cast&lt;std::size_t&gt;(sizeD.height)};
      std::size_t localsize[2] = { static_cast&lt;std::size_t&gt;(block_size), static_cast&lt;std::size_t&gt;(block_size)};
      return k.run(2, globalsize, block_size!=1 ? localsize : nullptr, false);
  
  }
  
  
  
  
  
  void CovarMatrix( cv::Mat_&lt;float&gt;&amp; src,cv::Mat_&lt;float&gt;&amp; covar)
  {
  
      if(!covar.empty())
          covar.release();
  
  
      cv::Mat_&lt;float&gt; o = cv::Mat_&lt;float&gt;::ones(src.rows,src.rows);  
  
      if(!cov_src)
          init_cov();
  
      ocl_gemm(o,src,covar);
  
  
  }
  </code></pre>
  
  <p>covarianceMatrix.cl :</p>
  
  <pre><code>#include "gemm.cl"
  #include "transpose.cl"
  
  
  
  __kernel void covarianceMatrix
  (
                     __global const uchar * A_ptr, int A_step, int A_offset,
                     __global const uchar * B_ptr, int B_step, int B_offset,
                     __global uchar * D_ptr, int D_step, int D_offset, int D_rows, int D_cols,
                     int n,
                     __global uchar * E_ptr, int E_step, int E_offset, int E_rows, int E_cols,
                     __global uchar * F_ptr, int F_step, int F_offset, int F_rows, int F_cols
                     )
  
  {
  
  
  //    cv::gemm(src2,src1,-0.5,src1,1.f,src2);
  //    cv::gemm(src2.t(),src2,1.f,cv::noArray(),0.f,dest);
  
  
  
      gemm(A_ptr,A_step,A_offset,
           B_ptr,B_step,B_offset,
           D_ptr,D_step,D_offset,D_rows,D_cols,
           n,-0.5f,1.f);
  
      transpose(D_ptr,D_step,D_offset,D_rows,D_cols*sizeof(float),
                E_ptr,E_step,E_offset);
  
  
      gemm(E_ptr,E_step,E_offset,
           D_ptr,D_step,D_offset,
           F_ptr,F_step,F_offset,F_rows,F_cols,
           n,1.f,0.f);
  
  
  }
  </code></pre>
  
  <p>If the size of the matrix is fewer than 6 is work perfectly :).
  Otherwise ... not really.
  It can be check with this code :</p>
  
  <pre><code>cv::RNG tutu(std::time(nullptr));
  
  int sz = tutu.uniform(1,20);
  
  cv::Mat_&lt;float&gt; a1(1,sz);
  cv::Mat_&lt;float&gt; a2(a1.size());
  
  for(std::size_t i=0;i&lt;sz;i++)
  {
      a1(i) = tutu.uniform(0,300);
      a2(i) = tutu.uniform(300,600);
  }
  
  cv::Mat_&lt;float&gt; f;
  
  cv::vconcat(a1,a2,f);
  
  cv::Mat_&lt;float&gt; c1;
  cv::Mat_&lt;float&gt; c2;
  
      cv::Mat_&lt;float&gt; mean;
  // reference
      cv::calcCovarMatrix(f,c1,mean,cv::COVAR_ROWS | cv::COVAR_NORMAL,CV_32F);
  // check
  test::ocl::CovarMatrix(f,c2);
  
  std::size_t cnt(0.f);
  
  for(auto it = c1.begin(),it2 = c2.begin();it != c1.end();it++,it2++)
      if(*it == *it2)
          cnt++;
  
  std::cout&lt;&lt;"check "&lt;&lt;cnt&lt;&lt;" "&lt;&lt;c1.total()&lt;&lt;std::endl;
  </code></pre>
  
  <p>I am still new in OpenCL and I am interrested to know what I did wrong.</p>
  
  <p>Does someone already implemented an OpenCL kernel for process the covariance matrix with OpenCV ?</p>
  
  <p>Thank in advance for any help.</p>