cudev: add feature to convert FP32(float) from/to FP16(half) on GPU

* add feature of Fp16 on GPU (cudev)
  * add test
  * leave template function as unimplemented to raise error

modules/cudev/include/opencv2/cudev/functional/functional.hpp

@@ -668,6 +668,27 @@ template <typename T, typename D> struct saturate_cast_func : unary_function<T,
     }
 };
 
+// Convert Fp16 dummy
+template <typename T, typename D> struct saturate_cast_fp16_func;
+
+// Convert Fp16 from Fp32
+template <> struct saturate_cast_fp16_func<float, short> : unary_function<float, short>
+{
+    __device__ __forceinline__ short operator ()(float v) const
+    {
+        return cast_fp16<float, short>(v);
+    }
+};
+
+// Convert Fp16 to Fp32
+template <> struct saturate_cast_fp16_func<short, float> : unary_function<short, float>
+{
+    __device__ __forceinline__ float operator ()(short v) const
+    {
+        return cast_fp16<short, float>(v);
+    }
+};
+
 // Threshold Functors
 
 template <typename T> struct ThreshBinaryFunc : unary_function<T, T>

modules/cudev/include/opencv2/cudev/util/saturate_cast.hpp

@@ -270,6 +270,17 @@ template <> __device__ __forceinline__ uint saturate_cast<uint>(double v)
 #endif
 }
 
+template <typename T, typename D> __device__ __forceinline__ D cast_fp16(T v);
+
+template <> __device__ __forceinline__ float cast_fp16<short, float>(short v)
+{
+    return __half2float(v);
+}
+
+template <> __device__ __forceinline__ short cast_fp16<float, short>(float v)
+{
+    return (short)__float2half_rn(v);
+}
 //! @}
 
 }}

modules/cudev/test/test_cvt.cu

@@ -49,6 +49,7 @@ using namespace cv::cudev;
 using namespace cvtest;
 
 typedef ::testing::Types<uchar, ushort, short, int, float> AllTypes;
+typedef ::testing::Types<short, float> Fp16Types;
 
 ////////////////////////////////////////////////////////////////////////////////
 // CvtTest
@@ -75,9 +76,75 @@ public:
     }
 };
 
+// dummy class
+template <typename T>
+class CvFp16Test : public ::testing::Test
+{
+public:
+    void test_gpumat()
+    {
+    }
+};
+
+template <>
+class CvFp16Test <short> : public ::testing::Test
+{
+public:
+    void test_gpumat()
+    {
+        const Size size = randomSize(100, 400);
+        const int type = DataType<float>::type;
+
+        Mat src = randomMat(size, type), dst, ref;
+
+        GpuMat_<float> g_src(src);
+        GpuMat g_dst;
+
+        // Fp32 -> Fp16
+        convertFp16Cuda(g_src, g_dst);
+        convertFp16Cuda(g_dst.clone(), g_dst);
+        // Fp16 -> Fp32
+        convertFp16(src, dst);
+        convertFp16(dst, ref);
+
+        g_dst.download(dst);
+        EXPECT_MAT_NEAR(dst, ref, 0.0);
+    }
+};
+
+template <>
+class CvFp16Test <float> : public ::testing::Test
+{
+public:
+    void test_gpumat()
+    {
+        const Size size = randomSize(100, 400);
+        const int type = DataType<float>::type;
+
+        Mat src = randomMat(size, type), dst, ref;
+
+        GpuMat_<float> g_src(src);
+        GpuMat g_dst;
+
+        // Fp32 -> Fp16
+        convertFp16Cuda(g_src, g_dst);
+        convertFp16(src, ref);
+
+        g_dst.download(dst);
+        EXPECT_MAT_NEAR(dst, ref, 0.0);
+    }
+};
+
 TYPED_TEST_CASE(CvtTest, AllTypes);
 
 TYPED_TEST(CvtTest, GpuMat)
 {
     CvtTest<TypeParam>::test_gpumat();
 }
+
+TYPED_TEST_CASE(CvFp16Test, Fp16Types);
+
+TYPED_TEST(CvFp16Test, GpuMat)
+{
+    CvFp16Test<TypeParam>::test_gpumat();
+}