Addressing CUDA9 shfl deprecation

borisfom · borisfom · commit dde04d5d3e89 · 2017-08-29T13:27:21.000-07:00
diff --git a/modules/core/include/opencv2/core/cuda/warp_shuffle.hpp b/modules/core/include/opencv2/core/cuda/warp_shuffle.hpp
@@ -51,6 +51,11 @@
 
 namespace cv { namespace cuda { namespace device
 {
+#if __CUDACC_VER_MAJOR__ >= 9
+#  define __shfl(x, y, z) __shfl_sync(0xFFFFFFFFU, x, y, z)
+#  define __shfl_up(x, y, z) __shfl_up_sync(0xFFFFFFFFU, x, y, z)
+#  define __shfl_down(x, y, z) __shfl_down_sync(0xFFFFFFFFU, x, y, z)
+#endif
     template <typename T>
     __device__ __forceinline__ T shfl(T val, int srcLane, int width = warpSize)
     {
@@ -148,6 +153,10 @@ namespace cv { namespace cuda { namespace device
     }
 }}}
 
+#  undef __shfl
+#  undef __shfl_up
+#  undef __shfl_down
+
 //! @endcond
 
 #endif // OPENCV_CUDA_WARP_SHUFFLE_HPP

Original file line number	Diff line number	Diff line change
`@@ -51,6 +51,11 @@`
`51`	`51`
`52`	`52`	`namespace cv { namespace cuda { namespace device`
`53`	`53`	`{`
	`54`	`+#if __CUDACC_VER_MAJOR__ >= 9`
	`55`	`+# define __shfl(x, y, z) __shfl_sync(0xFFFFFFFFU, x, y, z)`
	`56`	`+# define __shfl_up(x, y, z) __shfl_up_sync(0xFFFFFFFFU, x, y, z)`
	`57`	`+# define __shfl_down(x, y, z) __shfl_down_sync(0xFFFFFFFFU, x, y, z)`
	`58`	`+#endif`
`54`	`59`	`template <typename T>`
`55`	`60`	`__device__ __forceinline__ T shfl(T val, int srcLane, int width = warpSize)`
`56`	`61`	`{`
`@@ -148,6 +153,10 @@ namespace cv { namespace cuda { namespace device`
`148`	`153`	`}`
`149`	`154`	`}}}`
`150`	`155`
	`156`	`+# undef __shfl`
	`157`	`+# undef __shfl_up`
	`158`	`+# undef __shfl_down`
	`159`	`+`
`151`	`160`	`//! @endcond`
`152`	`161`
`153`	`162`	`#endif // OPENCV_CUDA_WARP_SHUFFLE_HPP`