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
accc4ccd
Commit
accc4ccd
authored
Aug 04, 2010
by
Mark Friedrichs
Browse files
Mods to quiet gcc (Ubuntu 4.3.3-5ubuntu4)
parent
545a83ee
Changes
7
Expand all
Hide whitespace changes
Inline
Side-by-side
Showing
7 changed files
with
89 additions
and
67 deletions
+89
-67
plugins/amoeba/platforms/cuda/src/kernels/AmoebaGpu.cpp
plugins/amoeba/platforms/cuda/src/kernels/AmoebaGpu.cpp
+6
-6
plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaElectrostatic.cu
...rms/cuda/src/kernels/kCalculateAmoebaCudaElectrostatic.cu
+1
-1
plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaKirkwood.cu
...latforms/cuda/src/kernels/kCalculateAmoebaCudaKirkwood.cu
+1
-1
plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaKirkwoodEDiff.cu
...rms/cuda/src/kernels/kCalculateAmoebaCudaKirkwoodEDiff.cu
+4
-2
plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaMutualInducedAndGkFields.cu
...c/kernels/kCalculateAmoebaCudaMutualInducedAndGkFields.cu
+2
-0
plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaWcaDispersion.cu
...rms/cuda/src/kernels/kCalculateAmoebaCudaWcaDispersion.cu
+1
-1
plugins/amoeba/platforms/cuda/tests/AmoebaTinkerParameterFile.cpp
...amoeba/platforms/cuda/tests/AmoebaTinkerParameterFile.cpp
+74
-56
No files found.
plugins/amoeba/platforms/cuda/src/kernels/AmoebaGpu.cpp
View file @
accc4ccd
...
...
@@ -504,7 +504,7 @@ void gpuSetAmoebaAngleParameters(amoebaGpuContext amoebaGpu, const std::vector<i
#define DUMP_PARAMETERS 5
#if (DUMP_PARAMETERS > 0 )
if
(
(
i
<
DUMP_PARAMETERS
||
i
>
bond_angles
-
(
DUMP_PARAMETERS
+
1
))
&&
amoebaGpu
->
log
)
fprintf
(
amoebaGpu
->
log
,
"Angles: %5d [%5d %5d %5d
] [%5d
%5d %5d] A=%15.7e k
[%15.7e
%15.7e
]
[%5d %5d %5d]
\n
"
,
i
,
fprintf
(
amoebaGpu
->
log
,
"Angles: %5d [%5d %5d %5d
%5d] [
%5d %5d] A=%15.7e k
=
%15.7e [%5d %5d %5d]
\n
"
,
i
,
(
*
psAngleID1
)[
i
].
x
,
(
*
psAngleID1
)[
i
].
y
,
(
*
psAngleID1
)[
i
].
z
,
(
*
psAngleID1
)[
i
].
w
,
(
*
psAngleID2
)[
i
].
x
,
(
*
psAngleID2
)[
i
].
y
,
(
*
psAngleParameter
)[
i
].
x
,
(
*
psAngleParameter
)[
i
].
y
,
...
...
@@ -564,7 +564,7 @@ void gpuSetAmoebaInPlaneAngleParameters(amoebaGpuContext amoebaGpu, const std::v
#define DUMP_PARAMETERS 5
#if (DUMP_PARAMETERS > 0 )
if
(
(
i
<
DUMP_PARAMETERS
||
i
>
bond_angles
-
(
DUMP_PARAMETERS
+
1
))
&&
amoebaGpu
->
log
)
fprintf
(
amoebaGpu
->
log
,
"InPlaneAngles: %5d [%5d %5d %5d %5d] [%5d %5d %5d %5d] A=%15.7e k
[%15.7e %15.7e %15.7e %15.7e %15.7e]
[%5d %5d %5d %5d]
\n
"
,
i
,
fprintf
(
amoebaGpu
->
log
,
"InPlaneAngles: %5d [%5d %5d %5d %5d] [%5d %5d %5d %5d] A=%15.7e k
=%15.7e
[%5d %5d %5d %5d]
\n
"
,
i
,
(
*
psAngleID1
)[
i
].
x
,
(
*
psAngleID1
)[
i
].
y
,
(
*
psAngleID1
)[
i
].
z
,
(
*
psAngleID1
)[
i
].
w
,
(
*
psAngleID2
)[
i
].
x
,
(
*
psAngleID2
)[
i
].
y
,
(
*
psAngleID2
)[
i
].
z
,
(
*
psAngleID2
)[
i
].
w
,
(
*
psAngleParameter
)[
i
].
x
,
(
*
psAngleParameter
)[
i
].
y
,
...
...
@@ -2100,7 +2100,7 @@ void gpuSetAmoebaVdwParameters( amoebaGpuContext amoebaGpu,
psVdwReductionID
->
_pSysStream
[
0
][
count
].
w
=
ii
;
}
if
(
ivMapping
[
ii
].
size
()
>
3
){
(
void
)
fprintf
(
stderr
,
"Atom %u has %u reductions -- invalid -- aborting"
,
ii
,
ivMapping
[
ii
].
size
()
);
(
void
)
fprintf
(
stderr
,
"Atom %u has %u reductions -- invalid -- aborting"
,
ii
,
static_cast
<
unsigned
int
>
(
ivMapping
[
ii
].
size
()
)
);
exit
(
1
);
}
count
++
;
...
...
@@ -2335,7 +2335,7 @@ void amoebaGpuBuildVdwExclusionList( amoebaGpuContext amoebaGpu, const std::vec
for
(
unsigned
int
ii
=
0
;
ii
<
actualAtoms
;
ii
++
){
bool
error
=
false
;
if
(
exclusions
[
ii
].
size
()
!=
echoExclusions
[
ii
].
size
()
){
(
void
)
fprintf
(
amoebaGpu
->
log
,
"
\n
Atom %6d sz %6u %6u XX
\n
"
,
ii
,
exclusions
[
ii
].
size
(),
echoExclusions
[
ii
].
size
()
);
(
void
)
fprintf
(
amoebaGpu
->
log
,
"
\n
Atom %6d sz %6u %6u XX
\n
"
,
ii
,
static_cast
<
unsigned
int
>
(
exclusions
[
ii
].
size
()
)
,
static_cast
<
unsigned
int
>
(
echoExclusions
[
ii
].
size
()
)
);
error
=
true
;
totalErrors
++
;
}
...
...
@@ -2736,7 +2736,7 @@ static int matchMaps( std::string idString, MapIntFloat* map1, MapIntFloat* map
int
equalSizes
=
1
;
int
error
=
0
;
if
(
map1
->
size
()
!=
map2
->
size
()
){
(
void
)
fprintf
(
log
,
"%s sizes unequal: %u %u
\n
"
,
idString
.
c_str
(),
map1
->
size
(),
map2
->
size
()
);
(
void
)
fprintf
(
log
,
"%s sizes unequal: %u %u
\n
"
,
idString
.
c_str
(),
static_cast
<
unsigned
int
>
(
map1
->
size
()
)
,
static_cast
<
unsigned
int
>
(
map2
->
size
()
)
);
equalSizes
=
0
;
error
++
;
}
...
...
@@ -3713,7 +3713,7 @@ void cudaWriteVectorOfDoubleVectorsToFile( char* fname, std::vector<int>& fileId
// ---------------------------------------------------------------------------------------
FILE
*
filePtr
=
getWriteToFilePtrV
(
fname
,
fileId
);
(
void
)
fprintf
(
filePtr
,
"%u
\n
"
,
outputVector
.
size
()
);
(
void
)
fprintf
(
filePtr
,
"%u
\n
"
,
static_cast
<
unsigned
int
>
(
outputVector
.
size
()
)
);
float
values
[
50
];
for
(
unsigned
int
ii
=
0
;
ii
<
outputVector
.
size
();
ii
++
){
...
...
plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaElectrostatic.cu
View file @
accc4ccd
...
...
@@ -829,7 +829,7 @@ void cudaComputeAmoebaElectrostatic( amoebaGpuContext amoebaGpu )
if
(
gpu
->
bOutputBufferPerWarp
){
(
void
)
fprintf
(
amoebaGpu
->
log
,
"kCalculateAmoebaCudaElectrostaticN2Forces warp: numBlocks=%u numThreads=%u bufferPerWarp=%u atm=%u shrd=%u Ebuf=%u ixnCt=%u workUnits=%u
\n
"
,
(
void
)
fprintf
(
amoebaGpu
->
log
,
"kCalculateAmoebaCudaElectrostaticN2Forces warp: numBlocks=%u numThreads=%u bufferPerWarp=%u atm=%
l
u shrd=%
l
u Ebuf=%u ixnCt=%
l
u workUnits=%u
\n
"
,
amoebaGpu
->
nonbondBlocks
,
threadsPerBlock
,
amoebaGpu
->
bOutputBufferPerWarp
,
sizeof
(
ElectrostaticParticle
),
sizeof
(
ElectrostaticParticle
)
*
threadsPerBlock
,
amoebaGpu
->
energyOutputBuffers
,
(
*
gpu
->
psInteractionCount
)[
0
],
gpu
->
sim
.
workUnits
);
(
void
)
fflush
(
amoebaGpu
->
log
);
...
...
plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaKirkwood.cu
View file @
accc4ccd
...
...
@@ -1971,7 +1971,7 @@ void kCalculateAmoebaKirkwood( amoebaGpuContext amoebaGpu )
if
(
amoebaGpu
->
log
){
(
void
)
fprintf
(
amoebaGpu
->
log
,
"kCalculateAmoebaCudaKirkwood: blcks=%u tds=%u %u bPrWrp=%u atm=%u shrd=%u Ebuf=%u ixnCt=%u workUnits=%u
\n
"
,
(
void
)
fprintf
(
amoebaGpu
->
log
,
"kCalculateAmoebaCudaKirkwood: blcks=%u tds=%u %u bPrWrp=%u atm=%
l
u shrd=%
l
u Ebuf=%u ixnCt=%
l
u workUnits=%u
\n
"
,
amoebaGpu
->
nonbondBlocks
,
threadsPerBlock
,
maxThreads
,
amoebaGpu
->
bOutputBufferPerWarp
,
sizeof
(
KirkwoodParticle
),
sizeof
(
KirkwoodParticle
)
*
threadsPerBlock
,
amoebaGpu
->
energyOutputBuffers
,
(
*
gpu
->
psInteractionCount
)[
0
],
gpu
->
sim
.
workUnits
);
...
...
plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaKirkwoodEDiff.cu
View file @
accc4ccd
...
...
@@ -1031,6 +1031,7 @@ __device__ void calculateKirkwoodEDiffPairIxn_kernel( float4 atomCoordinatesI,
}
#ifdef AMOEBA_DEBUG
__device__
static
int
debugAccumulate
(
unsigned
int
index
,
float4
*
debugArray
,
float
*
field
,
unsigned
int
addMask
,
float
idLabel
)
{
index
+=
cAmoebaSim
.
paddedNumberOfAtoms
;
...
...
@@ -1041,6 +1042,7 @@ __device__ static int debugAccumulate( unsigned int index, float4* debugArray, f
return
index
;
}
#endif
__device__
void
zeroKirkwoodEDiffParticleSharedField
(
struct
KirkwoodEDiffParticle
*
sA
)
{
...
...
@@ -1154,11 +1156,11 @@ void kCalculateAmoebaKirkwoodEDiff( amoebaGpuContext amoebaGpu )
// ---------------------------------------------------------------------------------------
static
const
char
*
methodName
=
"kCalculateAmoebaKirkwoodEDiff"
;
static
unsigned
int
threadsPerBlock
=
0
;
static
int
timestep
=
0
;
timestep
++
;
#ifdef AMOEBA_DEBUG
static
const
char
*
methodName
=
"kCalculateAmoebaKirkwoodEDiff"
;
std
::
vector
<
int
>
fileId
;
fileId
.
resize
(
2
);
fileId
[
0
]
=
timestep
;
...
...
@@ -1201,7 +1203,7 @@ void kCalculateAmoebaKirkwoodEDiff( amoebaGpuContext amoebaGpu )
}
if
(
amoebaGpu
->
log
&&
timestep
==
1
){
(
void
)
fprintf
(
amoebaGpu
->
log
,
"kCalculateAmoebaCudaKirkwoodEDiffN2Forces: numBlocks=%u numThreads=%u bufferPerWarp=%u atm=%u shrd=%u Ebuf=%u ixnCt=%u workUnits=%u
\n
"
,
(
void
)
fprintf
(
amoebaGpu
->
log
,
"kCalculateAmoebaCudaKirkwoodEDiffN2Forces: numBlocks=%u numThreads=%u bufferPerWarp=%u atm=%
l
u shrd=%
l
u Ebuf=%u ixnCt=%
l
u workUnits=%u
\n
"
,
amoebaGpu
->
nonbondBlocks
,
threadsPerBlock
,
amoebaGpu
->
bOutputBufferPerWarp
,
sizeof
(
KirkwoodEDiffParticle
),
sizeof
(
KirkwoodEDiffParticle
)
*
threadsPerBlock
,
amoebaGpu
->
energyOutputBuffers
,
(
*
gpu
->
psInteractionCount
)[
0
],
gpu
->
sim
.
workUnits
);
...
...
plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaMutualInducedAndGkFields.cu
View file @
accc4ccd
...
...
@@ -267,6 +267,7 @@ __device__ void calculateMutualInducedAndGkFieldsGkPairIxn_kernel( float4 atomCo
}
#ifdef AMOEBA_DEBUG
__device__
static
int
debugAccumulate
(
int
index
,
float4
*
debugArray
,
float
*
field
,
unsigned
int
addMask
,
float
idLabel
)
{
index
+=
cAmoebaSim
.
paddedNumberOfAtoms
;
...
...
@@ -277,6 +278,7 @@ __device__ static int debugAccumulate( int index, float4* debugArray, float* fie
return
index
;
}
#endif
#define GK
#include "kCalculateAmoebaCudaMutualInducedParticle.h"
...
...
plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaWcaDispersion.cu
View file @
accc4ccd
...
...
@@ -31,7 +31,7 @@ void GetCalculateAmoebaCudaWcaDispersionSim(amoebaGpuContext amoebaGpu)
gpuContext
gpu
=
amoebaGpu
->
gpuContext
;
status
=
cudaMemcpyFromSymbol
(
&
gpu
->
sim
,
cSim
,
sizeof
(
cudaGmxSimulation
));
fprintf
(
stderr
,
"In GetCalculateAmoebaCudaWcaDispersionSim: %p %u %u
\n
"
,
fprintf
(
stderr
,
"In GetCalculateAmoebaCudaWcaDispersionSim: %p %
l
u %u
\n
"
,
gpu
->
psInteractionCount
->
_pSysStream
[
0
],
gpu
->
psInteractionCount
->
_pSysStream
[
0
][
0
],
gpu
->
sim
.
workUnits
);
RTERROR
(
status
,
"GetCalculateAmoebaCudaWcaDispersionSim: cudaMemcpyFromSymbol: SetSim copy from cSim failed"
);
...
...
plugins/amoeba/platforms/cuda/tests/AmoebaTinkerParameterFile.cpp
View file @
accc4ccd
This diff is collapsed.
Click to expand it.
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