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
c6e15f01
Commit
c6e15f01
authored
Oct 04, 2017
by
Benedikt Zoennchen
Browse files
Options
Browse Files
Download
Email Patches
Plain Diff
merge the OpenCV to LWJGL changes from the private branch into the develop and resolve conflicts
parent
fddbac91
Changes
20
Hide whitespace changes
Inline
Side-by-side
Showing
20 changed files
with
569 additions
and
660 deletions
+569
-660
VadereGui/src/org/vadere/gui/components/utils/Messages.java
VadereGui/src/org/vadere/gui/components/utils/Messages.java
+0
-2
VadereSimulator/src/org/vadere/simulator/models/density/CLGaussianFilter.java
...org/vadere/simulator/models/density/CLGaussianFilter.java
+1
-1
VadereSimulator/src/org/vadere/simulator/models/density/IGaussianFilter.java
.../org/vadere/simulator/models/density/IGaussianFilter.java
+13
-21
VadereUtils/pom.xml
VadereUtils/pom.xml
+58
-77
VadereUtils/resources/FFT.cl
VadereUtils/resources/FFT.cl
+0
-0
VadereUtils/src/org/vadere/util/math/CLConvolution.java
VadereUtils/src/org/vadere/util/math/CLConvolution.java
+268
-123
VadereUtils/src/org/vadere/util/math/CLFFT.java
VadereUtils/src/org/vadere/util/math/CLFFT.java
+1
-36
VadereUtils/src/org/vadere/util/math/CLFourierTransformation.java
...ils/src/org/vadere/util/math/CLFourierTransformation.java
+0
-62
VadereUtils/src/org/vadere/util/math/CLUtils.java
VadereUtils/src/org/vadere/util/math/CLUtils.java
+18
-12
VadereUtils/src/org/vadere/util/opencl/IOUtil.java
VadereUtils/src/org/vadere/util/opencl/IOUtil.java
+67
-0
VadereUtils/src/org/vadere/util/opencl/InfoUtils.java
VadereUtils/src/org/vadere/util/opencl/InfoUtils.java
+132
-0
VadereUtils/src/org/vadere/util/opencl/kernels/Convolve.cl
VadereUtils/src/org/vadere/util/opencl/kernels/Convolve.cl
+0
-107
VadereUtils/src/org/vadere/util/opencl/kernels/Convolve.java
VadereUtils/src/org/vadere/util/opencl/kernels/Convolve.java
+0
-43
VadereUtils/src/org/vadere/util/opencl/kernels/DFT.cl
VadereUtils/src/org/vadere/util/opencl/kernels/DFT.cl
+0
-44
VadereUtils/src/org/vadere/util/opencl/kernels/DFT.java
VadereUtils/src/org/vadere/util/opencl/kernels/DFT.java
+0
-29
VadereUtils/src/org/vadere/util/opencl/kernels/FFT.java
VadereUtils/src/org/vadere/util/opencl/kernels/FFT.java
+0
-30
VadereUtils/src/org/vadere/util/potential/CellGrid.java
VadereUtils/src/org/vadere/util/potential/CellGrid.java
+0
-5
VadereUtils/tests/org/vadere/util/math/TestConvolution.java
VadereUtils/tests/org/vadere/util/math/TestConvolution.java
+10
-15
VadereUtils/tests/org/vadere/util/math/TestDFT.java
VadereUtils/tests/org/vadere/util/math/TestDFT.java
+1
-27
VadereUtils/tests/org/vadere/util/opencl/TestJavaCL.java
VadereUtils/tests/org/vadere/util/opencl/TestJavaCL.java
+0
-26
No files found.
VadereGui/src/org/vadere/gui/components/utils/Messages.java
View file @
c6e15f01
package
org.vadere.gui.components.utils
;
import
com.android.dx.gen.Local
;
import
javax.swing.*
;
import
org.vadere.gui.projectview.VadereApplication
;
...
...
VadereSimulator/src/org/vadere/simulator/models/density/CLGaussianFilter.java
View file @
c6e15f01
...
...
@@ -18,6 +18,6 @@ class CLGaussianFilter extends GaussianFilter {
@Override
public
void
filterImage
()
{
outputMatrix
=
this
.
convolution
.
convolveS
pe
rate
(
inputMatrix
,
matrixWidth
,
matrixHeight
,
kernel
,
kernelWidth
);
outputMatrix
=
this
.
convolution
.
convolveS
epa
rate
(
inputMatrix
,
matrixWidth
,
matrixHeight
,
kernel
,
kernelWidth
);
}
}
VadereSimulator/src/org/vadere/simulator/models/density/IGaussianFilter.java
View file @
c6e15f01
...
...
@@ -20,9 +20,8 @@ import org.vadere.state.scenario.Topography;
public
interface
IGaussianFilter
{
enum
Type
{
OpenCV
,
// old alternative
OpenCL
,
// default
NativeJava
;
// not jet implemented
NativeJava
// not jet implemented
}
/**
...
...
@@ -87,14 +86,6 @@ public interface IGaussianFilter {
/
(
2
*
Math
.
PI
*
standardDerivation
*
standardDerivation
);
switch
(
type
)
{
case
OpenCV:
{
throw
new
UnsupportedOperationException
();
/*
* return new PedestrianCVGaussianFilter(scenarioBounds, pedestrians, scale,
* scaleFactor, standardDerivation,
* loadingStrategy);
*/
}
case
OpenCL:
{
try
{
BiFunction
<
Integer
,
Integer
,
Float
>
f
=
...
...
@@ -108,8 +99,12 @@ public interface IGaussianFilter {
}
}
default
:
throw
new
IllegalArgumentException
(
type
+
" is not jet supported"
);
}
BiFunction
<
Integer
,
Integer
,
Float
>
f
=
(
centerI
,
i
)
->
(
float
)
(
Math
.
sqrt
(
scaleFactor
)
*
Math
.
exp
(-((
centerI
-
i
)
/
scale
)
*
((
centerI
-
i
)
/
scale
)
/
(
2
*
standardDerivation
*
standardDerivation
)));
IGaussianFilter
clFilter
=
new
JGaussianFilter
(
scenarioBounds
,
scale
,
f
,
false
);
return
new
PedestrianGaussianFilter
(
pedestrians
,
clFilter
,
loadingStrategy
);
}
}
static
IGaussianFilter
create
(
...
...
@@ -122,13 +117,6 @@ public interface IGaussianFilter {
final
Topography
scenario
,
final
double
scale
,
final
boolean
scenarioHasBoundary
,
final
double
standardDerivation
,
final
Type
type
)
{
switch
(
type
)
{
case
OpenCV:
{
throw
new
UnsupportedOperationException
();
/*
* return new ObstacleCVGaussianFilter(scenario, scale, scenarioHasBoundary,
* standardDerivation);
*/
}
case
OpenCL:
{
try
{
double
varianz
=
standardDerivation
*
standardDerivation
;
...
...
@@ -142,8 +130,12 @@ public interface IGaussianFilter {
}
}
default
:
throw
new
IllegalArgumentException
(
type
+
" is not jet supported"
);
}
double
varianz
=
standardDerivation
*
standardDerivation
;
BiFunction
<
Integer
,
Integer
,
Float
>
f
=
(
centerI
,
i
)
->
(
float
)
((
1.0
/
(
2
*
Math
.
PI
*
varianz
))
*
Math
.
exp
(-((
centerI
-
i
)
/
scale
)
*
((
centerI
-
i
)
/
scale
)
/
(
2
*
varianz
)));
IGaussianFilter
clFilter
=
new
JGaussianFilter
(
scenario
.
getBounds
(),
scale
,
f
,
true
);
return
new
ObstacleGaussianFilter
(
scenario
,
clFilter
);
}
}
static
IGaussianFilter
create
(
...
...
VadereUtils/pom.xml
View file @
c6e15f01
...
...
@@ -25,22 +25,6 @@
</resource>
</resources>
<plugins>
<plugin>
<groupId>
com.nativelibs4java
</groupId>
<artifactId>
maven-javacl-plugin
</artifactId>
<version>
1.0.0-RC4
</version>
<configuration>
<sourcesDirectory>
${project.build.directory}/../src/org/vadere/util/opencl
</sourcesDirectory>
</configuration>
<executions>
<execution>
<phase>
generate-sources
</phase>
<goals>
<goal>
compile
</goal>
</goals>
</execution>
</executions>
</plugin>
<plugin>
<groupId>
org.apache.maven.plugins
</groupId>
<artifactId>
maven-compiler-plugin
</artifactId>
...
...
@@ -51,75 +35,72 @@
</configuration>
</plugin>
</plugins>
<pluginManagement>
<plugins>
<!--This plugin's configuration is used to store Eclipse m2e settings only. It has no influence on the Maven build itself.-->
<plugin>
<groupId>
org.eclipse.m2e
</groupId>
<artifactId>
lifecycle-mapping
</artifactId>
<version>
1.0.0
</version>
<configuration>
<lifecycleMappingMetadata>
<pluginExecutions>
<pluginExecution>
<pluginExecutionFilter>
<groupId>
com.nativelibs4java
</groupId>
<artifactId>
maven-javacl-plugin
</artifactId>
<versionRange>
[1.0.0-RC4,)
</versionRange>
<goals>
<goal>
compile
</goal>
</goals>
</pluginExecutionFilter>
<action>
<execute>
<runOnIncremental>
false
</runOnIncremental>
</execute>
</action>
</pluginExecution>
</pluginExecutions>
</lifecycleMappingMetadata>
</configuration>
</plugin>
</plugins>
</pluginManagement>
</build>
<repositories>
<repository>
<id>
nativelibs4java-repo
</id>
<name>
nativelibs4java Maven2 Repository
</name>
<url>
http://nativelibs4java.sourceforge.net/maven
</url>
</repository>
</repositories>
<!-- generated by lwjgl -->
<properties>
<project.build.sourceEncoding>
UTF-8
</project.build.sourceEncoding>
<maven.compiler.source>
1.8
</maven.compiler.source>
<maven.compiler.target>
1.8
</maven.compiler.target>
<lwjgl.version>
3.1.2
</lwjgl.version>
</properties>
<!-- end generated by lwjgl -->
<dependencies>
<dependency>
<groupId>
com.nativelibs4java
</groupId>
<artifactId>
bridj
</artifactId>
<version>
0.7.0
</version>
<scope>
compile
</scope>
<groupId>
com.vividsolutions
</groupId>
<artifactId>
jts
</artifactId>
<version>
1.13
</version>
</dependency>
<!-- generated by lwjgl -->
<dependency>
<groupId>
com.nativelibs4java
</groupId>
<artifactId>
javacl
</artifactId>
<version>
1.0.0-RC4
</version>
<scope>
compile
</scope>
<groupId>
org.lwjgl
</groupId>
<artifactId>
lwjgl
</artifactId>
<version>
${lwjgl.version}
</version>
</dependency>
<dependency>
<groupId>
com.vividsolutions
</groupId>
<artifactId>
jts
</artifactId>
<version>
1.13
</version>
<groupId>
org.lwjgl
</groupId>
<artifactId>
lwjgl-opencl
</artifactId>
<version>
${lwjgl.version}
</version>
</dependency>
<dependency>
<groupId>
org.lwjgl
</groupId>
<artifactId>
lwjgl
</artifactId>
<version>
${lwjgl.version}
</version>
<classifier>
${lwjgl.natives}
</classifier>
<scope>
runtime
</scope>
</dependency>
</dependencies>
<properties>
<project.build.sourceEncoding>
UTF-8
</project.build.sourceEncoding>
</properties>
<!--<packaging>pom</packaging> -->
</dependencies>
<!-- end generated by lwjgl -->
<profiles>
<profile>
<id>
lwjgl-natives-linux
</id>
<activation>
<os><family>
unix
</family></os>
</activation>
<properties>
<lwjgl.natives>
natives-linux
</lwjgl.natives>
</properties>
</profile>
<profile>
<id>
lwjgl-natives-macos
</id>
<activation>
<os><family>
mac
</family></os>
</activation>
<properties>
<lwjgl.natives>
natives-macos
</lwjgl.natives>
</properties>
</profile>
<profile>
<id>
lwjgl-natives-windows
</id>
<activation>
<os><family>
windows
</family></os>
</activation>
<properties>
<lwjgl.natives>
natives-windows
</lwjgl.natives>
</properties>
</profile>
</profiles>
<!--<packaging>pom</packaging> -->
</project>
\ No newline at end of file
VadereUtils/
src/org/vadere/util/opencl/kernel
s/FFT.cl
→
VadereUtils/
resource
s/FFT.cl
View file @
c6e15f01
File moved
VadereUtils/src/org/vadere/util/math/CLConvolution.java
View file @
c6e15f01
package
org.vadere.util.math
;
import
com.nativelibs4java.opencl.*
;
import
org.bridj.Pointer
;
import
org.vadere.util.opencl.kernels.Convolve
;
import
org.apache.log4j.LogManager
;
import
org.apache.log4j.Logger
;
import
org.lwjgl.BufferUtils
;
import
org.lwjgl.PointerBuffer
;
import
org.lwjgl.opencl.CLContextCallback
;
import
org.lwjgl.opencl.CLProgramCallback
;
import
org.lwjgl.system.MemoryStack
;
import
org.vadere.util.opencl.IOUtil
;
import
org.vadere.util.opencl.InfoUtils
;
import
java.io.IOException
;
import
java.nio.ByteBuffer
;
import
java.nio.FloatBuffer
;
import
java.nio.IntBuffer
;
import
static
org
.
lwjgl
.
opencl
.
CL10
.*;
import
static
org
.
lwjgl
.
system
.
MemoryUtil
.
NULL
;
import
static
org
.
lwjgl
.
system
.
MemoryUtil
.
memUTF8
;
/**
* @author Benedikt Zoennchen
*/
public
class
CLConvolution
{
private
static
Logger
log
=
LogManager
.
getLogger
(
CLConvolution
.
class
);
private
Convolution
gaussianFilter
;
// CL ids
private
MemoryStack
stack
;
private
long
clPlatform
;
private
long
clDevice
;
private
long
clContext
;
private
long
clQueue
;
private
long
clProgram
;
// error code buffer
private
IntBuffer
errcode_ret
;
// CL Memory
private
long
clInput
;
private
long
clOutput
;
private
long
clGaussianKernel
;
private
long
clTmp
;
// Host Memory
private
FloatBuffer
hostScenario
;
private
FloatBuffer
hostGaussKernel
;
private
FloatBuffer
output
;
// CL callbacks
private
CLContextCallback
contextCB
;
private
CLProgramCallback
programCB
;
// CL kernel
private
long
clKernelConvolve
;
private
long
clKernelConvolveRow
;
private
long
clKernelConvolveCol
;
public
CLConvolution
()
{
this
.
stack
=
MemoryStack
.
stackPush
();
}
public
void
init
()
{
initCallbacks
();
initCL
();
buildProgram
();
}
public
float
[]
convolve
(
final
float
[]
input
,
final
int
matrixWidth
,
final
int
matrixHeight
,
final
float
[]
kernel
,
final
int
kernelWidth
)
{
init
();
float
[]
result
=
convolve
(
input
,
matrixWidth
,
matrixHeight
,
kernel
,
kernelWidth
,
clKernelConvolve
);
clearCL
();
clReleaseKernel
(
clKernelConvolve
);
return
result
;
}
public
float
[]
convolveRow
(
final
float
[]
input
,
final
int
matrixWidth
,
final
int
matrixHeight
,
final
float
[]
kernel
,
final
int
kernelWidth
)
{
init
();
float
[]
result
=
convolve
(
input
,
matrixWidth
,
matrixHeight
,
kernel
,
kernelWidth
,
clKernelConvolveRow
);
clearCL
();
clReleaseKernel
(
clKernelConvolveRow
);
return
result
;
}
public
float
[]
convolveCol
(
final
float
[]
input
,
final
int
matrixWidth
,
final
int
matrixHeight
,
final
float
[]
kernel
,
final
int
kernelWidth
)
{
init
();
float
[]
result
=
convolve
(
input
,
matrixWidth
,
matrixHeight
,
kernel
,
kernelWidth
,
clKernelConvolveCol
);
clearCL
();
clReleaseKernel
(
clKernelConvolveCol
);
return
result
;
}
public
float
[]
convolveSeparate
(
final
float
[]
input
,
final
int
matrixWidth
,
final
int
matrixHeight
,
final
float
[]
kernel
,
final
int
kernelWidth
)
{
assert
matrixWidth
*
matrixHeight
==
input
.
length
;
init
();
hostScenario
=
CLUtils
.
toFloatBuffer
(
input
);
output
=
CLUtils
.
toFloatBuffer
(
input
);
hostGaussKernel
=
CLUtils
.
toFloatBuffer
(
kernel
);
// host memory to gpu memory
clInput
=
clCreateBuffer
(
clContext
,
CL_MEM_READ_ONLY
|
CL_MEM_COPY_HOST_PTR
,
hostScenario
,
errcode_ret
);
clTmp
=
clCreateBuffer
(
clContext
,
CL_MEM_READ_WRITE
,
4
*
input
.
length
,
errcode_ret
);
clOutput
=
clCreateBuffer
(
clContext
,
CL_MEM_WRITE_ONLY
,
4
*
input
.
length
,
errcode_ret
);
clGaussianKernel
=
clCreateBuffer
(
clContext
,
CL_MEM_READ_ONLY
|
CL_MEM_COPY_HOST_PTR
,
hostGaussKernel
,
errcode_ret
);
clSetKernelArg1p
(
clKernelConvolveCol
,
0
,
clInput
);
clSetKernelArg1p
(
clKernelConvolveCol
,
1
,
clGaussianKernel
);
clSetKernelArg1p
(
clKernelConvolveCol
,
2
,
clTmp
);
clSetKernelArg1i
(
clKernelConvolveCol
,
3
,
matrixWidth
);
clSetKernelArg1i
(
clKernelConvolveCol
,
4
,
matrixHeight
);
clSetKernelArg1i
(
clKernelConvolveCol
,
5
,
kernelWidth
);
clSetKernelArg1p
(
clKernelConvolveRow
,
0
,
clTmp
);
clSetKernelArg1p
(
clKernelConvolveRow
,
1
,
clGaussianKernel
);
clSetKernelArg1p
(
clKernelConvolveRow
,
2
,
clOutput
);
clSetKernelArg1i
(
clKernelConvolveRow
,
3
,
matrixWidth
);
clSetKernelArg1i
(
clKernelConvolveRow
,
4
,
matrixHeight
);
clSetKernelArg1i
(
clKernelConvolveRow
,
5
,
kernelWidth
);
PointerBuffer
clGlobalWorkSizeEdges
=
BufferUtils
.
createPointerBuffer
(
2
);
clGlobalWorkSizeEdges
.
put
(
0
,
matrixWidth
);
clGlobalWorkSizeEdges
.
put
(
1
,
matrixHeight
);
// run the kernel and read the result
clEnqueueNDRangeKernel
(
clQueue
,
clKernelConvolveCol
,
2
,
null
,
clGlobalWorkSizeEdges
,
null
,
null
,
null
);
clEnqueueNDRangeKernel
(
clQueue
,
clKernelConvolveRow
,
2
,
null
,
clGlobalWorkSizeEdges
,
null
,
null
,
null
);
clFinish
(
clQueue
);
clEnqueueReadBuffer
(
clQueue
,
clOutput
,
true
,
0
,
output
,
null
,
null
);
float
[]
foutput
=
CLUtils
.
toFloatArray
(
output
,
input
.
length
);
clearCL
();
clReleaseKernel
(
clTmp
);
clReleaseKernel
(
clKernelConvolve
);
clReleaseKernel
(
clKernelConvolveRow
);
clReleaseKernel
(
clKernelConvolveCol
);
return
foutput
;
}
private
float
[]
convolve
(
final
float
[]
input
,
final
int
matrixWidth
,
final
int
matrixHeight
,
final
float
[]
kernel
,
final
int
kernelWidth
,
final
long
clKernel
)
{
assert
matrixWidth
*
matrixHeight
==
input
.
length
;
setArguments
(
input
,
matrixWidth
,
matrixHeight
,
kernel
,
kernelWidth
,
clKernel
);
PointerBuffer
clGlobalWorkSizeEdges
=
BufferUtils
.
createPointerBuffer
(
2
);
clGlobalWorkSizeEdges
.
put
(
0
,
matrixWidth
);
clGlobalWorkSizeEdges
.
put
(
1
,
matrixHeight
);
// run the kernel and read the result
clEnqueueNDRangeKernel
(
clQueue
,
clKernel
,
2
,
null
,
clGlobalWorkSizeEdges
,
null
,
null
,
null
);
clFinish
(
clQueue
);
clEnqueueReadBuffer
(
clQueue
,
clOutput
,
true
,
0
,
output
,
null
,
null
);
float
[]
foutput
=
CLUtils
.
toFloatArray
(
output
,
input
.
length
);
return
foutput
;
}
private
void
setArguments
(
final
float
[]
input
,
final
int
matrixWidth
,
final
int
matrixHeight
,
final
float
[]
kernel
,
final
int
kernelWidth
,
final
long
clKernel
)
{
hostScenario
=
CLUtils
.
toFloatBuffer
(
input
);
output
=
CLUtils
.
toFloatBuffer
(
input
);
hostGaussKernel
=
CLUtils
.
toFloatBuffer
(
kernel
);
// host memory to gpu memory
clInput
=
clCreateBuffer
(
clContext
,
CL_MEM_READ_ONLY
|
CL_MEM_COPY_HOST_PTR
,
hostScenario
,
errcode_ret
);
clOutput
=
clCreateBuffer
(
clContext
,
CL_MEM_WRITE_ONLY
,
4
*
input
.
length
,
errcode_ret
);
clGaussianKernel
=
clCreateBuffer
(
clContext
,
CL_MEM_READ_ONLY
|
CL_MEM_COPY_HOST_PTR
,
hostGaussKernel
,
errcode_ret
);
clSetKernelArg1p
(
clKernel
,
0
,
clInput
);
clSetKernelArg1p
(
clKernel
,
1
,
clGaussianKernel
);
clSetKernelArg1p
(
clKernel
,
2
,
clOutput
);
clSetKernelArg1i
(
clKernel
,
3
,
matrixWidth
);
clSetKernelArg1i
(
clKernel
,
4
,
matrixHeight
);
clSetKernelArg1i
(
clKernel
,
5
,
kernelWidth
);
}
private
void
clearCL
()
{
// release memory and devices
contextCB
.
free
();
programCB
.
free
();
clReleaseMemObject
(
clInput
);
clReleaseMemObject
(
clOutput
);
clReleaseCommandQueue
(
clQueue
);
clReleaseProgram
(
clProgram
);
clReleaseContext
(
clContext
);
}
// private helpers
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
)
->
{
log
.
info
(
"The cl_program [0x"
+
program
+
"] was built "
+
(
InfoUtils
.
getProgramBuildInfoInt
(
program
,
clDevice
,
CL_PROGRAM_BUILD_STATUS
)
==
CL_SUCCESS
?
"successfully"
:
"unsuccessfully"
));
});
}
private
void
initCL
()
{
// helper for the memory allocation in java
//stack = MemoryStack.stackPush();
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
);
InfoUtils
.
checkCLError
(
errcode_ret
);
clQueue
=
clCreateCommandQueue
(
clContext
,
clDevice
,
0
,
errcode_ret
);
}
private
void
buildProgram
()
{
PointerBuffer
strings
=
BufferUtils
.
createPointerBuffer
(
1
);
PointerBuffer
lengths
=
BufferUtils
.
createPointerBuffer
(
1
);
ByteBuffer
source
;
try
{
source
=
IOUtil
.
ioResourceToByteBuffer
(
"Convolve.cl"
,
4096
);
}
catch
(
IOException
e
)
{
throw
new
RuntimeException
(
e
);
}
strings
.
put
(
0
,
source
);
lengths
.
put
(
0
,
source
.
remaining
());
clProgram
=
clCreateProgramWithSource
(
clContext
,
strings
,
lengths
,
errcode_ret
);
int
errcode
=
clBuildProgram
(
clProgram
,
clDevice
,
""
,
programCB
,
NULL
);
InfoUtils
.
checkCLError
(
errcode
);
clKernelConvolve
=
clCreateKernel
(
clProgram
,
"convolve"
,
errcode_ret
);
clKernelConvolveRow
=
clCreateKernel
(
clProgram
,
"convolveRow"
,
errcode_ret
);
clKernelConvolveCol
=
clCreateKernel
(
clProgram
,
"convolveCol"
,
errcode_ret
);
}
private
CLContext
context
;
private
Convolve
gaussianFilter
;
public
CLConvolution
()
throws
IOException
{
//context = JavaCL.createBestContext(CLPlatform.DeviceFeature.CPU, CLPlatform.DeviceFeature.MaxComputeUnits);
context
=
JavaCL
.
createBestContext
();
gaussianFilter
=
new
Convolve
(
context
);
}
public
CLConvolution
(
CLPlatform
.
DeviceFeature
...
deviceFeature
)
throws
IOException
{
context
=
JavaCL
.
createBestContext
(
deviceFeature
);
gaussianFilter
=
new
Convolve
(
context
);
}
public
float
[]
convolve
(
final
float
[]
input
,
final
int
matrixWidth
,
final
int
matrixHeight
,
final
float
[]
kernel
,
final
int
kernelWidth
)
throws
IOException
{
float
[]
output
=
new
float
[
input
.
length
];
CLQueue
queue
=
context
.
createDefaultQueue
();
CLBuffer
<
Float
>
clInput
=
doubleArrayToCLBuffer
(
input
);
CLBuffer
<
Float
>
clKernel
=
doubleArrayToCLBuffer
(
kernel
);
CLBuffer
<
Float
>
clOutput
=
doubleArrayToCLBuffer
(
output
);
// timer.start();
CLEvent
convolveEvt
=
gaussianFilter
.
convolve
(
queue
,
clInput
,
clKernel
,
clOutput
,
matrixWidth
,
matrixHeight
,
kernelWidth
,
new
int
[]
{
matrixWidth
,
input
.
length
/
matrixWidth
},
null
);
queue
.
finish
();
convolveEvt
.
waitFor
();
Pointer
<
Float
>
outPtr
=
clOutput
.
read
(
queue
,
convolveEvt
);
for
(
int
i
=
0
;
i
<
output
.
length
;
i
++)
{
output
[
i
]
=
outPtr
.
get
(
i
);
}
outPtr
.
release
();
return
output
;
}
public
float
[]
convolveSperate
(
final
float
[]
input
,
final
int
matrixWidth
,
final
int
matrixHeight
,
final
float
[]
kernel
,
final
int
kernelWidth
)
{
float
[]
output
=
new
float
[
input
.
length
];
float
[]
tmp
=
new
float
[
input
.
length
];
CLQueue
queue
=
context
.
createDefaultQueue
();
CLBuffer
<
Float
>
clInput
=
doubleArrayToCLBuffer
(
input
);
CLBuffer
<
Float
>
clKernel
=
doubleArrayToCLBuffer
(
kernel
);
CLBuffer
<
Float
>
clOutput
=
doubleArrayToCLBuffer
(
output
);
CLBuffer
<
Float
>
clTmp
=
doubleArrayToCLBuffer
(
tmp
);
CLEvent
convolveEvtCol
=
gaussianFilter
.
convolveCol
(
queue
,
clInput
,
clKernel
,
clTmp
,
matrixWidth
,
matrixHeight
,
kernelWidth
,
new
int
[]
{
matrixWidth
,
input
.
length
/
matrixWidth
},
null
);
convolveEvtCol
.
waitFor
();
CLEvent
convolveEvtRow
=
gaussianFilter
.
convolveRow
(
queue
,
clTmp
,
clKernel
,
clOutput
,
matrixWidth
,
matrixHeight
,
kernelWidth
,
new
int
[]
{
matrixWidth
,
input
.
length
/
matrixWidth
},
null
);
queue
.
finish
();
convolveEvtRow
.
waitFor
();
Pointer
<
Float
>
outPtr
=
clOutput
.
read
(
queue
,
convolveEvtRow
);
for
(
int
i
=
0
;
i
<
output
.
length
;
i
++)
{
output
[
i
]
=
outPtr
.
get
(
i
);
}
outPtr
.
release
();
return
output
;
}
public
float
[]
convolveCol
(
final
float
[]
input
,
final
int
matrixWidth
,
final
int
matrixHeight
,
final
float
[]
kernel
,
final
int
kernelWidth
)
{
float
[]
output
=
new
float
[
input
.
length
];
CLQueue
queue
=
context
.
createDefaultQueue
();
CLBuffer
<
Float
>
clInput
=
doubleArrayToCLBuffer
(
input
);
CLBuffer
<
Float
>
clKernel
=
doubleArrayToCLBuffer
(
kernel
);
CLBuffer
<
Float
>
clOutput
=
doubleArrayToCLBuffer
(
output
);
CLEvent
convolveEvtCol
=
gaussianFilter
.
convolveCol
(
queue
,
clInput
,
clKernel
,
clOutput
,
matrixWidth
,
matrixHeight
,
kernelWidth
,
new
int
[]
{
matrixWidth
,
input
.
length
/
matrixWidth
},
null
);
convolveEvtCol
.
waitFor
();
queue
.
finish
();
Pointer
<
Float
>
outPtr
=
clOutput
.
read
(
queue
,
convolveEvtCol
);
for
(
int
i
=
0
;
i
<
output
.
length
;
i
++)
{
output
[
i
]
=
outPtr
.
get
(
i
);
}
outPtr
.
release
();
return
output
;
}
public
float
[]
convolveRow
(
final
float
[]
input
,
final
int
matrixWidth
,
final
int
matrixHeight
,
final
float
[]
kernel
,
final
int
kernelWidth
)
{
float
[]
output
=
new
float
[
input
.
length
];
float
[]
tmp
=
new
float
[
input
.
length
];
CLQueue
queue
=
context
.
createDefaultQueue
();