Add int4 += custom implementation for CUDA completeness

mconcas · mconcas · commit a1583a708ba1 · 2021-10-08T18:57:01.000+02:00
diff --git a/GPU/GPUbenchmark/Shared/Utils.h b/GPU/GPUbenchmark/Shared/Utils.h
@@ -192,6 +192,19 @@ class BSDRnd : public LCGRnd
   __host__ __device__ int rnd() { return LCGRnd::rnd(); }
 };
 
+
+// CUDA does not support <type4> operations:
+// https://forums.developer.nvidia.com/t/swizzling-float4-arithmetic-support/217
+#ifndef __HIPCC__ 
+inline __host__ __device__ void operator+=(int4 &a, int4 b)
+{
+    a.x += b.x;
+    a.y += b.y;
+    a.z += b.z;
+    a.w += b.w;
+}
+#endif
+
 namespace o2
 {
 namespace benchmark
diff --git a/GPU/GPUbenchmark/cuda/Kernels.cu b/GPU/GPUbenchmark/cuda/Kernels.cu
@@ -49,7 +49,7 @@ __global__ void read_k(
 {
   chunk_t sink{0};
   for (size_t i = blockIdx.x * blockDim.x + threadIdx.x; i < chunkSize; i += blockDim.x * gridDim.x) {
-    sink = chunkPtr[i] + sink;
+    sink += chunkPtr[i];
   }
   chunkPtr[threadIdx.x] = sink;
 }

Original file line number	Diff line number	Diff line change
`@@ -49,7 +49,7 @@ __global__ void read_k(`
`49`	`49`	`{`
`50`	`50`	`chunk_t sink{0};`
`51`	`51`	`for (size_t i = blockIdx.x * blockDim.x + threadIdx.x; i < chunkSize; i += blockDim.x * gridDim.x) {`
`52`		`- sink = chunkPtr[i] + sink;`
	`52`	`+ sink += chunkPtr[i];`
`53`	`53`	`}`
`54`	`54`	`chunkPtr[threadIdx.x] = sink;`
`55`	`55`	`}`