#include "saturate_cast.hpp"\r
#include "safe_call.hpp"\r
#include "cuda_shared.hpp"\r
+#include "vecmath.hpp"\r
\r
using namespace cv::gpu;\r
\r
\r
namespace filter_krnls\r
{\r
- template <int BLOCK_DIM_X, int BLOCK_DIM_Y, int KERNEL_SIZE, typename T, typename D>\r
+ template <int BLOCK_DIM_X, int BLOCK_DIM_Y, int KERNEL_SIZE, int CN, typename T, typename D>\r
__global__ void linearRowFilter(const T* src, size_t src_step, D* dst, size_t dst_step, int anchor, int width, int height)\r
{\r
__shared__ T smem[BLOCK_DIM_Y * BLOCK_DIM_X * 3];\r
{\r
const T* rowSrc = src + threadY * src_step;\r
\r
- sDataRow[threadIdx.x + blockDim.x] = threadX < width ? rowSrc[threadX] : 0;\r
+ sDataRow[threadIdx.x + blockDim.x] = threadX < width ? rowSrc[threadX] : VecTraits<T>::all(0);\r
\r
- sDataRow[threadIdx.x] = prevThreadX >= 0 ? rowSrc[prevThreadX] : 0;\r
+ sDataRow[threadIdx.x] = prevThreadX >= 0 ? rowSrc[prevThreadX] : VecTraits<T>::all(0);\r
\r
- sDataRow[(blockDim.x << 1) + threadIdx.x] = nextThreadX < width ? rowSrc[nextThreadX] : 0;\r
+ sDataRow[(blockDim.x << 1) + threadIdx.x] = nextThreadX < width ? rowSrc[nextThreadX] : VecTraits<T>::all(0);\r
\r
__syncthreads();\r
\r
if (threadX < width)\r
{\r
- float sum = 0;\r
+ typedef typename TypeVec<float, CN>::vec_t sum_t;\r
+ sum_t sum = VecTraits<sum_t>::all(0);\r
\r
sDataRow += threadIdx.x + blockDim.x - anchor;\r
\r
#pragma unroll\r
for(int i = 0; i < KERNEL_SIZE; ++i)\r
- sum += cLinearKernel[i] * sDataRow[i];\r
+ sum = sum + sDataRow[i] * cLinearKernel[i];\r
\r
dst[threadY * dst_step + threadX] = saturate_cast<D>(sum);\r
}\r
\r
namespace cv { namespace gpu { namespace filters\r
{\r
- template <int KERNEL_SIZE, typename T, typename D>\r
+ template <int KERNEL_SIZE, int CN, typename T, typename D>\r
void linearRowFilter_caller(const DevMem2D_<T>& src, const DevMem2D_<D>& dst, int anchor)\r
{\r
const int BLOCK_DIM_X = 16;\r
dim3 threads(BLOCK_DIM_X, BLOCK_DIM_Y);\r
dim3 blocks(divUp(src.cols, BLOCK_DIM_X), divUp(src.rows, BLOCK_DIM_Y));\r
\r
- filter_krnls::linearRowFilter<BLOCK_DIM_X, BLOCK_DIM_Y, KERNEL_SIZE><<<blocks, threads>>>(src.ptr, src.elem_step, \r
+ filter_krnls::linearRowFilter<BLOCK_DIM_X, BLOCK_DIM_Y, KERNEL_SIZE, CN><<<blocks, threads>>>(src.ptr, src.elem_step, \r
dst.ptr, dst.elem_step, anchor, src.cols, src.rows);\r
\r
cudaSafeCall( cudaThreadSynchronize() );\r
}\r
\r
- template <typename T, typename D>\r
+ template <int CN, typename T, typename D>\r
inline void linearRowFilter_gpu(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
{\r
typedef void (*caller_t)(const DevMem2D_<T>& src, const DevMem2D_<D>& dst, int anchor);\r
static const caller_t callers[] = \r
- {linearRowFilter_caller<0 , T, D>, linearRowFilter_caller<1 , T, D>, \r
- linearRowFilter_caller<2 , T, D>, linearRowFilter_caller<3 , T, D>, \r
- linearRowFilter_caller<4 , T, D>, linearRowFilter_caller<5 , T, D>, \r
- linearRowFilter_caller<6 , T, D>, linearRowFilter_caller<7 , T, D>, \r
- linearRowFilter_caller<8 , T, D>, linearRowFilter_caller<9 , T, D>, \r
- linearRowFilter_caller<10, T, D>, linearRowFilter_caller<11, T, D>, \r
- linearRowFilter_caller<12, T, D>, linearRowFilter_caller<13, T, D>, \r
- linearRowFilter_caller<14, T, D>, linearRowFilter_caller<15, T, D>};\r
+ {linearRowFilter_caller<0 , CN, T, D>, linearRowFilter_caller<1 , CN, T, D>, \r
+ linearRowFilter_caller<2 , CN, T, D>, linearRowFilter_caller<3 , CN, T, D>, \r
+ linearRowFilter_caller<4 , CN, T, D>, linearRowFilter_caller<5 , CN, T, D>, \r
+ linearRowFilter_caller<6 , CN, T, D>, linearRowFilter_caller<7 , CN, T, D>, \r
+ linearRowFilter_caller<8 , CN, T, D>, linearRowFilter_caller<9 , CN, T, D>, \r
+ linearRowFilter_caller<10, CN, T, D>, linearRowFilter_caller<11, CN, T, D>, \r
+ linearRowFilter_caller<12, CN, T, D>, linearRowFilter_caller<13, CN, T, D>, \r
+ linearRowFilter_caller<14, CN, T, D>, linearRowFilter_caller<15, CN, T, D>};\r
\r
loadLinearKernel(kernel, ksize);\r
callers[ksize]((DevMem2D_<T>)src, (DevMem2D_<D>)dst, anchor);\r
}\r
\r
- void linearRowFilter_gpu_32s32s(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ void linearRowFilter_gpu_8u_8u_c4(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
{\r
- linearRowFilter_gpu<int, int>(src, dst, kernel, ksize, anchor);\r
+ linearRowFilter_gpu<4, uchar4, uchar4>(src, dst, kernel, ksize, anchor);\r
}\r
- void linearRowFilter_gpu_32s32f(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ void linearRowFilter_gpu_8u_8s_c4(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
{\r
- linearRowFilter_gpu<int, float>(src, dst, kernel, ksize, anchor);\r
+ linearRowFilter_gpu<4, uchar4, char4>(src, dst, kernel, ksize, anchor);\r
}\r
- void linearRowFilter_gpu_32f32s(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ void linearRowFilter_gpu_8s_8u_c4(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
{\r
- linearRowFilter_gpu<float, int>(src, dst, kernel, ksize, anchor);\r
+ linearRowFilter_gpu<4, char4, uchar4>(src, dst, kernel, ksize, anchor);\r
}\r
- void linearRowFilter_gpu_32f32f(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ void linearRowFilter_gpu_8s_8s_c4(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
{\r
- linearRowFilter_gpu<float, float>(src, dst, kernel, ksize, anchor);\r
+ linearRowFilter_gpu<4, char4, char4>(src, dst, kernel, ksize, anchor);\r
+ }\r
+ void linearRowFilter_gpu_16u_16u_c2(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ {\r
+ linearRowFilter_gpu<2, ushort2, ushort2>(src, dst, kernel, ksize, anchor);\r
+ }\r
+ void linearRowFilter_gpu_16u_16s_c2(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ {\r
+ linearRowFilter_gpu<2, ushort2, short2>(src, dst, kernel, ksize, anchor);\r
+ }\r
+ void linearRowFilter_gpu_16s_16u_c2(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ {\r
+ linearRowFilter_gpu<2, short2, ushort2>(src, dst, kernel, ksize, anchor);\r
+ }\r
+ void linearRowFilter_gpu_16s_16s_c2(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ {\r
+ linearRowFilter_gpu<2, short2, short2>(src, dst, kernel, ksize, anchor);\r
+ }\r
+ void linearRowFilter_gpu_32s_32s_c1(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ {\r
+ linearRowFilter_gpu<1, int, int>(src, dst, kernel, ksize, anchor);\r
+ }\r
+ void linearRowFilter_gpu_32s_32f_c1(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ {\r
+ linearRowFilter_gpu<1, int, float>(src, dst, kernel, ksize, anchor);\r
+ }\r
+ void linearRowFilter_gpu_32f_32s_c1(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ {\r
+ linearRowFilter_gpu<1, float, int>(src, dst, kernel, ksize, anchor);\r
+ }\r
+ void linearRowFilter_gpu_32f_32f_c1(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ {\r
+ linearRowFilter_gpu<1 ,float, float>(src, dst, kernel, ksize, anchor);\r
}\r
}}}\r
\r
namespace filter_krnls\r
{\r
- template <int BLOCK_DIM_X, int BLOCK_DIM_Y, int KERNEL_SIZE, typename T, typename D>\r
+ template <int BLOCK_DIM_X, int BLOCK_DIM_Y, int KERNEL_SIZE, int CN, typename T, typename D>\r
__global__ void linearColumnFilter(const T* src, size_t src_step, D* dst, size_t dst_step, int anchor, int width, int height)\r
{\r
__shared__ T smem[BLOCK_DIM_Y * BLOCK_DIM_X * 3];\r
{\r
const T* colSrc = src + threadX;\r
\r
- sDataColumn[(threadIdx.y + blockDim.y) * smem_step] = threadY < height ? colSrc[threadY * src_step] : 0;\r
+ sDataColumn[(threadIdx.y + blockDim.y) * smem_step] = threadY < height ? colSrc[threadY * src_step] : VecTraits<T>::all(0);\r
\r
- sDataColumn[threadIdx.y * smem_step] = prevThreadY >= 0 ? colSrc[prevThreadY * src_step] : 0;\r
+ sDataColumn[threadIdx.y * smem_step] = prevThreadY >= 0 ? colSrc[prevThreadY * src_step] : VecTraits<T>::all(0);\r
\r
- sDataColumn[(threadIdx.y + (blockDim.y << 1)) * smem_step] = nextThreadY < height ? colSrc[nextThreadY * src_step] : 0;\r
+ sDataColumn[(threadIdx.y + (blockDim.y << 1)) * smem_step] = nextThreadY < height ? colSrc[nextThreadY * src_step] : VecTraits<T>::all(0);\r
\r
__syncthreads();\r
\r
if (threadY < height)\r
{\r
- float sum = 0;\r
+ typedef typename TypeVec<float, CN>::vec_t sum_t;\r
+ sum_t sum = VecTraits<sum_t>::all(0);\r
\r
sDataColumn += (threadIdx.y + blockDim.y - anchor)* smem_step;\r
\r
#pragma unroll\r
for(int i = 0; i < KERNEL_SIZE; ++i)\r
- sum += cLinearKernel[i] * sDataColumn[i * smem_step];\r
+ sum = sum + sDataColumn[i * smem_step] * cLinearKernel[i];\r
\r
dst[threadY * dst_step + threadX] = saturate_cast<D>(sum);\r
}\r
\r
namespace cv { namespace gpu { namespace filters\r
{\r
- template <int KERNEL_SIZE, typename T, typename D>\r
+ template <int KERNEL_SIZE, int CN, typename T, typename D>\r
void linearColumnFilter_caller(const DevMem2D_<T>& src, const DevMem2D_<D>& dst, int anchor)\r
{\r
const int BLOCK_DIM_X = 16;\r
dim3 threads(BLOCK_DIM_X, BLOCK_DIM_Y);\r
dim3 blocks(divUp(src.cols, BLOCK_DIM_X), divUp(src.rows, BLOCK_DIM_Y));\r
\r
- filter_krnls::linearColumnFilter<BLOCK_DIM_X, BLOCK_DIM_Y, KERNEL_SIZE><<<blocks, threads>>>(src.ptr, src.elem_step, \r
+ filter_krnls::linearColumnFilter<BLOCK_DIM_X, BLOCK_DIM_Y, KERNEL_SIZE, CN><<<blocks, threads>>>(src.ptr, src.elem_step, \r
dst.ptr, dst.elem_step, anchor, src.cols, src.rows);\r
\r
cudaSafeCall( cudaThreadSynchronize() );\r
}\r
\r
- template <typename T, typename D>\r
+ template <int CN, typename T, typename D>\r
inline void linearColumnFilter_gpu(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
{\r
typedef void (*caller_t)(const DevMem2D_<T>& src, const DevMem2D_<D>& dst, int anchor);\r
static const caller_t callers[] = \r
- {linearColumnFilter_caller<0 , T, D>, linearColumnFilter_caller<1 , T, D>, \r
- linearColumnFilter_caller<2 , T, D>, linearColumnFilter_caller<3 , T, D>, \r
- linearColumnFilter_caller<4 , T, D>, linearColumnFilter_caller<5 , T, D>, \r
- linearColumnFilter_caller<6 , T, D>, linearColumnFilter_caller<7 , T, D>, \r
- linearColumnFilter_caller<8 , T, D>, linearColumnFilter_caller<9 , T, D>, \r
- linearColumnFilter_caller<10, T, D>, linearColumnFilter_caller<11, T, D>, \r
- linearColumnFilter_caller<12, T, D>, linearColumnFilter_caller<13, T, D>, \r
- linearColumnFilter_caller<14, T, D>, linearColumnFilter_caller<15, T, D>};\r
+ {linearColumnFilter_caller<0 , CN, T, D>, linearColumnFilter_caller<1 , CN, T, D>, \r
+ linearColumnFilter_caller<2 , CN, T, D>, linearColumnFilter_caller<3 , CN, T, D>, \r
+ linearColumnFilter_caller<4 , CN, T, D>, linearColumnFilter_caller<5 , CN, T, D>, \r
+ linearColumnFilter_caller<6 , CN, T, D>, linearColumnFilter_caller<7 , CN, T, D>, \r
+ linearColumnFilter_caller<8 , CN, T, D>, linearColumnFilter_caller<9 , CN, T, D>, \r
+ linearColumnFilter_caller<10, CN, T, D>, linearColumnFilter_caller<11, CN, T, D>, \r
+ linearColumnFilter_caller<12, CN, T, D>, linearColumnFilter_caller<13, CN, T, D>, \r
+ linearColumnFilter_caller<14, CN, T, D>, linearColumnFilter_caller<15, CN, T, D>};\r
\r
loadLinearKernel(kernel, ksize);\r
callers[ksize]((DevMem2D_<T>)src, (DevMem2D_<D>)dst, anchor);\r
}\r
\r
- void linearColumnFilter_gpu_32s32s(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ void linearColumnFilter_gpu_8u_8u_c4(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ {\r
+ linearColumnFilter_gpu<4, uchar4, uchar4>(src, dst, kernel, ksize, anchor);\r
+ }\r
+ void linearColumnFilter_gpu_8u_8s_c4(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ {\r
+ linearColumnFilter_gpu<4, uchar4, char4>(src, dst, kernel, ksize, anchor);\r
+ }\r
+ void linearColumnFilter_gpu_8s_8u_c4(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ {\r
+ linearColumnFilter_gpu<4, char4, uchar4>(src, dst, kernel, ksize, anchor);\r
+ }\r
+ void linearColumnFilter_gpu_8s_8s_c4(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ {\r
+ linearColumnFilter_gpu<4, char4, char4>(src, dst, kernel, ksize, anchor);\r
+ }\r
+ void linearColumnFilter_gpu_16u_16u_c2(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ {\r
+ linearColumnFilter_gpu<2, ushort2, ushort2>(src, dst, kernel, ksize, anchor);\r
+ }\r
+ void linearColumnFilter_gpu_16u_16s_c2(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ {\r
+ linearColumnFilter_gpu<2, ushort2, short2>(src, dst, kernel, ksize, anchor);\r
+ }\r
+ void linearColumnFilter_gpu_16s_16u_c2(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ {\r
+ linearColumnFilter_gpu<2, short2, ushort2>(src, dst, kernel, ksize, anchor);\r
+ }\r
+ void linearColumnFilter_gpu_16s_16s_c2(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ {\r
+ linearColumnFilter_gpu<2, short2, short2>(src, dst, kernel, ksize, anchor);\r
+ }\r
+ void linearColumnFilter_gpu_32s_32s_c1(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
{\r
- linearColumnFilter_gpu<int, int>(src, dst, kernel, ksize, anchor);\r
+ linearColumnFilter_gpu<1, int, int>(src, dst, kernel, ksize, anchor);\r
}\r
- void linearColumnFilter_gpu_32s32f(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ void linearColumnFilter_gpu_32s_32f_c1(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
{\r
- linearColumnFilter_gpu<int, float>(src, dst, kernel, ksize, anchor);\r
+ linearColumnFilter_gpu<1, int, float>(src, dst, kernel, ksize, anchor);\r
}\r
- void linearColumnFilter_gpu_32f32s(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ void linearColumnFilter_gpu_32f_32s_c1(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
{\r
- linearColumnFilter_gpu<float, int>(src, dst, kernel, ksize, anchor);\r
+ linearColumnFilter_gpu<1, float, int>(src, dst, kernel, ksize, anchor);\r
}\r
- void linearColumnFilter_gpu_32f32f(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
+ void linearColumnFilter_gpu_32f_32f_c1(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor)\r
{\r
- linearColumnFilter_gpu<float, float>(src, dst, kernel, ksize, anchor);\r
+ linearColumnFilter_gpu<1, float, float>(src, dst, kernel, ksize, anchor);\r
}\r
}}}\r
\r
return saturate_cast<uint>((float)v);\r
#endif\r
}\r
-\r
- template<typename _Tp> static __device__ _Tp saturate_cast(uchar4 v) { return _Tp(v); }\r
- template<typename _Tp> static __device__ _Tp saturate_cast(char4 v) { return _Tp(v); }\r
- template<typename _Tp> static __device__ _Tp saturate_cast(ushort4 v) { return _Tp(v); }\r
- template<typename _Tp> static __device__ _Tp saturate_cast(short4 v) { return _Tp(v); }\r
- template<typename _Tp> static __device__ _Tp saturate_cast(uint4 v) { return _Tp(v); }\r
- template<typename _Tp> static __device__ _Tp saturate_cast(int4 v) { return _Tp(v); }\r
- template<typename _Tp> static __device__ _Tp saturate_cast(float4 v) { return _Tp(v); }\r
-\r
- template<> static __device__ uchar4 saturate_cast<uchar4>(char4 v)\r
- { return make_uchar4(saturate_cast<uchar>(v.x), saturate_cast<uchar>(v.y), saturate_cast<uchar>(v.z), saturate_cast<uchar>(v.w)); }\r
- template<> static __device__ uchar4 saturate_cast<uchar4>(ushort4 v)\r
- { return make_uchar4(saturate_cast<uchar>(v.x), saturate_cast<uchar>(v.y), saturate_cast<uchar>(v.z), saturate_cast<uchar>(v.w)); }\r
- template<> static __device__ uchar4 saturate_cast<uchar4>(short4 v)\r
- { return make_uchar4(saturate_cast<uchar>(v.x), saturate_cast<uchar>(v.y), saturate_cast<uchar>(v.z), saturate_cast<uchar>(v.w)); }\r
- template<> static __device__ uchar4 saturate_cast<uchar4>(uint4 v)\r
- { return make_uchar4(saturate_cast<uchar>(v.x), saturate_cast<uchar>(v.y), saturate_cast<uchar>(v.z), saturate_cast<uchar>(v.w)); }\r
- template<> static __device__ uchar4 saturate_cast<uchar4>(int4 v)\r
- { return make_uchar4(saturate_cast<uchar>(v.x), saturate_cast<uchar>(v.y), saturate_cast<uchar>(v.z), saturate_cast<uchar>(v.w)); }\r
- template<> static __device__ uchar4 saturate_cast<uchar4>(float4 v)\r
- { return make_uchar4(saturate_cast<uchar>(v.x), saturate_cast<uchar>(v.y), saturate_cast<uchar>(v.z), saturate_cast<uchar>(v.w)); }\r
-\r
- template<> static __device__ char4 saturate_cast<char4>(uchar4 v)\r
- { return make_char4(saturate_cast<char>(v.x), saturate_cast<char>(v.y), saturate_cast<char>(v.z), saturate_cast<char>(v.w)); }\r
- template<> static __device__ char4 saturate_cast<char4>(ushort4 v)\r
- { return make_char4(saturate_cast<char>(v.x), saturate_cast<char>(v.y), saturate_cast<char>(v.z), saturate_cast<char>(v.w)); }\r
- template<> static __device__ char4 saturate_cast<char4>(short4 v)\r
- { return make_char4(saturate_cast<char>(v.x), saturate_cast<char>(v.y), saturate_cast<char>(v.z), saturate_cast<char>(v.w)); }\r
- template<> static __device__ char4 saturate_cast<char4>(uint4 v)\r
- { return make_char4(saturate_cast<char>(v.x), saturate_cast<char>(v.y), saturate_cast<char>(v.z), saturate_cast<char>(v.w)); }\r
- template<> static __device__ char4 saturate_cast<char4>(int4 v)\r
- { return make_char4(saturate_cast<char>(v.x), saturate_cast<char>(v.y), saturate_cast<char>(v.z), saturate_cast<char>(v.w)); }\r
- template<> static __device__ char4 saturate_cast<char4>(float4 v)\r
- { return make_char4(saturate_cast<char>(v.x), saturate_cast<char>(v.y), saturate_cast<char>(v.z), saturate_cast<char>(v.w)); }\r
-\r
- template<> static __device__ ushort4 saturate_cast<ushort4>(uchar4 v)\r
- { return make_ushort4(v.x, v.y, v.z, v.w); }\r
- template<> static __device__ ushort4 saturate_cast<ushort4>(char4 v)\r
- { return make_ushort4(saturate_cast<ushort>(v.x), saturate_cast<ushort>(v.y), saturate_cast<ushort>(v.z), saturate_cast<ushort>(v.w)); }\r
- template<> static __device__ ushort4 saturate_cast<ushort4>(short4 v)\r
- { return make_ushort4(saturate_cast<ushort>(v.x), saturate_cast<ushort>(v.y), saturate_cast<ushort>(v.z), saturate_cast<ushort>(v.w)); }\r
- template<> static __device__ ushort4 saturate_cast<ushort4>(uint4 v)\r
- { return make_ushort4(saturate_cast<ushort>(v.x), saturate_cast<ushort>(v.y), saturate_cast<ushort>(v.z), saturate_cast<ushort>(v.w)); }\r
- template<> static __device__ ushort4 saturate_cast<ushort4>(int4 v)\r
- { return make_ushort4(saturate_cast<ushort>(v.x), saturate_cast<ushort>(v.y), saturate_cast<ushort>(v.z), saturate_cast<ushort>(v.w)); }\r
- template<> static __device__ ushort4 saturate_cast<ushort4>(float4 v)\r
- { return make_ushort4(saturate_cast<ushort>(v.x), saturate_cast<ushort>(v.y), saturate_cast<ushort>(v.z), saturate_cast<ushort>(v.w)); }\r
-\r
- template<> static __device__ short4 saturate_cast<short4>(uchar4 v)\r
- { return make_short4(v.x, v.y, v.z, v.w); }\r
- template<> static __device__ short4 saturate_cast<short4>(char4 v)\r
- { return make_short4(v.x, v.y, v.z, v.w); }\r
- template<> static __device__ short4 saturate_cast<short4>(ushort4 v)\r
- { return make_short4(saturate_cast<short>(v.x), saturate_cast<short>(v.y), saturate_cast<short>(v.z), saturate_cast<short>(v.w)); }\r
- template<> static __device__ short4 saturate_cast<short4>(uint4 v)\r
- { return make_short4(saturate_cast<short>(v.x), saturate_cast<short>(v.y), saturate_cast<short>(v.z), saturate_cast<short>(v.w)); }\r
- template<> static __device__ short4 saturate_cast<short4>(int4 v)\r
- { return make_short4(saturate_cast<short>(v.x), saturate_cast<short>(v.y), saturate_cast<short>(v.z), saturate_cast<short>(v.w)); }\r
- template<> static __device__ short4 saturate_cast<short4>(float4 v)\r
- { return make_short4(saturate_cast<short>(v.x), saturate_cast<short>(v.y), saturate_cast<short>(v.z), saturate_cast<short>(v.w)); }\r
- \r
- template<> static __device__ uint4 saturate_cast<uint4>(uchar4 v)\r
- { return make_uint4(v.x, v.y, v.z, v.w); }\r
- template<> static __device__ uint4 saturate_cast<uint4>(char4 v)\r
- { return make_uint4(saturate_cast<uint>(v.x), saturate_cast<uint>(v.y), saturate_cast<uint>(v.z), saturate_cast<uint>(v.w)); }\r
- template<> static __device__ uint4 saturate_cast<uint4>(ushort4 v)\r
- { return make_uint4(v.x, v.y, v.z, v.w); }\r
- template<> static __device__ uint4 saturate_cast<uint4>(short4 v)\r
- { return make_uint4(saturate_cast<uint>(v.x), saturate_cast<uint>(v.y), saturate_cast<uint>(v.z), saturate_cast<uint>(v.w)); }\r
- template<> static __device__ uint4 saturate_cast<uint4>(int4 v)\r
- { return make_uint4(saturate_cast<uint>(v.x), saturate_cast<uint>(v.y), saturate_cast<uint>(v.z), saturate_cast<uint>(v.w)); }\r
- template<> static __device__ uint4 saturate_cast<uint4>(float4 v)\r
- { return make_uint4(saturate_cast<uint>(v.x), saturate_cast<uint>(v.y), saturate_cast<uint>(v.z), saturate_cast<uint>(v.w)); }\r
- \r
- template<> static __device__ int4 saturate_cast<int4>(uchar4 v)\r
- { return make_int4(v.x, v.y, v.z, v.w); }\r
- template<> static __device__ int4 saturate_cast<int4>(char4 v)\r
- { return make_int4(v.x, v.y, v.z, v.w); }\r
- template<> static __device__ int4 saturate_cast<int4>(ushort4 v)\r
- { return make_int4(v.x, v.y, v.z, v.w); }\r
- template<> static __device__ int4 saturate_cast<int4>(short4 v)\r
- { return make_int4(v.x, v.y, v.z, v.w); }\r
- template<> static __device__ int4 saturate_cast<int4>(uint4 v)\r
- { return make_int4(saturate_cast<int>(v.x), saturate_cast<int>(v.y), saturate_cast<int>(v.z), saturate_cast<int>(v.w)); }\r
- template<> static __device__ int4 saturate_cast<int4>(float4 v)\r
- { return make_int4(saturate_cast<int>(v.x), saturate_cast<int>(v.y), saturate_cast<int>(v.z), saturate_cast<int>(v.w)); }\r
}\r
}\r
\r
#define __OPENCV_GPU_VECMATH_HPP__\r
\r
#include "cuda_shared.hpp"\r
+#include "saturate_cast.hpp"\r
\r
namespace cv\r
{\r
namespace gpu\r
{\r
template<typename T, int N> struct TypeVec;\r
- template<typename T> struct TypeVec<T, 1> { typedef T vec_t; };\r
- template<> struct TypeVec<unsigned char, 2> { typedef uchar2 vec_t; };\r
+\r
+ template<> struct TypeVec<uchar, 1> { typedef uchar vec_t; };\r
+ template<> struct TypeVec<uchar1, 1> { typedef uchar1 vec_t; };\r
+ template<> struct TypeVec<uchar, 2> { typedef uchar2 vec_t; };\r
template<> struct TypeVec<uchar2, 2> { typedef uchar2 vec_t; };\r
- template<> struct TypeVec<unsigned char, 3> { typedef uchar3 vec_t; };;\r
+ template<> struct TypeVec<uchar, 3> { typedef uchar3 vec_t; };\r
template<> struct TypeVec<uchar3, 3> { typedef uchar3 vec_t; };\r
- template<> struct TypeVec<unsigned char, 4> { typedef uchar4 vec_t; };;\r
+ template<> struct TypeVec<uchar, 4> { typedef uchar4 vec_t; };\r
template<> struct TypeVec<uchar4, 4> { typedef uchar4 vec_t; };\r
+\r
+ template<> struct TypeVec<char, 1> { typedef char vec_t; };\r
+ template<> struct TypeVec<char1, 1> { typedef char1 vec_t; };\r
template<> struct TypeVec<char, 2> { typedef char2 vec_t; };\r
template<> struct TypeVec<char2, 2> { typedef char2 vec_t; };\r
template<> struct TypeVec<char, 3> { typedef char3 vec_t; };\r
template<> struct TypeVec<char3, 3> { typedef char3 vec_t; };\r
template<> struct TypeVec<char, 4> { typedef char4 vec_t; };\r
template<> struct TypeVec<char4, 4> { typedef char4 vec_t; };\r
- template<> struct TypeVec<unsigned short, 2> { typedef ushort2 vec_t; };\r
+\r
+ template<> struct TypeVec<ushort, 1> { typedef ushort vec_t; };\r
+ template<> struct TypeVec<ushort1, 1> { typedef ushort1 vec_t; };\r
+ template<> struct TypeVec<ushort, 2> { typedef ushort2 vec_t; };\r
template<> struct TypeVec<ushort2, 2> { typedef ushort2 vec_t; };\r
- template<> struct TypeVec<unsigned short, 3> { typedef ushort3 vec_t; };\r
+ template<> struct TypeVec<ushort, 3> { typedef ushort3 vec_t; };\r
template<> struct TypeVec<ushort3, 3> { typedef ushort3 vec_t; };\r
- template<> struct TypeVec<unsigned short, 4> { typedef ushort4 vec_t; };\r
+ template<> struct TypeVec<ushort, 4> { typedef ushort4 vec_t; };\r
template<> struct TypeVec<ushort4, 4> { typedef ushort4 vec_t; };\r
+\r
+ template<> struct TypeVec<short, 1> { typedef short vec_t; };\r
+ template<> struct TypeVec<short1, 1> { typedef short1 vec_t; };\r
template<> struct TypeVec<short, 2> { typedef short2 vec_t; };\r
template<> struct TypeVec<short2, 2> { typedef short2 vec_t; };\r
template<> struct TypeVec<short, 3> { typedef short3 vec_t; };\r
template<> struct TypeVec<short3, 3> { typedef short3 vec_t; };\r
template<> struct TypeVec<short, 4> { typedef short4 vec_t; };\r
template<> struct TypeVec<short4, 4> { typedef short4 vec_t; };\r
- template<> struct TypeVec<unsigned int, 2> { typedef uint2 vec_t; };\r
+\r
+ template<> struct TypeVec<uint, 1> { typedef uint vec_t; };\r
+ template<> struct TypeVec<uint1, 1> { typedef uint1 vec_t; };\r
+ template<> struct TypeVec<uint, 2> { typedef uint2 vec_t; };\r
template<> struct TypeVec<uint2, 2> { typedef uint2 vec_t; };\r
- template<> struct TypeVec<unsigned int, 3> { typedef uint3 vec_t; };\r
+ template<> struct TypeVec<uint, 3> { typedef uint3 vec_t; };\r
template<> struct TypeVec<uint3, 3> { typedef uint3 vec_t; };\r
- template<> struct TypeVec<unsigned int, 4> { typedef uint4 vec_t; };\r
+ template<> struct TypeVec<uint, 4> { typedef uint4 vec_t; };\r
template<> struct TypeVec<uint4, 4> { typedef uint4 vec_t; };\r
+\r
+ template<> struct TypeVec<int, 1> { typedef int vec_t; };\r
+ template<> struct TypeVec<int1, 1> { typedef int1 vec_t; };\r
template<> struct TypeVec<int, 2> { typedef int2 vec_t; };\r
template<> struct TypeVec<int2, 2> { typedef int2 vec_t; };\r
template<> struct TypeVec<int, 3> { typedef int3 vec_t; };\r
template<> struct TypeVec<int3, 3> { typedef int3 vec_t; };\r
template<> struct TypeVec<int, 4> { typedef int4 vec_t; };\r
template<> struct TypeVec<int4, 4> { typedef int4 vec_t; };\r
+\r
+ template<> struct TypeVec<float, 1> { typedef float vec_t; };\r
+ template<> struct TypeVec<float1, 1> { typedef float1 vec_t; };\r
template<> struct TypeVec<float, 2> { typedef float2 vec_t; };\r
template<> struct TypeVec<float2, 2> { typedef float2 vec_t; };\r
template<> struct TypeVec<float, 3> { typedef float3 vec_t; };\r
template<> struct TypeVec<float3, 3> { typedef float3 vec_t; };\r
template<> struct TypeVec<float, 4> { typedef float4 vec_t; };\r
- template<> struct TypeVec<float4, 4> { typedef float4 vec_t; }; \r
+ template<> struct TypeVec<float4, 4> { typedef float4 vec_t; };\r
+\r
+ template<typename T> struct VecTraits;\r
+\r
+ template<> struct VecTraits<uchar> \r
+ { \r
+ typedef uchar elem_t; \r
+ enum {cn=1};\r
+ static __device__ uchar all(uchar v) {return v;}\r
+ };\r
+ template<> struct VecTraits<uchar1> \r
+ { \r
+ typedef uchar elem_t; \r
+ enum {cn=1};\r
+ static __device__ uchar1 all(uchar v) {return make_uchar1(v);}\r
+ };\r
+ template<> struct VecTraits<uchar2> \r
+ { \r
+ typedef uchar elem_t; \r
+ enum {cn=2}; \r
+ static __device__ uchar2 all(uchar v) {return make_uchar2(v, v);}\r
+ };\r
+ template<> struct VecTraits<uchar3> \r
+ { \r
+ typedef uchar elem_t; \r
+ enum {cn=3}; \r
+ static __device__ uchar3 all(uchar v) {return make_uchar3(v, v, v);}\r
+ };\r
+ template<> struct VecTraits<uchar4> \r
+ { \r
+ typedef uchar elem_t; \r
+ enum {cn=4}; \r
+ static __device__ uchar4 all(uchar v) {return make_uchar4(v, v, v, v);}\r
+ };\r
+\r
+ template<> struct VecTraits<char> \r
+ { \r
+ typedef char elem_t; \r
+ enum {cn=1}; \r
+ static __device__ char all(char v) {return v;}\r
+ };\r
+ template<> struct VecTraits<char1> \r
+ { \r
+ typedef char elem_t; \r
+ enum {cn=1}; \r
+ static __device__ char1 all(char v) {return make_char1(v);}\r
+ };\r
+ template<> struct VecTraits<char2> \r
+ { \r
+ typedef char elem_t; \r
+ enum {cn=2}; \r
+ static __device__ char2 all(char v) {return make_char2(v, v);}\r
+ };\r
+ template<> struct VecTraits<char3> \r
+ { \r
+ typedef char elem_t; \r
+ enum {cn=3}; \r
+ static __device__ char3 all(char v) {return make_char3(v, v, v);}\r
+ };\r
+ template<> struct VecTraits<char4> \r
+ { \r
+ typedef char elem_t; \r
+ enum {cn=4}; \r
+ static __device__ char4 all(char v) {return make_char4(v, v, v, v);}\r
+ };\r
+\r
+ template<> struct VecTraits<ushort> \r
+ { \r
+ typedef ushort elem_t; \r
+ enum {cn=1}; \r
+ static __device__ ushort all(ushort v) {return v;}\r
+ };\r
+ template<> struct VecTraits<ushort1> \r
+ { \r
+ typedef ushort elem_t; \r
+ enum {cn=1}; \r
+ static __device__ ushort1 all(ushort v) {return make_ushort1(v);}\r
+ };\r
+ template<> struct VecTraits<ushort2> \r
+ { \r
+ typedef ushort elem_t; \r
+ enum {cn=2}; \r
+ static __device__ ushort2 all(ushort v) {return make_ushort2(v, v);}\r
+ };\r
+ template<> struct VecTraits<ushort3> \r
+ { \r
+ typedef ushort elem_t; \r
+ enum {cn=3}; \r
+ static __device__ ushort3 all(ushort v) {return make_ushort3(v, v, v);}\r
+ };\r
+ template<> struct VecTraits<ushort4> \r
+ { \r
+ typedef ushort elem_t; \r
+ enum {cn=4}; \r
+ static __device__ ushort4 all(ushort v) {return make_ushort4(v, v, v, v);}\r
+ };\r
+\r
+ template<> struct VecTraits<short> \r
+ { \r
+ typedef short elem_t; \r
+ enum {cn=1}; \r
+ static __device__ short all(short v) {return v;}\r
+ };\r
+ template<> struct VecTraits<short1> \r
+ { \r
+ typedef short elem_t; \r
+ enum {cn=1}; \r
+ static __device__ short1 all(short v) {return make_short1(v);}\r
+ };\r
+ template<> struct VecTraits<short2> \r
+ { \r
+ typedef short elem_t; \r
+ enum {cn=2}; \r
+ static __device__ short2 all(short v) {return make_short2(v, v);}\r
+ };\r
+ template<> struct VecTraits<short3> \r
+ { \r
+ typedef short elem_t; \r
+ enum {cn=3}; \r
+ static __device__ short3 all(short v) {return make_short3(v, v, v);}\r
+ };\r
+ template<> struct VecTraits<short4> \r
+ { \r
+ typedef short elem_t; \r
+ enum {cn=4}; \r
+ static __device__ short4 all(short v) {return make_short4(v, v, v, v);}\r
+ };\r
+\r
+ template<> struct VecTraits<uint> \r
+ { \r
+ typedef uint elem_t; \r
+ enum {cn=1}; \r
+ static __device__ uint all(uint v) {return v;}\r
+ };\r
+ template<> struct VecTraits<uint1> \r
+ { \r
+ typedef uint elem_t; \r
+ enum {cn=1}; \r
+ static __device__ uint1 all(uint v) {return make_uint1(v);}\r
+ };\r
+ template<> struct VecTraits<uint2> \r
+ { \r
+ typedef uint elem_t; \r
+ enum {cn=2}; \r
+ static __device__ uint2 all(uint v) {return make_uint2(v, v);}\r
+ };\r
+ template<> struct VecTraits<uint3> \r
+ { \r
+ typedef uint elem_t; \r
+ enum {cn=3}; \r
+ static __device__ uint3 all(uint v) {return make_uint3(v, v, v);}\r
+ };\r
+ template<> struct VecTraits<uint4> \r
+ { \r
+ typedef uint elem_t; \r
+ enum {cn=4}; \r
+ static __device__ uint4 all(uint v) {return make_uint4(v, v, v, v);}\r
+ };\r
+\r
+ template<> struct VecTraits<int> \r
+ { \r
+ typedef int elem_t; \r
+ enum {cn=1}; \r
+ static __device__ int all(int v) {return v;}\r
+ };\r
+ template<> struct VecTraits<int1> \r
+ { \r
+ typedef int elem_t; \r
+ enum {cn=1}; \r
+ static __device__ int1 all(int v) {return make_int1(v);}\r
+ };\r
+ template<> struct VecTraits<int2> \r
+ { \r
+ typedef int elem_t; \r
+ enum {cn=2}; \r
+ static __device__ int2 all(int v) {return make_int2(v, v);}\r
+ };\r
+ template<> struct VecTraits<int3> \r
+ { \r
+ typedef int elem_t; \r
+ enum {cn=3}; \r
+ static __device__ int3 all(int v) {return make_int3(v, v, v);}\r
+ };\r
+ template<> struct VecTraits<int4> \r
+ { \r
+ typedef int elem_t; \r
+ enum {cn=4}; \r
+ static __device__ int4 all(int v) {return make_int4(v, v, v, v);}\r
+ };\r
+\r
+ template<> struct VecTraits<float> \r
+ { \r
+ typedef float elem_t; \r
+ enum {cn=1}; \r
+ static __device__ float all(float v) {return v;}\r
+ };\r
+ template<> struct VecTraits<float1> \r
+ { \r
+ typedef float elem_t; \r
+ enum {cn=1}; \r
+ static __device__ float1 all(float v) {return make_float1(v);}\r
+ };\r
+ template<> struct VecTraits<float2> \r
+ { \r
+ typedef float elem_t; \r
+ enum {cn=2}; \r
+ static __device__ float2 all(float v) {return make_float2(v, v);}\r
+ };\r
+ template<> struct VecTraits<float3> \r
+ { \r
+ typedef float elem_t; \r
+ enum {cn=3}; \r
+ static __device__ float3 all(float v) {return make_float3(v, v, v);}\r
+ };\r
+ template<> struct VecTraits<float4> \r
+ { \r
+ typedef float elem_t;\r
+ enum {cn=4}; \r
+ static __device__ float4 all(float v) {return make_float4(v, v, v, v);}\r
+ };\r
+\r
+ template <int cn, typename VecD> struct SatCast;\r
+ template <typename VecD> struct SatCast<1, VecD>\r
+ {\r
+ template <typename VecS>\r
+ __device__ VecD operator()(const VecS& v)\r
+ {\r
+ VecD res; \r
+ res.x = saturate_cast< VecTraits<VecD>::elem_t >(v.x);\r
+ return res;\r
+ }\r
+ };\r
+ template <typename VecD> struct SatCast<2, VecD>\r
+ {\r
+ template <typename VecS>\r
+ __device__ VecD operator()(const VecS& v)\r
+ {\r
+ VecD res; \r
+ res.x = saturate_cast< VecTraits<VecD>::elem_t >(v.x);\r
+ res.y = saturate_cast< VecTraits<VecD>::elem_t >(v.y);\r
+ return res;\r
+ }\r
+ };\r
+ template <typename VecD> struct SatCast<3, VecD>\r
+ {\r
+ template <typename VecS>\r
+ __device__ VecD operator()(const VecS& v)\r
+ {\r
+ VecD res; \r
+ res.x = saturate_cast< VecTraits<VecD>::elem_t >(v.x);\r
+ res.y = saturate_cast< VecTraits<VecD>::elem_t >(v.y);\r
+ res.y = saturate_cast< VecTraits<VecD>::elem_t >(v.z);\r
+ return res;\r
+ }\r
+ };\r
+ template <typename VecD> struct SatCast<4, VecD>\r
+ {\r
+ template <typename VecS>\r
+ __device__ VecD operator()(const VecS& v)\r
+ {\r
+ VecD res; \r
+ res.x = saturate_cast< VecTraits<VecD>::elem_t >(v.x);\r
+ res.y = saturate_cast< VecTraits<VecD>::elem_t >(v.y);\r
+ res.y = saturate_cast< VecTraits<VecD>::elem_t >(v.z);\r
+ res.w = saturate_cast< VecTraits<VecD>::elem_t >(v.w);\r
+ return res;\r
+ }\r
+ };\r
+\r
+ template <typename VecD, typename VecS> static __device__ VecD saturate_cast_caller(const VecS& v)\r
+ {\r
+ SatCast<VecTraits<VecD>::cn, VecD> cast;\r
+ return cast(v);\r
+ }\r
+\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const uchar1& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const char1& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const ushort1& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const short1& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const uint1& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const int1& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const float1& v) {return saturate_cast_caller<_Tp>(v);}\r
+\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const uchar2& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const char2& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const ushort2& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const short2& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const uint2& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const int2& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const float2& v) {return saturate_cast_caller<_Tp>(v);}\r
+\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const uchar3& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const char3& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const ushort3& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const short3& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const uint3& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const int3& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const float3& v) {return saturate_cast_caller<_Tp>(v);}\r
+\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const uchar4& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const char4& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const ushort4& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const short4& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const uint4& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const int4& v) {return saturate_cast_caller<_Tp>(v);}\r
+ template<typename _Tp> static __device__ _Tp saturate_cast(const float4& v) {return saturate_cast_caller<_Tp>(v);}\r
+\r
+ static __device__ uchar1 operator+(const uchar1& a, const uchar1& b)\r
+ {\r
+ return make_uchar1(a.x + b.x);\r
+ }\r
+ static __device__ uchar1 operator-(const uchar1& a, const uchar1& b)\r
+ {\r
+ return make_uchar1(a.x - b.x);\r
+ }\r
+ static __device__ uchar1 operator*(const uchar1& a, const uchar1& b)\r
+ {\r
+ return make_uchar1(a.x * b.x);\r
+ }\r
+ static __device__ uchar1 operator/(const uchar1& a, const uchar1& b)\r
+ {\r
+ return make_uchar1(a.x / b.x);\r
+ }\r
+ static __device__ float1 operator*(const uchar1& a, float s)\r
+ {\r
+ return make_float1(a.x * s);\r
+ }\r
+\r
+ static __device__ uchar2 operator+(const uchar2& a, const uchar2& b)\r
+ {\r
+ return make_uchar2(a.x + b.x, a.y + b.y);\r
+ }\r
+ static __device__ uchar2 operator-(const uchar2& a, const uchar2& b)\r
+ {\r
+ return make_uchar2(a.x - b.x, a.y - b.y);\r
+ }\r
+ static __device__ uchar2 operator*(const uchar2& a, const uchar2& b)\r
+ {\r
+ return make_uchar2(a.x * b.x, a.y * b.y);\r
+ }\r
+ static __device__ uchar2 operator/(const uchar2& a, const uchar2& b)\r
+ {\r
+ return make_uchar2(a.x / b.x, a.y / b.y);\r
+ }\r
+ static __device__ float2 operator*(const uchar2& a, float s)\r
+ {\r
+ return make_float2(a.x * s, a.y * s);\r
+ }\r
\r
- static __device__ uchar4 operator+(const uchar4& a, const uchar4& b)\r
+ static __device__ uchar3 operator+(const uchar3& a, const uchar3& b)\r
+ {\r
+ return make_uchar3(a.x + b.x, a.y + b.y, a.z + b.z);\r
+ }\r
+ static __device__ uchar3 operator-(const uchar3& a, const uchar3& b)\r
+ {\r
+ return make_uchar3(a.x - b.x, a.y - b.y, a.z - b.z);\r
+ }\r
+ static __device__ uchar3 operator*(const uchar3& a, const uchar3& b)\r
+ {\r
+ return make_uchar3(a.x * b.x, a.y * b.y, a.z * b.z);\r
+ }\r
+ static __device__ uchar3 operator/(const uchar3& a, const uchar3& b)\r
+ {\r
+ return make_uchar3(a.x / b.x, a.y / b.y, a.z / b.z);\r
+ }\r
+ static __device__ float3 operator*(const uchar3& a, float s)\r
+ {\r
+ return make_float3(a.x * s, a.y * s, a.z * s);\r
+ }\r
+\r
+ static __device__ uchar4 operator+(const uchar4& a, const uchar4& b)\r
{\r
return make_uchar4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w);\r
}\r
- static __device__ uchar4 operator-(const uchar4& a, const uchar4& b)\r
+ static __device__ uchar4 operator-(const uchar4& a, const uchar4& b)\r
{\r
return make_uchar4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w);\r
}\r
- static __device__ uchar4 operator*(const uchar4& a, const uchar4& b)\r
+ static __device__ uchar4 operator*(const uchar4& a, const uchar4& b)\r
{\r
return make_uchar4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w);\r
}\r
- static __device__ uchar4 operator/(const uchar4& a, const uchar4& b)\r
+ static __device__ uchar4 operator/(const uchar4& a, const uchar4& b)\r
{\r
return make_uchar4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w);\r
}\r
- template <typename T>\r
- static __device__ uchar4 operator*(const uchar4& a, T s)\r
+ static __device__ float4 operator*(const uchar4& a, float s)\r
+ {\r
+ return make_float4(a.x * s, a.y * s, a.z * s, a.w * s);\r
+ }\r
+\r
+ static __device__ char1 operator+(const char1& a, const char1& b)\r
+ {\r
+ return make_char1(a.x + b.x);\r
+ }\r
+ static __device__ char1 operator-(const char1& a, const char1& b)\r
+ {\r
+ return make_char1(a.x - b.x);\r
+ }\r
+ static __device__ char1 operator*(const char1& a, const char1& b)\r
+ {\r
+ return make_char1(a.x * b.x);\r
+ }\r
+ static __device__ char1 operator/(const char1& a, const char1& b)\r
+ {\r
+ return make_char1(a.x / b.x);\r
+ }\r
+ static __device__ float1 operator*(const char1& a, float s)\r
+ {\r
+ return make_float1(a.x * s);\r
+ }\r
+\r
+ static __device__ char2 operator+(const char2& a, const char2& b)\r
+ {\r
+ return make_char2(a.x + b.x, a.y + b.y);\r
+ }\r
+ static __device__ char2 operator-(const char2& a, const char2& b)\r
+ {\r
+ return make_char2(a.x - b.x, a.y - b.y);\r
+ }\r
+ static __device__ char2 operator*(const char2& a, const char2& b)\r
+ {\r
+ return make_char2(a.x * b.x, a.y * b.y);\r
+ }\r
+ static __device__ char2 operator/(const char2& a, const char2& b)\r
+ {\r
+ return make_char2(a.x / b.x, a.y / b.y);\r
+ }\r
+ static __device__ float2 operator*(const char2& a, float s)\r
+ {\r
+ return make_float2(a.x * s, a.y * s);\r
+ }\r
+\r
+ static __device__ char3 operator+(const char3& a, const char3& b)\r
+ {\r
+ return make_char3(a.x + b.x, a.y + b.y, a.z + b.z);\r
+ }\r
+ static __device__ char3 operator-(const char3& a, const char3& b)\r
+ {\r
+ return make_char3(a.x - b.x, a.y - b.y, a.z - b.z);\r
+ }\r
+ static __device__ char3 operator*(const char3& a, const char3& b)\r
+ {\r
+ return make_char3(a.x * b.x, a.y * b.y, a.z * b.z);\r
+ }\r
+ static __device__ char3 operator/(const char3& a, const char3& b)\r
+ {\r
+ return make_char3(a.x / b.x, a.y / b.y, a.z / b.z);\r
+ }\r
+ static __device__ float3 operator*(const char3& a, float s)\r
+ {\r
+ return make_float3(a.x * s, a.y * s, a.z * s);\r
+ }\r
+\r
+ static __device__ char4 operator+(const char4& a, const char4& b)\r
+ {\r
+ return make_char4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w);\r
+ }\r
+ static __device__ char4 operator-(const char4& a, const char4& b)\r
+ {\r
+ return make_char4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w);\r
+ }\r
+ static __device__ char4 operator*(const char4& a, const char4& b)\r
+ {\r
+ return make_char4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w);\r
+ }\r
+ static __device__ char4 operator/(const char4& a, const char4& b)\r
+ {\r
+ return make_char4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w);\r
+ }\r
+ static __device__ float4 operator*(const char4& a, float s)\r
+ {\r
+ return make_float4(a.x * s, a.y * s, a.z * s, a.w * s);\r
+ }\r
+\r
+ static __device__ ushort1 operator+(const ushort1& a, const ushort1& b)\r
+ {\r
+ return make_ushort1(a.x + b.x);\r
+ }\r
+ static __device__ ushort1 operator-(const ushort1& a, const ushort1& b)\r
+ {\r
+ return make_ushort1(a.x - b.x);\r
+ }\r
+ static __device__ ushort1 operator*(const ushort1& a, const ushort1& b)\r
+ {\r
+ return make_ushort1(a.x * b.x);\r
+ }\r
+ static __device__ ushort1 operator/(const ushort1& a, const ushort1& b)\r
+ {\r
+ return make_ushort1(a.x / b.x);\r
+ }\r
+ static __device__ float1 operator*(const ushort1& a, float s)\r
+ {\r
+ return make_float1(a.x * s);\r
+ }\r
+\r
+ static __device__ ushort2 operator+(const ushort2& a, const ushort2& b)\r
+ {\r
+ return make_ushort2(a.x + b.x, a.y + b.y);\r
+ }\r
+ static __device__ ushort2 operator-(const ushort2& a, const ushort2& b)\r
+ {\r
+ return make_ushort2(a.x - b.x, a.y - b.y);\r
+ }\r
+ static __device__ ushort2 operator*(const ushort2& a, const ushort2& b)\r
+ {\r
+ return make_ushort2(a.x * b.x, a.y * b.y);\r
+ }\r
+ static __device__ ushort2 operator/(const ushort2& a, const ushort2& b)\r
+ {\r
+ return make_ushort2(a.x / b.x, a.y / b.y);\r
+ }\r
+ static __device__ float2 operator*(const ushort2& a, float s)\r
+ {\r
+ return make_float2(a.x * s, a.y * s);\r
+ }\r
+\r
+ static __device__ ushort3 operator+(const ushort3& a, const ushort3& b)\r
+ {\r
+ return make_ushort3(a.x + b.x, a.y + b.y, a.z + b.z);\r
+ }\r
+ static __device__ ushort3 operator-(const ushort3& a, const ushort3& b)\r
+ {\r
+ return make_ushort3(a.x - b.x, a.y - b.y, a.z - b.z);\r
+ }\r
+ static __device__ ushort3 operator*(const ushort3& a, const ushort3& b)\r
+ {\r
+ return make_ushort3(a.x * b.x, a.y * b.y, a.z * b.z);\r
+ }\r
+ static __device__ ushort3 operator/(const ushort3& a, const ushort3& b)\r
+ {\r
+ return make_ushort3(a.x / b.x, a.y / b.y, a.z / b.z);\r
+ }\r
+ static __device__ float3 operator*(const ushort3& a, float s)\r
+ {\r
+ return make_float3(a.x * s, a.y * s, a.z * s);\r
+ }\r
+\r
+ static __device__ ushort4 operator+(const ushort4& a, const ushort4& b)\r
+ {\r
+ return make_ushort4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w);\r
+ }\r
+ static __device__ ushort4 operator-(const ushort4& a, const ushort4& b)\r
+ {\r
+ return make_ushort4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w);\r
+ }\r
+ static __device__ ushort4 operator*(const ushort4& a, const ushort4& b)\r
+ {\r
+ return make_ushort4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w);\r
+ }\r
+ static __device__ ushort4 operator/(const ushort4& a, const ushort4& b)\r
+ {\r
+ return make_ushort4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w);\r
+ }\r
+ static __device__ float4 operator*(const ushort4& a, float s)\r
+ {\r
+ return make_float4(a.x * s, a.y * s, a.z * s, a.w * s);\r
+ }\r
+\r
+ static __device__ short1 operator+(const short1& a, const short1& b)\r
+ {\r
+ return make_short1(a.x + b.x);\r
+ }\r
+ static __device__ short1 operator-(const short1& a, const short1& b)\r
+ {\r
+ return make_short1(a.x - b.x);\r
+ }\r
+ static __device__ short1 operator*(const short1& a, const short1& b)\r
+ {\r
+ return make_short1(a.x * b.x);\r
+ }\r
+ static __device__ short1 operator/(const short1& a, const short1& b)\r
+ {\r
+ return make_short1(a.x / b.x);\r
+ }\r
+ static __device__ float1 operator*(const short1& a, float s)\r
+ {\r
+ return make_float1(a.x * s);\r
+ }\r
+\r
+ static __device__ short2 operator+(const short2& a, const short2& b)\r
+ {\r
+ return make_short2(a.x + b.x, a.y + b.y);\r
+ }\r
+ static __device__ short2 operator-(const short2& a, const short2& b)\r
+ {\r
+ return make_short2(a.x - b.x, a.y - b.y);\r
+ }\r
+ static __device__ short2 operator*(const short2& a, const short2& b)\r
+ {\r
+ return make_short2(a.x * b.x, a.y * b.y);\r
+ }\r
+ static __device__ short2 operator/(const short2& a, const short2& b)\r
+ {\r
+ return make_short2(a.x / b.x, a.y / b.y);\r
+ }\r
+ static __device__ float2 operator*(const short2& a, float s)\r
+ {\r
+ return make_float2(a.x * s, a.y * s);\r
+ }\r
+\r
+ static __device__ short3 operator+(const short3& a, const short3& b)\r
+ {\r
+ return make_short3(a.x + b.x, a.y + b.y, a.z + b.z);\r
+ }\r
+ static __device__ short3 operator-(const short3& a, const short3& b)\r
+ {\r
+ return make_short3(a.x - b.x, a.y - b.y, a.z - b.z);\r
+ }\r
+ static __device__ short3 operator*(const short3& a, const short3& b)\r
+ {\r
+ return make_short3(a.x * b.x, a.y * b.y, a.z * b.z);\r
+ }\r
+ static __device__ short3 operator/(const short3& a, const short3& b)\r
+ {\r
+ return make_short3(a.x / b.x, a.y / b.y, a.z / b.z);\r
+ }\r
+ static __device__ float3 operator*(const short3& a, float s)\r
+ {\r
+ return make_float3(a.x * s, a.y * s, a.z * s);\r
+ }\r
+\r
+ static __device__ short4 operator+(const short4& a, const short4& b)\r
+ {\r
+ return make_short4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w);\r
+ }\r
+ static __device__ short4 operator-(const short4& a, const short4& b)\r
+ {\r
+ return make_short4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w);\r
+ }\r
+ static __device__ short4 operator*(const short4& a, const short4& b)\r
+ {\r
+ return make_short4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w);\r
+ }\r
+ static __device__ short4 operator/(const short4& a, const short4& b)\r
+ {\r
+ return make_short4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w);\r
+ }\r
+ static __device__ float4 operator*(const short4& a, float s)\r
+ {\r
+ return make_float4(a.x * s, a.y * s, a.z * s, a.w * s);\r
+ }\r
+\r
+ static __device__ int1 operator+(const int1& a, const int1& b)\r
+ {\r
+ return make_int1(a.x + b.x);\r
+ }\r
+ static __device__ int1 operator-(const int1& a, const int1& b)\r
+ {\r
+ return make_int1(a.x - b.x);\r
+ }\r
+ static __device__ int1 operator*(const int1& a, const int1& b)\r
+ {\r
+ return make_int1(a.x * b.x);\r
+ }\r
+ static __device__ int1 operator/(const int1& a, const int1& b)\r
+ {\r
+ return make_int1(a.x / b.x);\r
+ }\r
+ static __device__ float1 operator*(const int1& a, float s)\r
+ {\r
+ return make_float1(a.x * s);\r
+ }\r
+\r
+ static __device__ int2 operator+(const int2& a, const int2& b)\r
+ {\r
+ return make_int2(a.x + b.x, a.y + b.y);\r
+ }\r
+ static __device__ int2 operator-(const int2& a, const int2& b)\r
+ {\r
+ return make_int2(a.x - b.x, a.y - b.y);\r
+ }\r
+ static __device__ int2 operator*(const int2& a, const int2& b)\r
+ {\r
+ return make_int2(a.x * b.x, a.y * b.y);\r
+ }\r
+ static __device__ int2 operator/(const int2& a, const int2& b)\r
+ {\r
+ return make_int2(a.x / b.x, a.y / b.y);\r
+ }\r
+ static __device__ float2 operator*(const int2& a, float s)\r
+ {\r
+ return make_float2(a.x * s, a.y * s);\r
+ }\r
+\r
+ static __device__ int3 operator+(const int3& a, const int3& b)\r
+ {\r
+ return make_int3(a.x + b.x, a.y + b.y, a.z + b.z);\r
+ }\r
+ static __device__ int3 operator-(const int3& a, const int3& b)\r
+ {\r
+ return make_int3(a.x - b.x, a.y - b.y, a.z - b.z);\r
+ }\r
+ static __device__ int3 operator*(const int3& a, const int3& b)\r
+ {\r
+ return make_int3(a.x * b.x, a.y * b.y, a.z * b.z);\r
+ }\r
+ static __device__ int3 operator/(const int3& a, const int3& b)\r
+ {\r
+ return make_int3(a.x / b.x, a.y / b.y, a.z / b.z);\r
+ }\r
+ static __device__ float3 operator*(const int3& a, float s)\r
+ {\r
+ return make_float3(a.x * s, a.y * s, a.z * s);\r
+ }\r
+\r
+ static __device__ int4 operator+(const int4& a, const int4& b)\r
+ {\r
+ return make_int4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w);\r
+ }\r
+ static __device__ int4 operator-(const int4& a, const int4& b)\r
+ {\r
+ return make_int4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w);\r
+ }\r
+ static __device__ int4 operator*(const int4& a, const int4& b)\r
+ {\r
+ return make_int4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w);\r
+ }\r
+ static __device__ int4 operator/(const int4& a, const int4& b)\r
+ {\r
+ return make_int4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w);\r
+ }\r
+ static __device__ float4 operator*(const int4& a, float s)\r
+ {\r
+ return make_float4(a.x * s, a.y * s, a.z * s, a.w * s);\r
+ }\r
+\r
+ static __device__ float1 operator+(const float1& a, const float1& b)\r
+ {\r
+ return make_float1(a.x + b.x);\r
+ }\r
+ static __device__ float1 operator-(const float1& a, const float1& b)\r
+ {\r
+ return make_float1(a.x - b.x);\r
+ }\r
+ static __device__ float1 operator*(const float1& a, const float1& b)\r
+ {\r
+ return make_float1(a.x * b.x);\r
+ }\r
+ static __device__ float1 operator/(const float1& a, const float1& b)\r
+ {\r
+ return make_float1(a.x / b.x);\r
+ }\r
+ static __device__ float1 operator*(const float1& a, float s)\r
+ {\r
+ return make_float1(a.x * s);\r
+ }\r
+\r
+ static __device__ float2 operator+(const float2& a, const float2& b)\r
+ {\r
+ return make_float2(a.x + b.x, a.y + b.y);\r
+ }\r
+ static __device__ float2 operator-(const float2& a, const float2& b)\r
+ {\r
+ return make_float2(a.x - b.x, a.y - b.y);\r
+ }\r
+ static __device__ float2 operator*(const float2& a, const float2& b)\r
+ {\r
+ return make_float2(a.x * b.x, a.y * b.y);\r
+ }\r
+ static __device__ float2 operator/(const float2& a, const float2& b)\r
+ {\r
+ return make_float2(a.x / b.x, a.y / b.y);\r
+ }\r
+ static __device__ float2 operator*(const float2& a, float s)\r
+ {\r
+ return make_float2(a.x * s, a.y * s);\r
+ }\r
+\r
+ static __device__ float3 operator+(const float3& a, const float3& b)\r
+ {\r
+ return make_float3(a.x + b.x, a.y + b.y, a.z + b.z);\r
+ }\r
+ static __device__ float3 operator-(const float3& a, const float3& b)\r
+ {\r
+ return make_float3(a.x - b.x, a.y - b.y, a.z - b.z);\r
+ }\r
+ static __device__ float3 operator*(const float3& a, const float3& b)\r
+ {\r
+ return make_float3(a.x * b.x, a.y * b.y, a.z * b.z);\r
+ }\r
+ static __device__ float3 operator/(const float3& a, const float3& b)\r
+ {\r
+ return make_float3(a.x / b.x, a.y / b.y, a.z / b.z);\r
+ }\r
+ static __device__ float3 operator*(const float3& a, float s)\r
+ {\r
+ return make_float3(a.x * s, a.y * s, a.z * s);\r
+ }\r
+\r
+ static __device__ float4 operator+(const float4& a, const float4& b)\r
+ {\r
+ return make_float4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w);\r
+ }\r
+ static __device__ float4 operator-(const float4& a, const float4& b)\r
+ {\r
+ return make_float4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w);\r
+ }\r
+ static __device__ float4 operator*(const float4& a, const float4& b)\r
+ {\r
+ return make_float4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w);\r
+ }\r
+ static __device__ float4 operator/(const float4& a, const float4& b)\r
{\r
- return make_uchar4(a.x * s, a.y * s, a.z * s, a.w * s);\r
+ return make_float4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w);\r
}\r
- template <typename T>\r
- static __device__ uchar4 operator*(T s, const uchar4& a)\r
+ static __device__ float4 operator*(const float4& a, float s)\r
{\r
- return a * s;\r
+ return make_float4(a.x * s, a.y * s, a.z * s, a.w * s);\r
}\r
}\r
}\r
void cv::gpu::Stream::enqueueDownload(const GpuMat& src, Mat& dst)\r
{\r
// if not -> allocation will be done, but after that dst will not point to page locked memory\r
- CV_Assert(src.cols == dst.cols && src.rows == dst.rows && src.type() == dst.type() )\r
+ CV_Assert(src.cols == dst.cols && src.rows == dst.rows && src.type() == dst.type() );\r
devcopy(src, dst, impl->stream, cudaMemcpyDeviceToHost);\r
}\r
void cv::gpu::Stream::enqueueDownload(const GpuMat& src, CudaMem& dst) { devcopy(src, dst, impl->stream, cudaMemcpyDeviceToHost); }\r
\r
namespace cv { namespace gpu { namespace filters\r
{\r
- void linearRowFilter_gpu_32s32s(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
- void linearRowFilter_gpu_32s32f(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
- void linearRowFilter_gpu_32f32s(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
- void linearRowFilter_gpu_32f32f(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
-\r
- void linearColumnFilter_gpu_32s32s(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
- void linearColumnFilter_gpu_32s32f(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
- void linearColumnFilter_gpu_32f32s(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
- void linearColumnFilter_gpu_32f32f(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearRowFilter_gpu_8u_8u_c4(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearRowFilter_gpu_8u_8s_c4(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearRowFilter_gpu_8s_8u_c4(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearRowFilter_gpu_8s_8s_c4(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearRowFilter_gpu_16u_16u_c2(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearRowFilter_gpu_16u_16s_c2(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearRowFilter_gpu_16s_16u_c2(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearRowFilter_gpu_16s_16s_c2(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearRowFilter_gpu_32s_32s_c1(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearRowFilter_gpu_32s_32f_c1(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearRowFilter_gpu_32f_32s_c1(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearRowFilter_gpu_32f_32f_c1(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+\r
+ void linearColumnFilter_gpu_8u_8u_c4(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearColumnFilter_gpu_8u_8s_c4(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearColumnFilter_gpu_8s_8u_c4(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearColumnFilter_gpu_8s_8s_c4(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearColumnFilter_gpu_16u_16u_c2(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearColumnFilter_gpu_16u_16s_c2(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearColumnFilter_gpu_16s_16u_c2(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearColumnFilter_gpu_16s_16s_c2(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearColumnFilter_gpu_32s_32s_c1(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearColumnFilter_gpu_32s_32f_c1(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearColumnFilter_gpu_32f_32s_c1(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
+ void linearColumnFilter_gpu_32f_32f_c1(const DevMem2D& src, const DevMem2D& dst, const float kernel[], int ksize, int anchor);\r
}}}\r
\r
namespace\r
static const nppFilter1D_t nppFilter1D_callers[] = {0, nppiFilterRow_8u_C1R, 0, 0, nppiFilterRow_8u_C4R};\r
static const gpuFilter1D_t gpuFilter1D_callers[6][6] =\r
{\r
- {0,0,0,0,0,0},\r
- {0,0,0,0,0,0},\r
- {0,0,0,0,0,0},\r
- {0,0,0,0,0,0},\r
- {0,0,0,0,linearRowFilter_gpu_32s32s, linearRowFilter_gpu_32s32f},\r
- {0,0,0,0,linearRowFilter_gpu_32f32s, linearRowFilter_gpu_32f32f}\r
+ {linearRowFilter_gpu_8u_8u_c4,linearRowFilter_gpu_8u_8s_c4,0,0,0,0},\r
+ {linearRowFilter_gpu_8s_8u_c4,linearRowFilter_gpu_8s_8s_c4,0,0,0,0},\r
+ {0,0,linearRowFilter_gpu_16u_16u_c2,linearRowFilter_gpu_16u_16s_c2,0,0},\r
+ {0,0,linearRowFilter_gpu_16s_16u_c2,linearRowFilter_gpu_16s_16s_c2,0,0},\r
+ {0,0,0,0,linearRowFilter_gpu_32s_32s_c1, linearRowFilter_gpu_32s_32f_c1},\r
+ {0,0,0,0,linearRowFilter_gpu_32f_32s_c1, linearRowFilter_gpu_32f_32f_c1}\r
};\r
\r
- if ((srcType == CV_8UC1 || srcType == CV_8UC4) && bufType == srcType)\r
+ if ((bufType == srcType) && (srcType == CV_8UC1 || srcType == CV_8UC4))\r
{\r
GpuMat gpu_row_krnl;\r
int nDivisor;\r
return Ptr<BaseRowFilter_GPU>(new NppLinearRowFilter(ksize, anchor, gpu_row_krnl, nDivisor,\r
nppFilter1D_callers[CV_MAT_CN(srcType)]));\r
}\r
- else if ((srcType == CV_32SC1 || srcType == CV_32FC1) && (bufType == CV_32SC1 || bufType == CV_32FC1))\r
- {\r
- Mat temp(rowKernel.size(), CV_32FC1);\r
- rowKernel.convertTo(temp, CV_32FC1);\r
- Mat cont_krnl = temp.reshape(1, 1);\r
\r
- int ksize = cont_krnl.cols;\r
- normalizeAnchor(anchor, ksize);\r
+ CV_Assert(srcType == CV_8UC4 || srcType == CV_8SC4 || srcType == CV_16UC2 || srcType == CV_16SC2 || srcType == CV_32SC1 || srcType == CV_32FC1);\r
+ CV_Assert(bufType == CV_8UC4 || bufType == CV_8SC4 || bufType == CV_16UC2 || bufType == CV_16SC2 || bufType == CV_32SC1 || bufType == CV_32FC1);\r
\r
- return Ptr<BaseRowFilter_GPU>(new GpuLinearRowFilter(ksize, anchor, cont_krnl, \r
- gpuFilter1D_callers[CV_MAT_DEPTH(srcType)][CV_MAT_DEPTH(bufType)]));\r
- }\r
+ Mat temp(rowKernel.size(), CV_32FC1);\r
+ rowKernel.convertTo(temp, CV_32FC1);\r
+ Mat cont_krnl = temp.reshape(1, 1);\r
\r
- CV_Assert(!"Unsupported types"); \r
- return Ptr<BaseRowFilter_GPU>(0);\r
+ int ksize = cont_krnl.cols;\r
+ normalizeAnchor(anchor, ksize);\r
+\r
+ return Ptr<BaseRowFilter_GPU>(new GpuLinearRowFilter(ksize, anchor, cont_krnl, \r
+ gpuFilter1D_callers[CV_MAT_DEPTH(srcType)][CV_MAT_DEPTH(bufType)]));\r
}\r
\r
namespace\r
static const nppFilter1D_t nppFilter1D_callers[] = {0, nppiFilterColumn_8u_C1R, 0, 0, nppiFilterColumn_8u_C4R};\r
static const gpuFilter1D_t gpuFilter1D_callers[6][6] =\r
{\r
- {0,0,0,0,0,0},\r
- {0,0,0,0,0,0},\r
- {0,0,0,0,0,0},\r
- {0,0,0,0,0,0},\r
- {0,0,0,0,linearColumnFilter_gpu_32s32s, linearColumnFilter_gpu_32s32f},\r
- {0,0,0,0,linearColumnFilter_gpu_32f32s, linearColumnFilter_gpu_32f32f}\r
+ {linearColumnFilter_gpu_8u_8u_c4,linearColumnFilter_gpu_8u_8s_c4,0,0,0,0},\r
+ {linearColumnFilter_gpu_8s_8u_c4,linearColumnFilter_gpu_8s_8s_c4,0,0,0,0},\r
+ {0,0,linearColumnFilter_gpu_16u_16u_c2,linearColumnFilter_gpu_16u_16s_c2,0,0},\r
+ {0,0,linearColumnFilter_gpu_16s_16u_c2,linearColumnFilter_gpu_16s_16s_c2,0,0},\r
+ {0,0,0,0,linearColumnFilter_gpu_32s_32s_c1, linearColumnFilter_gpu_32s_32f_c1},\r
+ {0,0,0,0,linearColumnFilter_gpu_32f_32s_c1, linearColumnFilter_gpu_32f_32f_c1}\r
};\r
\r
- if ((bufType == CV_8UC1 || bufType == CV_8UC4) && dstType == bufType)\r
+ double kernelMin;\r
+ minMaxLoc(columnKernel, &kernelMin);\r
+ \r
+ if ((bufType == dstType) && (bufType == CV_8UC1 || bufType == CV_8UC4))\r
{\r
GpuMat gpu_col_krnl;\r
int nDivisor;\r
return Ptr<BaseColumnFilter_GPU>(new NppLinearColumnFilter(ksize, anchor, gpu_col_krnl, nDivisor, \r
nppFilter1D_callers[CV_MAT_CN(bufType)]));\r
}\r
- else if ((bufType == CV_32SC1 || bufType == CV_32FC1) && (dstType == CV_32SC1 || dstType == CV_32FC1))\r
- {\r
- Mat temp(columnKernel.size(), CV_32FC1);\r
- columnKernel.convertTo(temp, CV_32FC1);\r
- Mat cont_krnl = temp.reshape(1, 1);\r
\r
- int ksize = cont_krnl.cols;\r
- normalizeAnchor(anchor, ksize);\r
+ CV_Assert(dstType == CV_8UC4 || dstType == CV_8SC4 || dstType == CV_16UC2 || dstType == CV_16SC2 || dstType == CV_32SC1 || dstType == CV_32FC1);\r
+ CV_Assert(bufType == CV_8UC4 || bufType == CV_8SC4 || bufType == CV_16UC2 || bufType == CV_16SC2 || bufType == CV_32SC1 || bufType == CV_32FC1);\r
\r
- return Ptr<BaseColumnFilter_GPU>(new GpuLinearColumnFilter(ksize, anchor, cont_krnl, \r
- gpuFilter1D_callers[CV_MAT_DEPTH(bufType)][CV_MAT_DEPTH(dstType)]));\r
- }\r
+ Mat temp(columnKernel.size(), CV_32FC1);\r
+ columnKernel.convertTo(temp, CV_32FC1);\r
+ Mat cont_krnl = temp.reshape(1, 1);\r
+\r
+ int ksize = cont_krnl.cols;\r
+ normalizeAnchor(anchor, ksize);\r
\r
- CV_Assert(!"Unsupported types"); \r
- return Ptr<BaseColumnFilter_GPU>(0);\r
+ return Ptr<BaseColumnFilter_GPU>(new GpuLinearColumnFilter(ksize, anchor, cont_krnl, \r
+ gpuFilter1D_callers[CV_MAT_DEPTH(bufType)][CV_MAT_DEPTH(dstType)]));\r
}\r
\r
Ptr<FilterEngine_GPU> cv::gpu::createSeparableLinearFilter_GPU(int srcType, int dstType, const Mat& rowKernel, const Mat& columnKernel, \r
\r
double cv::gpu::threshold(const GpuMat& src, GpuMat& dst, double thresh) \r
{ \r
- CV_Assert(src.type() == CV_32FC1)\r
+ CV_Assert(src.type() == CV_32FC1);\r
\r
dst.create( src.size(), src.type() );\r
\r
\r
int test(const Mat& img)\r
{\r
+ if (img.type() != CV_8UC1)\r
+ return CvTS::OK;\r
int ksizes[] = {3, 5, 7};\r
int ksizes_num = sizeof(ksizes) / sizeof(int);\r
\r
cv::Sobel(img, cpudst, -1, dx, dy, ksizes[i]);\r
\r
GpuMat gpu1(img);\r
+ gpu1.convertTo(gpu1, CV_32S);\r
GpuMat gpudst;\r
cv::gpu::Sobel(gpu1, gpudst, -1, dx, dy, ksizes[i]);\r
+ gpudst.convertTo(gpudst, CV_8U);\r
\r
if (CheckNorm(cpudst, gpudst, Size(ksizes[i], ksizes[i])) != CvTS::OK)\r
test_res = CvTS::FAIL_GENERIC;\r
\r
int test(const Mat& img)\r
{\r
+ if (img.type() != CV_8UC1)\r
+ return CvTS::OK;\r
+\r
int dx = 1, dy = 0;\r
\r
Mat cpudst;\r
cv::Scharr(img, cpudst, -1, dx, dy);\r
\r
GpuMat gpu1(img);\r
+ gpu1.convertTo(gpu1, CV_32S);\r
GpuMat gpudst;\r
cv::gpu::Scharr(gpu1, gpudst, -1, dx, dy);\r
-\r
+ gpudst.convertTo(gpudst, CV_8U);\r
+ \r
return CheckNorm(cpudst, gpudst, Size(3, 3));\r
}\r
};\r
projects / platform / upstream / opencv.git / commitdiff