Skip to content
GitLab
Menu
Projects
Groups
Snippets
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in
Toggle navigation
Menu
Open sidebar
tsoc
openmm
Commits
a87df3ad
Commit
a87df3ad
authored
Jul 02, 2014
by
peastman
Browse files
Merged upstream changes
parents
1c604fe4
9db70803
Changes
27
Show whitespace changes
Inline
Side-by-side
Showing
20 changed files
with
48 additions
and
47 deletions
+48
-47
.travis.yml
.travis.yml
+2
-1
platforms/cuda/src/kernels/customGBEnergyN2.cu
platforms/cuda/src/kernels/customGBEnergyN2.cu
+2
-2
platforms/cuda/src/kernels/customGBValueN2.cu
platforms/cuda/src/kernels/customGBValueN2.cu
+2
-2
platforms/cuda/src/kernels/gbsaObc1.cu
platforms/cuda/src/kernels/gbsaObc1.cu
+4
-4
platforms/cuda/src/kernels/nonbonded.cu
platforms/cuda/src/kernels/nonbonded.cu
+2
-2
platforms/opencl/src/kernels/customGBEnergyN2.cl
platforms/opencl/src/kernels/customGBEnergyN2.cl
+2
-2
platforms/opencl/src/kernels/customGBEnergyN2_cpu.cl
platforms/opencl/src/kernels/customGBEnergyN2_cpu.cl
+2
-2
platforms/opencl/src/kernels/customGBValueN2.cl
platforms/opencl/src/kernels/customGBValueN2.cl
+2
-2
platforms/opencl/src/kernels/customGBValueN2_cpu.cl
platforms/opencl/src/kernels/customGBValueN2_cpu.cl
+2
-2
platforms/opencl/src/kernels/gbsaObc.cl
platforms/opencl/src/kernels/gbsaObc.cl
+4
-4
platforms/opencl/src/kernels/gbsaObc_cpu.cl
platforms/opencl/src/kernels/gbsaObc_cpu.cl
+4
-4
platforms/opencl/src/kernels/nonbonded.cl
platforms/opencl/src/kernels/nonbonded.cl
+2
-2
platforms/opencl/src/kernels/nonbonded_cpu.cl
platforms/opencl/src/kernels/nonbonded_cpu.cl
+2
-2
plugins/amoeba/openmmapi/include/openmm/AmoebaMultipoleForce.h
...ns/amoeba/openmmapi/include/openmm/AmoebaMultipoleForce.h
+1
-1
plugins/amoeba/openmmapi/include/openmm/AmoebaVdwForce.h
plugins/amoeba/openmmapi/include/openmm/AmoebaVdwForce.h
+1
-1
plugins/amoeba/openmmapi/src/AmoebaMultipoleForce.cpp
plugins/amoeba/openmmapi/src/AmoebaMultipoleForce.cpp
+1
-1
plugins/amoeba/openmmapi/src/AmoebaVdwForce.cpp
plugins/amoeba/openmmapi/src/AmoebaVdwForce.cpp
+1
-1
plugins/amoeba/platforms/cuda/src/kernels/amoebaGk.cu
plugins/amoeba/platforms/cuda/src/kernels/amoebaGk.cu
+8
-8
plugins/amoeba/platforms/cuda/src/kernels/amoebaWcaForce.cu
plugins/amoeba/platforms/cuda/src/kernels/amoebaWcaForce.cu
+2
-2
plugins/amoeba/platforms/cuda/src/kernels/multipoleElectrostatics.cu
...eba/platforms/cuda/src/kernels/multipoleElectrostatics.cu
+2
-2
No files found.
.travis.yml
View file @
a87df3ad
...
@@ -6,7 +6,8 @@ before_install:
...
@@ -6,7 +6,8 @@ before_install:
-
sudo apt-get update -qq
-
sudo apt-get update -qq
-
sudo apt-get install -qq libpcre3 libpcre3-dev gromacs
-
sudo apt-get install -qq libpcre3 libpcre3-dev gromacs
-
sudo apt-get install -qq swig doxygen llvm-3.3
-
sudo apt-get install -qq swig doxygen llvm-3.3
-
sudo apt-get install -qq python-numpy python-scipy python-nose
-
sudo apt-get install -qq python-numpy python-scipy python-pip
-
sudo pip install nose
-
export ASAN_SYMBOLIZER_PATH=/usr/bin/llvm-symbolizer-3.3
-
export ASAN_SYMBOLIZER_PATH=/usr/bin/llvm-symbolizer-3.3
script
:
script
:
...
...
platforms/cuda/src/kernels/customGBEnergyN2.cu
View file @
a87df3ad
...
@@ -190,7 +190,7 @@ extern "C" __global__ void computeN2Energy(unsigned long long* __restrict__ forc
...
@@ -190,7 +190,7 @@ extern "C" __global__ void computeN2Energy(unsigned long long* __restrict__ forc
// Extract the coordinates of this tile.
// Extract the coordinates of this tile.
unsigned
int
x
,
y
;
int
x
,
y
;
bool
singlePeriodicCopy
=
false
;
bool
singlePeriodicCopy
=
false
;
#ifdef USE_CUTOFF
#ifdef USE_CUTOFF
if
(
numTiles
<=
maxTiles
)
{
if
(
numTiles
<=
maxTiles
)
{
...
@@ -203,7 +203,7 @@ extern "C" __global__ void computeN2Energy(unsigned long long* __restrict__ forc
...
@@ -203,7 +203,7 @@ extern "C" __global__ void computeN2Energy(unsigned long long* __restrict__ forc
else
else
#endif
#endif
{
{
y
=
(
unsigned
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
y
=
(
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
y
+=
(
x
<
y
?
-
1
:
1
);
y
+=
(
x
<
y
?
-
1
:
1
);
...
...
platforms/cuda/src/kernels/customGBValueN2.cu
View file @
a87df3ad
...
@@ -166,7 +166,7 @@ extern "C" __global__ void computeN2Value(const real4* __restrict__ posq, const
...
@@ -166,7 +166,7 @@ extern "C" __global__ void computeN2Value(const real4* __restrict__ posq, const
// Extract the coordinates of this tile.
// Extract the coordinates of this tile.
unsigned
int
x
,
y
;
int
x
,
y
;
bool
singlePeriodicCopy
=
false
;
bool
singlePeriodicCopy
=
false
;
#ifdef USE_CUTOFF
#ifdef USE_CUTOFF
if
(
numTiles
<=
maxTiles
)
{
if
(
numTiles
<=
maxTiles
)
{
...
@@ -179,7 +179,7 @@ extern "C" __global__ void computeN2Value(const real4* __restrict__ posq, const
...
@@ -179,7 +179,7 @@ extern "C" __global__ void computeN2Value(const real4* __restrict__ posq, const
else
else
#endif
#endif
{
{
y
=
(
unsigned
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
y
=
(
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
y
+=
(
x
<
y
?
-
1
:
1
);
y
+=
(
x
<
y
?
-
1
:
1
);
...
...
platforms/cuda/src/kernels/gbsaObc1.cu
View file @
a87df3ad
...
@@ -226,7 +226,7 @@ extern "C" __global__ void computeBornSum(unsigned long long* __restrict__ globa
...
@@ -226,7 +226,7 @@ extern "C" __global__ void computeBornSum(unsigned long long* __restrict__ globa
// Extract the coordinates of this tile.
// Extract the coordinates of this tile.
unsigned
int
x
,
y
;
int
x
,
y
;
bool
singlePeriodicCopy
=
false
;
bool
singlePeriodicCopy
=
false
;
#ifdef USE_CUTOFF
#ifdef USE_CUTOFF
if
(
numTiles
<=
maxTiles
)
{
if
(
numTiles
<=
maxTiles
)
{
...
@@ -239,7 +239,7 @@ extern "C" __global__ void computeBornSum(unsigned long long* __restrict__ globa
...
@@ -239,7 +239,7 @@ extern "C" __global__ void computeBornSum(unsigned long long* __restrict__ globa
else
else
#endif
#endif
{
{
y
=
(
unsigned
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
y
=
(
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
y
+=
(
x
<
y
?
-
1
:
1
);
y
+=
(
x
<
y
?
-
1
:
1
);
...
@@ -590,7 +590,7 @@ extern "C" __global__ void computeGBSAForce1(unsigned long long* __restrict__ fo
...
@@ -590,7 +590,7 @@ extern "C" __global__ void computeGBSAForce1(unsigned long long* __restrict__ fo
// Extract the coordinates of this tile.
// Extract the coordinates of this tile.
unsigned
int
x
,
y
;
int
x
,
y
;
bool
singlePeriodicCopy
=
false
;
bool
singlePeriodicCopy
=
false
;
#ifdef USE_CUTOFF
#ifdef USE_CUTOFF
if
(
numTiles
<=
maxTiles
)
{
if
(
numTiles
<=
maxTiles
)
{
...
@@ -603,7 +603,7 @@ extern "C" __global__ void computeGBSAForce1(unsigned long long* __restrict__ fo
...
@@ -603,7 +603,7 @@ extern "C" __global__ void computeGBSAForce1(unsigned long long* __restrict__ fo
else
else
#endif
#endif
{
{
y
=
(
unsigned
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
y
=
(
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
y
+=
(
x
<
y
?
-
1
:
1
);
y
+=
(
x
<
y
?
-
1
:
1
);
...
...
platforms/cuda/src/kernels/nonbonded.cu
View file @
a87df3ad
...
@@ -333,7 +333,7 @@ extern "C" __global__ void computeNonbonded(
...
@@ -333,7 +333,7 @@ extern "C" __global__ void computeNonbonded(
bool
includeTile
=
true
;
bool
includeTile
=
true
;
// Extract the coordinates of this tile.
// Extract the coordinates of this tile.
unsigned
int
x
,
y
;
int
x
,
y
;
bool
singlePeriodicCopy
=
false
;
bool
singlePeriodicCopy
=
false
;
#ifdef USE_CUTOFF
#ifdef USE_CUTOFF
if
(
numTiles
<=
maxTiles
)
{
if
(
numTiles
<=
maxTiles
)
{
...
@@ -346,7 +346,7 @@ extern "C" __global__ void computeNonbonded(
...
@@ -346,7 +346,7 @@ extern "C" __global__ void computeNonbonded(
else
else
#endif
#endif
{
{
y
=
(
unsigned
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
y
=
(
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
y
+=
(
x
<
y
?
-
1
:
1
);
y
+=
(
x
<
y
?
-
1
:
1
);
...
...
platforms/opencl/src/kernels/customGBEnergyN2.cl
View file @
a87df3ad
...
@@ -200,7 +200,7 @@ __kernel void computeN2Energy(
...
@@ -200,7 +200,7 @@ __kernel void computeN2Energy(
// Extract the coordinates of this tile.
// Extract the coordinates of this tile.
unsigned
int x, y;
int x, y;
bool singlePeriodicCopy = false;
bool singlePeriodicCopy = false;
#ifdef USE_CUTOFF
#ifdef USE_CUTOFF
if (numTiles <= maxTiles) {
if (numTiles <= maxTiles) {
...
@@ -213,7 +213,7 @@ __kernel void computeN2Energy(
...
@@ -213,7 +213,7 @@ __kernel void computeN2Energy(
else
else
#endif
#endif
{
{
y = (
unsigned
int) floor(NUM_BLOCKS+0.5f-SQRT((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
y = (int) floor(NUM_BLOCKS+0.5f-SQRT((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
if (x < y |
|
x
>=
NUM_BLOCKS
)
{
//
Occasionally
happens
due
to
roundoff
error.
if (x < y |
|
x
>=
NUM_BLOCKS
)
{
//
Occasionally
happens
due
to
roundoff
error.
y
+=
(
x
<
y
?
-1
:
1
)
;
y
+=
(
x
<
y
?
-1
:
1
)
;
...
...
platforms/opencl/src/kernels/customGBEnergyN2_cpu.cl
View file @
a87df3ad
...
@@ -216,7 +216,7 @@ __kernel void computeN2Energy(
...
@@ -216,7 +216,7 @@ __kernel void computeN2Energy(
// Extract the coordinates of this tile.
// Extract the coordinates of this tile.
unsigned
int x, y;
int x, y;
bool singlePeriodicCopy = false;
bool singlePeriodicCopy = false;
#ifdef USE_CUTOFF
#ifdef USE_CUTOFF
if (numTiles <= maxTiles) {
if (numTiles <= maxTiles) {
...
@@ -229,7 +229,7 @@ __kernel void computeN2Energy(
...
@@ -229,7 +229,7 @@ __kernel void computeN2Energy(
else
else
#endif
#endif
{
{
y = (
unsigned
int) floor(NUM_BLOCKS+0.5f-SQRT((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
y = (int) floor(NUM_BLOCKS+0.5f-SQRT((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
if (x < y |
|
x
>=
NUM_BLOCKS
)
{
//
Occasionally
happens
due
to
roundoff
error.
if (x < y |
|
x
>=
NUM_BLOCKS
)
{
//
Occasionally
happens
due
to
roundoff
error.
y
+=
(
x
<
y
?
-1
:
1
)
;
y
+=
(
x
<
y
?
-1
:
1
)
;
...
...
platforms/opencl/src/kernels/customGBValueN2.cl
View file @
a87df3ad
...
@@ -174,7 +174,7 @@ __kernel void computeN2Value(__global const real4* restrict posq, __local real4*
...
@@ -174,7 +174,7 @@ __kernel void computeN2Value(__global const real4* restrict posq, __local real4*
// Extract the coordinates of this tile.
// Extract the coordinates of this tile.
unsigned
int x, y;
int x, y;
bool singlePeriodicCopy = false;
bool singlePeriodicCopy = false;
#ifdef USE_CUTOFF
#ifdef USE_CUTOFF
if (numTiles <= maxTiles) {
if (numTiles <= maxTiles) {
...
@@ -187,7 +187,7 @@ __kernel void computeN2Value(__global const real4* restrict posq, __local real4*
...
@@ -187,7 +187,7 @@ __kernel void computeN2Value(__global const real4* restrict posq, __local real4*
else
else
#endif
#endif
{
{
y = (
unsigned
int) floor(NUM_BLOCKS+0.5f-SQRT((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
y = (int) floor(NUM_BLOCKS+0.5f-SQRT((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
if (x < y |
|
x
>=
NUM_BLOCKS
)
{
//
Occasionally
happens
due
to
roundoff
error.
if (x < y |
|
x
>=
NUM_BLOCKS
)
{
//
Occasionally
happens
due
to
roundoff
error.
y
+=
(
x
<
y
?
-1
:
1
)
;
y
+=
(
x
<
y
?
-1
:
1
)
;
...
...
platforms/opencl/src/kernels/customGBValueN2_cpu.cl
View file @
a87df3ad
...
@@ -184,7 +184,7 @@ __kernel void computeN2Value(__global const real4* restrict posq, __local real4*
...
@@ -184,7 +184,7 @@ __kernel void computeN2Value(__global const real4* restrict posq, __local real4*
// Extract the coordinates of this tile.
// Extract the coordinates of this tile.
unsigned
int x, y;
int x, y;
bool singlePeriodicCopy = false;
bool singlePeriodicCopy = false;
#ifdef USE_CUTOFF
#ifdef USE_CUTOFF
if (numTiles <= maxTiles) {
if (numTiles <= maxTiles) {
...
@@ -197,7 +197,7 @@ __kernel void computeN2Value(__global const real4* restrict posq, __local real4*
...
@@ -197,7 +197,7 @@ __kernel void computeN2Value(__global const real4* restrict posq, __local real4*
else
else
#endif
#endif
{
{
y = (
unsigned
int) floor(NUM_BLOCKS+0.5f-SQRT((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
y = (int) floor(NUM_BLOCKS+0.5f-SQRT((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
if (x < y |
|
x
>=
NUM_BLOCKS
)
{
//
Occasionally
happens
due
to
roundoff
error.
if (x < y |
|
x
>=
NUM_BLOCKS
)
{
//
Occasionally
happens
due
to
roundoff
error.
y
+=
(
x
<
y
?
-1
:
1
)
;
y
+=
(
x
<
y
?
-1
:
1
)
;
...
...
platforms/opencl/src/kernels/gbsaObc.cl
View file @
a87df3ad
...
@@ -186,7 +186,7 @@ __kernel void computeBornSum(
...
@@ -186,7 +186,7 @@ __kernel void computeBornSum(
//
Extract
the
coordinates
of
this
tile.
//
Extract
the
coordinates
of
this
tile.
unsigned
int
x,
y
;
int
x,
y
;
bool
singlePeriodicCopy
=
false
;
bool
singlePeriodicCopy
=
false
;
#
ifdef
USE_CUTOFF
#
ifdef
USE_CUTOFF
if
(
numTiles
<=
maxTiles
)
{
if
(
numTiles
<=
maxTiles
)
{
...
@@ -199,7 +199,7 @@ __kernel void computeBornSum(
...
@@ -199,7 +199,7 @@ __kernel void computeBornSum(
else
else
#
endif
#
endif
{
{
y
=
(
unsigned
int
)
floor
(
NUM_BLOCKS+0.5f-SQRT
((
NUM_BLOCKS+0.5f
)
*
(
NUM_BLOCKS+0.5f
)
-2*pos
))
;
y
=
(
int
)
floor
(
NUM_BLOCKS+0.5f-SQRT
((
NUM_BLOCKS+0.5f
)
*
(
NUM_BLOCKS+0.5f
)
-2*pos
))
;
x
=
(
pos-y*NUM_BLOCKS+y*
(
y+1
)
/2
)
;
x
=
(
pos-y*NUM_BLOCKS+y*
(
y+1
)
/2
)
;
if
(
x
<
y
|
| x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
if
(
x
<
y
|
| x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
y += (x < y ? -1 : 1);
y += (x < y ? -1 : 1);
...
@@ -574,7 +574,7 @@ __kernel void computeGBSAForce1(
...
@@ -574,7 +574,7 @@ __kernel void computeGBSAForce1(
// Extract the coordinates of this tile.
// Extract the coordinates of this tile.
unsigned
int x, y;
int x, y;
bool singlePeriodicCopy = false;
bool singlePeriodicCopy = false;
#ifdef USE_CUTOFF
#ifdef USE_CUTOFF
if (numTiles <= maxTiles) {
if (numTiles <= maxTiles) {
...
@@ -587,7 +587,7 @@ __kernel void computeGBSAForce1(
...
@@ -587,7 +587,7 @@ __kernel void computeGBSAForce1(
else
else
#endif
#endif
{
{
y = (
unsigned
int) floor(NUM_BLOCKS+0.5f-SQRT((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
y = (int) floor(NUM_BLOCKS+0.5f-SQRT((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
if (x < y |
|
x
>=
NUM_BLOCKS
)
{
//
Occasionally
happens
due
to
roundoff
error.
if (x < y |
|
x
>=
NUM_BLOCKS
)
{
//
Occasionally
happens
due
to
roundoff
error.
y
+=
(
x
<
y
?
-1
:
1
)
;
y
+=
(
x
<
y
?
-1
:
1
)
;
...
...
platforms/opencl/src/kernels/gbsaObc_cpu.cl
View file @
a87df3ad
...
@@ -192,7 +192,7 @@ __kernel void computeBornSum(
...
@@ -192,7 +192,7 @@ __kernel void computeBornSum(
//
Extract
the
coordinates
of
this
tile.
//
Extract
the
coordinates
of
this
tile.
unsigned
int
x,
y
;
int
x,
y
;
bool
singlePeriodicCopy
=
false
;
bool
singlePeriodicCopy
=
false
;
#
ifdef
USE_CUTOFF
#
ifdef
USE_CUTOFF
if
(
numTiles
<=
maxTiles
)
{
if
(
numTiles
<=
maxTiles
)
{
...
@@ -205,7 +205,7 @@ __kernel void computeBornSum(
...
@@ -205,7 +205,7 @@ __kernel void computeBornSum(
else
else
#
endif
#
endif
{
{
y
=
(
unsigned
int
)
floor
(
NUM_BLOCKS+0.5f-SQRT
((
NUM_BLOCKS+0.5f
)
*
(
NUM_BLOCKS+0.5f
)
-2*pos
))
;
y
=
(
int
)
floor
(
NUM_BLOCKS+0.5f-SQRT
((
NUM_BLOCKS+0.5f
)
*
(
NUM_BLOCKS+0.5f
)
-2*pos
))
;
x
=
(
pos-y*NUM_BLOCKS+y*
(
y+1
)
/2
)
;
x
=
(
pos-y*NUM_BLOCKS+y*
(
y+1
)
/2
)
;
if
(
x
<
y
|
| x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
if
(
x
<
y
|
| x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
y += (x < y ? -1 : 1);
y += (x < y ? -1 : 1);
...
@@ -607,7 +607,7 @@ __kernel void computeGBSAForce1(
...
@@ -607,7 +607,7 @@ __kernel void computeGBSAForce1(
// Extract the coordinates of this tile.
// Extract the coordinates of this tile.
unsigned
int x, y;
int x, y;
bool singlePeriodicCopy = false;
bool singlePeriodicCopy = false;
#ifdef USE_CUTOFF
#ifdef USE_CUTOFF
if (numTiles <= maxTiles) {
if (numTiles <= maxTiles) {
...
@@ -620,7 +620,7 @@ __kernel void computeGBSAForce1(
...
@@ -620,7 +620,7 @@ __kernel void computeGBSAForce1(
else
else
#endif
#endif
{
{
y = (
unsigned
int) floor(NUM_BLOCKS+0.5f-SQRT((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
y = (int) floor(NUM_BLOCKS+0.5f-SQRT((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
if (x < y |
|
x
>=
NUM_BLOCKS
)
{
//
Occasionally
happens
due
to
roundoff
error.
if (x < y |
|
x
>=
NUM_BLOCKS
)
{
//
Occasionally
happens
due
to
roundoff
error.
y
+=
(
x
<
y
?
-1
:
1
)
;
y
+=
(
x
<
y
?
-1
:
1
)
;
...
...
platforms/opencl/src/kernels/nonbonded.cl
View file @
a87df3ad
...
@@ -213,7 +213,7 @@ __kernel void computeNonbonded(
...
@@ -213,7 +213,7 @@ __kernel void computeNonbonded(
// Extract the coordinates of this tile.
// Extract the coordinates of this tile.
unsigned
int x, y;
int x, y;
bool singlePeriodicCopy = false;
bool singlePeriodicCopy = false;
#ifdef USE_CUTOFF
#ifdef USE_CUTOFF
if (numTiles <= maxTiles) {
if (numTiles <= maxTiles) {
...
@@ -226,7 +226,7 @@ __kernel void computeNonbonded(
...
@@ -226,7 +226,7 @@ __kernel void computeNonbonded(
else
else
#endif
#endif
{
{
y = (
unsigned
int) floor(NUM_BLOCKS+0.5f-SQRT((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
y = (int) floor(NUM_BLOCKS+0.5f-SQRT((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
y += (x < y ? -1 : 1);
y += (x < y ? -1 : 1);
...
...
platforms/opencl/src/kernels/nonbonded_cpu.cl
View file @
a87df3ad
...
@@ -230,7 +230,7 @@ __kernel void computeNonbonded(
...
@@ -230,7 +230,7 @@ __kernel void computeNonbonded(
// Extract the coordinates of this tile.
// Extract the coordinates of this tile.
unsigned
int x, y;
int x, y;
bool singlePeriodicCopy = false;
bool singlePeriodicCopy = false;
#ifdef USE_CUTOFF
#ifdef USE_CUTOFF
if (numTiles <= maxTiles) {
if (numTiles <= maxTiles) {
...
@@ -243,7 +243,7 @@ __kernel void computeNonbonded(
...
@@ -243,7 +243,7 @@ __kernel void computeNonbonded(
else
else
#endif
#endif
{
{
y = (
unsigned
int) floor(NUM_BLOCKS+0.5f-SQRT((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
y = (int) floor(NUM_BLOCKS+0.5f-SQRT((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
y += (x < y ? -1 : 1);
y += (x < y ? -1 : 1);
...
...
plugins/amoeba/openmmapi/include/openmm/AmoebaMultipoleForce.h
View file @
a87df3ad
...
@@ -126,7 +126,7 @@ public:
...
@@ -126,7 +126,7 @@ public:
*
*
* @return the cutoff distance, measured in nm
* @return the cutoff distance, measured in nm
*/
*/
double
getCutoffDistance
(
void
)
const
;
double
getCutoffDistance
()
const
;
/**
/**
* Set the cutoff distance (in nm) being used for nonbonded interactions. If the NonbondedMethod in use
* Set the cutoff distance (in nm) being used for nonbonded interactions. If the NonbondedMethod in use
...
...
plugins/amoeba/openmmapi/include/openmm/AmoebaVdwForce.h
View file @
a87df3ad
...
@@ -191,7 +191,7 @@ public:
...
@@ -191,7 +191,7 @@ public:
/**
/**
* Get the cutoff distance.
* Get the cutoff distance.
*/
*/
double
getCutoff
(
void
)
const
;
double
getCutoff
()
const
;
/**
/**
* Get the method used for handling long range nonbonded interactions.
* Get the method used for handling long range nonbonded interactions.
...
...
plugins/amoeba/openmmapi/src/AmoebaMultipoleForce.cpp
View file @
a87df3ad
...
@@ -61,7 +61,7 @@ void AmoebaMultipoleForce::setPolarizationType( AmoebaMultipoleForce::Polarizati
...
@@ -61,7 +61,7 @@ void AmoebaMultipoleForce::setPolarizationType( AmoebaMultipoleForce::Polarizati
polarizationType
=
type
;
polarizationType
=
type
;
}
}
double
AmoebaMultipoleForce
::
getCutoffDistance
(
void
)
const
{
double
AmoebaMultipoleForce
::
getCutoffDistance
()
const
{
return
cutoffDistance
;
return
cutoffDistance
;
}
}
...
...
plugins/amoeba/openmmapi/src/AmoebaVdwForce.cpp
View file @
a87df3ad
...
@@ -106,7 +106,7 @@ void AmoebaVdwForce::setCutoff( double inputCutoff ){
...
@@ -106,7 +106,7 @@ void AmoebaVdwForce::setCutoff( double inputCutoff ){
cutoff
=
inputCutoff
;
cutoff
=
inputCutoff
;
}
}
double
AmoebaVdwForce
::
getCutoff
(
void
)
const
{
double
AmoebaVdwForce
::
getCutoff
()
const
{
return
cutoff
;
return
cutoff
;
}
}
...
...
plugins/amoeba/platforms/cuda/src/kernels/amoebaGk.cu
View file @
a87df3ad
...
@@ -97,11 +97,11 @@ extern "C" __global__ void computeBornSum(unsigned long long* __restrict__ bornS
...
@@ -97,11 +97,11 @@ extern "C" __global__ void computeBornSum(unsigned long long* __restrict__ bornS
// Extract the coordinates of this tile
// Extract the coordinates of this tile
const
unsigned
int
tgx
=
threadIdx
.
x
&
(
TILE_SIZE
-
1
);
const
unsigned
int
tgx
=
threadIdx
.
x
&
(
TILE_SIZE
-
1
);
const
unsigned
int
tbx
=
threadIdx
.
x
-
tgx
;
const
unsigned
int
tbx
=
threadIdx
.
x
-
tgx
;
unsigned
int
x
,
y
;
int
x
,
y
;
AtomData1
data
;
AtomData1
data
;
data
.
bornSum
=
0
;
data
.
bornSum
=
0
;
if
(
pos
<
end
)
{
if
(
pos
<
end
)
{
y
=
(
unsigned
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
sqrt
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
y
=
(
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
sqrt
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
y
+=
(
x
<
y
?
-
1
:
1
);
y
+=
(
x
<
y
?
-
1
:
1
);
...
@@ -232,10 +232,10 @@ extern "C" __global__ void computeGKForces(
...
@@ -232,10 +232,10 @@ extern "C" __global__ void computeGKForces(
// Extract the coordinates of this tile
// Extract the coordinates of this tile
const
unsigned
int
tgx
=
threadIdx
.
x
&
(
TILE_SIZE
-
1
);
const
unsigned
int
tgx
=
threadIdx
.
x
&
(
TILE_SIZE
-
1
);
const
unsigned
int
tbx
=
threadIdx
.
x
-
tgx
;
const
unsigned
int
tbx
=
threadIdx
.
x
-
tgx
;
unsigned
int
x
,
y
;
int
x
,
y
;
AtomData2
data
;
AtomData2
data
;
if
(
pos
<
end
)
{
if
(
pos
<
end
)
{
y
=
(
unsigned
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
y
=
(
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
y
+=
(
x
<
y
?
-
1
:
1
);
y
+=
(
x
<
y
?
-
1
:
1
);
...
@@ -471,10 +471,10 @@ extern "C" __global__ void computeChainRuleForce(
...
@@ -471,10 +471,10 @@ extern "C" __global__ void computeChainRuleForce(
// Extract the coordinates of this tile
// Extract the coordinates of this tile
const
unsigned
int
tgx
=
threadIdx
.
x
&
(
TILE_SIZE
-
1
);
const
unsigned
int
tgx
=
threadIdx
.
x
&
(
TILE_SIZE
-
1
);
const
unsigned
int
tbx
=
threadIdx
.
x
-
tgx
;
const
unsigned
int
tbx
=
threadIdx
.
x
-
tgx
;
unsigned
int
x
,
y
;
int
x
,
y
;
AtomData3
data
;
AtomData3
data
;
if
(
pos
<
end
)
{
if
(
pos
<
end
)
{
y
=
(
unsigned
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
y
=
(
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
y
+=
(
x
<
y
?
-
1
:
1
);
y
+=
(
x
<
y
?
-
1
:
1
);
...
@@ -766,8 +766,8 @@ extern "C" __global__ void computeEDiffForce(
...
@@ -766,8 +766,8 @@ extern "C" __global__ void computeEDiffForce(
while
(
pos
<
end
)
{
while
(
pos
<
end
)
{
// Extract the coordinates of this tile.
// Extract the coordinates of this tile.
unsigned
int
x
,
y
;
int
x
,
y
;
y
=
(
unsigned
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
y
=
(
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
y
+=
(
x
<
y
?
-
1
:
1
);
y
+=
(
x
<
y
?
-
1
:
1
);
...
...
plugins/amoeba/platforms/cuda/src/kernels/amoebaWcaForce.cu
View file @
a87df3ad
...
@@ -207,10 +207,10 @@ extern "C" __global__ void computeWCAForce(unsigned long long* __restrict__ forc
...
@@ -207,10 +207,10 @@ extern "C" __global__ void computeWCAForce(unsigned long long* __restrict__ forc
const
unsigned
int
tgx
=
threadIdx
.
x
&
(
TILE_SIZE
-
1
);
const
unsigned
int
tgx
=
threadIdx
.
x
&
(
TILE_SIZE
-
1
);
const
unsigned
int
tbx
=
threadIdx
.
x
-
tgx
;
const
unsigned
int
tbx
=
threadIdx
.
x
-
tgx
;
const
unsigned
int
localGroupIndex
=
threadIdx
.
x
/
TILE_SIZE
;
const
unsigned
int
localGroupIndex
=
threadIdx
.
x
/
TILE_SIZE
;
unsigned
int
x
,
y
;
int
x
,
y
;
AtomData
data
;
AtomData
data
;
if
(
pos
<
end
)
{
if
(
pos
<
end
)
{
y
=
(
unsigned
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
y
=
(
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
y
+=
(
x
<
y
?
-
1
:
1
);
y
+=
(
x
<
y
?
-
1
:
1
);
...
...
plugins/amoeba/platforms/cuda/src/kernels/multipoleElectrostatics.cu
View file @
a87df3ad
...
@@ -234,14 +234,14 @@ extern "C" __global__ void computeElectrostatics(
...
@@ -234,14 +234,14 @@ extern "C" __global__ void computeElectrostatics(
// Extract the coordinates of this tile.
// Extract the coordinates of this tile.
unsigned
int
x
,
y
;
int
x
,
y
;
#ifdef USE_CUTOFF
#ifdef USE_CUTOFF
if
(
numTiles
<=
maxTiles
)
if
(
numTiles
<=
maxTiles
)
x
=
tiles
[
pos
];
x
=
tiles
[
pos
];
else
else
#endif
#endif
{
{
y
=
(
unsigned
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
y
=
(
int
)
floor
(
NUM_BLOCKS
+
0.5
f
-
SQRT
((
NUM_BLOCKS
+
0.5
f
)
*
(
NUM_BLOCKS
+
0.5
f
)
-
2
*
pos
));
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
x
=
(
pos
-
y
*
NUM_BLOCKS
+
y
*
(
y
+
1
)
/
2
);
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
if
(
x
<
y
||
x
>=
NUM_BLOCKS
)
{
// Occasionally happens due to roundoff error.
y
+=
(
x
<
y
?
-
1
:
1
);
y
+=
(
x
<
y
?
-
1
:
1
);
...
...
Prev
1
2
Next
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
.
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment