Fix units of time measurements

brian-team · May 26, 2016 · dee9bf7 · dee9bf7
1 parent a45d196
commit dee9bf7
Show file tree

Hide file tree

Showing 2 changed files with 23 additions and 13 deletions.
diff --git a/...es/issue9_spikespace/global_atomicAdds/code_objects/neurongroup_thresholder_codeobject.cu b/...es/issue9_spikespace/global_atomicAdds/code_objects/neurongroup_thresholder_codeobject.cu
@@ -90,15 +90,21 @@ __global__ void kernel_neurongroup_thresholder_codeobject(
  const int32_t i = _ptr_array_neurongroup_i[_idx];
  const double _cond = true;//(i / 2) == ((i + 1) / 2);
 
+	int32_t spike_index;
+
  	__syncthreads();
  	clock_t start = clock();
 	if(_cond) {
-		int32_t spike_index = atomicAdd(&_ptr_array_neurongroup__spikespace[10000], 1);
-		_ptr_array_neurongroup__spikespace[spike_index] = _idx;
+		spike_index = atomicAdd(&_ptr_array_neurongroup__spikespace[10000], 1);
 	}
  	clock_t end = clock();
 	glob_time[tid] = (int)(end - start);
 
+	__syncthreads();
+	if(_cond) {
+		_ptr_array_neurongroup__spikespace[spike_index] = _idx;
+	}
+
 }
 
 void _run_neurongroup_thresholder_codeobject()
@@ -137,10 +143,11 @@ kernel_neurongroup_thresholder_codeobject<<<num_blocks(10000), num_threads(10000
 		glob_sum += h_glob_time[i];
 	}
 
+	// clock cycles / clock rate(MHz) = time(us)
 	double glob_time = (double)(glob_sum)/980;
 
-	std::cout << "Average time per global atomicAdd:\n" << glob_time/num_threads(10000) << "ms\n\n" << std::endl;
-	std::cout << "Total time in global atomicAdd:\n" << glob_time << "ms\n\n" << std::endl;
+	std::cout << "Average time per global atomicAdd:\n" << glob_time/num_threads(10000) << "us\n\n" << std::endl;
+	std::cout << "Total time in global atomicAdd:\n" << glob_time << "us\n\n" << std::endl;
 }
 
 
diff --git a/...es/issue9_spikespace/shared_atomicAdds/code_objects/neurongroup_thresholder_codeobject.cu b/...es/issue9_spikespace/shared_atomicAdds/code_objects/neurongroup_thresholder_codeobject.cu
@@ -103,19 +103,22 @@ __global__ void kernel_neurongroup_thresholder_codeobject(
  const double _cond = true;//(i / 2) == ((i + 1) / 2);
 
 
-
+	int32_t spike_index;
 
  	__syncthreads();
  	clock_t block_start = clock();
 	if(_cond) {
-		int32_t spike_index = atomicAdd(&spike_counter_block, 1);
-		_ptr_array_neurongroup__spikespace[bid * THREADS_PER_BLOCK + spike_index] = _idx;
+		spike_index = atomicAdd(&spike_counter_block, 1);
 	}
  	clock_t block_end = clock();
 	block_time[tid] = (int)(block_end - block_start);
 
-	__syncthreads();
 
+	if(_cond) {
+		_ptr_array_neurongroup__spikespace[bid * THREADS_PER_BLOCK + spike_index] = _idx;
+	}
+
+	__syncthreads();
 
  	clock_t glob_start = clock();
 	if (tid == 0) {
@@ -176,11 +179,11 @@ kernel_neurongroup_thresholder_codeobject<<<num_blocks(10000), num_threads(10000
 	double block_time = (double)(block_sum)/980;
 	double glob_time = (double)(glob_sum)/980;
 
-	std::cout << "Average time per shared atomicAdd:\n" << block_time/num_threads(10000) << "ms\n\n" << std::endl;
-	std::cout << "Total time in shared atomicAdd:\n" << block_time << "ms\n\n" << std::endl;
-	std::cout << "Average time per global atomicAdd:\n" << glob_time/num_blocks(10000) << "ms\n\n" << std::endl;
-	std::cout << "Total time in global atomicAdd:\n" << glob_time << "ms\n\n" << std::endl;
-	std::cout << "Total time in all atomicAdd:\n" << (block_time + glob_time) << "ms\n\n" << std::endl;
+	std::cout << "Average time per shared atomicAdd:\n" << block_time/num_threads(10000) << "us\n\n" << std::endl;
+	std::cout << "Total time in shared atomicAdd:\n" << block_time << "us\n\n" << std::endl;
+	std::cout << "Average time per global atomicAdd:\n" << glob_time/num_blocks(10000) << "us\n\n" << std::endl;
+	std::cout << "Total time in global atomicAdd:\n" << glob_time << "us\n\n" << std::endl;
+	std::cout << "Total time in all atomicAdd:\n" << (block_time + glob_time) << "us\n\n" << std::endl;
 
 }