//! computes exponent of each matrix element (b = e**a)\r
//! supports only CV_32FC1 type\r
CV_EXPORTS void exp(const GpuMat& a, GpuMat& b, Stream& stream = Stream::Null());\r
+ \r
+ //! computes power of each matrix element:\r
+ // (dst(i,j) = pow( src(i,j) , power), if src.type() is integer\r
+ // (dst(i,j) = pow(fabs(src(i,j)), power), otherwise\r
+ //! supports all, except depth == CV_64F\r
+ CV_EXPORTS void pow(const GpuMat& src, double power, GpuMat& dst, Stream& stream = Stream::Null());\r
\r
//! computes natural logarithm of absolute value of each matrix element: b = log(abs(a))\r
//! supports only CV_32FC1 type\r
\r
#include "opencv2/gpu/device/vecmath.hpp"\r
#include "opencv2/gpu/device/transform.hpp"\r
+#include "opencv2/gpu/device/limits_gpu.hpp"\r
#include "opencv2/gpu/device/saturate_cast.hpp"\r
#include "internal_shared.hpp"\r
\r
}\r
\r
template void subtractCaller<short>(const DevMem2D src1, const DevMem2D src2, DevMem2D dst, cudaStream_t stream);\r
+\r
+\r
+ //////////////////////////////////////////////////////////////////////////\r
+ // pow\r
+ \r
+ template<typename T, bool Signed = device::numeric_limits_gpu<T>::is_signed>\r
+ struct PowOp\r
+ { \r
+ float power;\r
+ PowOp(float power_) : power(power_) {}\r
+ \r
+ template<typename T>\r
+ __device__ __forceinline__ T operator()(const T& e) const\r
+ { \r
+ return saturate_cast<T>(__powf((float)e, power));\r
+ } \r
+ };\r
+\r
+ template<typename T>\r
+ struct PowOp<T, true>\r
+ {\r
+ float power;\r
+ PowOp(float power_) : power(power_) {}\r
+\r
+ __device__ __forceinline__ float operator()(const T& e)\r
+ {\r
+ T res = saturate_cast<T>(__powf((float)e, power)); \r
+ \r
+ if ( (e < 0) && (1 & (int)power) )\r
+ res *= -1; \r
+ return res; \r
+ }\r
+ };\r
+\r
+ template<>\r
+ struct PowOp<float>\r
+ {\r
+ float power;\r
+ PowOp(float power_) : power(power_) {}\r
+\r
+ __device__ __forceinline__ float operator()(const float& e)\r
+ {\r
+ return __powf(fabs(e), power);\r
+ }\r
+ };\r
+\r
+ template<typename T>\r
+ void pow_caller(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream)\r
+ {\r
+ transform((DevMem2D_<T>)src, (DevMem2D_<T>)dst, PowOp<T>(power), stream);\r
+ } \r
+\r
+ template void pow_caller<uchar>(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream);\r
+ template void pow_caller<schar>(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream);\r
+ template void pow_caller<short>(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream);\r
+ template void pow_caller<ushort>(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream);\r
+ template void pow_caller<int>(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream);\r
+ template void pow_caller<uint>(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream);\r
+ template void pow_caller<float>(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream);\r
}}}\r
void cv::gpu::max(const GpuMat&, double, GpuMat&, Stream&) { throw_nogpu(); }\r
double cv::gpu::threshold(const GpuMat&, GpuMat&, double, double, int, Stream&) {throw_nogpu(); return 0.0;}\r
\r
+void cv::gpu::pow(const GpuMat&, double, GpuMat&, Stream&) { throw_nogpu(); }\r
+\r
#else\r
\r
////////////////////////////////////////////////////////////////////////\r
return thresh;\r
}\r
\r
+////////////////////////////////////////////////////////////////////////\r
+// pow\r
+\r
+namespace cv\r
+{\r
+ namespace gpu\r
+ {\r
+ namespace mathfunc\r
+ {\r
+ template<typename T>\r
+ void pow_caller(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream);\r
+ }\r
+ }\r
+}\r
+\r
+void cv::gpu::pow(const GpuMat& src, double power, GpuMat& dst, Stream& stream)\r
+{ \r
+ CV_Assert( src.depth() != CV_64F );\r
+ dst.create(src.size(), src.type());\r
+\r
+ typedef void (*caller_t)(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream);\r
+\r
+ static const caller_t callers[] = \r
+ {\r
+ mathfunc::pow_caller<unsigned char>, mathfunc::pow_caller<signed char>, \r
+ mathfunc::pow_caller<unsigned short>, mathfunc::pow_caller<short>, \r
+ mathfunc::pow_caller<int>, mathfunc::pow_caller<float>\r
+ };\r
+\r
+ callers[src.depth()](src.reshape(1), (float)power, dst.reshape(1), StreamAccessor::getStream(stream)); \r
+}\r
+\r
#endif\r
static const bool is_signed = (char)-1 == -1;\r
};\r
\r
+ template<> struct numeric_limits_gpu<signed char>\r
+ {\r
+ typedef char type;\r
+ __device__ __forceinline__ static type min() { return CHAR_MIN; };\r
+ __device__ __forceinline__ static type max() { return CHAR_MAX; };\r
+ __device__ __forceinline__ static type epsilon();\r
+ __device__ __forceinline__ static type round_error();\r
+ __device__ __forceinline__ static type denorm_min();\r
+ __device__ __forceinline__ static type infinity();\r
+ __device__ __forceinline__ static type quiet_NaN();\r
+ __device__ __forceinline__ static type signaling_NaN();\r
+ static const bool is_signed = (signed char)-1 == -1;\r
+ };\r
+\r
template<> struct numeric_limits_gpu<unsigned char>\r
{\r
typedef unsigned char type;\r
levels = (int)::log(static_cast<double>(mm)) * 2 / 3;\r
if (levels == 0) levels++;\r
\r
- nr_plane = (int) ((float) ndisp / pow(2.0, levels + 1));\r
+ nr_plane = (int) ((float) ndisp / std::pow(2.0, levels + 1));\r
}\r
\r
cv::gpu::StereoConstantSpaceBP::StereoConstantSpaceBP(int ndisp_, int iters_, int levels_, int nr_plane_,\r
\r
INSTANTIATE_TEST_CASE_P(Arithm, Exp, testing::ValuesIn(devices()));\r
\r
+\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// pow\r
+\r
+struct Pow : testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, int> >\r
+{\r
+ cv::gpu::DeviceInfo devInfo;\r
+ int type;\r
+\r
+ double power;\r
+ cv::Size size;\r
+ cv::Mat mat;\r
+\r
+ cv::Mat dst_gold;\r
+\r
+ virtual void SetUp() \r
+ { \r
+ devInfo = std::tr1::get<0>(GetParam());\r
+ type = std::tr1::get<1>(GetParam()); \r
+\r
+ cv::gpu::setDevice(devInfo.deviceID());\r
+\r
+ cv::RNG& rng = cvtest::TS::ptr()->get_rng();\r
+\r
+ size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200));\r
+ //size = cv::Size(2, 2);\r
+\r
+ mat = cvtest::randomMat(rng, size, type, 0.0, 100.0, false); \r
+\r
+ if (mat.depth() == CV_32F)\r
+ power = rng.uniform(1.2f, 3.f);\r
+ else\r
+ {\r
+ int ipower = rng.uniform(2, 8);\r
+ power = (float)ipower;\r
+ }\r
+ cv::pow(mat, power, dst_gold);\r
+ }\r
+};\r
+\r
+TEST_P(Pow, Accuracy) \r
+{\r
+ PRINT_PARAM(devInfo);\r
+ PRINT_TYPE(type);\r
+ PRINT_PARAM(size);\r
+ PRINT_PARAM(power);\r
+\r
+ cv::Mat dst;\r
+ \r
+ ASSERT_NO_THROW(\r
+ cv::gpu::GpuMat gpu_res;\r
+\r
+ cv::gpu::pow(cv::gpu::GpuMat(mat), power, gpu_res);\r
+\r
+ gpu_res.download(dst);\r
+ );\r
+\r
+ /*std::cout << mat << std::endl << std::endl;\r
+ std::cout << dst << std::endl << std::endl;\r
+ std::cout << dst_gold << std::endl;*/\r
+ EXPECT_MAT_NEAR(dst_gold, dst, 1);\r
+}\r
+\r
+INSTANTIATE_TEST_CASE_P(Arithm, Pow, testing::Combine(\r
+ testing::ValuesIn(devices()),\r
+ testing::Values(CV_32F, CV_32FC3)));\r
+\r
////////////////////////////////////////////////////////////////////////////////\r
// log\r
\r
#endif\r
\r
int deviceCount = cv::gpu::getCudaEnabledDeviceCount();\r
+ int driver;\r
+ cudaDriverGetVersion(&driver);\r
\r
- printf("CUDA version: %d\n", CUDART_VERSION); \r
+ printf("CUDA Driver version: %d\n", driver); \r
+ printf("CUDA Runtime version: %d\n", CUDART_VERSION); \r
printf("CUDA device count: %d\n\n", deviceCount);\r
+ \r
\r
for (int i = 0; i < deviceCount; ++i)\r
{\r