         {
             maxval[tid] = val;
             maxloc[tid] = maxloc[tid + offset];
+        }
+    }
+
+
+    template <int size, typename T>
+    __device__ void find_min_max_loc_in_smem(volatile T* minval, volatile T* maxval, volatile unsigned int* minloc, 
+                                             volatile unsigned int* maxloc, const unsigned int tid)
+    {
+        if (size >= 512) { if (tid < 256) { merge(tid, 256, minval, maxval, minloc, maxloc); } __syncthreads(); }
+        if (size >= 256) { if (tid < 128) { merge(tid, 128, minval, maxval, minloc, maxloc); }  __syncthreads(); }
+        if (size >= 128) { if (tid < 64) { merge(tid, 64, minval, maxval, minloc, maxloc); } __syncthreads(); }
+
+        if (tid < 32)
+        {
+            if (size >= 64) merge(tid, 32, minval, maxval, minloc, maxloc);
+            if (size >= 32) merge(tid, 16, minval, maxval, minloc, maxloc);
+            if (size >= 16) merge(tid, 8, minval, maxval, minloc, maxloc);
+            if (size >= 8) merge(tid, 4, minval, maxval, minloc, maxloc);
+            if (size >= 4) merge(tid, 2, minval, maxval, minloc, maxloc);
+            if (size >= 2) merge(tid, 1, minval, maxval, minloc, maxloc);
         }
     }
 



    // This kernel will be used only when compute capability is 1.0
    template <typename T>
    __global__ void min_max_loc_kernel_2ndstep(T* minval, T* maxval, unsigned int* minloc, unsigned int* maxloc, int size)
    {
        T val;
        T mymin = minval[0];
        T mymax = maxval[0];
        unsigned int imin  = 0, imax = 0;
        for (unsigned int i = 1; i < size; ++i)
        {     
            val = minval[i]; if (val < mymin) { mymin = val; imin = i; }
            val = maxval[i]; if (val > mymax) { mymax = val; imax = i; }
        }
        minval[0] = mymin;
        maxval[0] = mymax;
        minloc[0] = minloc[imin];
        maxloc[0] = maxloc[imax];
    }



        unsigned int tid = threadIdx.y * blockDim.x + threadIdx.x;

        scount[tid] = tid < size ? count[tid] : 0;
		sum_is_smem<nthreads, unsigned int>(scount, tid);

		if (tid == 0) count[0] = scount[0];
    }
