         }
 
 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 110
-        
-        // Process partial results in the first thread of the last block      
-        if ((gridDim.x > 1 || gridDim.y > 1) && tid == 0)
-        {
-            __threadfence();
-            if (atomicInc(&blocks_finished, gridDim.x * gridDim.y) == gridDim.x * gridDim.y - 1)
+		__shared__ bool is_last;
+
+		if (tid == 0)
+		{
+			minval[blockIdx.y * gridDim.x + blockIdx.x] = (T)sminval[0];
+            maxval[blockIdx.y * gridDim.x + blockIdx.x] = (T)smaxval[0];
+			__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);
+
+            sminval[tid] = minval[idx];
+            smaxval[tid] = maxval[idx];
+            __syncthreads();
+
+			find_min_max_in_smem<nthreads, best_type>(sminval, smaxval, tid);
+
+            if (tid == 0) 
             {
-                mymin = minval[0];
-                mymax = maxval[0];
-                for (unsigned int i = 1; i < gridDim.x * gridDim.y; ++i)
-                {                    
-                    mymin = min(mymin, minval[i]);
-                    mymax = max(mymax, maxval[i]);
-                }
-                minval[0] = mymin;
-                maxval[0] = mymax;
+                minval[0] = (T)sminval[0];
+                maxval[0] = (T)smaxval[0];
             }
-        }
-
+		}
+#else
+        if (tid == 0) 
+        {
+            minval[blockIdx.y * gridDim.x + blockIdx.x] = (T)sminval[0];
+            maxval[blockIdx.y * gridDim.x + blockIdx.x] = (T)smaxval[0];
+        }
 #endif
     }
 


        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_;


        sminval[tid] = mymin;
        smaxval[tid] = mymax;
        __syncthreads();

        find_min_max_in_smem<nthreads, best_type>(sminval, smaxval, tid);

        if (tid == 0) 
        {
