#include <stdio.h>
#include "cuda_shared.hpp"
#include "cuda_runtime.h"
+#include "saturate_cast.hpp"
using namespace cv::gpu;
using namespace cv::gpu::impl;
//////////////////////////////// ConvertTo ////////////////////////////////
///////////////////////////////////////////////////////////////////////////
- template <typename T, typename DT>
- struct ScaleTraits
- {
- __device__ static DT scale(T src, double alpha, double beta)
- {
- return (DT)__double2int_rn(alpha * src + beta);
- }
- };
- template <typename T>
- struct ScaleTraits<T, float>
- {
- __device__ static float scale(T src, double alpha, double beta)
- {
- return (float)(alpha * src + beta);
- }
- };
- template <typename T>
- struct ScaleTraits<T, double>
- {
- __device__ static double scale(T src, double alpha, double beta)
- {
- return alpha * src + beta;
- }
- };
-
template <typename T, typename DT, size_t src_elem_size, size_t dst_elem_size>
struct ReadWriteTraits
{
DT* dst1_el = (DT*) &dstn_el;
for (int i = 0; i < shift; ++i)
- dst1_el[i] = ScaleTraits<T, DT>::scale(src1_el[i], alpha, beta);
+ dst1_el[i] = saturate_cast<DT>(alpha * src1_el[i] + beta);
((write_type*)dst)[x] = dstn_el;
}
{
for (int i = 0; i < shift - 1; ++i)
if ((x * shift) + i < width)
- dst[(x * shift) + i] = ScaleTraits<T, DT>::scale(src[(x * shift) + i], alpha, beta);
+ dst[(x * shift) + i] = saturate_cast<DT>(alpha * src[(x * shift) + i] + beta);
}
}
}
#ifndef __OPENCV_GPU_SATURATE_CAST_HPP__
#define __OPENCV_GPU_SATURATE_CAST_HPP__
+#include "cuda_shared.hpp"
-template<class F,T>
-__device__ void saturate_cast<T>(F)
+namespace cv
{
-
+ namespace gpu
+ {
+ template<typename _Tp> __device__ _Tp saturate_cast(uchar v) { return _Tp(v); }\r
+ template<typename _Tp> __device__ _Tp saturate_cast(schar v) { return _Tp(v); }\r
+ template<typename _Tp> __device__ _Tp saturate_cast(ushort v) { return _Tp(v); }\r
+ template<typename _Tp> __device__ _Tp saturate_cast(short v) { return _Tp(v); }\r
+ template<typename _Tp> __device__ _Tp saturate_cast(uint v) { return _Tp(v); }\r
+ template<typename _Tp> __device__ _Tp saturate_cast(int v) { return _Tp(v); }\r
+ template<typename _Tp> __device__ _Tp saturate_cast(float v) { return _Tp(v); }\r
+ template<typename _Tp> __device__ _Tp saturate_cast(double v) { return _Tp(v); }\r
+\r
+ template<> __device__ uchar saturate_cast<uchar>(schar v)\r
+ { return (uchar)max((int)v, 0); }\r
+ template<> __device__ uchar saturate_cast<uchar>(ushort v)\r
+ { return (uchar)min((uint)v, (uint)UCHAR_MAX); }\r
+ template<> __device__ uchar saturate_cast<uchar>(int v)\r
+ { return (uchar)((uint)v <= UCHAR_MAX ? v : v > 0 ? UCHAR_MAX : 0); }\r
+ template<> __device__ uchar saturate_cast<uchar>(uint v)\r
+ { return (uchar)min(v, (uint)UCHAR_MAX); }\r
+ template<> __device__ uchar saturate_cast<uchar>(short v)\r
+ { return saturate_cast<uchar>((uint)v); }\r
+\r
+ template<> __device__ uchar saturate_cast<uchar>(float v)\r
+ { int iv = __float2int_rn(v); return saturate_cast<uchar>(iv); }\r
+ template<> __device__ uchar saturate_cast<uchar>(double v)\r
+ {\r
+ #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130\r
+ int iv = __double2int_rn(v); return saturate_cast<uchar>(iv);\r
+ #else\r
+ return saturate_cast<uchar>((float)v);\r
+ #endif\r
+ }\r
+\r
+ template<> __device__ schar saturate_cast<schar>(uchar v)\r
+ { return (schar)min((int)v, SCHAR_MAX); }\r
+ template<> __device__ schar saturate_cast<schar>(ushort v)\r
+ { return (schar)min((uint)v, (uint)SCHAR_MAX); }\r
+ template<> __device__ schar saturate_cast<schar>(int v)\r
+ {\r
+ return (schar)((uint)(v-SCHAR_MIN) <= (uint)UCHAR_MAX ?\r
+ v : v > 0 ? SCHAR_MAX : SCHAR_MIN);\r
+ }\r
+ template<> __device__ schar saturate_cast<schar>(short v)\r
+ { return saturate_cast<schar>((int)v); }\r
+ template<> __device__ schar saturate_cast<schar>(uint v)\r
+ { return (schar)min(v, (uint)SCHAR_MAX); }\r
+\r
+ template<> __device__ schar saturate_cast<schar>(float v)\r
+ { int iv = __float2int_rn(v); return saturate_cast<schar>(iv); }\r
+ template<> __device__ schar saturate_cast<schar>(double v)\r
+ { \r
+ #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130\r
+ int iv = __double2int_rn(v); return saturate_cast<schar>(iv);\r
+ #else\r
+ return saturate_cast<schar>((float)v);\r
+ #endif\r
+ }\r
+\r
+ template<> __device__ ushort saturate_cast<ushort>(schar v)\r
+ { return (ushort)max((int)v, 0); }\r
+ template<> __device__ ushort saturate_cast<ushort>(short v)\r
+ { return (ushort)max((int)v, 0); }\r
+ template<> __device__ ushort saturate_cast<ushort>(int v)\r
+ { return (ushort)((uint)v <= (uint)USHRT_MAX ? v : v > 0 ? USHRT_MAX : 0); }\r
+ template<> __device__ ushort saturate_cast<ushort>(uint v)\r
+ { return (ushort)min(v, (uint)USHRT_MAX); }\r
+ template<> __device__ ushort saturate_cast<ushort>(float v)\r
+ { int iv = __float2int_rn(v); return saturate_cast<ushort>(iv); }\r
+ template<> __device__ ushort saturate_cast<ushort>(double v)\r
+ { \r
+ #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130\r
+ int iv = __double2int_rn(v); return saturate_cast<ushort>(iv);\r
+ #else\r
+ return saturate_cast<ushort>((float)v);\r
+ #endif\r
+ }\r
+\r
+ template<> __device__ short saturate_cast<short>(ushort v)\r
+ { return (short)min((int)v, SHRT_MAX); }\r
+ template<> __device__ short saturate_cast<short>(int v)\r
+ {\r
+ return (short)((uint)(v - SHRT_MIN) <= (uint)USHRT_MAX ?\r
+ v : v > 0 ? SHRT_MAX : SHRT_MIN);\r
+ }\r
+ template<> __device__ short saturate_cast<short>(uint v)\r
+ { return (short)min(v, (uint)SHRT_MAX); }\r
+ template<> __device__ short saturate_cast<short>(float v)\r
+ { int iv = __float2int_rn(v); return saturate_cast<short>(iv); }\r
+ template<> __device__ short saturate_cast<short>(double v)\r
+ { \r
+ #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130\r
+ int iv = __double2int_rn(v); return saturate_cast<short>(iv);\r
+ #else\r
+ return saturate_cast<short>((float)v);\r
+ #endif\r
+ }\r
+\r
+ template<> __device__ int saturate_cast<int>(float v) { return __float2int_rn(v); }\r
+ template<> __device__ int saturate_cast<int>(double v) \r
+ {\r
+ #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130 \r
+ return __double2int_rn(v);\r
+ #else\r
+ return saturate_cast<int>((float)v);\r
+ #endif\r
+ }\r
+\r
+ template<> __device__ uint saturate_cast<uint>(float v){ return __float2uint_rn(v); }\r
+ template<> __device__ uint saturate_cast<uint>(double v)
+ { \r
+ #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
+ return __double2uint_rn(v);\r
+ #else\r
+ return saturate_cast<uint>((float)v);\r
+ #endif
+ }
+ }
}
#endif /* __OPENCV_GPU_SATURATE_CAST_HPP__ */
\ No newline at end of file
const char* types_str[] = {"CV_8U", "CV_8S", "CV_16U", "CV_16S", "CV_32S", "CV_32F", "CV_64F"};
bool passed = true;
-
- for (int i = 0; i < types_num && passed; ++i)
+ try
{
- for (int j = 0; j < types_num && passed; ++j)
+ for (int i = 0; i < types_num && passed; ++i)
{
- for (int c = 1; c < 2 && passed; ++c)
+ for (int j = 0; j < types_num && passed; ++j)
{
- const int src_type = CV_MAKETYPE(types[i], c);
- const int dst_type = types[j];
- const double alpha = (double)rand() / RAND_MAX * 10.0;
- const double beta = (double)rand() / RAND_MAX * 10.0;
-
- cv::RNG rng(*ts->get_rng());
+ for (int c = 1; c < 2 && passed; ++c)
+ {
+ const int src_type = CV_MAKETYPE(types[i], c);
+ const int dst_type = types[j];
+ const double alpha = (double)rand() / RAND_MAX * 2.0;
+ const double beta = (double)rand() / RAND_MAX * 150.0 - 75;
- Mat cpumatsrc(img_size, src_type);
+ cv::RNG rng(*ts->get_rng());
- rng.fill(cpumatsrc, RNG::UNIFORM, Scalar::all(0), Scalar::all(10));
+ Mat cpumatsrc(img_size, src_type);
- GpuMat gpumatsrc(cpumatsrc);
- Mat cpumatdst;
- GpuMat gpumatdst;
+ rng.fill(cpumatsrc, RNG::UNIFORM, Scalar::all(0), Scalar::all(300));
- cpumatsrc.convertTo(cpumatdst, dst_type, alpha, beta);
+ GpuMat gpumatsrc(cpumatsrc);
+ Mat cpumatdst;
+ GpuMat gpumatdst;
- try
- {
+ cpumatsrc.convertTo(cpumatdst, dst_type, alpha, beta);
gpumatsrc.convertTo(gpumatdst, dst_type, alpha, beta);
- }
- catch(cv::Exception& e)
- {
- cout << "ERROR: " << e.err << endl;
- passed = false;
- break;
- }
-
- double r = norm(cpumatdst, gpumatdst, NORM_INF);
- if (r > 1)
- {
- cout << "FAILED: " << "SRC_TYPE=" << types_str[i] << "C" << c << " DST_TYPE=" << types_str[j] << " NORM = " << r << endl;
-
- passed = false;
+
+ double r = norm(cpumatdst, gpumatdst, NORM_INF);
+ if (r > 1)
+ {
+ cout << "FAILED: " << "SRC_TYPE=" << types_str[i] << "C" << c << " DST_TYPE=" << types_str[j] << " NORM = " << r << endl;
+
+ passed = false;
+ }
}
}
}
}
+ catch(cv::Exception& e)
+ {
+ cout << "ERROR: " << e.err << endl;
+ }
ts->set_failed_test_info(passed ? CvTS::OK : CvTS::FAIL_GENERIC);
}