plots

romnn · Nov 14, 2023 · 0ff9984 · 0ff9984
1 parent 0579a31
commit 0ff9984
Show file tree

Hide file tree

Showing 23 changed files with 284 additions and 152 deletions.
diff --git a/gpucachesim/microbench/pchase.py b/gpucachesim/microbench/pchase.py
diff --git a/gpucachesim/remote/__init__.py b/gpucachesim/remote/__init__.py
@@ -420,6 +420,8 @@ def submit_pchase(
             slurm_script += "export {}={}\n".format(k, v)
         slurm_script += "{} {}\n".format(executable, " ".join(args))
 
+        print([str(executable)] + args)
+
         # upload slurm script
         self.upload_data(data=StringIO(slurm_script), remote_path=remote_slurm_job_path)
 

diff --git a/plot/A4000/cache_line_size-l1data-native.pdf b/plot/A4000/cache_line_size-l1data-native.pdf
diff --git a/plot/A4000/cache_line_size-l2-native.pdf b/plot/A4000/cache_line_size-l2-native.pdf
diff --git a/plot/A4000/cache_set_mapping-l1data-native-cc86-random.pdf b/plot/A4000/cache_set_mapping-l1data-native-cc86-random.pdf
diff --git a/plot/A4000/cache_set_mapping-l1data-native-random (copy).pdf b/plot/A4000/cache_set_mapping-l1data-native-random (copy).pdf
diff --git a/plot/A4000/cache_set_mapping-l1data-native-random.pdf b/plot/A4000/cache_set_mapping-l1data-native-random.pdf
diff --git a/plot/A4000/cache_sets.l1data.native (copy).pdf b/plot/A4000/cache_sets.l1data.native (copy).pdf
diff --git a/plot/A4000/cache_size-l1data-native.pdf b/plot/A4000/cache_size-l1data-native.pdf
diff --git a/plot/A4000/cache_size-l2-native.pdf b/plot/A4000/cache_size-l2-native.pdf
diff --git a/plot/A4000/l2_prefetch_size-l2-native.pdf b/plot/A4000/l2_prefetch_size-l2-native.pdf
diff --git a/plot/A4000/latency_distribution-l1data-native.pdf b/plot/A4000/latency_distribution-l1data-native.pdf
diff --git a/plot/GTX980/cache_size-l1data-native.pdf b/plot/GTX980/cache_size-l1data-native.pdf
diff --git a/plot/GTX980/cache_size-l2-native.pdf b/plot/GTX980/cache_size-l2-native.pdf
diff --git a/plot/TitanX-Pascal/cache_line_size-l2-native (copy).pdf b/plot/TitanX-Pascal/cache_line_size-l2-native (copy).pdf
diff --git a/plot/TitanX-Pascal/cache_line_size-l2-native.pdf b/plot/TitanX-Pascal/cache_line_size-l2-native.pdf
diff --git a/plot/TitanX-Pascal/cache_size-l1data-native.pdf b/plot/TitanX-Pascal/cache_size-l1data-native.pdf
diff --git a/plot/TitanX-Pascal/cache_size-l2-native.pdf b/plot/TitanX-Pascal/cache_size-l2-native.pdf
diff --git a/plot/TitanX-Pascal/l2_prefetch_size-l2-native.pdf b/plot/TitanX-Pascal/l2_prefetch_size-l2-native.pdf
diff --git a/plot/cache_set_mapping-l1data-native-random.pdf b/plot/cache_set_mapping-l1data-native-random.pdf
diff --git a/plot/cache_size-l2-native.pdf b/plot/cache_size-l2-native.pdf
diff --git a/test-apps/microbenches/chxw/pchase.cu b/test-apps/microbenches/chxw/pchase.cu
@@ -504,7 +504,7 @@ int parametric_measure_global(unsigned int *h_a, unsigned int *d_a, memory mem,
     for (size_t k = 0; k < iter_size; k++) {
       unsigned int index = indexof(h_a, N, h_index[k]);
       assert(index == (N + h_index[k] - stride) % N);
-      if (k == 0) {
+      if (k == 0 && warmup_iterations < 1) {
         assert(index == 0);
       }
       unsigned int latency = (int)h_latency[k] - (int)clock_overhead;

diff --git a/test-apps/microbenches/chxw/random_set_mapping.cu b/test-apps/microbenches/chxw/random_set_mapping.cu
@@ -6,41 +6,10 @@
 #include <stdint.h>
 #include <stdio.h>
 #include <unordered_set>
-// #include <random>
 
 #include "common.hpp"
 #include "cuda_runtime.h"
 
-// __global__ __noinline__ void
-// global_latency_l1_random_set_mapping_host_mapped(
-//     unsigned int *array, int array_length, unsigned int *latency,
-//     unsigned int *index, int iter_size, size_t warmup_iterations,
-//     unsigned int overflow_index) {
-//   unsigned int start_time, end_time;
-//   volatile uint32_t j = 0;
-//
-//   for (int k = (int)warmup_iterations * -iter_size; k < iter_size; k++) {
-//     if (k >= 0 && j == 0) {
-//       // overflow the cache now
-//       index[k] = array[array_length + overflow_index];
-//     }
-//     if (k >= 0) {
-//       start_time = clock();
-//       j = array[j];
-//       index[k] = j;
-//       end_time = clock();
-//
-//       latency[k] = end_time - start_time;
-//     } else {
-//       j = array[j];
-//     }
-//   }
-//
-//   // store to avoid caching in readonly?
-//   array[array_length] = j;
-//   array[array_length + 1] = array[j];
-// }
-
 __global__ __noinline__ void global_latency_l1_random_set_mapping_host_mapped(
     unsigned int *array, int array_length, unsigned int *latency,
     unsigned int *index, int iter_size, size_t warmup_iterations,
@@ -51,10 +20,8 @@ __global__ __noinline__ void global_latency_l1_random_set_mapping_host_mapped(
   volatile uint32_t j = start_j;
 
   for (int k = (int)warmup_iterations * -iter_size; k < iter_size; k++) {
-    // if (k >= 0 && j == 0) {
     if (k == round_size) {
       // overflow the cache now
-      // index[k] = array[array_length + overflow_index];
       index[k] = array[(array_length + overflow_index) % (2 * array_length)];
     }
     if (k >= 0) {
@@ -84,7 +51,6 @@ __global__ __noinline__ void global_latency_l2_random_set_mapping_host_mapped(
   volatile uint32_t j = start_j;
 
   for (int k = (int)warmup_iterations * -iter_size; k < iter_size; k++) {
-    // if (k >= 0 && j == 0) {
     if (k == round_size) {
       // overflow the cache now
       index[k] = array[(array_length + overflow_index) % (2 * array_length)];