Continuing converting AMOEBA to the new CUDA platform:...

Continuing converting AMOEBA to the new CUDA platform: AmoebaOutOfPlaneBendForce, AmoebaStretchBendForce, AmoebaTorsionTorsionForce

plugins/amoeba/platforms/cuda2/src/AmoebaCudaKernelFactory.cpp

@@ -88,15 +88,15 @@ KernelImpl* AmoebaCudaKernelFactory::createKernelImpl(std::string name, const Pl
     if (name == CalcAmoebaPiTorsionForceKernel::Name())
         return new CudaCalcAmoebaPiTorsionForceKernel(name, platform, cu, context.getSystem());
 
-//    if (name == CalcAmoebaStretchBendForceKernel::Name())
-//        return new CudaCalcAmoebaStretchBendForceKernel(name, platform, cu, context.getSystem());
-//
-//    if (name == CalcAmoebaOutOfPlaneBendForceKernel::Name())
-//        return new CudaCalcAmoebaOutOfPlaneBendForceKernel(name, platform, cu, context.getSystem());
-//
-//    if (name == CalcAmoebaTorsionTorsionForceKernel::Name())
-//        return new CudaCalcAmoebaTorsionTorsionForceKernel(name, platform, cu, context.getSystem());
-//
+    if (name == CalcAmoebaStretchBendForceKernel::Name())
+        return new CudaCalcAmoebaStretchBendForceKernel(name, platform, cu, context.getSystem());
+
+    if (name == CalcAmoebaOutOfPlaneBendForceKernel::Name())
+        return new CudaCalcAmoebaOutOfPlaneBendForceKernel(name, platform, cu, context.getSystem());
+
+    if (name == CalcAmoebaTorsionTorsionForceKernel::Name())
+        return new CudaCalcAmoebaTorsionTorsionForceKernel(name, platform, cu, context.getSystem());
+
 //    if (name == CalcAmoebaMultipoleForceKernel::Name())
 //        return new CudaCalcAmoebaMultipoleForceKernel(name, platform, cu, context.getSystem());
 //

plugins/amoeba/platforms/cuda2/src/AmoebaCudaKernels.h

@@ -253,6 +253,7 @@ private:
     int numStretchBends;
     CudaContext& cu;
     System& system;
+    CudaArray* params;
 };
 
 /**
@@ -283,6 +284,7 @@ private:
     int numOutOfPlaneBends;
     CudaContext& cu;
     System& system;
+    CudaArray* params;
 };
 
 /**
@@ -314,6 +316,9 @@ private:
     int numTorsionTorsionGrids;
     CudaContext& cu;
     System& system;
+    CudaArray* gridValues;
+    CudaArray* gridParams;
+    CudaArray* torsionParams;
 };
 
 /**

plugins/amoeba/platforms/cuda2/src/kernels/amoebaOutOfPlaneBendForce.cu

+// compute the value of the bond angle
+
+real xab = pos1.x - pos2.x;
+real yab = pos1.y - pos2.y;
+real zab = pos1.z - pos2.z;
+
+real xcb = pos3.x - pos2.x;
+real ycb = pos3.y - pos2.y;
+real zcb = pos3.z - pos2.z;
+
+// compute the out-of-plane bending angle
+
+real xdb = pos4.x - pos2.x;
+real ydb = pos4.y - pos2.y;
+real zdb = pos4.z - pos2.z;
+
+real xad = pos1.x - pos4.x;
+real yad = pos1.y - pos4.y;
+real zad = pos1.z - pos4.z;
+
+real xcd = pos3.x - pos4.x;
+real ycd = pos3.y - pos4.y;
+real zcd = pos3.z - pos4.z;
+
+real rdb2 = xdb*xdb + ydb*ydb + zdb*zdb;
+real rad2 = xad*xad + yad*yad + zad*zad;
+real rcd2 = xcd*xcd + ycd*ycd + zcd*zcd;
+
+real ee = xab*(ycb*zdb-zcb*ydb) + yab*(zcb*xdb-xcb*zdb) + zab*(xcb*ydb-ycb*xdb);
+
+real dot = xad*xcd + yad*ycd + zad*zcd;
+real cc = rad2*rcd2 - dot*dot;
+real bkk2 = (cc != 0 ? (ee*ee)/(cc) : (real) 0);
+bkk2 = rdb2 - bkk2;
+
+real adXcd_0 = yad*zcd - zad*ycd;
+real adXcd_1 = zad*xcd - xad*zcd;
+real adXcd_2 = xad*ycd - yad*xcd;
+real adXcd_nrm2 = adXcd_0*adXcd_0 + adXcd_1*adXcd_1 + adXcd_2*adXcd_2;
+
+real adXcd_dot_db = xdb*adXcd_0 + ydb*adXcd_1 + zdb*adXcd_2;
+adXcd_dot_db /= SQRT(rdb2*adXcd_nrm2);
+
+real angle = abs(ASIN(adXcd_dot_db));
+
+// find the out-of-plane energy and master chain rule terms
+
+real dt = RAD_TO_DEG*angle;
+real dt2 = dt * dt;
+real dt3 = dt2 * dt;
+real dt4 = dt2 * dt2;
+float k = (rdb2 != 0 && cc != 0) ? PARAMS[index] : 0.0f;
+
+energy += k*dt2*(1.0f + CUBIC_K*dt + QUARTIC_K*dt2 + PENTIC_K*dt3 + SEXTIC_K*dt4);
+
+real deddt = k*dt*RAD_TO_DEG*(2.0f + 3.0f*CUBIC_K*dt + 4.0f*QUARTIC_K*dt2 + 5.0f*PENTIC_K*dt3 + 6.0f*SEXTIC_K*dt4);
+
+real eeSign = (ee >= 0 ? 1 : -1);
+real dedcos = -deddt*eeSign/SQRT(cc*bkk2);
+
+// chain rule terms for first derivative components
+
+real term = ee / cc;
+
+real dccdxia = (xad*rcd2-xcd*dot) * term;
+real dccdyia = (yad*rcd2-ycd*dot) * term;
+real dccdzia = (zad*rcd2-zcd*dot) * term;
+
+real dccdxic = (xcd*rad2-xad*dot) * term;
+real dccdyic = (ycd*rad2-yad*dot) * term;
+real dccdzic = (zcd*rad2-zad*dot) * term;
+
+real dccdxid = -dccdxia - dccdxic;
+real dccdyid = -dccdyia - dccdyic;
+real dccdzid = -dccdzia - dccdzic;
+
+term = ee / rdb2;
+
+real deedxia = ydb*zcb - zdb*ycb;
+real deedyia = zdb*xcb - xdb*zcb;
+real deedzia = xdb*ycb - ydb*xcb;
+
+real deedxic = yab*zdb - zab*ydb;
+real deedyic = zab*xdb - xab*zdb;
+real deedzic = xab*ydb - yab*xdb;
+
+real deedxid = ycb*zab - zcb*yab + xdb*term;
+real deedyid = zcb*xab - xcb*zab + ydb*term;
+real deedzid = xcb*yab - ycb*xab + zdb*term;
+
+// compute first derivative components for this angle
+
+real3 force1 = make_real3(-dedcos*(dccdxia+deedxia), -dedcos*(dccdyia+deedyia), -dedcos*(dccdzia+deedzia));
+real3 force3 = make_real3(-dedcos*(dccdxic+deedxic), -dedcos*(dccdyic+deedyic), -dedcos*(dccdzic+deedzic));
+real3 force4 = make_real3(-dedcos*(dccdxid+deedxid), -dedcos*(dccdyid+deedyid), -dedcos*(dccdzid+deedzid));
+real3 force2 = make_real3(-force1.x-force3.x-force4.x, -force1.y-force3.y-force4.y, -force1.z-force3.z-force4.z);

plugins/amoeba/platforms/cuda2/src/kernels/amoebaStretchBendForce.cu

+// compute the value of the bond angle
+
+real xab = pos1.x - pos2.x;
+real yab = pos1.y - pos2.y;
+real zab = pos1.z - pos2.z;
+
+real xcb = pos3.x - pos2.x;
+real ycb = pos3.y - pos2.y;
+real zcb = pos3.z - pos2.z;
+
+real rab = SQRT(xab*xab + yab*yab + zab*zab);
+real rcb = SQRT(xcb*xcb + ycb*ycb + zcb*zcb);
+
+real xp = ycb*zab - zcb*yab;
+real yp = zcb*xab - xcb*zab;
+real zp = xcb*yab - ycb*xab;
+
+real rp = SQRT(xp*xp + yp*yp + zp*zp);
+
+real dot = xab*xcb + yab*ycb + zab*zcb;
+real cosine = rab*rcb > 0 ? (dot / (rab*rcb)) : (real) 1;
+cosine = (cosine > 1 ? (real) 1 : cosine);
+cosine = (cosine < -1 ? -(real) 1 : cosine);
+real angle = ACOS(cosine);
+
+// find chain rule terms for the bond angle deviation
+
+float4 parameters = PARAMS[index];
+
+real dt = RAD_TO_DEG*(angle - parameters.z);
+real terma = rab*rp != 0 ? (-RAD_TO_DEG/(rab*rab*rp)) : (real) 0;
+real termc = rcb*rp != 0 ? (RAD_TO_DEG/(rcb*rcb*rp)) : (real) 0;
+
+real ddtdxia = terma * (yab*zp-zab*yp);
+real ddtdyia = terma * (zab*xp-xab*zp);
+real ddtdzia = terma * (xab*yp-yab*xp);
+
+real ddtdxic = termc * (ycb*zp-zcb*yp);
+real ddtdyic = termc * (zcb*xp-xcb*zp);
+real ddtdzic = termc * (xcb*yp-ycb*xp);
+
+// find chain rule terms for the bond length deviations
+
+real dr = (parameters.x > 0 ? (rab - parameters.x) : (real) 0);
+terma = (parameters.x > 0 ? RECIP(rab) : (real) 0);
+
+dr += (parameters.y > 0 ? (rcb - parameters.y) : (real) 0);
+termc = (parameters.y > 0 ? RECIP(rcb) : (real) 0);
+
+real ddrdxia = terma * xab;
+real ddrdyia = terma * yab;
+real ddrdzia = terma * zab;
+
+real ddrdxic = termc * xcb;
+real ddrdyic = termc * ycb;
+real ddrdzic = termc * zcb;
+
+// get the energy and master chain rule terms for derivatives
+
+real term = ((rp != 0) ? parameters.w : (real) 0);
+energy += term*dt*dr;
+
+real3 force1 = make_real3(-term*(dt*ddrdxia+ddtdxia*dr), -term*(dt*ddrdyia+ddtdyia*dr), -term*(dt*ddrdzia+ddtdzia*dr));
+real3 force3 = make_real3(-term*(dt*ddrdxic+ddtdxic*dr), -term*(dt*ddrdyic+ddtdyic*dr), -term*(dt*ddrdzic+ddtdzic*dr));
+real3 force2 = make_real3(-force1.x-force3.x, -force1.y-force3.y, -force1.z-force3.z);

plugins/amoeba/platforms/cuda2/src/kernels/amoebaTorsionTorsionForce.cu

+int2 torsionParams = TORSION_PARAMS[index];
+
+real xba = pos2.x - pos1.x;
+real yba = pos2.y - pos1.y;
+real zba = pos2.z - pos1.z;
+
+real xcb = pos3.x - pos2.x;
+real ycb = pos3.y - pos2.y;
+real zcb = pos3.z - pos2.z;
+
+real xdc = pos4.x - pos3.x;
+real ydc = pos4.y - pos3.y;
+real zdc = pos4.z - pos3.z;
+
+real xed = pos5.x - pos4.x;
+real yed = pos5.y - pos4.y;
+real zed = pos5.z - pos4.z;
+
+real xt = yba*zcb - ycb*zba;
+real yt = zba*xcb - zcb*xba;
+real zt = xba*ycb - xcb*yba;
+
+real xu = ycb*zdc - ydc*zcb;
+real yu = zcb*xdc - zdc*xcb;
+real zu = xcb*ydc - xdc*ycb;
+
+real rt2 = xt*xt + yt*yt + zt*zt;
+real ru2 = xu*xu + yu*yu + zu*zu;
+
+real rtru = SQRT(rt2 * ru2);
+
+real xv = ydc*zed - yed*zdc;
+real yv = zdc*xed - zed*xdc;
+real zv = xdc*yed - xed*ydc;
+
+real rv2 = xv*xv + yv*yv + zv*zv;
+real rurv = SQRT(ru2 * rv2);
+
+real rcb = SQRT(xcb*xcb + ycb*ycb + zcb*zcb);
+real cosine1 = (rtru != 0 ? (xt*xu+yt*yu+zt*zu)/rtru : (real) 0);
+cosine1 = (cosine1 > 1 ? (real) 1 : cosine1);
+cosine1 = (cosine1 < -1 ? (real) -1 : cosine1);
+
+real angle1 = RAD_TO_DEG * ACOS(cosine1);
+real sign = xba*xu + yba*yu + zba*zu;
+angle1 = (sign < 0 ? -angle1 : angle1);
+real value1 = angle1;
+
+real rdc = SQRT(xdc*xdc + ydc*ydc + zdc*zdc);
+real cosine2 = (xu*xv + yu*yv + zu*zv) / rurv;
+cosine2 = (cosine2 > 1 ? (real) 1 : cosine2);
+cosine2 = (cosine2 < -1 ? (real) -1 : cosine2);
+real angle2 = RAD_TO_DEG * ACOS(cosine2);
+
+sign = xcb*xv + ycb*yv + zcb*zv;
+angle2 = (sign < 0 ? -angle2 : angle2);
+real value2 = angle2;
+
+// check for inverted chirality at the central atom
+
+// if atom2.y < 0, then no chiral check required
+// sign is set to 1.0 in this case
+// use atom5 for the atom index to avoid warp divergence
+
+int chiralAtomIndex = (torsionParams.x > -1 ? torsionParams.x : atom5);
+real4 pos6 = posq[chiralAtomIndex];
+
+real xac = pos6.x - pos3.x;
+real yac = pos6.y - pos3.y;
+real zac = pos6.z - pos3.z;
+
+real xbc = pos2.x - pos3.x;
+real ybc = pos2.y - pos3.y;
+real zbc = pos2.z - pos3.z;
+
+// xdc, ydc, zdc appear above
+
+real xdc1 = pos4.x - pos3.x;
+real ydc1 = pos4.y - pos3.y;
+real zdc1 = pos4.z - pos3.z;
+
+real c1 = ybc*zdc1 - zbc*ydc1;
+real c2 = ydc1*zac - zdc1*yac;
+real c3 = yac*zbc - zac*ybc;
+real vol = xac*c1 + xbc*c2 + xdc1*c3;
+sign = (vol > 0 ? (real) 1 : (real) -1);
+sign = (torsionParams.x < 0 ? (real) 1 : sign);
+value1 *= sign;
+value2 *= sign;
+
+// use bicubic interpolation to compute spline values
+// compute indices into grid based on angles
+
+float4 gridParams = GRID_PARAMS[torsionParams.y];
+int index1 = (int) ((value1 - gridParams.y)/gridParams.z + 1.0e-05f);
+real fIndex = (real) index1;
+real x1l = gridParams.z*fIndex + gridParams.y;
+real x1u = x1l + gridParams.z;
+
+int index2 = (int) ((value2 - gridParams.y)/gridParams.z + 1.0e-05f);
+fIndex = (real) index2;
+real x2l = gridParams.z*fIndex + gridParams.y;
+real x2u = x2l + gridParams.z;
+
+int posIndex1 = index2 + index1*(int) gridParams.w;
+posIndex1 += (int) gridParams.x;
+
+int posIndex2 = index2 + (index1+1)*(int) gridParams.w;
+posIndex2 += (int) gridParams.x;
+
+// load grid points surrounding angle
+
+real4 y;
+real4 y1;
+real4 y2;
+real4 y12;
+
+int localIndex = posIndex1;
+y.x = GRID_VALUES[localIndex].x;
+y1.x = GRID_VALUES[localIndex].y;
+y2.x = GRID_VALUES[localIndex].z;
+y12.x = GRID_VALUES[localIndex].w;
+
+localIndex = posIndex2;
+y.y = GRID_VALUES[localIndex].x;
+y1.y = GRID_VALUES[localIndex].y;
+y2.y = GRID_VALUES[localIndex].z;
+y12.y = GRID_VALUES[localIndex].w;
+
+localIndex = posIndex2 + 1;
+y.z = GRID_VALUES[localIndex].x;
+y1.z = GRID_VALUES[localIndex].y;
+y2.z = GRID_VALUES[localIndex].z;
+y12.z = GRID_VALUES[localIndex].w;
+
+localIndex = posIndex1 + 1;
+y.w = GRID_VALUES[localIndex].x;
+y1.w = GRID_VALUES[localIndex].y;
+y2.w = GRID_VALUES[localIndex].z;
+y12.w = GRID_VALUES[localIndex].w;
+
+// perform interpolation
+
+real e;
+real dedang1;
+real dedang2;
+
+bicubic(y, y1, y2, y12, value1, x1l, x1u, value2, x2l, x2u, &e, &dedang1, &dedang2);
+energy += e;
+dedang1 *= sign * RAD_TO_DEG;
+dedang2 *= sign * RAD_TO_DEG;
+
+// chain rule terms for first angle derivative components
+
+real xca = pos3.x - pos1.x;
+real yca = pos3.y - pos1.y;
+real zca = pos3.z - pos1.z;
+
+real xdb = pos4.x - pos2.x;
+real ydb = pos4.y - pos2.y;
+real zdb = pos4.z - pos2.z;
+
+real dedxt = dedang1 * (yt*zcb - ycb*zt) / (rt2*rcb);
+real dedyt = dedang1 * (zt*xcb - zcb*xt) / (rt2*rcb);
+real dedzt = dedang1 * (xt*ycb - xcb*yt) / (rt2*rcb);
+real dedxu = -dedang1 * (yu*zcb - ycb*zu) / (ru2*rcb);
+real dedyu = -dedang1 * (zu*xcb - zcb*xu) / (ru2*rcb);
+real dedzu = -dedang1 * (xu*ycb - xcb*yu) / (ru2*rcb);
+
+// compute first derivative components for first angle
+
+real dedxia = zcb*dedyt - ycb*dedzt;
+real dedyia = xcb*dedzt - zcb*dedxt;
+real dedzia = ycb*dedxt - xcb*dedyt;
+
+real dedxib = yca*dedzt - zca*dedyt + zdc*dedyu - ydc*dedzu;
+real dedyib = zca*dedxt - xca*dedzt + xdc*dedzu - zdc*dedxu;
+real dedzib = xca*dedyt - yca*dedxt + ydc*dedxu - xdc*dedyu;
+
+real dedxic = zba*dedyt - yba*dedzt + ydb*dedzu - zdb*dedyu;
+real dedyic = xba*dedzt - zba*dedxt + zdb*dedxu - xdb*dedzu;
+real dedzic = yba*dedxt - xba*dedyt + xdb*dedyu - ydb*dedxu;
+
+real dedxid = zcb*dedyu - ycb*dedzu;
+real dedyid = xcb*dedzu - zcb*dedxu;
+real dedzid = ycb*dedxu - xcb*dedyu;
+
+// chain rule terms for second angle derivative components
+
+real xec = pos5.x - pos3.x;
+real yec = pos5.y - pos3.y;
+real zec = pos5.z - pos3.z;
+
+real dedxu2 = dedang2 * (yu*zdc - ydc*zu) / (ru2*rdc);
+real dedyu2 = dedang2 * (zu*xdc - zdc*xu) / (ru2*rdc);
+real dedzu2 = dedang2 * (xu*ydc - xdc*yu) / (ru2*rdc);
+real dedxv2 = -dedang2 * (yv*zdc - ydc*zv) / (rv2*rdc);
+real dedyv2 = -dedang2 * (zv*xdc - zdc*xv) / (rv2*rdc);
+real dedzv2 = -dedang2 * (xv*ydc - xdc*yv) / (rv2*rdc);
+
+// compute first derivative components for second angle
+
+real dedxib2 = zdc*dedyu2 - ydc*dedzu2;
+real dedyib2 = xdc*dedzu2 - zdc*dedxu2;
+real dedzib2 = ydc*dedxu2 - xdc*dedyu2;
+real dedxic2 = ydb*dedzu2 - zdb*dedyu2 + zed*dedyv2 - yed*dedzv2;
+real dedyic2 = zdb*dedxu2 - xdb*dedzu2 + xed*dedzv2 - zed*dedxv2;
+real dedzic2 = xdb*dedyu2 - ydb*dedxu2 + yed*dedxv2 - xed*dedyv2;
+real dedxid2 = zcb*dedyu2 - ycb*dedzu2 + yec*dedzv2 - zec*dedyv2;
+real dedyid2 = xcb*dedzu2 - zcb*dedxu2 + zec*dedxv2 - xec*dedzv2;
+real dedzid2 = ycb*dedxu2 - xcb*dedyu2 + xec*dedyv2 - yec*dedxv2;
+real dedxie2 = zdc*dedyv2 - ydc*dedzv2;
+real dedyie2 = xdc*dedzv2 - zdc*dedxv2;
+real dedzie2 = ydc*dedxv2 - xdc*dedyv2;
+
+real3 force1 = make_real3(-dedxia, -dedyia, -dedzia);
+real3 force2 = make_real3(-dedxib-dedxib2, -dedyib-dedyib2, -dedzib-dedzib2);
+real3 force3 = make_real3(-dedxic-dedxic2, -dedyic-dedyic2, -dedzic-dedzic2);
+real3 force4 = make_real3(-dedxid-dedxid2, -dedyid-dedyid2, -dedzid-dedzid2);
+real3 force5 = make_real3(-dedxie2, -dedyie2, -dedzie2);
\ No newline at end of file

plugins/amoeba/platforms/cuda2/src/kernels/bicubic.cu

+__device__ void bicubic(real4 y, real4 y1i, real4 y2i, real4 y12i, real x1, real x1l, real x1u,
+                        real x2, real x2l, real x2u, real* energyOut, real* dang1Out, real* dang2Out) {
+    real c[4][4];
+    real d1 = x1u - x1l;
+    real d2 = x2u - x2l;
+    real d12 = d1*d2;
+    real4 y1 = d1*y1i;
+    real4 y2 = d2*y2i;
+    real4 y12 = d12*y12i;
+
+    c[0][0] = y.x;
+    c[0][1] = y2.x;
+    c[0][2] = 3.0f*(y.w - y.x) - (2.0f*y2.x + y2.w);
+    c[0][3] = 2.0f*(y.x - y.w) + y2.x + y2.w;
+    c[1][0] = y1.x;
+    c[1][1] = y12.x;
+    c[1][2] = 3.0f*(y1.w - y1.x) - (2.0f*y12.x + y12.w);
+    c[1][3] = 2.0f*(y1.x - y1.w) + y12.x + y12.w;
+    c[2][0] = 3.0f*(y.y - y.x) - (2.0f*y1.x + y1.y);
+    c[2][1] = 3.0f*(y2.y - y2.x) - (2.0f*y12.x + y12.y);
+    c[2][2] = 9.0f*(y.x - y.y + y.z - y.w) + 6.0f* y1.x  +  3.0f* y1.y  -  3.0f* y1.z  -  6.0f* y1.w +
+                                             6.0f* y2.x  -  6.0f* y2.y  -  3.0f* y2.z  +  3.0f* y2.w +
+                                             4.0f*y12.x  +  2.0f*y12.y  +       y12.z  +  2.0f*y12.w;
+    c[2][3] = 6.0f*(y.y - y.x + y.w - y.z) + -4.0f* y1.x  -  2.0f* y1.y  +  2.0f* y1.z  +  4.0f* y1.w +
+                                             -3.0f* y2.x  +  3.0f* y2.y  +  3.0f* y2.z  -  3.0f* y2.w
+                                             -2.0f*y12.x  -       y12.y  -       y12.z  -  2.0f*y12.w;
+    c[3][0] = 2.0f*(y.x - y.y) + y1.x + y1.y;
+    c[3][1] = 2.0f*(y2.x - y2.y) + y12.x + y12.y;
+    c[3][2] = 6.0f*( y.y -  y.x +  y.w -  y.z) +
+              3.0f*(y1.z + y1.w - y1.x - y1.y) +
+              2.0f*( 2.0f*(y2.y - y2.x) + y2.z - y2.w) +
+             -2.0f*(y12.x + y12.y) - y12.z - y12.w;
+    c[3][3] = 4.0f*(  y.x -    y.y  +   y.z -  y.w)  +
+              2.0f*( y1.x +   y1.y  -  y1.z -  y1.w) +
+              2.0f*( y2.x -   y2.y  -  y2.z +  y2.w) +
+                    y12.x +  y12.y  + y12.z + y12.w;
+
+    real t = (x1-x1l) / (x1u-x1l);
+    real u = (x2-x2l) / (x2u-x2l);
+
+    real energy =       ((c[3][3]*u + c[3][2])*u + c[3][1])*u + c[3][0];
+    energy = t*energy + ((c[2][3]*u + c[2][2])*u + c[2][1])*u + c[2][0];
+    energy = t*energy + ((c[1][3]*u + c[1][2])*u + c[1][1])*u + c[1][0];
+    energy = t*energy + ((c[0][3]*u + c[0][2])*u + c[0][1])*u + c[0][0];
+
+    real dang1 =           (3.0f*c[3][3]*t + 2.0f*c[2][3])*t + c[1][3];
+         dang1 = u*dang1 + (3.0f*c[3][2]*t + 2.0f*c[2][2])*t + c[1][2];
+         dang1 = u*dang1 + (3.0f*c[3][1]*t + 2.0f*c[2][1])*t + c[1][1];
+         dang1 = u*dang1 + (3.0f*c[3][0]*t + 2.0f*c[2][0])*t + c[1][0];
+
+    real dang2 =           (3.0f*c[3][3]*u + 2.0f*c[3][2])*u + c[3][1];
+         dang2 = t*dang2 + (3.0f*c[2][3]*u + 2.0f*c[2][2])*u + c[2][1];
+         dang2 = t*dang2 + (3.0f*c[1][3]*u + 2.0f*c[1][2])*u + c[1][1];
+         dang2 = t*dang2 + (3.0f*c[0][3]*u + 2.0f*c[0][2])*u + c[0][1];
+
+    dang1 = dang1 / (x1u-x1l);
+    dang2 = dang2 / (x2u-x2l);
+
+    *energyOut = energy;
+    *dang1Out = dang1;
+    *dang2Out = dang2;
+}
\ No newline at end of file

plugins/amoeba/platforms/cuda2/tests/TestCudaAmoebaStretchBendForce.cpp

+/* -------------------------------------------------------------------------- *
+ *                                   OpenMMAmoeba                             *
+ * -------------------------------------------------------------------------- *
+ * This is part of the OpenMM molecular simulation toolkit originating from   *
+ * Simbios, the NIH National Center for Physics-Based Simulation of           *
+ * Biological Structures at Stanford, funded under the NIH Roadmap for        *
+ * Medical Research, grant U54 GM072970. See https://simtk.org.               *
+ *                                                                            *
+ * Portions copyright (c) 2008 Stanford University and the Authors.           *
+ * Authors: Mark Friedrichs                                                   *
+ * Contributors:                                                              *
+ *                                                                            *
+ * Permission is hereby granted, free of charge, to any person obtaining a    *
+ * copy of this software and associated documentation files (the "Software"), *
+ * to deal in the Software without restriction, including without limitation  *
+ * the rights to use, copy, modify, merge, publish, distribute, sublicense,   *
+ * and/or sell copies of the Software, and to permit persons to whom the      *
+ * Software is furnished to do so, subject to the following conditions:       *
+ *                                                                            *
+ * The above copyright notice and this permission notice shall be included in *
+ * all copies or substantial portions of the Software.                        *
+ *                                                                            *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR *
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,   *
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL    *
+ * THE AUTHORS, CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,    *
+ * DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR      *
+ * OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE  *
+ * USE OR OTHER DEALINGS IN THE SOFTWARE.                                     *
+ * -------------------------------------------------------------------------- */
+
+/**
+ * This tests the CUDA implementation of CudaAmoebaStretchBendForce.
+ */
+
+#include "openmm/internal/AssertionUtilities.h"
+#include "AmoebaTinkerParameterFile.h"
+#include "openmm/Context.h"
+#include "OpenMMAmoeba.h"
+#include "openmm/System.h"
+#include "openmm/LangevinIntegrator.h"
+#include <iostream>
+#include <vector>
+
+using namespace OpenMM;
+
+const double TOL = 1e-5;
+#define PI_M               3.141592653589
+#define RADIAN            57.29577951308
+
+/* ---------------------------------------------------------------------------------------
+
+   Compute cross product of two 3-vectors and place in 3rd vector
+
+   vectorZ = vectorX x vectorY
+
+   @param vectorX             x-vector
+   @param vectorY             y-vector
+   @param vectorZ             z-vector
+
+   @return vector is vectorZ
+
+   --------------------------------------------------------------------------------------- */
+     
+static void crossProductVector3( double* vectorX, double* vectorY, double* vectorZ ){
+
+    vectorZ[0]  = vectorX[1]*vectorY[2] - vectorX[2]*vectorY[1];
+    vectorZ[1]  = vectorX[2]*vectorY[0] - vectorX[0]*vectorY[2];
+    vectorZ[2]  = vectorX[0]*vectorY[1] - vectorX[1]*vectorY[0];
+
+    return;
+}
+
+static double dotVector3( double* vectorX, double* vectorY ){
+    return vectorX[0]*vectorY[0] + vectorX[1]*vectorY[1] + vectorX[2]*vectorY[2];
+}
+
+
+static void computeAmoebaStretchBendForce(int bondIndex,  std::vector<Vec3>& positions, AmoebaStretchBendForce& amoebaStretchBendForce,
+                                          std::vector<Vec3>& forces, double* energy, FILE* log ) {
+
+    int particle1, particle2, particle3;
+    double abBondLength, cbBondLength, angleStretchBend, kStretchBend;
+
+    amoebaStretchBendForce.getStretchBendParameters(bondIndex, particle1, particle2, particle3, abBondLength, cbBondLength, angleStretchBend, kStretchBend);
+    angleStretchBend *= RADIAN;
+#ifdef AMOEBA_DEBUG
+    if( log ){
+        (void) fprintf( log, "computeAmoebaStretchBendForce: bond %d [%d %d %d] ab=%10.3e cb=%10.3e angle=%10.3e k=%10.3e\n", 
+                             bondIndex, particle1, particle2, particle3, abBondLength, cbBondLength, angleStretchBend, kStretchBend );
+        (void) fflush( log );
+    }
+#endif
+
+    enum { A, B, C, LastAtomIndex };
+    enum { AB, CB, CBxAB, ABxP, CBxP, LastDeltaIndex };
+ 
+    // ---------------------------------------------------------------------------------------
+ 
+    // get deltaR between various combinations of the 3 atoms
+    // and various intermediate terms
+ 
+    double deltaR[LastDeltaIndex][3];
+    double rAB2 = 0.0;
+    double rCB2 = 0.0;
+    for( int ii = 0; ii < 3; ii++ ){
+         deltaR[AB][ii]  = positions[particle1][ii] - positions[particle2][ii];
+         rAB2           += deltaR[AB][ii]*deltaR[AB][ii];
+
+         deltaR[CB][ii]  = positions[particle3][ii] - positions[particle2][ii];
+         rCB2           += deltaR[CB][ii]*deltaR[CB][ii];
+    }
+    double rAB   = sqrt( rAB2 );
+    double rCB   = sqrt( rCB2 );
+
+    crossProductVector3( deltaR[CB], deltaR[AB], deltaR[CBxAB] );
+    double  rP   = dotVector3( deltaR[CBxAB], deltaR[CBxAB] );
+            rP   = sqrt( rP );
+ 
+    if( rP <= 0.0 ){
+       return;
+    }
+    double dot    = dotVector3( deltaR[CB], deltaR[AB] );
+    double cosine = dot/(rAB*rCB);
+ 
+    double angle;
+    if( cosine >= 1.0 ){
+       angle = 0.0;
+    } else if( cosine <= -1.0 ){
+       angle = PI_M;
+    } else {
+       angle = RADIAN*acos(cosine);
+    }
+ 
+    double termA = -RADIAN/(rAB2*rP);
+    double termC =  RADIAN/(rCB2*rP);
+ 
+    // P = CBxAB
+ 
+    crossProductVector3( deltaR[AB], deltaR[CBxAB], deltaR[ABxP] );
+    crossProductVector3( deltaR[CB], deltaR[CBxAB], deltaR[CBxP] );
+    for( int ii = 0; ii < 3; ii++ ){
+       deltaR[ABxP][ii] *= termA;
+       deltaR[CBxP][ii] *= termC;
+    }
+ 
+    double dr    = rAB - abBondLength + rCB - cbBondLength;
+ 
+    termA        = 1.0/rAB;
+    termC        = 1.0/rCB;
+ 
+    double term  = kStretchBend;
+ 
+    // ---------------------------------------------------------------------------------------
+ 
+    // forces
+ 
+    // calculate forces for atoms a, b, c
+    // the force for b is then -( a + c)
+ 
+    double subForce[LastAtomIndex][3];
+    double dt = angle - angleStretchBend;
+    for( int jj = 0; jj < 3; jj++ ){
+        subForce[A][jj] = term*(dt*termA*deltaR[AB][jj] + dr*deltaR[ABxP][jj] );
+        subForce[C][jj] = term*(dt*termC*deltaR[CB][jj] + dr*deltaR[CBxP][jj] );
+        subForce[B][jj] = -( subForce[A][jj] + subForce[C][jj] );
+    }
+ 
+    // ---------------------------------------------------------------------------------------
+ 
+    // accumulate forces and energy
+ 
+    forces[particle1][0]       -= subForce[0][0];
+    forces[particle1][1]       -= subForce[0][1];
+    forces[particle1][2]       -= subForce[0][2];
+
+    forces[particle2][0]       -= subForce[1][0];
+    forces[particle2][1]       -= subForce[1][1];
+    forces[particle2][2]       -= subForce[1][2];
+
+    forces[particle3][0]       -= subForce[2][0];
+    forces[particle3][1]       -= subForce[2][1];
+    forces[particle3][2]       -= subForce[2][2];
+
+    *energy                    += term*dt*dr;
+
+#ifdef AMOEBA_DEBUG
+    if( log ){
+        (void) fprintf( log, "computeAmoebaStretchBendForce: angle=%10.3e dt=%10.3e dr=%10.3e\n", angle, dt, dr ); 
+        (void) fflush( log );
+    }
+#endif
+
+    return;
+}
+ 
+static void computeAmoebaStretchBendForces( Context& context, AmoebaStretchBendForce& amoebaStretchBendForce,
+                                          std::vector<Vec3>& expectedForces, double* expectedEnergy, FILE* log ) {
+
+    // get positions and zero forces
+
+    State state                 = context.getState(State::Positions);
+    std::vector<Vec3> positions = state.getPositions();
+    expectedForces.resize( positions.size() );
+    
+    for( unsigned int ii = 0; ii < expectedForces.size(); ii++ ){
+        expectedForces[ii][0] = expectedForces[ii][1] = expectedForces[ii][2] = 0.0;
+    }
+
+    // calculates forces/energy
+
+    *expectedEnergy = 0.0;
+    for( int ii = 0; ii < amoebaStretchBendForce.getNumStretchBends(); ii++ ){
+        computeAmoebaStretchBendForce(ii, positions, amoebaStretchBendForce, expectedForces, expectedEnergy, log );
+    }
+
+#ifdef AMOEBA_DEBUG
+    if( log ){
+        (void) fprintf( log, "computeAmoebaStretchBendForces: expected energy=%14.7e\n", *expectedEnergy );
+        for( unsigned int ii = 0; ii < positions.size(); ii++ ){
+            (void) fprintf( log, "%6u [%14.7e %14.7e %14.7e]\n", ii, expectedForces[ii][0], expectedForces[ii][1], expectedForces[ii][2] );
+        }
+        (void) fflush( log );
+    }
+#endif
+    return;
+
+}
+
+void compareWithExpectedForceAndEnergy( Context& context, AmoebaStretchBendForce& amoebaStretchBendForce,
+                                        double tolerance, const std::string& idString, FILE* log) {
+
+    std::vector<Vec3> expectedForces;
+    double expectedEnergy;
+    computeAmoebaStretchBendForces( context, amoebaStretchBendForce, expectedForces, &expectedEnergy, log );
+   
+    State state                      = context.getState(State::Forces | State::Energy);
+    const std::vector<Vec3> forces   = state.getForces();
+
+#ifdef AMOEBA_DEBUG
+    if( log ){
+        (void) fprintf( log, "computeAmoebaStretchBendForces: expected energy=%14.7e %14.7e\n", expectedEnergy, state.getPotentialEnergy() );
+        for( unsigned int ii = 0; ii < forces.size(); ii++ ){
+            (void) fprintf( log, "%6u [%14.7e %14.7e %14.7e]   [%14.7e %14.7e %14.7e]\n", ii,
+                            expectedForces[ii][0], expectedForces[ii][1], expectedForces[ii][2], forces[ii][0], forces[ii][1], forces[ii][2] );
+        }
+        (void) fflush( log );
+    }
+#endif
+
+    for( unsigned int ii = 0; ii < forces.size(); ii++ ){
+        ASSERT_EQUAL_VEC( expectedForces[ii], forces[ii], tolerance );
+    }
+    ASSERT_EQUAL_TOL( expectedEnergy, state.getPotentialEnergy(), tolerance );
+}
+
+void testOneStretchBend( FILE* log ) {
+
+    System system;
+    int numberOfParticles = 3;
+    for( int ii = 0; ii < numberOfParticles; ii++ ){
+        system.addParticle(1.0);
+    }
+
+    LangevinIntegrator integrator(0.0, 0.1, 0.01);
+
+    AmoebaStretchBendForce* amoebaStretchBendForce = new AmoebaStretchBendForce();
+
+    double abLength         = 0.144800000E+01;
+    double cbLength         = 0.101500000E+01;
+    double angleStretchBend = 0.108500000E+03*DegreesToRadians;
+    //double kStretchBend     = 0.750491578E-01;
+    double kStretchBend     = 1.0;
+
+    amoebaStretchBendForce->addStretchBend(0, 1, 2, abLength, cbLength, angleStretchBend, kStretchBend );
+
+    system.addForce(amoebaStretchBendForce);
+    Context context(system, integrator, Platform::getPlatformByName( "CUDA"));
+
+    std::vector<Vec3> positions(numberOfParticles);
+
+    positions[0] = Vec3( 0.262660000E+02,  0.254130000E+02,  0.284200000E+01 );
+    positions[1] = Vec3( 0.273400000E+02,  0.244300000E+02,  0.261400000E+01 );
+    positions[2] = Vec3( 0.269573220E+02,  0.236108860E+02,  0.216376800E+01 );
+
+    context.setPositions(positions);
+    compareWithExpectedForceAndEnergy( context, *amoebaStretchBendForce, TOL, "testOneStretchBend", log );
+
+}
+
+int main( int numberOfArguments, char* argv[] ) {
+
+    try {
+        std::cout << "TestCudaAmoebaStretchBendForce running test..." << std::endl;
+        registerAmoebaCudaKernelFactories();
+
+        FILE* log = NULL;
+        testOneStretchBend( log );
+    } catch(const std::exception& e) {
+        std::cout << "exception: " << e.what() << std::endl;
+        std::cout << "FAIL - ERROR.  Test failed." << std::endl;
+        return 1;
+    }
+    std::cout << "Done" << std::endl;
+    return 0;
+}
+
+
+