finished documentation

jme3-core/src/main/java/com/jme3/opencl/Buffer.java

@@ -382,5 +382,6 @@ public abstract class Buffer implements OpenCLObject {
      * @return the event object
      */
     public abstract Event releaseBufferForSharingAsync(CommandQueue queue);
+    
     //TODO: add variants of the above two methods that don't create the event object, but release the event immediately
 }

jme3-core/src/main/java/com/jme3/opencl/Kernel.java

@@ -35,13 +35,61 @@ import com.jme3.math.Quaternion;
 import com.jme3.math.Vector2f;
 import com.jme3.math.Vector4f;
 import java.nio.ByteBuffer;
+import java.util.Arrays;
 
 /**
+ * Wrapper for an OpenCL kernel, a piece of executable code on the GPU.
+ * <p>
+ * Terminology:<br>
+ * A Kernel is executed in parallel. In total number of parallel threads, 
+ * called work items, are specified by the <i>global work size</i> (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}.
+ * <br>
+ * Not all threads can always be executed in parallel because there simply might
+ * not be enough processor cores.
+ * Therefore, the concept of a <i>work group</i> 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: <i> shared memory</i>
+ * 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.<br>
+ * 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.
+ * 
+ * <p>
+ * There are two ways to launch a kernel:<br>
+ * 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) }.<br>
+ * 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,54 +97,119 @@ 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);
 
     public abstract void setArg(int index, LocalMem 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.<br>
+     * 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);
@@ -182,4 +343,204 @@ 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;
+        }
+    }
+    
 }

jme3-core/src/main/java/com/jme3/opencl/LocalMem.java

@@ -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;
-	}
-	
-}

jme3-core/src/main/java/com/jme3/opencl/LocalMemPerElement.java

@@ -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;
-	}
-	
-}

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
 }

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
 }

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;
 	}

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}.<br>
+ * This class is mostly used within {@link PlatformChooser}.
  * 
  * @author Sebastian Weiss
  */
 public interface Platform {
 
+    /**
+     * @return the list of available devices for this platform
+     */
     List<? extends Device> 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<? extends String> getExtensions();
     
-    
 }

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.
+ * <p>
+ * 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();
     
 }

jme3-core/src/main/java/com/jme3/opencl/WorkSize.java

@@ -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;
-	}
-	
-	
-}

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

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

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);

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]);