 // Min max
 
     enum { MIN, MAX };  
+
+    template <typename T> struct MinMaxTypeTraits {};
+    template <> struct MinMaxTypeTraits<unsigned char> { typedef int best_type; };
+    template <> struct MinMaxTypeTraits<signed char> { typedef int best_type; };
+    template <> struct MinMaxTypeTraits<unsigned short> { typedef int best_type; };
+    template <> struct MinMaxTypeTraits<signed short> { typedef int best_type; };
+    template <> struct MinMaxTypeTraits<int> { typedef int best_type; };
+    template <> struct MinMaxTypeTraits<float> { typedef float best_type; };
+    template <> struct MinMaxTypeTraits<double> { typedef double best_type; };
+
     template <typename T, int op> struct Cmp {};
     
     template <typename T>


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

   

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

        opt_kernel<256, MIN, T><<<grid, threads>>>(cursize.x, cursize.y, src, minval_buf[curbuf]);
        opt_kernel<256, 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, MIN, T><<<grid, threads>>>(cursize.x, cursize.y, minval_buf[curbuf], minval_buf[1 - curbuf]);
            opt_kernel<256, MAX, T><<<grid, threads>>>(cursize.x, cursize.y, maxval_buf[curbuf], maxval_buf[1 - curbuf]);
            curbuf = 1 - curbuf;
            cursize = grid;
        }
