diff --git a/jme3-android/src/main/java/com/jme3/system/android/OGLESContext.java b/jme3-android/src/main/java/com/jme3/system/android/OGLESContext.java
index 991ad25c0..35c3c3911 100644
--- a/jme3-android/src/main/java/com/jme3/system/android/OGLESContext.java
+++ b/jme3-android/src/main/java/com/jme3/system/android/OGLESContext.java
@@ -457,4 +457,10 @@ public class OGLESContext implements JmeContext, GLSurfaceView.Renderer, SoftTex
}
});
}
+
+ @Override
+ public com.jme3.opencl.Context getOpenCLContext() {
+ logger.warning("OpenCL is not yet supported on android");
+ return null;
+ }
}
diff --git a/jme3-core/src/main/java/com/jme3/opencl/AbstractOpenCLObject.java b/jme3-core/src/main/java/com/jme3/opencl/AbstractOpenCLObject.java
new file mode 100644
index 000000000..fce0f481d
--- /dev/null
+++ b/jme3-core/src/main/java/com/jme3/opencl/AbstractOpenCLObject.java
@@ -0,0 +1,62 @@
+/*
+ * Copyright (c) 2009-2016 jMonkeyEngine
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are
+ * met:
+ *
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ *
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ *
+ * * Neither the name of 'jMonkeyEngine' nor the names of its contributors
+ * may be used to endorse or promote products derived from this software
+ * without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+ * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+ * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
+ * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
+ * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+package com.jme3.opencl;
+
+/**
+ * Abstract implementation of {@link OpenCLObject} providing the release
+ * mechanisms.
+ * @author Sebastian Weiss
+ */
+public abstract class AbstractOpenCLObject implements OpenCLObject {
+
+ protected final ObjectReleaser releaser;
+ protected AbstractOpenCLObject(ObjectReleaser releaser) {
+ this.releaser = releaser;
+ }
+ @Override
+ public void register() {
+ OpenCLObjectManager.getInstance().registerObject(this);
+ }
+ @Override
+ public void release() {
+ releaser.release();
+ }
+ @Override
+ @SuppressWarnings("FinalizeDeclaration")
+ protected void finalize() throws Throwable {
+ release();
+ }
+ @Override
+ public ObjectReleaser getReleaser() {
+ return releaser;
+ }
+}
diff --git a/jme3-core/src/main/java/com/jme3/opencl/Buffer.java b/jme3-core/src/main/java/com/jme3/opencl/Buffer.java
new file mode 100644
index 000000000..ebb7d2bb2
--- /dev/null
+++ b/jme3-core/src/main/java/com/jme3/opencl/Buffer.java
@@ -0,0 +1,427 @@
+/*
+ * Copyright (c) 2009-2016 jMonkeyEngine
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are
+ * met:
+ *
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ *
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ *
+ * * Neither the name of 'jMonkeyEngine' nor the names of its contributors
+ * may be used to endorse or promote products derived from this software
+ * without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+ * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+ * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
+ * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
+ * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+package com.jme3.opencl;
+
+import java.nio.ByteBuffer;
+
+/**
+ * Wrapper for an OpenCL buffer object.
+ * A buffer object stores a one-dimensional collection of elements. Elements of a buffer object can
+ * be a scalar data type (such as an int, float), vector data type, or a user-defined structure.
+ *
+ * Buffers are created by the {@link Context}.
+ *
+ * All access methods (read/write/copy/map) are available in both sychronized/blocking versions
+ * and in async/non-blocking versions. The later ones always return an {@link Event} object
+ * and have the prefix -Async in their name.
+ *
+ * @see Context#createBuffer(long, com.jme3.opencl.MemoryAccess)
+ * @author shaman
+ */
+public abstract class Buffer extends AbstractOpenCLObject {
+
+ protected Buffer(ObjectReleaser releaser) {
+ super(releaser);
+ }
+
+ /**
+ * @return the size of the buffer in bytes.
+ * @see Context#createBuffer(long)
+ */
+ public abstract long getSize();
+
+ /**
+ * @return the memory access flags set on creation.
+ * @see Context#createBuffer(long, com.jme3.opencl.MemoryAccess)
+ */
+ public abstract MemoryAccess getMemoryAccessFlags();
+
+ /**
+ * Performs a blocking read of the buffer.
+ * The target buffer must have at least {@code size} bytes remaining.
+ * This method may set the limit to the last byte read.
+ * @param queue the command queue
+ * @param dest the target buffer
+ * @param size the size in bytes being read
+ * @param offset the offset in bytes in the buffer to read from
+ */
+ public abstract void read(CommandQueue queue, ByteBuffer dest, long size, long offset);
+
+ /**
+ * Alternative version of {@link #read(com.jme3.opencl.CommandQueue, java.nio.ByteBuffer, long, long) },
+ * sets {@code offset} to zero.
+ */
+ public void read(CommandQueue queue, ByteBuffer dest, long size) {
+ read(queue, dest, size, 0);
+ }
+
+ /**
+ * Alternative version of {@link #read(com.jme3.opencl.CommandQueue, java.nio.ByteBuffer, long) },
+ * sets {@code size} to {@link #getSize() }.
+ */
+ public void read(CommandQueue queue, ByteBuffer dest) {
+ read(queue, dest, getSize());
+ }
+
+ /**
+ * Performs an async/non-blocking read of the buffer.
+ * The target buffer must have at least {@code size} bytes remaining.
+ * This method may set the limit to the last byte read.
+ * @param queue the command queue
+ * @param dest the target buffer
+ * @param size the size in bytes being read
+ * @param offset the offset in bytes in the buffer to read from
+ * @return the event indicating when the memory has been fully read into the provided buffer
+ */
+ public abstract Event readAsync(CommandQueue queue, ByteBuffer dest, long size, long offset);
+
+ /**
+ * Alternative version of {@link #readAsync(com.jme3.opencl.CommandQueue, java.nio.ByteBuffer, long, long) },
+ * sets {@code offset} to zero.
+ */
+ public Event readAsync(CommandQueue queue, ByteBuffer dest, long size) {
+ return readAsync(queue, dest, size, 0);
+ }
+
+ /**
+ * Alternative version of {@link #readAsync(com.jme3.opencl.CommandQueue, java.nio.ByteBuffer, long) },
+ * sets {@code size} to {@link #getSize() }
+ */
+ public Event readAsync(CommandQueue queue, ByteBuffer dest) {
+ return readAsync(queue, dest, getSize());
+ }
+
+ /**
+ * Performs a blocking write to the buffer.
+ * The target buffer must have at least {@code size} bytes remaining.
+ * This method may set the limit to the last byte that will be written.
+ * @param queue the command queue
+ * @param src the source buffer, its data is written to this buffer
+ * @param size the size in bytes to write
+ * @param offset the offset into the target buffer
+ */
+ public abstract void write(CommandQueue queue, ByteBuffer src, long size, long offset);
+
+ /**
+ * Alternative version of {@link #write(com.jme3.opencl.CommandQueue, java.nio.ByteBuffer, long, long) },
+ * sets {@code offset} to zero.
+ */
+ public void write(CommandQueue queue, ByteBuffer src, long size) {
+ write(queue, src, size, 0);
+ }
+
+ /**
+ * Alternative version of {@link #write(com.jme3.opencl.CommandQueue, java.nio.ByteBuffer, long) },
+ * sets {@code size} to {@link #getSize() }.
+ */
+ public void write(CommandQueue queue, ByteBuffer src) {
+ write(queue, src, getSize());
+ }
+
+ /**
+ * Performs an async/non-blocking write to the buffer.
+ * The target buffer must have at least {@code size} bytes remaining.
+ * This method may set the limit to the last byte that will be written.
+ * @param queue the command queue
+ * @param src the source buffer, its data is written to this buffer
+ * @param size the size in bytes to write
+ * @param offset the offset into the target buffer
+ * @return the event object indicating when the write operation is completed
+ */
+ public abstract Event writeAsync(CommandQueue queue, ByteBuffer src, long size, long offset);
+
+ /**
+ * Alternative version of {@link #writeAsync(com.jme3.opencl.CommandQueue, java.nio.ByteBuffer, long, long) },
+ * sets {@code offset} to zero.
+ */
+ public Event writeAsync(CommandQueue queue, ByteBuffer src, long size) {
+ return writeAsync(queue, src, size, 0);
+ }
+
+ /**
+ * Alternative version of {@link #writeAsync(com.jme3.opencl.CommandQueue, java.nio.ByteBuffer, long) },
+ * sets {@code size} to {@link #getSize() }.
+ */
+ public Event writeAsync(CommandQueue queue, ByteBuffer src) {
+ return writeAsync(queue, src, getSize());
+ }
+
+ /**
+ * Performs a blocking copy operation from this buffer to the specified buffer.
+ * @param queue the command queue
+ * @param dest the target buffer
+ * @param size the size in bytes to copy
+ * @param srcOffset offset in bytes into this buffer
+ * @param destOffset offset in bytes into the target buffer
+ */
+ public abstract void copyTo(CommandQueue queue, Buffer dest, long size, long srcOffset, long destOffset);
+
+ /**
+ * Alternative version of {@link #copyTo(com.jme3.opencl.CommandQueue, com.jme3.opencl.Buffer, long, long, long) },
+ * sets {@code srcOffset} and {@code destOffset} to zero.
+ */
+ public void copyTo(CommandQueue queue, Buffer dest, long size) {
+ copyTo(queue, dest, size, 0, 0);
+ }
+
+ /**
+ * Alternative version of {@link #copyTo(com.jme3.opencl.CommandQueue, com.jme3.opencl.Buffer, long) },
+ * sets {@code size} to {@code this.getSize()}.
+ */
+ public void copyTo(CommandQueue queue, Buffer dest) {
+ copyTo(queue, dest, getSize());
+ }
+
+ /**
+ * Performs an async/non-blocking copy operation from this buffer to the specified buffer.
+ * @param queue the command queue
+ * @param dest the target buffer
+ * @param size the size in bytes to copy
+ * @param srcOffset offset in bytes into this buffer
+ * @param destOffset offset in bytes into the target buffer
+ * @return the event object indicating when the copy operation is finished
+ */
+ public abstract Event copyToAsync(CommandQueue queue, Buffer dest, long size, long srcOffset, long destOffset);
+
+ /**
+ * Alternative version of {@link #copyToAsync(com.jme3.opencl.CommandQueue, com.jme3.opencl.Buffer, long, long, long) },
+ * sets {@code srcOffset} and {@code destOffset} to zero.
+ */
+ public Event copyToAsync(CommandQueue queue, Buffer dest, long size) {
+ return copyToAsync(queue, dest, size, 0, 0);
+ }
+
+ /**
+ * Alternative version of {@link #copyToAsync(com.jme3.opencl.CommandQueue, com.jme3.opencl.Buffer, long) },
+ * sets {@code size} to {@code this.getSize()}.
+ */
+ public Event copyToAsync(CommandQueue queue, Buffer dest) {
+ return copyToAsync(queue, dest, getSize());
+ }
+
+ /**
+ * Maps this buffer directly into host memory. This might be the fastest method
+ * to access the contents of the buffer since the OpenCL implementation directly
+ * provides the memory.
+ * Important: The mapped memory MUST be released by calling
+ * {@link #unmap(com.jme3.opencl.CommandQueue, java.nio.ByteBuffer) }.
+ * @param queue the command queue
+ * @param size the size in bytes to map
+ * @param offset the offset into this buffer
+ * @param access specifies the possible access to the memory: READ_ONLY, WRITE_ONLY, READ_WRITE
+ * @return the byte buffer directly reflecting the buffer contents
+ */
+ public abstract ByteBuffer map(CommandQueue queue, long size, long offset, MappingAccess access);
+
+ /**
+ * Alternative version of {@link #map(com.jme3.opencl.CommandQueue, long, long, com.jme3.opencl.MappingAccess) },
+ * sets {@code offset} to zero.
+ * Important: The mapped memory MUST be released by calling
+ * {@link #unmap(com.jme3.opencl.CommandQueue, java.nio.ByteBuffer) }.
+ */
+ public ByteBuffer map(CommandQueue queue, long size, MappingAccess access) {
+ return map(queue, size, 0, access);
+ }
+
+ /**
+ * Alternative version of {@link #map(com.jme3.opencl.CommandQueue, long, com.jme3.opencl.MappingAccess) },
+ * sets {@code size} to {@link #getSize() }.
+ * Important: The mapped memory MUST be released by calling
+ * {@link #unmap(com.jme3.opencl.CommandQueue, java.nio.ByteBuffer) }.
+ */
+ public ByteBuffer map(CommandQueue queue, MappingAccess access) {
+ return map(queue, getSize(), access);
+ }
+
+ /**
+ * Unmaps a previously mapped memory.
+ * This releases the native resources and for WRITE_ONLY or READ_WRITE access,
+ * the memory content is sent back to the GPU.
+ * @param queue the command queue
+ * @param ptr the buffer that was previously mapped
+ */
+ public abstract void unmap(CommandQueue queue, ByteBuffer ptr);
+
+ /**
+ * Maps this buffer asynchronously into host memory. This might be the fastest method
+ * to access the contents of the buffer since the OpenCL implementation directly
+ * provides the memory.
+ * Important: The mapped memory MUST be released by calling
+ * {@link #unmap(com.jme3.opencl.CommandQueue, java.nio.ByteBuffer) }.
+ * @param queue the command queue
+ * @param size the size in bytes to map
+ * @param offset the offset into this buffer
+ * @param access specifies the possible access to the memory: READ_ONLY, WRITE_ONLY, READ_WRITE
+ * @return the byte buffer directly reflecting the buffer contents
+ * and the event indicating when the buffer contents are available
+ */
+ public abstract AsyncMapping mapAsync(CommandQueue queue, long size, long offset, MappingAccess access);
+ /**
+ * Alternative version of {@link #mapAsync(com.jme3.opencl.CommandQueue, long, long, com.jme3.opencl.MappingAccess) },
+ * sets {@code offset} to zero.
+ * Important: The mapped memory MUST be released by calling
+ * {@link #unmap(com.jme3.opencl.CommandQueue, java.nio.ByteBuffer) }.
+ */
+ public AsyncMapping mapAsync(CommandQueue queue, long size, MappingAccess access) {
+ return mapAsync(queue, size, 0, access);
+ }
+ /**
+ * Alternative version of {@link #mapAsync(com.jme3.opencl.CommandQueue, long, com.jme3.opencl.MappingAccess) },
+ * sets {@code size} to {@link #getSize() }.
+ * Important: The mapped memory MUST be released by calling
+ * {@link #unmap(com.jme3.opencl.CommandQueue, java.nio.ByteBuffer) }.
+ */
+ public AsyncMapping mapAsync(CommandQueue queue, MappingAccess access) {
+ return mapAsync(queue, getSize(), 0, access);
+ }
+
+ /**
+ * Enqueues a fill operation. This method can be used to initialize or clear
+ * a buffer with a certain value.
+ * @param queue the command queue
+ * @param pattern the buffer containing the filling pattern.
+ * The remaining bytes specify the pattern length
+ * @param size the size in bytes to fill, must be a multiple of the pattern length
+ * @param offset the offset in bytes into the buffer, must be a multiple of the pattern length
+ * @return an event indicating when this operation is finished
+ */
+ public abstract Event fillAsync(CommandQueue queue, ByteBuffer pattern, long size, long offset);
+
+ /**
+ * Result of an async mapping operation, contains the event and the target byte buffer.
+ * This is a work-around since no generic pair-structure is avaiable.
+ *
+ * @author shaman
+ */
+ public static class AsyncMapping {
+
+ public final Event event;
+ public final ByteBuffer buffer;
+
+ public AsyncMapping(Event event, ByteBuffer buffer) {
+ super();
+ this.event = event;
+ this.buffer = buffer;
+ }
+
+ /**
+ * @return the event object indicating when the data in the mapped buffer
+ * is available
+ */
+ public Event getEvent() {
+ return event;
+ }
+
+ /**
+ * @return the mapped buffer, only valid when the event object signals completion
+ */
+ public ByteBuffer getBuffer() {
+ return buffer;
+ }
+ }
+
+ /**
+ * Copies this buffer to the specified image.
+ * Note that no format conversion is done.
+ *
+ * For detailed description of the origin and region paramenter, see the
+ * documentation of the {@link Image} class.
+ *
+ * @param queue the command queue
+ * @param dest the target image
+ * @param srcOffset the offset in bytes into this buffer
+ * @param destOrigin the origin of the copied area
+ * @param destRegion the size of the copied area
+ * @return the event object
+ */
+ public abstract Event copyToImageAsync(CommandQueue queue, Image dest, long srcOffset, long[] destOrigin, long[] destRegion);
+
+ /**
+ * Aquires this buffer object for using. Only call this method if this buffer
+ * represents a shared object from OpenGL, created with e.g.
+ * {@link Context#bindVertexBuffer(com.jme3.scene.VertexBuffer, com.jme3.opencl.MemoryAccess) }.
+ * This method must be called before the buffer is used. After the work is
+ * done, the buffer must be released by calling
+ * {@link #releaseBufferForSharingAsync(com.jme3.opencl.CommandQueue) }
+ * so that OpenGL can use the VertexBuffer again.
+ * @param queue the command queue
+ * @return the event object
+ */
+ public abstract Event acquireBufferForSharingAsync(CommandQueue queue);
+
+ /**
+ * Aquires this buffer object for using. Only call this method if this buffer
+ * represents a shared object from OpenGL, created with e.g.
+ * {@link Context#bindVertexBuffer(com.jme3.scene.VertexBuffer, com.jme3.opencl.MemoryAccess) }.
+ * This method must be called before the buffer is used. After the work is
+ * done, the buffer must be released by calling
+ * {@link #releaseBufferForSharingAsync(com.jme3.opencl.CommandQueue) }
+ * so that OpenGL can use the VertexBuffer again.
+ *
+ * The generated event object is directly released.
+ * This brings a performance improvement when the resource is e.g. directly
+ * used by a kernel afterwards on the same queue (this implicitly waits for
+ * this action). If you need the event, use
+ * {@link #acquireBufferForSharingAsync(com.jme3.opencl.CommandQueue) } instead.
+ *
+ * @param queue the command queue
+ */
+ public void acquireBufferForSharingNoEvent(CommandQueue queue) {
+ //default implementation, overwrite for better performance
+ acquireBufferForSharingAsync(queue).release();
+ }
+
+ /**
+ * Releases a shared buffer object.
+ * Call this method after the buffer object was acquired by
+ * {@link #acquireBufferForSharingAsync(com.jme3.opencl.CommandQueue) }
+ * to hand the control back to OpenGL.
+ * @param queue the command queue
+ * @return the event object
+ */
+ public abstract Event releaseBufferForSharingAsync(CommandQueue queue);
+
+ /**
+ * Releases a shared buffer object.
+ * Call this method after the buffer object was acquired by
+ * {@link #acquireBufferForSharingAsync(com.jme3.opencl.CommandQueue) }
+ * to hand the control back to OpenGL.
+ * The generated event object is directly released, resulting in
+ * performance improvements.
+ * @param queue the command queue
+ */
+ public void releaseBufferForSharingNoEvent(CommandQueue queue) {
+ //default implementation, overwrite for better performance
+ releaseBufferForSharingAsync(queue).release();
+ }
+
+}
diff --git a/jme3-core/src/main/java/com/jme3/opencl/CommandQueue.java b/jme3-core/src/main/java/com/jme3/opencl/CommandQueue.java
new file mode 100644
index 000000000..6d54237f9
--- /dev/null
+++ b/jme3-core/src/main/java/com/jme3/opencl/CommandQueue.java
@@ -0,0 +1,68 @@
+/*
+ * Copyright (c) 2009-2016 jMonkeyEngine
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are
+ * met:
+ *
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ *
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ *
+ * * Neither the name of 'jMonkeyEngine' nor the names of its contributors
+ * may be used to endorse or promote products derived from this software
+ * without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+ * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+ * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
+ * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
+ * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+package com.jme3.opencl;
+
+/**
+ * Wrapper for an OpenCL command queue.
+ * The command queue serializes every GPU function call: By passing the same
+ * queue to OpenCL function (buffer, image operations, kernel calls), it is
+ * ensured that they are executed in the order in which they are passed.
+ *
+ * Each command queue is associtated with exactly one device: that device
+ * is specified on creation ({@link Context#createQueue(com.jme3.opencl.Device) })
+ * and all commands are sent to this device.
+ * @author shaman
+ */
+public abstract class CommandQueue extends AbstractOpenCLObject {
+
+ protected CommandQueue(ObjectReleaser releaser) {
+ super(releaser);
+ }
+
+ /**
+ * Issues all previously queued OpenCL commands in command_queue to the
+ * device associated with command queue. Flush only guarantees that all
+ * queued commands to command_queue will eventually be submitted to the
+ * appropriate device. There is no guarantee that they will be complete
+ * after flush returns.
+ */
+ public abstract void flush();
+
+ /**
+ * Blocks until all previously queued OpenCL commands in command queue are
+ * issued to the associated device and have completed. Finish does not
+ * return until all previously queued commands in command queue have been
+ * processed and completed. Finish is also a synchronization point.
+ */
+ public abstract void finish();
+
+}
diff --git a/jme3-core/src/main/java/com/jme3/opencl/Context.java b/jme3-core/src/main/java/com/jme3/opencl/Context.java
new file mode 100644
index 000000000..fed3cd2a3
--- /dev/null
+++ b/jme3-core/src/main/java/com/jme3/opencl/Context.java
@@ -0,0 +1,439 @@
+/*
+ * Copyright (c) 2009-2016 jMonkeyEngine
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are
+ * met:
+ *
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ *
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ *
+ * * Neither the name of 'jMonkeyEngine' nor the names of its contributors
+ * may be used to endorse or promote products derived from this software
+ * without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+ * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+ * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
+ * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
+ * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+package com.jme3.opencl;
+
+import com.jme3.asset.AssetInfo;
+import com.jme3.asset.AssetKey;
+import com.jme3.asset.AssetManager;
+import com.jme3.asset.AssetNotFoundException;
+import com.jme3.opencl.Image.ImageDescriptor;
+import com.jme3.opencl.Image.ImageFormat;
+import com.jme3.opencl.Image.ImageType;
+import com.jme3.scene.VertexBuffer;
+import com.jme3.texture.FrameBuffer;
+import com.jme3.texture.Texture;
+import java.io.BufferedReader;
+import java.io.IOException;
+import java.io.InputStreamReader;
+import java.io.StringReader;
+import java.nio.ByteBuffer;
+import java.util.Arrays;
+import java.util.List;
+import java.util.logging.Level;
+import java.util.logging.Logger;
+
+/**
+ * The central OpenCL context. Every action starts from here.
+ * The context can be obtained by {@link com.jme3.system.JmeContext#getOpenCLContext() }.
+ *
+ * The context is used to:
+ *
+ *
Query the available devices
+ *
Create a command queue
+ *
Create buffers and images
+ *
Created buffers and images shared with OpenGL vertex buffers, textures and renderbuffers
+ *
Create program objects from source code and source files
+ *
+ * @author shaman
+ */
+public abstract class Context extends AbstractOpenCLObject {
+ private static final Logger LOG = Logger.getLogger(Context.class.getName());
+
+ protected Context(ObjectReleaser releaser) {
+ super(releaser);
+ }
+
+ /**
+ * Returns all available devices for this context.
+ * These devices all belong to the same {@link Platform}.
+ * They are used to create a command queue sending commands to a particular
+ * device, see {@link #createQueue(com.jme3.opencl.Device) }.
+ * Also, device capabilities, like the supported OpenCL version, extensions,
+ * memory size and so on, are queried over the Device instances.
+ *
+ * The available devices were specified by a {@link PlatformChooser}.
+ * @return
+ */
+ public abstract List getDevices();
+
+ /**
+ * Alternative version of {@link #createQueue(com.jme3.opencl.Device) },
+ * just uses the first device returned by {@link #getDevices() }.
+ * @return the command queue
+ */
+ public CommandQueue createQueue() {
+ return createQueue(getDevices().get(0));
+ }
+ /**
+ * Creates a command queue sending commands to the specified device.
+ * The device must be an entry of {@link #getDevices() }.
+ * @param device the target device
+ * @return the command queue
+ */
+ public abstract CommandQueue createQueue(Device device);
+
+ /**
+ * Allocates a new buffer of the specific size and access type on the device.
+ * @param size the size of the buffer in bytes
+ * @param access the allowed access of this buffer from kernel code
+ * @return the new buffer
+ */
+ public abstract Buffer createBuffer(long size, MemoryAccess access);
+ /**
+ * Alternative version of {@link #createBuffer(long, com.jme3.opencl.MemoryAccess) },
+ * creates a buffer with read and write access.
+ * @param size the size of the buffer in bytes
+ * @return the new buffer
+ */
+ public Buffer createBuffer(long size) {
+ return createBuffer(size, MemoryAccess.READ_WRITE);
+ }
+
+ /**
+ * Creates a new buffer wrapping the specific host memory. This host memory
+ * specified by a ByteBuffer can then be used directly by kernel code,
+ * although the access might be slower than with native buffers
+ * created by {@link #createBuffer(long, com.jme3.opencl.MemoryAccess) }.
+ * @param data the host buffer to use
+ * @param access the allowed access of this buffer from kernel code
+ * @return the new buffer
+ */
+ public abstract Buffer createBufferFromHost(ByteBuffer data, MemoryAccess access);
+ /**
+ * Alternative version of {@link #createBufferFromHost(java.nio.ByteBuffer, com.jme3.opencl.MemoryAccess) },
+ * creates a buffer with read and write access.
+ * @param data the host buffer to use
+ * @return the new buffer
+ */
+ public Buffer createBufferFromHost(ByteBuffer data) {
+ return createBufferFromHost(data, MemoryAccess.READ_WRITE);
+ }
+
+ /**
+ * Creates a new 1D, 2D, 3D image.
+ * {@code ImageFormat} specifies the element type and order, like RGBA of floats.
+ * {@code ImageDescriptor} specifies the dimension of the image.
+ * Furthermore, a ByteBuffer can be specified in the ImageDescriptor together
+ * with row and slice pitches. This buffer is then used to store the image.
+ * If no ByteBuffer is specified, a new buffer is allocated (this is the
+ * normal behaviour).
+ * @param access the allowed access of this image from kernel code
+ * @param format the image format
+ * @param descr the image descriptor
+ * @return the new image object
+ */
+ public abstract Image createImage(MemoryAccess access, ImageFormat format, ImageDescriptor descr);
+ //TODO: add simplified methods for 1D, 2D, 3D textures
+
+ /**
+ * Queries all supported image formats for a specified memory access and
+ * image type.
+ *
+ * Note that the returned array may contain {@code ImageFormat} objects
+ * where {@code ImageChannelType} or {@code ImageChannelOrder} are {@code null}
+ * (or both). This is the case when the device supports new formats that
+ * are not included in this wrapper yet.
+ * @param access the memory access type
+ * @param type the image type (1D, 2D, 3D, ...)
+ * @return an array of all supported image formats
+ */
+ public abstract ImageFormat[] querySupportedFormats(MemoryAccess access, ImageType type);
+
+ //Interop
+ /**
+ * Creates a shared buffer from a VertexBuffer.
+ * The returned buffer and the vertex buffer operate on the same memory,
+ * changes in one view are visible in the other view.
+ * This can be used to modify meshes directly from OpenCL (e.g. for particle systems).
+ *
+ * Note: The vertex buffer must already been uploaded to the GPU,
+ * i.e. it must be used at least once for drawing.
+ *
+ * Before the returned buffer can be used, it must be acquried explicitly
+ * by {@link Buffer#acquireBufferForSharingAsync(com.jme3.opencl.CommandQueue) }
+ * and after modifying it, released by {@link Buffer#releaseBufferForSharingAsync(com.jme3.opencl.CommandQueue) }.
+ * This is needed so that OpenGL and OpenCL operations do not interfer with each other.
+ * @param vb the vertex buffer to share
+ * @param access the memory access for the kernel
+ * @return the new buffer
+ */
+ public abstract Buffer bindVertexBuffer(VertexBuffer vb, MemoryAccess access);
+
+ /**
+ * Creates a shared image object from a jME3-image.
+ * The returned image shares the same memory with the jME3-image, changes
+ * in one view are visible in the other view.
+ * This can be used to modify textures and images directly from OpenCL
+ * (e.g. for post processing effects and other texture effects).
+ *
+ * Note: The image must already been uploaded to the GPU,
+ * i.e. it must be used at least once for drawing.
+ *
+ * Before the returned image can be used, it must be acquried explicitly
+ * by {@link Image#acquireImageForSharingAsync(com.jme3.opencl.CommandQueue) }
+ * and after modifying it, released by {@link Image#releaseImageForSharingAsync(com.jme3.opencl.CommandQueue) }
+ * This is needed so that OpenGL and OpenCL operations do not interfer with each other.
+ *
+ * @param image the jME3 image object
+ * @param textureType the texture type (1D, 2D, 3D), since this is not stored in the image
+ * @param miplevel the mipmap level that should be shared
+ * @param access the allowed memory access for kernels
+ * @return the OpenCL image
+ */
+ public abstract Image bindImage(com.jme3.texture.Image image, Texture.Type textureType, int miplevel, MemoryAccess access);
+ /**
+ * Creates a shared image object from a jME3 texture.
+ * The returned image shares the same memory with the jME3 texture, changes
+ * in one view are visible in the other view.
+ * This can be used to modify textures and images directly from OpenCL
+ * (e.g. for post processing effects and other texture effects).
+ *
+ * Note: The image must already been uploaded to the GPU,
+ * i.e. it must be used at least once for drawing.
+ *
+ * Before the returned image can be used, it must be acquried explicitly
+ * by {@link Image#acquireImageForSharingAsync(com.jme3.opencl.CommandQueue) }
+ * and after modifying it, released by {@link Image#releaseImageForSharingAsync(com.jme3.opencl.CommandQueue) }
+ * This is needed so that OpenGL and OpenCL operations do not interfer with each other.
+ *
+ * This method is equivalent to calling
+ * {@code bindImage(texture.getImage(), texture.getType(), miplevel, access)}.
+ *
+ * @param texture the jME3 texture
+ * @param miplevel the mipmap level that should be shared
+ * @param access the allowed memory access for kernels
+ * @return the OpenCL image
+ */
+ public Image bindImage(Texture texture, int miplevel, MemoryAccess access) {
+ return bindImage(texture.getImage(), texture.getType(), miplevel, access);
+ }
+ /**
+ * Alternative version to {@link #bindImage(com.jme3.texture.Texture, int, com.jme3.opencl.MemoryAccess) },
+ * uses {@code miplevel=0}.
+ * @param texture the jME3 texture
+ * @param access the allowed memory access for kernels
+ * @return the OpenCL image
+ */
+ public Image bindImage(Texture texture, MemoryAccess access) {
+ return bindImage(texture, 0, access);
+ }
+ /**
+ * Creates a shared image object from a jME3 render buffer.
+ * The returned image shares the same memory with the jME3 render buffer, changes
+ * in one view are visible in the other view.
+ *
+ * This can be used as an alternative to post processing effects
+ * (e.g. reduce sum operations, needed e.g. for tone mapping).
+ *
+ * Note: The renderbuffer must already been uploaded to the GPU,
+ * i.e. it must be used at least once for drawing.
+ *
+ * Before the returned image can be used, it must be acquried explicitly
+ * by {@link Image#acquireImageForSharingAsync(com.jme3.opencl.CommandQueue) }
+ * and after modifying it, released by {@link Image#releaseImageForSharingAsync(com.jme3.opencl.CommandQueue) }
+ * This is needed so that OpenGL and OpenCL operations do not interfer with each other.
+ *
+ * @param buffer
+ * @param access
+ * @return
+ */
+ public Image bindRenderBuffer(FrameBuffer.RenderBuffer buffer, MemoryAccess access) {
+ if (buffer.getTexture() == null) {
+ return bindPureRenderBuffer(buffer, access);
+ } else {
+ return bindImage(buffer.getTexture(), access);
+ }
+ }
+ protected abstract Image bindPureRenderBuffer(FrameBuffer.RenderBuffer buffer, MemoryAccess access);
+
+ /**
+ * Creates a program object from the provided source code.
+ * The program still needs to be compiled using {@link Program#build() }.
+ *
+ * @param sourceCode the source code
+ * @return the program object
+ */
+ public abstract Program createProgramFromSourceCode(String sourceCode);
+
+ /**
+ * Resolves dependencies (using {@code #include } in the source code)
+ * and delegates the combined source code to
+ * {@link #createProgramFromSourceCode(java.lang.String) }.
+ * Important: only absolute paths are allowed.
+ * @param sourceCode the original source code
+ * @param assetManager the asset manager to load the files
+ * @return the created program object
+ * @throws AssetNotFoundException if a dependency could not be loaded
+ */
+ public Program createProgramFromSourceCodeWithDependencies(String sourceCode, AssetManager assetManager) {
+ StringBuilder builder = new StringBuilder(sourceCode.length());
+ BufferedReader reader = new BufferedReader(new StringReader(sourceCode));
+ try {
+ buildSourcesRec(reader, builder, assetManager);
+ } catch (IOException ex) {
+ throw new AssetNotFoundException("Unable to read a dependency file", ex);
+ }
+ return createProgramFromSourceCode(builder.toString());
+ }
+ private void buildSourcesRec(BufferedReader reader, StringBuilder builder, AssetManager assetManager) throws IOException {
+ String ln;
+ while ((ln = reader.readLine()) != null) {
+ if (ln.trim().startsWith("#import ")) {
+ ln = ln.trim().substring(8).trim();
+ if (ln.startsWith("\"")) {
+ ln = ln.substring(1);
+ }
+ if (ln.endsWith("\"")) {
+ ln = ln.substring(0, ln.length()-1);
+ }
+ AssetInfo info = assetManager.locateAsset(new AssetKey(ln));
+ if (info == null) {
+ throw new AssetNotFoundException("Unable to load source file \""+ln+"\"");
+ }
+ try (BufferedReader r = new BufferedReader(new InputStreamReader(info.openStream()))) {
+ builder.append("//-- begin import ").append(ln).append(" --\n");
+ buildSourcesRec(r, builder, assetManager);
+ builder.append("//-- end import ").append(ln).append(" --\n");
+ }
+ } else {
+ builder.append(ln).append('\n');
+ }
+ }
+ }
+
+ /**
+ * Creates a program object from the provided source code and files.
+ * The source code is made up from the specified include string first,
+ * then all files specified by the resource array (array of asset paths)
+ * are loaded by the provided asset manager and appended to the source code.
+ *
+ * The typical use case is:
+ *
+ *
The include string contains some compiler constants like the grid size
+ *
Some common OpenCL files used as libraries (Convention: file names end with {@code .clh}
+ *
One main OpenCL file containing the actual kernels (Convention: file name ends with {@code .cl})
+ *
+ *
+ * After the files were combined, additional include statements are resolved
+ * by {@link #createProgramFromSourceCodeWithDependencies(java.lang.String, com.jme3.asset.AssetManager) }.
+ *
+ * @param assetManager the asset manager used to load the files
+ * @param include an additional include string
+ * @param resources an array of asset paths pointing to OpenCL source files
+ * @return the new program objects
+ * @throws AssetNotFoundException if a file could not be loaded
+ */
+ public Program createProgramFromSourceFilesWithInclude(AssetManager assetManager, String include, String... resources) {
+ return createProgramFromSourceFilesWithInclude(assetManager, include, Arrays.asList(resources));
+ }
+
+ /**
+ * Creates a program object from the provided source code and files.
+ * The source code is made up from the specified include string first,
+ * then all files specified by the resource array (array of asset paths)
+ * are loaded by the provided asset manager and appended to the source code.
+ *
+ * The typical use case is:
+ *
+ *
The include string contains some compiler constants like the grid size
+ *
Some common OpenCL files used as libraries (Convention: file names end with {@code .clh}
+ *
One main OpenCL file containing the actual kernels (Convention: file name ends with {@code .cl})
+ *
+ *
+ * After the files were combined, additional include statements are resolved
+ * by {@link #createProgramFromSourceCodeWithDependencies(java.lang.String, com.jme3.asset.AssetManager) }.
+ *
+ * @param assetManager the asset manager used to load the files
+ * @param include an additional include string
+ * @param resources an array of asset paths pointing to OpenCL source files
+ * @return the new program objects
+ * @throws AssetNotFoundException if a file could not be loaded
+ */
+ public Program createProgramFromSourceFilesWithInclude(AssetManager assetManager, String include, List resources) {
+ StringBuilder str = new StringBuilder();
+ str.append(include);
+ for (String res : resources) {
+ AssetInfo info = assetManager.locateAsset(new AssetKey(res));
+ if (info == null) {
+ throw new AssetNotFoundException("Unable to load source file \""+res+"\"");
+ }
+ try (BufferedReader reader = new BufferedReader(new InputStreamReader(info.openStream()))) {
+ while (true) {
+ String line = reader.readLine();
+ if (line == null) {
+ break;
+ }
+ str.append(line).append('\n');
+ }
+ } catch (IOException ex) {
+ LOG.log(Level.WARNING, "unable to load source file '"+res+"'", ex);
+ }
+ }
+ return createProgramFromSourceCodeWithDependencies(str.toString(), assetManager);
+ }
+
+ /**
+ * Alternative version of {@link #createProgramFromSourceFilesWithInclude(com.jme3.asset.AssetManager, java.lang.String, java.lang.String...) }
+ * with an empty include string
+ * @throws AssetNotFoundException if a file could not be loaded
+ */
+ public Program createProgramFromSourceFiles(AssetManager assetManager, String... resources) {
+ return createProgramFromSourceFilesWithInclude(assetManager, "", resources);
+ }
+
+ /**
+ * Alternative version of {@link #createProgramFromSourceFilesWithInclude(com.jme3.asset.AssetManager, java.lang.String, java.util.List) }
+ * with an empty include string
+ * @throws AssetNotFoundException if a file could not be loaded
+ */
+ public Program createProgramFromSourceFiles(AssetManager assetManager, List resources) {
+ return createProgramFromSourceFilesWithInclude(assetManager, "", resources);
+ }
+
+ /**
+ * Creates a program from the specified binaries.
+ * The binaries are created by {@link Program#getBinary(com.jme3.opencl.Device) }.
+ * The returned program still needs to be build using
+ * {@link Program#build(java.lang.String, com.jme3.opencl.Device...) }.
+ * Important:The device passed to {@code Program.getBinary(..)},
+ * this method and {@code Program#build(..)} must be the same.
+ *
+ * The binaries are used to build a program cache across multiple launches
+ * of the application. The programs build mach faster from binaries than
+ * from sources.
+ *
+ * @param binaries the binaries
+ * @param device the device to use
+ * @return the new program
+ */
+ public abstract Program createProgramFromBinary(ByteBuffer binaries, Device device);
+}
diff --git a/jme3-core/src/main/java/com/jme3/opencl/DefaultPlatformChooser.java b/jme3-core/src/main/java/com/jme3/opencl/DefaultPlatformChooser.java
new file mode 100644
index 000000000..8f8a53505
--- /dev/null
+++ b/jme3-core/src/main/java/com/jme3/opencl/DefaultPlatformChooser.java
@@ -0,0 +1,91 @@
+/*
+ * Copyright (c) 2009-2016 jMonkeyEngine
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are
+ * met:
+ *
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ *
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ *
+ * * Neither the name of 'jMonkeyEngine' nor the names of its contributors
+ * may be used to endorse or promote products derived from this software
+ * without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+ * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+ * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
+ * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
+ * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+package com.jme3.opencl;
+
+import java.util.ArrayList;
+import java.util.List;
+import java.util.logging.Logger;
+
+/**
+ * A default implementation of {@link PlatformChooser}.
+ * It favors GPU devices with OpenGL sharing, then any devices with OpenGL sharing,
+ * then any possible device.
+ * @author shaman
+ */
+public class DefaultPlatformChooser implements PlatformChooser {
+ private static final Logger LOG = Logger.getLogger(DefaultPlatformChooser.class.getName());
+
+ @Override
+ public List chooseDevices(List platforms) {
+ ArrayList result = new ArrayList();
+ for (Platform p : platforms) {
+ if (!p.hasOpenGLInterop()) {
+ continue; //must support interop
+ }
+ for (Device d : p.getDevices()) {
+ if (d.hasOpenGLInterop() && d.getDeviceType()==Device.DeviceType.GPU) {
+ result.add(d); //GPU prefered
+ }
+ }
+ if (!result.isEmpty()) {
+ return result;
+ }
+ }
+ //no GPU devices found, try all
+ for (Platform p : platforms) {
+ if (!p.hasOpenGLInterop()) {
+ continue; //must support interop
+ }
+ for (Device d : p.getDevices()) {
+ if (d.hasOpenGLInterop()) {
+ result.add(d); //just interop needed
+ }
+ }
+ if (!result.isEmpty()) {
+ return result;
+ }
+ }
+ //still no one found, try without interop
+ LOG.warning("No device with OpenCL-OpenGL-interop found, try without");
+ for (Platform p : platforms) {
+ for (Device d : p.getDevices()) {
+ result.add(d);
+ }
+ if (!result.isEmpty()) {
+ return result;
+ }
+ }
+ //no devices available at all!
+ return result; //result is empty
+ }
+
+}
diff --git a/jme3-core/src/main/java/com/jme3/opencl/Device.java b/jme3-core/src/main/java/com/jme3/opencl/Device.java
new file mode 100644
index 000000000..d34854c0c
--- /dev/null
+++ b/jme3-core/src/main/java/com/jme3/opencl/Device.java
@@ -0,0 +1,311 @@
+/*
+ * Copyright (c) 2009-2016 jMonkeyEngine
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are
+ * met:
+ *
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ *
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ *
+ * * Neither the name of 'jMonkeyEngine' nor the names of its contributors
+ * may be used to endorse or promote products derived from this software
+ * without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+ * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+ * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
+ * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
+ * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+package com.jme3.opencl;
+
+import java.util.Collection;
+
+/**
+ * Represents a hardware device actually running the OpenCL kernels.
+ * A {@link Context} can be accociated with multiple {@code Devices}
+ * that all belong to the same {@link Platform}.
+ * For execution, a single device must be chosen and passed to a command
+ * queue ({@link Context#createQueue(com.jme3.opencl.Device) }).
+ *
+ * This class is used to query the capabilities of the underlying device.
+ *
+ * @author shaman
+ */
+public interface Device {
+
+ /**
+ * @return the platform accociated with this device
+ */
+ Platform getPlatform();
+
+ /**
+ * The device type
+ */
+ public static enum DeviceType {
+ DEFAULT,
+ CPU,
+ GPU,
+ ACCELEARTOR,
+ ALL
+ }
+ /**
+ * @return queries the device type
+ */
+ DeviceType getDeviceType();
+ /**
+ * @return the vendor id
+ */
+ int getVendorId();
+ /**
+ * checks if this device is available at all, must always be tested
+ * @return checks if this device is available at all, must always be tested
+ */
+ boolean isAvailable();
+
+ /**
+ * @return if this device has a compiler for kernel code
+ */
+ boolean hasCompiler();
+ /**
+ * @return supports double precision floats (64 bit)
+ */
+ boolean hasDouble();
+ /**
+ * @return supports half precision floats (16 bit)
+ */
+ boolean hasHalfFloat();
+ /**
+ * @return supports error correction for every access to global or constant memory
+ */
+ boolean hasErrorCorrectingMemory();
+ /**
+ * @return supports unified virtual memory (OpenCL 2.0)
+ */
+ boolean hasUnifiedMemory();
+ /**
+ * @return supports images
+ */
+ boolean hasImageSupport();
+ /**
+ * @return supports writes to 3d images (this is an extension)
+ */
+ boolean hasWritableImage3D();
+ /**
+ * @return supports sharing with OpenGL
+ */
+ boolean hasOpenGLInterop();
+ /**
+ * Explicetly tests for the availability of the specified extension
+ * @param extension the name of the extension
+ * @return {@code true} iff this extension is supported
+ */
+ boolean hasExtension(String extension);
+ /**
+ * Lists all available extensions
+ * @return all available extensions
+ */
+ Collection getExtensions();
+
+ /**
+ * Returns the number of parallel compute units on
+ * the OpenCL device. A work-group
+ * executes on a single compute unit. The
+ * minimum value is 1.
+ * @return the number of parallel compute units
+ * @see #getMaximumWorkItemDimensions()
+ * @see #getMaximumWorkItemSizes()
+ */
+ int getComputeUnits();
+ /**
+ * @return maximum clock frequency of the device in MHz
+ */
+ int getClockFrequency();
+ /**
+ * Returns the default compute device address space
+ * size specified as an unsigned integer value
+ * in bits. Currently supported values are 32
+ * or 64 bits.
+ * @return the size of an adress
+ */
+ int getAddressBits();
+ /**
+ * @return {@code true} if this device is little endian
+ */
+ boolean isLittleEndian();
+
+ /**
+ * The maximum dimension that specify the local and global work item ids.
+ * You can always assume to be this at least 3.
+ * Therefore, the ids are always three integers x,y,z.
+ * @return the maximum dimension of work item ids
+ */
+ long getMaximumWorkItemDimensions();
+ /**
+ * Maximum number of work-items that can be specified in each dimension of the
+ * work-group to {@link Kernel#Run2(com.jme3.opencl.CommandQueue, com.jme3.opencl.WorkSize, com.jme3.opencl.WorkSize, java.lang.Object...) }.
+ * The array has a length of at least 3.
+ * @return the maximum size of the work group in each dimension
+ */
+ long[] getMaximumWorkItemSizes();
+ /**
+ * Maximum number of work-items in a
+ * work-group executing a kernel on a single
+ * compute unit, using the data parallel
+ * execution model.
+ * @return maximum number of work-items in a work-group
+ */
+ long getMaxiumWorkItemsPerGroup();
+
+ /**
+ * @return the maximum number of samples that can be used in a kernel
+ */
+ int getMaximumSamplers();
+ /**
+ * @return the maximum number of images that can be used for reading in a kernel
+ */
+ int getMaximumReadImages();
+ /**
+ * @return the maximum number of images that can be used for writing in a kernel
+ */
+ int getMaximumWriteImages();
+ /**
+ * Queries the maximal size of a 2D image
+ * @return an array of length 2 with the maximal size of a 2D image
+ */
+ long[] getMaximumImage2DSize();
+ /**
+ * Queries the maximal size of a 3D image
+ * @return an array of length 3 with the maximal size of a 3D image
+ */
+ long[] getMaximumImage3DSize();
+
+ /**
+ * @return the maximal size of a memory object (buffer and image) in bytes
+ */
+ long getMaximumAllocationSize();
+ /**
+ * @return the total available global memory in bytes
+ */
+ long getGlobalMemorySize();
+ /**
+ * @return the total available local memory in bytes
+ */
+ long getLocalMemorySize();
+ /**
+ * Returns the maximal size of a constant buffer.
+ *
+ * Constant buffers are normal buffer objects, but passed to the kernel
+ * with the special declaration {@code __constant BUFFER_TYPE* BUFFER_NAME}.
+ * Because they have a special caching, their size is usually very limited.
+ *
+ * @return the maximal size of a constant buffer
+ */
+ long getMaximumConstantBufferSize();
+ /**
+ * @return the maximal number of constant buffer arguments in a kernel call
+ */
+ int getMaximumConstantArguments();
+
+ //TODO: cache, prefered sizes properties
+ /**
+ * OpenCL profile string. Returns the profile name supported by the device.
+ * The profile name returned can be one of the following strings:
+ * FULL_PROFILE – if the device supports the OpenCL specification
+ * (functionality defined as part of the core specification and does not
+ * require any extensions to be supported).
+ * EMBEDDED_PROFILE - if the device supports the OpenCL embedded profile.
+ *
+ * @return the profile string
+ */
+ String getProfile();
+ /**
+ * OpenCL version string. Returns the OpenCL version supported by the
+ * device. This version string has the following format: OpenCL space
+ * major_version.minor_version space vendor-specific information.
+ *
+ * E.g. OpenCL 1.1, OpenCL 1.2, OpenCL 2.0
+ *
+ * @return the version string
+ */
+ String getVersion();
+ /**
+ * Extracts the major version from the version string
+ * @return the major version
+ * @see #getVersion()
+ */
+ int getVersionMajor();
+ /**
+ * Extracts the minor version from the version string
+ * @return the minor version
+ * @see #getVersion() }
+ */
+ int getVersionMinor();
+
+ /**
+ * OpenCL C version string. Returns the highest OpenCL C version supported
+ * by the compiler for this device that is not of type
+ * CL_DEVICE_TYPE_CUSTOM. This version string has the following format:
+ * OpenCL space C space major_version.minor_version space vendor-specific
+ * information.
+ * The major_version.minor_version value returned must be 1.2 if
+ * CL_DEVICE_VERSION is OpenCL 1.2. The major_version.minor_version value
+ * returned must be 1.1 if CL_DEVICE_VERSION is OpenCL 1.1. The
+ * major_version.minor_version value returned can be 1.0 or 1.1 if
+ * CL_DEVICE_VERSION is OpenCL 1.0.
+ *
+ * @return the compiler version
+ */
+ String getCompilerVersion();
+ /**
+ * Extracts the major version from the compiler version
+ * @return the major compiler version
+ * @see #getCompilerVersion()
+ */
+ int getCompilerVersionMajor();
+ /**
+ * Extracts the minor version from the compiler version
+ * @return the minor compiler version
+ * @see #getCompilerVersion()
+ */
+ int getCompilerVersionMinor();
+ /**
+ * @return the OpenCL software driver version string in the form
+ * major_number.minor_number
+ */
+ String getDriverVersion();
+ /**
+ * Extracts the major version from the driver version
+ * @return the major driver version
+ * @see #getDriverVersion()
+ */
+ int getDriverVersionMajor();
+ /**
+ * Extracts the minor version from the driver version
+ * @return the minor driver version
+ * @see #getDriverVersion()
+ */
+ int getDriverVersionMinor();
+
+ /**
+ * @return the device name
+ */
+ String getName();
+ /**
+ * @return the vendor
+ */
+ String getVendor();
+
+}
diff --git a/jme3-core/src/main/java/com/jme3/opencl/Event.java b/jme3-core/src/main/java/com/jme3/opencl/Event.java
new file mode 100644
index 000000000..44ea3da16
--- /dev/null
+++ b/jme3-core/src/main/java/com/jme3/opencl/Event.java
@@ -0,0 +1,59 @@
+/*
+ * Copyright (c) 2009-2016 jMonkeyEngine
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are
+ * met:
+ *
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ *
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ *
+ * * Neither the name of 'jMonkeyEngine' nor the names of its contributors
+ * may be used to endorse or promote products derived from this software
+ * without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+ * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+ * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
+ * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
+ * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+package com.jme3.opencl;
+
+/**
+ * Wrapper for an OpenCL Event object.
+ * Events are returned from kernel launches and all asynchronous operations.
+ * They allow to test if the action has completed and to block until the operation
+ * is done.
+ * @author shaman
+ */
+public abstract class Event extends AbstractOpenCLObject {
+
+ protected Event(ObjectReleaser releaser) {
+ super(releaser);
+ }
+
+ /**
+ * Waits until the action has finished (blocking).
+ * This automatically releases the event.
+ */
+ public abstract void waitForFinished();
+
+ /**
+ * Tests if the action is completed.
+ * If the action is completed, the event is released.
+ * @return {@code true} if the action is completed
+ */
+ public abstract boolean isCompleted();
+}
diff --git a/jme3-core/src/main/java/com/jme3/opencl/Image.java b/jme3-core/src/main/java/com/jme3/opencl/Image.java
new file mode 100644
index 000000000..f9d9d9f28
--- /dev/null
+++ b/jme3-core/src/main/java/com/jme3/opencl/Image.java
@@ -0,0 +1,537 @@
+/*
+ * Copyright (c) 2009-2016 jMonkeyEngine
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are
+ * met:
+ *
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ *
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ *
+ * * Neither the name of 'jMonkeyEngine' nor the names of its contributors
+ * may be used to endorse or promote products derived from this software
+ * without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+ * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+ * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
+ * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
+ * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+package com.jme3.opencl;
+
+import com.jme3.math.ColorRGBA;
+import java.nio.ByteBuffer;
+import java.util.Objects;
+
+/**
+ * Wrapper for an OpenCL image.
+ *
+ * An image object is similar to a {@link Buffer}, but with a specific element
+ * format and buffer structure.
+ *
+ * The image is specified by the {@link ImageDescriptor}, specifying
+ * the extend and dimension of the image, and {@link ImageFormat}, specifying
+ * the type of each pixel.
+ *
+ * An image is created from scratch using
+ * {@link Context#createImage(com.jme3.opencl.MemoryAccess, com.jme3.opencl.Image.ImageFormat, com.jme3.opencl.Image.ImageDescriptor) }
+ * or from OpenGL by
+ * {@link Context#bindImage(com.jme3.texture.Image, com.jme3.texture.Texture.Type, int, com.jme3.opencl.MemoryAccess) }
+ * (and alternative versions).
+ *
+ *
+ * Most methods take long arrays as input: {@code long[] origin} and {@code long[] region}.
+ * Both are arrays of length 3.
+ *
+ * origin defines the (x, y, z) offset in pixels in the 1D, 2D or 3D
+ * image, the (x, y) offset and the image index in the 2D image array or the (x)
+ * offset and the image index in the 1D image array. If image is a 2D image
+ * object, origin[2] must be 0. If image is a 1D image or 1D image buffer
+ * object, origin[1] and origin[2] must be 0. If image is a 1D image array
+ * object, origin[2] must be 0. If image is a 1D image array object, origin[1]
+ * describes the image index in the 1D image array. If image is a 2D image array
+ * object, origin[2] describes the image index in the 2D image array.
+ *
+ * region defines the (width, height, depth) in pixels of the 1D, 2D or
+ * 3D rectangle, the (width, height) in pixels of the 2D rectangle and the
+ * number of images of a 2D image array or the (width) in pixels of the 1D
+ * rectangle and the number of images of a 1D image array. If image is a 2D
+ * image object, region[2] must be 1. If image is a 1D image or 1D image buffer
+ * object, region[1] and region[2] must be 1. If image is a 1D image array
+ * object, region[2] must be 1. The values in region cannot be 0.
+ *
+ * @author shaman
+ */
+public abstract class Image extends AbstractOpenCLObject {
+
+ /**
+ * {@code ImageChannelType} describes the size of the channel data type.
+ */
+ public static enum ImageChannelType {
+ SNORM_INT8,
+ SNORM_INT16,
+ UNORM_INT8,
+ UNORM_INT16,
+ UNORM_SHORT_565,
+ UNORM_SHORT_555,
+ UNORM_INT_101010,
+ SIGNED_INT8,
+ SIGNED_INT16,
+ SIGNED_INT32,
+ UNSIGNED_INT8,
+ UNSIGNED_INT16,
+ UNSIGNED_INT32,
+ HALF_FLOAT,
+ FLOAT
+ }
+
+ /**
+ * {@code ImageChannelOrder} specifies the number of channels and the channel layout i.e. the
+memory layout in which channels are stored in the image.
+ */
+ public static enum ImageChannelOrder {
+ R, Rx, A,
+ INTENSITY,
+ LUMINANCE,
+ RG, RGx, RA,
+ RGB, RGBx,
+ RGBA,
+ ARGB, BGRA
+ }
+
+ /**
+ * Describes the image format, consisting of
+ * {@link ImageChannelOrder} and {@link ImageChannelType}.
+ */
+ public static class ImageFormat { //Struct
+ public ImageChannelOrder channelOrder;
+ public ImageChannelType channelType;
+
+ public ImageFormat() {
+ }
+
+ public ImageFormat(ImageChannelOrder channelOrder, ImageChannelType channelType) {
+ this.channelOrder = channelOrder;
+ this.channelType = channelType;
+ }
+
+ @Override
+ public String toString() {
+ return "ImageFormat{" + "channelOrder=" + channelOrder + ", channelType=" + channelType + '}';
+ }
+
+ @Override
+ public int hashCode() {
+ int hash = 5;
+ hash = 61 * hash + Objects.hashCode(this.channelOrder);
+ hash = 61 * hash + Objects.hashCode(this.channelType);
+ return hash;
+ }
+
+ @Override
+ public boolean equals(Object obj) {
+ if (obj == null) {
+ return false;
+ }
+ if (getClass() != obj.getClass()) {
+ return false;
+ }
+ final ImageFormat other = (ImageFormat) obj;
+ if (this.channelOrder != other.channelOrder) {
+ return false;
+ }
+ if (this.channelType != other.channelType) {
+ return false;
+ }
+ return true;
+ }
+
+ }
+
+ /**
+ * The possible image types / dimensions.
+ */
+ public static enum ImageType {
+ IMAGE_1D,
+ IMAGE_1D_BUFFER,
+ IMAGE_2D,
+ IMAGE_3D,
+ IMAGE_1D_ARRAY,
+ IMAGE_2D_ARRAY
+ }
+
+ /**
+ * The image descriptor structure describes the type and dimensions of the image or image array.
+ *
+ * There exists two constructors:
+ * {@link #ImageDescriptor(com.jme3.opencl.Image.ImageType, long, long, long, long) }
+ * is used when an image with new memory should be created (used most often).
+ * {@link #ImageDescriptor(com.jme3.opencl.Image.ImageType, long, long, long, long, long, long, java.nio.ByteBuffer) }
+ * creates an image using the provided {@code ByteBuffer} as source.
+ */
+ public static class ImageDescriptor { //Struct
+ public ImageType type;
+ public long width;
+ public long height;
+ public long depth;
+ public long arraySize;
+ public long rowPitch;
+ public long slicePitch;
+ public ByteBuffer hostPtr;
+ /*
+ public int numMipLevels; //They must always be set to zero
+ public int numSamples;
+ */
+
+ public ImageDescriptor() {
+ }
+
+ /**
+ * Used to specify an image with the provided ByteBuffer as soruce
+ * @param type the image type
+ * @param width the width
+ * @param height the height, unused for image types {@code ImageType.IMAGE_1D*}
+ * @param depth the depth of the image, only used for image type {@code ImageType.IMAGE_3D}
+ * @param arraySize the number of array elements for image type {@code ImageType.IMAGE_1D_ARRAY} and {@code ImageType.IMAGE_2D_ARRAY}
+ * @param rowPitch the row pitch of the provided buffer
+ * @param slicePitch the slice pitch of the provided buffer
+ * @param hostPtr host buffer used as image memory
+ */
+ public ImageDescriptor(ImageType type, long width, long height, long depth, long arraySize, long rowPitch, long slicePitch, ByteBuffer hostPtr) {
+ this.type = type;
+ this.width = width;
+ this.height = height;
+ this.depth = depth;
+ this.arraySize = arraySize;
+ this.rowPitch = rowPitch;
+ this.slicePitch = slicePitch;
+ this.hostPtr = hostPtr;
+ }
+ /**
+ * Specifies an image without a host buffer, a new chunk of memory
+ * will be allocated.
+ * @param type the image type
+ * @param width the width
+ * @param height the height, unused for image types {@code ImageType.IMAGE_1D*}
+ * @param depth the depth of the image, only used for image type {@code ImageType.IMAGE_3D}
+ * @param arraySize the number of array elements for image type {@code ImageType.IMAGE_1D_ARRAY} and {@code ImageType.IMAGE_2D_ARRAY}
+ */
+ public ImageDescriptor(ImageType type, long width, long height, long depth, long arraySize) {
+ this.type = type;
+ this.width = width;
+ this.height = height;
+ this.depth = depth;
+ this.arraySize = arraySize;
+ this.rowPitch = 0;
+ this.slicePitch = 0;
+ hostPtr = null;
+ }
+
+ @Override
+ public String toString() {
+ return "ImageDescriptor{" + "type=" + type + ", width=" + width + ", height=" + height + ", depth=" + depth + ", arraySize=" + arraySize + ", rowPitch=" + rowPitch + ", slicePitch=" + slicePitch + '}';
+ }
+
+ }
+
+ protected Image(ObjectReleaser releaser) {
+ super(releaser);
+ }
+
+ /**
+ * @return the width of the image
+ */
+ public abstract long getWidth();
+ /**
+ * @return the height of the image
+ */
+ public abstract long getHeight();
+ /**
+ * @return the depth of the image
+ */
+ public abstract long getDepth();
+ /**
+ * @return the row pitch when the image was created from a host buffer
+ * @see ImageDescriptor#ImageDescriptor(com.jme3.opencl.Image.ImageType, long, long, long, long, long, long, java.nio.ByteBuffer)
+ */
+ public abstract long getRowPitch();
+ /**
+ * @return the slice pitch when the image was created from a host buffer
+ * @see ImageDescriptor#ImageDescriptor(com.jme3.opencl.Image.ImageType, long, long, long, long, long, long, java.nio.ByteBuffer)
+ */
+ public abstract long getSlicePitch();
+ /**
+ * @return the number of elements in the image array
+ * @see ImageType#IMAGE_1D_ARRAY
+ * @see ImageType#IMAGE_2D_ARRAY
+ */
+ public abstract long getArraySize();
+ /**
+ * @return the image format
+ */
+ public abstract ImageFormat getImageFormat();
+ /**
+ * @return the image type
+ */
+ public abstract ImageType getImageType();
+ /**
+ * @return the number of bytes per pixel
+ */
+ public abstract int getElementSize();
+
+ /**
+ * Performs a blocking read of the image into the specified byte buffer.
+ * @param queue the command queue
+ * @param dest the target byte buffer
+ * @param origin the image origin location, see class description for the format
+ * @param region the copied region, see class description for the format
+ * @param rowPitch the row pitch of the target buffer, must be set to 0 if the image is 1D.
+ * If set to 0 for 2D and 3D image, the row pitch is calculated as {@code bytesPerElement * width}
+ * @param slicePitch the slice pitch of the target buffer, must be set to 0 for 1D and 2D images.
+ * If set to 0 for 3D images, the slice pitch is calculated as {@code rowPitch * height}
+ */
+ public abstract void readImage(CommandQueue queue, ByteBuffer dest, long[] origin, long[] region, long rowPitch, long slicePitch);
+ /**
+ * Performs an async/non-blocking read of the image into the specified byte buffer.
+ * @param queue the command queue
+ * @param dest the target byte buffer
+ * @param origin the image origin location, see class description for the format
+ * @param region the copied region, see class description for the format
+ * @param rowPitch the row pitch of the target buffer, must be set to 0 if the image is 1D.
+ * If set to 0 for 2D and 3D image, the row pitch is calculated as {@code bytesPerElement * width}
+ * @param slicePitch the slice pitch of the target buffer, must be set to 0 for 1D and 2D images.
+ * If set to 0 for 3D images, the slice pitch is calculated as {@code rowPitch * height}
+ * @return the event object indicating the status of the operation
+ */
+ public abstract Event readImageAsync(CommandQueue queue, ByteBuffer dest, long[] origin, long[] region, long rowPitch, long slicePitch);
+
+ /**
+ * Performs a blocking write from the specified byte buffer into the image.
+ * @param queue the command queue
+ * @param src the source buffer
+ * @param origin the image origin location, see class description for the format
+ * @param region the copied region, see class description for the format
+ * @param rowPitch the row pitch of the target buffer, must be set to 0 if the image is 1D.
+ * If set to 0 for 2D and 3D image, the row pitch is calculated as {@code bytesPerElement * width}
+ * @param slicePitch the slice pitch of the target buffer, must be set to 0 for 1D and 2D images.
+ * If set to 0 for 3D images, the slice pitch is calculated as {@code rowPitch * height}
+ */
+ public abstract void writeImage(CommandQueue queue, ByteBuffer src, long[] origin, long[] region, long rowPitch, long slicePitch);
+ /**
+ * Performs an async/non-blocking write from the specified byte buffer into the image.
+ * @param queue the command queue
+ * @param src the source buffer
+ * @param origin the image origin location, see class description for the format
+ * @param region the copied region, see class description for the format
+ * @param rowPitch the row pitch of the target buffer, must be set to 0 if the image is 1D.
+ * If set to 0 for 2D and 3D image, the row pitch is calculated as {@code bytesPerElement * width}
+ * @param slicePitch the slice pitch of the target buffer, must be set to 0 for 1D and 2D images.
+ * If set to 0 for 3D images, the slice pitch is calculated as {@code rowPitch * height}
+ * @return the event object indicating the status of the operation
+ */
+ public abstract Event writeImageAsync(CommandQueue queue, ByteBuffer src, long[] origin, long[] region, long rowPitch, long slicePitch);
+
+ /**
+ * Performs a blocking copy operation from one image to another.
+ * Important: Both images must have the same format!
+ * @param queue the command queue
+ * @param dest the target image
+ * @param srcOrigin the source image origin, see class description for the format
+ * @param destOrigin the target image origin, see class description for the format
+ * @param region the copied region, see class description for the format
+ */
+ public abstract void copyTo(CommandQueue queue, Image dest, long[] srcOrigin, long[] destOrigin, long[] region);
+ /**
+ * Performs an async/non-blocking copy operation from one image to another.
+ * Important: Both images must have the same format!
+ * @param queue the command queue
+ * @param dest the target image
+ * @param srcOrigin the source image origin, see class description for the format
+ * @param destOrigin the target image origin, see class description for the format
+ * @param region the copied region, see class description for the format
+ * @return the event object indicating the status of the operation
+ */
+ public abstract Event copyToAsync(CommandQueue queue, Image dest, long[] srcOrigin, long[] destOrigin, long[] region);
+
+ /**
+ * Maps the image into host memory.
+ * The returned structure contains the mapped byte buffer and row and slice pitch.
+ * The event object is set to {@code null}, it is needed for the asnyc
+ * version {@link #mapAsync(com.jme3.opencl.CommandQueue, long[], long[], com.jme3.opencl.MappingAccess) }.
+ * @param queue the command queue
+ * @param origin the image origin, see class description for the format
+ * @param region the mapped region, see class description for the format
+ * @param access the allowed memory access to the mapped memory
+ * @return a structure describing the mapped memory
+ * @see #unmap(com.jme3.opencl.CommandQueue, com.jme3.opencl.Image.ImageMapping)
+ */
+ public abstract ImageMapping map(CommandQueue queue, long[] origin, long[] region, MappingAccess access);
+ /**
+ * Non-blocking version of {@link #map(com.jme3.opencl.CommandQueue, long[], long[], com.jme3.opencl.MappingAccess) }.
+ * The returned structure contains the mapped byte buffer and row and slice pitch.
+ * The event object is used to detect when the mapped memory is available.
+ * @param queue the command queue
+ * @param origin the image origin, see class description for the format
+ * @param region the mapped region, see class description for the format
+ * @param access the allowed memory access to the mapped memory
+ * @return a structure describing the mapped memory
+ * @see #unmap(com.jme3.opencl.CommandQueue, com.jme3.opencl.Image.ImageMapping)
+ */
+ public abstract ImageMapping mapAsync(CommandQueue queue, long[] origin, long[] region, MappingAccess access);
+ /**
+ * Unmaps the mapped memory
+ * @param queue the command queue
+ * @param mapping the mapped memory
+ */
+ public abstract void unmap(CommandQueue queue, ImageMapping mapping);
+
+ /**
+ * Describes a mapped region of the image
+ */
+ public static class ImageMapping {
+ /**
+ * The raw byte buffer
+ */
+ public final ByteBuffer buffer;
+ /**
+ * The row pitch in bytes.
+ * This value is at least {@code bytesPerElement * width}
+ */
+ public final long rowPitch;
+ /**
+ * The slice pitch in bytes.
+ * This value is at least {@code rowPitch * height}
+ */
+ public final long slicePitch;
+ /**
+ * The event object used to detect when the memory is available.
+ * @see #mapAsync(com.jme3.opencl.CommandQueue, long[], long[], com.jme3.opencl.MappingAccess)
+ */
+ public final Event event;
+
+ public ImageMapping(ByteBuffer buffer, long rowPitch, long slicePitch, Event event) {
+ this.buffer = buffer;
+ this.rowPitch = rowPitch;
+ this.slicePitch = slicePitch;
+ this.event = event;
+ }
+ public ImageMapping(ByteBuffer buffer, long rowPitch, long slicePitch) {
+ this.buffer = buffer;
+ this.rowPitch = rowPitch;
+ this.slicePitch = slicePitch;
+ this.event = null;
+ }
+
+ }
+
+ /**
+ * Fills the image with the specified color.
+ * Does only work if the image channel is {@link ImageChannelType#FLOAT}
+ * or {@link ImageChannelType#HALF_FLOAT}.
+ * @param queue the command queue
+ * @param origin the image origin, see class description for the format
+ * @param region the size of the region, see class description for the format
+ * @param color the color to fill
+ * @return an event object to detect for the completion
+ */
+ public abstract Event fillAsync(CommandQueue queue, long[] origin, long[] region, ColorRGBA color);
+ /**
+ * Fills the image with the specified color given as four integer variables.
+ * Does not work if the image channel is {@link ImageChannelType#FLOAT}
+ * or {@link ImageChannelType#HALF_FLOAT}.
+ * @param queue the command queue
+ * @param origin the image origin, see class description for the format
+ * @param region the size of the region, see class description for the format
+ * @param color the color to fill, must be an array of length 4
+ * @return an event object to detect for the completion
+ */
+ public abstract Event fillAsync(CommandQueue queue, long[] origin, long[] region, int[] color);
+
+ /**
+ * Copies this image into the specified buffer, no format conversion is done.
+ * This is the dual function to
+ * {@link Buffer#copyToImageAsync(com.jme3.opencl.CommandQueue, com.jme3.opencl.Image, long, long[], long[]) }.
+ * @param queue the command queue
+ * @param dest the target buffer
+ * @param srcOrigin the image origin, see class description for the format
+ * @param srcRegion the copied region, see class description for the format
+ * @param destOffset an offset into the target buffer
+ * @return the event object to detect the completion of the operation
+ */
+ public abstract Event copyToBufferAsync(CommandQueue queue, Buffer dest, long[] srcOrigin, long[] srcRegion, long destOffset);
+
+ /**
+ * Aquires this image object for using. Only call this method if this image
+ * represents a shared object from OpenGL, created with e.g.
+ * {@link Context#bindImage(com.jme3.texture.Image, com.jme3.texture.Texture.Type, int, com.jme3.opencl.MemoryAccess) }
+ * or variations.
+ * This method must be called before the image is used. After the work is
+ * done, the image must be released by calling
+ * {@link #releaseImageForSharingAsync(com.jme3.opencl.CommandQueue) }
+ * so that OpenGL can use the image/texture/renderbuffer again.
+ * @param queue the command queue
+ * @return the event object
+ */
+ public abstract Event acquireImageForSharingAsync(CommandQueue queue);
+
+ /**
+ * Aquires this image object for using. Only call this method if this image
+ * represents a shared object from OpenGL, created with e.g.
+ * {@link Context#bindImage(com.jme3.texture.Image, com.jme3.texture.Texture.Type, int, com.jme3.opencl.MemoryAccess) }
+ * or variations.
+ * This method must be called before the image is used. After the work is
+ * done, the image must be released by calling
+ * {@link #releaseImageForSharingAsync(com.jme3.opencl.CommandQueue) }
+ * so that OpenGL can use the image/texture/renderbuffer again.
+ *
+ * The generated event object is directly released.
+ * This brings a performance improvement when the resource is e.g. directly
+ * used by a kernel afterwards on the same queue (this implicitly waits for
+ * this action). If you need the event, use
+ * {@link #acquireImageForSharingAsync(com.jme3.opencl.CommandQueue) }.
+ *
+ * @param queue the command queue
+ */
+ public void acquireImageForSharingNoEvent(CommandQueue queue) {
+ //Default implementation, overwrite for performance
+ acquireImageForSharingAsync(queue).release();
+ }
+
+ /**
+ * Releases a shared image object.
+ * Call this method after the image object was acquired by
+ * {@link #acquireImageForSharingAsync(com.jme3.opencl.CommandQueue) }
+ * to hand the control back to OpenGL.
+ * @param queue the command queue
+ * @return the event object
+ */
+ public abstract Event releaseImageForSharingAsync(CommandQueue queue);
+
+ /**
+ * Releases a shared image object.
+ * Call this method after the image object was acquired by
+ * {@link #acquireImageForSharingAsync(com.jme3.opencl.CommandQueue) }
+ * to hand the control back to OpenGL.
+ * The generated event object is directly released, resulting in
+ * performance improvements.
+ * @param queue the command queue
+ */
+ public void releaseImageForSharingNoEvent(CommandQueue queue) {
+ //default implementation, overwrite it for performance improvements
+ releaseImageForSharingAsync(queue).release();
+ }
+
+ //TODO: add variants of the above two methods that don't create the event object, but release the event immediately
+}
diff --git a/jme3-core/src/main/java/com/jme3/opencl/Kernel.java b/jme3-core/src/main/java/com/jme3/opencl/Kernel.java
new file mode 100644
index 000000000..09f22a813
--- /dev/null
+++ b/jme3-core/src/main/java/com/jme3/opencl/Kernel.java
@@ -0,0 +1,628 @@
+/*
+ * Copyright (c) 2009-2016 jMonkeyEngine
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are
+ * met:
+ *
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ *
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ *
+ * * Neither the name of 'jMonkeyEngine' nor the names of its contributors
+ * may be used to endorse or promote products derived from this software
+ * without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+ * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+ * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
+ * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
+ * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+package com.jme3.opencl;
+
+import com.jme3.math.*;
+import com.jme3.util.TempVars;
+import java.nio.ByteBuffer;
+import java.util.Arrays;
+
+/**
+ * Wrapper for an OpenCL kernel, a piece of executable code on the GPU.
+ *
+ * Terminology:
+ * A Kernel is executed in parallel. In total number of parallel threads,
+ * called work items, are specified by the global work size (of type
+ * {@link WorkSize}. These threads are organized in a 1D, 2D or 3D grid
+ * (of coarse, this is only a logical view). Inside each kernel,
+ * the id of each thread (i.e. the index inside this grid) can be requested
+ * by {@code get_global_id(dimension)} with {@code dimension=0,1,2}.
+ *
+ * Not all threads can always be executed in parallel because there simply might
+ * not be enough processor cores.
+ * Therefore, the concept of a work group is introduced. The work group
+ * specifies the actual number of threads that are executed in parallel.
+ * The maximal size of it can be queried by {@link Device#getMaxiumWorkItemsPerGroup() }.
+ * Again, the threads inside the work group can be organized in a 1D, 2D or 3D
+ * grid, but this is also just a logical view (specifying how the threads are
+ * indexed).
+ * The work group is imporatant for another concept: shared memory
+ * Unlike the normal global or constant memory (passing a {@link Buffer} object
+ * as argument), shared memory can't be set from outside. Shared memory is
+ * allocated by the kernel and is only valid within the kernel. It is used
+ * to quickly share data between threads within a work group.
+ * The size of the shared memory is specified by setting an instance of
+ * {@link LocalMem} or {@link LocalMemPerElement} as argument.
+ * Due to heavy register usage or other reasons, a kernel might not be able
+ * to utilize a whole work group. Therefore, the actual number of threads
+ * that can be executed in a work group can be queried by
+ * {@link #getMaxWorkGroupSize(com.jme3.opencl.Device) }, which might differ from the
+ * value returned from the Device.
+ *
+ *
+ * There are two ways to launch a kernel:
+ * First, arguments and the work group sizes can be set in advance
+ * ({@code setArg(index, ...)}, {@code setGlobalWorkSize(...)} and {@code setWorkGroupSize(...)}.
+ * Then a kernel is launched by {@link #Run(com.jme3.opencl.CommandQueue) }.
+ * Second, two convenient functions are provided that set the arguments
+ * and work sizes in one call:
+ * {@link #Run1(com.jme3.opencl.CommandQueue, com.jme3.opencl.Kernel.WorkSize, java.lang.Object...) }
+ * and {@link #Run2(com.jme3.opencl.CommandQueue, com.jme3.opencl.Kernel.WorkSize, com.jme3.opencl.Kernel.WorkSize, java.lang.Object...) }.
+ *
+ * @author shaman
+ * @see Program#createKernel(java.lang.String)
+ */
+public abstract class Kernel extends AbstractOpenCLObject {
+ /**
+ * The current global work size
+ */
+ protected final WorkSize globalWorkSize;
+ /**
+ * The current local work size
+ */
+ protected final WorkSize workGroupSize;
+
+ protected Kernel(ObjectReleaser releaser) {
+ super(releaser);
+ this.globalWorkSize = new WorkSize(0);
+ this.workGroupSize = new WorkSize(0);
+ }
+
+ /**
+ * @return the name of the kernel as defined in the program source code
+ */
+ public abstract String getName();
+
+ /**
+ * @return the number of arguments
+ */
+ public abstract int getArgCount();
+
+ /**
+ * @return the current global work size
+ */
+ public WorkSize getGlobalWorkSize() {
+ return globalWorkSize;
+ }
+
+ /**
+ * Sets the global work size.
+ * @param ws the work size to set
+ */
+ public void setGlobalWorkSize(WorkSize ws) {
+ globalWorkSize.set(ws);
+ }
+
+ /**
+ * Sets the global work size to a 1D grid
+ * @param size the size in 1D
+ */
+ public void setGlobalWorkSize(int size) {
+ globalWorkSize.set(1, size);
+ }
+
+ /**
+ * Sets the global work size to be a 2D grid
+ * @param width the width
+ * @param height the height
+ */
+ public void setGlobalWorkSize(int width, int height) {
+ globalWorkSize.set(2, width, height);
+ }
+
+ /**
+ * Sets the global work size to be a 3D grid
+ * @param width the width
+ * @param height the height
+ * @param depth the depth
+ */
+ public void setGlobalWorkSize(int width, int height, int depth) {
+ globalWorkSize.set(3, width, height, depth);
+ }
+
+ /**
+ * @return the current work group size
+ */
+ public WorkSize getWorkGroupSize() {
+ return workGroupSize;
+ }
+
+ /**
+ * Sets the work group size
+ * @param ws the work group size to set
+ */
+ public void setWorkGroupSize(WorkSize ws) {
+ workGroupSize.set(ws);
+ }
+
+ /**
+ * Sets the work group size to be a 1D grid
+ * @param size the size to set
+ */
+ public void setWorkGroupSize(int size) {
+ workGroupSize.set(1, size);
+ }
+
+ /**
+ * Sets the work group size to be a 2D grid
+ * @param width the width
+ * @param height the height
+ */
+ public void setWorkGroupSize(int width, int height) {
+ workGroupSize.set(2, width, height);
+ }
+
+ /**
+ * Sets the work group size to be a 3D grid
+ * @param width the width
+ * @param height the height
+ * @param depth the depth
+ */
+ public void setWorkGroupSdize(int width, int height, int depth) {
+ workGroupSize.set(3, width, height, depth);
+ }
+
+ /**
+ * Tells the driver to figure out the work group size on their own.
+ * Use this if you do not rely on specific work group layouts, i.e.
+ * because shared memory is not used.
+ * {@link #Run1(com.jme3.opencl.CommandQueue, com.jme3.opencl.Kernel.WorkSize, java.lang.Object...) }
+ * implicetly calls this mehtod.
+ */
+ public void setWorkGroupSizeToNull() {
+ workGroupSize.set(1, 0, 0, 0);
+ }
+
+ /**
+ * Returns the maximal work group size when this kernel is executed on
+ * the specified device
+ * @param device the device
+ * @return the maximal work group size
+ */
+ public abstract long getMaxWorkGroupSize(Device device);
+
+ public abstract void setArg(int index, LocalMemPerElement t);
+
+ public abstract void setArg(int index, LocalMem t);
+
+ public abstract void setArg(int index, Buffer t);
+
+ public abstract void setArg(int index, Image i);
+
+ public abstract void setArg(int index, byte b);
+
+ public abstract void setArg(int index, short s);
+
+ public abstract void setArg(int index, int i);
+
+ public abstract void setArg(int index, long l);
+
+ public abstract void setArg(int index, float f);
+
+ public abstract void setArg(int index, double d);
+
+ public abstract void setArg(int index, Vector2f v);
+
+ public abstract void setArg(int index, Vector4f v);
+
+ public abstract void setArg(int index, Quaternion q);
+
+ public abstract void setArg(int index, Matrix4f mat);
+
+ public void setArg(int index, Matrix3f mat) {
+ TempVars vars = TempVars.get();
+ try {
+ Matrix4f m = vars.tempMat4;
+ m.zero();
+ for (int i=0; i<3; ++i) {
+ for (int j=0; j<3; ++j) {
+ m.set(i, j, mat.get(i, j));
+ }
+ }
+ setArg(index, m);
+ } finally {
+ vars.release();
+ }
+ }
+
+ /**
+ * Raw version to set an argument.
+ * {@code size} bytes of the provided byte buffer are copied to the kernel
+ * argument. The size in bytes must match exactly the argument size
+ * as defined in the kernel code.
+ * Use this method to send custom structures to the kernel
+ * @param index the index of the argument
+ * @param buffer the raw buffer
+ * @param size the size in bytes
+ */
+ public abstract void setArg(int index, ByteBuffer buffer, long size);
+
+ /**
+ * Sets the kernel argument at the specified index.
+ * The argument must be a known type:
+ * {@code LocalMemPerElement, LocalMem, Image, Buffer, byte, short, int,
+ * long, float, double, Vector2f, Vector4f, Quaternion, Matrix3f, Matrix4f}.
+ *
+ * Note: Matrix3f and Matrix4f will be mapped to a {@code float16} (row major).
+ * @param index the index of the argument, from 0 to {@link #getArgCount()}-1
+ * @param arg the argument
+ * @throws IllegalArgumentException if the argument type is not one of the listed ones
+ */
+ public void setArg(int index, Object arg) {
+ if (arg instanceof Byte) {
+ setArg(index, (byte) arg);
+ } else if (arg instanceof Short) {
+ setArg(index, (short) arg);
+ } else if (arg instanceof Integer) {
+ setArg(index, (int) arg);
+ } else if (arg instanceof Long) {
+ setArg(index, (long) arg);
+ } else if (arg instanceof Float) {
+ setArg(index, (float) arg);
+ } else if (arg instanceof Double) {
+ setArg(index, (double) arg);
+ } else if (arg instanceof Vector2f) {
+ setArg(index, (Vector2f) arg);
+ } else if (arg instanceof Vector4f) {
+ setArg(index, (Vector4f) arg);
+ } else if (arg instanceof Quaternion) {
+ setArg(index, (Quaternion) arg);
+ } else if (arg instanceof Matrix3f) {
+ setArg(index, (Matrix3f) arg);
+ } else if (arg instanceof Matrix4f) {
+ setArg(index, (Matrix4f) arg);
+ } else if (arg instanceof LocalMemPerElement) {
+ setArg(index, (LocalMemPerElement) arg);
+ } else if (arg instanceof LocalMem) {
+ setArg(index, (LocalMem) arg);
+ } else if (arg instanceof Buffer) {
+ setArg(index, (Buffer) arg);
+ } else if (arg instanceof Image) {
+ setArg(index, (Image) arg);
+ } else {
+ throw new IllegalArgumentException("unknown kernel argument type: " + arg);
+ }
+ }
+
+ private void setArgs(Object... args) {
+ for (int i = 0; i < args.length; ++i) {
+ setArg(i, args[i]);
+ }
+ }
+
+ /**
+ * Launches the kernel with the current global work size, work group size
+ * and arguments.
+ * If the returned event object is not needed and would otherwise be
+ * released immediately, {@link #RunNoEvent(com.jme3.opencl.CommandQueue) }
+ * might bring a better performance.
+ * @param queue the command queue
+ * @return an event object indicating when the kernel is finished
+ * @see #setGlobalWorkSize(com.jme3.opencl.Kernel.WorkSize)
+ * @see #setWorkGroupSize(com.jme3.opencl.Kernel.WorkSize)
+ * @see #setArg(int, java.lang.Object)
+ */
+ public abstract Event Run(CommandQueue queue);
+
+ /**
+ * Launches the kernel with the current global work size, work group size
+ * and arguments without returning an event object.
+ * The generated event is directly released. Therefore, the performance
+ * is better, but there is no way to detect when the kernel execution
+ * has finished. For this purpose, use {@link #Run(com.jme3.opencl.CommandQueue) }.
+ * @param queue the command queue
+ * @see #setGlobalWorkSize(com.jme3.opencl.Kernel.WorkSize)
+ * @see #setWorkGroupSize(com.jme3.opencl.Kernel.WorkSize)
+ * @see #setArg(int, java.lang.Object)
+ */
+ public void RunNoEvent(CommandQueue queue) {
+ //Default implementation, overwrite to not allocate the event object
+ Run(queue).release();
+ }
+
+ /**
+ * Sets the work sizes and arguments in one call and launches the kernel.
+ * The global work size is set to the specified size. The work group
+ * size is automatically determined by the driver.
+ * Each object in the argument array is sent to the kernel by
+ * {@link #setArg(int, java.lang.Object) }.
+ * @param queue the command queue
+ * @param globalWorkSize the global work size
+ * @param args the kernel arguments
+ * @return an event object indicating when the kernel is finished
+ * @see #Run2(com.jme3.opencl.CommandQueue, com.jme3.opencl.Kernel.WorkSize, com.jme3.opencl.Kernel.WorkSize, java.lang.Object...)
+ */
+ public Event Run1(CommandQueue queue, WorkSize globalWorkSize, Object... args) {
+ setGlobalWorkSize(globalWorkSize);
+ setWorkGroupSizeToNull();
+ setArgs(args);
+ return Run(queue);
+ }
+
+ /**
+ * Sets the work sizes and arguments in one call and launches the kernel.
+ * The global work size is set to the specified size. The work group
+ * size is automatically determined by the driver.
+ * Each object in the argument array is sent to the kernel by
+ * {@link #setArg(int, java.lang.Object) }.
+ * The generated event is directly released. Therefore, the performance
+ * is better, but there is no way to detect when the kernel execution
+ * has finished. For this purpose, use
+ * {@link #Run1(com.jme3.opencl.CommandQueue, com.jme3.opencl.Kernel.WorkSize, java.lang.Object...) }.
+ * @param queue the command queue
+ * @param globalWorkSize the global work size
+ * @param args the kernel arguments
+ * @see #Run2(com.jme3.opencl.CommandQueue, com.jme3.opencl.Kernel.WorkSize, com.jme3.opencl.Kernel.WorkSize, java.lang.Object...)
+ */
+ public void Run1NoEvent(CommandQueue queue, WorkSize globalWorkSize, Object... args) {
+ setGlobalWorkSize(globalWorkSize);
+ setWorkGroupSizeToNull();
+ setArgs(args);
+ RunNoEvent(queue);
+ }
+
+ /**
+ * Sets the work sizes and arguments in one call and launches the kernel.
+ * @param queue the command queue
+ * @param globalWorkSize the global work size
+ * @param workGroupSize the work group size
+ * @param args the kernel arguments
+ * @return an event object indicating when the kernel is finished
+ */
+ public Event Run2(CommandQueue queue, WorkSize globalWorkSize,
+ WorkSize workGroupSize, Object... args) {
+ setGlobalWorkSize(globalWorkSize);
+ setWorkGroupSize(workGroupSize);
+ setArgs(args);
+ return Run(queue);
+ }
+
+ /**
+ * Sets the work sizes and arguments in one call and launches the kernel.
+ * The generated event is directly released. Therefore, the performance
+ * is better, but there is no way to detect when the kernel execution
+ * has finished. For this purpose, use
+ * {@link #Run2(com.jme3.opencl.CommandQueue, com.jme3.opencl.Kernel.WorkSize, com.jme3.opencl.Kernel.WorkSize, java.lang.Object...) }.
+ * @param queue the command queue
+ * @param globalWorkSize the global work size
+ * @param workGroupSize the work group size
+ * @param args the kernel arguments
+ */
+ public void Run2NoEvent(CommandQueue queue, WorkSize globalWorkSize,
+ WorkSize workGroupSize, Object... args) {
+ setGlobalWorkSize(globalWorkSize);
+ setWorkGroupSize(workGroupSize);
+ setArgs(args);
+ RunNoEvent(queue);
+ }
+
+ /**
+ * A placeholder for kernel arguments representing local kernel memory.
+ * This defines the size of available shared memory of a {@code __shared} kernel
+ * argument
+ */
+ public static final class LocalMem {
+
+ private int size;
+
+ /**
+ * Creates a new LocalMem instance
+ * @param size the size of the available shared memory in bytes
+ */
+ public LocalMem(int size) {
+ super();
+ this.size = size;
+ }
+
+ public int getSize() {
+ return size;
+ }
+
+ @Override
+ public int hashCode() {
+ int hash = 3;
+ hash = 79 * hash + this.size;
+ return hash;
+ }
+
+ @Override
+ public boolean equals(Object obj) {
+ if (obj == null) {
+ return false;
+ }
+ if (getClass() != obj.getClass()) {
+ return false;
+ }
+ final LocalMem other = (LocalMem) obj;
+ if (this.size != other.size) {
+ return false;
+ }
+ return true;
+ }
+ }
+
+ /**
+ * A placeholder for a kernel argument representing local kernel memory per thread.
+ * This effectively computes {@code SharedMemoryPerElement * WorkGroupSize}
+ * and uses this value as the size of shared memory available in the kernel.
+ * Therefore, an instance of this class must be set as an argument AFTER
+ * the work group size has been specified. This is
+ * ensured by {@link #Run2(com.jme3.opencl.CommandQueue, com.jme3.opencl.Kernel.WorkSize, com.jme3.opencl.Kernel.WorkSize, java.lang.Object...) }.
+ * This argument can't be used when no work group size was defined explicetly
+ * (e.g. by {@link #setWorkGroupSizeToNull()} or {@link #Run1(com.jme3.opencl.CommandQueue, com.jme3.opencl.Kernel.WorkSize, java.lang.Object...) }.
+ */
+ public static final class LocalMemPerElement {
+
+ private int size;
+
+ /**
+ * Creates a new LocalMemPerElement instance
+ * @param size the number of bytes available for each thread within
+ * a work group
+ */
+ public LocalMemPerElement(int size) {
+ super();
+ this.size = size;
+ }
+
+ public int getSize() {
+ return size;
+ }
+
+ @Override
+ public int hashCode() {
+ int hash = 3;
+ hash = 79 * hash + this.size;
+ return hash;
+ }
+
+ @Override
+ public boolean equals(Object obj) {
+ if (obj == null) {
+ return false;
+ }
+ if (getClass() != obj.getClass()) {
+ return false;
+ }
+ final LocalMemPerElement other = (LocalMemPerElement) obj;
+ if (this.size != other.size) {
+ return false;
+ }
+ return true;
+ }
+ }
+
+ /**
+ * The work size (global and local) for executing a kernel
+ * @author shaman
+ */
+ public static final class WorkSize {
+
+ private int dimension;
+ private long[] sizes;
+
+ /**
+ * Creates a new work size object
+ * @param dimension the dimension (1,2,3)
+ * @param sizes the sizes in each dimension, the length must match the specified dimension
+ */
+ public WorkSize(int dimension, long... sizes) {
+ super();
+ set(dimension, sizes);
+ }
+
+ /**
+ * Creates a work size of dimension 1 and extend 1,1,1 (only one thread).
+ */
+ public WorkSize() {
+ this(1, 1, 1, 1);
+ }
+
+ /**
+ * Creates a 1D work size of the specified extend
+ * @param size the size
+ */
+ public WorkSize(long size) {
+ this(1, size, 1, 1);
+ }
+
+ /**
+ * Creates a 2D work size of the specified extend
+ * @param width the width
+ * @param height the height
+ */
+ public WorkSize(long width, long height) {
+ this(2, width, height, 1);
+ }
+
+ /**
+ * Creates a 3D work size of the specified extend.
+ * @param width the width
+ * @param height the height
+ * @param depth the depth
+ */
+ public WorkSize(long width, long height, long depth) {
+ this(3, width, height, depth);
+ }
+
+ public int getDimension() {
+ return dimension;
+ }
+
+ public long[] getSizes() {
+ return sizes;
+ }
+
+ public void set(int dimension, long... sizes) {
+ if (sizes == null || sizes.length != 3) {
+ throw new IllegalArgumentException("sizes must be an array of length 3");
+ }
+ if (dimension <= 0 || dimension > 3) {
+ throw new IllegalArgumentException("dimension must be between 1 and 3");
+ }
+ this.dimension = dimension;
+ this.sizes = sizes;
+ }
+
+ public void set(WorkSize ws) {
+ this.dimension = ws.dimension;
+ this.sizes = ws.sizes;
+ }
+
+ @Override
+ public int hashCode() {
+ int hash = 5;
+ hash = 47 * hash + this.dimension;
+ hash = 47 * hash + Arrays.hashCode(this.sizes);
+ return hash;
+ }
+
+ @Override
+ public boolean equals(Object obj) {
+ if (obj == null) {
+ return false;
+ }
+ if (getClass() != obj.getClass()) {
+ return false;
+ }
+ final WorkSize other = (WorkSize) obj;
+ if (this.dimension != other.dimension) {
+ return false;
+ }
+ if (!Arrays.equals(this.sizes, other.sizes)) {
+ return false;
+ }
+ return true;
+ }
+ }
+
+}
diff --git a/jme3-core/src/main/java/com/jme3/opencl/KernelCompilationException.java b/jme3-core/src/main/java/com/jme3/opencl/KernelCompilationException.java
new file mode 100644
index 000000000..a0e20c9e5
--- /dev/null
+++ b/jme3-core/src/main/java/com/jme3/opencl/KernelCompilationException.java
@@ -0,0 +1,58 @@
+/*
+ * Copyright (c) 2009-2016 jMonkeyEngine
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are
+ * met:
+ *
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ *
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ *
+ * * Neither the name of 'jMonkeyEngine' nor the names of its contributors
+ * may be used to endorse or promote products derived from this software
+ * without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+ * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+ * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
+ * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
+ * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+package com.jme3.opencl;
+
+/**
+ * This exception is thrown by {@link Program#build() } and {@link Program#build(java.lang.String) }
+ * when the compilation failed.
+ * The error log returned by {@link #getLog() } contains detailed information
+ * where the error occured.
+ * @author shaman
+ */
+public class KernelCompilationException extends OpenCLException {
+
+ private final String log;
+
+ public KernelCompilationException(String msg, int errorCode, String log) {
+ super(msg, errorCode);
+ this.log = log;
+ }
+
+ /**
+ * The output of the compiler
+ * @return
+ */
+ public String getLog() {
+ return log;
+ }
+
+}
\ No newline at end of file
diff --git a/jme3-core/src/main/java/com/jme3/opencl/MappingAccess.java b/jme3-core/src/main/java/com/jme3/opencl/MappingAccess.java
new file mode 100644
index 000000000..121c6c8cb
--- /dev/null
+++ b/jme3-core/src/main/java/com/jme3/opencl/MappingAccess.java
@@ -0,0 +1,58 @@
+/*
+ * Copyright (c) 2009-2016 jMonkeyEngine
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are
+ * met:
+ *
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ *
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ *
+ * * Neither the name of 'jMonkeyEngine' nor the names of its contributors
+ * may be used to endorse or promote products derived from this software
+ * without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+ * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+ * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
+ * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
+ * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+package com.jme3.opencl;
+
+/**
+ * Specifies the access flags when mapping a {@link Buffer} or {@link Image} object.
+ * @see Buffer#map(com.jme3.opencl.CommandQueue, long, long, com.jme3.opencl.MappingAccess)
+ * @see Image#map(com.jme3.opencl.CommandQueue, long[], long[], com.jme3.opencl.MappingAccess)
+ * @author shaman
+ */
+public enum MappingAccess {
+ /**
+ * Only read access is allowed to the mapped memory.
+ */
+ MAP_READ_ONLY,
+ /**
+ * Only write access is allowed to the mapped memory.
+ */
+ MAP_WRITE_ONLY,
+ /**
+ * Both read and write access is allowed.
+ */
+ MAP_READ_WRITE,
+ /**
+ * The old memory content is completely discarded and the buffer is filled
+ * completely with new data. This might be faster than {@link #MAP_WRITE_ONLY}
+ */
+ MAP_WRITE_INVALIDATE
+}
diff --git a/jme3-core/src/main/java/com/jme3/opencl/MemoryAccess.java b/jme3-core/src/main/java/com/jme3/opencl/MemoryAccess.java
new file mode 100644
index 000000000..d5071665e
--- /dev/null
+++ b/jme3-core/src/main/java/com/jme3/opencl/MemoryAccess.java
@@ -0,0 +1,52 @@
+/*
+ * Copyright (c) 2009-2016 jMonkeyEngine
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are
+ * met:
+ *
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ *
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ *
+ * * Neither the name of 'jMonkeyEngine' nor the names of its contributors
+ * may be used to endorse or promote products derived from this software
+ * without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+ * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+ * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
+ * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
+ * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+package com.jme3.opencl;
+
+/**
+ * Specifies how a buffer object can be accessed by the kernel.
+ * @author shaman
+ * @see Buffer
+ */
+public enum MemoryAccess {
+ /**
+ * A kernel can both read and write the buffer.
+ */
+ READ_WRITE,
+ /**
+ * A kernel can only write this buffer.
+ */
+ WRITE_ONLY,
+ /**
+ * A kernel can only read this buffer
+ */
+ READ_ONLY
+}
diff --git a/jme3-core/src/main/java/com/jme3/opencl/OpenCLException.java b/jme3-core/src/main/java/com/jme3/opencl/OpenCLException.java
new file mode 100644
index 000000000..fd0836ac9
--- /dev/null
+++ b/jme3-core/src/main/java/com/jme3/opencl/OpenCLException.java
@@ -0,0 +1,78 @@
+/*
+ * Copyright (c) 2009-2016 jMonkeyEngine
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are
+ * met:
+ *
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ *
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ *
+ * * Neither the name of 'jMonkeyEngine' nor the names of its contributors
+ * may be used to endorse or promote products derived from this software
+ * without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+ * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+ * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
+ * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
+ * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+package com.jme3.opencl;
+
+/**
+ * Generic OpenCL exception, can be thrown in every method of this package.
+ * The error code and its name is reported in the message string as well as the OpenCL call that
+ * causes this exception. Please refer to the official OpenCL specification
+ * to see what might cause this exception.
+ * @author shaman
+ */
+public class OpenCLException extends RuntimeException {
+ private static final long serialVersionUID = 8471229972153694848L;
+
+ private final int errorCode;
+
+ /**
+ * Creates a new instance of OpenCLExceptionn without detail
+ * message.
+ */
+ public OpenCLException() {
+ errorCode = 0;
+ }
+
+ /**
+ * Constructs an instance of OpenCLExceptionn with the
+ * specified detail message.
+ *
+ * @param msg the detail message.
+ */
+ public OpenCLException(String msg) {
+ super(msg);
+ errorCode = 0;
+ }
+
+ public OpenCLException(String msg, int errorCode) {
+ super(msg);
+ this.errorCode = errorCode;
+ }
+
+ /**
+ * @return the error code
+ */
+ public int getErrorCode() {
+ return errorCode;
+ }
+
+
+}
diff --git a/jme3-core/src/main/java/com/jme3/opencl/OpenCLObject.java b/jme3-core/src/main/java/com/jme3/opencl/OpenCLObject.java
new file mode 100644
index 000000000..f9d631346
--- /dev/null
+++ b/jme3-core/src/main/java/com/jme3/opencl/OpenCLObject.java
@@ -0,0 +1,75 @@
+/*
+ * Copyright (c) 2009-2016 jMonkeyEngine
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are
+ * met:
+ *
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ *
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ *
+ * * Neither the name of 'jMonkeyEngine' nor the names of its contributors
+ * may be used to endorse or promote products derived from this software
+ * without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+ * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+ * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
+ * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
+ * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+package com.jme3.opencl;
+
+/**
+ * Base interface of all native OpenCL objects.
+ * This interface provides the functionality for savely release the object.
+ * @author shaman
+ */
+public interface OpenCLObject {
+
+ /**
+ * Releaser for an {@link OpenCLObject}.
+ * Implementations of this interface must not hold a reference to the
+ * {@code OpenCLObject} directly.
+ */
+ public static interface ObjectReleaser {
+ /**
+ * Releases the native resources of the associated {@link OpenCLObject}.
+ * This method must be guarded against multiple calls: only the first
+ * call should release, the next ones must not throw an exception.
+ */
+ void release();
+ }
+ /**
+ * Returns the releaser object. Multiple calls should return the same object.
+ * The ObjectReleaser is used to release the OpenCLObject when it is garbage
+ * collected. Therefore, the returned object must not hold a reference to
+ * the OpenCLObject.
+ * @return the object releaser
+ */
+ ObjectReleaser getReleaser();
+ /**
+ * Releases this native object.
+ * Should delegate to {@code getReleaser().release()}.
+ */
+ void release();
+ /**
+ * Registers this object for automatic releasing on garbage collection.
+ * By default, OpenCLObjects are not registered in the
+ * {@link OpenCLObjectManager}, you have to release it manually
+ * by calling {@link #release() }.
+ * Without registering or releasing, a memory leak might occur.
+ */
+ void register();
+}
diff --git a/jme3-core/src/main/java/com/jme3/opencl/OpenCLObjectManager.java b/jme3-core/src/main/java/com/jme3/opencl/OpenCLObjectManager.java
new file mode 100644
index 000000000..4b761da5b
--- /dev/null
+++ b/jme3-core/src/main/java/com/jme3/opencl/OpenCLObjectManager.java
@@ -0,0 +1,123 @@
+/*
+ * Copyright (c) 2009-2016 jMonkeyEngine
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are
+ * met:
+ *
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ *
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ *
+ * * Neither the name of 'jMonkeyEngine' nor the names of its contributors
+ * may be used to endorse or promote products derived from this software
+ * without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+ * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+ * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
+ * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
+ * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+package com.jme3.opencl;
+
+import java.lang.ref.PhantomReference;
+import java.lang.ref.ReferenceQueue;
+import java.util.HashSet;
+import java.util.logging.Level;
+import java.util.logging.Logger;
+
+/**
+ *
+ * @author shaman
+ */
+public class OpenCLObjectManager {
+ private static final Logger LOG = Logger.getLogger(OpenCLObjectManager.class.getName());
+ private static final Level LOG_LEVEL1 = Level.FINER;
+ private static final Level LOG_LEVEL2 = Level.FINE;
+ /**
+ * Call Runtime.getRuntime().gc() every these frames
+ */
+ private static final int GC_FREQUENCY = 10;
+
+ private static final OpenCLObjectManager INSTANCE = new OpenCLObjectManager();
+ private OpenCLObjectManager() {}
+
+ public static OpenCLObjectManager getInstance() {
+ return INSTANCE;
+ }
+
+ private ReferenceQueue