ZFP: add tests for ZFP+CUDA

ornladios · Oct 26, 2021 · 32ff1b6 · 32ff1b6
1 parent d39e521
commit 32ff1b6
Show file tree

Hide file tree

Showing 5 changed files with 302 additions and 83 deletions.
diff --git a/examples/cuda/cudaWriteRead.cu b/examples/cuda/cudaWriteRead.cu
@@ -5,99 +5,111 @@
  */
 
 #include <ios>
-#include <vector>
 #include <iostream>
+#include <vector>
 
 #include <adios2.h>
 
 #include <cuda.h>
 #include <cuda_runtime.h>
 
-__global__ void update_array(float *vect, int val) {
-    vect[blockIdx.x] += val;
-}
+__global__ void update_array(float *vect, int val) { vect[blockIdx.x] += val; }
 
-int BPWrite(const std::string fname, const size_t N, int nSteps){
-  // Initialize the simulation data
-  float *gpuSimData;
-  cudaMalloc(&gpuSimData, N * sizeof(float));
-  cudaMemset(gpuSimData, 0, N);
-
-  // Set up the ADIOS structures
-  adios2::ADIOS adios;
-  adios2::IO io = adios.DeclareIO("WriteIO");
-
-  // Declare an array for the ADIOS data of size (NumOfProcesses * N)
-  const adios2::Dims shape{static_cast<size_t>(N)};
-  const adios2::Dims start{static_cast<size_t>(0)};
-  const adios2::Dims count{N};
-  auto data = io.DefineVariable<float>("data", shape, start, count);
-
-  adios2::Engine bpWriter = io.Open(fname, adios2::Mode::Write);
-
-  // Simulation steps
-  for (size_t step = 0; step < nSteps; ++step)
-  {
-      // Make a 1D selection to describe the local dimensions of the
-      // variable we write and its offsets in the global spaces
-      adios2::Box<adios2::Dims> sel({0}, {N});
-      data.SetSelection(sel);
-
-      // Start IO step every write step
-      bpWriter.BeginStep();
-      data.SetMemorySpace(adios2::MemorySpace::CUDA);
-      bpWriter.Put(data, gpuSimData);
-      bpWriter.EndStep();
-
-      // Update values in the simulation data
-      update_array<<<N,1>>>(gpuSimData, 10);
-  }
-
-  bpWriter.Close();
-  return 0;
-}
+int BPWrite(const std::string fname, const size_t N, int nSteps)
+{
+    // Initialize the simulation data
+    float *gpuSimData;
+    cudaMalloc(&gpuSimData, N * sizeof(float));
+    cudaMemset(gpuSimData, 0, N);
+
+    // Set up the ADIOS structures
+    adios2::ADIOS adios;
+    adios2::IO io = adios.DeclareIO("WriteIO");
 
-int BPRead(const std::string fname, const size_t N, int nSteps){
-  // Create ADIOS structures
-  adios2::ADIOS adios;
-  adios2::IO io = adios.DeclareIO("ReadIO");
-
-  adios2::Engine bpReader = io.Open(fname, adios2::Mode::Read);
-
-  auto data = io.InquireVariable<float>("data");
-  std::cout << "Steps expected by the reader: " << bpReader.Steps() << std::endl;
-  std::cout << "Expecting data per step: " << data.Shape()[0];
-  std::cout  << " elements" << std::endl;
-
-  int write_step = bpReader.Steps();
-  // Create the local buffer and initialize the access point in the ADIOS file
-  std::vector<float> simData(N); //set size to N
-  const adios2::Dims start{0};
-  const adios2::Dims count{N};
-  const adios2::Box<adios2::Dims> sel(start, count);
-  data.SetSelection(sel);
-
-  // Read the data in each of the ADIOS steps
-  for (size_t step = 0; step < write_step; step++)
-  {
-      data.SetStepSelection({step, 1});
-      bpReader.Get(data, simData.data());
-      bpReader.PerformGets();
-      std::cout << "Simualation step " << step << " : ";
-      std::cout << simData.size() << " elements: " << simData[1] << std::endl;
-  }
-  bpReader.Close();
-  return 0;
+    // Declare an array for the ADIOS data of size (NumOfProcesses * N)
+    const adios2::Dims shape{static_cast<size_t>(N)};
+    const adios2::Dims start{static_cast<size_t>(0)};
+    const adios2::Dims count{N};
+    auto data = io.DefineVariable<float>("data", shape, start, count);
+
+/*
+ * TODO: Check if ZPF has CUDA (a PR impl this is waiting to be merged)
+ */
+#ifdef ADIOS2_HAVE_ZFP
+    adios2::Operator ZFPOp =
+        adios.DefineOperator("ZFPCompressor", adios2::ops::LossyZFP);
+
+    data.AddOperation(ZFPOp, {{"rate", "8"}});
+#endif
+
+    adios2::Engine bpWriter = io.Open(fname, adios2::Mode::Write);
+
+    // Simulation steps
+    for (size_t step = 0; step < nSteps; ++step)
+    {
+        // Make a 1D selection to describe the local dimensions of the
+        // variable we write and its offsets in the global spaces
+        adios2::Box<adios2::Dims> sel({0}, {N});
+        data.SetSelection(sel);
+
+        // Start IO step every write step
+        bpWriter.BeginStep();
+        data.SetMemorySpace(adios2::MemorySpace::CUDA);
+        bpWriter.Put(data, gpuSimData);
+        bpWriter.EndStep();
+
+        // Update values in the simulation data
+        update_array<<<N, 1>>>(gpuSimData, 10);
+    }
+
+    bpWriter.Close();
+    return 0;
 }
 
-int main(int argc, char **argv){
-  const std::string fname("GPUWriteRead.bp");
-  const int device_id = 1;
-  cudaSetDevice(device_id);
-  const size_t N = 6000;
-  int nSteps = 10, ret = 0;
+int BPRead(const std::string fname, const size_t N, int nSteps)
+{
+    // Create ADIOS structures
+    adios2::ADIOS adios;
+    adios2::IO io = adios.DeclareIO("ReadIO");
+
+    adios2::Engine bpReader = io.Open(fname, adios2::Mode::Read);
+
+    auto data = io.InquireVariable<float>("data");
+    std::cout << "Steps expected by the reader: " << bpReader.Steps()
+              << std::endl;
+    std::cout << "Expecting data per step: " << data.Shape()[0];
+    std::cout << " elements" << std::endl;
+
+    int write_step = bpReader.Steps();
+    // Create the local buffer and initialize the access point in the ADIOS file
+    std::vector<float> simData(N); // set size to N
+    const adios2::Dims start{0};
+    const adios2::Dims count{N};
+    const adios2::Box<adios2::Dims> sel(start, count);
+    data.SetSelection(sel);
+
+    // Read the data in each of the ADIOS steps
+    for (size_t step = 0; step < write_step; step++)
+    {
+        data.SetStepSelection({step, 1});
+        bpReader.Get(data, simData.data());
+        bpReader.PerformGets();
+        std::cout << "Simualation step " << step << " : ";
+        std::cout << simData.size() << " elements: " << simData[1] << std::endl;
+    }
+    bpReader.Close();
+    return 0;
+}
 
-  ret += BPWrite(fname, N, nSteps);
-  ret += BPRead(fname, N, nSteps);
-  return ret;
+int main(int argc, char **argv)
+{
+    const std::string fname("GPUWriteRead.bp");
+    const int device_id = 1;
+    cudaSetDevice(device_id);
+    const size_t N = 6000;
+    int nSteps = 10, ret = 0;
+
+    ret += BPWrite(fname, N, nSteps);
+    ret += BPRead(fname, N, nSteps);
+    return ret;
 }
diff --git a/testing/adios2/engine/bp/operations/CMakeLists.txt b/testing/adios2/engine/bp/operations/CMakeLists.txt
@@ -22,13 +22,26 @@ if(ADIOS2_HAVE_ZFP)
 
   gtest_add_tests_helper(ZfpComplex MPI_ALLOW BPWriteRead Engine. "")
 
+
   foreach(tgt
       ${Test.Engine.BP.WriteReadZfpConfig-TARGETS}
       )
     target_compile_definitions(${tgt} PRIVATE
       "XML_CONFIG_DIR=${CMAKE_CURRENT_SOURCE_DIR}"
       )
   endforeach()
+
+  if(ADIOS2_HAVE_CUDA)
+    enable_language(CUDA)
+
+    gtest_add_tests_helper(WriteReadZfpCuda MPI_ALLOW BP Engine.BP. .BP4
+      WORKING_DIRECTORY ${BP4_DIR} EXTRA_ARGS "BP4"
+      )
+    set_source_files_properties(CudaRoutines.cu PROPERTIES LANGUAGE CUDA)
+    foreach(tgt ${Test.Engine.BP.WriteReadZfpCuda-TARGETS})
+      target_sources(${tgt} PRIVATE CudaRoutines.cu)
+    endforeach()
+  endif()
 endif()
 
 if(ADIOS2_HAVE_MGARD)

diff --git a/testing/adios2/engine/bp/operations/CudaRoutines.cu b/testing/adios2/engine/bp/operations/CudaRoutines.cu
@@ -0,0 +1,11 @@
+#include "CudaRoutines.h"
+
+__global__ void __cuda_increment(int offset, float *vec, float val)
+{
+    vec[blockIdx.x + offset] += val;
+}
+
+void cuda_increment(int M, int N, int offset, float *vec, float val)
+{
+    __cuda_increment<<<M, N>>>(offset, vec, val);
+}
diff --git a/testing/adios2/engine/bp/operations/CudaRoutines.h b/testing/adios2/engine/bp/operations/CudaRoutines.h
@@ -0,0 +1,9 @@
+#ifndef __TESTING_ADIOS2_CUDA_ROUTINES_H__
+#define __TESTING_ADIOS2_CUDA_ROUTINES_H__
+
+#include <cuda.h>
+#include <cuda_runtime.h>
+
+void cuda_increment(int M, int N, int offset, float *vec, float val);
+
+#endif