Use undocumented event-based APIHEADmaster

1 files changed, 18 insertions, 0 deletions

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;