From ca1b1cea796bcfaeb86f201cf35065a606921cc1 Mon Sep 17 00:00:00 2001
From: root <root@ipepdvdev1.ipe.kit.edu>
Date: Tue, 24 May 2016 00:04:46 +0200
Subject: Read hw counter

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

(limited to 'kernels.cu')

diff --git a/kernels.cu b/kernels.cu
index 341bb59..3b5e1d7 100644
--- a/kernels.cu
+++ b/kernels.cu
@@ -17,8 +17,20 @@ __global__ void ipedma(volatile void *bar, uintptr_t bus_addr, volatile uint64_t
 
 	desc[1] = 0;
 
+	    // Reset counter
+#ifdef USE_HW_CONTER
+        WR32 (REG_DMA, 0);
+	WR32 (REG_PERF_COUNTER, 0);
+
+	WR64 (REG_DESCRIPTOR_ADDRESS, bus_addr);
+
+        t1 = clock64();
+	WR32 (REG_DMA, 1);
+#else
         t1 = clock64();
 	WR64 (REG_DESCRIPTOR_ADDRESS, bus_addr);
+#endif
+
 
         do {
 	    if (++wait > 0x10000) break;
@@ -26,6 +38,10 @@ __global__ void ipedma(volatile void *bar, uintptr_t bus_addr, volatile uint64_t
 
 	t2 = clock64();
 
+#ifdef USE_HW_CONTER
+	WR32 (REG_PERF_COUNTER, 1);
+#endif
+
 	null<<<1,1>>>(data);
 	cudaDeviceSynchronize();
 	t3 = clock64();
-- 
cgit v1.2.3