 
         __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];
     }
 
    

        dim3 cursize(src.cols, src.rows);
        dim3 grid(divUp(cursize.x, threads.x), divUp(cursize.y, threads.y));

        opt_kernel<256, Cmp<T, MIN>, T><<<grid, threads>>>(cursize.x, cursize.y, src, minval_buf[curbuf]);
        opt_kernel<256, Cmp<T, MAX>, T><<<grid, threads>>>(cursize.x, cursize.y, src, maxval_buf[curbuf]);
        cursize = grid;

        while (cursize.x > 1 || cursize.y > 1)
        {
            grid.x = divUp(cursize.x, threads.x); 
            grid.y = divUp(cursize.y, threads.y);  
            opt_kernel<256, Cmp<T, MIN>, T><<<grid, threads>>>(cursize.x, cursize.y, minval_buf[curbuf], minval_buf[1 - curbuf]);
            opt_kernel<256, Cmp<T, MAX>, T><<<grid, threads>>>(cursize.x, cursize.y, maxval_buf[curbuf], maxval_buf[1 - curbuf]);
            curbuf = 1 - curbuf;
            cursize = grid;
        }

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

