 #define BLOCKS_NUM 1
 #define TOTAL_THREADS (THREADS_PER_BLOCK*BLOCKS_NUM)
 #define WARP_SIZE 32
-#define REPEAT_TIMES 1
+#define REPEAT_TIMES 1024
 
 // GPU error check
 #define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }

int main(){
	uint32_t *startClk = (uint32_t*) malloc(TOTAL_THREADS*sizeof(uint32_t));
	uint32_t *stopClk = (uint32_t*) malloc(TOTAL_THREADS*sizeof(uint32_t));
	int32_t *data1 = (int32_t*) malloc(TOTAL_THREADS*sizeof(int32_t));
	//int32_t *data2 = (int32_t*) malloc(TOTAL_THREADS*sizeof(int32_t));
	int32_t *res = (int32_t*) malloc(TOTAL_THREADS*sizeof(int32_t));


	//register T s1 = data1[gid];
	//register T s2 = data2[gid];
	//register T result = 0;
	uint32_t index = 0;
	int32_t offset = 10;
	// synchronize all threads
	asm volatile ("bar.sync 0;");
	
	// start timing
	uint32_t start = 0;
	asm volatile ("mov.u32 %0, %%clock;" : "=r"(start) :: "memory");
	//printf("%ld \n", &data1[0]);
	for (int j=0 ; j<REPEAT_TIMES ; ++j) {
		index = atomicAdd(&data1[index], offset);
		//printf("index = %d", index);
	}
	// synchronize all threads
	asm volatile("bar.sync 0;");
