Async memcpy's and cleanup

Grid/lattice/Lattice_slicesum_gpu.h

@@ -4,9 +4,8 @@
 #include <cub/cub.cuh>
 #define gpucub cub
 #define gpuMalloc cudaMalloc
-#define gpuMemcpy cudaMemcpy
+#define gpuMemcpyDtoHAsync cudaMemcpyDtoHAsync
-#define gpuMemcpyDeviceToHost cudaMemcpyDeviceToHost
+#define gpuMemcpyHtoDAsync cudaMemcpyHtoDAsync
-#define gpuMemcpyHostToDevice cudaMemcpyHostToDevice
 #define gpuError_t cudaError_t
 #define gpuSuccess cudaSuccess
 
@@ -15,13 +14,11 @@
 #include <hipcub/hipcub.hpp>
 #define gpucub hipcub
 #define gpuMalloc hipMalloc
-#define gpuMemcpy hipMemcpy
+#define gpuMemcpyDtoHAsync hipMemcpyDtoHAsync
-#define gpuMemcpyDeviceToHost hipMemcpyDeviceToHost
+#define gpuMemcpyHtoDAsync hipMemcpyHtoDAsync
-#define gpuMemcpyHostToDevice hipMemcpyHostToDevice
 #define gpuError_t hipError_t
 #define gpuSuccess hipSuccess
 
-// extern hipStream_t computeStream;
 #endif
 
 
@@ -49,78 +46,101 @@ template<class vobj> inline void sliceSumGpu(const Lattice<vobj> &Data,std::vect
   int e2=    grid->_slice_block [orthogdim];
   int stride=grid->_slice_stride[orthogdim];
   int ostride=grid->_ostride[orthogdim];
+  size_t subvol_size = e1*e2;
+
   Vector<vobj> lvSum(rd); 
   Vector<sobj> lsSum(ld,Zero());                    
   commVector<vobj> reduction_buffer(rd*e1*e2);
   ExtractBuffer<sobj> extracted(Nsimd);                  
   
   result.resize(fd);
+
   for(int r=0;r<rd;r++){
     lvSum[r]=Zero();
   }
-  vobj identity;
+
-  zeroit(identity);
+  vobj vobj_zero; //Need to provide initial value for reduction operation
+  zeroit(vobj_zero);
   
   autoView( Data_v, Data, AcceleratorRead);
-  auto rb_p = &reduction_buffer[0];
 
+  auto rb_p = &reduction_buffer[0];
   void *helperArray = NULL;
   vobj *d_out;
   size_t temp_storage_bytes = 0;
-  size_t size = e1*e2;
-  std::vector<int> offsets(rd+1,0);
-  for (int i = 0; i < offsets.size(); i++) {
-    offsets[i] = i*size;
-  }
   int* d_offsets;
 
-  gpuError_t gpuErr = gpuMalloc(&d_out,rd*sizeof(vobj));
+  std::vector<int> offsets(rd+1,0);
-   if (gpuErr != gpuSuccess) {
+  
-    std::cout << "Lattice_slicesum_gpu.h: Encountered error during gpuMalloc(1) Error: " << gpuErr <<std::endl;
+  for (int i = 0; i < offsets.size(); i++) {
-  }
+    offsets[i] = i*subvol_size;
-  gpuErr = gpuMalloc(&d_offsets,sizeof(int)*(rd+1));
-   if (gpuErr != gpuSuccess) {
-    std::cout << "Lattice_slicesum_gpu.h: Encountered error during gpuMalloc(2) Error: " << gpuErr <<std::endl;
-  }
-  gpuErr = gpuMemcpy(d_offsets,&offsets[0],sizeof(int)*(rd+1),gpuMemcpyHostToDevice);
-  if (gpuErr != gpuSuccess) {
-    std::cout << "Lattice_slicesum_gpu.h: Encountered error during gpuMemcpy(1) Error: " << gpuErr <<std::endl;
-  }
-
-  gpuErr = gpucub::DeviceSegmentedReduce::Reduce(helperArray, temp_storage_bytes, rb_p,d_out, rd, d_offsets, d_offsets+1, ::gpucub::Sum(), identity, computeStream);
-  if (gpuErr!=gpuSuccess) {
-    std::cout << "Lattice_slicesum_gpu.h: Encountered error during cub::DeviceReduce::Sum(1)! Error: " << gpuErr <<std::endl;
   }
   
+  //Allocate memory for output and offset arrays on device
+  gpuError_t gpuErr = gpuMalloc(&d_out,rd*sizeof(vobj));
+  if (gpuErr != gpuSuccess) {
+    std::cout << GridLogError << "Lattice_slicesum_gpu.h: Encountered error during gpuMalloc (d_out)! Error: " << gpuErr <<std::endl;
+    exit(EXIT_FAILURE);
+  }
+
+  gpuErr = gpuMalloc(&d_offsets,sizeof(int)*(rd+1));
+  if (gpuErr != gpuSuccess) {
+    std::cout << GridLogError << "Lattice_slicesum_gpu.h: Encountered error during gpuMalloc (d_offsets)! Error: " << gpuErr <<std::endl;
+    exit(EXIT_FAILURE);
+  }
+
+  //copy offsets to device
+  gpuErr = gpuMemcpyHtoDAsync(d_offsets,&offsets[0],sizeof(int)*(rd+1),computeStream);
+  if (gpuErr != gpuSuccess) {
+    std::cout << GridLogError << "Lattice_slicesum_gpu.h: Encountered error during gpuMemcpy (d_offsets)! Error: " << gpuErr <<std::endl;
+    exit(EXIT_FAILURE);
+  }
+
+  //determine helperArray size
+  gpuErr = gpucub::DeviceSegmentedReduce::Reduce(helperArray, temp_storage_bytes, rb_p,d_out, rd, d_offsets, d_offsets+1, ::gpucub::Sum(), vobj_zero, computeStream);
+  if (gpuErr!=gpuSuccess) {
+    std::cout << GridLogError << "Lattice_slicesum_gpu.h: Encountered error during gpucub::DeviceSegmentedReduce::Reduce (setup)! Error: " << gpuErr <<std::endl;
+    exit(EXIT_FAILURE);
+  }
+
+  //allocate memory for helperArray  
   gpuErr = gpuMalloc(&helperArray,temp_storage_bytes);
   if (gpuErr!=gpuSuccess) {
-    std::cout << "Lattice_slicesum_gpu.h: Encountered error during gpuMalloc Error: " << gpuErr <<std::endl;
+    std::cout << GridLogError << "Lattice_slicesum_gpu.h: Encountered error during gpuMalloc (helperArray)! Error: " << gpuErr <<std::endl;
+    exit(EXIT_FAILURE);
   }
 
   //prepare buffer for reduction
-  accelerator_for2dNB( s,e1*e2, r,rd, grid->Nsimd(),{ //use non-blocking accelerator_for to avoid syncs (ok because we submit to same computeStream)
+  //use non-blocking accelerator_for to avoid syncs (ok because we submit to same computeStream)
-                                                      //use 2d accelerator_for to avoid launch latencies found when looping over rd 
+  //use 2d accelerator_for to avoid launch latencies found when serially looping over rd 
+  
+  accelerator_for2dNB( s,subvol_size, r,rd, grid->Nsimd(),{ 
+  
     int n = s / e2;
     int b = s % e2;
     int so=r*ostride; // base offset for start of plane 
     int ss= so+n*stride+b;
 
-    coalescedWrite(rb_p[r*e1*e2+s], coalescedRead(Data_v[ss]));
+    coalescedWrite(rb_p[r*subvol_size+s], coalescedRead(Data_v[ss]));
 
   });
   
-  //issue reductions in computeStream
+  //issue segmented reductions in computeStream
-  gpuErr =gpucub::DeviceSegmentedReduce::Reduce(helperArray, temp_storage_bytes, rb_p, d_out, rd, d_offsets, d_offsets+1,::gpucub::Sum(), identity, computeStream);
+  gpuErr = gpucub::DeviceSegmentedReduce::Reduce(helperArray, temp_storage_bytes, rb_p, d_out, rd, d_offsets, d_offsets+1,::gpucub::Sum(), vobj_zero, computeStream);
   if (gpuErr!=gpuSuccess) {
-    std::cout << "Lattice_slicesum_gpu.h: Encountered error during cub::DeviceReduce::Sum(2)! Error: " << gpuErr <<std::endl;
+    std::cout << GridLogError << "Lattice_slicesum_gpu.h: Encountered error during gpucub::DeviceSegmentedReduce::Reduce! Error: " << gpuErr <<std::endl;
+    exit(EXIT_FAILURE);
   }
   
-  //sync before copy
+  gpuErr = gpuMemcpyDtoHAsync(&lvSum[0],d_out,rd*sizeof(vobj),computeStream);
-  accelerator_barrier();
-  gpuErr = gpuMemcpy(&lvSum[0],d_out,rd*sizeof(vobj),gpuMemcpyDeviceToHost);
   if (gpuErr!=gpuSuccess) {
-    std::cout << "Lattice_slicesum_gpu.h: Encountered error during gpuMemcpy(2) Error: " << gpuErr <<std::endl;
+    std::cout << GridLogError << "Lattice_slicesum_gpu.h: Encountered error during gpuMemcpy (d_out)! Error: " << gpuErr <<std::endl;
+    exit(EXIT_FAILURE);
   }
+
+  //sync after copy
+  accelerator_barrier();
+ 
   // Sum across simd lanes in the plane, breaking out orthog dir.
   Coordinate icoor(Nd);
 

tests/core/Test_sliceSum.cc

@@ -29,29 +29,42 @@ int main (int argc, char ** argv) {
       sliceSumGpu(test_data,reduction_result,0);
     }
 
-
+    int trace_id = traceStart("sliceSum benchmark");
     for (int i = 0; i < Nd; i++) {
-        RealD t=-usecond();
-        sliceSum(test_data,reduction_reference,i);
-        t+=usecond();
-        std::cout << " sliceSum took "<<t<<" usecs"<<std::endl;
-        
-        RealD tgpu=-usecond();
-        tracePush("sliceSumGpu");
-        sliceSumGpu(test_data,reduction_result,i);
-        tracePop("sliceSumGpu");
-        tgpu+=usecond();
-        std::cout <<" sliceSumGpu took "<<tgpu<<" usecs"<<std::endl;
 
-    for(int t=0;t<reduction_reference.size();t++){
+      RealD t=-usecond();
-      auto diff = reduction_reference[t]-reduction_result[t];
-      // std::cout << "Difference = " << diff <<std::endl;
 
-      assert(abs(TensorRemove(diff)) < 1e-8 );
+      tracePush("sliceSum");
-    }
+      sliceSum(test_data,reduction_reference,i);
+      tracePop("sliceSum");
+
+      t+=usecond();
+
+      std::cout << GridLogMessage << " sliceSum took "<<t<<" usecs"<<std::endl;
+      
+      
+      RealD tgpu=-usecond();
+
+      tracePush("sliceSumGpu");
+      sliceSumGpu(test_data,reduction_result,i);
+      tracePop("sliceSumGpu");
+
+      tgpu+=usecond();
+
+      std::cout << GridLogMessage <<" sliceSumGpu took "<<tgpu<<" usecs"<<std::endl;
+
+
+      for(int t=0;t<reduction_reference.size();t++) {
+
+        auto diff = reduction_reference[t]-reduction_result[t];
+        assert(abs(TensorRemove(diff)) < 1e-8 );
+
+      }
 
     
     }
+    traceStop(trace_id);
+    
     Grid_finalize();
     return 0;
 }