From c24916b5e06c109a1be4837d3e64dbd744c29568 Mon Sep 17 00:00:00 2001
From: Hansung Kim <hansung_kim@berkeley.edu>
Date: Wed, 5 Jul 2023 21:24:53 -0700
Subject: [PATCH] [tests] Add compute-bound variant of vecadd

This loops 1000 times over `sum += A[i] + B[i]`, making every memory op
hit at L1 cache.
---
 tests/opencl/vecadd/kernel.alll1hit.loop1000.cl | 12 ++++++++++++
 1 file changed, 12 insertions(+)
 create mode 100644 tests/opencl/vecadd/kernel.alll1hit.loop1000.cl

diff --git a/tests/opencl/vecadd/kernel.alll1hit.loop1000.cl b/tests/opencl/vecadd/kernel.alll1hit.loop1000.cl
new file mode 100644
index 00000000..a610b183
--- /dev/null
+++ b/tests/opencl/vecadd/kernel.alll1hit.loop1000.cl
@@ -0,0 +1,12 @@
+__kernel void vecadd (__global const float *A,
+	                    __global const float *B,
+	                    __global float *C)
+{
+  int gid = get_global_id(0);
+  float sum = 0.;
+  for (int i = 0; i < 1000; i++) {
+	  int addr = gid + (i % 2);
+	  sum += A[addr] + B[addr];
+  }
+  C[gid] = sum;
+}