+ * 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 Sebastian Weiss
+ * @see Program#createKernel(java.lang.String)
*/
public abstract class Kernel implements OpenCLObject {
+ /**
+ * The current global work size
+ */
protected final WorkSize globalWorkSize;
+ /**
+ * The current local work size
+ */
protected final WorkSize workGroupSize;
protected Kernel() {
@@ -49,53 +97,118 @@ public abstract class Kernel implements OpenCLObject {
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);
}
- public void setWorkGroupSize(int width, int height, int depth) {
+ /**
+ * 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);
@@ -123,8 +236,27 @@ public abstract class Kernel implements OpenCLObject {
public abstract void setArg(int index, Quaternion q);
+ /**
+ * 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}
+ * @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);
@@ -163,8 +295,29 @@ public abstract class Kernel implements OpenCLObject {
}
}
+ /**
+ * Launches the kernel with the current global work size, work group size
+ * and arguments.
+ * @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);
+ /**
+ * 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();
@@ -172,6 +325,14 @@ public abstract class Kernel implements OpenCLObject {
return Run(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);
@@ -181,5 +342,205 @@ public abstract class Kernel implements OpenCLObject {
}
//TODO: add variants of the above three methods that don't create the event object, but release the event immediately
+
+ /**
+ * 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 Sebastian Weiss
+ */
+ 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/LocalMem.java b/jme3-core/src/main/java/com/jme3/opencl/LocalMem.java
deleted file mode 100644
index 71a2c8876..000000000
--- a/jme3-core/src/main/java/com/jme3/opencl/LocalMem.java
+++ /dev/null
@@ -1,71 +0,0 @@
-/*
- * 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;
-
-/**
- * A placeholder for kernel launches representing local kernel memory
- * @author Sebastian Weiss
- */
-public final class LocalMem {
- private int size;
-
- public LocalMem(int size) {
- 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;
- }
-
-}
diff --git a/jme3-core/src/main/java/com/jme3/opencl/LocalMemPerElement.java b/jme3-core/src/main/java/com/jme3/opencl/LocalMemPerElement.java
deleted file mode 100644
index 287eb1f4d..000000000
--- a/jme3-core/src/main/java/com/jme3/opencl/LocalMemPerElement.java
+++ /dev/null
@@ -1,71 +0,0 @@
-/*
- * 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;
-
-/**
- * A placeholder for kernel launches representing local kernel memory per thread
- * @author Sebastian Weiss
- */
-public final class LocalMemPerElement {
- private int size;
-
- public LocalMemPerElement(int size) {
- 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;
- }
-
-}
diff --git a/jme3-core/src/main/java/com/jme3/opencl/MappingAccess.java b/jme3-core/src/main/java/com/jme3/opencl/MappingAccess.java
index 0f1d3bced..07b89970d 100644
--- a/jme3-core/src/main/java/com/jme3/opencl/MappingAccess.java
+++ b/jme3-core/src/main/java/com/jme3/opencl/MappingAccess.java
@@ -32,12 +32,27 @@
package com.jme3.opencl;
/**
- * Specifies the access flags when mapping a {@link Buffer} object.
+ * 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 Sebastian Weiss
*/
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
index 0d03791ef..47bf9f75e 100644
--- a/jme3-core/src/main/java/com/jme3/opencl/MemoryAccess.java
+++ b/jme3-core/src/main/java/com/jme3/opencl/MemoryAccess.java
@@ -32,12 +32,21 @@
package com.jme3.opencl;
/**
- * Specifies how a buffer object can be accessed
+ * Specifies how a buffer object can be accessed by the kernel.
* @author Sebastian Weiss
* @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
index 5cfe47a9a..75ffe52ca 100644
--- a/jme3-core/src/main/java/com/jme3/opencl/OpenCLException.java
+++ b/jme3-core/src/main/java/com/jme3/opencl/OpenCLException.java
@@ -32,10 +32,14 @@
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 Sebastian Weiss
*/
public class OpenCLException extends RuntimeException {
+ private static final long serialVersionUID = 8471229972153694848L;
private final int errorCode;
@@ -63,6 +67,9 @@ public class OpenCLException extends RuntimeException {
this.errorCode = errorCode;
}
+ /**
+ * @return the error code
+ */
public int getErrorCode() {
return errorCode;
}
diff --git a/jme3-core/src/main/java/com/jme3/opencl/Platform.java b/jme3-core/src/main/java/com/jme3/opencl/Platform.java
index 0d5b636ac..fae1bf57f 100644
--- a/jme3-core/src/main/java/com/jme3/opencl/Platform.java
+++ b/jme3-core/src/main/java/com/jme3/opencl/Platform.java
@@ -35,26 +35,73 @@ import java.util.Collection;
import java.util.List;
/**
- *
+ * A wrapper for an OpenCL platform. A platform is the highest object in the
+ * object hierarchy, it creates the {@link Device}s which are then used to
+ * create the {@link Context}.
+ * This class is mostly used within {@link PlatformChooser}.
+ *
* @author Sebastian Weiss
*/
public interface Platform {
+ /**
+ * @return the list of available devices for this platform
+ */
List getDevices();
+ /**
+ * @return The profile string
+ */
String getProfile();
+ /**
+ * @return {@code true} if this platform implements the full profile
+ */
boolean isFullProfile();
+ /**
+ * @return {@code true} if this platform implements the embedded profile
+ */
boolean isEmbeddedProfile();
+ /**
+ * @return the version string
+ */
String getVersion();
+ /**
+ * Extracts the major version from the version string
+ * @return the major version
+ */
int getVersionMajor();
+ /**
+ * Extracts the minor version from the version string
+ * @return the minor version
+ */
int getVersionMinor();
+ /**
+ * @return the name of the platform
+ */
String getName();
+ /**
+ * @return the vendor of the platform
+ */
String getVendor();
+ /**
+ * Queries if this platform supports OpenGL interop at all.
+ * This value has also to be tested for every device.
+ * @return {@code true} if OpenGL interop is supported
+ */
boolean hasOpenGLInterop();
+ /**
+ * Queries if the specified extension is available.
+ * This value has to be tested also for every device.
+ * @param extension the extension string
+ * @return {@code true} if this extension is supported by the platform
+ * (however, not all devices might support it as well)
+ */
boolean hasExtension(String extension);
+ /**
+ * @return All available extensions
+ */
Collection getExtensions();
-
}
diff --git a/jme3-core/src/main/java/com/jme3/opencl/Program.java b/jme3-core/src/main/java/com/jme3/opencl/Program.java
index 91f588e79..f9ecc5516 100644
--- a/jme3-core/src/main/java/com/jme3/opencl/Program.java
+++ b/jme3-core/src/main/java/com/jme3/opencl/Program.java
@@ -32,15 +32,48 @@
package com.jme3.opencl;
/**
- *
+ * A wrapper for an OpenCL program. A program is created from kernel source code,
+ * manages the build process and creates the kernels.
+ *
+ * Warning: Creating the same kernel more than one leads to undefined behaviour, + * this is especially important for {@link #createAllKernels() } + * + * @see Context#createProgramFromSourceCode(java.lang.String) + * @see #createKernel(java.lang.String) * @author Sebastian Weiss */ public interface Program extends OpenCLObject { + /** + * Builds this program with the specified argument string. + * Please see the official OpenCL specification for a definition of + * all supported arguments + * @param args the compilation arguments + * @throws KernelCompilationException if the compilation fails + * @see #build() + */ void build(String args) throws KernelCompilationException; + /** + * Builds this program without additional arguments + * @throws KernelCompilationException if the compilation fails + * @see #build(java.lang.String) + */ void build() throws KernelCompilationException; + /** + * Creates the kernel with the specified name. + * @param name the name of the kernel as defined in the source code + * @return the kernel object + * @throws OpenCLException if the kernel was not found or some other + * error occured + */ Kernel createKernel(String name); + + /** + * Creates all available kernels in this program. + * The names of the kernels can then by queried by {@link Kernel#getName() }. + * @return an array of all kernels + */ Kernel[] createAllKernels(); } diff --git a/jme3-core/src/main/java/com/jme3/opencl/WorkSize.java b/jme3-core/src/main/java/com/jme3/opencl/WorkSize.java deleted file mode 100644 index 9a923a42f..000000000 --- a/jme3-core/src/main/java/com/jme3/opencl/WorkSize.java +++ /dev/null @@ -1,111 +0,0 @@ -/* - * 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.Arrays; - -/** - * The work size (global and local) for executing a kernel - * @author Sebastian Weiss - */ -public final class WorkSize { - private int dimension; - private long[] sizes; - - public WorkSize(int dimension, long... sizes) - { - set(dimension, sizes); - } - - public WorkSize() { - this(1, 1, 1, 1); - } - public WorkSize(long size) { - this(1, size, 1, 1); - } - public WorkSize(long width, long height) { - this(2, width, height, 1); - } - 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-examples/src/main/java/jme3test/opencl/HelloOpenCL.java b/jme3-examples/src/main/java/jme3test/opencl/HelloOpenCL.java index e697dc584..f4427bb4f 100644 --- a/jme3-examples/src/main/java/jme3test/opencl/HelloOpenCL.java +++ b/jme3-examples/src/main/java/jme3test/opencl/HelloOpenCL.java @@ -181,7 +181,7 @@ public class HelloOpenCL extends SimpleApplication { int size = 256+128; Buffer buffer = clContext.createBuffer(size*4); float value = 5; - Event event = kernel.Run1(clQueue, new WorkSize(buffer.getSize() / 4), buffer, value); + Event event = kernel.Run1(clQueue, new com.jme3.opencl.Kernel.WorkSize(buffer.getSize() / 4), buffer, value); event.waitForFinished(); //check if filled diff --git a/jme3-examples/src/main/java/jme3test/opencl/TestVertexBufferSharing.java b/jme3-examples/src/main/java/jme3test/opencl/TestVertexBufferSharing.java index 6b251641c..1f175bd2f 100644 --- a/jme3-examples/src/main/java/jme3test/opencl/TestVertexBufferSharing.java +++ b/jme3-examples/src/main/java/jme3test/opencl/TestVertexBufferSharing.java @@ -56,7 +56,7 @@ public class TestVertexBufferSharing extends SimpleApplication { private Geometry geom; private Buffer buffer; private Kernel kernel; - private WorkSize ws; + private com.jme3.opencl.Kernel.WorkSize ws; private float time; public static void main(String[] args){ @@ -123,7 +123,7 @@ public class TestVertexBufferSharing extends SimpleApplication { //bind vertex buffer to OpenCL VertexBuffer vb = geom.getMesh().getBuffer(VertexBuffer.Type.Position); buffer = clContext.bindVertexBuffer(vb, MemoryAccess.READ_WRITE); - ws = new WorkSize(geom.getMesh().getVertexCount()); + ws = new com.jme3.opencl.Kernel.WorkSize(geom.getMesh().getVertexCount()); } private void updateOpenCL(float tpf) { //advect time diff --git a/jme3-examples/src/main/java/jme3test/opencl/TestWriteToTexture.java b/jme3-examples/src/main/java/jme3test/opencl/TestWriteToTexture.java index 2f583ea61..296f2062c 100644 --- a/jme3-examples/src/main/java/jme3test/opencl/TestWriteToTexture.java +++ b/jme3-examples/src/main/java/jme3test/opencl/TestWriteToTexture.java @@ -137,7 +137,7 @@ public class TestWriteToTexture extends SimpleApplication implements AnalogListe //no need to wait for the returned event, since the kernel implicitely waits for it (same command queue) //execute kernel - kernel.Run1(clQueue, new WorkSize(settings.getWidth(), settings.getHeight()), texCL, C, 16); + kernel.Run1(clQueue, new com.jme3.opencl.Kernel.WorkSize(settings.getWidth(), settings.getHeight()), texCL, C, 16); //release resource texCL.releaseImageForSharingAsync(clQueue); diff --git a/jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglKernel.java b/jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglKernel.java index 19d536e2a..e1260d09b 100644 --- a/jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglKernel.java +++ b/jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglKernel.java @@ -40,6 +40,7 @@ import java.nio.*; import org.lwjgl.PointerBuffer; import org.lwjgl.opencl.CL10; import org.lwjgl.opencl.CLCommandQueue; +import org.lwjgl.opencl.CLDevice; import org.lwjgl.opencl.CLKernel; /** @@ -69,6 +70,12 @@ public class LwjglKernel extends Kernel { return kernel.getInfoInt(CL10.CL_KERNEL_NUM_ARGS); } + @Override + public long getMaxWorkGroupSize(Device device) { + CLDevice d = ((LwjglDevice) device).getDevice(); + return kernel.getWorkGroupInfoSize(d, CL10.CL_KERNEL_WORK_GROUP_SIZE); + } + @Override public void setArg(int index, LocalMemPerElement t) { int ret = CL10.clSetKernelArg (kernel, index, t.getSize() * workGroupSize.getSizes()[0] * workGroupSize.getSizes()[1] * workGroupSize.getSizes()[2]);