summaryrefslogtreecommitdiffstats
path: root/kernels.cu
diff options
context:
space:
mode:
authorSuren A. Chilingaryan <csa@suren.me>2016-06-17 21:14:09 +0200
committerSuren A. Chilingaryan <csa@suren.me>2016-06-17 21:14:09 +0200
commit8a59e1d17a83e4744071dfa790db974c296c206e (patch)
tree935fed6d263103b2249286d5fcd9b28470e45317 /kernels.cu
parenta6f3e96f2cafc183ab29e53007a86bb968d654b8 (diff)
downloadgdrtest-master.tar.gz
gdrtest-master.tar.bz2
gdrtest-master.tar.xz
gdrtest-master.zip
Use undocumented event-based APIHEADmaster
Diffstat (limited to 'kernels.cu')
-rw-r--r--kernels.cu18
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;