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
2db43eae
Commit
2db43eae
authored
Apr 05, 2010
by
Mark Friedrichs
Browse files
Modified pForce4a -> pForce4
parent
aeb5dc66
Changes
4
Hide whitespace changes
Inline
Side-by-side
Showing
4 changed files
with
23 additions
and
23 deletions
+23
-23
plugins/freeEnergy/platforms/cuda/src/kernels/kCalculateCDLJObcGbsaSoftcoreForces1.h
...s/cuda/src/kernels/kCalculateCDLJObcGbsaSoftcoreForces1.h
+12
-12
plugins/freeEnergy/platforms/cuda/src/kernels/kCalculateGBVISoftcoreForces2.cu
...atforms/cuda/src/kernels/kCalculateGBVISoftcoreForces2.cu
+1
-1
plugins/freeEnergy/platforms/cuda/src/kernels/kCalculateLocalSoftcoreForces.cu
...atforms/cuda/src/kernels/kCalculateLocalSoftcoreForces.cu
+1
-1
plugins/freeEnergy/platforms/cuda/src/kernels/kCalculateNonbondedSoftcore.h
.../platforms/cuda/src/kernels/kCalculateNonbondedSoftcore.h
+9
-9
No files found.
plugins/freeEnergy/platforms/cuda/src/kernels/kCalculateCDLJObcGbsaSoftcoreForces1.h
View file @
2db43eae
...
@@ -298,13 +298,13 @@ __global__ void METHOD_NAME(kCalculateCDLJObcGbsaSoftcore, Forces1_kernel)(unsig
...
@@ -298,13 +298,13 @@ __global__ void METHOD_NAME(kCalculateCDLJObcGbsaSoftcore, Forces1_kernel)(unsig
unsigned
int
offset
=
x
+
tgx
+
(
x
>>
GRIDBITS
)
*
cSim
.
stride
;
unsigned
int
offset
=
x
+
tgx
+
(
x
>>
GRIDBITS
)
*
cSim
.
stride
;
#endif
#endif
float4
of
=
cSim
.
pForce4
a
[
offset
];
float4
of
=
cSim
.
pForce4
[
offset
];
float
bf
=
cSim
.
pBornForce
[
offset
];
float
bf
=
cSim
.
pBornForce
[
offset
];
of
.
x
+=
af
.
x
;
of
.
x
+=
af
.
x
;
of
.
y
+=
af
.
y
;
of
.
y
+=
af
.
y
;
of
.
z
+=
af
.
z
;
of
.
z
+=
af
.
z
;
bf
+=
af
.
w
;
bf
+=
af
.
w
;
cSim
.
pForce4
a
[
offset
]
=
of
;
cSim
.
pForce4
[
offset
]
=
of
;
cSim
.
pBornForce
[
offset
]
=
bf
;
cSim
.
pBornForce
[
offset
]
=
bf
;
}
}
...
@@ -691,13 +691,13 @@ __global__ void METHOD_NAME(kCalculateCDLJObcGbsaSoftcore, Forces1_kernel)(unsig
...
@@ -691,13 +691,13 @@ __global__ void METHOD_NAME(kCalculateCDLJObcGbsaSoftcore, Forces1_kernel)(unsig
#else
#else
unsigned
int
offset
=
x
+
tgx
+
(
y
>>
GRIDBITS
)
*
cSim
.
stride
;
unsigned
int
offset
=
x
+
tgx
+
(
y
>>
GRIDBITS
)
*
cSim
.
stride
;
#endif
#endif
float4
of
=
cSim
.
pForce4
a
[
offset
];
float4
of
=
cSim
.
pForce4
[
offset
];
float
bf
=
cSim
.
pBornForce
[
offset
];
float
bf
=
cSim
.
pBornForce
[
offset
];
of
.
x
+=
af
.
x
;
of
.
x
+=
af
.
x
;
of
.
y
+=
af
.
y
;
of
.
y
+=
af
.
y
;
of
.
z
+=
af
.
z
;
of
.
z
+=
af
.
z
;
bf
+=
af
.
w
;
bf
+=
af
.
w
;
cSim
.
pForce4
a
[
offset
]
=
of
;
cSim
.
pForce4
[
offset
]
=
of
;
cSim
.
pBornForce
[
offset
]
=
bf
;
cSim
.
pBornForce
[
offset
]
=
bf
;
#ifdef USE_OUTPUT_BUFFER_PER_WARP
#ifdef USE_OUTPUT_BUFFER_PER_WARP
...
@@ -705,44 +705,44 @@ __global__ void METHOD_NAME(kCalculateCDLJObcGbsaSoftcore, Forces1_kernel)(unsig
...
@@ -705,44 +705,44 @@ __global__ void METHOD_NAME(kCalculateCDLJObcGbsaSoftcore, Forces1_kernel)(unsig
#else
#else
offset
=
y
+
tgx
+
(
x
>>
GRIDBITS
)
*
cSim
.
stride
;
offset
=
y
+
tgx
+
(
x
>>
GRIDBITS
)
*
cSim
.
stride
;
#endif
#endif
of
=
cSim
.
pForce4
a
[
offset
];
of
=
cSim
.
pForce4
[
offset
];
bf
=
cSim
.
pBornForce
[
offset
];
bf
=
cSim
.
pBornForce
[
offset
];
of
.
x
+=
sA
[
threadIdx
.
x
].
fx
;
of
.
x
+=
sA
[
threadIdx
.
x
].
fx
;
of
.
y
+=
sA
[
threadIdx
.
x
].
fy
;
of
.
y
+=
sA
[
threadIdx
.
x
].
fy
;
of
.
z
+=
sA
[
threadIdx
.
x
].
fz
;
of
.
z
+=
sA
[
threadIdx
.
x
].
fz
;
bf
+=
sA
[
threadIdx
.
x
].
fb
;
bf
+=
sA
[
threadIdx
.
x
].
fb
;
cSim
.
pForce4
a
[
offset
]
=
of
;
cSim
.
pForce4
[
offset
]
=
of
;
cSim
.
pBornForce
[
offset
]
=
bf
;
cSim
.
pBornForce
[
offset
]
=
bf
;
#if 0
#if 0
#ifdef USE_OUTPUT_BUFFER_PER_WARP
#ifdef USE_OUTPUT_BUFFER_PER_WARP
unsigned int offset = x + tgx + warp*cSim.stride;
unsigned int offset = x + tgx + warp*cSim.stride;
float4 of = cSim.pForce4
a
[offset];
float4 of = cSim.pForce4[offset];
of.x += af.x;
of.x += af.x;
of.y += af.y;
of.y += af.y;
of.z += af.z;
of.z += af.z;
of.w += af.w;
of.w += af.w;
cSim.pForce4
a
[offset] = of;
cSim.pForce4[offset]
= of;
cSim.pBornForce[offset] = af.w;
cSim.pBornForce[offset] = af.w;
offset = y + tgx + warp*cSim.stride;
offset = y + tgx + warp*cSim.stride;
of = cSim.pForce4
a
[offset];
of = cSim.pForce4[offset];
of.x += sA[threadIdx.x].fx;
of.x += sA[threadIdx.x].fx;
of.y += sA[threadIdx.x].fy;
of.y += sA[threadIdx.x].fy;
of.z += sA[threadIdx.x].fz;
of.z += sA[threadIdx.x].fz;
of.w += sA[threadIdx.x].fb;
of.w += sA[threadIdx.x].fb;
cSim.pForce4
a
[offset] = of;
cSim.pForce4[offset]
= of;
cSim.pBornForce[offset] = af.w;
cSim.pBornForce[offset] = af.w;
#else
#else
unsigned int offset = x + tgx + (y >> GRIDBITS) * cSim.stride;
unsigned int offset = x + tgx + (y >> GRIDBITS) * cSim.stride;
cSim.pForce4
a
[offset] = af;
cSim.pForce4[offset]
= af;
cSim.pBornForce[offset] = af.w;
cSim.pBornForce[offset] = af.w;
af.x = sA[threadIdx.x].fx;
af.x = sA[threadIdx.x].fx;
af.y = sA[threadIdx.x].fy;
af.y = sA[threadIdx.x].fy;
af.z = sA[threadIdx.x].fz;
af.z = sA[threadIdx.x].fz;
af.w = sA[threadIdx.x].fb;
af.w = sA[threadIdx.x].fb;
offset = y + tgx + (x >> GRIDBITS) * cSim.stride;
offset = y + tgx + (x >> GRIDBITS) * cSim.stride;
cSim.pForce4
a
[offset] = af;
cSim.pForce4[offset]
= af;
cSim.pBornForce[offset] = af.w;
cSim.pBornForce[offset] = af.w;
#endif
#endif
...
...
plugins/freeEnergy/platforms/cuda/src/kernels/kCalculateGBVISoftcoreForces2.cu
View file @
2db43eae
...
@@ -130,7 +130,7 @@ __global__ void kCalculateGBVISoftcoreForces2a_kernel()
...
@@ -130,7 +130,7 @@ __global__ void kCalculateGBVISoftcoreForces2a_kernel()
}
}
// Write results
// Write results
cSim
.
pForce4
a
[
pos
]
=
force
;
cSim
.
pForce4
[
pos
]
=
force
;
}
}
...
...
plugins/freeEnergy/platforms/cuda/src/kernels/kCalculateLocalSoftcoreForces.cu
View file @
2db43eae
...
@@ -26,7 +26,7 @@
...
@@ -26,7 +26,7 @@
#include "GpuLJ14Softcore.h"
#include "GpuLJ14Softcore.h"
#include "GpuFreeEnergyCudaKernels.h"
#include "GpuFreeEnergyCudaKernels.h"
#include <cuda.h>
//
#include <cuda.h>
static
__constant__
cudaGmxSimulation
cSim
;
static
__constant__
cudaGmxSimulation
cSim
;
static
__constant__
cudaFreeEnergySimulationNonbonded14
feSim
;
static
__constant__
cudaFreeEnergySimulationNonbonded14
feSim
;
...
...
plugins/freeEnergy/platforms/cuda/src/kernels/kCalculateNonbondedSoftcore.h
View file @
2db43eae
...
@@ -250,18 +250,18 @@ __global__ void METHOD_NAME(kCalculateCDLJSoftcore, Forces_kernel)(unsigned int*
...
@@ -250,18 +250,18 @@ __global__ void METHOD_NAME(kCalculateCDLJSoftcore, Forces_kernel)(unsigned int*
float4
of
;
float4
of
;
#ifdef USE_OUTPUT_BUFFER_PER_WARP
#ifdef USE_OUTPUT_BUFFER_PER_WARP
unsigned
int
offset
=
x
+
tgx
+
warp
*
cSim
.
stride
;
unsigned
int
offset
=
x
+
tgx
+
warp
*
cSim
.
stride
;
of
=
cSim
.
pForce4
a
[
offset
];
of
=
cSim
.
pForce4
[
offset
];
of
.
x
+=
af
.
x
;
of
.
x
+=
af
.
x
;
of
.
y
+=
af
.
y
;
of
.
y
+=
af
.
y
;
of
.
z
+=
af
.
z
;
of
.
z
+=
af
.
z
;
cSim
.
pForce4
a
[
offset
]
=
of
;
cSim
.
pForce4
[
offset
]
=
of
;
#else
#else
of
.
x
=
af
.
x
;
of
.
x
=
af
.
x
;
of
.
y
=
af
.
y
;
of
.
y
=
af
.
y
;
of
.
z
=
af
.
z
;
of
.
z
=
af
.
z
;
of
.
w
=
0
.
0
f
;
of
.
w
=
0
.
0
f
;
unsigned
int
offset
=
x
+
tgx
+
(
x
>>
GRIDBITS
)
*
cSim
.
stride
;
unsigned
int
offset
=
x
+
tgx
+
(
x
>>
GRIDBITS
)
*
cSim
.
stride
;
cSim
.
pForce4
a
[
offset
]
=
of
;
cSim
.
pForce4
[
offset
]
=
of
;
#endif
#endif
}
}
else
// 100% utilization
else
// 100% utilization
...
@@ -564,29 +564,29 @@ __global__ void METHOD_NAME(kCalculateCDLJSoftcore, Forces_kernel)(unsigned int*
...
@@ -564,29 +564,29 @@ __global__ void METHOD_NAME(kCalculateCDLJSoftcore, Forces_kernel)(unsigned int*
float4
of
;
float4
of
;
#ifdef USE_OUTPUT_BUFFER_PER_WARP
#ifdef USE_OUTPUT_BUFFER_PER_WARP
unsigned
int
offset
=
x
+
tgx
+
warp
*
cSim
.
stride
;
unsigned
int
offset
=
x
+
tgx
+
warp
*
cSim
.
stride
;
of
=
cSim
.
pForce4
a
[
offset
];
of
=
cSim
.
pForce4
[
offset
];
of
.
x
+=
af
.
x
;
of
.
x
+=
af
.
x
;
of
.
y
+=
af
.
y
;
of
.
y
+=
af
.
y
;
of
.
z
+=
af
.
z
;
of
.
z
+=
af
.
z
;
cSim
.
pForce4
a
[
offset
]
=
of
;
cSim
.
pForce4
[
offset
]
=
of
;
offset
=
y
+
tgx
+
warp
*
cSim
.
stride
;
offset
=
y
+
tgx
+
warp
*
cSim
.
stride
;
of
=
cSim
.
pForce4
a
[
offset
];
of
=
cSim
.
pForce4
[
offset
];
of
.
x
+=
sA
[
threadIdx
.
x
].
fx
;
of
.
x
+=
sA
[
threadIdx
.
x
].
fx
;
of
.
y
+=
sA
[
threadIdx
.
x
].
fy
;
of
.
y
+=
sA
[
threadIdx
.
x
].
fy
;
of
.
z
+=
sA
[
threadIdx
.
x
].
fz
;
of
.
z
+=
sA
[
threadIdx
.
x
].
fz
;
cSim
.
pForce4
a
[
offset
]
=
of
;
cSim
.
pForce4
[
offset
]
=
of
;
#else
#else
of
.
x
=
af
.
x
;
of
.
x
=
af
.
x
;
of
.
y
=
af
.
y
;
of
.
y
=
af
.
y
;
of
.
z
=
af
.
z
;
of
.
z
=
af
.
z
;
of
.
w
=
0
.
0
f
;
of
.
w
=
0
.
0
f
;
unsigned
int
offset
=
x
+
tgx
+
(
y
>>
GRIDBITS
)
*
cSim
.
stride
;
unsigned
int
offset
=
x
+
tgx
+
(
y
>>
GRIDBITS
)
*
cSim
.
stride
;
cSim
.
pForce4
a
[
offset
]
=
of
;
cSim
.
pForce4
[
offset
]
=
of
;
of
.
x
=
sA
[
threadIdx
.
x
].
fx
;
of
.
x
=
sA
[
threadIdx
.
x
].
fx
;
of
.
y
=
sA
[
threadIdx
.
x
].
fy
;
of
.
y
=
sA
[
threadIdx
.
x
].
fy
;
of
.
z
=
sA
[
threadIdx
.
x
].
fz
;
of
.
z
=
sA
[
threadIdx
.
x
].
fz
;
offset
=
y
+
tgx
+
(
x
>>
GRIDBITS
)
*
cSim
.
stride
;
offset
=
y
+
tgx
+
(
x
>>
GRIDBITS
)
*
cSim
.
stride
;
cSim
.
pForce4
a
[
offset
]
=
of
;
cSim
.
pForce4
[
offset
]
=
of
;
#endif
#endif
lasty
=
y
;
lasty
=
y
;
}
}
...
...
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
