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
53539843
Commit
53539843
authored
Jan 05, 2010
by
Peter Eastman
Browse files
Fixed compilation warnings under Windows
parent
c553fe60
Changes
7
Expand all
Hide whitespace changes
Inline
Side-by-side
Showing
7 changed files
with
436 additions
and
433 deletions
+436
-433
platforms/opencl/src/OpenCLCompact.cpp
platforms/opencl/src/OpenCLCompact.cpp
+1
-1
platforms/opencl/src/OpenCLContext.cpp
platforms/opencl/src/OpenCLContext.cpp
+9
-9
platforms/opencl/src/OpenCLContext.h
platforms/opencl/src/OpenCLContext.h
+365
-365
platforms/opencl/src/OpenCLExpressionUtilities.cpp
platforms/opencl/src/OpenCLExpressionUtilities.cpp
+5
-2
platforms/opencl/src/OpenCLIntegrationUtilities.cpp
platforms/opencl/src/OpenCLIntegrationUtilities.cpp
+5
-5
platforms/opencl/src/OpenCLKernels.cpp
platforms/opencl/src/OpenCLKernels.cpp
+38
-38
platforms/opencl/src/OpenCLNonbondedUtilities.cpp
platforms/opencl/src/OpenCLNonbondedUtilities.cpp
+13
-13
No files found.
platforms/opencl/src/OpenCLCompact.cpp
View file @
53539843
...
...
@@ -42,7 +42,7 @@ OpenCLCompact::~OpenCLCompact() {
void
OpenCLCompact
::
compactStream
(
OpenCLArray
<
cl_uint
>&
dOut
,
OpenCLArray
<
cl_uint
>&
dIn
,
OpenCLArray
<
cl_uint
>&
dValid
,
OpenCLArray
<
cl_uint
>&
numValid
)
{
// Figure out # elements per block
int
len
=
dIn
.
getSize
();
unsigned
int
len
=
dIn
.
getSize
();
unsigned
int
numBlocks
=
context
.
getNumThreadBlocks
();
if
(
numBlocks
*
128
>
len
)
numBlocks
=
(
len
+
127
)
/
128
;
...
...
platforms/opencl/src/OpenCLContext.cpp
View file @
53539843
...
...
@@ -46,11 +46,11 @@ OpenCLContext::OpenCLContext(int numParticles, int deviceIndex) : time(0.0), ste
context
=
cl
::
Context
(
CL_DEVICE_TYPE_ALL
);
vector
<
cl
::
Device
>
devices
=
context
.
getInfo
<
CL_CONTEXT_DEVICES
>
();
const
int
minThreadBlockSize
=
32
;
if
(
deviceIndex
<
0
||
deviceIndex
>=
devices
.
size
())
{
if
(
deviceIndex
<
0
||
deviceIndex
>=
(
int
)
devices
.
size
())
{
// Try to figure out which device is the fastest.
int
bestSpeed
=
0
;
for
(
int
i
=
0
;
i
<
devices
.
size
();
i
++
)
{
for
(
int
i
=
0
;
i
<
(
int
)
devices
.
size
();
i
++
)
{
int
maxSize
=
devices
[
i
].
getInfo
<
CL_DEVICE_MAX_WORK_ITEM_SIZES
>
()[
0
];
int
speed
=
devices
[
i
].
getInfo
<
CL_DEVICE_MAX_COMPUTE_UNITS
>
()
*
devices
[
i
].
getInfo
<
CL_DEVICE_MAX_CLOCK_FREQUENCY
>
();
if
(
maxSize
>=
minThreadBlockSize
&&
speed
>
bestSpeed
)
...
...
@@ -245,12 +245,12 @@ void OpenCLContext::findMoleculeGroups(const System& system) {
atomBonds
[
particle1
].
push_back
(
particle2
);
atomBonds
[
particle2
].
push_back
(
particle1
);
}
for
(
int
i
=
0
;
i
<
forces
.
size
();
i
++
)
{
for
(
int
i
=
0
;
i
<
(
int
)
forces
.
size
();
i
++
)
{
for
(
int
j
=
0
;
j
<
forces
[
i
]
->
getNumParticleGroups
();
j
++
)
{
vector
<
int
>
particles
;
forces
[
i
]
->
getParticlesInGroup
(
j
,
particles
);
for
(
int
k
=
0
;
k
<
particles
.
size
();
k
++
)
for
(
int
m
=
0
;
m
<
particles
.
size
();
m
++
)
for
(
int
k
=
0
;
k
<
(
int
)
particles
.
size
();
k
++
)
for
(
int
m
=
0
;
m
<
(
int
)
particles
.
size
();
m
++
)
if
(
k
!=
m
)
atomBonds
[
particles
[
k
]].
push_back
(
particles
[
m
]);
}
...
...
@@ -280,7 +280,7 @@ void OpenCLContext::findMoleculeGroups(const System& system) {
system
.
getConstraintParameters
(
i
,
particle1
,
particle2
,
distance
);
molecules
[
atomMolecule
[
particle1
]].
constraints
.
push_back
(
i
);
}
for
(
int
i
=
0
;
i
<
forces
.
size
();
i
++
)
for
(
int
i
=
0
;
i
<
(
int
)
forces
.
size
();
i
++
)
for
(
int
j
=
0
;
j
<
forces
[
i
]
->
getNumParticleGroups
();
j
++
)
{
vector
<
int
>
particles
;
forces
[
i
]
->
getParticlesInGroup
(
j
,
particles
);
...
...
@@ -307,7 +307,7 @@ void OpenCLContext::findMoleculeGroups(const System& system) {
for
(
int
i
=
0
;
i
<
(
int
)
mol
.
atoms
.
size
()
&&
identical
;
i
++
)
{
if
(
mol
.
atoms
[
i
]
!=
mol2
.
atoms
[
i
]
-
atomOffset
||
system
.
getParticleMass
(
mol
.
atoms
[
i
])
!=
system
.
getParticleMass
(
mol2
.
atoms
[
i
]))
identical
=
false
;
for
(
int
k
=
0
;
k
<
forces
.
size
();
k
++
)
for
(
int
k
=
0
;
k
<
(
int
)
forces
.
size
();
k
++
)
if
(
!
forces
[
k
]
->
areParticlesIdentical
(
mol
.
atoms
[
i
],
mol2
.
atoms
[
i
]))
identical
=
false
;
}
...
...
@@ -325,10 +325,10 @@ void OpenCLContext::findMoleculeGroups(const System& system) {
// See if the force groups are identical.
for
(
int
i
=
0
;
i
<
forces
.
size
()
&&
identical
;
i
++
)
{
for
(
int
i
=
0
;
i
<
(
int
)
forces
.
size
()
&&
identical
;
i
++
)
{
if
(
mol
.
groups
[
i
].
size
()
!=
mol2
.
groups
[
i
].
size
())
identical
=
false
;
for
(
int
k
=
0
;
k
<
mol
.
groups
[
i
].
size
()
&&
identical
;
k
++
)
for
(
int
k
=
0
;
k
<
(
int
)
mol
.
groups
[
i
].
size
()
&&
identical
;
k
++
)
if
(
!
forces
[
i
]
->
areGroupsIdentical
(
mol
.
groups
[
i
][
k
],
mol2
.
groups
[
i
][
k
]))
identical
=
false
;
}
...
...
platforms/opencl/src/OpenCLContext.h
View file @
53539843
This diff is collapsed.
Click to expand it.
platforms/opencl/src/OpenCLExpressionUtilities.cpp
View file @
53539843
...
...
@@ -260,7 +260,7 @@ void OpenCLExpressionUtilities::processExpression(stringstream& out, const Expre
bool
done
=
false
;
while
(
!
done
)
{
done
=
true
;
for
(
int
i
=
0
;
i
<
exponents
.
size
();
i
++
)
{
for
(
int
i
=
0
;
i
<
(
int
)
exponents
.
size
();
i
++
)
{
if
(
exponents
[
i
]
%
2
==
1
)
{
if
(
!
hasAssigned
[
i
])
out
<<
names
[
i
]
<<
" = multiplier;
\n
"
;
...
...
@@ -313,7 +313,10 @@ void OpenCLExpressionUtilities::findRelatedTabulatedFunctions(const ExpressionTr
void
OpenCLExpressionUtilities
::
findRelatedPowers
(
const
ExpressionTreeNode
&
node
,
const
ExpressionTreeNode
&
searchNode
,
map
<
int
,
const
ExpressionTreeNode
*>&
powers
)
{
if
(
searchNode
.
getOperation
().
getId
()
==
Operation
::
POWER_CONSTANT
&&
node
.
getChildren
()[
0
]
==
searchNode
.
getChildren
()[
0
])
{
int
power
=
dynamic_cast
<
const
Operation
::
PowerConstant
*>
(
&
searchNode
.
getOperation
())
->
getValue
();
double
realPower
=
dynamic_cast
<
const
Operation
::
PowerConstant
*>
(
&
searchNode
.
getOperation
())
->
getValue
();
int
power
=
(
int
)
realPower
;
if
(
power
!=
realPower
)
return
;
// We are only interested in integer powers.
if
(
powers
.
find
(
power
)
!=
powers
.
end
())
return
;
// This power is already in the map.
if
(
powers
.
begin
()
->
first
*
power
<
0
)
...
...
platforms/opencl/src/OpenCLIntegrationUtilities.cpp
View file @
53539843
...
...
@@ -103,8 +103,8 @@ OpenCLIntegrationUtilities::OpenCLIntegrationUtilities(OpenCLContext& context, c
vector
<
map
<
int
,
float
>
>
settleConstraints
(
system
.
getNumParticles
());
for
(
int
i
=
0
;
i
<
(
int
)
atom1
.
size
();
i
++
)
{
if
(
constraintCount
[
atom1
[
i
]]
==
2
&&
constraintCount
[
atom2
[
i
]]
==
2
)
{
settleConstraints
[
atom1
[
i
]][
atom2
[
i
]]
=
distance
[
i
];
settleConstraints
[
atom2
[
i
]][
atom1
[
i
]]
=
distance
[
i
];
settleConstraints
[
atom1
[
i
]][
atom2
[
i
]]
=
(
float
)
distance
[
i
];
settleConstraints
[
atom2
[
i
]][
atom1
[
i
]]
=
(
float
)
distance
[
i
];
}
}
...
...
@@ -232,7 +232,7 @@ OpenCLIntegrationUtilities::OpenCLIntegrationUtilities(OpenCLContext& context, c
if
(
!
cluster
.
valid
)
continue
;
atoms
.
push_back
(
mm_int4
(
cluster
.
centralID
,
cluster
.
peripheralID
[
0
],
(
cluster
.
size
>
1
?
cluster
.
peripheralID
[
1
]
:
-
1
),
(
cluster
.
size
>
2
?
cluster
.
peripheralID
[
2
]
:
-
1
)));
params
.
push_back
(
mm_float4
(
cluster
.
centralInvMass
,
0.5
f
/
(
cluster
.
centralInvMass
+
cluster
.
peripheralInvMass
),
cluster
.
distance
*
cluster
.
distance
,
cluster
.
peripheralInvMass
));
params
.
push_back
(
mm_float4
(
(
cl_float
)
cluster
.
centralInvMass
,
(
cl_float
)
(
0.5
/
(
cluster
.
centralInvMass
+
cluster
.
peripheralInvMass
)
)
,
(
cl_float
)
(
cluster
.
distance
*
cluster
.
distance
),
(
cl_float
)
cluster
.
peripheralInvMass
));
isShakeAtom
[
cluster
.
centralID
]
=
true
;
isShakeAtom
[
cluster
.
peripheralID
[
0
]]
=
true
;
if
(
cluster
.
size
>
1
)
...
...
@@ -270,7 +270,7 @@ OpenCLIntegrationUtilities::~OpenCLIntegrationUtilities() {
void
OpenCLIntegrationUtilities
::
applyConstraints
(
double
tol
)
{
if
(
settleAtoms
!=
NULL
)
{
settleKernel
.
setArg
<
cl_int
>
(
0
,
settleAtoms
->
getSize
());
settleKernel
.
setArg
<
cl_float
>
(
1
,
tol
);
settleKernel
.
setArg
<
cl_float
>
(
1
,
(
cl_float
)
tol
);
settleKernel
.
setArg
<
cl
::
Buffer
>
(
2
,
context
.
getPosq
().
getDeviceBuffer
());
settleKernel
.
setArg
<
cl
::
Buffer
>
(
3
,
posDelta
->
getDeviceBuffer
());
settleKernel
.
setArg
<
cl
::
Buffer
>
(
4
,
posDelta
->
getDeviceBuffer
());
...
...
@@ -281,7 +281,7 @@ void OpenCLIntegrationUtilities::applyConstraints(double tol) {
}
if
(
shakeAtoms
!=
NULL
)
{
shakeKernel
.
setArg
<
cl_int
>
(
0
,
shakeAtoms
->
getSize
());
shakeKernel
.
setArg
<
cl_float
>
(
1
,
tol
);
shakeKernel
.
setArg
<
cl_float
>
(
1
,
(
cl_float
)
tol
);
shakeKernel
.
setArg
<
cl
::
Buffer
>
(
2
,
context
.
getPosq
().
getDeviceBuffer
());
shakeKernel
.
setArg
<
cl
::
Buffer
>
(
3
,
posDelta
->
getDeviceBuffer
());
shakeKernel
.
setArg
<
cl
::
Buffer
>
(
4
,
posDelta
->
getDeviceBuffer
());
...
...
platforms/opencl/src/OpenCLKernels.cpp
View file @
53539843
...
...
@@ -120,12 +120,12 @@ void OpenCLUpdateStateDataKernel::setPositions(ContextImpl& context, const std::
for
(
int
i
=
0
;
i
<
numParticles
;
++
i
)
{
mm_float4
&
pos
=
posq
[
i
];
const
Vec3
&
p
=
positions
[
order
[
i
]];
pos
.
x
=
p
[
0
];
pos
.
y
=
p
[
1
];
pos
.
z
=
p
[
2
];
pos
.
x
=
(
cl_float
)
p
[
0
];
pos
.
y
=
(
cl_float
)
p
[
1
];
pos
.
z
=
(
cl_float
)
p
[
2
];
}
posq
.
upload
();
for
(
int
i
=
0
;
i
<
cl
.
getPosCellOffsets
().
size
();
i
++
)
for
(
int
i
=
0
;
i
<
(
int
)
cl
.
getPosCellOffsets
().
size
();
i
++
)
cl
.
getPosCellOffsets
()[
i
]
=
mm_int4
(
0
,
0
,
0
,
0
);
}
...
...
@@ -148,9 +148,9 @@ void OpenCLUpdateStateDataKernel::setVelocities(ContextImpl& context, const std:
for
(
int
i
=
0
;
i
<
numParticles
;
++
i
)
{
mm_float4
&
vel
=
velm
[
i
];
const
Vec3
&
p
=
velocities
[
order
[
i
]];
vel
.
x
=
p
[
0
];
vel
.
y
=
p
[
1
];
vel
.
z
=
p
[
2
];
vel
.
x
=
(
cl_float
)
p
[
0
];
vel
.
y
=
(
cl_float
)
p
[
1
];
vel
.
z
=
(
cl_float
)
p
[
2
];
}
velm
.
upload
();
}
...
...
@@ -219,7 +219,7 @@ void OpenCLCalcHarmonicBondForceKernel::initialize(const System& system, const H
params
->
upload
(
paramVector
);
indices
->
upload
(
indicesVector
);
int
maxBuffers
=
1
;
for
(
int
i
=
0
;
i
<
forceBufferCounter
.
size
();
i
++
)
for
(
int
i
=
0
;
i
<
(
int
)
forceBufferCounter
.
size
();
i
++
)
maxBuffers
=
max
(
maxBuffers
,
forceBufferCounter
[
i
]);
cl
.
addForce
(
new
OpenCLBondForceInfo
(
maxBuffers
,
force
));
cl
::
Program
program
=
cl
.
createProgram
(
cl
.
loadSourceFromFile
(
"harmonicBondForce.cl"
));
...
...
@@ -304,14 +304,14 @@ void OpenCLCalcCustomBondForceKernel::initialize(const System& system, const Cus
vector
<
double
>
parameters
;
force
.
getBondParameters
(
i
,
particle1
,
particle2
,
parameters
);
paramVector
[
i
].
resize
(
parameters
.
size
());
for
(
int
j
=
0
;
j
<
parameters
.
size
();
j
++
)
for
(
int
j
=
0
;
j
<
(
int
)
parameters
.
size
();
j
++
)
paramVector
[
i
][
j
]
=
(
cl_float
)
parameters
[
j
];
indicesVector
[
i
]
=
mm_int4
(
particle1
,
particle2
,
forceBufferCounter
[
particle1
]
++
,
forceBufferCounter
[
particle2
]
++
);
}
params
->
setParameterValues
(
paramVector
);
indices
->
upload
(
indicesVector
);
int
maxBuffers
=
1
;
for
(
int
i
=
0
;
i
<
forceBufferCounter
.
size
();
i
++
)
for
(
int
i
=
0
;
i
<
(
int
)
forceBufferCounter
.
size
();
i
++
)
maxBuffers
=
max
(
maxBuffers
,
forceBufferCounter
[
i
]);
cl
.
addForce
(
new
OpenCLCustomBondForceInfo
(
maxBuffers
,
force
));
...
...
@@ -367,7 +367,7 @@ void OpenCLCalcCustomBondForceKernel::executeForces(ContextImpl& context) {
return
;
if
(
globals
!=
NULL
)
{
bool
changed
=
false
;
for
(
int
i
=
0
;
i
<
globalParamNames
.
size
();
i
++
)
{
for
(
int
i
=
0
;
i
<
(
int
)
globalParamNames
.
size
();
i
++
)
{
cl_float
value
=
(
cl_float
)
context
.
getParameter
(
globalParamNames
[
i
]);
if
(
value
!=
globalParamValues
[
i
])
changed
=
true
;
...
...
@@ -455,7 +455,7 @@ void OpenCLCalcHarmonicAngleForceKernel::initialize(const System& system, const
params
->
upload
(
paramVector
);
indices
->
upload
(
indicesVector
);
int
maxBuffers
=
1
;
for
(
int
i
=
0
;
i
<
forceBufferCounter
.
size
();
i
++
)
for
(
int
i
=
0
;
i
<
(
int
)
forceBufferCounter
.
size
();
i
++
)
maxBuffers
=
max
(
maxBuffers
,
forceBufferCounter
[
i
]);
cl
.
addForce
(
new
OpenCLAngleForceInfo
(
maxBuffers
,
force
));
cl
::
Program
program
=
cl
.
createProgram
(
cl
.
loadSourceFromFile
(
"harmonicAngleForce.cl"
));
...
...
@@ -539,7 +539,7 @@ void OpenCLCalcPeriodicTorsionForceKernel::initialize(const System& system, cons
params
->
upload
(
paramVector
);
indices
->
upload
(
indicesVector
);
int
maxBuffers
=
1
;
for
(
int
i
=
0
;
i
<
forceBufferCounter
.
size
();
i
++
)
for
(
int
i
=
0
;
i
<
(
int
)
forceBufferCounter
.
size
();
i
++
)
maxBuffers
=
max
(
maxBuffers
,
forceBufferCounter
[
i
]);
cl
.
addForce
(
new
OpenCLPeriodicTorsionForceInfo
(
maxBuffers
,
force
));
cl
::
Program
program
=
cl
.
createProgram
(
cl
.
loadSourceFromFile
(
"periodicTorsionForce.cl"
));
...
...
@@ -623,7 +623,7 @@ void OpenCLCalcRBTorsionForceKernel::initialize(const System& system, const RBTo
params
->
upload
(
paramVector
);
indices
->
upload
(
indicesVector
);
int
maxBuffers
=
1
;
for
(
int
i
=
0
;
i
<
forceBufferCounter
.
size
();
i
++
)
for
(
int
i
=
0
;
i
<
(
int
)
forceBufferCounter
.
size
();
i
++
)
maxBuffers
=
max
(
maxBuffers
,
forceBufferCounter
[
i
]);
cl
.
addForce
(
new
OpenCLRBTorsionForceInfo
(
maxBuffers
,
force
));
cl
::
Program
program
=
cl
.
createProgram
(
cl
.
loadSourceFromFile
(
"rbTorsionForce.cl"
));
...
...
@@ -809,7 +809,7 @@ void OpenCLCalcNonbondedForceKernel::initialize(const System& system, const Nonb
}
exceptionParams
->
upload
(
exceptionParamsVector
);
exceptionIndices
->
upload
(
exceptionIndicesVector
);
for
(
int
i
=
0
;
i
<
forceBufferCounter
.
size
();
i
++
)
for
(
int
i
=
0
;
i
<
(
int
)
forceBufferCounter
.
size
();
i
++
)
maxBuffers
=
max
(
maxBuffers
,
forceBufferCounter
[
i
]);
}
cl
.
addForce
(
new
OpenCLNonbondedForceInfo
(
maxBuffers
,
force
));
...
...
@@ -829,7 +829,7 @@ void OpenCLCalcNonbondedForceKernel::executeForces(ContextImpl& context) {
int
numExceptions
=
exceptionIndices
->
getSize
();
exceptionsKernel
.
setArg
<
cl_int
>
(
0
,
cl
.
getPaddedNumAtoms
());
exceptionsKernel
.
setArg
<
cl_int
>
(
1
,
numExceptions
);
exceptionsKernel
.
setArg
<
cl_float
>
(
2
,
cutoffSquared
);
exceptionsKernel
.
setArg
<
cl_float
>
(
2
,
(
cl_float
)
cutoffSquared
);
exceptionsKernel
.
setArg
<
mm_float4
>
(
3
,
cl
.
getNonbondedUtilities
().
getPeriodicBoxSize
());
exceptionsKernel
.
setArg
<
cl
::
Buffer
>
(
4
,
cl
.
getForceBuffers
().
getDeviceBuffer
());
exceptionsKernel
.
setArg
<
cl
::
Buffer
>
(
5
,
cl
.
getEnergyBuffer
().
getDeviceBuffer
());
...
...
@@ -868,7 +868,7 @@ public:
vector
<
double
>
params2
;
force
.
getParticleParameters
(
particle1
,
params1
);
force
.
getParticleParameters
(
particle2
,
params2
);
for
(
int
i
=
0
;
i
<
params1
.
size
();
i
++
)
for
(
int
i
=
0
;
i
<
(
int
)
params1
.
size
();
i
++
)
if
(
params1
[
i
]
!=
params2
[
i
])
return
false
;
return
true
;
...
...
@@ -919,7 +919,7 @@ void OpenCLCalcCustomNonbondedForceKernel::initialize(const System& system, cons
vector
<
double
>
parameters
;
force
.
getParticleParameters
(
i
,
parameters
);
paramVector
[
i
].
resize
(
parameters
.
size
());
for
(
int
j
=
0
;
j
<
parameters
.
size
();
j
++
)
for
(
int
j
=
0
;
j
<
(
int
)
parameters
.
size
();
j
++
)
paramVector
[
i
][
j
]
=
(
cl_float
)
parameters
[
j
];
exclusionList
[
i
].
push_back
(
i
);
}
...
...
@@ -1013,7 +1013,7 @@ void OpenCLCalcCustomNonbondedForceKernel::initialize(const System& system, cons
void
OpenCLCalcCustomNonbondedForceKernel
::
executeForces
(
ContextImpl
&
context
)
{
if
(
globals
!=
NULL
)
{
bool
changed
=
false
;
for
(
int
i
=
0
;
i
<
globalParamNames
.
size
();
i
++
)
{
for
(
int
i
=
0
;
i
<
(
int
)
globalParamNames
.
size
();
i
++
)
{
cl_float
value
=
(
cl_float
)
context
.
getParameter
(
globalParamNames
[
i
]);
if
(
value
!=
globalParamValues
[
i
])
changed
=
true
;
...
...
@@ -1187,7 +1187,7 @@ public:
vector
<
double
>
params2
;
force
.
getParticleParameters
(
particle1
,
params1
);
force
.
getParticleParameters
(
particle2
,
params2
);
for
(
int
i
=
0
;
i
<
params1
.
size
();
i
++
)
for
(
int
i
=
0
;
i
<
(
int
)
params1
.
size
();
i
++
)
if
(
params1
[
i
]
!=
params2
[
i
])
return
false
;
return
true
;
...
...
@@ -1260,7 +1260,7 @@ void OpenCLCalcCustomGBForceKernel::initialize(const System& system, const Custo
vector
<
double
>
parameters
;
force
.
getParticleParameters
(
i
,
parameters
);
paramVector
[
i
].
resize
(
parameters
.
size
());
for
(
int
j
=
0
;
j
<
parameters
.
size
();
j
++
)
for
(
int
j
=
0
;
j
<
(
int
)
parameters
.
size
();
j
++
)
paramVector
[
i
][
j
]
=
(
cl_float
)
parameters
[
j
];
exclusionList
[
i
].
push_back
(
i
);
}
...
...
@@ -1807,7 +1807,7 @@ void OpenCLCalcCustomGBForceKernel::executeForces(ContextImpl& context) {
}
if
(
globals
!=
NULL
)
{
bool
changed
=
false
;
for
(
int
i
=
0
;
i
<
globalParamNames
.
size
();
i
++
)
{
for
(
int
i
=
0
;
i
<
(
int
)
globalParamNames
.
size
();
i
++
)
{
cl_float
value
=
(
cl_float
)
context
.
getParameter
(
globalParamNames
[
i
]);
if
(
value
!=
globalParamValues
[
i
])
changed
=
true
;
...
...
@@ -1854,7 +1854,7 @@ public:
vector
<
double
>
params2
;
force
.
getParticleParameters
(
particle1
,
temp
,
params1
);
force
.
getParticleParameters
(
particle2
,
temp
,
params2
);
for
(
int
i
=
0
;
i
<
params1
.
size
();
i
++
)
for
(
int
i
=
0
;
i
<
(
int
)
params1
.
size
();
i
++
)
if
(
params1
[
i
]
!=
params2
[
i
])
return
false
;
return
true
;
...
...
@@ -1888,7 +1888,7 @@ void OpenCLCalcCustomExternalForceKernel::initialize(const System& system, const
vector
<
double
>
parameters
;
force
.
getParticleParameters
(
i
,
indicesVector
[
i
],
parameters
);
paramVector
[
i
].
resize
(
parameters
.
size
());
for
(
int
j
=
0
;
j
<
parameters
.
size
();
j
++
)
for
(
int
j
=
0
;
j
<
(
int
)
parameters
.
size
();
j
++
)
paramVector
[
i
][
j
]
=
(
cl_float
)
parameters
[
j
];
}
params
->
setParameterValues
(
paramVector
);
...
...
@@ -1951,7 +1951,7 @@ void OpenCLCalcCustomExternalForceKernel::initialize(const System& system, const
void
OpenCLCalcCustomExternalForceKernel
::
executeForces
(
ContextImpl
&
context
)
{
if
(
globals
!=
NULL
)
{
bool
changed
=
false
;
for
(
int
i
=
0
;
i
<
globalParamNames
.
size
();
i
++
)
{
for
(
int
i
=
0
;
i
<
(
int
)
globalParamNames
.
size
();
i
++
)
{
cl_float
value
=
(
cl_float
)
context
.
getParameter
(
globalParamNames
[
i
]);
if
(
value
!=
globalParamValues
[
i
])
changed
=
true
;
...
...
@@ -2014,7 +2014,7 @@ void OpenCLIntegrateVerletStepKernel::execute(ContextImpl& context, const Verlet
}
if
(
dt
!=
prevStepSize
)
{
vector
<
mm_float2
>
stepSizeVec
(
1
);
stepSizeVec
[
0
]
=
mm_float2
(
dt
,
dt
);
stepSizeVec
[
0
]
=
mm_float2
(
(
cl_float
)
dt
,
(
cl_float
)
dt
);
cl
.
getIntegrationUtilities
().
getStepSize
().
upload
(
stepSizeVec
);
prevStepSize
=
dt
;
}
...
...
@@ -2136,17 +2136,17 @@ void OpenCLIntegrateLangevinStepKernel::execute(ContextImpl& context, const Lang
double
Yv
=
sqrt
(
kT
*
B
/
C
);
double
Yx
=
tau
*
sqrt
(
kT
*
B
/
(
1.0
-
EM
));
vector
<
cl_float
>
p
(
params
->
getSize
());
p
[
0
]
=
EM
;
p
[
1
]
=
EM
;
p
[
2
]
=
DOverTauC
;
p
[
3
]
=
TauOneMinusEM
;
p
[
4
]
=
TauDOverEMMinusOne
;
p
[
5
]
=
V
;
p
[
6
]
=
X
;
p
[
7
]
=
Yv
;
p
[
8
]
=
Yx
;
p
[
9
]
=
fix1
;
p
[
10
]
=
oneOverFix1
;
p
[
0
]
=
(
cl_float
)
EM
;
p
[
1
]
=
(
cl_float
)
EM
;
p
[
2
]
=
(
cl_float
)
DOverTauC
;
p
[
3
]
=
(
cl_float
)
TauOneMinusEM
;
p
[
4
]
=
(
cl_float
)
TauDOverEMMinusOne
;
p
[
5
]
=
(
cl_float
)
V
;
p
[
6
]
=
(
cl_float
)
X
;
p
[
7
]
=
(
cl_float
)
Yv
;
p
[
8
]
=
(
cl_float
)
Yx
;
p
[
9
]
=
(
cl_float
)
fix1
;
p
[
10
]
=
(
cl_float
)
oneOverFix1
;
params
->
upload
(
p
);
prevTemp
=
temperature
;
prevFriction
=
friction
;
...
...
@@ -2278,7 +2278,7 @@ void OpenCLIntegrateVariableVerletStepKernel::execute(ContextImpl& context, cons
float
maxStepSize
=
(
float
)(
maxTime
-
cl
.
getTime
());
selectSizeKernel
.
setArg
<
cl_float
>
(
1
,
maxStepSize
);
selectSizeKernel
.
setArg
<
cl_float
>
(
2
,
integrator
.
getErrorTolerance
());
selectSizeKernel
.
setArg
<
cl_float
>
(
2
,
(
cl_float
)
integrator
.
getErrorTolerance
());
cl
.
executeKernel
(
selectSizeKernel
,
blockSize
,
blockSize
);
// Call the first integration kernel.
...
...
platforms/opencl/src/OpenCLNonbondedUtilities.cpp
View file @
53539843
...
...
@@ -79,10 +79,10 @@ void OpenCLNonbondedUtilities::addInteraction(bool usesCutoff, bool usesPeriodic
}
if
(
usesExclusions
&&
atomExclusions
.
size
()
!=
0
)
{
bool
sameExclusions
=
(
exclusionList
.
size
()
==
atomExclusions
.
size
());
for
(
int
i
=
0
;
i
<
exclusionList
.
size
()
&&
sameExclusions
;
i
++
)
{
for
(
int
i
=
0
;
i
<
(
int
)
exclusionList
.
size
()
&&
sameExclusions
;
i
++
)
{
if
(
exclusionList
[
i
].
size
()
!=
atomExclusions
[
i
].
size
())
sameExclusions
=
false
;
for
(
int
j
=
0
;
j
<
exclusionList
[
i
].
size
();
j
++
)
for
(
int
j
=
0
;
j
<
(
int
)
exclusionList
[
i
].
size
();
j
++
)
if
(
exclusionList
[
i
][
j
]
!=
atomExclusions
[
i
][
j
])
sameExclusions
=
false
;
}
...
...
@@ -115,7 +115,7 @@ void OpenCLNonbondedUtilities::initialize(const System& system) {
// No exclusions were specifically requested, so just mark every atom as not interacting with itself.
atomExclusions
.
resize
(
context
.
getNumAtoms
());
for
(
int
i
=
0
;
i
<
atomExclusions
.
size
();
i
++
)
for
(
int
i
=
0
;
i
<
(
int
)
atomExclusions
.
size
();
i
++
)
atomExclusions
[
i
].
push_back
(
i
);
}
...
...
@@ -126,8 +126,8 @@ void OpenCLNonbondedUtilities::initialize(const System& system) {
tiles
=
new
OpenCLArray
<
cl_uint
>
(
context
,
numTiles
,
"tiles"
);
vector
<
cl_uint
>
tileVec
(
tiles
->
getSize
());
unsigned
int
count
=
0
;
for
(
unsigned
int
y
=
0
;
y
<
numAtomBlocks
;
y
++
)
for
(
unsigned
int
x
=
y
;
x
<
numAtomBlocks
;
x
++
)
for
(
unsigned
int
y
=
0
;
y
<
(
unsigned
int
)
numAtomBlocks
;
y
++
)
for
(
unsigned
int
x
=
y
;
x
<
(
unsigned
int
)
numAtomBlocks
;
x
++
)
tileVec
[
count
++
]
=
(
x
<<
17
)
|
(
y
<<
2
);
// Mark which tiles have exclusions.
...
...
@@ -242,14 +242,14 @@ void OpenCLNonbondedUtilities::initialize(const System& system) {
findBlockBoundsKernel
.
setArg
<
cl
::
Buffer
>
(
4
,
blockBoundingBox
->
getDeviceBuffer
());
findInteractingBlocksKernel
=
cl
::
Kernel
(
interactingBlocksProgram
,
"findBlocksWithInteractions"
);
findInteractingBlocksKernel
.
setArg
<
cl_int
>
(
0
,
tiles
->
getSize
());
findInteractingBlocksKernel
.
setArg
<
cl_float
>
(
1
,
cutoff
*
cutoff
);
findInteractingBlocksKernel
.
setArg
<
cl_float
>
(
1
,
(
cl_float
)
(
cutoff
*
cutoff
)
)
;
findInteractingBlocksKernel
.
setArg
<
mm_float4
>
(
2
,
periodicBoxSize
);
findInteractingBlocksKernel
.
setArg
<
cl
::
Buffer
>
(
3
,
tiles
->
getDeviceBuffer
());
findInteractingBlocksKernel
.
setArg
<
cl
::
Buffer
>
(
4
,
blockCenter
->
getDeviceBuffer
());
findInteractingBlocksKernel
.
setArg
<
cl
::
Buffer
>
(
5
,
blockBoundingBox
->
getDeviceBuffer
());
findInteractingBlocksKernel
.
setArg
<
cl
::
Buffer
>
(
6
,
interactionFlags
->
getDeviceBuffer
());
findInteractionsWithinBlocksKernel
=
cl
::
Kernel
(
interactingBlocksProgram
,
"findInteractionsWithinBlocks"
);
findInteractionsWithinBlocksKernel
.
setArg
<
cl_float
>
(
0
,
cutoff
*
cutoff
);
findInteractionsWithinBlocksKernel
.
setArg
<
cl_float
>
(
0
,
(
cl_float
)
(
cutoff
*
cutoff
)
)
;
findInteractionsWithinBlocksKernel
.
setArg
<
mm_float4
>
(
1
,
periodicBoxSize
);
findInteractionsWithinBlocksKernel
.
setArg
<
cl
::
Buffer
>
(
2
,
context
.
getPosq
().
getDeviceBuffer
());
findInteractionsWithinBlocksKernel
.
setArg
<
cl
::
Buffer
>
(
3
,
interactingTiles
->
getDeviceBuffer
());
...
...
@@ -282,7 +282,7 @@ cl::Kernel OpenCLNonbondedUtilities::createInteractionKernel(const string& sourc
map
<
string
,
string
>
replacements
;
replacements
[
"COMPUTE_INTERACTION"
]
=
source
;
stringstream
args
;
for
(
int
i
=
0
;
i
<
params
.
size
();
i
++
)
{
for
(
int
i
=
0
;
i
<
(
int
)
params
.
size
();
i
++
)
{
args
<<
", __global "
;
args
<<
params
[
i
].
getType
();
args
<<
"* global_"
;
...
...
@@ -292,7 +292,7 @@ cl::Kernel OpenCLNonbondedUtilities::createInteractionKernel(const string& sourc
args
<<
"* local_"
;
args
<<
params
[
i
].
getName
();
}
for
(
int
i
=
0
;
i
<
arguments
.
size
();
i
++
)
{
for
(
int
i
=
0
;
i
<
(
int
)
arguments
.
size
();
i
++
)
{
if
((
arguments
[
i
].
getBuffer
().
getInfo
<
CL_MEM_FLAGS
>
()
&
CL_MEM_READ_ONLY
)
==
0
)
args
<<
", __global "
;
else
...
...
@@ -303,7 +303,7 @@ cl::Kernel OpenCLNonbondedUtilities::createInteractionKernel(const string& sourc
}
replacements
[
"PARAMETER_ARGUMENTS"
]
=
args
.
str
();
stringstream
loadLocal1
;
for
(
int
i
=
0
;
i
<
params
.
size
();
i
++
)
{
for
(
int
i
=
0
;
i
<
(
int
)
params
.
size
();
i
++
)
{
loadLocal1
<<
"local_"
;
loadLocal1
<<
params
[
i
].
getName
();
loadLocal1
<<
"[get_local_id(0)] = "
;
...
...
@@ -312,7 +312,7 @@ cl::Kernel OpenCLNonbondedUtilities::createInteractionKernel(const string& sourc
}
replacements
[
"LOAD_LOCAL_PARAMETERS_FROM_1"
]
=
loadLocal1
.
str
();
stringstream
loadLocal2
;
for
(
int
i
=
0
;
i
<
params
.
size
();
i
++
)
{
for
(
int
i
=
0
;
i
<
(
int
)
params
.
size
();
i
++
)
{
loadLocal2
<<
"local_"
;
loadLocal2
<<
params
[
i
].
getName
();
loadLocal2
<<
"[get_local_id(0)] = global_"
;
...
...
@@ -321,7 +321,7 @@ cl::Kernel OpenCLNonbondedUtilities::createInteractionKernel(const string& sourc
}
replacements
[
"LOAD_LOCAL_PARAMETERS_FROM_GLOBAL"
]
=
loadLocal2
.
str
();
stringstream
load1
;
for
(
int
i
=
0
;
i
<
params
.
size
();
i
++
)
{
for
(
int
i
=
0
;
i
<
(
int
)
params
.
size
();
i
++
)
{
load1
<<
params
[
i
].
getType
();
load1
<<
" "
;
load1
<<
params
[
i
].
getName
();
...
...
@@ -331,7 +331,7 @@ cl::Kernel OpenCLNonbondedUtilities::createInteractionKernel(const string& sourc
}
replacements
[
"LOAD_ATOM1_PARAMETERS"
]
=
load1
.
str
();
stringstream
load2j
;
for
(
int
i
=
0
;
i
<
params
.
size
();
i
++
)
{
for
(
int
i
=
0
;
i
<
(
int
)
params
.
size
();
i
++
)
{
load2j
<<
params
[
i
].
getType
();
load2j
<<
" "
;
load2j
<<
params
[
i
].
getName
();
...
...
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