 
         __syncthreads();
 
-        if (nthreads >= 512) if (tid < 256) { Cmp::call(tid, 256, soptval); __syncthreads(); }
-        if (nthreads >= 256) if (tid < 128) { Cmp::call(tid, 128, soptval); __syncthreads(); }
-        if (nthreads >= 128) if (tid < 64) { Cmp::call(tid, 64, soptval); __syncthreads(); }
+        if (nthreads >= 512) if (tid < 256) { Cmp<best_type, op>::call(tid, 256, soptval); __syncthreads(); }
+        if (nthreads >= 256) if (tid < 128) { Cmp<best_type, op>::call(tid, 128, soptval); __syncthreads(); }
+        if (nthreads >= 128) if (tid < 64) { Cmp<best_type, op>::call(tid, 64, soptval); __syncthreads(); }
 
         if (tid < 32)
         {
-            if (nthreads >= 64) Cmp::call(tid, 32, soptval);
-            if (nthreads >= 32) Cmp::call(tid, 16, soptval);
-            if (nthreads >= 16) Cmp::call(tid, 8, soptval);
-            if (nthreads >= 8) Cmp::call(tid, 4, soptval);
-            if (nthreads >= 4) Cmp::call(tid, 2, soptval);
-            if (nthreads >= 2) Cmp::call(tid, 1, soptval);
-        }
-
-        if (tid == 0) ((T*)optval.ptr(blockIdx.y))[blockIdx.x] = soptval[0];
+            if (nthreads >= 64) Cmp<best_type, op>::call(tid, 32, soptval);
+            if (nthreads >= 32) Cmp<best_type, op>::call(tid, 16, soptval);
+            if (nthreads >= 16) Cmp<best_type, op>::call(tid, 8, soptval);
+            if (nthreads >= 8) Cmp<best_type, op>::call(tid, 4, soptval);
+            if (nthreads >= 4) Cmp<best_type, op>::call(tid, 2, soptval);
+            if (nthreads >= 2) Cmp<best_type, op>::call(tid, 1, soptval);
+        }
+
+        if (tid == 0) ((T*)optval.ptr(blockIdx.y))[blockIdx.x] = (T)soptval[0];
     }
 
    

    {
        static __device__ void call(unsigned int tid, unsigned int offset, volatile T* optval)
        {
            T val = optval[tid + offset];
            if (val < optval[tid]) optval[tid] = val;
            //optval[tid] = min(optval[tid], optval[tid + offset]); 
        }
    };


    {
        static __device__ void call(unsigned int tid, unsigned int offset, volatile T* optval)
        {
            optval[tid] = max(optval[tid], optval[tid + offset]);
        }
    };


    template <int nthreads, int op, typename T>
    __global__ void opt_kernel(int cols, int rows, const PtrStep src, PtrStep optval)
    {
        typedef typename MinMaxTypeTraits<T>::best_type best_type;
        __shared__ best_type soptval[nthreads];

        unsigned int x0 = blockIdx.x * blockDim.x;
        unsigned int y0 = blockIdx.y * blockDim.y;
