diff options
| author | root <root@ipepdvdev1.ipe.kit.edu> | 2016-05-24 00:04:46 +0200 |
|---|---|---|
| committer | root <root@ipepdvdev1.ipe.kit.edu> | 2016-05-24 00:04:46 +0200 |
| commit | ca1b1cea796bcfaeb86f201cf35065a606921cc1 (patch) | |
| tree | 9dbe7a85bad655dbbc7852149c749d9e85573601 /kernels.cu | |
| parent | 16e0aeeed527f8452e336685f664d7aa848702d3 (diff) | |
| download | gdrtest-ca1b1cea796bcfaeb86f201cf35065a606921cc1.tar.gz gdrtest-ca1b1cea796bcfaeb86f201cf35065a606921cc1.tar.bz2 gdrtest-ca1b1cea796bcfaeb86f201cf35065a606921cc1.tar.xz gdrtest-ca1b1cea796bcfaeb86f201cf35065a606921cc1.zip | |
Read hw counter
Diffstat (limited to 'kernels.cu')
| -rw-r--r-- | kernels.cu | 16 |
1 files changed, 16 insertions, 0 deletions
@@ -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(); |