cuda + opencv非法内存访问

Question

我是新的Cuda编程。

我想要做的正是实现矩阵反正切函数。 为了知道如何与OpenCV交互，我从OpenCV的双边过滤器实现（cudaimgproc之一）中提取了我的代码。

所以我写了atan.hpp：

#ifndef ATAN_HPP 
#define ATAN_HPP 

#include <opencv2/core.hpp> 
#include <opencv2/core/cuda.hpp> 


namespace support 
{ 

namespace cuda 
{ 
void atan(cv::InputArray _src,cv::OutputArray _dst,cv::cuda::Stream& stream = cv::cuda::Stream::Null()); 
} 

} 

#endif // ATAN_HPP 

然后atan.cpp

#include "atan.hpp" 

#include <opencv2/core/cuda.hpp> 
#include <opencv2/cudev/common.hpp> 



namespace support 
{ 

namespace cuda 
{ 

namespace device 
{ 
template<class _Ty> 
void atan_(const cv::cuda::PtrStepSzb& src, cv::cuda::PtrStepSzb dst,cudaStream_t); 

} 

void atan(cv::InputArray _src, cv::OutputArray _dst, cv::cuda::Stream &stream) 
{ 

    CV_Assert(
       (_src.depth() == _src.type()) && 
       (_src.isMat() || _src.isUMat() || ((_src.kind() & cv::_InputArray::CUDA_GPU_MAT) == cv::_InputArray::CUDA_GPU_MAT)) && 
       (_dst.isMat() || _dst.isUMat() || ((_dst.kind() & cv::_OutputArray::CUDA_GPU_MAT) == cv::_OutputArray::CUDA_GPU_MAT)) 
       ); 

    cv::cuda::GpuMat src; 
    cv::cuda::GpuMat buf; 

    int type = _src.isMat() ? 0 : _src.isUMat() ? 1 : 2; 

    switch(type) 
    { 
    case 0: 
    { 
     cv::Mat tmp = _src.getMat(); 

     src.upload(tmp); 

    } 
     break; 

    case 1: 
    { 
     cv::UMat tmp = _src.getUMat(); 
     cv::Mat tmp2; 

     tmp.copyTo(tmp2); 

     src.upload(tmp2); 
    } 
     break; 

    case 2: 
     src = _src.getGpuMat(); 
     break; 

    } 



    buf.create(src.size(),src.type()); 
// buf.upload(cv::Mat::zeros(src.size(),src.type())); 

    switch (buf.depth()) 
    { 
    case CV_8U: 
     device::atan_<uchar>(src,buf, cv::cuda::StreamAccessor::getStream(stream)); 
     break; 
    case CV_8S: 
     device::atan_<char>(src,buf, cv::cuda::StreamAccessor::getStream(stream)); 
     break; 
    case CV_16U: 
     device::atan_<ushort>(src,buf, cv::cuda::StreamAccessor::getStream(stream)); 
     break; 
    case CV_16S: 
     device::atan_<short>(src,buf, cv::cuda::StreamAccessor::getStream(stream)); 
     break; 
    case CV_32S: 
     device::atan_<int>(src,buf, cv::cuda::StreamAccessor::getStream(stream)); 
     break; 
    case CV_32F: 
     device::atan_<float>(src,buf, cv::cuda::StreamAccessor::getStream(stream)); 
     break; 
    case CV_64F: 
     device::atan_<double>(src,buf, cv::cuda::StreamAccessor::getStream(stream)); 
     break; 
    } 




    type = _dst.isMat() ? 0 : _dst.isUMat() ? 1 : 2; 

    switch(type) 
    { 
    case 0: 
    { 
     cv::Mat tmp; 

     buf.download(tmp); 
    } 
     break; 

    case 1: 
    { 
     cv::Mat tmp; 
     cv::UMat tmp2; 

     buf.download(tmp); 

     tmp.copyTo(tmp2); 
    } 
     break; 

    case 2: 
     buf.copyTo(_dst); 
     break; 

    } 

} 

} 

} 

，并在.CU

#include <opencv2/core/cuda/common.hpp> 


typedef unsigned char uchar; 
typedef unsigned short ushort; 

namespace support 
{ 

namespace cuda 
{ 

namespace device 
{ 

template<class _Ty> 
__global__ void katan(const cv::cuda::PtrStepSz<_Ty>& src, cv::cuda::PtrStep<_Ty> dst) 
{ 

    int x = threadIdx.x + blockIdx.x * blockDim.x; 
    int y = threadIdx.y + blockIdx.y * blockDim.y; 

    if((y>=src.rows) && (x>=src.cols)) 
     return; 


     dst(y,x) = ::atan(static_cast<double>(src(y,x))); 

} 


template<class _Ty> 
void atan_(const cv::cuda::PtrStepSzb& src, cv::cuda::PtrStepSzb dst,cudaStream_t stream) 
{ 

    dim3 block (32, 8); 
    dim3 grid (cv::cuda::device::divUp (src.cols, block.x), cv::cuda::device::divUp (src.rows, block.y)); 


    cudaSafeCall(cudaFuncSetCacheConfig (katan<_Ty>, cudaFuncCachePreferL1)); 
    katan<<<grid, block, 0, stream>>>((cv::cuda::PtrStepSz<_Ty>)src, (cv::cuda::PtrStepSz<_Ty>)dst); 
    cudaSafeCall (cudaGetLastError()); 

    if (stream == 0) 
     cudaSafeCall(cudaDeviceSynchronize()); 

} 

} 

} 

} 

#define INSTANTIATE_ATAN(T) \ 
    template void support::cuda::device::atan_<T>(const cv::cuda::PtrStepSzb&, cv::cuda::PtrStepSzb, cudaStream_t); 

INSTANTIATE_ATAN(uchar) 
INSTANTIATE_ATAN(char) 
INSTANTIATE_ATAN(ushort) 
INSTANTIATE_ATAN(short) 
INSTANTIATE_ATAN(int) 
INSTANTIATE_ATAN(float) 
INSTANTIATE_ATAN(double) 

我为了创建一个最低为例去年ATAN检查它是否有效：

main.cpp中：

#include <iostream> 

#include <opencv2/core.hpp> 
#include <opencv2/core/cuda.hpp> 



#include "atan.hpp" 


int main(int argc,char* argv[]) 
{ 

    typedef float type; 

    cv::Mat_<type> input(32,32); 
    cv::Mat_<type> output(input.size()); 

    std::for_each(input.begin(),input.end(),[&](type& v){ v = cv::theRNG().uniform(1.,100.);}); 

    std::transform(input.begin(),input.end(),output.begin(),[&](const type& v){ return std::atan(v); }); 


    cv::cuda::GpuMat buf; 

    buf.upload(input); 

    support::cuda::atan(buf,buf); 

    cv::Mat tmp; 

    buf.download(tmp); 

    std::cout<<tmp<<std::endl; 

    std::cout << "Hello World!" << std::endl; 
    return EXIT_SUCCESS; 
} 

它编译良好，但是当我尝试执行我有这样的例外：

OpenCV的错误：GPU API调用（遇到非法内存访问）在atan_，文件.. ../bilateral/atan.cu:51：错误：（-217）一个非法的内存访问是/bilateral/atan.cu 51行 抛出“CV ::例外” 什么（）的一个实例后终止叫在功能上遇到atan_

对谷歌的快速调研让我明白这个错误的根源可能是不同的。

我想了解我的代码有什么问题。

在此先感谢您的帮助。

Answer 1

在katan内核，src应由值而不是通过const引用传递。请记住，在主叫侧，src驻留在主机存储器侧，因此通过引用传递不会招致任何主机到设备的存储器复制的，这意味着在katan内核，src也是在主机侧。您不能在内核代码中访问主机内存。按值传递意味着隐式主机到设备的内存拷贝，所以它会没事的。

而且我相信它应该是在if((y>=src.rows) && (x>=src.cols))||？