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) 

создать минимальный Exemple для того, чтобы проверьте, не работает ли он:

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 Error: вызов API графического процессора (незаконный доступ к памяти встречалось) в atan_, файл .. /bilateral/atan.cu, строка 51 завершение вызова после вызова экземпляра 'cv :: Exception' what(): ../bilateral/atan.cu:51: ошибка: (-217) незаконный доступ к памяти был встречается в функции atan_

Быстрое исследование на google позволило мне понять, что происхождение этой ошибки может быть различным.

Я хотел бы понять, что не так с моим кодом.

Заранее благодарим за любую помощь.

Answer 1

В ядре katan значение src должно передаваться по значению вместо ссылки на const. Помните, что со стороны вызывающего абонента src находится в памяти хоста, поэтому при передаче по ссылке не будет никакого хоста для копии памяти устройства, а это значит, что в ядре katansrc также находится на стороне хоста. Вы не можете получить доступ к памяти хоста в коде ядра. Передача по значению подразумевает неявный хост для копии памяти устройства, поэтому все будет в порядке.

И я считаю, что это должно быть || в if((y>=src.rows) && (x>=src.cols))?