namespace mat_operators
{
- //////////////////////////////////////////////////////////
- // CopyTo
- //////////////////////////////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
+////////////////////////////////// CopyTo /////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
template<typename T>
__global__ void kernel_copy_to_with_mask(T * mat_src, T * mat_dst, const unsigned char * mask, int cols, int rows, int step_mat, int step_mask, int channels)
}
}
- //////////////////////////////////////////////////////////
- // SetTo
- //////////////////////////////////////////////////////////
+
+///////////////////////////////////////////////////////////////////////////
+////////////////////////////////// SetTo //////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
template<typename T>
__global__ void kernel_set_to_without_mask(T * mat, int cols, int rows, int step, int channels)
}
- //////////////////////////////////////////////////////////
- // ConvertTo
- //////////////////////////////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
+//////////////////////////////// ConvertTo ////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
- template <typename T, typename DT, size_t src_elem_size, size_t dst_elem_size>
- struct Converter
+ template <typename T, typename DT>
+ struct CalcTraits
{
- __device__ static void convert(uchar* srcmat, size_t src_step, uchar* dstmat, size_t dst_step, size_t width, size_t height, double alpha, double beta)
+ __device__ static DT calc(T src, double alpha, double beta)
{
- size_t x = threadIdx.x + blockIdx.x * blockDim.x;
- size_t y = threadIdx.y + blockIdx.y * blockDim.y;
- if (x < width && y < height)
- {
- const T* src = (const T*)(srcmat + src_step * y);
- DT* dst = (DT*)(dstmat + dst_step * y);
-
- dst[x] = (DT)__double2int_rn(alpha * src[x] + beta);
- }
+ return (DT)__double2int_rn(alpha * src + beta);
}
- __host__ static inline dim3 calcGrid(size_t width, size_t height, dim3 block)
+ };
+ template <typename T>
+ struct CalcTraits<T, float>
+ {
+ __device__ static float calc(T src, double alpha, double beta)
{
- return dim3(divUp(width, block.x), divUp(height, block.y));
+ return (float)(alpha * src + beta);
}
};
-
- template <typename T, typename DT>
- struct Converter<T, DT, 1, 1>
+ template <typename T>
+ struct CalcTraits<T, double>
{
- __device__ static void convert(uchar* srcmat, size_t src_step, uchar* dstmat, size_t dst_step, size_t width, size_t height, double alpha, double beta)
+ __device__ static double calc(T src, double alpha, double beta)
{
- size_t x = threadIdx.x + blockIdx.x * blockDim.x;
- size_t y = threadIdx.y + blockIdx.y * blockDim.y;
- if (y < height)
- {
- const T* src = (const T*)(srcmat + src_step * y);
- DT* dst = (DT*)(dstmat + dst_step * y);
- if ((x << 2) + 3 < width)
- {
- uchar4 src4b = ((const uchar4*)src)[x];
- uchar4 dst4b;
-
- const T* src1b = (const T*) &src4b.x;
- DT* dst1b = (DT*) &dst4b.x;
-
- dst1b[0] = (DT)__double2int_rn(alpha * src1b[0] + beta);
- dst1b[1] = (DT)__double2int_rn(alpha * src1b[1] + beta);
- dst1b[2] = (DT)__double2int_rn(alpha * src1b[2] + beta);
- dst1b[3] = (DT)__double2int_rn(alpha * src1b[3] + beta);
-
- ((uchar4*)dst)[x] = dst4b;
- }
- else
- {
- if ((x << 2) + 0 < width)
- dst[(x << 2) + 0] = (DT)__double2int_rn(alpha * src[(x << 2) + 0] + beta);
+ return alpha * src + beta;
+ }
+ };
- if ((x << 2) + 1 < width)
- dst[(x << 2) + 1] = (DT)__double2int_rn(alpha * src[(x << 2) + 1] + beta);
+ template <typename T, typename DT, size_t src_elem_size, size_t dst_elem_size>
+ struct ConverterTraits
+ {
+ enum {shift=1};
- if ((x << 2) + 2 < width)
- dst[(x << 2) + 2] = (DT)__double2int_rn(alpha * src[(x << 2) + 2] + beta);
- }
- }
- }
- __host__ static inline dim3 calcGrid(size_t width, size_t height, dim3 block)
- {
- return dim3(divUp(width, block.x << 2), divUp(height, block.y));
- }
- };/**/
+ typedef T read_type;
+ typedef DT write_type;
+ };
+ template <typename T, typename DT>
+ struct ConverterTraits<T, DT, 1, 1>
+ {
+ enum {shift=4};
+ typedef char4 read_type;
+ typedef char4 write_type;
+ };
template <typename T, typename DT>
- struct Converter<T, DT, 1, 2>
+ struct ConverterTraits<T, DT, 2, 1>
{
- __device__ static void convert(uchar* srcmat, size_t src_step, uchar* dstmat, size_t dst_step, size_t width, size_t height, double alpha, double beta)
- {
- size_t x = threadIdx.x + blockIdx.x * blockDim.x;
- size_t y = threadIdx.y + blockIdx.y * blockDim.y;
- if (y < height)
- {
- const T* src = (const T*)(srcmat + src_step * y);
- DT* dst = (DT*)(dstmat + dst_step * y);
- if ((x << 1) + 1 < width)
- {
- uchar2 src2b = ((const uchar2*)src)[x];
- ushort2 dst2s;
+ enum {shift=4};
- const T* src1b = (const T*) &src2b;
- DT* dst1s = (DT*) &dst2s;
- dst1s[0] = (DT)__double2int_rn(alpha * src1b[0] + beta);
- dst1s[1] = (DT)__double2int_rn(alpha * src1b[1] + beta);
+ typedef short4 read_type;
+ typedef char4 write_type;
+ };
+ template <typename T, typename DT>
+ struct ConverterTraits<T, DT, 4, 1>
+ {
+ enum {shift=4};
- ((ushort2*)(dst))[x] = dst2s;
- }
- else
- {
- if ((x << 1) < width)
- dst[(x << 1)] = (DT)__double2int_rn(alpha * src[(x << 1)] + beta);
- }
- }
- }
- __host__ static inline dim3 calcGrid(size_t width, size_t height, dim3 block)
- {
- return dim3(divUp(width, block.x << 1), divUp(height, block.y));
- }
- };/**/
+ typedef int4 read_type;
+ typedef char4 write_type;
+ };
+ template <typename T, typename DT>
+ struct ConverterTraits<T, DT, 1, 2>
+ {
+ enum {shift=2};
+ typedef char2 read_type;
+ typedef short2 write_type;
+ };
template <typename T, typename DT>
- struct Converter<T, DT, 2, 1>
+ struct ConverterTraits<T, DT, 2, 2>
{
- __device__ static void convert(uchar* srcmat, size_t src_step, uchar* dstmat, size_t dst_step, size_t width, size_t height, double alpha, double beta)
- {
- size_t x = threadIdx.x + blockIdx.x * blockDim.x;
- size_t y = threadIdx.y + blockIdx.y * blockDim.y;
- if (y < height)
- {
- const T* src = (const T*)(srcmat + src_step * y);
- DT* dst = (DT*)(dstmat + dst_step * y);
- if ((x << 2) + 3 < width)
- {
- ushort4 src4s = ((const ushort4*)src)[x];
- uchar4 dst4b;
+ enum {shift=2};
- const T* src1s = (const T*) &src4s.x;
- DT* dst1b = (DT*) &dst4b.x;
- dst1b[0] = (DT)__double2int_rn(alpha * src1s[0] + beta);
- dst1b[1] = (DT)__double2int_rn(alpha * src1s[1] + beta);
- dst1b[2] = (DT)__double2int_rn(alpha * src1s[2] + beta);
- dst1b[3] = (DT)__double2int_rn(alpha * src1s[3] + beta);
+ typedef short2 read_type;
+ typedef short2 write_type;
+ };
+ template <typename T, typename DT>
+ struct ConverterTraits<T, DT, 4, 2>
+ {
+ enum {shift=2};
- ((uchar4*)(dst))[x] = dst4b;
- }
- else
- {
- if ((x << 2) + 0 < width)
- dst[(x << 2) + 0] = (DT)__double2int_rn(alpha * src[(x << 2) + 0] + beta);
- if ((x << 2) + 1 < width)
- dst[(x << 2) + 1] = (DT)__double2int_rn(alpha * src[(x << 2) + 1] + beta);
- if ((x << 2) + 2 < width)
- dst[(x << 2) + 2] = (DT)__double2int_rn(alpha * src[(x << 2) + 2] + beta);
- }
- }
- }
- __host__ static inline dim3 calcGrid(size_t width, size_t height, dim3 block)
- {
- return dim3(divUp(width, block.x << 2), divUp(height, block.y));
- }
- };/**/
+ typedef int2 read_type;
+ typedef short2 write_type;
+ };
template <typename T, typename DT>
- struct Converter<T, DT, 2, 2>
+ struct Converter
{
__device__ static void convert(uchar* srcmat, size_t src_step, uchar* dstmat, size_t dst_step, size_t width, size_t height, double alpha, double beta)
{
{
const T* src = (const T*)(srcmat + src_step * y);
DT* dst = (DT*)(dstmat + dst_step * y);
- if ((x << 1) + 1 < width)
+ if ((x * ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift) + ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift - 1 < width)
{
- ushort2 src2s = ((const ushort2*)src)[x];
- ushort2 dst2s;
+ typename ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::read_type srcn_el = ((const typename ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::read_type*)src)[x];
+ typename ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::write_type dstn_el;
- const T* src1s = (const T*) &src2s.x;
- DT* dst1s = (DT*) &dst2s.x;
- dst1s[0] = (DT)__double2int_rn(alpha * src1s[0] + beta);
- dst1s[1] = (DT)__double2int_rn(alpha * src1s[1] + beta);
+ const T* src1_el = (const T*) &srcn_el;
+ DT* dst1_el = (DT*) &dstn_el;
- ((ushort2*)dst)[x] = dst2s;
+ for (int i = 0; i < ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift; ++i)
+ dst1_el[i] = CalcTraits<T, DT>::calc(src1_el[i], alpha, beta);
+
+ ((typename ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::write_type*)dst)[x] = dstn_el;
}
else
- {
- if ((x << 1) < width)
- dst[(x << 1)] = (DT)__double2int_rn(alpha * src[(x << 1)] + beta);
+ {
+ for (int i = 0; i < ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift - 1; ++i)
+ if ((x * ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift) + i < width)
+ dst[(x * ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift) + i] = CalcTraits<T, DT>::calc(src[(x * ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift) + i], alpha, beta);
}
}
}
__host__ static inline dim3 calcGrid(size_t width, size_t height, dim3 block)
{
- return dim3(divUp(width, block.x << 1), divUp(height, block.y));
- }
- };/**/
-
- template <typename T, size_t src_elem_size, size_t dst_elem_size>
- struct Converter<T, float, src_elem_size, dst_elem_size>
- {
- __device__ static void convert(uchar* srcmat, size_t src_step, uchar* dstmat, size_t dst_step, size_t width, size_t height, double alpha, double beta)
- {
- size_t x = threadIdx.x + blockIdx.x * blockDim.x;
- size_t y = threadIdx.y + blockIdx.y * blockDim.y;
- if (x < width && y < height)
- {
- const T* src = (const T*)(srcmat + src_step * y);
- float* dst = (float*)(dstmat + dst_step * y);
-
- dst[x] = (float)(alpha * src[x] + beta);
- }
- }
- __host__ static inline dim3 calcGrid(size_t width, size_t height, dim3 block)
- {
- return dim3(divUp(width, block.x), divUp(height, block.y));
- }
- };
-
- template <typename T, size_t src_elem_size, size_t dst_elem_size>
- struct Converter<T, double, src_elem_size, dst_elem_size>
- {
- __device__ static void convert(uchar* srcmat, size_t src_step, uchar* dstmat, size_t dst_step, size_t width, size_t height, double alpha, double beta)
- {
- size_t x = threadIdx.x + blockIdx.x * blockDim.x;
- size_t y = threadIdx.y + blockIdx.y * blockDim.y;
- if (x < width && y < height)
- {
- const T* src = (const T*)(srcmat + src_step * y);
- double* dst = (double*)(dstmat + dst_step * y);
-
- dst[x] = (double)(alpha * src[x] + beta);
- }
- }
- __host__ static inline dim3 calcGrid(size_t width, size_t height, dim3 block)
- {
- return dim3(divUp(width, block.x), divUp(height, block.y));
+ return dim3(divUp(width, block.x * ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift), divUp(height, block.y));
}
};
-
- template <typename T, typename DT>
+
+ template <typename T, typename DT>
__global__ static void kernel_convert_to(uchar* srcmat, size_t src_step, uchar* dstmat, size_t dst_step, size_t width, size_t height, double alpha, double beta)
{
- Converter<T, DT, sizeof(T), sizeof(DT)>::convert(srcmat, src_step, dstmat, dst_step, width, height, alpha, beta);
+ Converter<T, DT>::convert(srcmat, src_step, dstmat, dst_step, width, height, alpha, beta);
}
} // namespace mat_operators
namespace impl
{
- //////////////////////////////////////////////////////////////
- // CopyTo
- //////////////////////////////////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
+////////////////////////////////// CopyTo /////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
typedef void (*CopyToFunc)(const DevMem2D& mat_src, const DevMem2D& mat_dst, const DevMem2D& mask, int channels);
}
- //////////////////////////////////////////////////////////////
- // SetTo
- //////////////////////////////////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
+////////////////////////////////// SetTo //////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
typedef void (*SetToFunc_with_mask)(const DevMem2D& mat, const DevMem2D& mask, int channels);
typedef void (*SetToFunc_without_mask)(const DevMem2D& mat, int channels);
func(mat, mask, channels);
}
- //////////////////////////////////////////////////////////////
- // ConvertTo
- //////////////////////////////////////////////////////////////
-
- typedef void (*CvtFunc)(const DevMem2D& src, DevMem2D& dst, size_t width, size_t height, double alpha, double beta);
-
- //#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 130)
-
- template<typename T, typename DT>
- void cvt_(const DevMem2D& src, DevMem2D& dst, size_t width, size_t height, double alpha, double beta)
- {
- dim3 block(32, 8);
- dim3 grid = ::mat_operators::Converter<T, DT, sizeof(T), sizeof(DT)>::calcGrid(width, height, block);
- ::mat_operators::kernel_convert_to<T, DT><<<grid, block>>>(src.ptr, src.step, dst.ptr, dst.step, width, height, alpha, beta);
- cudaSafeCall( cudaThreadSynchronize() );
- }
- //#endif
-
- extern "C" void convert_to(const DevMem2D& src, int sdepth, DevMem2D dst, int ddepth, size_t width, size_t height, double alpha, double beta)
- {
- static CvtFunc tab[8][8] =
- {
- {cvt_<uchar, uchar>, cvt_<uchar, schar>, cvt_<uchar, ushort>, cvt_<uchar, short>,
- cvt_<uchar, int>, cvt_<uchar, float>, cvt_<uchar, double>, 0},
-
- {cvt_<schar, uchar>, cvt_<schar, schar>, cvt_<schar, ushort>, cvt_<schar, short>,
- cvt_<schar, int>, cvt_<schar, float>, cvt_<schar, double>, 0},
-
- {cvt_<ushort, uchar>, cvt_<ushort, schar>, cvt_<ushort, ushort>, cvt_<ushort, short>,
- cvt_<ushort, int>, cvt_<ushort, float>, cvt_<ushort, double>, 0},
-
- {cvt_<short, uchar>, cvt_<short, schar>, cvt_<short, ushort>, cvt_<short, short>,
- cvt_<short, int>, cvt_<short, float>, cvt_<short, double>, 0},
-
- {cvt_<int, uchar>, cvt_<int, schar>, cvt_<int, ushort>,
- cvt_<int, short>, cvt_<int, int>, cvt_<int, float>, cvt_<int, double>, 0},
-
- {cvt_<float, uchar>, cvt_<float, schar>, cvt_<float, ushort>,
- cvt_<float, short>, cvt_<float, int>, cvt_<float, float>, cvt_<float, double>, 0},
-
- {cvt_<double, uchar>, cvt_<double, schar>, cvt_<double, ushort>,
- cvt_<double, short>, cvt_<double, int>, cvt_<double, float>, cvt_<double, double>, 0},
-
- {0,0,0,0,0,0,0,0}
- };
-
- CvtFunc func = tab[sdepth][ddepth];
- if (func == 0)
- error("Operation \'ConvertTo\' doesn't supported on your GPU model", __FILE__, __LINE__);
- func(src, dst, width, height, alpha, beta);
- }
- }
-
-
- }
-}
+
+///////////////////////////////////////////////////////////////////////////
+//////////////////////////////// ConvertTo ////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
+
+ typedef void (*CvtFunc)(const DevMem2D& src, DevMem2D& dst, size_t width, size_t height, double alpha, double beta);
+
+ template<typename T, typename DT>
+ void cvt_(const DevMem2D& src, DevMem2D& dst, size_t width, size_t height, double alpha, double beta)
+ {
+ dim3 block(32, 8);
+ dim3 grid = ::mat_operators::Converter<T, DT>::calcGrid(width, height, block);
+ ::mat_operators::kernel_convert_to<T, DT><<<grid, block>>>(src.ptr, src.step, dst.ptr, dst.step, width, height, alpha, beta);
+ cudaSafeCall( cudaThreadSynchronize() );
+ }
+
+ extern "C" void convert_to(const DevMem2D& src, int sdepth, DevMem2D dst, int ddepth, size_t width, size_t height, double alpha, double beta)
+ {
+ static CvtFunc tab[8][8] =
+ {
+ {cvt_<uchar, uchar>, cvt_<uchar, schar>, cvt_<uchar, ushort>, cvt_<uchar, short>,
+ cvt_<uchar, int>, cvt_<uchar, float>, cvt_<uchar, double>, 0},
+
+ {cvt_<schar, uchar>, cvt_<schar, schar>, cvt_<schar, ushort>, cvt_<schar, short>,
+ cvt_<schar, int>, cvt_<schar, float>, cvt_<schar, double>, 0},
+
+ {cvt_<ushort, uchar>, cvt_<ushort, schar>, cvt_<ushort, ushort>, cvt_<ushort, short>,
+ cvt_<ushort, int>, cvt_<ushort, float>, cvt_<ushort, double>, 0},
+
+ {cvt_<short, uchar>, cvt_<short, schar>, cvt_<short, ushort>, cvt_<short, short>,
+ cvt_<short, int>, cvt_<short, float>, cvt_<short, double>, 0},
+
+ {cvt_<int, uchar>, cvt_<int, schar>, cvt_<int, ushort>,
+ cvt_<int, short>, cvt_<int, int>, cvt_<int, float>, cvt_<int, double>, 0},
+
+ {cvt_<float, uchar>, cvt_<float, schar>, cvt_<float, ushort>,
+ cvt_<float, short>, cvt_<float, int>, cvt_<float, float>, cvt_<float, double>, 0},
+
+ {cvt_<double, uchar>, cvt_<double, schar>, cvt_<double, ushort>,
+ cvt_<double, short>, cvt_<double, int>, cvt_<double, float>, cvt_<double, double>, 0},
+
+ {0,0,0,0,0,0,0,0}
+ };
+
+ CvtFunc func = tab[sdepth][ddepth];
+ if (func == 0)
+ cv::gpu::error("Operation \'ConvertTo\' doesn't supported on your GPU model", __FILE__, __LINE__);
+ func(src, dst, width, height, alpha, beta);
+ }
+ } // namespace impl
+ } // namespace gpu
+} // namespace cv