From f3ff9baba8618b3e87f06d435ff954b965da4dd6 Mon Sep 17 00:00:00 2001 From: Loren Schwiebert Date: Tue, 3 Sep 2024 12:03:55 -0400 Subject: [PATCH] Convert #defines to const variables; change from #pragma once to #ifdef for single inclusion --- src/GPU/CUDAMemoryManager.cu | 3 ++- src/GPU/CUDAMemoryManager.cuh | 7 +++++-- src/GPU/CalculateEnergyCUDAKernel.cu | 4 ++-- src/GPU/CalculateEnergyCUDAKernel.cuh | 5 ++++- src/GPU/CalculateEwaldCUDAKernel.cu | 13 ++++++++----- src/GPU/CalculateForceCUDAKernel.cu | 6 +++--- src/GPU/CalculateMinImageCUDAKernel.cuh | 8 +++++--- src/GPU/ConstantDefinitionsCUDAKernel.cuh | 10 +++++----- src/GPU/TransformParticlesCUDAKernel.cu | 4 ++-- src/GPU/TransformParticlesCUDAKernel.cuh | 7 +++++-- src/GPU/VariablesCUDA.cuh | 9 ++++++--- 11 files changed, 47 insertions(+), 29 deletions(-) diff --git a/src/GPU/CUDAMemoryManager.cu b/src/GPU/CUDAMemoryManager.cu index d0c7be074..d9f00b6e5 100644 --- 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 . ********************************************************************************/ +#ifdef GOMC_CUDA + #include "CUDAMemoryManager.cuh" -#ifdef GOMC_CUDA long long CUDAMemoryManager::totalAllocatedBytes = 0; std::unordered_map> CUDAMemoryManager::allocatedPointers; diff --git a/src/GPU/CUDAMemoryManager.cuh b/src/GPU/CUDAMemoryManager.cuh index 06e476a9d..08d428c73 100644 --- 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 . ********************************************************************************/ -#pragma once +#ifndef CUDA_MEMORY_MANAGER +#define CUDA_MEMORY_MANAGER + #ifdef GOMC_CUDA #include #include @@ -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 index 221710a39..84750b946 100644 --- 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 index 314148e61..c7dbe2c63 100644 --- 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 . ********************************************************************************/ -#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 index 580a544ef..00d6273c4 100644 --- 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<<>>( 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<<>>( 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<<>>( 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<<>>( 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; + using BlockReduce = cub::BlockReduce; // Allocate shared memory for BlockReduce __shared__ typename BlockReduce::TempStorage sumR_temp_storage; diff --git a/src/GPU/CalculateForceCUDAKernel.cu b/src/GPU/CalculateForceCUDAKernel.cu index 219fd7c06..b565a5393 100644 --- 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 index b8209ff89..a014485e6 100644 --- 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 . ********************************************************************************/ -#pragma once -#ifdef GOMC_CUDA +#ifndef CALCULATE_MINIMAGE_CUDA_KERNEL +#define CALCULATE_MINIMAGE_CUDA_KERNEL +#ifdef GOMC_CUDA #include "ConstantDefinitionsCUDAKernel.cuh" #include #include @@ -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 index 55db277f0..31469eff8 100644 --- a/src/GPU/ConstantDefinitionsCUDAKernel.cuh +++ b/src/GPU/ConstantDefinitionsCUDAKernel.cuh @@ -15,11 +15,11 @@ along with this program, also can be found at #include #include -#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 index 8830a0b06..1b3360fd2 100644 --- 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 index 74a67c3eb..a205d0edd 100644 --- 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 . ********************************************************************************/ -#pragma once +#ifndef TRANSFORM_PARTICLES_CUDA_KERNEL +#define TRANSFORM_PARTICLES_CUDA_KERNEL + #ifdef GOMC_CUDA #include "Random123/philox.h" #include @@ -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 index 2f7d04741..507d6da53 100644 --- 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 . ********************************************************************************/ -#pragma once -#ifdef GOMC_CUDA +#ifndef VARIABLES_CUDA +#define VARIABLES_CUDA +#ifdef GOMC_CUDA #include "EnsemblePreprocessor.h" #include "NumLib.h" #include @@ -160,4 +161,6 @@ public: // new pair interaction calculation done on GPU int *gpu_cellVector, *gpu_mapParticleToCell; }; -#endif + +#endif /*GOMC_CUDA*/ +#endif /*VARIABLES_CUDA*/