 
 #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;

        }
    }

    // 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];

        const int xIndex = xBlock * blockDim.x + threadIdx.x;
        const int yIndex = yBlock * blockDim.y + threadIdx.y;

        s_PATCH[threadIdx.y][threadIdx.x] = calcPATCH(yIndex, xIndex, centerX, centerY, win_offset, cos_dir, sin_dir, win_size);
        if (threadIdx.x == 0)
            s_PATCH[threadIdx.y][5] = calcPATCH(yIndex, xBlock * blockDim.x + 5, centerX, centerY, win_offset, cos_dir, sin_dir, win_size);
        if (threadIdx.y == 0)
            s_PATCH[5][threadIdx.x] = calcPATCH(yBlock * blockDim.y + 5, xIndex, centerX, centerY, win_offset, cos_dir, sin_dir, win_size);
        if (threadIdx.x == 0 && threadIdx.y == 0)
            s_PATCH[5][5] = calcPATCH(yBlock * blockDim.y + 5, xBlock * blockDim.x + 5, centerX, centerY, win_offset, cos_dir, sin_dir, win_size);
        __syncthreads();

        const float dw = c_DW[yIndex * PATCH_SZ + xIndex];
