Update reading kernels

mconcas · mconcas · commit c9dc62dfd7d4 · 2021-09-23T18:01:59.000+02:00
diff --git a/GPU/GPUbenchmark/cuda/Kernels.cu b/GPU/GPUbenchmark/cuda/Kernels.cu
@@ -91,12 +91,11 @@ __global__ void readChunkMBKernel(
   chunk_type* results,
   size_t chunkSize)
 {
+  chunk_type sink{0};
   for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < chunkSize; i += blockDim.x * gridDim.x) {
-    if (chunkPtr[i] == static_cast<chunk_type>(1)) { // actual read operation is performed here
-      results[chunkId] += chunkPtr[i];               // this case should never happen and waves should be always in sync
-      printf("Should never happen\n");
-    }
+    sink += chunkPtr[i];
   }
+  results[chunkId] = sink;
 }
 
 // Write

Original file line number	Diff line number	Diff line change
`@@ -91,12 +91,11 @@ __global__ void readChunkMBKernel(`
`91`	`91`	`chunk_type* results,`
`92`	`92`	`size_t chunkSize)`
`93`	`93`	`{`
	`94`	`+ chunk_type sink{0};`
`94`	`95`	`for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < chunkSize; i += blockDim.x * gridDim.x) {`
`95`		`- if (chunkPtr[i] == static_cast<chunk_type>(1)) { // actual read operation is performed here`
`96`		`- results[chunkId] += chunkPtr[i]; // this case should never happen and waves should be always in sync`
`97`		`- printf("Should never happen\n");`
`98`		`- }`
	`96`	`+ sink += chunkPtr[i];`
`99`	`97`	`}`
	`98`	`+ results[chunkId] = sink;`
`100`	`99`	`}`
`101`	`100`
`102`	`101`	`// Write`