Merge branch 'gpu-device-layer'

modules/core/CMakeLists.txt

@@ -3,14 +3,14 @@ ocv_add_module(core ${ZLIB_LIBRARIES})
 ocv_module_include_directories(${ZLIB_INCLUDE_DIR})
 
 if(HAVE_CUDA)
-  ocv_source_group("Src\\Cuda" GLOB "src/cuda/*.cu")  
-  ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/gpu/src" "${OpenCV_SOURCE_DIR}/modules/gpu/src/cuda" ${CUDA_INCLUDE_DIRS})
+  ocv_source_group("Src\\Cuda" GLOB "src/cuda/*.cu")
+  ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/gpu/include" ${CUDA_INCLUDE_DIRS})
   ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef)
-  
+
   file(GLOB lib_cuda "src/cuda/*.cu")
   ocv_cuda_compile(cuda_objs ${lib_cuda})
 
-  
+
   set(cuda_link_libs ${CUDA_LIBRARIES} ${CUDA_npp_LIBRARY})
 else()
   set(lib_cuda "")

modules/gpu/CMakeLists.txt

@@ -7,11 +7,11 @@ ocv_add_module(gpu opencv_imgproc opencv_calib3d opencv_objdetect opencv_video o
 
 ocv_module_include_directories("${CMAKE_CURRENT_SOURCE_DIR}/src/cuda" "${CMAKE_CURRENT_SOURCE_DIR}/../highgui/src")
 
-file(GLOB lib_hdrs               "include/opencv2/${name}/*.hpp" "include/opencv2/${name}/*.h")
-file(GLOB lib_int_hdrs           "src/*.hpp" "src/*.h")
+file(GLOB lib_hdrs               "include/opencv2/${name}/*.hpp"               "include/opencv2/${name}/*.h")
+file(GLOB lib_device_hdrs        "include/opencv2/${name}/device/*.hpp"        "include/opencv2/${name}/device/*.h")
+file(GLOB lib_device_hdrs_detail "include/opencv2/${name}/device/detail/*.hpp" "include/opencv2/${name}/device/detail/*.h")
+file(GLOB lib_int_hdrs           "src/*.hpp"      "src/*.h")
 file(GLOB lib_cuda_hdrs          "src/cuda/*.hpp" "src/cuda/*.h")
-file(GLOB lib_device_hdrs        "src/opencv2/gpu/device/*.hpp" "src/opencv2/gpu/device/*.h")
-file(GLOB lib_device_hdrs_detail "src/opencv2/gpu/device/detail/*.hpp" "src/opencv2/gpu/device/detail/*.h")
 file(GLOB lib_srcs               "src/*.cpp")
 file(GLOB lib_cuda               "src/cuda/*.cu*")
 
@@ -74,8 +74,8 @@ else()
 endif()
 
 ocv_set_module_sources(
-  HEADERS ${lib_hdrs}
-  SOURCES ${lib_int_hdrs} ${lib_cuda_hdrs} ${lib_device_hdrs} ${lib_device_hdrs_detail} ${lib_srcs} ${lib_cuda} ${ncv_files} ${cuda_objs}
+  HEADERS ${lib_hdrs} ${lib_device_hdrs} ${lib_device_hdrs_detail}
+  SOURCES ${lib_int_hdrs} ${lib_cuda_hdrs} ${lib_srcs} ${lib_cuda} ${ncv_files} ${cuda_objs}
   )
 
 ocv_create_module(${cuda_link_libs})

modules/gpu/src/opencv2/gpu/device/detail/color_detail.hpp → modules/gpu/include/opencv2/gpu/device/detail/color_detail.hpp

@@ -1535,6 +1535,8 @@ namespace cv { namespace gpu { namespace device
             return functor_type(); \
         } \
     };
+
+    #undef CV_DESCALE
 }}} // namespace cv { namespace gpu { namespace device
 
 #endif // __OPENCV_GPU_COLOR_DETAIL_HPP__

modules/gpu/src/opencv2/gpu/device/detail/transform_detail.hpp → modules/gpu/include/opencv2/gpu/device/detail/transform_detail.hpp

@@ -47,7 +47,7 @@
 #include "../vec_traits.hpp"
 #include "../functional.hpp"
 
-namespace cv { namespace gpu { namespace device 
+namespace cv { namespace gpu { namespace device
 {
     namespace transform_detail
     {
@@ -203,7 +203,7 @@ namespace cv { namespace gpu { namespace device
         };
 
         template <typename T, typename D, typename UnOp, typename Mask>
-        __global__ static void transformSmart(const PtrStepSz<T> src_, PtrStep<D> dst_, const Mask mask, const UnOp op)
+        static __global__ void transformSmart(const PtrStepSz<T> src_, PtrStep<D> dst_, const Mask mask, const UnOp op)
         {
             typedef TransformFunctorTraits<UnOp> ft;
             typedef typename UnaryReadWriteTraits<T, D, ft::smart_shift>::read_type read_type;
@@ -239,10 +239,10 @@ namespace cv { namespace gpu { namespace device
         }
 
         template <typename T, typename D, typename UnOp, typename Mask>
-        static __global__ void transformSimple(const PtrStepSz<T> src, PtrStep<D> dst, const Mask mask, const UnOp op)
+        __global__ static void transformSimple(const PtrStepSz<T> src, PtrStep<D> dst, const Mask mask, const UnOp op)
         {
-	        const int x = blockDim.x * blockIdx.x + threadIdx.x;
-	        const int y = blockDim.y * blockIdx.y + threadIdx.y;
+            const int x = blockDim.x * blockIdx.x + threadIdx.x;
+            const int y = blockDim.y * blockIdx.y + threadIdx.y;
 
             if (x < src.cols && y < src.rows && mask(y, x))
             {
@@ -251,7 +251,7 @@ namespace cv { namespace gpu { namespace device
         }
 
         template <typename T1, typename T2, typename D, typename BinOp, typename Mask>
-        __global__ static void transformSmart(const PtrStepSz<T1> src1_, const PtrStep<T2> src2_, PtrStep<D> dst_, 
+        static __global__ void transformSmart(const PtrStepSz<T1> src1_, const PtrStep<T2> src2_, PtrStep<D> dst_,
             const Mask mask, const BinOp op)
         {
             typedef TransformFunctorTraits<BinOp> ft;
@@ -274,7 +274,7 @@ namespace cv { namespace gpu { namespace device
                     const read_type1 src1_n_el = ((const read_type1*)src1)[x];
                     const read_type2 src2_n_el = ((const read_type2*)src2)[x];
                     write_type dst_n_el = ((const write_type*)dst)[x];
-                    
+
                     OpUnroller<ft::smart_shift>::unroll(src1_n_el, src2_n_el, dst_n_el, mask, op, x_shifted, y);
 
                     ((write_type*)dst)[x] = dst_n_el;
@@ -291,11 +291,11 @@ namespace cv { namespace gpu { namespace device
         }
 
         template <typename T1, typename T2, typename D, typename BinOp, typename Mask>
-        static __global__ void transformSimple(const PtrStepSz<T1> src1, const PtrStep<T2> src2, PtrStep<D> dst, 
+        static __global__ void transformSimple(const PtrStepSz<T1> src1, const PtrStep<T2> src2, PtrStep<D> dst,
             const Mask mask, const BinOp op)
         {
-	        const int x = blockDim.x * blockIdx.x + threadIdx.x;
-	        const int y = blockDim.y * blockIdx.y + threadIdx.y;
+            const int x = blockDim.x * blockIdx.x + threadIdx.x;
+            const int y = blockDim.y * blockIdx.y + threadIdx.y;
 
             if (x < src1.cols && y < src1.rows && mask(y, x))
             {
@@ -314,13 +314,13 @@ namespace cv { namespace gpu { namespace device
                 typedef TransformFunctorTraits<UnOp> ft;
 
                 const dim3 threads(ft::simple_block_dim_x, ft::simple_block_dim_y, 1);
-                const dim3 grid(divUp(src.cols, threads.x), divUp(src.rows, threads.y), 1);     
+                const dim3 grid(divUp(src.cols, threads.x), divUp(src.rows, threads.y), 1);
 
                 transformSimple<T, D><<<grid, threads, 0, stream>>>(src, dst, mask, op);
                 cudaSafeCall( cudaGetLastError() );
 
                 if (stream == 0)
-                    cudaSafeCall( cudaDeviceSynchronize() ); 
+                    cudaSafeCall( cudaDeviceSynchronize() );
             }
 
             template <typename T1, typename T2, typename D, typename BinOp, typename Mask>
@@ -329,13 +329,13 @@ namespace cv { namespace gpu { namespace device
                 typedef TransformFunctorTraits<BinOp> ft;
 
                 const dim3 threads(ft::simple_block_dim_x, ft::simple_block_dim_y, 1);
-                const dim3 grid(divUp(src1.cols, threads.x), divUp(src1.rows, threads.y), 1);     
+                const dim3 grid(divUp(src1.cols, threads.x), divUp(src1.rows, threads.y), 1);
 
                 transformSimple<T1, T2, D><<<grid, threads, 0, stream>>>(src1, src2, dst, mask, op);
                 cudaSafeCall( cudaGetLastError() );
 
                 if (stream == 0)
-                    cudaSafeCall( cudaDeviceSynchronize() );            
+                    cudaSafeCall( cudaDeviceSynchronize() );
             }
         };
         template<> struct TransformDispatcher<true>
@@ -347,7 +347,7 @@ namespace cv { namespace gpu { namespace device
 
                 StaticAssert<ft::smart_shift != 1>::check();
 
-                if (!isAligned(src.data, ft::smart_shift * sizeof(T)) || !isAligned(src.step, ft::smart_shift * sizeof(T)) || 
+                if (!isAligned(src.data, ft::smart_shift * sizeof(T)) || !isAligned(src.step, ft::smart_shift * sizeof(T)) ||
                     !isAligned(dst.data, ft::smart_shift * sizeof(D)) || !isAligned(dst.step, ft::smart_shift * sizeof(D)))
                 {
                     TransformDispatcher<false>::call(src, dst, op, mask, stream);
@@ -355,7 +355,7 @@ namespace cv { namespace gpu { namespace device
                 }
 
                 const dim3 threads(ft::smart_block_dim_x, ft::smart_block_dim_y, 1);
-                const dim3 grid(divUp(src.cols, threads.x * ft::smart_shift), divUp(src.rows, threads.y), 1);      
+                const dim3 grid(divUp(src.cols, threads.x * ft::smart_shift), divUp(src.rows, threads.y), 1);
 
                 transformSmart<T, D><<<grid, threads, 0, stream>>>(src, dst, mask, op);
                 cudaSafeCall( cudaGetLastError() );
@@ -380,15 +380,15 @@ namespace cv { namespace gpu { namespace device
                 }
 
                 const dim3 threads(ft::smart_block_dim_x, ft::smart_block_dim_y, 1);
-                const dim3 grid(divUp(src1.cols, threads.x * ft::smart_shift), divUp(src1.rows, threads.y), 1);    
+                const dim3 grid(divUp(src1.cols, threads.x * ft::smart_shift), divUp(src1.rows, threads.y), 1);
 
                 transformSmart<T1, T2, D><<<grid, threads, 0, stream>>>(src1, src2, dst, mask, op);
                 cudaSafeCall( cudaGetLastError() );
 
                 if (stream == 0)
-                    cudaSafeCall( cudaDeviceSynchronize() );            
+                    cudaSafeCall( cudaDeviceSynchronize() );
             }
-        };        
+        };
     } // namespace transform_detail
 }}} // namespace cv { namespace gpu { namespace device
 

modules/gpu/src/opencv2/gpu/device/functional.hpp → modules/gpu/include/opencv2/gpu/device/functional.hpp

@@ -43,7 +43,7 @@
 #ifndef __OPENCV_GPU_FUNCTIONAL_HPP__
 #define __OPENCV_GPU_FUNCTIONAL_HPP__
 
-#include <thrust/functional.h>
+#include <functional>
 #include "saturate_cast.hpp"
 #include "vec_traits.hpp"
 #include "type_traits.hpp"
@@ -52,9 +52,8 @@
 namespace cv { namespace gpu { namespace device
 {
     // Function Objects
-
-    using thrust::unary_function;
-    using thrust::binary_function;
+    template<typename Argument, typename Result> struct unary_function : public std::unary_function<Argument, Result> {};
+    template<typename Argument1, typename Argument2, typename Result> struct binary_function : public std::binary_function<Argument1, Argument2, Result> {};
 
     // Arithmetic Operations
     template <typename T> struct plus : binary_function<T, T, T>

modules/gpu/include/opencv2/gpu/device/scan.hpp

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#ifndef __OPENCV_GPU_SCAN_HPP__
+#define __OPENCV_GPU_SCAN_HPP__
+
+#include "common.hpp"
+
+namespace cv { namespace gpu { namespace device
+{
+    enum ScanKind { EXCLUSIVE = 0,  INCLUSIVE = 1 };
+
+    template <ScanKind Kind, typename T, typename F> struct WarpScan
+    {
+        __device__ __forceinline__ WarpScan() {}
+        __device__ __forceinline__ WarpScan(const WarpScan& other) { (void)other; }
+
+        __device__ __forceinline__ T operator()( volatile T *ptr , const unsigned int idx)
+        {
+            const unsigned int lane = idx & 31;
+            F op;
+
+            if ( lane >=  1) ptr [idx ] = op(ptr [idx -  1], ptr [idx]);
+            if ( lane >=  2) ptr [idx ] = op(ptr [idx -  2], ptr [idx]);
+            if ( lane >=  4) ptr [idx ] = op(ptr [idx -  4], ptr [idx]);
+            if ( lane >=  8) ptr [idx ] = op(ptr [idx -  8], ptr [idx]);
+            if ( lane >= 16) ptr [idx ] = op(ptr [idx - 16], ptr [idx]);
+
+            if( Kind == INCLUSIVE )
+                return ptr [idx];
+            else
+                return (lane > 0) ? ptr [idx - 1] : 0;
+        }
+
+        __device__ __forceinline__ unsigned int index(const unsigned int tid)
+        {
+            return tid;
+        }
+
+        __device__ __forceinline__ void init(volatile T *ptr){}
+
+        static const int warp_offset      = 0;
+
+        typedef WarpScan<INCLUSIVE, T, F>  merge;
+    };
+
+    template <ScanKind Kind , typename T, typename F> struct WarpScanNoComp
+    {
+        __device__ __forceinline__ WarpScanNoComp() {}
+        __device__ __forceinline__ WarpScanNoComp(const WarpScanNoComp& other) { (void)other; }
+
+        __device__ __forceinline__ T operator()( volatile T *ptr , const unsigned int idx)
+        {
+            const unsigned int lane = threadIdx.x & 31;
+            F op;
+
+            ptr [idx ] = op(ptr [idx -  1], ptr [idx]);
+            ptr [idx ] = op(ptr [idx -  2], ptr [idx]);
+            ptr [idx ] = op(ptr [idx -  4], ptr [idx]);
+            ptr [idx ] = op(ptr [idx -  8], ptr [idx]);
+            ptr [idx ] = op(ptr [idx - 16], ptr [idx]);
+
+            if( Kind == INCLUSIVE )
+                return ptr [idx];
+            else
+                return (lane > 0) ? ptr [idx - 1] : 0;
+        }
+
+        __device__ __forceinline__ unsigned int index(const unsigned int tid)
+        {
+            return (tid >> warp_log) * warp_smem_stride + 16 + (tid & warp_mask);
+        }
+
+        __device__ __forceinline__ void init(volatile T *ptr)
+        {
+            ptr[threadIdx.x] = 0;
+        }
+
+        static const int warp_smem_stride = 32 + 16 + 1;
+        static const int warp_offset      = 16;
+        static const int warp_log         = 5;
+        static const int warp_mask        = 31;
+
+        typedef WarpScanNoComp<INCLUSIVE, T, F> merge;
+    };
+
+    template <ScanKind Kind , typename T, typename Sc, typename F> struct BlockScan
+    {
+        __device__ __forceinline__ BlockScan() {}
+        __device__ __forceinline__ BlockScan(const BlockScan& other) { (void)other; }
+
+        __device__ __forceinline__ T operator()(volatile T *ptr)
+        {
+            const unsigned int tid  = threadIdx.x;
+            const unsigned int lane = tid & warp_mask;
+            const unsigned int warp = tid >> warp_log;
+
+            Sc scan;
+            typename Sc::merge merge_scan;
+            const unsigned int idx = scan.index(tid);
+
+            T val = scan(ptr, idx);
+            __syncthreads ();
+
+            if( warp == 0)
+                scan.init(ptr);
+            __syncthreads ();
+
+            if( lane == 31 )
+                ptr [scan.warp_offset + warp ] = (Kind == INCLUSIVE) ? val : ptr [idx];
+            __syncthreads ();
+
+            if( warp == 0 )
+                merge_scan(ptr, idx);
+            __syncthreads();
+
+            if ( warp > 0)
+                val = ptr [scan.warp_offset + warp - 1] + val;
+            __syncthreads ();
+
+            ptr[idx] = val;
+            __syncthreads ();
+
+            return val ;
+        }
+
+        static const int warp_log  = 5;
+        static const int warp_mask = 31;
+    };
+}}}
+
+#endif // __OPENCV_GPU_SCAN_HPP__

modules/gpu/src/opencv2/gpu/device/static_check.hpp → modules/gpu/include/opencv2/gpu/device/static_check.hpp

@@ -60,10 +60,8 @@ namespace cv { namespace gpu
             __OPENCV_GPU_HOST_DEVICE__ static void check() {};
         };
     }
-
-    using ::cv::gpu::device::Static;
 }}
 
 #undef __OPENCV_GPU_HOST_DEVICE__
 
-#endif /* __OPENCV_GPU_GPU_DEVICE_STATIC_CHECK_HPP__ */
+#endif /* __OPENCV_GPU_GPU_DEVICE_STATIC_CHECK_HPP__ */
\ No newline at end of file

modules/gpu/src/opencv2/gpu/device/utility.hpp → modules/gpu/include/opencv2/gpu/device/utility.hpp

@@ -45,7 +45,7 @@
 
 #include "saturate_cast.hpp"
 #include "datamov_utils.hpp"
-#include "detail/utility_detail.hpp"
+#include "detail/reduction_detail.hpp"
 
 namespace cv { namespace gpu { namespace device
 {

modules/gpu/src/cuda/lbp.cu

@@ -42,9 +42,9 @@
 
 #if !defined CUDA_DISABLER
 
-#include <opencv2/gpu/device/lbp.hpp>
-#include <opencv2/gpu/device/vec_traits.hpp>
-#include <opencv2/gpu/device/saturate_cast.hpp>
+#include "lbp.hpp"
+#include "opencv2/gpu/device/vec_traits.hpp"
+#include "opencv2/gpu/device/saturate_cast.hpp"
 
 namespace cv { namespace gpu { namespace device
 {
@@ -299,4 +299,4 @@ namespace cv { namespace gpu { namespace device
     }
 }}}
 
-#endif /* CUDA_DISABLER */
+#endif /* CUDA_DISABLER */
\ No newline at end of file

modules/gpu/src/opencv2/gpu/device/scan.hpp

-/*M///////////////////////////////////////////////////////////////////////////////////////
-//
-//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
-//
-//  By downloading, copying, installing or using the software you agree to this license.
-//  If you do not agree to this license, do not download, install,
-//  copy or use the software.
-//
-//
-//                           License Agreement
-//                For Open Source Computer Vision Library
-//
-// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
-// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
-// Third party copyrights are property of their respective owners.
-//
-// Redistribution and use in source and binary forms, with or without modification,
-// are permitted provided that the following conditions are met:
-//
-//   * Redistribution's of source code must retain the above copyright notice,
-//     this list of conditions and the following disclaimer.
-//
-//   * Redistribution's in binary form must reproduce the above copyright notice,
-//     this list of conditions and the following disclaimer in the documentation
-//     and/or other materials provided with the distribution.
-//
-//   * The name of the copyright holders may not be used to endorse or promote products
-//     derived from this software without specific prior written permission.
-//
-// This software is provided by the copyright holders and contributors "as is" and
-// any express or implied warranties, including, but not limited to, the implied
-// warranties of merchantability and fitness for a particular purpose are disclaimed.
-// In no event shall the Intel Corporation or contributors be liable for any direct,
-// indirect, incidental, special, exemplary, or consequential damages
-// (including, but not limited to, procurement of substitute goods or services;
-// loss of use, data, or profits; or business interruption) however caused
-// and on any theory of liability, whether in contract, strict liability,
-// or tort (including negligence or otherwise) arising in any way out of
-// the use of this software, even if advised of the possibility of such damage.
-//
-//M*/
-
-#ifndef __OPENCV_GPU_SCAN_HPP__
-#define __OPENCV_GPU_SCAN_HPP__
-
-        enum ScanKind { EXCLUSIVE = 0,  INCLUSIVE = 1 };
-
-        template <ScanKind Kind, typename T, typename F> struct WarpScan
-        {
-            __device__ __forceinline__ WarpScan() {}
-            __device__ __forceinline__ WarpScan(const WarpScan& other) { (void)other; }
-
-            __device__ __forceinline__ T operator()( volatile T *ptr , const unsigned int idx)
-            {
-                const unsigned int lane = idx & 31;
-                F op;
-
-                if ( lane >=  1) ptr [idx ] = op(ptr [idx -  1], ptr [idx]);
-                if ( lane >=  2) ptr [idx ] = op(ptr [idx -  2], ptr [idx]);
-                if ( lane >=  4) ptr [idx ] = op(ptr [idx -  4], ptr [idx]);
-                if ( lane >=  8) ptr [idx ] = op(ptr [idx -  8], ptr [idx]);
-                if ( lane >= 16) ptr [idx ] = op(ptr [idx - 16], ptr [idx]);
-
-                if( Kind == INCLUSIVE )
-                    return ptr [idx];
-                else
-                    return (lane > 0) ? ptr [idx - 1] : 0;
-            }
-
-            __device__ __forceinline__ unsigned int index(const unsigned int tid)
-            {
-                return tid;
-            }
-
-            __device__ __forceinline__ void init(volatile T *ptr){}
-
-            static const int warp_offset      = 0;
-
-            typedef WarpScan<INCLUSIVE, T, F>  merge;
-        };
-
-        template <ScanKind Kind , typename T, typename F> struct WarpScanNoComp
-        {
-            __device__ __forceinline__ WarpScanNoComp() {}
-            __device__ __forceinline__ WarpScanNoComp(const WarpScanNoComp& other) { (void)other; }
-
-            __device__ __forceinline__ T operator()( volatile T *ptr , const unsigned int idx)
-            {
-                const unsigned int lane = threadIdx.x & 31;
-                F op;
-
-                ptr [idx ] = op(ptr [idx -  1], ptr [idx]);
-                ptr [idx ] = op(ptr [idx -  2], ptr [idx]);
-                ptr [idx ] = op(ptr [idx -  4], ptr [idx]);
-                ptr [idx ] = op(ptr [idx -  8], ptr [idx]);
-                ptr [idx ] = op(ptr [idx - 16], ptr [idx]);
-
-                if( Kind == INCLUSIVE )
-                    return ptr [idx];
-                else
-                    return (lane > 0) ? ptr [idx - 1] : 0;
-            }
-
-            __device__ __forceinline__ unsigned int index(const unsigned int tid)
-            {
-                return (tid >> warp_log) * warp_smem_stride + 16 + (tid & warp_mask);
-            }
-
-            __device__ __forceinline__ void init(volatile T *ptr)
-            {
-                ptr[threadIdx.x] = 0;
-            }
-
-            static const int warp_smem_stride = 32 + 16 + 1;
-            static const int warp_offset      = 16;
-            static const int warp_log         = 5;
-            static const int warp_mask        = 31;
-
-            typedef WarpScanNoComp<INCLUSIVE, T, F> merge;
-        };
-
-        template <ScanKind Kind , typename T, typename Sc, typename F> struct BlockScan
-        {
-            __device__ __forceinline__ BlockScan() {}
-            __device__ __forceinline__ BlockScan(const BlockScan& other) { (void)other; }
-
-            __device__ __forceinline__ T operator()(volatile T *ptr)
-            {
-                const unsigned int tid  = threadIdx.x;
-                const unsigned int lane = tid & warp_mask;
-                const unsigned int warp = tid >> warp_log;
-
-                Sc scan;
-                typename Sc::merge merge_scan;
-                const unsigned int idx = scan.index(tid);
-
-                T val = scan(ptr, idx);
-                __syncthreads ();
-
-                if( warp == 0)
-                    scan.init(ptr);
-                __syncthreads ();
-
-                if( lane == 31 )
-                    ptr [scan.warp_offset + warp ] = (Kind == INCLUSIVE) ? val : ptr [idx];
-                __syncthreads ();
-
-                if( warp == 0 )
-                    merge_scan(ptr, idx);
-                __syncthreads();
-
-                if ( warp > 0)
-                    val = ptr [scan.warp_offset + warp - 1] + val;
-                __syncthreads ();
-
-                ptr[idx] = val;
-                __syncthreads ();
-
-                return val ;
-            }
-
-            static const int warp_log  = 5;
-            static const int warp_mask = 31;
-        };
-
-#endif
\ No newline at end of file