From 151541b83d540c3476965368f819e48a7b289cad Mon Sep 17 00:00:00 2001
From: "Suren A. Chilingaryan" <csa@suren.me>
Date: Sun, 19 Jun 2016 05:11:03 +0200
Subject: Initial

---
 kernel.cl | 19 +++++++++++++++++++
 1 file changed, 19 insertions(+)
 create mode 100644 kernel.cl

(limited to 'kernel.cl')

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);
+}
-- 
cgit v1.2.3