 
 
     // This kernel will be used only when compute capability is 1.0
-    template <typename T>
+    template <int nthreads, 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];
+        typedef typename MinMaxTypeTraits<T>::best_type best_type;
+        __shared__ best_type sminval[nthreads];
+        __shared__ best_type smaxval[nthreads];
+        __shared__ unsigned int sminloc[nthreads];
+        __shared__ unsigned int smaxloc[nthreads];
+
+        unsigned int tid = threadIdx.y * blockDim.x + threadIdx.x;
+        unsigned int idx = min(tid, gridDim.x * gridDim.y - 1);
+
+        sminval[tid] = minval[idx];
+        smaxval[tid] = maxval[idx];
+        sminloc[tid] = minloc[idx];
+        smaxloc[tid] = maxloc[idx];
+        __syncthreads();
+
+		find_min_max_loc_in_smem<nthreads, best_type>(sminval, smaxval, sminloc, smaxloc, tid);
+
+        if (tid == 0) 
+        {
+            minval[0] = (T)sminval[0];
+            maxval[0] = (T)smaxval[0];
+            minloc[0] = sminloc[0];
+            maxloc[0] = smaxloc[0];
+        }
     }
 
 


        cudaSafeCall(cudaMemcpyToSymbol(blocks_finished, &czero, sizeof(blocks_finished)));
        min_max_kernel<256, T><<<grid, threads>>>(src, minval_buf, maxval_buf);
        min_max_kernel_2ndstep<T><<<1, 1>>>(minval_buf, maxval_buf, grid.x * grid.y);
        cudaSafeCall(cudaThreadSynchronize());

        T minval_, maxval_;

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