Convert #defines to const variables; change from #pragma once to #ifd…

…ef for single inclusion
GOMC-WSU · Sep 3, 2024 · f3ff9ba · f3ff9ba
1 parent 9f7a9c0
commit f3ff9ba
Show file tree

Hide file tree

Showing 11 changed files with 47 additions and 29 deletions.
diff --git a/src/GPU/CUDAMemoryManager.cu b/src/GPU/CUDAMemoryManager.cu
@@ -5,9 +5,10 @@ A copy of the MIT License can be found in License.txt
 along with this program, also can be found at
 <https://opensource.org/licenses/MIT>.
 ********************************************************************************/
+#ifdef GOMC_CUDA
+
 #include "CUDAMemoryManager.cuh"
 
-#ifdef GOMC_CUDA
 long long CUDAMemoryManager::totalAllocatedBytes = 0;
 std::unordered_map<void *, std::pair<unsigned int, std::string>>
     CUDAMemoryManager::allocatedPointers;

diff --git a/src/GPU/CUDAMemoryManager.cuh b/src/GPU/CUDAMemoryManager.cuh
@@ -5,7 +5,9 @@ A copy of the MIT License can be found in License.txt
 along with this program, also can be found at
 <https://opensource.org/licenses/MIT>.
 ********************************************************************************/
-#pragma once
+#ifndef CUDA_MEMORY_MANAGER
+#define CUDA_MEMORY_MANAGER
+
 #ifdef GOMC_CUDA
 #include <cuda.h>
 #include <cuda_runtime.h>
@@ -29,4 +31,5 @@ private:
       allocatedPointers;
 };
 
-#endif
+#endif /*GOMC_CUDA*/
+#endif /*CUDA_MEMORY_MANAGER*/
diff --git a/src/GPU/CalculateEnergyCUDAKernel.cu b/src/GPU/CalculateEnergyCUDAKernel.cu
@@ -17,8 +17,8 @@ along with this program, also can be found at
 #include "ConstantDefinitionsCUDAKernel.cuh"
 #include "cub/cub.cuh"
 
-#define NUMBER_OF_NEIGHBOR_CELL 27
-#define THREADS_PER_BLOCK 128
+const int NUMBER_OF_NEIGHBOR_CELL = 27;
+const int THREADS_PER_BLOCK = 128;
 
 using namespace cub;
 

diff --git a/src/GPU/CalculateEnergyCUDAKernel.cuh b/src/GPU/CalculateEnergyCUDAKernel.cuh
@@ -5,7 +5,9 @@ A copy of the MIT License can be found in License.txt
 along with this program, also can be found at
 <https://opensource.org/licenses/MIT>.
 ********************************************************************************/
-#pragma once
+#ifndef CALCULATE_ENERGY_CUDA_KERNEL
+#define CALCULATE_ENERGY_CUDA_KERNEL
+
 #ifdef GOMC_CUDA
 #include "BoxDimensions.h"
 #include "VariablesCUDA.cuh"
@@ -155,3 +157,4 @@ __device__ double CalcEnSwitchGPUNoLambda(double distSq, int index,
                                           double gpu_rCut, double gpu_rOn);
 
 #endif /*GOMC_CUDA*/
+#endif /*CALCULATE_ENERGY_CUDA_KERNEL*/
diff --git a/src/GPU/CalculateEwaldCUDAKernel.cu b/src/GPU/CalculateEwaldCUDAKernel.cu
@@ -21,7 +21,8 @@ along with this program, also can be found at
 
 using namespace cub;
 
-#define THREADS_PER_BLOCK 128
+const int THREADS_PER_BLOCK = 128;
+const int THREADS_PER_BLOCK_SM = 64;
 
 // Use this function when calculating the reciprocal terms
 // for a new volume. Such as a change in box dimensions.
@@ -55,7 +56,7 @@ void CallBoxReciprocalSetupGPU(VariablesCUDA *vars, XYZArray const &coords,
   checkLastErrorCUDA(__FILE__, __LINE__);
 #endif
 
-  int threadsPerBlock = THREADS_PER_BLOCK;
+  int threadsPerBlock = THREADS_PER_BLOCK_SM;
   int blocksPerGrid = imageSize;
   BoxReciprocalSumsGPU<<<blocksPerGrid, threadsPerBlock>>>(
       vars->gpu_x, vars->gpu_y, vars->gpu_z, vars->gpu_kx[box],
@@ -67,6 +68,7 @@ void CallBoxReciprocalSetupGPU(VariablesCUDA *vars, XYZArray const &coords,
 #endif
 
   // Fewer blocks are needed since we are doing one computation per image
+  threadsPerBlock = THREADS_PER_BLOCK;
   blocksPerGrid = (imageSize + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
   BoxReciprocalGPU<<<blocksPerGrid, threadsPerBlock>>>(
       vars->gpu_prefact[box], vars->gpu_sumRnew[box], vars->gpu_sumInew[box],
@@ -102,7 +104,7 @@ void CallBoxReciprocalSumsGPU(VariablesCUDA *vars, XYZArray const &coords,
   checkLastErrorCUDA(__FILE__, __LINE__);
 #endif
 
-  int threadsPerBlock = THREADS_PER_BLOCK;
+  int threadsPerBlock = THREADS_PER_BLOCK_SM;
   int blocksPerGrid = imageSize;
   BoxReciprocalSumsGPU<<<blocksPerGrid, threadsPerBlock>>>(
       vars->gpu_x, vars->gpu_y, vars->gpu_z, vars->gpu_kxRef[box],
@@ -114,6 +116,7 @@ void CallBoxReciprocalSumsGPU(VariablesCUDA *vars, XYZArray const &coords,
 #endif
 
   // Fewer blocks are needed since we are doing one computation per image
+  threadsPerBlock = THREADS_PER_BLOCK;
   blocksPerGrid = (imageSize + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
   BoxReciprocalGPU<<<blocksPerGrid, threadsPerBlock>>>(
       vars->gpu_prefactRef[box], vars->gpu_sumRnew[box], vars->gpu_sumInew[box],
@@ -138,7 +141,7 @@ __global__ void BoxReciprocalSumsGPU(double *gpu_x, double *gpu_y,
   int image = blockIdx.x;
   double sumR = 0.0, sumI = 0.0;
 #pragma unroll 16
-  for (int particleID = threadIdx.x; particleID < atomNumber; particleID += THREADS_PER_BLOCK) {
+  for (int particleID = threadIdx.x; particleID < atomNumber; particleID += THREADS_PER_BLOCK_SM) {
     double dot = DotProductGPU(gpu_kx[image], gpu_ky[image],
                                gpu_kz[image], gpu_x[particleID],
                                gpu_y[particleID],
@@ -151,7 +154,7 @@ __global__ void BoxReciprocalSumsGPU(double *gpu_x, double *gpu_y,
   __syncthreads();
 
   // Specialize BlockReduce for a 1D block of threads of type double
-  using BlockReduce = cub::BlockReduce<double, THREADS_PER_BLOCK>;
+  using BlockReduce = cub::BlockReduce<double, THREADS_PER_BLOCK_SM>;
 
   // Allocate shared memory for BlockReduce
   __shared__ typename BlockReduce::TempStorage sumR_temp_storage;

diff --git a/src/GPU/CalculateForceCUDAKernel.cu b/src/GPU/CalculateForceCUDAKernel.cu
@@ -16,9 +16,9 @@ along with this program, also can be found at
 #include "ConstantDefinitionsCUDAKernel.cuh"
 #include "cub/cub.cuh"
 
-#define NUMBER_OF_NEIGHBOR_CELLS 27
-#define PARTICLES_PER_BLOCK 32
-#define THREADS_PER_BLOCK 128
+const int NUMBER_OF_NEIGHBOR_CELLS = 27;
+const int PARTICLES_PER_BLOCK = 32;
+const int THREADS_PER_BLOCK = 128;
 
 using namespace cub;
 

diff --git a/src/GPU/CalculateMinImageCUDAKernel.cuh b/src/GPU/CalculateMinImageCUDAKernel.cuh
@@ -5,9 +5,10 @@ A copy of the MIT License can be found in License.txt
 along with this program, also can be found at
 <https://opensource.org/licenses/MIT>.
 ********************************************************************************/
-#pragma once
-#ifdef GOMC_CUDA
+#ifndef CALCULATE_MINIMAGE_CUDA_KERNEL
+#define CALCULATE_MINIMAGE_CUDA_KERNEL
 
+#ifdef GOMC_CUDA
 #include "ConstantDefinitionsCUDAKernel.cuh"
 #include <cuda.h>
 #include <cuda_runtime.h>
@@ -276,6 +277,7 @@ static __inline__ __device__ double atomicAdd(double *address, double val) {
   } while (assumed != old);
   return __longlong_as_double(old);
 }
-#endif
+#endif /*CUDA_ARCH*/
 
 #endif /*GOMC_CUDA*/
+#endif /*CALCULATE_MINIMAGE_CUDA_KERNEL*/
diff --git a/src/GPU/ConstantDefinitionsCUDAKernel.cuh b/src/GPU/ConstantDefinitionsCUDAKernel.cuh
@@ -15,11 +15,11 @@ along with this program, also can be found at
 #include <cuda.h>
 #include <cuda_runtime.h>
 
-#define GPU_VDW_STD_KIND 0
-#define GPU_VDW_SHIFT_KIND 1
-#define GPU_VDW_SWITCH_KIND 2
-#define GPU_VDW_EXP6_KIND 3
-#define MAX_PAIR_SIZE 10000000
+const int GPU_VDW_STD_KIND = 0;
+const int GPU_VDW_SHIFT_KIND = 1;
+const int GPU_VDW_SWITCH_KIND = 2;
+const int GPU_VDW_EXP6_KIND = 3;
+const int MAX_PAIR_SIZE = 10000000;
 
 void UpdateGPULambda(VariablesCUDA *vars, int *molIndex, double *lambdaVDW,
                      double *lambdaCoulomb, bool *isFraction);

diff --git a/src/GPU/TransformParticlesCUDAKernel.cu b/src/GPU/TransformParticlesCUDAKernel.cu
@@ -11,8 +11,8 @@ along with this program, also can be found at
 #include "Random123/boxmuller.hpp"
 #include "TransformParticlesCUDAKernel.cuh"
 
-#define MIN_FORCE 1E-12
-#define MAX_FORCE 30
+const double MIN_FORCE = 1.0e-12;
+const double MAX_FORCE = 30.0;
 
 __device__ inline double randomGPU(unsigned int counter, ulong step,
                                    ulong seed) {

diff --git a/src/GPU/TransformParticlesCUDAKernel.cuh b/src/GPU/TransformParticlesCUDAKernel.cuh
@@ -5,7 +5,9 @@ A copy of the MIT License can be found in License.txt
 along with this program, also can be found at
 <https://opensource.org/licenses/MIT>.
 ********************************************************************************/
-#pragma once
+#ifndef TRANSFORM_PARTICLES_CUDA_KERNEL
+#define TRANSFORM_PARTICLES_CUDA_KERNEL
+
 #ifdef GOMC_CUDA
 #include "Random123/philox.h"
 #include <vector>
@@ -94,4 +96,5 @@ __global__ void BrownianMotionTranslateKernel(
     double3 axis, double3 halfAx, int atomCount, double t_max, ulong step,
     unsigned int key, ulong seed, double BETA);
 
-#endif
+#endif /*GOMC_CUDA*/
+#endif /*TRANSFORM_PARTICLES_CUDA_KERNEL*/
diff --git a/src/GPU/VariablesCUDA.cuh b/src/GPU/VariablesCUDA.cuh
@@ -5,9 +5,10 @@ A copy of the MIT License can be found in License.txt
 along with this program, also can be found at
 <https://opensource.org/licenses/MIT>.
 ********************************************************************************/
-#pragma once
-#ifdef GOMC_CUDA
+#ifndef VARIABLES_CUDA
+#define VARIABLES_CUDA
 
+#ifdef GOMC_CUDA
 #include "EnsemblePreprocessor.h"
 #include "NumLib.h"
 #include <cuda.h>
@@ -160,4 +161,6 @@ public:
   // new pair interaction calculation done on GPU
   int *gpu_cellVector, *gpu_mapParticleToCell;
 };
-#endif
+
+#endif /*GOMC_CUDA*/
+#endif /*VARIABLES_CUDA*/