 	printf("\t--stream\tWhether or not threads also stream 32-bit words \n");
 	printf("\t\t\tthrough global memory concurrently as well (to simulate)\n");
 	printf("\n");
+	printf("\t--local\tWhether or not to perform atomics on local smem \n");
+	printf("\t\t\tcounters versus a global gmem counter\n");
+	printf("\n");
 }
 
 


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


		timer.Start();
		for (int i = 0; i < g_iterations; i++) {

			if (g_local) {
				LocalAtomicKernel<<<grid_size, CTA_THREADS>>>(
					d_in,
					d_out,
					stride,
					tiles_per_cta,
					extra_tiles);
			} else {
				GlobalAtomicKernel<<<grid_size, CTA_THREADS>>>(
					d_counter,
					d_in,
					d_out,
					stride,
					tiles_per_cta,
					extra_tiles);
			}
		}
		timer.Stop();

