Increased the allowed size of custom expressions. Also fixed one failing test case.

platforms/cuda/src/kernels/cudaKernels.h

@@ -111,12 +111,12 @@ extern void SetBrownianUpdateSim(gpuContext gpu);
 extern void GetBrownianUpdateSim(gpuContext gpu);
 extern void SetRandomSim(gpuContext gpu);
 extern void GetRandomSim(gpuContext gpu);
-extern void SetCustomBondForceExpression(const Expression<128>& expression);
-extern void SetCustomBondEnergyExpression(const Expression<128>& expression);
+extern void SetCustomBondForceExpression(const Expression<256>& expression);
+extern void SetCustomBondEnergyExpression(const Expression<256>& expression);
 extern void SetCustomBondGlobalParams(const std::vector<float>&  paramValues);
-extern void SetCustomExternalForceExpressions(const Expression<128>& expressionX, const Expression<128>& expressionY, const Expression<128>& expressionZ);
-extern void SetCustomExternalEnergyExpression(const Expression<128>& expression);
+extern void SetCustomExternalForceExpressions(const Expression<256>& expressionX, const Expression<256>& expressionY, const Expression<256>& expressionZ);
+extern void SetCustomExternalEnergyExpression(const Expression<256>& expression);
 extern void SetCustomExternalGlobalParams(const std::vector<float>& paramValues);
-extern void SetCustomNonbondedForceExpression(const Expression<128>& expression);
-extern void SetCustomNonbondedEnergyExpression(const Expression<128>& expression);
+extern void SetCustomNonbondedForceExpression(const Expression<256>& expression);
+extern void SetCustomNonbondedEnergyExpression(const Expression<256>& expression);
 extern void SetCustomNonbondedGlobalParams(const std::vector<float>& paramValues);

platforms/cuda/src/kernels/gpu.cpp

@@ -691,8 +691,8 @@ void gpuSetCustomBondParameters(gpuContext gpu, const vector<int>& bondAtom1, co
     variables.push_back("r");
     for (int i = 0; i < (int) paramNames.size(); i++)
         variables.push_back(paramNames[i]);
-    SetCustomBondEnergyExpression(createExpression<128>(gpu, energyExp, Lepton::Parser::parse(energyExp).optimize().createProgram(), variables, globalParamNames, gpu->sim.customExpressionStackSize));
-    SetCustomBondForceExpression(createExpression<128>(gpu, energyExp, Lepton::Parser::parse(energyExp).differentiate("r").optimize().createProgram(), variables, globalParamNames, gpu->sim.customExpressionStackSize));
+    SetCustomBondEnergyExpression(createExpression<256>(gpu, energyExp, Lepton::Parser::parse(energyExp).optimize().createProgram(), variables, globalParamNames, gpu->sim.customExpressionStackSize));
+    SetCustomBondForceExpression(createExpression<256>(gpu, energyExp, Lepton::Parser::parse(energyExp).differentiate("r").optimize().createProgram(), variables, globalParamNames, gpu->sim.customExpressionStackSize));
 }
 
 extern "C"
@@ -733,10 +733,10 @@ void gpuSetCustomExternalParameters(gpuContext gpu, const vector<int>& atomIndex
     variables.push_back("z");
     for (int i = 0; i < (int) paramNames.size(); i++)
         variables.push_back(paramNames[i]);
-    SetCustomExternalEnergyExpression(createExpression<128>(gpu, energyExp, Lepton::Parser::parse(energyExp).optimize().createProgram(), variables, globalParamNames, gpu->sim.customExpressionStackSize));
-    SetCustomExternalForceExpressions(createExpression<128>(gpu, energyExp, Lepton::Parser::parse(energyExp).differentiate("x").optimize().createProgram(), variables, globalParamNames, gpu->sim.customExpressionStackSize),
-                                  createExpression<128>(gpu, energyExp, Lepton::Parser::parse(energyExp).differentiate("y").optimize().createProgram(), variables, globalParamNames, gpu->sim.customExpressionStackSize),
-                                  createExpression<128>(gpu, energyExp, Lepton::Parser::parse(energyExp).differentiate("z").optimize().createProgram(), variables, globalParamNames, gpu->sim.customExpressionStackSize));
+    SetCustomExternalEnergyExpression(createExpression<256>(gpu, energyExp, Lepton::Parser::parse(energyExp).optimize().createProgram(), variables, globalParamNames, gpu->sim.customExpressionStackSize));
+    SetCustomExternalForceExpressions(createExpression<256>(gpu, energyExp, Lepton::Parser::parse(energyExp).differentiate("x").optimize().createProgram(), variables, globalParamNames, gpu->sim.customExpressionStackSize),
+                                  createExpression<256>(gpu, energyExp, Lepton::Parser::parse(energyExp).differentiate("y").optimize().createProgram(), variables, globalParamNames, gpu->sim.customExpressionStackSize),
+                                  createExpression<256>(gpu, energyExp, Lepton::Parser::parse(energyExp).differentiate("z").optimize().createProgram(), variables, globalParamNames, gpu->sim.customExpressionStackSize));
 }
 
 extern "C"
@@ -815,8 +815,8 @@ void gpuSetCustomNonbondedParameters(gpuContext gpu, const vector<vector<double>
             variables.push_back("");
     }
     variables.push_back("r");
-    SetCustomNonbondedEnergyExpression(createExpression<128>(gpu, energyExp, Lepton::Parser::parse(energyExp, functions).optimize().createProgram(), variables, globalParamNames, gpu->sim.customExpressionStackSize));
-    SetCustomNonbondedForceExpression(createExpression<128>(gpu, energyExp, Lepton::Parser::parse(energyExp, functions).differentiate("r").optimize().createProgram(), variables, globalParamNames, gpu->sim.customExpressionStackSize));
+    SetCustomNonbondedEnergyExpression(createExpression<256>(gpu, energyExp, Lepton::Parser::parse(energyExp, functions).optimize().createProgram(), variables, globalParamNames, gpu->sim.customExpressionStackSize));
+    SetCustomNonbondedForceExpression(createExpression<256>(gpu, energyExp, Lepton::Parser::parse(energyExp, functions).differentiate("r").optimize().createProgram(), variables, globalParamNames, gpu->sim.customExpressionStackSize));
     delete fp;
 }
 
@@ -1843,6 +1843,7 @@ void* gpuInit(int numAtoms, unsigned int device, bool useBlockingSync)
     gpu->psShakeParameter           = NULL;
     gpu->psSettleID                 = NULL;
     gpu->psSettleParameter          = NULL;
+    gpu->psNonShakeID               = NULL;
     gpu->psExclusion                = NULL;
     gpu->psExclusionIndex           = NULL;
     gpu->psWorkUnit                 = NULL;
@@ -2044,6 +2045,7 @@ void gpuShutDown(gpuContext gpu)
     delete gpu->psShakeParameter;
     delete gpu->psSettleID;
     delete gpu->psSettleParameter;
+    if (gpu->psNonShakeID != NULL)
         delete gpu->psNonShakeID;
     delete gpu->psExclusion;
     delete gpu->psExclusionIndex;

platforms/cuda/src/kernels/kCalculateCustomBondForces.cu

@@ -37,8 +37,8 @@ using namespace std;
 #include "cudatypes.h"
 
 static __constant__ cudaGmxSimulation cSim;
-static __constant__ Expression<128> forceExp;
-static __constant__ Expression<128> energyExp;
+static __constant__ Expression<256> forceExp;
+static __constant__ Expression<256> energyExp;
 
 #include "kEvaluateExpression.h"
 
@@ -56,14 +56,14 @@ void GetCalculateCustomBondForcesSim(gpuContext gpu)
     RTERROR(status, "cudaMemcpyFromSymbol: SetSim copy from cSim failed");
 }
 
-void SetCustomBondForceExpression(const Expression<128>& expression)
+void SetCustomBondForceExpression(const Expression<256>& expression)
 {
     cudaError_t status;
     status = cudaMemcpyToSymbol(forceExp, &expression, sizeof(forceExp));
     RTERROR(status, "SetCustomBondForceExpression: cudaMemcpyToSymbol failed");
 }
 
-void SetCustomBondEnergyExpression(const Expression<128>& expression)
+void SetCustomBondEnergyExpression(const Expression<256>& expression)
 {
     cudaError_t status;
     status = cudaMemcpyToSymbol(energyExp, &expression, sizeof(energyExp));

platforms/cuda/src/kernels/kCalculateCustomExternalForces.cu

@@ -37,10 +37,10 @@ using namespace std;
 #include "cudatypes.h"
 
 static __constant__ cudaGmxSimulation cSim;
-static __constant__ Expression<128> forceExpX;
-static __constant__ Expression<128> forceExpY;
-static __constant__ Expression<128> forceExpZ;
-static __constant__ Expression<128> energyExp;
+static __constant__ Expression<256> forceExpX;
+static __constant__ Expression<256> forceExpY;
+static __constant__ Expression<256> forceExpZ;
+static __constant__ Expression<256> energyExp;
 
 #include "kEvaluateExpression.h"
 
@@ -58,7 +58,7 @@ void GetCalculateCustomExternalForcesSim(gpuContext gpu)
     RTERROR(status, "cudaMemcpyFromSymbol: SetSim copy from cSim failed");
 }
 
-void SetCustomExternalForceExpressions(const Expression<128>& expressionX, const Expression<128>& expressionY, const Expression<128>& expressionZ)
+void SetCustomExternalForceExpressions(const Expression<256>& expressionX, const Expression<256>& expressionY, const Expression<256>& expressionZ)
 {
     cudaError_t status;
     status = cudaMemcpyToSymbol(forceExpX, &expressionX, sizeof(forceExpX));
@@ -67,7 +67,7 @@ void SetCustomExternalForceExpressions(const Expression<128>& expressionX, const
     RTERROR(status, "SetCustomExternalForceExpression: cudaMemcpyToSymbol failed");
 }
 
-void SetCustomExternalEnergyExpression(const Expression<128>& expression)
+void SetCustomExternalEnergyExpression(const Expression<256>& expression)
 {
     cudaError_t status;
     status = cudaMemcpyToSymbol(energyExp, &expression, sizeof(energyExp));

platforms/cuda/src/kernels/kCalculateCustomNonbondedForces.cu

@@ -50,8 +50,8 @@ struct Atom {
 };
 
 static __constant__ cudaGmxSimulation cSim;
-static __constant__ Expression<128> forceExp;
-static __constant__ Expression<128> energyExp;
+static __constant__ Expression<256> forceExp;
+static __constant__ Expression<256> energyExp;
 
 #include "kEvaluateExpression.h"
 
@@ -69,14 +69,14 @@ void GetCalculateCustomNonbondedForcesSim(gpuContext gpu)
     RTERROR(status, "cudaMemcpyFromSymbol: SetSim copy from cSim failed");
 }
 
-void SetCustomNonbondedForceExpression(const Expression<128>& expression)
+void SetCustomNonbondedForceExpression(const Expression<256>& expression)
 {
     cudaError_t status;
     status = cudaMemcpyToSymbol(forceExp, &expression, sizeof(forceExp));
     RTERROR(status, "SetCustomNonbondedForceExpression: cudaMemcpyToSymbol failed");
 }
 
-void SetCustomNonbondedEnergyExpression(const Expression<128>& expression)
+void SetCustomNonbondedEnergyExpression(const Expression<256>& expression)
 {
     cudaError_t status;
     status = cudaMemcpyToSymbol(energyExp, &expression, sizeof(energyExp));