         s_dy_bin[tid] = vy;
     }
 
-    __device__ void reduce_sum25(volatile float* sdata1, volatile float* sdata2,
-                                 volatile float* sdata3, volatile float* sdata4, int tid)
+    __device__ void reduce_sum25(volatile float* sdata1, volatile float* sdata2, volatile float* sdata3, volatile float* sdata4, int tid)
     {
         // first step is to reduce from 25 to 16
         if (tid < 9) // use 9 threads

        }
    }

    // Spawn 16 blocks per interest point
    // - computes unnormalized 64 dimensional descriptor, puts it into d_descriptors in the correct location
    __global__ void compute_descriptors64(PtrStepf descriptors, const KeyPoint_GPU* features)
    {
        // 2 floats (dx,dy) for each thread (5x5 sample points in each sub-region)
        __shared__ float s_PATCH[6][6];
        __shared__ float sdx[25];
        __shared__ float sdy[25];
        __shared__ float sdxabs[25];


#include "internal_shared.hpp"
#include "opencv2/gpu/device/limits_gpu.hpp"
#include "opencv2/gpu/device/saturate_cast.hpp"

using namespace cv::gpu;
using namespace cv::gpu::device;
