Change Valid Name for GatherScatter Kernel

README.md

@@ -87,7 +87,7 @@ Usage: ./spatter
 -d (--delta) Delta (default 8)
 -e (--boundary) Set Boundary (limits max value of pattern using modulo)
 -f (--file) Input File
--g (--pattern-gather) Set Inner Gather Pattern (Valid with kernel-name: sg, multigather)
+-g (--pattern-gather) Set Inner Gather Pattern (Valid with kernel-name: gs, multigather)
 -h (--help) Print Help Message
 -j (--pattern-size) Set Pattern Size (truncates pattern to pattern-size)
 -k (--kernel) Kernel (default gather)
@@ -98,7 +98,7 @@ Usage: ./spatter
 -r (--runs) Set Number of Runs (default 10)
 -s (--random) Set Random Seed (default random)
 -t (--omp-threads) Set Number of Threads (default 1 if !USE_OPENMP or backend != openmp or OMP_MAX_THREADS if USE_OPENMP)
--u (--pattern-scatter) Set Inner Scatter Pattern (Valid with kernel-name: sg, multiscatter)
+-u (--pattern-scatter) Set Inner Scatter Pattern (Valid with kernel-name: gs, multiscatter)
 -v (--verbosity) Set Verbosity Level (default 1)
 -w (--wrap) Set Wrap (default 1)
 -x (--delta-gather) Delta (default 8)

notebooks/spatter_util.py

@@ -58,7 +58,7 @@
 
 EXPERIMENTS = {'ustride':'Uniform Stride', 'stream': 'Stream', 'nekbone':'Nekbone', 'lulesh':'LULESH', 'amg':'AMG', 'pennant':'PENNANT'}
 
-KERNELS = {'gather':'Gather', 'scatter':'Scatter', 'sg':'GS', 'multigather':'MultiGather', 'multiscatter':'MultiScatter'}
+KERNELS = {'gather':'Gather', 'scatter':'Scatter', 'gs':'GS', 'multigather':'MultiGather', 'multiscatter':'MultiScatter'}
 
 #################################################################
 # NO EDITING IS REQUIRED BEYOND THIS POINT TO ADD NEW PLATFORMS #

src/Spatter/Configuration.cc

@@ -49,8 +49,8 @@ int ConfigurationBase::run(bool timed, unsigned long run_id) {
     gather(timed, run_id);
   else if (kernel.compare("scatter") == 0)
     scatter(timed, run_id);
-  else if (kernel.compare("sg") == 0)
-    scatter_gather(timed, run_id);
+  else if (kernel.compare("gs") == 0)
+    gather_scatter(timed, run_id);
   else if (kernel.compare("multigather") == 0)
     multi_gather(timed, run_id);
   else if (kernel.compare("multiscatter") == 0)
@@ -69,7 +69,7 @@ void ConfigurationBase::report() {
   if (kernel.compare("gather") == 0 || kernel.compare("scatter") == 0)
     bytes_moved = pattern.size() * count * sizeof(size_t);
 
-  if (kernel.compare("sg") == 0)
+  if (kernel.compare("gs") == 0)
     bytes_moved = (pattern_scatter.size() + pattern_gather.size()) * count * sizeof(size_t);
 
   if (kernel.compare("multiscatter") == 0)
@@ -142,7 +142,7 @@ void ConfigurationBase::setup() {
                 << std::endl;
       exit(1);
     }
-  } else if (kernel.compare("sg") == 0) {
+  } else if (kernel.compare("gs") == 0) {
     if (pattern_gather.size() == 0) {
       std::cerr << "Pattern-Gather needs to have length of at least 1"
                 << std::endl;
@@ -185,7 +185,7 @@ void ConfigurationBase::setup() {
   // sparse size = max_pattern_val + delta * (count - 1) + 1
   // assert(pattern.size() > max_pattern_scatter_val + 1)
 
-  if (kernel.compare("sg") == 0) {
+  if (kernel.compare("gs") == 0) {
     size_t max_pattern_scatter_val = *(std::max_element(
         std::cbegin(pattern_scatter), std::cend(pattern_scatter)));
     size_t max_pattern_gather_val = *(std::max_element(
@@ -442,7 +442,7 @@ void Configuration<Spatter::Serial>::scatter(bool timed, unsigned long run_id) {
   }
 }
 
-void Configuration<Spatter::Serial>::scatter_gather(
+void Configuration<Spatter::Serial>::gather_scatter(
     bool timed, unsigned long run_id) {
   assert(pattern_scatter.size() == pattern_gather.size());
   size_t pattern_length = pattern_scatter.size();
@@ -615,7 +615,7 @@ void Configuration<Spatter::OpenMP>::scatter(bool timed, unsigned long run_id) {
   }
 }
 
-void Configuration<Spatter::OpenMP>::scatter_gather(
+void Configuration<Spatter::OpenMP>::gather_scatter(
     bool timed, unsigned long run_id) {
   assert(pattern_scatter.size() == pattern_gather.size());
   size_t pattern_length = pattern_scatter.size();
@@ -815,7 +815,7 @@ void Configuration<Spatter::CUDA>::scatter(bool timed, unsigned long run_id) {
     time_seconds[run_id] = ((double)time_ms / 1000.0);
 }
 
-void Configuration<Spatter::CUDA>::scatter_gather(
+void Configuration<Spatter::CUDA>::gather_scatter(
     bool timed, unsigned long run_id) {
   assert(pattern_scatter.size() == pattern_gather.size());
   int pattern_length = static_cast<int>(pattern_scatter.size());
@@ -827,11 +827,11 @@ void Configuration<Spatter::CUDA>::scatter_gather(
   float time_ms = 0.0;
 
   if (atomic)
-    time_ms = cuda_scatter_gather_atomic_wrapper(dev_pattern_scatter,
+    time_ms = cuda_gather_scatter_atomic_wrapper(dev_pattern_scatter,
         dev_sparse_scatter, dev_pattern_gather, dev_sparse_gather,
         pattern_length, delta_scatter, delta_gather, wrap, count);
   else
-    time_ms = cuda_scatter_gather_wrapper(dev_pattern_scatter,
+    time_ms = cuda_gather_scatter_wrapper(dev_pattern_scatter,
         dev_sparse_scatter, dev_pattern_gather, dev_sparse_gather,
         pattern_length, delta_scatter, delta_gather, wrap, count);
 

src/Spatter/Configuration.hh

@@ -77,7 +77,7 @@ public:
 
   virtual void gather(bool timed, unsigned long run_id) = 0;
   virtual void scatter(bool timed, unsigned long run_id) = 0;
-  virtual void scatter_gather(bool timed, unsigned long run_id) = 0;
+  virtual void gather_scatter(bool timed, unsigned long run_id) = 0;
   virtual void multi_gather(bool timed, unsigned long run_id) = 0;
   virtual void multi_scatter(bool timed, unsigned long run_id) = 0;
 
@@ -168,7 +168,7 @@ public:
 
   void gather(bool timed, unsigned long run_id);
   void scatter(bool timed, unsigned long run_id);
-  void scatter_gather(bool timed, unsigned long run_id);
+  void gather_scatter(bool timed, unsigned long run_id);
   void multi_gather(bool timed, unsigned long run_id);
   void multi_scatter(bool timed, unsigned long run_id);
 };
@@ -197,7 +197,7 @@ public:
 
   void gather(bool timed, unsigned long run_id);
   void scatter(bool timed, unsigned long run_id);
-  void scatter_gather(bool timed, unsigned long run_id);
+  void gather_scatter(bool timed, unsigned long run_id);
   void multi_gather(bool timed, unsigned long run_id);
   void multi_scatter(bool timed, unsigned long run_id);
 };
@@ -228,7 +228,7 @@ public:
   int run(bool timed, unsigned long run_id);
   void gather(bool timed, unsigned long run_id);
   void scatter(bool timed, unsigned long run_id);
-  void scatter_gather(bool timed, unsigned long run_id);
+  void gather_scatter(bool timed, unsigned long run_id);
   void multi_gather(bool timed, unsigned long run_id);
   void multi_scatter(bool timed, unsigned long run_id);
   void setup();

src/Spatter/CudaBackend.cu

@@ -48,7 +48,7 @@ __global__ void cuda_scatter_atomic(const size_t *pattern, double *sparse,
         __double_as_longlong(dense[j + pattern_length * (i % wrap)]));
 }
 
-__global__ void cuda_scatter_gather(const size_t *pattern_scatter,
+__global__ void cuda_gather_scatter(const size_t *pattern_scatter,
     double *sparse_scatter, const size_t *pattern_gather,
     const double *sparse_gather, const size_t pattern_length,
     const size_t delta_scatter, const size_t delta_gather, const size_t wrap,
@@ -64,7 +64,7 @@ __global__ void cuda_scatter_gather(const size_t *pattern_scatter,
         sparse_gather[pattern_gather[j] + delta_gather * i];
 }
 
-__global__ void cuda_scatter_gather_atomic(const size_t *pattern_scatter,
+__global__ void cuda_gather_scatter_atomic(const size_t *pattern_scatter,
     double *sparse_scatter, const size_t *pattern_gather,
     const double *sparse_gather, const size_t pattern_length,
     const size_t delta_scatter, const size_t delta_gather, const size_t wrap,
@@ -224,7 +224,7 @@ float cuda_scatter_atomic_wrapper(const size_t *pattern, double *sparse,
   return time_ms;
 }
 
-float cuda_scatter_gather_wrapper(const size_t *pattern_scatter,
+float cuda_gather_scatter_wrapper(const size_t *pattern_scatter,
     double *sparse_scatter, const size_t *pattern_gather,
     const double *sparse_gather, const size_t pattern_length,
     const size_t delta_scatter, const size_t delta_gather, const size_t wrap,
@@ -241,7 +241,7 @@ float cuda_scatter_gather_wrapper(const size_t *pattern_scatter,
   checkCudaErrors(cudaDeviceSynchronize());
   checkCudaErrors(cudaEventRecord(start));
 
-  cuda_scatter_gather<<<blocks_per_grid, threads_per_block>>>(pattern_scatter,
+  cuda_gather_scatter<<<blocks_per_grid, threads_per_block>>>(pattern_scatter,
       sparse_scatter, pattern_gather, sparse_gather, pattern_length,
       delta_scatter, delta_gather, wrap, count);
   checkCudaErrors(cudaGetLastError());
@@ -258,7 +258,7 @@ float cuda_scatter_gather_wrapper(const size_t *pattern_scatter,
   return time_ms;
 }
 
-float cuda_scatter_gather_atomic_wrapper(const size_t *pattern_scatter,
+float cuda_gather_scatter_atomic_wrapper(const size_t *pattern_scatter,
     double *sparse_scatter, const size_t *pattern_gather,
     const double *sparse_gather, const size_t pattern_length,
     const size_t delta_scatter, const size_t delta_gather, const size_t wrap,
@@ -275,7 +275,7 @@ float cuda_scatter_gather_atomic_wrapper(const size_t *pattern_scatter,
   checkCudaErrors(cudaDeviceSynchronize());
   checkCudaErrors(cudaEventRecord(start));
 
-  cuda_scatter_gather_atomic<<<blocks_per_grid, threads_per_block>>>(
+  cuda_gather_scatter_atomic<<<blocks_per_grid, threads_per_block>>>(
       pattern_scatter, sparse_scatter, pattern_gather, sparse_gather,
       pattern_length, delta_scatter, delta_gather, wrap, count);
   checkCudaErrors(cudaGetLastError());

src/Spatter/CudaBackend.hh

@@ -12,12 +12,12 @@ float cuda_scatter_wrapper(const size_t *pattern, double *sparse,
 float cuda_scatter_atomic_wrapper(const size_t *pattern, double *sparse,
     const double *dense, const size_t pattern_length, const size_t delta,
     const size_t wrap, const size_t count);
-float cuda_scatter_gather_wrapper(const size_t *pattern_scatter,
+float cuda_gather_scatter_wrapper(const size_t *pattern_scatter,
     double *sparse_scatter, const size_t *pattern_gather,
     const double *sparse_gather, const size_t pattern_length,
     const size_t delta_scatter, const size_t delta_gather, const size_t wrap,
     const size_t count);
-float cuda_scatter_gather_atomic_wrapper(const size_t *pattern_scatter,
+float cuda_gather_scatter_atomic_wrapper(const size_t *pattern_scatter,
     double *sparse_scatter, const size_t *pattern_gather,
     const double *sparse_gather, const size_t pattern_length,
     const size_t delta_scatter, const size_t delta_gather, const size_t wrap,

src/Spatter/Input.hh

@@ -147,7 +147,7 @@ void help(char *progname) {
             << "Input File" << std::left << "\n";
   std::cout
       << std::left << std::setw(10) << "-g (--pattern-gather)" << std::setw(40)
-      << "Set Inner Gather Pattern (Valid with kernel-name: sg, multigather)"
+      << "Set Inner Gather Pattern (Valid with kernel-name: gs, multigather)"
       << std::left << "\n";
   std::cout << std::left << std::setw(10) << "-h (--help)" << std::setw(40)
             << "Print Help Message" << std::left << "\n";
@@ -182,7 +182,7 @@ void help(char *progname) {
   std::cout << std::left << std::setw(10)
             << "-u (--pattern-scatter)" << std::setw(4)
             << "Set Inner Scatter Pattern "
-            << "(Valid with kernel-name: sg, multiscatter)"
+            << "(Valid with kernel-name: gs, multiscatter)"
             << std::left << "\n";
   std::cout << std::left << std::setw(10) << "-v (--verbosity)" << std::setw(40)
             << "Set Verbosity Level (default 1)" << std::left << "\n";
@@ -405,9 +405,9 @@ int parse_input(const int argc, char **argv, ClArgs &cl) {
           [](unsigned char c) { return std::tolower(c); });
 
       if ((kernel.compare("gather") != 0) && (kernel.compare("scatter") != 0) &&
-          (kernel.compare("sg") != 0) && (kernel.compare("multigather") != 0) &&
+          (kernel.compare("gs") != 0) && (kernel.compare("multigather") != 0) &&
           (kernel.compare("multiscatter") != 0)) {
-        std::cerr << "Valid Kernels are: gather, scatter, sg, multigather, "
+        std::cerr << "Valid Kernels are: gather, scatter, gs, multigather, "
                      "multiscatter"
                   << std::endl;
         return -1;

src/Spatter/JSONParser.cc

@@ -59,11 +59,7 @@ JSONParser::JSONParser(std::string filename, aligned_vector<double> &sparse,
       std::transform(kernel.begin(), kernel.end(), kernel.begin(),
           [](unsigned char c) { return std::tolower(c); });
 
-      // The kernel may be specified as 'GS' instead of 'sg'
-      kernel = (kernel.compare("gs") == 0) ? "sg" : kernel;
-      v["kernel"] = kernel;
-
-      if (kernel.compare("sg") == 0) {
+      if (kernel.compare("gs") == 0) {
         // This kernel does not require --pattern to be specified
         assert(v.contains("pattern-gather") && v.contains("pattern-scatter"));
       } else {

tests/concurrent.cc

@@ -5,7 +5,7 @@ int uniform_concurrent_test() {
     char *command;
 
     int ret = asprintf(&command,
-        "../spatter -kSG -uUNIFORM:%d:%d -gUNIFORM:%d:%d ", i, i, i,
+        "../spatter -kGS -uUNIFORM:%d:%d -gUNIFORM:%d:%d ", i, i, i,
         i);
     if (ret == -1 || system(command) != EXIT_SUCCESS) {
       std::cerr << "Test failure on " << command << std::endl;

tests/parse_concurrent.cc

@@ -14,7 +14,7 @@ int main(int argc, char **argv) {
   asprintf(&argv_[0], "./spatter");
   asprintf(&argv_[1], "-uUNIFORM:8:1");
   asprintf(&argv_[2], "-gUNIFORM:8:1");
-  asprintf(&argv_[3], "-kSG");
+  asprintf(&argv_[3], "-kGS");
 
   Spatter::ClArgs cl;
 

tests/parse_json_suite.cc

@@ -62,8 +62,8 @@ int json_test(size_t run_count, std::string kernel,
       return EXIT_FAILURE;
     }
 
-    if (!kernel.compare("sg") && cl.configs[i]->kernel.compare("sg")) {
-      std::cerr << "Test failure on JSON Parse: user requested kernel SG but "
+    if (!kernel.compare("gs") && cl.configs[i]->kernel.compare("gs")) {
+      std::cerr << "Test failure on JSON Parse: user requested kernel GS but "
                    "instead got other kernel "
                 << cl.configs[i]->kernel << "." << std::endl;
       return EXIT_FAILURE;

tests/parse_run_config_suite_cpu.cc

@@ -31,19 +31,19 @@ int parse_check(int argc_, char **argv_, Spatter::ClArgs &cl) {
 }
 
 int k_tests(int argc_, char **argv_) {
-  int sg_argc_ = 4;
-  char **sg_argv_ = (char **)malloc(sizeof(char *) * sg_argc_);
+  int gs_argc_ = 4;
+  char **gs_argv_ = (char **)malloc(sizeof(char *) * gs_argc_);
 
   int ret;
-  ret = asprintf(&sg_argv_[0], "./spatter");
+  ret = asprintf(&gs_argv_[0], "./spatter");
   if (ret == -1)
     return EXIT_FAILURE;
 
-  ret = asprintf(&sg_argv_[1], "-g1,2,3,4");
+  ret = asprintf(&gs_argv_[1], "-g1,2,3,4");
   if (ret == -1)
     return EXIT_FAILURE;
 
-  ret = asprintf(&sg_argv_[2], "-u1,2,3,4");
+  ret = asprintf(&gs_argv_[2], "-u1,2,3,4");
   if (ret == -1)
     return EXIT_FAILURE;
 
@@ -76,15 +76,15 @@ int k_tests(int argc_, char **argv_) {
     return EXIT_FAILURE;
   }
 
-  asprintf(&sg_argv_[3], "-kSG");
+  asprintf(&gs_argv_[3], "-kGS");
 
   Spatter::ClArgs cl3;
-  if (parse_check(sg_argc_, sg_argv_, cl3) == EXIT_FAILURE)
+  if (parse_check(gs_argc_, gs_argv_, cl3) == EXIT_FAILURE)
     return EXIT_FAILURE;
 
-  free(sg_argv_[3]);
+  free(gs_argv_[3]);
 
-  if (cl3.configs[0]->kernel.compare("sg") != 0) {
+  if (cl3.configs[0]->kernel.compare("gs") != 0) {
     std::cerr << "Test failure on Run_Config Suite: POSIX-style k with "
                  "argument GS resulted in kernel "
               << cl3.configs[0]->kernel << "." << std::endl;
@@ -121,16 +121,16 @@ int k_tests(int argc_, char **argv_) {
     return EXIT_FAILURE;
   }
 
-  asprintf(&sg_argv_[3], "--kernel=SG");
+  asprintf(&gs_argv_[3], "--kernel=GS");
 
   Spatter::ClArgs cl6;
-  if (parse_check(sg_argc_, sg_argv_, cl6) == EXIT_FAILURE)
+  if (parse_check(gs_argc_, gs_argv_, cl6) == EXIT_FAILURE)
     return EXIT_FAILURE;
 
-  free(sg_argv_[3]);
+  free(gs_argv_[3]);
 
-  if (cl6.configs[0]->kernel.compare("sg") != 0) {
-    std::cerr << "Test failure on Run_Config Suite: --kernel with argument SG "
+  if (cl6.configs[0]->kernel.compare("gs") != 0) {
+    std::cerr << "Test failure on Run_Config Suite: --kernel with argument GS "
                  "resulted in kernel "
               << cl6.configs[0]->kernel << "." << std::endl;
     return EXIT_FAILURE;