1 files changed, 16 insertions, 0 deletions
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();