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
a39126ab
Commit
a39126ab
authored
Jan 06, 2010
by
Peter Eastman
Browse files
Fixed compilation warnings under Windows
parent
53539843
Changes
4
Show whitespace changes
Inline
Side-by-side
Showing
4 changed files
with
41 additions
and
41 deletions
+41
-41
platforms/cuda/src/CudaKernels.cpp
platforms/cuda/src/CudaKernels.cpp
+9
-9
platforms/cuda/src/kernels/bbsort.cu
platforms/cuda/src/kernels/bbsort.cu
+2
-2
platforms/cuda/src/kernels/gpu.cpp
platforms/cuda/src/kernels/gpu.cpp
+29
-29
platforms/cuda/src/kernels/kCalculatePME.cu
platforms/cuda/src/kernels/kCalculatePME.cu
+1
-1
No files found.
platforms/cuda/src/CudaKernels.cpp
View file @
a39126ab
...
...
@@ -130,12 +130,12 @@ void CudaUpdateStateDataKernel::setPositions(ContextImpl& context, const std::ve
for
(
int
i
=
0
;
i
<
numParticles
;
++
i
)
{
float4
&
pos
=
(
*
gpu
->
psPosq4
)[
i
];
const
Vec3
&
p
=
positions
[
order
[
i
]];
pos
.
x
=
p
[
0
];
pos
.
y
=
p
[
1
];
pos
.
z
=
p
[
2
];
pos
.
x
=
(
float
)
p
[
0
];
pos
.
y
=
(
float
)
p
[
1
];
pos
.
z
=
(
float
)
p
[
2
];
}
gpu
->
psPosq4
->
Upload
();
for
(
int
i
=
0
;
i
<
gpu
->
posCellOffsets
.
size
();
i
++
)
for
(
int
i
=
0
;
i
<
(
int
)
gpu
->
posCellOffsets
.
size
();
i
++
)
gpu
->
posCellOffsets
[
i
]
=
make_int3
(
0
,
0
,
0
);
}
...
...
@@ -158,9 +158,9 @@ void CudaUpdateStateDataKernel::setVelocities(ContextImpl& context, const std::v
for
(
int
i
=
0
;
i
<
numParticles
;
++
i
)
{
float4
&
vel
=
(
*
gpu
->
psVelm4
)[
i
];
const
Vec3
&
v
=
velocities
[
order
[
i
]];
vel
.
x
=
v
[
0
];
vel
.
y
=
v
[
1
];
vel
.
z
=
v
[
2
];
vel
.
x
=
(
float
)
v
[
0
];
vel
.
y
=
(
float
)
v
[
1
];
vel
.
z
=
(
float
)
v
[
2
];
}
gpu
->
psVelm4
->
Upload
();
}
...
...
@@ -392,7 +392,7 @@ void CudaCalcNonbondedForceKernel::initialize(const System& system, const Nonbon
gpuSetPeriodicBoxSize
(
gpu
,
(
float
)
boxVectors
[
0
][
0
],
(
float
)
boxVectors
[
1
][
1
],
(
float
)
boxVectors
[
2
][
2
]);
CudaNonbondedMethod
method
=
NO_CUTOFF
;
if
(
force
.
getNonbondedMethod
()
!=
NonbondedForce
::
NoCutoff
)
{
gpuSetNonbondedCutoff
(
gpu
,
(
float
)
force
.
getCutoffDistance
(),
force
.
getReactionFieldDielectric
());
gpuSetNonbondedCutoff
(
gpu
,
(
float
)
force
.
getCutoffDistance
(),
(
float
)
force
.
getReactionFieldDielectric
());
method
=
CUTOFF
;
}
if
(
force
.
getNonbondedMethod
()
==
NonbondedForce
::
CutoffPeriodic
)
{
...
...
@@ -576,7 +576,7 @@ void CudaCalcGBVIForceKernel::initialize(const System& system, const GBVIForce&
vector
<
float
>
gammas
(
numParticles
);
for
(
int
i
=
0
;
i
<
numParticles
;
i
++
)
{
double
charge
,
particleRadius
,
gamma
,
bornRadiusScaleFactor
;
double
charge
,
particleRadius
,
gamma
;
force
.
getParticleParameters
(
i
,
charge
,
particleRadius
,
gamma
);
particle
[
i
]
=
i
;
radius
[
i
]
=
(
float
)
particleRadius
;
...
...
platforms/cuda/src/kernels/bbsort.cu
View file @
a39126ab
...
...
@@ -189,8 +189,8 @@ void bbSortBody(T* dData,int size,int listOrder/*,float sliceStep,int sliceSize,
evaluateDisorder
(
dData
,
size
,
maxValue
,
minValue
,
listOrder
);
}
float
sliceStep
=
(
50.0
*
((
double
)(
maxValue
-
minValue
)
/
(
double
)
size
));
int
sliceSize
=
(
maxValue
-
minValue
)
/
sliceStep
+
10
;
float
sliceStep
=
(
float
)
(
50.0
*
((
double
)(
maxValue
-
minValue
)
/
(
double
)
size
));
int
sliceSize
=
(
int
)
(
(
maxValue
-
minValue
)
/
sliceStep
+
10
)
;
int
blockCount
;
...
...
platforms/cuda/src/kernels/gpu.cpp
View file @
a39126ab
...
...
@@ -148,14 +148,14 @@ static Expression<SIZE> createExpression(gpuContext gpu, const string& expressio
throw
OpenMMException
(
"Expression contains too many operations: "
+
expression
);
exp
.
length
=
program
.
getNumOperations
();
exp
.
stackSize
=
program
.
getStackSize
();
if
(
exp
.
stackSize
>
maxStackSize
)
if
(
exp
.
stackSize
>
(
int
)
maxStackSize
)
maxStackSize
=
exp
.
stackSize
;
for
(
int
i
=
0
;
i
<
program
.
getNumOperations
();
i
++
)
{
const
Operation
&
op
=
program
.
getOperation
(
i
);
switch
(
op
.
getId
())
{
case
Operation
::
CONSTANT
:
exp
.
op
[
i
]
=
CONSTANT
;
exp
.
arg
[
i
]
=
dynamic_cast
<
const
Operation
::
Constant
*>
(
&
op
)
->
getValue
();
exp
.
arg
[
i
]
=
(
float
)
dynamic_cast
<
const
Operation
::
Constant
*>
(
&
op
)
->
getValue
();
break
;
case
Operation
::
VARIABLE
:
if
(
variables
.
size
()
>
0
&&
op
.
getName
()
==
variables
[
0
])
...
...
@@ -178,18 +178,18 @@ static Expression<SIZE> createExpression(gpuContext gpu, const string& expressio
exp
.
op
[
i
]
=
VARIABLE8
;
else
{
int
j
;
for
(
j
=
0
;
j
<
globalParamNames
.
size
()
&&
op
.
getName
()
!=
globalParamNames
[
j
];
j
++
);
for
(
j
=
0
;
j
<
(
int
)
globalParamNames
.
size
()
&&
op
.
getName
()
!=
globalParamNames
[
j
];
j
++
);
if
(
j
==
globalParamNames
.
size
())
throw
OpenMMException
(
"Unknown variable '"
+
op
.
getName
()
+
"' in expression: "
+
expression
);
exp
.
op
[
i
]
=
GLOBAL
;
exp
.
arg
[
i
]
=
j
;
exp
.
arg
[
i
]
=
(
float
)
j
;
}
break
;
case
Operation
::
CUSTOM
:
exp
.
op
[
i
]
=
dynamic_cast
<
const
Operation
::
Custom
*>
(
&
op
)
->
getDerivOrder
()[
0
]
==
0
?
CUSTOM
:
CUSTOM_DERIV
;
for
(
int
j
=
0
;
j
<
MAX_TABULATED_FUNCTIONS
;
j
++
)
if
(
op
.
getName
()
==
gpu
->
tabulatedFunctions
[
j
].
name
)
{
exp
.
arg
[
i
]
=
j
;
exp
.
arg
[
i
]
=
(
float
)
j
;
break
;
}
break
;
...
...
@@ -276,15 +276,15 @@ static Expression<SIZE> createExpression(gpuContext gpu, const string& expressio
break
;
case
Operation
::
ADD_CONSTANT
:
exp
.
op
[
i
]
=
ADD_CONSTANT
;
exp
.
arg
[
i
]
=
dynamic_cast
<
const
Operation
::
AddConstant
*>
(
&
op
)
->
getValue
();
exp
.
arg
[
i
]
=
(
float
)
dynamic_cast
<
const
Operation
::
AddConstant
*>
(
&
op
)
->
getValue
();
break
;
case
Operation
::
MULTIPLY_CONSTANT
:
exp
.
op
[
i
]
=
MULTIPLY_CONSTANT
;
exp
.
arg
[
i
]
=
dynamic_cast
<
const
Operation
::
MultiplyConstant
*>
(
&
op
)
->
getValue
();
exp
.
arg
[
i
]
=
(
float
)
dynamic_cast
<
const
Operation
::
MultiplyConstant
*>
(
&
op
)
->
getValue
();
break
;
case
Operation
::
POWER_CONSTANT
:
exp
.
op
[
i
]
=
POWER_CONSTANT
;
exp
.
arg
[
i
]
=
dynamic_cast
<
const
Operation
::
PowerConstant
*>
(
&
op
)
->
getValue
();
exp
.
arg
[
i
]
=
(
float
)
dynamic_cast
<
const
Operation
::
PowerConstant
*>
(
&
op
)
->
getValue
();
break
;
}
}
...
...
@@ -540,11 +540,11 @@ void gpuSetLJ14Parameters(gpuContext gpu, float epsfac, float fudge, const vecto
extern
"C"
void
setExclusions
(
gpuContext
gpu
,
const
vector
<
vector
<
int
>
>&
exclusions
)
{
if
(
gpu
->
exclusions
.
size
()
>
0
)
{
bool
ok
=
(
exclusions
.
size
()
==
gpu
->
exclusions
.
size
());
for
(
int
i
=
0
;
i
<
exclusions
.
size
()
&&
ok
;
i
++
)
{
for
(
int
i
=
0
;
i
<
(
int
)
exclusions
.
size
()
&&
ok
;
i
++
)
{
if
(
exclusions
[
i
].
size
()
!=
gpu
->
exclusions
[
i
].
size
())
ok
=
false
;
else
{
for
(
int
j
=
0
;
j
<
exclusions
[
i
].
size
();
j
++
)
for
(
int
j
=
0
;
j
<
(
int
)
exclusions
[
i
].
size
();
j
++
)
if
(
find
(
gpu
->
exclusions
[
i
].
begin
(),
gpu
->
exclusions
[
i
].
end
(),
exclusions
[
i
][
j
])
==
gpu
->
exclusions
[
i
].
end
())
ok
=
false
;
}
...
...
@@ -638,16 +638,16 @@ void gpuSetTabulatedFunction(gpuContext gpu, int index, const string& name, cons
for
(
int
i
=
0
;
i
<
(
int
)
values
.
size
()
-
1
;
i
++
)
{
float4
c
;
if
(
interpolating
)
{
c
.
x
=
padded
[
i
+
1
];
c
.
y
=
0.5
*
(
-
padded
[
i
]
+
padded
[
i
+
2
]);
c
.
z
=
0.5
*
(
2.0
*
padded
[
i
]
-
5.0
*
padded
[
i
+
1
]
+
4.0
*
padded
[
i
+
2
]
-
padded
[
i
+
3
]);
c
.
w
=
0.5
*
(
-
padded
[
i
]
+
3.0
*
padded
[
i
+
1
]
-
3.0
*
padded
[
i
+
2
]
+
padded
[
i
+
3
]);
c
.
x
=
(
float
)
padded
[
i
+
1
];
c
.
y
=
(
float
)
(
0.5
*
(
-
padded
[
i
]
+
padded
[
i
+
2
])
)
;
c
.
z
=
(
float
)
(
0.5
*
(
2.0
*
padded
[
i
]
-
5.0
*
padded
[
i
+
1
]
+
4.0
*
padded
[
i
+
2
]
-
padded
[
i
+
3
])
)
;
c
.
w
=
(
float
)
(
0.5
*
(
-
padded
[
i
]
+
3.0
*
padded
[
i
+
1
]
-
3.0
*
padded
[
i
+
2
]
+
padded
[
i
+
3
])
)
;
}
else
{
c
.
x
=
(
padded
[
i
]
+
4.0
*
padded
[
i
+
1
]
+
padded
[
i
+
2
])
/
6.0
;
c
.
y
=
(
-
3.0
*
padded
[
i
]
+
3.0
*
padded
[
i
+
2
])
/
6.0
;
c
.
z
=
(
3.0
*
padded
[
i
]
-
6.0
*
padded
[
i
+
1
]
+
3.0
*
padded
[
i
+
2
])
/
6.0
;
c
.
w
=
(
-
padded
[
i
]
+
3.0
*
padded
[
i
+
1
]
-
3.0
*
padded
[
i
+
2
]
+
padded
[
i
+
3
])
/
6.0
;
c
.
x
=
(
float
)
(
(
padded
[
i
]
+
4.0
*
padded
[
i
+
1
]
+
padded
[
i
+
2
])
/
6.0
)
;
c
.
y
=
(
float
)
(
(
-
3.0
*
padded
[
i
]
+
3.0
*
padded
[
i
+
2
])
/
6.0
)
;
c
.
z
=
(
float
)
(
(
3.0
*
padded
[
i
]
-
6.0
*
padded
[
i
+
1
]
+
3.0
*
padded
[
i
+
2
])
/
6.0
)
;
c
.
w
=
(
float
)
(
(
-
padded
[
i
]
+
3.0
*
padded
[
i
+
1
]
-
3.0
*
padded
[
i
+
2
]
+
padded
[
i
+
3
])
/
6.0
)
;
}
(
*
coeff
)[
i
]
=
c
;
}
...
...
@@ -677,18 +677,18 @@ void gpuSetCustomBondParameters(gpuContext gpu, const vector<int>& bondAtom1, co
(
*
gpu
->
psCustomBondID
)[
i
].
z
=
forceBufferCounter
[
bondAtom1
[
i
]]
++
;
(
*
gpu
->
psCustomBondID
)[
i
].
w
=
forceBufferCounter
[
bondAtom2
[
i
]]
++
;
if
(
bondParams
[
i
].
size
()
>
0
)
(
*
gpu
->
psCustomBondParams
)[
i
].
x
=
bondParams
[
i
][
0
];
(
*
gpu
->
psCustomBondParams
)[
i
].
x
=
(
float
)
bondParams
[
i
][
0
];
if
(
bondParams
[
i
].
size
()
>
1
)
(
*
gpu
->
psCustomBondParams
)[
i
].
y
=
bondParams
[
i
][
1
];
(
*
gpu
->
psCustomBondParams
)[
i
].
y
=
(
float
)
bondParams
[
i
][
1
];
if
(
bondParams
[
i
].
size
()
>
2
)
(
*
gpu
->
psCustomBondParams
)[
i
].
z
=
bondParams
[
i
][
2
];
(
*
gpu
->
psCustomBondParams
)[
i
].
z
=
(
float
)
bondParams
[
i
][
2
];
if
(
bondParams
[
i
].
size
()
>
3
)
(
*
gpu
->
psCustomBondParams
)[
i
].
w
=
bondParams
[
i
][
3
];
(
*
gpu
->
psCustomBondParams
)[
i
].
w
=
(
float
)
bondParams
[
i
][
3
];
}
gpu
->
psCustomBondID
->
Upload
();
gpu
->
psCustomBondParams
->
Upload
();
for
(
int
i
=
0
;
i
<
(
int
)
forceBufferCounter
.
size
();
i
++
)
if
(
forceBufferCounter
[
i
]
>
gpu
->
pOutputBufferCounter
[
i
])
if
(
forceBufferCounter
[
i
]
>
(
int
)
gpu
->
pOutputBufferCounter
[
i
])
gpu
->
pOutputBufferCounter
[
i
]
=
forceBufferCounter
[
i
];
// Create the Expressions.
...
...
@@ -720,13 +720,13 @@ void gpuSetCustomExternalParameters(gpuContext gpu, const vector<int>& atomIndex
for
(
int
i
=
0
;
i
<
(
int
)
atomIndex
.
size
();
i
++
)
{
(
*
gpu
->
psCustomExternalID
)[
i
]
=
atomIndex
[
i
];
if
(
atomParams
[
i
].
size
()
>
0
)
(
*
gpu
->
psCustomExternalParams
)[
i
].
x
=
atomParams
[
i
][
0
];
(
*
gpu
->
psCustomExternalParams
)[
i
].
x
=
(
float
)
atomParams
[
i
][
0
];
if
(
atomParams
[
i
].
size
()
>
1
)
(
*
gpu
->
psCustomExternalParams
)[
i
].
y
=
atomParams
[
i
][
1
];
(
*
gpu
->
psCustomExternalParams
)[
i
].
y
=
(
float
)
atomParams
[
i
][
1
];
if
(
atomParams
[
i
].
size
()
>
2
)
(
*
gpu
->
psCustomExternalParams
)[
i
].
z
=
atomParams
[
i
][
2
];
(
*
gpu
->
psCustomExternalParams
)[
i
].
z
=
(
float
)
atomParams
[
i
][
2
];
if
(
atomParams
[
i
].
size
()
>
3
)
(
*
gpu
->
psCustomExternalParams
)[
i
].
w
=
atomParams
[
i
][
3
];
(
*
gpu
->
psCustomExternalParams
)[
i
].
w
=
(
float
)
atomParams
[
i
][
3
];
}
gpu
->
psCustomExternalID
->
Upload
();
gpu
->
psCustomExternalParams
->
Upload
();
...
...
@@ -834,7 +834,7 @@ static void tabulateErfc(gpuContext gpu)
gpu
->
psTabulatedErfc
=
new
CUDAStream
<
float
>
(
tableSize
,
1
,
"TabulatedErfc"
);
gpu
->
sim
.
pTabulatedErfc
=
gpu
->
psTabulatedErfc
->
_pDevData
;
for
(
int
i
=
0
;
i
<
tableSize
;
++
i
)
(
*
gpu
->
psTabulatedErfc
)[
i
]
=
erfc
(
i
*
(
gpu
->
sim
.
alphaEwald
*
gpu
->
sim
.
nonbondedCutoff
)
/
tableSize
);
(
*
gpu
->
psTabulatedErfc
)[
i
]
=
(
float
)
erfc
(
i
*
(
gpu
->
sim
.
alphaEwald
*
gpu
->
sim
.
nonbondedCutoff
)
/
tableSize
);
gpu
->
psTabulatedErfc
->
Upload
();
}
...
...
@@ -996,7 +996,7 @@ void gpuSetGBVIParameters(gpuContext gpu, float innerDielectric, float solventDi
{
(
*
gpu
->
psGBVIData
)[
i
].
x
=
radius
[
i
];
(
*
gpu
->
psGBVIData
)[
i
].
y
=
scaledRadii
[
i
];
(
*
gpu
->
psGBVIData
)[
i
].
z
=
tau
*
gamma
[
i
];
(
*
gpu
->
psGBVIData
)[
i
].
z
=
(
float
)
(
tau
*
gamma
[
i
]
)
;
(
*
gpu
->
psGBVIData
)[
i
].
w
=
1.0
f
;
(
*
gpu
->
psObcData
)[
i
].
x
=
radius
[
i
];
...
...
platforms/cuda/src/kernels/kCalculatePME.cu
View file @
a39126ab
...
...
@@ -89,7 +89,7 @@ inline __host__ __device__ float4 operator*(float4 a, float4 b)
}
inline
__host__
__device__
float4
make_float4
(
int3
a
)
{
return
make_float4
(
a
.
x
,
a
.
y
,
a
.
z
,
0
);
return
make_float4
(
(
float
)
a
.
x
,
(
float
)
a
.
y
,
(
float
)
a
.
z
,
0
);
}
__global__
void
kUpdateGridIndexAndFraction_kernel
()
...
...
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