Initial

1 files changed, 19 insertions, 0 deletions

diff --git a/kernel.cl b/kernel.cl
new file mode 100644
index 0000000..26a0009
--- /dev/null
+++ b/kernel.cl
@@ -0,0 +1,19 @@
+#define REG_PERF_COUNTER 0x20
+#define REG_DESCRIPTOR_ADDRESS 0x50
+
+#define RD32(addr) ((__global volatile uint*)bar)[addr / 4];
+#define WR32(addr, value) ((__global volatile uint*)bar)[addr / 4] = value;
+#define WR64(addr, value) ((__global volatile ulong*)bar)[addr / 8] = value;
+
+__kernel void process(uint iter, __global uint *input, __global uint *output, __global volatile uint *bar, ulong bus_addr)
+{
+    WR32 (REG_PERF_COUNTER, 1);
+}
+
+
+__kernel void measure(uint iter, __global uint *input, __global uint *output, __global volatile uint *bar, ulong bus_addr)
+{
+    output[0] = iter;
+    output[iter + 1] = RD32(REG_PERF_COUNTER);
+    WR64 (REG_DESCRIPTOR_ADDRESS, bus_addr);
+}