Skip to content
GitLab
Projects
Groups
Snippets
/
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Sign in
Toggle navigation
Menu
9.2.2023: Due to updates GitLab will be unavailable for some minutes between 9:00 and 11:00.
Open sidebar
vadere
vadere
Commits
49b9dc91
Commit
49b9dc91
authored
Jun 01, 2020
by
Benedikt Zoennchen
Browse files
extend CLDistMesh from CLOperation.
parent
27da0fa3
Pipeline
#264552
failed with stages
in 1 minute and 50 seconds
Changes
5
Pipelines
1
Hide whitespace changes
Inline
Side-by-side
VadereMeshing/src/org/vadere/meshing/mesh/triangulation/improver/eikmesh/opencl/CLEikMesh.java
View file @
49b9dc91
...
...
@@ -196,7 +196,12 @@ public class CLEikMesh implements IMeshImprover<AVertex, AHalfEdge, AFace>, ITri
}
numberOfIterations
++;
return
clDistMesh
.
step
(
flipAll
);
try
{
return
clDistMesh
.
step
(
flipAll
);
}
catch
(
OpenCLException
e
)
{
e
.
printStackTrace
();
return
false
;
}
/*log.info("#illegalMovementTests: " + numberOfIllegalMovementTests);
log.info("#retriangulations: " + numberOfRetriangulations);
log.info("#steps: " + numberOfIterations);
...
...
@@ -204,7 +209,11 @@ public class CLEikMesh implements IMeshImprover<AVertex, AHalfEdge, AFace>, ITri
}
public
void
finish
()
{
clDistMesh
.
finish
();
try
{
clDistMesh
.
finish
();
}
catch
(
OpenCLException
e
)
{
e
.
printStackTrace
();
}
}
public
void
refresh
()
{
...
...
VadereMeshing/src/org/vadere/meshing/mesh/triangulation/plots/VisualTestGPUEdgeBased.java
View file @
49b9dc91
...
...
@@ -56,21 +56,23 @@ public class VisualTestGPUEdgeBased {
overAllTime
.
resume
();
meshGenerator
.
improve
();
overAllTime
.
suspend
();
meshGenerator
.
refresh
();
synchronized
(
meshGenerator
.
getMesh
())
{
meshGenerator
.
refresh
();
}
try
{
Thread
.
sleep
(
10
0
);
Thread
.
sleep
(
10
);
}
catch
(
InterruptedException
e
)
{
e
.
printStackTrace
();
}
distmeshPanel
.
repaint
();
}
overAllTime
.
stop
();
log
.
info
(
"#vertices: "
+
meshGenerator
.
getMesh
().
getVertices
().
size
());
log
.
info
(
"#edges: "
+
meshGenerator
.
getMesh
().
getEdges
().
size
());
log
.
info
(
"#faces: "
+
meshGenerator
.
getMesh
().
getFaces
().
size
());
log
.
info
(
"quality: "
+
meshGenerator
.
getQuality
());
log
.
info
(
"overall time: "
+
overAllTime
.
getTime
()
+
"[ms]"
);
meshGenerator
.
finish
();
System
.
out
.
println
(
"#vertices: "
+
meshGenerator
.
getMesh
().
getVertices
().
size
());
System
.
out
.
println
(
"#edges: "
+
meshGenerator
.
getMesh
().
getEdges
().
size
());
System
.
out
.
println
(
"#faces: "
+
meshGenerator
.
getMesh
().
getFaces
().
size
());
System
.
out
.
println
(
"quality: "
+
meshGenerator
.
getQuality
());
System
.
out
.
println
(
"overall time: "
+
overAllTime
.
getTime
()
+
"[ms]"
);
}
...
...
VadereMeshing/src/org/vadere/meshing/mesh/triangulation/plots/VisualTestGPUVertexBased.java
View file @
49b9dc91
...
...
@@ -66,11 +66,11 @@ public class VisualTestGPUVertexBased {
overAllTime
.
stop
();
meshGenerator
.
finish
();
log
.
info
(
"#vertices: "
+
meshGenerator
.
getMesh
().
getVertices
().
size
());
log
.
info
(
"#edges: "
+
meshGenerator
.
getMesh
().
getEdges
().
size
());
log
.
info
(
"#faces: "
+
meshGenerator
.
getMesh
().
getFaces
().
size
());
log
.
info
(
"quality: "
+
meshGenerator
.
getQuality
());
log
.
info
(
"overall time: "
+
overAllTime
.
getTime
()
+
"[ms]"
);
System
.
out
.
println
(
"#vertices: "
+
meshGenerator
.
getMesh
().
getVertices
().
size
());
System
.
out
.
println
(
"#edges: "
+
meshGenerator
.
getMesh
().
getEdges
().
size
());
System
.
out
.
println
(
"#faces: "
+
meshGenerator
.
getMesh
().
getFaces
().
size
());
System
.
out
.
println
(
"quality: "
+
meshGenerator
.
getQuality
());
System
.
out
.
println
(
"overall time: "
+
overAllTime
.
getTime
()
+
"[ms]"
);
}
public
static
void
main
(
String
[]
args
)
throws
OpenCLException
{
...
...
VadereMeshing/src/org/vadere/meshing/opencl/CLDistMesh.java
View file @
49b9dc91
...
...
@@ -2,6 +2,7 @@ package org.vadere.meshing.opencl;
import
org.apache.commons.lang3.time.StopWatch
;
import
org.lwjgl.system.Configuration
;
import
org.vadere.util.logging.Logger
;
import
org.jetbrains.annotations.NotNull
;
import
org.lwjgl.PointerBuffer
;
...
...
@@ -16,6 +17,7 @@ import org.vadere.meshing.mesh.gen.AVertex;
import
org.vadere.meshing.mesh.gen.CLGatherer
;
import
org.vadere.util.geometry.shapes.IPoint
;
import
org.vadere.util.opencl.CLInfo
;
import
org.vadere.util.opencl.CLOperation
;
import
org.vadere.util.opencl.CLUtils
;
import
org.vadere.util.opencl.OpenCLException
;
...
...
@@ -28,28 +30,19 @@ import java.util.*;
import
java.util.stream.Collectors
;
import
static
org
.
lwjgl
.
opencl
.
CL10
.*;
import
static
org
.
lwjgl
.
opencl
.
CL10
.
clEnqueueNDRangeKernel
;
import
static
org
.
lwjgl
.
opencl
.
CL11
.
CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE
;
import
static
org
.
lwjgl
.
system
.
MemoryStack
.
stackPush
;
import
static
org
.
lwjgl
.
system
.
MemoryUtil
.
NULL
;
import
static
org
.
lwjgl
.
system
.
MemoryUtil
.
memUTF8
;
/**
* @author Benedikt Zoennchen
*
* DistMesh GPU implementation.
*/
public
class
CLDistMesh
{
public
class
CLDistMesh
extends
CLOperation
{
private
static
Logger
log
=
Logger
.
getLogger
(
CLDistMesh
.
class
);
// CL ids
private
long
clPlatform
;
private
long
clDevice
;
private
long
clContext
;
private
long
clQueue
;
private
long
clProgram
;
// CL kernel ids
private
long
clKernelForces
;
private
long
clKernelMove
;
...
...
@@ -70,10 +63,6 @@ public class CLDistMesh {
//private long clKernelRemoveTriangles;
//private long clKernelCheckTriangles;
// CL callbacks
private
CLContextCallback
contextCB
;
private
CLProgramCallback
programCB
;
// data on the host
private
DoubleBuffer
vD
;
...
...
@@ -139,19 +128,20 @@ public class CLDistMesh {
private
AMesh
mesh
;
private
boolean
doublePrecision
=
true
;
private
boolean
profiling
=
true
;
private
boolean
doublePrecision
=
false
;
private
List
<
IPoint
>
result
;
private
boolean
hasToRead
=
false
;
public
CLDistMesh
(
@NotNull
AMesh
mesh
)
{
/*if(profiling) {
public
CLDistMesh
(
@NotNull
final
AMesh
mesh
)
{
super
(
CL_DEVICE_TYPE_CPU
);
profiling
=
true
;
if
(
profiling
)
{
Configuration
.
DEBUG
.
set
(
true
);
Configuration
.
DEBUG_MEMORY_ALLOCATOR
.
set
(
true
);
Configuration
.
DEBUG_STACK
.
set
(
true
);
Configuration
.
DEBUG_STACK
.
set
(
true
);
}
*/
}
this
.
mesh
=
mesh
;
this
.
mesh
.
garbageCollection
();
...
...
@@ -188,126 +178,21 @@ public class CLDistMesh {
this
.
result
=
mesh
.
streamPoints
().
collect
(
Collectors
.
toList
());
}
private
void
initCallbacks
()
{
contextCB
=
CLContextCallback
.
create
((
errinfo
,
private_info
,
cb
,
user_data
)
->
{
log
.
warn
(
"[LWJGL] cl_context_callback"
);
log
.
warn
(
"\tInfo: "
+
memUTF8
(
errinfo
));
});
programCB
=
CLProgramCallback
.
create
((
program
,
user_data
)
->
{
try
{
log
.
info
(
"The cl_program [0x"
+
program
+
"] was built "
+
(
CLInfo
.
getProgramBuildInfoInt
(
program
,
clDevice
,
CL_PROGRAM_BUILD_STATUS
)
==
CL_SUCCESS
?
"successfully"
:
"unsuccessfully"
));
String
message
=
CLInfo
.
getProgramBuildInfoStringASCII
(
program
,
clDevice
,
CL_PROGRAM_BUILD_LOG
);
if
(!
message
.
isEmpty
())
{
log
.
info
(
"BUILD LOG:\n----\n"
+
message
+
"\n-----"
);
}
}
catch
(
OpenCLException
e
)
{
log
.
error
(
e
.
getMessage
());
}
});
}
private
void
initCL
()
throws
OpenCLException
{
try
(
MemoryStack
stack
=
stackPush
())
{
// helper for the memory allocation in java
//stack = MemoryStack.stackPush();
IntBuffer
errcode_ret
=
stack
.
callocInt
(
1
);
IntBuffer
numberOfPlatforms
=
stack
.
mallocInt
(
1
);
clGetPlatformIDs
(
null
,
numberOfPlatforms
);
PointerBuffer
platformIDs
=
stack
.
mallocPointer
(
numberOfPlatforms
.
get
(
0
));
clGetPlatformIDs
(
platformIDs
,
numberOfPlatforms
);
clPlatform
=
platformIDs
.
get
(
0
);
IntBuffer
numberOfDevices
=
stack
.
mallocInt
(
1
);
clGetDeviceIDs
(
clPlatform
,
CL_DEVICE_TYPE_GPU
,
null
,
numberOfDevices
);
PointerBuffer
deviceIDs
=
stack
.
mallocPointer
(
numberOfDevices
.
get
(
0
));
clGetDeviceIDs
(
clPlatform
,
CL_DEVICE_TYPE_GPU
,
deviceIDs
,
numberOfDevices
);
clDevice
=
deviceIDs
.
get
(
0
);
printDeviceInfo
(
clDevice
,
"CL_DEVICE_NAME"
,
CL_DEVICE_NAME
);
PointerBuffer
ctxProps
=
stack
.
mallocPointer
(
3
);
ctxProps
.
put
(
CL_CONTEXT_PLATFORM
)
.
put
(
clPlatform
)
.
put
(
NULL
)
.
flip
();
clContext
=
clCreateContext
(
ctxProps
,
clDevice
,
contextCB
,
NULL
,
errcode_ret
);
CLInfo
.
checkCLError
(
errcode_ret
);
if
(
profiling
)
{
clQueue
=
clCreateCommandQueue
(
clContext
,
clDevice
,
CL_QUEUE_PROFILING_ENABLE
,
errcode_ret
);
}
else
{
clQueue
=
clCreateCommandQueue
(
clContext
,
clDevice
,
0
,
errcode_ret
);
}
CLInfo
.
checkCLError
(
errcode_ret
);
PointerBuffer
pp
=
stack
.
mallocPointer
(
1
);
clGetDeviceInfo
(
clDevice
,
CL_DEVICE_MAX_WORK_GROUP_SIZE
,
pp
,
null
);
maxGroupSize
=
pp
.
get
(
0
);
clGetDeviceInfo
(
clDevice
,
CL_DEVICE_MAX_COMPUTE_UNITS
,
pp
,
null
);
maxComputeUnits
=
pp
.
get
(
0
);
log
.
info
(
"MAX_GRP_SIZE = "
+
maxGroupSize
);
log
.
info
(
"MAX_COMPUTE_UNITS = "
+
maxComputeUnits
);
}
}
private
void
buildProgram
()
throws
OpenCLException
{
try
(
MemoryStack
stack
=
stackPush
())
{
// helper for the memory allocation in java
//stack = MemoryStack.stackPush();
IntBuffer
errcode_ret
=
stack
.
mallocInt
(
1
);
IntBuffer
numberOfPlatforms
=
stack
.
mallocInt
(
1
);
clGetPlatformIDs
(
null
,
numberOfPlatforms
);
PointerBuffer
platformIDs
=
stack
.
mallocPointer
(
numberOfPlatforms
.
get
(
0
));
clGetPlatformIDs
(
platformIDs
,
numberOfPlatforms
);
clPlatform
=
platformIDs
.
get
(
0
);
IntBuffer
numberOfDevices
=
stack
.
mallocInt
(
1
);
clGetDeviceIDs
(
clPlatform
,
CL_DEVICE_TYPE_GPU
,
null
,
numberOfDevices
);
PointerBuffer
deviceIDs
=
stack
.
mallocPointer
(
numberOfDevices
.
get
(
0
));
clGetDeviceIDs
(
clPlatform
,
CL_DEVICE_TYPE_GPU
,
deviceIDs
,
numberOfDevices
);
clDevice
=
deviceIDs
.
get
(
0
);
printDeviceInfo
(
clDevice
,
"CL_DEVICE_NAME"
,
CL_DEVICE_NAME
);
PointerBuffer
ctxProps
=
stack
.
mallocPointer
(
3
);
ctxProps
.
put
(
CL_CONTEXT_PLATFORM
)
.
put
(
clPlatform
)
.
put
(
NULL
)
.
flip
();
clContext
=
clCreateContext
(
ctxProps
,
clDevice
,
contextCB
,
NULL
,
errcode_ret
);
CLInfo
.
checkCLError
(
errcode_ret
);
if
(
profiling
)
{
clQueue
=
clCreateCommandQueue
(
clContext
,
clDevice
,
CL_QUEUE_PROFILING_ENABLE
,
errcode_ret
);
}
else
{
clQueue
=
clCreateCommandQueue
(
clContext
,
clDevice
,
0
,
errcode_ret
);
}
CLInfo
.
checkCLError
(
errcode_ret
);
PointerBuffer
pp
=
stack
.
mallocPointer
(
1
);
clGetDeviceInfo
(
clDevice
,
CL_DEVICE_MAX_WORK_GROUP_SIZE
,
pp
,
null
);
maxGroupSize
=
pp
.
get
(
0
);
clGetDeviceInfo
(
clDevice
,
CL_DEVICE_MAX_COMPUTE_UNITS
,
pp
,
null
);
maxComputeUnits
=
pp
.
get
(
0
);
log
.
info
(
"MAX_GRP_SIZE = "
+
maxGroupSize
);
log
.
info
(
"MAX_COMPUTE_UNITS = "
+
maxComputeUnits
);
PointerBuffer
clProgramStrings
=
stack
.
mallocPointer
(
1
);
PointerBuffer
clProgramLengths
=
stack
.
mallocPointer
(
1
);
PointerBuffer
pp
=
stack
.
mallocPointer
(
1
);
clGetDeviceInfo
(
clDevice
,
CL_DEVICE_MAX_WORK_GROUP_SIZE
,
pp
,
null
);
maxGroupSize
=
pp
.
get
(
0
);
clGetDeviceInfo
(
clDevice
,
CL_DEVICE_MAX_COMPUTE_UNITS
,
pp
,
null
);
maxComputeUnits
=
pp
.
get
(
0
);
log
.
info
(
"MAX_GRP_SIZE = "
+
maxGroupSize
);
log
.
info
(
"MAX_COMPUTE_UNITS = "
+
maxComputeUnits
);
PointerBuffer
clProgramStrings
=
stack
.
mallocPointer
(
1
);
PointerBuffer
clProgramLengths
=
stack
.
mallocPointer
(
1
);
try
{
if
(
doublePrecision
)
{
...
...
@@ -319,7 +204,6 @@ public class CLDistMesh {
throw
new
RuntimeException
(
e
);
}
clProgramStrings
.
put
(
0
,
source
);
clProgramLengths
.
put
(
0
,
source
.
remaining
());
...
...
@@ -554,12 +438,12 @@ public class CLDistMesh {
}
}*/
public
boolean
step
()
{
public
boolean
step
()
throws
OpenCLException
{
return
step
(
true
);
}
// TODO: think about the use of only 1 work-group!!! It might be bad! solution: use global barrier? force computation? flip hangs after some time?
public
boolean
step
(
final
boolean
flipAll
)
{
public
boolean
step
(
final
boolean
flipAll
)
throws
OpenCLException
{
try
(
MemoryStack
stack
=
stackPush
())
{
/*
* DistMesh-Loop
...
...
@@ -646,25 +530,6 @@ public class CLDistMesh {
}
}
private
long
enqueueNDRangeKernel
(
final
String
name
,
long
command_queue
,
long
kernel
,
int
work_dim
,
PointerBuffer
global_work_offset
,
PointerBuffer
global_work_size
,
PointerBuffer
local_work_size
,
PointerBuffer
event_wait_list
,
PointerBuffer
event
)
{
if
(
profiling
)
{
long
result
=
clEnqueueNDRangeKernel
(
command_queue
,
kernel
,
work_dim
,
global_work_offset
,
global_work_size
,
local_work_size
,
event_wait_list
,
clEvent
);
clWaitForEvents
(
clEvent
);
long
eventAddr
=
clEvent
.
get
();
clGetEventProfilingInfo
(
eventAddr
,
CL_PROFILING_COMMAND_START
,
startTime
,
retSize
);
clGetEventProfilingInfo
(
eventAddr
,
CL_PROFILING_COMMAND_END
,
endTime
,
retSize
);
clEvent
.
clear
();
// in nanaSec
log
.
info
(
name
+
" event time "
+
"0x"
+
eventAddr
+
": "
+
(
endTime
.
getLong
()
-
startTime
.
getLong
())
+
" ns"
);
endTime
.
clear
();
startTime
.
clear
();
return
result
;
}
else
{
return
clEnqueueNDRangeKernel
(
command_queue
,
kernel
,
work_dim
,
global_work_offset
,
global_work_size
,
local_work_size
,
event_wait_list
,
event
);
}
}
private
void
checkTriLocks
()
{
for
(
int
i
=
0
;
i
<
numberOfFaces
;
i
++)
{
int
lock
=
triLocks
.
get
(
i
);
...
...
@@ -861,7 +726,8 @@ public class CLDistMesh {
return
tmp
;
}
private
void
clearCL
()
{
@Override
protected
void
clearCL
()
throws
OpenCLException
{
clReleaseMemObject
(
clVertices
);
clReleaseMemObject
(
clEdges
);
clReleaseMemObject
(
clTwins
);
...
...
@@ -888,13 +754,7 @@ public class CLDistMesh {
clReleaseKernel
(
clKernelRepair
);
clReleaseKernel
(
clKernelLabelEdges
);
clReleaseKernel
(
clKernelLabelEdgesUpdate
);
clReleaseProgram
(
clProgram
);
clReleaseCommandQueue
(
clQueue
);
clReleaseContext
(
clContext
);
contextCB
.
close
();
programCB
.
close
();
super
.
clearCL
();
}
private
void
clearHost
()
{
...
...
@@ -955,11 +815,10 @@ public class CLDistMesh {
//printResult();
}
public
void
finish
()
{
public
void
finish
()
throws
OpenCLException
{
refresh
();
clearHost
();
clearCL
();
}
private
void
printTri
()
{
...
...
VadereMeshing/src/org/vadere/meshing/opencl/CLDistMeshHE.java
View file @
49b9dc91
...
...
@@ -124,7 +124,7 @@ public class CLDistMeshHE {
private
AMesh
mesh
;
private
boolean
doublePrecision
=
tru
e
;
private
boolean
doublePrecision
=
fals
e
;
private
boolean
profiling
=
false
;
private
boolean
hasToRead
=
false
;
...
...
@@ -189,9 +189,9 @@ public class CLDistMeshHE {
clGetPlatformIDs
(
platformIDs
,
numberOfPlatforms
);
clPlatform
=
platformIDs
.
get
(
0
);
clGetDeviceIDs
(
clPlatform
,
CL_DEVICE_TYPE_
G
PU
,
null
,
numberOfDevices
);
clGetDeviceIDs
(
clPlatform
,
CL_DEVICE_TYPE_
C
PU
,
null
,
numberOfDevices
);
PointerBuffer
deviceIDs
=
stack
.
mallocPointer
(
numberOfDevices
.
get
(
0
));
clGetDeviceIDs
(
clPlatform
,
CL_DEVICE_TYPE_
G
PU
,
deviceIDs
,
numberOfDevices
);
clGetDeviceIDs
(
clPlatform
,
CL_DEVICE_TYPE_
C
PU
,
deviceIDs
,
numberOfDevices
);
clDevice
=
deviceIDs
.
get
(
0
);
printDeviceInfo
(
clDevice
,
"CL_DEVICE_NAME"
,
CL_DEVICE_NAME
);
...
...
Write
Preview
Supports
Markdown
0%
Try again
or
attach a new 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