 }
 
 
+/**
+ * Updates local counter
+ */
+__global__ void LocalAtomicKernel(
+	int *d_in,
+	int *d_out,
+	int stride,
+	int tiles_per_cta,
+	int extra_tiles)
+{
+	__shared__ int s_counter;
+
+	int data;
+	int tid = (blockIdx.x * blockDim.x * tiles_per_cta) + threadIdx.x;
+	int mask = stride - 1;
+
+	if (blockIdx.x < extra_tiles) {
+
+		// We get an extra tile
+		tiles_per_cta++;
+		tid += blockIdx.x * blockDim.x;
+
+	} else if (extra_tiles > 0) {
+
+		// We don't get an extra tile, but others did
+		tid += blockDim.x * extra_tiles;
+	}
+
+	// Iterate over tiles
+	for (int i = 0; i < tiles_per_cta; i++) {
+
+		if (d_in) {
+			data = d_in[tid];
+		}
+
+		if ((tid & mask) == mask) {
+			if (stride > 0) {
+
+#if __B40C_CUDA_ARCH__ >= 120
+				// access shared counter
+				atomicAdd(&s_counter, 1);
+#endif
+			}
+		}
+
+		if (d_out) {
+			d_out[tid] = data;
+		}
+
+		tid += blockDim.x;
+	}
+}
+
 
 /******************************************************************************
  * Main


    // Compute kernel params
    util::CudaProperties cuda_props;
    const int CTA_THREADS = 128;
    const int MAX_CTAS = 1024 * 16;

    if (g_max_ctas <= 0) {
        g_max_ctas = MAX_CTAS;
    }

	}

    g_stream = args.CheckCmdLineFlag("stream");
    g_local = args.CheckCmdLineFlag("local");
    args.GetCmdLineArgument("log-stride", g_log_stride);
    args.GetCmdLineArgument("i", g_iterations);
    args.GetCmdLineArgument("n", g_num_elements);
