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
0d34441d
Commit
0d34441d
authored
Oct 25, 2015
by
Robert McGibbon
Browse files
Fix thread synchronization in customCentroidBond.cl
parent
8eaf3c9c
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
32 additions
and
28 deletions
+32
-28
platforms/opencl/src/kernels/customCentroidBond.cl
platforms/opencl/src/kernels/customCentroidBond.cl
+32
-28
No files found.
platforms/opencl/src/kernels/customCentroidBond.cl
View file @
0d34441d
...
@@ -27,37 +27,41 @@ __kernel void computeGroupCenters(__global const real4* restrict posq, __global
...
@@ -27,37 +27,41 @@ __kernel void computeGroupCenters(__global const real4* restrict posq, __global
temp[thread].x
=
center.x
;
temp[thread].x
=
center.x
;
temp[thread].y
=
center.y
;
temp[thread].y
=
center.y
;
temp[thread].z
=
center.z
;
temp[thread].z
=
center.z
;
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
thread
<
32
)
{
if
(
thread
<
32
)
{
temp[thread].x
+=
temp[thread+32].x
;
temp[thread].x
+=
temp[thread+32].x
;
temp[thread].y
+=
temp[thread+32].y
;
temp[thread].y
+=
temp[thread+32].y
;
temp[thread].z
+=
temp[thread+32].z
;
temp[thread].z
+=
temp[thread+32].z
;
}
SYNC_WARPS
;
SYNC_WARPS
;
if
(
thread
<
16
)
{
if
(
thread
<
16
)
{
temp[thread].x
+=
temp[thread+16].x
;
temp[thread].x
+=
temp[thread+16].x
;
temp[thread].y
+=
temp[thread+16].y
;
temp[thread].y
+=
temp[thread+16].y
;
temp[thread].z
+=
temp[thread+16].z
;
temp[thread].z
+=
temp[thread+16].z
;
SYNC_WARPS
;
}
}
SYNC_WARPS
;
if
(
thread
<
8
)
{
if
(
thread
<
8
)
{
temp[thread].x
+=
temp[thread+8].x
;
temp[thread].x
+=
temp[thread+8].x
;
temp[thread].y
+=
temp[thread+8].y
;
temp[thread].y
+=
temp[thread+8].y
;
temp[thread].z
+=
temp[thread+8].z
;
temp[thread].z
+=
temp[thread+8].z
;
SYNC_WARPS
;
}
}
SYNC_WARPS
;
if
(
thread
<
4
)
{
if
(
thread
<
4
)
{
temp[thread].x
+=
temp[thread+4].x
;
temp[thread].x
+=
temp[thread+4].x
;
temp[thread].y
+=
temp[thread+4].y
;
temp[thread].y
+=
temp[thread+4].y
;
temp[thread].z
+=
temp[thread+4].z
;
temp[thread].z
+=
temp[thread+4].z
;
SYNC_WARPS
;
}
}
SYNC_WARPS
;
if
(
thread
<
2
)
{
if
(
thread
<
2
)
{
temp[thread].x
+=
temp[thread+2].x
;
temp[thread].x
+=
temp[thread+2].x
;
temp[thread].y
+=
temp[thread+2].y
;
temp[thread].y
+=
temp[thread+2].y
;
temp[thread].z
+=
temp[thread+2].z
;
temp[thread].z
+=
temp[thread+2].z
;
SYNC_WARPS
;
}
}
}
SYNC_WARPS
;
if
(
thread
==
0
)
if
(
thread
==
0
)
centerPositions[group]
=
(
real4
)
(
temp[0].x+temp[1].x,
temp[0].y+temp[1].y,
temp[0].z+temp[1].z,
0
)
;
centerPositions[group]
=
(
real4
)
(
temp[0].x+temp[1].x,
temp[0].y+temp[1].y,
temp[0].z+temp[1].z,
0
)
;
}
}
...
...
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