         *minval_buf = numeric_limits<int>::max();
         *maxval_buf = numeric_limits<int>::min();
     }
+    __global__ void setDefaultKernel(float* minval_buf, float* maxval_buf)
+    {
+        *minval_buf = numeric_limits<float>::max();
+        *maxval_buf = -numeric_limits<float>::max();
+    }
+    __global__ void setDefaultKernel(double* minval_buf, double* maxval_buf)
+    {
+        *minval_buf = numeric_limits<double>::max();
+        *maxval_buf = -numeric_limits<double>::max();
+    }
 
     template <typename R>
-    void setDefault(R*, R*)
-    {
-    }
-    void setDefault(int* minval_buf, int* maxval_buf)
+    void setDefault(R* minval_buf, R* maxval_buf)
     {
         setDefaultKernel<<<1, 1>>>(minval_buf, maxval_buf);
     }

            maxval[bid] = (T) mymax;
            minloc[bid] = myminloc;
            maxloc[bid] = mymaxloc;

            __threadfence();

            unsigned int ticket = ::atomicInc(&blocks_finished, gridDim.x * gridDim.y);
            is_last = (ticket == gridDim.x * gridDim.y - 1);
        }

        __syncthreads();

        if (is_last)
        {
            unsigned int idx = ::min(tid, gridDim.x * gridDim.y - 1);

            mymin = minval[idx];
            mymax = maxval[idx];
            myminloc = minloc[idx];
            mymaxloc = maxloc[idx];

            reduceKeyVal<BLOCK_SIZE>(smem_tuple(sminval, smaxval), thrust::tie(mymin, mymax),
                                     smem_tuple(sminloc, smaxloc), thrust::tie(myminloc, mymaxloc),
                                     tid,
                                     thrust::make_tuple(less<work_type>(), greater<work_type>()));

            if (tid == 0)
            {
                minval[0] = (T) mymin;
                maxval[0] = (T) mymax;
                minloc[0] = myminloc;
                maxloc[0] = mymaxloc;

                blocks_finished = 0;
            }
        }
    }


        maxloc[1] = maxloc_ / src.cols; maxloc[0] = maxloc_ - maxloc[1] * src.cols;
    }

    template void run<unsigned char >(const PtrStepSzb src, const PtrStepb mask, double* minval, double* maxval, int* minloc, int* maxloc, PtrStepb valbuf, PtrStep<unsigned int> locbuf);
    template void run<signed char >(const PtrStepSzb src, const PtrStepb mask, double* minval, double* maxval, int* minloc, int* maxloc, PtrStepb valbuf, PtrStep<unsigned int> locbuf);
    template void run<unsigned short>(const PtrStepSzb src, const PtrStepb mask, double* minval, double* maxval, int* minloc, int* maxloc, PtrStepb valbuf, PtrStep<unsigned int> locbuf);
    template void run<short >(const PtrStepSzb src, const PtrStepb mask, double* minval, double* maxval, int* minloc, int* maxloc, PtrStepb valbuf, PtrStep<unsigned int> locbuf);
    template void run<int   >(const PtrStepSzb src, const PtrStepb mask, double* minval, double* maxval, int* minloc, int* maxloc, PtrStepb valbuf, PtrStep<unsigned int> locbuf);
    template void run<float >(const PtrStepSzb src, const PtrStepb mask, double* minval, double* maxval, int* minloc, int* maxloc, PtrStepb valbuf, PtrStep<unsigned int> locbuf);
