summaryrefslogtreecommitdiffstats
path: root/kernels.cu
diff options
context:
space:
mode:
Diffstat (limited to 'kernels.cu')
-rw-r--r--kernels.cu16
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();