cuda + opencv 非法内存访问

Question

我是 Cuda 编程新手。

我想做的只是一个实现矩阵arctan的函数。为了知道如何与 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;

    }

}

}

}

最后是 atan .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)

我创建了一个最小示例以检查它是否有效：

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 行在抛出 'cv::Exception' 的实例后终止调用 what(): ../bilateral/atan.cu:51: 错误: (-217) 在函数 atan_

中遇到非法内存访问

对 google 的快速研究让我明白了这个错误的来源可能是多种多样的。

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

在此先感谢您的帮助。

Answer 1

在katan 内核中，src 应该按值而不是const 引用传递。请记住，在调用方，src 驻留在主机内存方，因此通过引用传递不会导致任何主机到设备内存复制，这意味着在 katan 内核中，src也在主机端。您无法在内核代码中访问主机内存。按值传递将暗示主机到设备内存的隐式复制，所以没问题。

而且我认为 if( (y>=src.rows) && (x>=src.cols) ) 中应该是 || ?

cuda + opencv 非法内存访问

cuda + opencv Illegal memory access

c++

opencv

cuda

opencv3.0