Skip to content
GitLab
Projects
Groups
Snippets
Help
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Sign in
Toggle navigation
V
vadere
Project overview
Project overview
Details
Activity
Releases
Repository
Repository
Files
Commits
Branches
Tags
Contributors
Graph
Compare
Locked Files
Issues
110
Issues
110
List
Boards
Labels
Service Desk
Milestones
Iterations
Merge Requests
3
Merge Requests
3
Requirements
Requirements
List
CI / CD
CI / CD
Pipelines
Jobs
Schedules
Security & Compliance
Security & Compliance
Dependency List
License Compliance
Operations
Operations
Incidents
Environments
Analytics
Analytics
CI / CD
Code Review
Insights
Issue
Repository
Value Stream
Wiki
Wiki
Members
Members
Collapse sidebar
Close sidebar
Activity
Graph
Create a new issue
Jobs
Commits
Issue Boards
Open sidebar
vadere
vadere
Commits
4ddf575d
Commit
4ddf575d
authored
Jul 17, 2018
by
Benedikt Zoennchen
Browse files
Options
Browse Files
Download
Email Patches
Plain Diff
fix an error in CLLinkedCell: wrong memory allocation.
parent
019afe2b
Changes
3
Hide whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
149 additions
and
20 deletions
+149
-20
VadereUtils/resources/Particles.cl
VadereUtils/resources/Particles.cl
+63
-3
VadereUtils/src/org/vadere/util/opencl/CLLinkedCell.java
VadereUtils/src/org/vadere/util/opencl/CLLinkedCell.java
+5
-6
VadereUtils/tests/org/vadere/util/math/TestCLLinkedList.java
VadereUtils/tests/org/vadere/util/math/TestCLLinkedList.java
+81
-11
No files found.
VadereUtils/resources/Particles.cl
View file @
4ddf575d
...
...
@@ -15,6 +15,7 @@
//
Common
definitions
////////////////////////////////////////////////////////////////////////////////
#
define
UMAD
(
a,
b,
c
)
(
(
a
)
*
(
b
)
+
(
c
)
)
//#define
LOCAL_SIZE_LIMIT
512U
typedef
struct{
float
x
;
...
...
@@ -321,7 +322,7 @@ __kernel void collide(
////////////////////////////////////////////////////////////////////////////////
//
Monolithic
bitonic
sort
kernel
for
short
arrays
fitting
into
local
memory
////////////////////////////////////////////////////////////////////////////////
__kernel
void
bitonicSortLocal
(
/*__kernel
void
bitonicSortLocal1
(
__global
uint
*d_DstKey,
__global
uint
*d_DstVal,
__global
uint
*d_SrcKey,
...
...
@@ -332,8 +333,7 @@ __kernel void bitonicSortLocal(
__local
uint
*l_val
)
{
uint
LOCAL_SIZE_LIMIT
=
get_local_size
(
0
)
*
2
;
//Offset
to
the
beginning
of
subbatch
and
load
data
//Offset
to
the
beginning
of
subarray
and
load
data
d_SrcKey
+=
get_group_id
(
0
)
*
LOCAL_SIZE_LIMIT
+
get_local_id
(
0
)
;
d_SrcVal
+=
get_group_id
(
0
)
*
LOCAL_SIZE_LIMIT
+
get_local_id
(
0
)
;
d_DstKey
+=
get_group_id
(
0
)
*
LOCAL_SIZE_LIMIT
+
get_local_id
(
0
)
;
...
...
@@ -343,6 +343,66 @@ __kernel void bitonicSortLocal(
l_key[get_local_id
(
0
)
+
(
LOCAL_SIZE_LIMIT
/
2
)
]
=
d_SrcKey[
(
LOCAL_SIZE_LIMIT
/
2
)
]
;
l_val[get_local_id
(
0
)
+
(
LOCAL_SIZE_LIMIT
/
2
)
]
=
d_SrcVal[
(
LOCAL_SIZE_LIMIT
/
2
)
]
;
uint
comparatorI
=
get_global_id
(
0
)
&
((
LOCAL_SIZE_LIMIT
/
2
)
-
1
)
;
for
(
uint
size
=
2
; size < LOCAL_SIZE_LIMIT; size <<= 1){
//Bitonic
merge
uint
ddd
=
(
comparatorI
&
(
size
/
2
))
!=
0
;
for
(
uint
stride
=
size
/
2
; stride > 0; stride >>= 1){
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
uint
pos
=
2
*
get_local_id
(
0
)
-
(
get_local_id
(
0
)
&
(
stride
-
1
))
;
ComparatorLocal
(
&l_key[pos
+
0],
&l_val[pos
+
0],
&l_key[pos
+
stride],
&l_val[pos
+
stride],
ddd
)
;
}
}
//Odd
/
even
arrays
of
LOCAL_SIZE_LIMIT
elements
//sorted
in
opposite
directions
{
uint
ddd
=
(
get_group_id
(
0
)
&
1
)
;
for
(
uint
stride
=
LOCAL_SIZE_LIMIT
/
2
; stride > 0; stride >>= 1){
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
uint
pos
=
2
*
get_local_id
(
0
)
-
(
get_local_id
(
0
)
&
(
stride
-
1
))
;
ComparatorLocal
(
&l_key[pos
+
0],
&l_val[pos
+
0],
&l_key[pos
+
stride],
&l_val[pos
+
stride],
ddd
)
;
}
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
d_DstKey[
0]
=
l_key[get_local_id
(
0
)
+
0]
;
d_DstVal[
0]
=
l_val[get_local_id
(
0
)
+
0]
;
d_DstKey[
(
LOCAL_SIZE_LIMIT
/
2
)
]
=
l_key[get_local_id
(
0
)
+
(
LOCAL_SIZE_LIMIT
/
2
)
]
;
d_DstVal[
(
LOCAL_SIZE_LIMIT
/
2
)
]
=
l_val[get_local_id
(
0
)
+
(
LOCAL_SIZE_LIMIT
/
2
)
]
;
}*/
__kernel
void
bitonicSortLocal
(
__global
uint
*d_DstKey,
__global
uint
*d_DstVal,
__global
uint
*d_SrcKey,
__global
uint
*d_SrcVal,
uint
arrayLength,
uint
dir,
__local
uint
*l_key,
__local
uint
*l_val
)
{
uint
LOCAL_SIZE_LIMIT
=
get_local_size
(
0
)
*
2
;
//Offset
to
the
beginning
of
subbatch
and
load
data
d_SrcKey
+=
get_local_id
(
0
)
;
d_SrcVal
+=
get_local_id
(
0
)
;
d_DstKey
+=
get_local_id
(
0
)
;
d_DstVal
+=
get_local_id
(
0
)
;
l_key[get_local_id
(
0
)
+
0]
=
d_SrcKey[
0]
;
l_val[get_local_id
(
0
)
+
0]
=
d_SrcVal[
0]
;
l_key[get_local_id
(
0
)
+
(
LOCAL_SIZE_LIMIT
/
2
)
]
=
d_SrcKey[
(
LOCAL_SIZE_LIMIT
/
2
)
]
;
l_val[get_local_id
(
0
)
+
(
LOCAL_SIZE_LIMIT
/
2
)
]
=
d_SrcVal[
(
LOCAL_SIZE_LIMIT
/
2
)
]
;
for
(
uint
size
=
2
; size < arrayLength; size <<= 1){
//Bitonic
merge
uint
ddd
=
dir
^
(
(
get_local_id
(
0
)
&
(
size
/
2
))
!=
0
)
;
...
...
VadereUtils/src/org/vadere/util/opencl/CLLinkedCell.java
View file @
4ddf575d
...
...
@@ -424,8 +424,8 @@ public class CLLinkedCell {
CLInfo
.
checkCLError
(
clSetKernelArg
(
clFindCellBoundsAndReorder
,
6
,
(
max_work_group_size
+
1
)
*
4
));
// local memory
CLInfo
.
checkCLError
(
clSetKernelArg1i
(
clFindCellBoundsAndReorder
,
7
,
numberOfElements
));
clGlobalWorkSize
.
put
(
0
,
numberOfElements
);
clLocalWorkSize
.
put
(
0
,
max_work_group_size
);
clGlobalWorkSize
.
put
(
0
,
Math
.
min
(
max_work_group_size
,
numberOfElements
)
);
clLocalWorkSize
.
put
(
0
,
Math
.
min
(
max_work_group_size
,
numberOfElements
)
);
//TODO: local work size? + check 2^n constrain!
CLInfo
.
checkCLError
(
clEnqueueNDRangeKernel
(
clQueue
,
clFindCellBoundsAndReorder
,
1
,
null
,
clGlobalWorkSize
,
clLocalWorkSize
,
null
,
null
));
}
...
...
@@ -451,10 +451,9 @@ public class CLLinkedCell {
CLInfo
.
checkCLError
(
clSetKernelArg1p
(
clBitonicSortLocal
,
2
,
clKeysIn
));
CLInfo
.
checkCLError
(
clSetKernelArg1p
(
clBitonicSortLocal
,
3
,
clValuesIn
));
CLInfo
.
checkCLError
(
clSetKernelArg1i
(
clBitonicSortLocal
,
4
,
numberOfElements
));
//TODO: check the hard coded 1, and the waiting of the queue
CLInfo
.
checkCLError
(
clSetKernelArg1i
(
clBitonicSortLocal
,
5
,
1
));
CLInfo
.
checkCLError
(
clSetKernelArg
(
clBitonicSortLocal
,
6
,
keys
.
length
*
4
));
// local memory
CLInfo
.
checkCLError
(
clSetKernelArg
(
clBitonicSortLocal
,
7
,
keys
.
length
*
4
));
// local memory
CLInfo
.
checkCLError
(
clSetKernelArg1i
(
clBitonicSortLocal
,
5
,
dir
));
CLInfo
.
checkCLError
(
clSetKernelArg
(
clBitonicSortLocal
,
6
,
numberOfElements
*
4
));
// local memory
CLInfo
.
checkCLError
(
clSetKernelArg
(
clBitonicSortLocal
,
7
,
numberOfElements
*
4
));
// local memory
clGlobalWorkSize
.
put
(
0
,
numberOfElements
/
2
);
clLocalWorkSize
.
put
(
0
,
numberOfElements
/
2
);
...
...
VadereUtils/tests/org/vadere/util/math/TestCLLinkedList.java
View file @
4ddf575d
...
...
@@ -29,29 +29,51 @@ public class TestCLLinkedList {
public
void
setUp
()
throws
Exception
{}
@Test
public
void
testCalcHash
()
throws
IOException
,
OpenCLException
{
CLLinkedCell
clUniformHashedGrid
=
new
CLLinkedCell
(
1024
,
new
VRectangle
(
0
,
0
,
10
,
10
),
1
);
public
void
testCalcHashSmall
()
throws
IOException
,
OpenCLException
{
int
size
=
8
;
CLLinkedCell
clUniformHashedGrid
=
new
CLLinkedCell
(
size
,
new
VRectangle
(
0
,
0
,
10
,
10
),
1
);
ArrayList
<
VPoint
>
positions
=
new
ArrayList
<>();
for
(
int
i
=
0
;
i
<
1024
;
i
++)
{
for
(
int
i
=
0
;
i
<
size
;
i
++)
{
positions
.
add
(
new
VPoint
(
random
.
nextDouble
()
*
10
,
random
.
nextDouble
()
*
10
));
}
int
[]
hasehs
=
clUniformHashedGrid
.
calcHashes
(
positions
);
assertEquals
(
hasehs
.
length
,
positions
.
size
());
logger
.
info
(
"number of cells = "
+
clUniformHashedGrid
.
getGridSize
()[
0
]
*
clUniformHashedGrid
.
getGridSize
()[
1
]);
//
logger.info("number of cells = " + clUniformHashedGrid.getGridSize()[0] * clUniformHashedGrid.getGridSize()[1]);
for
(
int
i
=
0
;
i
<
hasehs
.
length
;
i
++)
{
int
hash
=
getGridHash
(
getGridPosition
(
positions
.
get
(
i
),
clUniformHashedGrid
.
getCellSize
(),
clUniformHashedGrid
.
getWorldOrign
()),
clUniformHashedGrid
.
getGridSize
());
assertEquals
(
hasehs
[
i
],
hash
);
logger
.
info
(
"hash = "
+
hash
);
//
logger.info("hash = " + hash);
}
}
@Test
public
void
testCalcAndSortHash
()
throws
IOException
,
OpenCLException
{
CLLinkedCell
clUniformHashedGrid
=
new
CLLinkedCell
(
1024
,
new
VRectangle
(
0
,
0
,
10
,
10
),
1
);
public
void
testCalcHashLarge
()
throws
IOException
,
OpenCLException
{
int
size
=
32768
;
// 2^15
CLLinkedCell
clUniformHashedGrid
=
new
CLLinkedCell
(
size
,
new
VRectangle
(
0
,
0
,
10
,
10
),
1
);
ArrayList
<
VPoint
>
positions
=
new
ArrayList
<>();
for
(
int
i
=
0
;
i
<
1024
;
i
++)
{
for
(
int
i
=
0
;
i
<
size
;
i
++)
{
positions
.
add
(
new
VPoint
(
random
.
nextDouble
()
*
10
,
random
.
nextDouble
()
*
10
));
}
int
[]
hasehs
=
clUniformHashedGrid
.
calcHashes
(
positions
);
assertEquals
(
hasehs
.
length
,
positions
.
size
());
//logger.info("number of cells = " + clUniformHashedGrid.getGridSize()[0] * clUniformHashedGrid.getGridSize()[1]);
for
(
int
i
=
0
;
i
<
hasehs
.
length
;
i
++)
{
int
hash
=
getGridHash
(
getGridPosition
(
positions
.
get
(
i
),
clUniformHashedGrid
.
getCellSize
(),
clUniformHashedGrid
.
getWorldOrign
()),
clUniformHashedGrid
.
getGridSize
());
assertEquals
(
hasehs
[
i
],
hash
);
//logger.info("hash = " + hash);
}
}
@Test
public
void
testCalcAndSortHashSmall
()
throws
IOException
,
OpenCLException
{
int
size
=
8
;
CLLinkedCell
clUniformHashedGrid
=
new
CLLinkedCell
(
size
,
new
VRectangle
(
0
,
0
,
10
,
10
),
1
);
ArrayList
<
VPoint
>
positions
=
new
ArrayList
<>();
for
(
int
i
=
0
;
i
<
size
;
i
++)
{
positions
.
add
(
new
VPoint
(
random
.
nextDouble
()
*
10
,
random
.
nextDouble
()
*
10
));
}
int
[]
hasehs
=
clUniformHashedGrid
.
calcSortedHashes
(
positions
);
...
...
@@ -71,10 +93,58 @@ public class TestCLLinkedList {
}
@Test
public
void
testGridCell
()
throws
IOException
,
OpenCLException
{
CLLinkedCell
clUniformHashedGrid
=
new
CLLinkedCell
(
1024
,
new
VRectangle
(
0
,
0
,
10
,
10
),
1
);
public
void
testCalcAndSortHashLarge
()
throws
IOException
,
OpenCLException
{
int
size
=
32768
;
// 2^15
CLLinkedCell
clUniformHashedGrid
=
new
CLLinkedCell
(
size
,
new
VRectangle
(
0
,
0
,
10
,
10
),
1
);
ArrayList
<
VPoint
>
positions
=
new
ArrayList
<>();
for
(
int
i
=
0
;
i
<
size
;
i
++)
{
positions
.
add
(
new
VPoint
(
random
.
nextDouble
()
*
10
,
random
.
nextDouble
()
*
10
));
}
int
[]
hasehs
=
clUniformHashedGrid
.
calcSortedHashes
(
positions
);
assertEquals
(
hasehs
.
length
,
positions
.
size
());
int
[]
expectedHashes
=
new
int
[
positions
.
size
()];
for
(
int
i
=
0
;
i
<
hasehs
.
length
;
i
++)
{
int
hash
=
getGridHash
(
getGridPosition
(
positions
.
get
(
i
),
clUniformHashedGrid
.
getCellSize
(),
clUniformHashedGrid
.
getWorldOrign
()),
clUniformHashedGrid
.
getGridSize
());
expectedHashes
[
i
]
=
hash
;
}
Arrays
.
sort
(
expectedHashes
);
for
(
int
i
=
0
;
i
<
hasehs
.
length
;
i
++)
{
assertEquals
(
hasehs
[
i
],
expectedHashes
[
i
]);
}
}
@Test
public
void
testGridCellSmall
()
throws
IOException
,
OpenCLException
{
int
size
=
8
;
CLLinkedCell
clUniformHashedGrid
=
new
CLLinkedCell
(
size
,
new
VRectangle
(
0
,
0
,
10
,
10
),
1
);
ArrayList
<
VPoint
>
positions
=
new
ArrayList
<>();
for
(
int
i
=
0
;
i
<
size
;
i
++)
{
positions
.
add
(
new
VPoint
(
random
.
nextDouble
()
*
10
,
random
.
nextDouble
()
*
10
));
}
CLLinkedCell
.
LinkedCell
gridCells
=
clUniformHashedGrid
.
calcLinkedCell
(
positions
);
int
numberOfCells
=
clUniformHashedGrid
.
getGridSize
()[
0
]
*
clUniformHashedGrid
.
getGridSize
()[
1
];
for
(
int
cell
=
0
;
cell
<
numberOfCells
;
cell
++)
{
int
cellStart
=
gridCells
.
cellStarts
[
cell
];
int
cellEnd
=
gridCells
.
cellEnds
[
cell
];
for
(
int
i
=
cellStart
;
i
<
cellEnd
;
i
++)
{
VPoint
point
=
new
VPoint
(
gridCells
.
reorderedPositions
[
i
*
2
],
gridCells
.
reorderedPositions
[
i
*
2
+
1
]);
int
[]
gridPosition
=
getGridPosition
(
point
,
clUniformHashedGrid
.
getCellSize
(),
clUniformHashedGrid
.
getWorldOrign
());
int
gridHash
=
getGridHash
(
gridPosition
,
clUniformHashedGrid
.
getGridSize
());
assertEquals
(
gridHash
,
cell
);
}
}
}
@Test
public
void
testGridCellLarge
()
throws
IOException
,
OpenCLException
{
int
size
=
32768
;
CLLinkedCell
clUniformHashedGrid
=
new
CLLinkedCell
(
size
,
new
VRectangle
(
0
,
0
,
10
,
10
),
1
);
ArrayList
<
VPoint
>
positions
=
new
ArrayList
<>();
for
(
int
i
=
0
;
i
<
1024
;
i
++)
{
for
(
int
i
=
0
;
i
<
size
;
i
++)
{
positions
.
add
(
new
VPoint
(
random
.
nextDouble
()
*
10
,
random
.
nextDouble
()
*
10
));
}
CLLinkedCell
.
LinkedCell
gridCells
=
clUniformHashedGrid
.
calcLinkedCell
(
positions
);
...
...
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