title: Custom Kernel GpuMat with float
body: <p>I'm trying to write a custom kernel using <code>GpuMat</code> data to find the arc cosine of an image's pixels. I can upload, download, and change values when I upload data when the GPU has <code>CV_8UC1</code> data but chars cannot be used to calculate arc cosines. However, when I try to convert my GPU to <code>CV_32FC1</code> type (floats) I get an illegal memory access error during the download part. Here is my code:</p>
  
  <pre><code>//.cu code 
  #include &lt;cuda_runtime.h&gt;
  #include &lt;stdlib.h&gt;
  #include &lt;iostream&gt;
  #include &lt;stdio.h&gt;
  __global__ void funcKernel(const float* srcptr, float* dstptr, size_t srcstep, const     size_t dststep, int cols, int rows){
      int rowInd = blockIdx.y*blockDim.y+threadIdx.y;
      int colInd = blockIdx.x*blockDim.x+threadIdx.x;
      if(rowInd &gt;= rows || colInd &gt;= cols)
              return;
      const float* rowsrcptr=srcptr+rowInd*srcstep;
      float* rowdstPtr=  dstptr+rowInd*dststep;
      float val = rowsrcptr[colInd];
      if((int) val % 90 == 0)
              rowdstPtr[colInd] = -1 ;
      else{
              float acos_val = acos(val);
              rowdstPtr[colInd] = acos_val;
      }
  }
  
  int divUp(int a, int b){
      return (a+b-1)/b;
  }
  
  extern "C"
  {
  void func(const float* srcptr, float* dstptr, size_t srcstep, const size_t dststep, int cols, int rows){
      dim3 blDim(32,8);
      dim3 grDim(divUp(cols, blDim.x), divUp(rows,blDim.y));
      std::cout &lt;&lt; "calling kernel from func\n";
      funcKernel&lt;&lt;&lt;grDim,blDim&gt;&gt;&gt;(srcptr,dstptr,srcstep,dststep,cols,rows);
      std::cout &lt;&lt; "done with kernel call\n";
       cudaDeviceSynchronize();
  }
  
  //.cpp code
  void callKernel(const GpuMat &amp;src, GpuMat &amp;dst){
      float* p = (float*)src.data;
      float* p2 =(float*) dst.data;
      func(p,p2,src.step,dst.step,src.cols,src.rows);
  }
  
  int main(){
      Mat input = imread("cat.jpg",0);
      Mat float_input;
      input.convertTo(float_input,CV_32FC1);
      GpuMat d_frame,d_output;
      Size size = float_input.size();
      d_frame.upload(float_input);
      d_output.create(size,CV_32FC1);
      callKernel(d_frame,d_output);
      Mat output(d_output);
      return 0;
  }
  </code></pre>
  
  <p>When I run the program my compiler tells me this:</p>
  
  <blockquote>
    <p>OpenCV Error: Gpu API call (an illegal memory access was encountered)
    in copy, file
    /home/mobile/opencv-2.4.9/modules/dynamicuda/include/opencv2/dynamicuda/dynamicuda.hpp,
    line 882 terminate called after throwing an instance of
    'cv::Exception'   what(): 
    /home/mobile/opencv-2.4.9/modules/dynamicuda/include/opencv2/dynamicuda/dynamicuda.hpp:882:
    error: (-217) an illegal memory access was encountered in function
    copy</p>
  </blockquote>