Mods for Windows compile

Renamed files to be consistent w/ other file names

plugins/amoeba/platforms/cuda/sharedTarget/CMakeLists.txt

@@ -64,7 +64,8 @@ CUDA_INCLUDE_DIRECTORIES(${OPENMM_BUILD_AMOEBA_PATH}/platforms/cuda/../src
                          ${OPENMM_DIR}/platforms/cuda/src
                          ${OPENMM_DIR}/platforms/cuda/include
                          ${OPENMM_DIR}/platforms/cuda/src/kernels
-                         ${OPENMM_DIR}/openmmapi/include ) 
+                         ${OPENMM_DIR}/openmmapi/include 
+                         ${OPENMM_DIR}/olla/include ) 
 
 CUDA_ADD_LIBRARY(${SHARED_TARGET} SHARED ${SOURCE_FILES} ${SOURCE_INCLUDE_FILES} ${API_ABS_INCLUDE_FILES})
 

plugins/amoeba/platforms/cuda/src/kernels/AmoebaGpu.cpp → plugins/amoeba/platforms/cuda/src/kernels/amoebaCudaGpu.cpp

@@ -2749,6 +2749,7 @@ void amoebaGpuSetConstants(amoebaGpuContext amoebaGpu)
     SetCalculateAmoebaPmeDirectElectrostaticSim( amoebaGpu );
     SetCalculateAmoebaCudaMapTorquesSim( amoebaGpu );
     SetCalculateAmoebaKirkwoodSim( amoebaGpu );
+    SetCalculateAmoebaCudaUtilitiesSim( amoebaGpu );
     SetCalculateAmoebaKirkwoodEDiffSim( amoebaGpu );
     SetCalculateAmoebaCudaFixedEAndGKFieldsSim( amoebaGpu );
     SetCalculateAmoebaCudaMutualInducedAndGkFieldsSim( amoebaGpu );

plugins/amoeba/platforms/cuda/src/kernels/amoebaCudaKernels.h

@@ -168,5 +168,7 @@ extern void kCalculateAmoebaPMEFixedMultipoles(amoebaGpuContext amoebaGpu);
 extern void kCalculateAmoebaPMEInducedDipoleField(amoebaGpuContext amoebaGpu);
 extern void kCalculateAmoebaPMEInducedDipoleForces(amoebaGpuContext amoebaGpu);
 
+
+extern void SetCalculateAmoebaCudaUtilitiesSim( amoebaGpuContext amoebaGpu );
 #endif //__AMOEBA_GPU_TYPES_H__
 

plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaRotateFrame.cu → plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaRotateFrame.cu

@@ -4,6 +4,7 @@
 
 #include "cudaKernels.h"
 #include "amoebaCudaKernels.h"
+#include "kCalculateAmoebaCudaUtilities.h"
 
 #include <stdio.h>
 #include <cuda.h>
@@ -49,6 +50,8 @@ __device__ static float normVector3( float* vector )
     return returnNorm;
 }
 
+#undef AMOEBA_DEBUG
+
 __global__
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_THREADS_PER_BLOCK, 1)

plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaUtilities.cu

@@ -6,6 +6,36 @@
 #include "amoebaCudaKernels.h"
 //#define AMOEBA_DEBUG
 
+static __constant__ cudaGmxSimulation cSim;
+static __constant__ cudaAmoebaGmxSimulation cAmoebaSim;
+
+void SetCalculateAmoebaCudaUtilitiesSim(amoebaGpuContext amoebaGpu)
+{
+    cudaError_t status;
+    gpuContext gpu = amoebaGpu->gpuContext;
+    status         = cudaMemcpyToSymbol(cSim, &gpu->sim, sizeof(cudaGmxSimulation));    
+    RTERROR(status, "SetCalculateAmoebaCudaUtilitiesSim: cudaMemcpyToSymbol: SetSim copy to cSim failed");
+    status         = cudaMemcpyToSymbol(cAmoebaSim, &amoebaGpu->amoebaSim, sizeof(cudaAmoebaGmxSimulation));    
+    RTERROR(status, "SetCalculateAmoebaCudaUtilitiesSim: cudaMemcpyToSymbol: SetSim copy to cAmoebaSim failed");
+}
+
+void GetCalculateAmoebaCudaUtilitiesSim(amoebaGpuContext amoebaGpu)
+{
+    cudaError_t status;
+    gpuContext gpu = amoebaGpu->gpuContext;
+    status = cudaMemcpyFromSymbol(&gpu->sim, cSim, sizeof(cudaGmxSimulation));    
+    RTERROR(status, "GetCalculateAmoebaCudaUtilitiesSim: cudaMemcpyFromSymbol: SetSim copy from cSim failed");
+    status = cudaMemcpyFromSymbol(&amoebaGpu->amoebaSim, cAmoebaSim, sizeof(cudaAmoebaGmxSimulation));    
+    RTERROR(status, "GetCalculateAmoebaCudaUtilitiesSim: cudaMemcpyFromSymbol: SetSim copy from cAmoebaSim failed");
+}
+
+#undef METHOD_NAME
+#define USE_PERIODIC
+#define METHOD_NAME(a, b) a##Periodic##b
+#include "kFindInteractingBlocks.h"
+#undef METHOD_NAME
+#undef USE_PERIODIC
+
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_THREADS_PER_BLOCK, 1)

plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaUtilities.h

@@ -7,4 +7,9 @@ __global__ void kReduceFields_kernel( unsigned int fieldComponents, unsigned int
 __global__ void kReduceAndCombineFields_kernel( unsigned int fieldComponents, unsigned int outputBuffers, float* fieldIn1, float* fieldIn2, float* fieldOut );
 __global__ void kReduceFieldsToFloat4_kernel( unsigned int fieldComponents, unsigned int outputBuffers, float* fieldIn, float4* fieldOut );
 
+extern __global__ void kFindBlockBoundsPeriodic_kernel();
+extern __global__ void kFindBlocksWithInteractionsPeriodic_kernel();
+extern __global__ void kFindInteractionsWithinBlocksPeriodic_kernel(unsigned int*);
+
+
 #endif