I am new in Cuda programming.
What I want to do is just a function what realize the arctan of a matrix. In order to know how to interface with OpenCV I inpire my code from OpenCV's bilateral filter implementation (the one of cudaimgproc).
So I wrote 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
Then 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;
}
}
}
}
and at last 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)
I create a minimum exemple in order to check if it work or not :
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;
}
It compile well but when I try to execute I have this exception :
OpenCV Error: Gpu API call (an illegal memory access was encountered) in atan_, file ../bilateral/atan.cu, line 51 terminate called after throwing an instance of 'cv::Exception' what(): ../bilateral/atan.cu:51: error: (-217) an illegal memory access was encountered in function atan_
A fast research on google let me understood the origin of this error could be various.
I would like to understand what's wrong with my code.
Thanks in advance for any help.
