Reorder cuda initialization of periodic BCs until after cuda initialization of component

Author: Martin D. Weinberg

src/Cube.cc

@@ -388,7 +388,7 @@ void * Cube::determine_acceleration_and_potential_thread(void * arg)
     auto stepy = std::exp(kfac*y);
     auto stepz = std::exp(kfac*z);
 
-    // Initial values (note sign change)
+    // Initial values (note sign change from coefficient accumulation)
     auto startx = std::exp(-kfac*(x*nmaxx));
     auto starty = std::exp(-kfac*(y*nmaxy));
     auto startz = std::exp(-kfac*(z*nmaxz));

src/PeriodicBC.H

@@ -28,7 +28,7 @@ private:
 #if HAVE_LIBCUDA==1
   //! Cuda implementation
   void determine_acceleration_and_potential_cuda();
-  void cuda_initialize();
+  bool cuda_initialized = false;
 #endif
 
   std::vector<double> offset, L;

src/PeriodicBC.cc

@@ -145,14 +145,6 @@ void PeriodicBC::initialize()
 			   << std::string(60, '-') << std::endl;
     throw std::runtime_error("PeriodicBC::initialize: error parsing YAML");
   }
-    
-
-  // Cuda initialization
-#if HAVE_LIBCUDA==1
-  if (use_cuda) {
-    cuda_initialize();
-  }
-#endif
 }
 
 

src/cudaPeriodicBC.cu

@@ -104,43 +104,45 @@ void testConstantsPeriodicBC(cuFP_t tnow)
 	 cudaBCside[0], cudaBCside[1], cudaBCside[2] );
   printf("Offset = %e, %e, %e\n",
 	 cudaBCoffset[0], cudaBCoffset[1], cudaBCoffset[2] );
-  printf(" BCs   = %c, %c, %c\n", cudaBC[0], cudaBC[1], cudaBC[2] );
+  printf("   BCs = %c, %c, %c\n", cudaBC[0], cudaBC[1], cudaBC[2] );
   printf("-------------------------\n");
 }
 
-void PeriodicBC::cuda_initialize()
+
+void PeriodicBC::determine_acceleration_and_potential_cuda()
 {
-  cuFP_t vec[3];
-  for (int k=0; k<3; k++) vec[k] = offset[k];
+  if (not cuda_initialized) {
+    cuda_initialized = true;
 
-  cuda_safe_call(cudaMemcpyToSymbol(cudaBCoffset, vec, sizeof(cuFP_t)*3,
-				    size_t(0), cudaMemcpyHostToDevice),
-		 __FILE__, __LINE__, "Error copying cudaBCoffset");
-  
-  for (int k=0; k<3; k++) vec[k] = L[k];
+    // Copy the offset, side length, and BC type to device constant memory
+    cuFP_t vec[3];
+    for (int k=0; k<3; k++) vec[k] = offset[k];
 
-  cuda_safe_call(cudaMemcpyToSymbol(cudaBCside, vec, sizeof(cuFP_t)*3,
-				    size_t(0), cudaMemcpyHostToDevice),
-		 __FILE__, __LINE__, "Error copying cudaBCside");
+    cuda_safe_call(cudaMemcpyToSymbol(cudaBCoffset, vec, sizeof(cuFP_t)*3,
+				      size_t(0), cudaMemcpyHostToDevice),
+		   __FILE__, __LINE__, "Error copying cudaBCoffset");
 
-  char h_cudaBC[3];
-  for (int k=0; k<3; k++) h_cudaBC[k] = bc[k];
-
-  cuda_safe_call(cudaMemcpyToSymbol(cudaBC, h_cudaBC, sizeof(char)*3,
-				    size_t(0), cudaMemcpyHostToDevice),
-		 __FILE__, __LINE__, "Error copying cudaBC");
+    for (int k=0; k<3; k++) vec[k] = L[k];
+    
+    cuda_safe_call(cudaMemcpyToSymbol(cudaBCside, vec, sizeof(cuFP_t)*3,
+				      size_t(0), cudaMemcpyHostToDevice),
+		   __FILE__, __LINE__, "Error copying cudaBCside");
+  
+    char h_cudaBC[3];
+    for (int k=0; k<3; k++) h_cudaBC[k] = bc[k];
+    
+    cuda_safe_call(cudaMemcpyToSymbol(cudaBC, h_cudaBC, sizeof(char)*3,
+				      size_t(0), cudaMemcpyHostToDevice),
+		   __FILE__, __LINE__, "Error copying cudaBC");
 
-  if (myid==0 and VERBOSE>4) {
-    auto cr = cC->cuStream;
-    testConstantsPeriodicBC<<<1, 1, 0, cr->stream>>>(tnow);
-    cudaDeviceSynchronize();
-    cuda_check_last_error_mpi("cudaDeviceSynchronize", __FILE__, __LINE__, myid);
+    if (myid==0 and VERBOSE>4) {
+      auto cr = cC->cuStream;
+      testConstantsPeriodicBC<<<1, 1, 0, cr->stream>>>(tnow);
+      cudaDeviceSynchronize();
+      cuda_check_last_error_mpi("cudaDeviceSynchronize", __FILE__, __LINE__, myid);
+    }
   }
-}
 
-
-void PeriodicBC::determine_acceleration_and_potential_cuda()
-{
   // Sanity check
   //
   int nbodies = cC->Number();