From 8a59e1d17a83e4744071dfa790db974c296c206e Mon Sep 17 00:00:00 2001
From: "Suren A. Chilingaryan" <csa@suren.me>
Date: Fri, 17 Jun 2016 21:14:09 +0200
Subject: Use undocumented event-based API

---
 kernels.cu | 18 ++++++++++++++++++
 1 file changed, 18 insertions(+)

(limited to 'kernels.cu')

diff --git a/kernels.cu b/kernels.cu
index f4ea114..7e7e689 100644
--- a/kernels.cu
+++ b/kernels.cu
@@ -7,6 +7,24 @@
 __global__ void null(uint32_t *data) {
 }
 
+__global__ void feedback(volatile void *bar, uint32_t *data) {
+    WR32 (REG_PERF_COUNTER, 1);
+}
+
+__global__ void measure(int iter, volatile void *bar, uint64_t bus_addr, uint64_t *res, uint32_t *data) {
+	// Clocks are incorrect as not running while waiting.
+    //res[iter] = clock64();
+	// Thats does not work either (no RD support in kernels?)
+    //res[iter] = RD32(REG_PERF_COUNTER);
+
+    WR64 (REG_DESCRIPTOR_ADDRESS, bus_addr);
+
+    if ((iter + 1) == GPU_ITERS) {
+	WR32 (REG_PERF_COUNTER, 1);
+    }
+}
+
+
 __device__ void ksleep(uint32_t clocks) {
     clock_t start = clock64(), now;
 
-- 
cgit v1.2.3