Improve read test

mconcas · mconcas · commit dcec171d6d50 · 2021-10-05T17:14:16.000+02:00
diff --git a/GPU/GPUbenchmark/cuda/Kernels.cu b/GPU/GPUbenchmark/cuda/Kernels.cu
@@ -74,11 +74,10 @@ __global__ void readChunkSBKernel(
   size_t chunkSize)
 {
   chunk_t sink{0}; // local memory -> excluded from bandwidth accounting
-  size_t last{0};
-  for (last = threadIdx.x; last < chunkSize; last += blockDim.x) {
-    sink += chunkPtr[last]; // 1 read operation, performed "chunkSize" times
+  for (size_t i = threadIdx.x; i < chunkSize; i += blockDim.x) {
+    sink = chunkPtr[i]; // 1 read operation, performed "chunkSize" times
   }
-  chunkPtr[last] = sink; // writing done once
+  chunkPtr[threadIdx.x] = sink; // writing done once
 }
 
 template <class chunk_t>
@@ -87,11 +86,10 @@ __global__ void readChunkMBKernel(
   size_t chunkSize)
 {
   chunk_t sink{0}; // local memory -> excluded from bandwidth accounting
-  size_t last{0};
-  for (last = blockIdx.x * blockDim.x + threadIdx.x; last < chunkSize; last += blockDim.x * gridDim.x) {
-    sink += chunkPtr[last]; // 1 read operation, performed "chunkSize" times
+  for (size_t i = blockIdx.x * blockDim.x + threadIdx.x; i < chunkSize; i += blockDim.x * gridDim.x) {
+    sink = chunkPtr[i]; // 1 read operation, performed "chunkSize" times
   }
-  chunkPtr[last] = sink;
+  chunkPtr[threadIdx.x] = sink;
 }
 
 // Write
@@ -853,7 +851,7 @@ void GPUbenchmark<chunk_t>::run()
 template class GPUbenchmark<char>;
 template class GPUbenchmark<size_t>;
 template class GPUbenchmark<int>;
-// template class GPUbenchmark<uint4>;
+// template class GPUbenchmark<int4>;
 
 } // namespace benchmark
 } // namespace o2

Original file line number	Diff line number	Diff line change
`@@ -74,11 +74,10 @@ __global__ void readChunkSBKernel(`
`74`	`74`	`size_t chunkSize)`
`75`	`75`	`{`
`76`	`76`	`chunk_t sink{0}; // local memory -> excluded from bandwidth accounting`
`77`		`- size_t last{0};`
`78`		`- for (last = threadIdx.x; last < chunkSize; last += blockDim.x) {`
`79`		`- sink += chunkPtr[last]; // 1 read operation, performed "chunkSize" times`
	`77`	`+ for (size_t i = threadIdx.x; i < chunkSize; i += blockDim.x) {`
	`78`	`+ sink = chunkPtr[i]; // 1 read operation, performed "chunkSize" times`
`80`	`79`	`}`
`81`		`- chunkPtr[last] = sink; // writing done once`
	`80`	`+ chunkPtr[threadIdx.x] = sink; // writing done once`
`82`	`81`	`}`
`83`	`82`
`84`	`83`	`template <class chunk_t>`
`@@ -87,11 +86,10 @@ __global__ void readChunkMBKernel(`
`87`	`86`	`size_t chunkSize)`
`88`	`87`	`{`
`89`	`88`	`chunk_t sink{0}; // local memory -> excluded from bandwidth accounting`
`90`		`- size_t last{0};`
`91`		`- for (last = blockIdx.x * blockDim.x + threadIdx.x; last < chunkSize; last += blockDim.x * gridDim.x) {`
`92`		`- sink += chunkPtr[last]; // 1 read operation, performed "chunkSize" times`
	`89`	`+ for (size_t i = blockIdx.x * blockDim.x + threadIdx.x; i < chunkSize; i += blockDim.x * gridDim.x) {`
	`90`	`+ sink = chunkPtr[i]; // 1 read operation, performed "chunkSize" times`
`93`	`91`	`}`
`94`		`- chunkPtr[last] = sink;`
	`92`	`+ chunkPtr[threadIdx.x] = sink;`
`95`	`93`	`}`
`96`	`94`
`97`	`95`	`// Write`
`@@ -853,7 +851,7 @@ void GPUbenchmark<chunk_t>::run()`
`853`	`851`	`template class GPUbenchmark<char>;`
`854`	`852`	`template class GPUbenchmark<size_t>;`
`855`	`853`	`template class GPUbenchmark<int>;`
`856`		`-// template class GPUbenchmark<uint4>;`
	`854`	`+// template class GPUbenchmark<int4>;`
`857`	`855`
`858`	`856`	`} // namespace benchmark`
`859`	`857`	`} // namespace o2`