How to pass cv::cuda::GpuMat to a kernel as cv::cuda::PtrStepSz?

cudawarped · May 10, 2024, 5:30am

Nearly all of the internal CUDA code is passed a GpuMat which is implicitly converted to a PtrStep or PtrStepSz because of

template <typename _Tp> operator PtrStepSz<_Tp>() const;

github.com

opencv/opencv/blob/1d9a4120fa7dbeab6316c255d283c16cd006a35e/modules/core/include/opencv2/core/cuda.hpp#L277


      
          CV_WRAP void assignTo(GpuMat& m, int type = -1) const;
          
          //! returns pointer to y-th row
          uchar* ptr(int y = 0);
          const uchar* ptr(int y = 0) const;
          
          //! template version of the above method
          template<typename _Tp> _Tp* ptr(int y = 0);
          template<typename _Tp> const _Tp* ptr(int y = 0) const;
          
          template <typename _Tp> operator PtrStepSz<_Tp>() const;
          template <typename _Tp> operator PtrStep<_Tp>() const;
          
          //! returns a new GpuMat header for the specified row
          CV_WRAP GpuMat row(int y) const;
          
          //! returns a new GpuMat header for the specified column
          CV_WRAP GpuMat col(int x) const;
          
          //! ... for the specified row span
          CV_WRAP GpuMat rowRange(int startrow, int endrow) const;

See resize for an example. The kernel is passed a PtrStepSz for dst

template <class Ptr2D, typename T> __global__ void resize(Ptr2D src, PtrStepSz<T> dst, const float fy, const float fx)

github.com

opencv/opencv_contrib/blob/3c2bcbfe8374edaf3eb756b560374244538d57b6/modules/cudawarping/src/cuda/resize.cu#L109


      
                  src_reg = src(y2_read, x1);
                  out = out + src_reg * ((x2 - src_x) * (src_y - y1));
          
                  src_reg = src(y2_read, x2_read);
                  out = out + src_reg * ((src_x - x1) * (src_y - y1));
          
                  dst(dst_y, dst_x) = saturate_cast<T>(out);
              }
          }
          
          template <class Ptr2D, typename T> __global__ void resize(Ptr2D src, PtrStepSz<T> dst, const float fy, const float fx)
          {
              const int dst_x = blockDim.x * blockIdx.x + threadIdx.x;
              const int dst_y = blockDim.y * blockIdx.y + threadIdx.y;
          
              if (dst_x < dst.cols && dst_y < dst.rows)
              {
                  const float src_x = dst_x * fx;
                  const float src_y = dst_y * fy;
          
                  dst(dst_y, dst_x) = src(src_y, src_x);

but is orginally called with a GpuMat

github.com

opencv/opencv_contrib/blob/3c2bcbfe8374edaf3eb756b560374244538d57b6/modules/cudawarping/src/resize.cpp#L87


      
          {
              dsize = Size(saturate_cast<int>(src.cols * fx), saturate_cast<int>(src.rows * fy));
          }
          else
          {
              fx = static_cast<double>(dsize.width) / src.cols;
              fy = static_cast<double>(dsize.height) / src.rows;
          }
          
          _dst.create(dsize, src.type());
          GpuMat dst = _dst.getGpuMat();
          
          if (dsize == src.size())
          {
              src.copyTo(dst, stream);
              return;
          }
          
          const func_t func = funcs[src.depth()][src.channels() - 1];
          
          if (!func)

Topic		Replies	Views
How to address a cv::cuda::GpuMat variable, passed to kernel as cv::cudaPtrStepSz? C++ cuda	1	292	May 30, 2024
Copy cv::cuda::GpuMat in Cuda Kernel C++ cuda	4	1663	October 5, 2022
How to use cv::cudev::GpuMat_ correctly? C++ cuda	4	867	May 26, 2021
Manually download cv::cuda::GpuMat to cv::Mat using cudaMemcpyAsync C++ cuda	1	110	August 14, 2024
cv::Matx in cuda kernel C++ cuda , core	6	982	November 4, 2022

How to pass cv::cuda::GpuMat to a kernel as cv::cuda::PtrStepSz?

Related topics