renamed package com.mbien.* in com.jogamp.* JOCL is now officially a JogAmp team player ;).

36 files changed, 7914 insertions, 0 deletions

diff --git a/src/com/jogamp/opencl/BuildProgramCallback.java b/src/com/jogamp/opencl/BuildProgramCallback.java
new file mode 100644
index 00000000..fe2b29b9
--- /dev/null
+++ b/src/com/jogamp/opencl/BuildProgramCallback.java
@@ -0,0 +1,10 @@
+package com.jogamp.opencl;
+
+/**
+ * @author Michael Bien
+ */
+// TODO implement callbacks
+public interface BuildProgramCallback {
+    
+    public void buildProgramCallback(long cl_program, Object user_data);
+}
diff --git a/src/com/jogamp/opencl/CLBuffer.java b/src/com/jogamp/opencl/CLBuffer.java
new file mode 100644
index 00000000..5a2a37e7
--- /dev/null
+++ b/src/com/jogamp/opencl/CLBuffer.java
@@ -0,0 +1,60 @@
+package com.jogamp.opencl;
+
+import java.nio.Buffer;
+
+import static com.jogamp.opencl.CLException.*;
+
+/**
+ *
+ * @author Michael Bien
+ */
+public class CLBuffer<B extends Buffer> extends CLMemory<B> {
+
+    protected CLBuffer(CLContext context, long id, int flags) {
+        super(context, id, flags);
+    }
+
+    protected CLBuffer(CLContext context, B directBuffer, long id, int flags) {
+        super(context, directBuffer, id, flags);
+    }
+
+    @SuppressWarnings("unchecked")
+    static CLBuffer<?> create(CLContext context, int size, int flags) {
+
+        CL cl = context.cl;
+        int[] result = new int[1];
+
+        if(isHostPointerFlag(flags)) {
+            throw new IllegalArgumentException("no host pointer defined");
+        }
+
+        long id = cl.clCreateBuffer(context.ID, flags, size, null, result, 0);
+        checkForError(result[0], "can not create cl buffer");
+
+        return new CLBuffer(context, id, flags);
+    }
+
+    static <B extends Buffer> CLBuffer<B> create(CLContext context, B directBuffer, int flags) {
+
+        if(!directBuffer.isDirect())
+            throw new IllegalArgumentException("buffer is not direct");
+
+        B host_ptr = null;
+        CL cl = context.cl;
+        int[] result = new int[1];
+
+        if(isHostPointerFlag(flags)) {
+            host_ptr = directBuffer;
+        }
+        long id = cl.clCreateBuffer(context.ID, flags, sizeOfBufferElem(directBuffer)*directBuffer.capacity(), host_ptr, result, 0);
+        checkForError(result[0], "can not create cl buffer");
+        
+        return new CLBuffer<B>(context, directBuffer, id, flags);
+    }
+    
+    @Override
+    public <T extends Buffer> CLBuffer<T> cloneWith(T directBuffer) {
+        return new CLBuffer<T>(context, directBuffer, ID, FLAGS);
+    }
+
+}
diff --git a/src/com/jogamp/opencl/CLCommandQueue.java b/src/com/jogamp/opencl/CLCommandQueue.java
new file mode 100644
index 00000000..1781a8ae
--- /dev/null
+++ b/src/com/jogamp/opencl/CLCommandQueue.java
@@ -0,0 +1,1585 @@
+package com.jogamp.opencl;
+
+import com.jogamp.common.nio.Int64Buffer;
+import com.jogamp.opencl.gl.CLGLI;
+import com.jogamp.common.nio.PointerBuffer;
+import java.nio.ByteBuffer;
+import java.nio.IntBuffer;
+import java.util.ArrayList;
+import java.util.Arrays;
+import java.util.EnumSet;
+import java.util.List;
+
+import static com.jogamp.opencl.CLException.*;
+import static com.jogamp.opencl.CL.*;
+import static com.jogamp.opencl.util.CLUtil.*;
+
+/**
+ * The command queue is used to queue a set of operations for a specific {@link CLDevice}.
+ * Having multiple command-queues allows applications to queue multiple independent commands without
+ * requiring synchronization. Note that this should work as long as these objects are
+ * not being shared.<br/>
+ * Sharing of objects across multiple queues or using a CLCommandQueue
+ * form multiple Threads will require the application to perform appropriate synchronization.
+ * @see CLDevice#createCommandQueue(com.jogamp.opencl.CLCommandQueue.Mode...)
+ * @author Michael Bien
+ */
+public class CLCommandQueue extends CLObject implements CLResource {
+
+    private final CLDevice device;
+    private long properties;
+
+    /*
+     * Those direct memory buffers are used to move data between the JVM and OpenCL.
+     */
+    private final PointerBuffer pbA;
+    private final Int64Buffer ibA;
+    private final Int64Buffer ibB;
+    private final Int64Buffer ibC;
+
+    private CLCommandQueue(CLContext context, long id, CLDevice device, long properties) {
+        super(context, id);
+
+        this.device = device;
+        this.properties = properties;
+
+        this.ibA = Int64Buffer.allocateDirect(3);
+        this.ibB = Int64Buffer.allocateDirect(3);
+        this.ibC = Int64Buffer.allocateDirect(3);
+
+        this.pbA = PointerBuffer.allocateDirect(1);
+
+    }
+
+    static CLCommandQueue create(CLContext context, CLDevice device, long properties) {
+        int[] status = new int[1];
+        long id = context.cl.clCreateCommandQueue(context.ID, device.ID, properties, status, 0);
+
+        if(status[0] != CL_SUCCESS) {
+            throw newException(status[0], "can not create command queue on " + device +" with properties: " + Mode.valuesOf(properties));
+        }
+
+        return new CLCommandQueue(context, id, device, properties);
+    }
+
+    /**
+     * Calls {@native clEnqueueWriteBuffer}.
+     */
+    public CLCommandQueue putWriteBuffer(CLBuffer<?> writeBuffer, boolean blockingRead) {
+        return putWriteBuffer(writeBuffer, blockingRead, null, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueWriteBuffer}.
+     */
+    public CLCommandQueue putWriteBuffer(CLBuffer<?> writeBuffer, boolean blockingRead, CLEventList events) {
+        return putWriteBuffer(writeBuffer, blockingRead, null, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueWriteBuffer}.
+     */
+    public CLCommandQueue putWriteBuffer(CLBuffer<?> writeBuffer, boolean blockingWrite, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        int ret = cl.clEnqueueWriteBuffer(
+                ID, writeBuffer.ID, clBoolean(blockingWrite),
+                0, writeBuffer.getSize(), writeBuffer.buffer,
+                conditions, conditionIDs, events==null ? null : events.IDs);
+
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not enqueue WriteBuffer: " + writeBuffer + "with " + toStr(condition, events));
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueReadBuffer}.
+     */
+    public CLCommandQueue putReadBuffer(CLBuffer<?> readBuffer, boolean blockingRead) {
+        putReadBuffer(readBuffer, blockingRead, null, null);
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueReadBuffer}.
+     */
+    public CLCommandQueue putReadBuffer(CLBuffer<?> readBuffer, boolean blockingRead, CLEventList events) {
+        putReadBuffer(readBuffer, blockingRead, null, events);
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueReadBuffer}.
+     */
+    public CLCommandQueue putReadBuffer(CLBuffer<?> readBuffer, boolean blockingRead, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        int ret = cl.clEnqueueReadBuffer(
+                ID, readBuffer.ID, clBoolean(blockingRead),
+                0, readBuffer.getSize(), readBuffer.buffer,
+                conditions, conditionIDs, events==null ? null : events.IDs);
+
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not enqueue ReadBuffer: " + readBuffer + "with " + toStr(condition, events));
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueCopyBuffer}.
+     */
+    public CLCommandQueue putCopyBuffer(CLBuffer<?> src, CLBuffer<?> dest) {
+        return putCopyBuffer(src, dest, 0, 0, src.getCLSize(), null, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueCopyBuffer}.
+     */
+    public CLCommandQueue putCopyBuffer(CLBuffer<?> src, CLBuffer<?> dest, long bytesToCopy) {
+        return putCopyBuffer(src, dest, 0, 0, bytesToCopy, null, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueCopyBuffer}.
+     */
+    public CLCommandQueue putCopyBuffer(CLBuffer<?> src, CLBuffer<?> dest, int srcOffset, int destOffset, long bytesToCopy, CLEventList events) {
+        return putCopyBuffer(src, dest, 0, 0, bytesToCopy, null, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueCopyBuffer}.
+     */
+    public CLCommandQueue putCopyBuffer(CLBuffer<?> src, CLBuffer<?> dest, int srcOffset, int destOffset, long bytesToCopy, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        int ret = cl.clEnqueueCopyBuffer(
+                        ID, src.ID, dest.ID, srcOffset, destOffset, bytesToCopy,
+                        conditions, conditionIDs, events==null ? null : events.IDs);
+
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not copy Buffer from " + src + " to " + dest + " with srcOffset: "+ srcOffset
+                    + " dstOffset: " + destOffset + " bytesToCopy: " + bytesToCopy + toStr(condition, events));
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+
+        return this;
+    }
+
+    //2D
+    /**
+     * Calls {@native clEnqueueWriteImage}.
+     */
+    public CLCommandQueue putWriteImage(CLImage2d<?> writeImage, boolean blockingWrite) {
+        return putWriteImage(writeImage, 0, 0, 0, writeImage.width, writeImage.height, blockingWrite, null, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueWriteImage}.
+     */
+    public CLCommandQueue putWriteImage(CLImage2d<?> writeImage, boolean blockingWrite, CLEventList events) {
+        return putWriteImage(writeImage, 0, 0, 0, writeImage.width, writeImage.height, blockingWrite, null, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueWriteImage}.
+     */
+    public CLCommandQueue putWriteImage(CLImage2d<?> writeImage, boolean blockingWrite, CLEventList condition, CLEventList events) {
+        return putWriteImage(writeImage, 0, 0, 0, writeImage.width, writeImage.height, blockingWrite, condition, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueWriteImage}.
+     */
+    public CLCommandQueue putWriteImage(CLImage2d<?> writeImage, int inputRowPitch,
+            int originX, int originY, int rangeX, int rangeY, boolean blockingWrite) {
+        return putWriteImage(writeImage, inputRowPitch, originX, originY, rangeX, rangeY, blockingWrite, null, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueWriteImage}.
+     */
+    public CLCommandQueue putWriteImage(CLImage2d<?> writeImage, int inputRowPitch,
+            int originX, int originY, int rangeX, int rangeY, boolean blockingWrite, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        // spec: CL_INVALID_VALUE if image is a 2D image object and origin[2] is not equal to 0
+        // or region[2] is not equal to 1 or slice_pitch is not equal to 0.
+        copy2NIO(ibA, originX, originY, 0);
+        copy2NIO(ibB, rangeX, rangeY, 1);
+
+        int ret = cl.clEnqueueWriteImage(ID, writeImage.ID, clBoolean(blockingWrite),
+                                         ibA, ibB, inputRowPitch, 0, writeImage.buffer,
+                                         conditions, conditionIDs, events==null ? null : events.IDs);
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not enqueue WriteImage " + writeImage + " with inputRowPitch: " + inputRowPitch
+                       + " origin: " + toStr(originX, originY)+ " range: " + toStr(rangeX, rangeY) + toStr(condition, events));
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+        return this;
+    }
+
+    //3D
+    /**
+     * Calls {@native clEnqueueWriteImage}.
+     */
+    public CLCommandQueue putWriteImage(CLImage3d<?> writeImage, boolean blockingWrite) {
+        return putWriteImage(writeImage, 0, 0, 0, 0, 0, writeImage.width, writeImage.height, writeImage.depth, blockingWrite, null, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueWriteImage}.
+     */
+    public CLCommandQueue putWriteImage(CLImage3d<?> writeImage, boolean blockingWrite, CLEventList events) {
+        return putWriteImage(writeImage, 0, 0, 0, 0, 0, writeImage.width, writeImage.height, writeImage.depth, blockingWrite, null, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueWriteImage}.
+     */
+    public CLCommandQueue putWriteImage(CLImage3d<?> writeImage, boolean blockingWrite, CLEventList condition, CLEventList events) {
+        return putWriteImage(writeImage, 0, 0, 0, 0, 0, writeImage.width, writeImage.height, writeImage.depth, blockingWrite, condition, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueWriteImage}.
+     */
+    public CLCommandQueue putWriteImage(CLImage3d<?> writeImage, int inputRowPitch, int inputSlicePitch,
+            int originX, int originY, int originZ, int rangeX, int rangeY, int rangeZ, boolean blockingWrite) {
+        return putWriteImage(writeImage, inputRowPitch, inputSlicePitch, originX, originY, originZ, rangeX, rangeY, rangeZ, blockingWrite, null, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueWriteImage}.
+     */
+    public CLCommandQueue putWriteImage(CLImage3d<?> writeImage, int inputRowPitch, int inputSlicePitch,
+            int originX, int originY, int originZ, int rangeX, int rangeY, int rangeZ, boolean blockingWrite, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        copy2NIO(ibA, originX, originY, originZ);
+        copy2NIO(ibB, rangeX, rangeY, rangeZ);
+
+        int ret = cl.clEnqueueWriteImage(ID, writeImage.ID, clBoolean(blockingWrite),
+                                         ibA, ibB, inputRowPitch, inputSlicePitch, writeImage.buffer,
+                                         conditions, conditionIDs, events==null ? null : events.IDs);
+
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not enqueue WriteImage " + writeImage + " with inputRowPitch: " + inputRowPitch + " inputSlicePitch: " + inputSlicePitch
+                       + " origin: " + toStr(originX, originY, originZ)+ " range: " + toStr(rangeX, rangeY, rangeZ) + toStr(condition, events));
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+        return this;
+    }
+
+    //2D
+    /**
+     * Calls {@native clEnqueueReadImage}.
+     */
+    public CLCommandQueue putReadImage(CLImage2d<?> readImage, boolean blockingRead) {
+        return putReadImage(readImage, 0, 0, 0, readImage.width, readImage.height, blockingRead, null, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueReadImage}.
+     */
+    public CLCommandQueue putReadImage(CLImage2d<?> readImage, boolean blockingRead, CLEventList events) {
+        return putReadImage(readImage, 0, 0, 0, readImage.width, readImage.height, blockingRead, null, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueReadImage}.
+     */
+    public CLCommandQueue putReadImage(CLImage2d<?> readImage, boolean blockingRead, CLEventList condition, CLEventList events) {
+        return putReadImage(readImage, 0, 0, 0, readImage.width, readImage.height, blockingRead, condition, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueReadImage}.
+     */
+    public CLCommandQueue putReadImage(CLImage2d<?> readImage, int inputRowPitch,
+            int originX, int originY, int rangeX, int rangeY, boolean blockingRead) {
+        return putReadImage(readImage, inputRowPitch, originX, originY, rangeX, rangeY, blockingRead, null, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueReadImage}.
+     */
+    public CLCommandQueue putReadImage(CLImage2d<?> readImage, int inputRowPitch,
+            int originX, int originY, int rangeX, int rangeY, boolean blockingRead, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        // spec: CL_INVALID_VALUE if image is a 2D image object and origin[2] is not equal to 0
+        // or region[2] is not equal to 1 or slice_pitch is not equal to 0.
+        copy2NIO(ibA, originX, originY, 0);
+        copy2NIO(ibB, rangeX, rangeY, 1);
+
+        int ret = cl.clEnqueueReadImage(ID, readImage.ID, clBoolean(blockingRead),
+                                         ibA, ibB, inputRowPitch, 0, readImage.buffer,
+                                         conditions, conditionIDs, events==null ? null : events.IDs);
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not enqueue ReadImage " + readImage + " with inputRowPitch: " + inputRowPitch
+                       + " origin: " + toStr(originX, originY)+ " range: " + toStr(rangeX, rangeY) + toStr(condition, events));
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+        return this;
+    }
+
+    //3D
+    /**
+     * Calls {@native clEnqueueReadImage}.
+     */
+    public CLCommandQueue putReadImage(CLImage3d<?> readImage, boolean blockingRead) {
+        return putReadImage(readImage, 0, 0, 0, 0, 0, readImage.width, readImage.height, readImage.depth, blockingRead, null, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueReadImage}.
+     */
+    public CLCommandQueue putReadImage(CLImage3d<?> readImage, boolean blockingRead, CLEventList events) {
+        return putReadImage(readImage, 0, 0, 0, 0, 0, readImage.width, readImage.height, readImage.depth, blockingRead, null, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueReadImage}.
+     */
+    public CLCommandQueue putReadImage(CLImage3d<?> readImage, boolean blockingRead, CLEventList condition, CLEventList events) {
+        return putReadImage(readImage, 0, 0, 0, 0, 0, readImage.width, readImage.height, readImage.depth, blockingRead, condition, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueReadImage}.
+     */
+    public CLCommandQueue putReadImage(CLImage3d<?> readImage, int inputRowPitch, int inputSlicePitch,
+            int originX, int originY, int originZ, int rangeX, int rangeY, int rangeZ, boolean blockingRead) {
+        return putReadImage(readImage, inputRowPitch, inputSlicePitch, originX, originY, originZ, rangeX, rangeY, rangeZ, blockingRead, null, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueReadImage}.
+     */
+    public CLCommandQueue putReadImage(CLImage3d<?> readImage, int inputRowPitch, int inputSlicePitch,
+            int originX, int originY, int originZ, int rangeX, int rangeY, int rangeZ, boolean blockingRead, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        copy2NIO(ibA, originX, originY, originZ);
+        copy2NIO(ibB, rangeX, rangeY, rangeZ);
+
+        int ret = cl.clEnqueueReadImage(ID, readImage.ID, clBoolean(blockingRead),
+                                        ibA, ibB, inputRowPitch, inputSlicePitch, readImage.buffer,
+                                        conditions, conditionIDs, events==null ? null : events.IDs);
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not enqueue ReadImage " + readImage + " with inputRowPitch: " + inputRowPitch + " inputSlicePitch: " + inputSlicePitch
+                       + " origin: " + toStr(originX, originY, originZ)+ " range: " + toStr(rangeX, rangeY, rangeZ) + toStr(condition, events));
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+        return this;
+    }
+
+    //2D
+    /**
+     * Calls {@native clEnqueueCopyImage}.
+     */
+    public CLCommandQueue putCopyImage(CLImage2d<?> srcImage, CLImage2d<?> dstImage) {
+        return putCopyImage(srcImage, dstImage, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueCopyImage}.
+     */
+    public CLCommandQueue putCopyImage(CLImage2d<?> srcImage, CLImage2d<?> dstImage, CLEventList events) {
+        return putCopyImage(srcImage, dstImage, 0, 0, 0, 0, srcImage.width, srcImage.height, null, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueCopyImage}.
+     */
+    public CLCommandQueue putCopyImage(CLImage2d<?> srcImage, CLImage2d<?> dstImage, CLEventList condition, CLEventList events) {
+        return putCopyImage(srcImage, dstImage, 0, 0, 0, 0, srcImage.width, srcImage.height, condition, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueCopyImage}.
+     */
+    public CLCommandQueue putCopyImage(CLImage2d<?> srcImage, CLImage2d<?> dstImage,
+                                        int srcOriginX, int srcOriginY,
+                                        int dstOriginX, int dstOriginY,
+                                        int rangeX, int rangeY) {
+        return putCopyImage(srcImage, dstImage, srcOriginX, srcOriginY, dstOriginX, dstOriginY, rangeX, rangeY, null, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueCopyImage}.
+     */
+    public CLCommandQueue putCopyImage(CLImage2d<?> srcImage, CLImage2d<?> dstImage,
+                                        int srcOriginX, int srcOriginY,
+                                        int dstOriginX, int dstOriginY,
+                                        int rangeX, int rangeY, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        //spec: CL_INVALID_VALUE if src_image is a 2D image object and origin[2] or dst_origin[2] is not equal to 0
+        // or region[2] is not equal to 1.
+        copy2NIO(ibA, srcOriginX, srcOriginY, 0);
+        copy2NIO(ibB, dstOriginX, dstOriginY, 0);
+        copy2NIO(ibC, rangeX, rangeY, 1);
+
+        int ret = cl.clEnqueueCopyImage(ID, srcImage.ID, dstImage.ID, ibA, ibB, ibC,
+                                         conditions, conditionIDs, events==null ? null : events.IDs);
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not CopyImage " + srcImage +" to "+ dstImage
+                    + " with srcOrigin: " + toStr(srcOriginX, srcOriginY) + " dstOrigin: " + toStr(dstOriginX, dstOriginY)
+                    + " range:  " + toStr(rangeX, rangeY) + toStr(condition, events));
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+        return this;
+    }
+
+    //3D
+    /**
+     * Calls {@native clEnqueueCopyImage}.
+     */
+    public CLCommandQueue putCopyImage(CLImage3d<?> srcImage, CLImage3d<?> dstImage) {
+        return putCopyImage(srcImage, dstImage, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueCopyImage}.
+     */
+    public CLCommandQueue putCopyImage(CLImage3d<?> srcImage, CLImage3d<?> dstImage, CLEventList events) {
+        return putCopyImage(srcImage, dstImage, 0, 0, 0, 0, 0, 0, srcImage.width, srcImage.height, srcImage.depth, null, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueCopyImage}.
+     */
+    public CLCommandQueue putCopyImage(CLImage3d<?> srcImage, CLImage3d<?> dstImage, CLEventList condition, CLEventList events) {
+        return putCopyImage(srcImage, dstImage, 0, 0, 0, 0, 0, 0, srcImage.width, srcImage.height, srcImage.depth, condition, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueCopyImage}.
+     */
+    public CLCommandQueue putCopyImage(CLImage3d<?> srcImage, CLImage3d<?> dstImage,
+                                        int srcOriginX, int srcOriginY, int srcOriginZ,
+                                        int dstOriginX, int dstOriginY, int dstOriginZ,
+                                        int rangeX, int rangeY, int rangeZ) {
+        return putCopyImage(srcImage, dstImage, srcOriginX, srcOriginY, srcOriginZ,
+                                                dstOriginX, dstOriginY, dstOriginZ,
+                                                rangeX, rangeY, rangeZ, null, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueCopyImage}.
+     */
+    public CLCommandQueue putCopyImage(CLImage3d<?> srcImage, CLImage3d<?> dstImage,
+                                        int srcOriginX, int srcOriginY, int srcOriginZ,
+                                        int dstOriginX, int dstOriginY, int dstOriginZ,
+                                        int rangeX, int rangeY, int rangeZ, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        copy2NIO(ibA, srcOriginX, srcOriginY, srcOriginZ);
+        copy2NIO(ibB, dstOriginX, dstOriginY, dstOriginZ);
+        copy2NIO(ibC, rangeX, rangeY, rangeZ);
+
+        int ret = cl.clEnqueueCopyImage(ID, srcImage.ID, dstImage.ID, ibA, ibB, ibC,
+                                         conditions, conditionIDs, events==null ? null : events.IDs);
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not CopyImage " + srcImage +" to "+ dstImage
+                    + " with srcOrigin: " + toStr(srcOriginX, srcOriginY, srcOriginZ) + " dstOrigin: " + toStr(dstOriginX, dstOriginY, dstOriginZ)
+                    + " range:  " + toStr(rangeX, rangeY, rangeZ) + toStr(condition, events));
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+        return this;
+    }
+    
+    //2D
+    /**
+     * Calls {@native clEnqueueCopyBufferToImage}.
+     */
+    public CLCommandQueue putCopyBufferToImage(CLBuffer<?> srcBuffer, CLImage2d<?> dstImage) {
+        return putCopyBufferToImage(srcBuffer, dstImage, null);
+    }
+    
+    /**
+     * Calls {@native clEnqueueCopyBufferToImage}.
+     */
+    public CLCommandQueue putCopyBufferToImage(CLBuffer<?> srcBuffer, CLImage2d<?> dstImage, CLEventList events) {
+        return putCopyBufferToImage(srcBuffer, dstImage, 0, 0, 0, dstImage.width, dstImage.height, null, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueCopyBufferToImage}.
+     */
+    public CLCommandQueue putCopyBufferToImage(CLBuffer<?> srcBuffer, CLImage2d<?> dstImage, CLEventList condition, CLEventList events) {
+        return putCopyBufferToImage(srcBuffer, dstImage, 0, 0, 0, dstImage.width, dstImage.height, condition, events);
+    }
+        
+    /**
+     * Calls {@native clEnqueueCopyBufferToImage}.
+     */
+    public CLCommandQueue putCopyBufferToImage(CLBuffer<?> srcBuffer, CLImage2d<?> dstImage,
+                                        long srcOffset, int dstOriginX, int dstOriginY,
+                                        int rangeX, int rangeY) {
+        return putCopyBufferToImage(srcBuffer, dstImage, 
+                srcOffset, dstOriginX, dstOriginY, rangeX, rangeY, null, null);
+    }
+    
+    /**
+     * Calls {@native clEnqueueCopyBufferToImage}.
+     */
+    public CLCommandQueue putCopyBufferToImage(CLBuffer<?> srcBuffer, CLImage2d<?> dstImage,
+                                        long srcOffset, int dstOriginX, int dstOriginY,
+                                        int rangeX, int rangeY, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        // spec: CL_INVALID_VALUE if dst_image is a 2D image object and dst_origin[2] is not equal to 0
+        // or region[2] is not equal to 1.
+        copy2NIO(ibA, dstOriginX, dstOriginY, 0);
+        copy2NIO(ibB, rangeX, rangeY, 1);
+
+        int ret = cl.clEnqueueCopyBufferToImage(ID, srcBuffer.ID, dstImage.ID,
+                                         srcOffset, ibA, ibB,
+                                         conditions, conditionIDs, events==null ? null : events.IDs);
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not copy " + srcBuffer +" to "+ dstImage
+                    + " with srcOffset: " + srcOffset + " dstOrigin: " + toStr(dstOriginX, dstOriginY)
+                    + " range:  " + toStr(rangeX, rangeY) + toStr(condition, events));
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+        return this;
+    }
+    
+    //3D
+    /**
+     * Calls {@native clEnqueueCopyBufferToImage}.
+     */
+    public CLCommandQueue putCopyBufferToImage(CLBuffer<?> srcBuffer, CLImage3d<?> dstImage) {
+        return putCopyBufferToImage(srcBuffer, dstImage, null);
+    }
+    
+    /**
+     * Calls {@native clEnqueueCopyBufferToImage}.
+     */
+    public CLCommandQueue putCopyBufferToImage(CLBuffer<?> srcBuffer, CLImage3d<?> dstImage, CLEventList events) {
+        return putCopyBufferToImage(srcBuffer, dstImage, 0, 0, 0, 0, dstImage.width, dstImage.height, dstImage.depth, null, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueCopyBufferToImage}.
+     */
+    public CLCommandQueue putCopyBufferToImage(CLBuffer<?> srcBuffer, CLImage3d<?> dstImage, CLEventList condition, CLEventList events) {
+        return putCopyBufferToImage(srcBuffer, dstImage, 0, 0, 0, 0, dstImage.width, dstImage.height, dstImage.depth, condition, events);
+    }
+        
+    /**
+     * Calls {@native clEnqueueCopyBufferToImage}.
+     */
+    public CLCommandQueue putCopyBufferToImage(CLBuffer<?> srcBuffer, CLImage3d<?> dstImage,
+                                        long srcOffset, int dstOriginX, int dstOriginY, int dstOriginZ,
+                                        int rangeX, int rangeY, int rangeZ) {
+        return putCopyBufferToImage(srcBuffer, dstImage, 
+                srcOffset, dstOriginX, dstOriginY, dstOriginZ, rangeX, rangeY, rangeZ, null, null);
+        
+    }
+    
+    /**
+     * Calls {@native clEnqueueCopyBufferToImage}.
+     */
+    public CLCommandQueue putCopyBufferToImage(CLBuffer<?> srcBuffer, CLImage3d<?> dstImage,
+                                        long srcOffset, int dstOriginX, int dstOriginY, int dstOriginZ,
+                                        int rangeX, int rangeY, int rangeZ, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        copy2NIO(ibA, dstOriginX, dstOriginY, dstOriginZ);
+        copy2NIO(ibB, rangeX, rangeY, rangeZ);
+
+        int ret = cl.clEnqueueCopyBufferToImage(ID, srcBuffer.ID, dstImage.ID,
+                                         srcOffset, ibA, ibB,
+                                         conditions, conditionIDs, events==null ? null : events.IDs);
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not copy " + srcBuffer +" to "+ dstImage
+                    + " with srcOffset: " + srcOffset + " dstOrigin: " + toStr(dstOriginX, dstOriginY, dstOriginZ)
+                    + " range:  " + toStr(rangeX, rangeY, dstOriginZ) + toStr(condition, events));
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+        return this;
+    }
+
+    //2D
+    /**
+     * Calls {@native clEnqueueCopyImageToBuffer}.
+     */
+    public CLCommandQueue putCopyImageToBuffer(CLImage2d<?> srcImage, CLBuffer<?> dstBuffer) {
+        return putCopyImageToBuffer(srcImage, dstBuffer, null);
+    }
+    
+    /**
+     * Calls {@native clEnqueueCopyImageToBuffer}.
+     */
+    public CLCommandQueue putCopyImageToBuffer(CLImage2d<?> srcImage, CLBuffer<?> dstBuffer, CLEventList events) {
+        return putCopyImageToBuffer(srcImage, dstBuffer, 0, 0, srcImage.width, srcImage.height, 0, null, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueCopyImageToBuffer}.
+     */
+    public CLCommandQueue putCopyImageToBuffer(CLImage2d<?> srcImage, CLBuffer<?> dstBuffer, CLEventList condition, CLEventList events) {
+        return putCopyImageToBuffer(srcImage, dstBuffer, 0, 0, srcImage.width, srcImage.height, 0, condition, events);
+    }
+        
+    /**
+     * Calls {@native clEnqueueCopyImageToBuffer}.
+     */
+    public CLCommandQueue putCopyImageToBuffer(CLImage2d<?> srcImage, CLBuffer<?> dstBuffer,
+                                        int srcOriginX, int srcOriginY,
+                                        int rangeX, int rangeY, long dstOffset) {
+        return putCopyImageToBuffer(srcImage, dstBuffer, 
+                srcOriginX, srcOriginY, rangeX, rangeY, dstOffset, null, null);
+    }
+    
+    /**
+     * Calls {@native clEnqueueCopyImageToBuffer}.
+     */
+    public CLCommandQueue putCopyImageToBuffer(CLImage2d<?> srcImage, CLBuffer<?> dstBuffer,
+                                        int srcOriginX, int srcOriginY,
+                                        int rangeX, int rangeY, long dstOffset, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        // spec: CL_INVALID_VALUE if src_image is a 2D image object and src_origin[2] is not equal to 0
+        // or region[2] is not equal to 1.
+        copy2NIO(ibA, srcOriginX, srcOriginY, 0);
+        copy2NIO(ibB, rangeX, rangeY, 1);
+
+        int ret = cl.clEnqueueCopyImageToBuffer(ID, dstBuffer.ID, srcImage.ID,
+                                         ibA, ibB, dstOffset,
+                                         conditions, conditionIDs, events==null ? null : events.IDs);
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not copy " + srcImage +" to "+ dstBuffer
+                    + " with srcOrigin: " + toStr(srcOriginX, srcOriginY) + " range: " + toStr(rangeX, rangeY)
+                    + " dstOffset: " + dstOffset + toStr(condition, events));
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+        return this;
+    }
+    
+    //3D
+    /**
+     * Calls {@native clEnqueueCopyImageToBuffer}.
+     */
+    public CLCommandQueue putCopyImageToBuffer(CLImage3d<?> srcImage, CLBuffer<?> dstBuffer) {
+        return putCopyImageToBuffer(srcImage, dstBuffer, 0, 0, 0, srcImage.width, srcImage.height, srcImage.depth, 0, null, null);
+    }
+    
+    /**
+     * Calls {@native clEnqueueCopyImageToBuffer}.
+     */
+    public CLCommandQueue putCopyImageToBuffer(CLImage3d<?> srcImage, CLBuffer<?> dstBuffer, CLEventList events) {
+        return putCopyImageToBuffer(srcImage, dstBuffer, 0, 0, 0, srcImage.width, srcImage.height, srcImage.depth, 0, null, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueCopyImageToBuffer}.
+     */
+    public CLCommandQueue putCopyImageToBuffer(CLImage3d<?> srcImage, CLBuffer<?> dstBuffer, CLEventList condition, CLEventList events) {
+        return putCopyImageToBuffer(srcImage, dstBuffer, 0, 0, 0, srcImage.width, srcImage.height, srcImage.depth, 0, condition, events);
+    }
+        
+    /**
+     * Calls {@native clEnqueueCopyImageToBuffer}.
+     */
+    public CLCommandQueue putCopyImageToBuffer(CLImage3d<?> srcImage, CLBuffer<?> dstBuffer,
+                                        int srcOriginX, int srcOriginY, int srcOriginZ,
+                                        int rangeX, int rangeY, int rangeZ, long dstOffset) {
+        return putCopyImageToBuffer(srcImage, dstBuffer, 
+                srcOriginX, srcOriginY, srcOriginZ, rangeX, rangeY, rangeZ, dstOffset, null, null);
+        
+    }
+    
+    /**
+     * Calls {@native clEnqueueCopyImageToBuffer}.
+     */
+    public CLCommandQueue putCopyImageToBuffer(CLImage3d<?> srcImage, CLBuffer<?> dstBuffer,
+                                        int srcOriginX, int srcOriginY, int srcOriginZ, 
+                                        int rangeX, int rangeY, int rangeZ, long dstOffset, CLEventList condition, CLEventList events) {
+        
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        copy2NIO(ibA, srcOriginX, srcOriginY, srcOriginZ);
+        copy2NIO(ibB, rangeX, rangeY, rangeZ);
+
+        int ret = cl.clEnqueueCopyImageToBuffer(ID, dstBuffer.ID, srcImage.ID,
+                                         ibA, ibB, dstOffset,
+                                         conditions, conditionIDs, events==null ? null : events.IDs);
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not copy " + srcImage +" to "+ dstBuffer
+                    + " with srcOrigin: " + toStr(srcOriginX, srcOriginY, srcOriginZ) + " range: " + toStr(rangeX, rangeY, rangeZ)
+                    + " dstOffset: " + dstOffset + toStr(condition, events));
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueMapBuffer}.
+     */
+    public ByteBuffer putMapBuffer(CLBuffer<?> buffer, CLMemory.Map flag, boolean blockingMap) {
+        return putMapBuffer(buffer, flag, blockingMap, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueMapBuffer}.
+     */
+    public ByteBuffer putMapBuffer(CLBuffer<?> buffer, CLMemory.Map flag, boolean blockingMap, CLEventList events) {
+        return putMapBuffer(buffer, flag, 0, buffer.getCLSize(), blockingMap, null, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueMapBuffer}.
+     */
+    public ByteBuffer putMapBuffer(CLBuffer<?> buffer, CLMemory.Map flag, boolean blockingMap, CLEventList condition, CLEventList events) {
+        return putMapBuffer(buffer, flag, 0, buffer.getCLSize(), blockingMap, condition, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueMapBuffer}.
+     */
+    public ByteBuffer putMapBuffer(CLBuffer<?> buffer, CLMemory.Map flag, long offset, long length, boolean blockingMap) {
+        return putMapBuffer(buffer, flag, offset, length, blockingMap, null, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueMapBuffer}.
+     */
+    public ByteBuffer putMapBuffer(CLBuffer<?> buffer, CLMemory.Map flag, long offset, long length, boolean blockingMap, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        IntBuffer error = pbA.position(0).getBuffer().asIntBuffer();
+        ByteBuffer mappedBuffer = cl.clEnqueueMapBuffer(ID, buffer.ID, clBoolean(blockingMap),
+                                         flag.FLAGS, offset, length,
+                                         conditions, conditionIDs, events==null ? null : events.IDs, error);
+        if(error.get(0) != CL_SUCCESS) {
+            throw newException(error.get(0), "can not map " + buffer + " with: " + flag
+                    + " offset: " + offset + " lenght: " + length + toStr(condition, events));
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+
+        return mappedBuffer;
+    }
+
+    // 2D
+    /**
+     * Calls {@native clEnqueueMapImage}.
+     */
+    public ByteBuffer putMapImage(CLImage2d<?> image, CLMemory.Map flag, boolean blockingMap) {
+        return putMapImage(image, flag, blockingMap, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueMapImage}.
+     */
+    public ByteBuffer putMapImage(CLImage2d<?> image, CLMemory.Map flag, boolean blockingMap, CLEventList events) {
+        return putMapImage(image, flag, 0, 0, image.width, image.height, blockingMap, null, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueMapImage}.
+     */
+    public ByteBuffer putMapImage(CLImage2d<?> image, CLMemory.Map flag, boolean blockingMap, CLEventList condition, CLEventList events) {
+        return putMapImage(image, flag, 0, 0, image.width, image.height, blockingMap, condition, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueMapImage}.
+     */
+    public ByteBuffer putMapImage(CLImage2d<?> buffer, CLMemory.Map flag, int offsetX, int offsetY,
+                                    int rangeX, int rangeY, boolean blockingMap) {
+        return putMapImage(buffer, flag, offsetX, offsetY, rangeX, rangeY, blockingMap, null, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueMapImage}.
+     */
+    public ByteBuffer putMapImage(CLImage2d<?> image, CLMemory.Map flag,
+                                    int offsetX, int offsetY,
+                                    int rangeX, int rangeY, boolean blockingMap, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        IntBuffer error = pbA.position(0).getBuffer().asIntBuffer();
+
+        // spec: CL_INVALID_VALUE if image is a 2D image object and origin[2] is not equal to 0 or region[2] is not equal to 1
+        copy2NIO(ibB, offsetX, offsetY, 0);
+        copy2NIO(ibC, rangeX, rangeY, 1);
+
+        ByteBuffer mappedImage = cl.clEnqueueMapImage(ID, image.ID, clBoolean(blockingMap),
+                                         flag.FLAGS, ibB, ibC, null, null,
+                                         conditions, conditionIDs, events==null ? null : events.IDs, error);
+        if(error.get(0) != CL_SUCCESS) {
+            throw newException(error.get(0), "can not map " + image + " with: " + flag
+                    + " offset: " + toStr(offsetX, offsetY) + " range: " + toStr(rangeX, rangeY) + toStr(condition, events));
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+
+        return mappedImage;
+    }
+
+    // 3D
+    /**
+     * Calls {@native clEnqueueMapImage}.
+     */
+    public ByteBuffer putMapImage(CLImage3d<?> image, CLMemory.Map flag, boolean blockingMap) {
+        return putMapImage(image, flag, blockingMap, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueMapImage}.
+     */
+    public ByteBuffer putMapImage(CLImage3d<?> image, CLMemory.Map flag, boolean blockingMap, CLEventList events) {
+        return putMapImage(image, flag, 0, 0, 0, image.width, image.height, image.depth, blockingMap, null, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueMapImage}.
+     */
+    public ByteBuffer putMapImage(CLImage3d<?> image, CLMemory.Map flag, boolean blockingMap, CLEventList condition, CLEventList events) {
+        return putMapImage(image, flag, 0, 0, 0, image.width, image.height, image.depth, blockingMap, condition, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueMapImage}.
+     */
+    public ByteBuffer putMapImage(CLImage3d<?> image, CLMemory.Map flag,
+                                    int offsetX, int offsetY, int offsetZ,
+                                    int rangeX, int rangeY, int rangeZ, boolean blockingMap) {
+        return putMapImage(image, flag, offsetX, offsetY, offsetZ, rangeX, rangeY, rangeZ, blockingMap, null, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueMapImage}.
+     */
+    public ByteBuffer putMapImage(CLImage3d<?> image, CLMemory.Map flag,
+                                    int offsetX, int offsetY, int offsetZ,
+                                    int rangeX, int rangeY, int rangeZ, boolean blockingMap, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        IntBuffer error = pbA.position(0).getBuffer().asIntBuffer();
+        copy2NIO(ibB, offsetX, offsetY, offsetZ);
+        copy2NIO(ibC, rangeX, rangeY, rangeZ);
+        ByteBuffer mappedImage = cl.clEnqueueMapImage(ID, image.ID, clBoolean(blockingMap),
+                                         flag.FLAGS, ibB, ibC, null, null,
+                                         conditions, conditionIDs, events==null ? null : events.IDs, error);
+        if(error.get(0) != CL_SUCCESS) {
+            throw newException(error.get(0), "can not map " + image + " with: " + flag
+                    + " offset: " + toStr(offsetX, offsetY, offsetZ) + " range: " + toStr(rangeX, rangeY, rangeZ) + toStr(condition, events));
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+
+        return mappedImage;
+    }
+
+    /**
+     * Calls {@native clEnqueueUnmapMemObject}.
+     */
+    public CLCommandQueue putUnmapMemory(CLMemory<?> memory) {
+        return putUnmapMemory(memory, null, null);
+    }
+
+    /**
+     * Calls {@native clEnqueueUnmapMemObject}.
+     */
+    public CLCommandQueue putUnmapMemory(CLMemory<?> memory, CLEventList events) {
+        return putUnmapMemory(memory, null, events);
+    }
+
+    /**
+     * Calls {@native clEnqueueUnmapMemObject}.
+     */
+    public CLCommandQueue putUnmapMemory(CLMemory<?> memory, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        int ret = cl.clEnqueueUnmapMemObject(ID, memory.ID, memory.getBuffer(),
+                                        conditions, conditionIDs, events==null ? null : events.IDs);
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not unmap " + memory + toStr(condition, events));
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueMarker}.
+     */
+    public CLCommandQueue putMarker(CLEventList events) {
+        int ret = cl.clEnqueueMarker(ID, events.IDs);
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not enqueue marker " + events);
+        }
+        return this;
+    }
+
+    /**
+     * Calls {@native clWaitForEvents} if blockingWait equals true otherwise {@native clEnqueueWaitForEvents}.
+     */
+    public CLCommandQueue putWaitForEvent(CLEventList list, int index, boolean blockingWait) {
+        int marker = list.IDs.position()-1;
+        list.IDs.position(index);
+        int ret = blockingWait ? cl.clWaitForEvents(1, list.IDs)
+                               : cl.clEnqueueWaitForEvents(ID, 1, list.IDs);
+        list.IDs.position(marker);
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not "+ (blockingWait?"blocking": "") +" wait for event #" + index+ " in "+list);
+        }
+        return this;
+    }
+
+    /**
+     * Calls {@native clWaitForEvents} if blockingWait equals true otherwise {@native clEnqueueWaitForEvents}.
+     */
+    public CLCommandQueue putWaitForEvents(CLEventList list, boolean blockingWait) {
+        list.IDs.rewind();
+        int ret = blockingWait ? cl.clWaitForEvents(list.size, list.IDs)
+                               : cl.clEnqueueWaitForEvents(ID, list.size, list.IDs);
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not "+ (blockingWait?"blocking": "") +" wait for events " + list);
+        }
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueBarrier}.
+     */
+    public CLCommandQueue putBarrier() {
+        int ret = cl.clEnqueueBarrier(ID);
+        checkForError(ret, "can not enqueue Barrier");
+        return this;
+    }
+
+    /**
+     * Equivalent to calling
+     * {@link #put1DRangeKernel(CLKernel kernel, long globalWorkOffset, long globalWorkSize, long localWorkSize)}
+     * with globalWorkOffset = null, globalWorkSize set to 1, and localWorkSize set to 1.
+     * <p>Calls {@native clEnqueueTask}.</p>
+     */
+    public CLCommandQueue putTask(CLKernel kernel) {
+        putTask(kernel, null, null);
+        return this;
+    }
+
+    /**
+     * <p>Calls {@native clEnqueueTask}.</p>
+     * @see #putTask(com.jogamp.opencl.CLKernel)
+     */
+    public CLCommandQueue putTask(CLKernel kernel, CLEventList events) {
+        putTask(kernel, null, events);
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueTask}.
+     * @see #putTask(com.jogamp.opencl.CLKernel)
+     */
+    public CLCommandQueue putTask(CLKernel kernel, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        int ret = cl.clEnqueueTask(ID, kernel.ID, conditions, conditionIDs, events==null ? null : events.IDs);
+        if(ret != CL_SUCCESS) {
+            checkForError(ret, "can not enqueue Task: " + kernel + toStr(condition, events));
+        }
+        if(events != null) {
+            events.createEvent(context);
+        }
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueNDRangeKernel}.
+     */
+    public CLCommandQueue put1DRangeKernel(CLKernel kernel, long globalWorkOffset, long globalWorkSize, long localWorkSize) {
+        this.put1DRangeKernel(kernel, globalWorkOffset, globalWorkSize, localWorkSize, null, null);
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueNDRangeKernel}.
+     */
+    public CLCommandQueue put1DRangeKernel(CLKernel kernel, long globalWorkOffset, long globalWorkSize, long localWorkSize, CLEventList events) {
+        this.put1DRangeKernel(kernel, globalWorkOffset, globalWorkSize, localWorkSize, null, events);
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueNDRangeKernel}.
+     */
+    public CLCommandQueue put1DRangeKernel(CLKernel kernel, long globalWorkOffset, long globalWorkSize, long localWorkSize, CLEventList condition, CLEventList events) {
+        Int64Buffer globWO = null;
+        Int64Buffer globWS = null;
+        Int64Buffer locWS = null;
+
+        if(globalWorkOffset != 0) {
+            globWO = copy2NIO(ibA, globalWorkOffset);
+        }
+        if(globalWorkSize != 0) {
+            globWS = copy2NIO(ibB, globalWorkSize);
+        }
+        if(localWorkSize != 0) {
+            locWS = copy2NIO(ibC, localWorkSize);
+        }
+
+        this.putNDRangeKernel(kernel, 1, globWO, globWS, locWS, condition, events);
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueNDRangeKernel}.
+     */
+    public CLCommandQueue put2DRangeKernel(CLKernel kernel, long globalWorkOffsetX, long globalWorkOffsetY,
+                                                            long globalWorkSizeX, long globalWorkSizeY,
+                                                            long localWorkSizeX, long localWorkSizeY) {
+        this.put2DRangeKernel(kernel,
+                globalWorkOffsetX, globalWorkOffsetY,
+                globalWorkSizeX, globalWorkSizeY,
+                localWorkSizeX, localWorkSizeY, null, null);
+
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueNDRangeKernel}.
+     */
+    public CLCommandQueue put2DRangeKernel(CLKernel kernel, long globalWorkOffsetX, long globalWorkOffsetY,
+                                                            long globalWorkSizeX, long globalWorkSizeY,
+                                                            long localWorkSizeX, long localWorkSizeY, CLEventList events) {
+        this.put2DRangeKernel(kernel,
+                globalWorkOffsetX, globalWorkOffsetY,
+                globalWorkSizeX, globalWorkSizeY,
+                localWorkSizeX, localWorkSizeY, null, events);
+        return this;
+    }
+    
+    /**
+     * Calls {@native clEnqueueNDRangeKernel}.
+     */
+    public CLCommandQueue put2DRangeKernel(CLKernel kernel, long globalWorkOffsetX, long globalWorkOffsetY,
+                                                            long globalWorkSizeX, long globalWorkSizeY,
+                                                            long localWorkSizeX, long localWorkSizeY, CLEventList condition, CLEventList events) {
+        Int64Buffer globalWorkOffset = null;
+        Int64Buffer globalWorkSize = null;
+        Int64Buffer localWorkSize = null;
+
+        if(globalWorkOffsetX != 0 && globalWorkOffsetY != 0) {
+            globalWorkOffset = copy2NIO(ibA, globalWorkOffsetX, globalWorkOffsetY);
+        }
+        if(globalWorkSizeX != 0 && globalWorkSizeY != 0) {
+            globalWorkSize = copy2NIO(ibB, globalWorkSizeX, globalWorkSizeY);
+        }
+        if(localWorkSizeX != 0 && localWorkSizeY !=0) {
+            localWorkSize = copy2NIO(ibC, localWorkSizeX, localWorkSizeY);
+        }
+        this.putNDRangeKernel(kernel, 2, globalWorkOffset, globalWorkSize, localWorkSize, condition, events);
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueNDRangeKernel}.
+     */
+    public CLCommandQueue putNDRangeKernel(CLKernel kernel, int workDimension, Int64Buffer globalWorkOffset, Int64Buffer globalWorkSize, Int64Buffer localWorkSize) {
+        this.putNDRangeKernel(kernel, workDimension, globalWorkOffset, globalWorkSize, localWorkSize, null, null);
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueNDRangeKernel}.
+     */
+    public CLCommandQueue putNDRangeKernel(CLKernel kernel, int workDimension, Int64Buffer globalWorkOffset, Int64Buffer globalWorkSize, Int64Buffer localWorkSize, CLEventList events) {
+        this.putNDRangeKernel(kernel, workDimension, globalWorkOffset, globalWorkSize, localWorkSize, null, events);
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueNDRangeKernel}.
+     */
+    public CLCommandQueue putNDRangeKernel(CLKernel kernel, int workDimension, Int64Buffer globalWorkOffset,
+            Int64Buffer globalWorkSize, Int64Buffer localWorkSize, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        int ret = cl.clEnqueueNDRangeKernel(
+                ID, kernel.ID, workDimension,
+                globalWorkOffset,
+                globalWorkSize, 
+                localWorkSize, 
+                conditions, conditionIDs,
+                events==null ? null : events.IDs);
+
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not enqueue NDRangeKernel: " + kernel + " with gWO: "+toStr(globalWorkOffset)
+                            + " gWS: "+toStr(globalWorkSize) + " lWS: " + toStr(localWorkSize) + toStr(condition, events));
+
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueAcquireGLObjects}.
+     */
+    public CLCommandQueue putAcquireGLObject(long glObject) {
+        this.putAcquireGLObject(glObject, null, null);
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueAcquireGLObjects}.
+     */
+    public CLCommandQueue putAcquireGLObject(long glObject, CLEventList events) {
+        this.putAcquireGLObject(glObject, null, events);
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueAcquireGLObjects}.
+     */
+    public CLCommandQueue putAcquireGLObject(long glObject, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        CLGLI xl = (CLGLI) cl;
+
+        PointerBuffer glObj = copy2NIO(pbA, glObject);
+        
+        int ret = xl.clEnqueueAcquireGLObjects(ID, 1, glObj,
+                    conditions, conditionIDs,
+                    events==null ? null : events.IDs);
+
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not aquire GLObject: " + glObject + "with " + toStr(condition, events));
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueReleaseGLObjects}.
+     */
+    public CLCommandQueue putReleaseGLObject(long glObject) {
+        this.putReleaseGLObject(glObject, null);
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueReleaseGLObjects}.
+     */
+    public CLCommandQueue putReleaseGLObject(long glObject, CLEventList events) {
+        this.putReleaseGLObject(glObject, null, events);
+        return this;
+    }
+
+    /**
+     * Calls {@native clEnqueueReleaseGLObjects}.
+     */
+    public CLCommandQueue putReleaseGLObject(long glObject, CLEventList condition, CLEventList events) {
+
+        PointerBuffer conditionIDs = null;
+        int conditions = 0;
+        if(condition != null) {
+            conditionIDs = condition.IDs;
+            conditions   = condition.size;
+        }
+
+        CLGLI xl = (CLGLI) cl;
+
+        PointerBuffer glObj = copy2NIO(pbA, glObject);
+
+        int ret = xl.clEnqueueReleaseGLObjects(ID, 1, glObj,
+                conditions, conditionIDs,
+                events==null ? null : events.IDs);
+
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "can not release GLObject: " + glObject + "with " + toStr(condition, events));
+        }
+
+        if(events != null) {
+            events.createEvent(context);
+        }
+
+        return this;
+    }
+
+    /**
+     * Calls {@native clFinish}.
+     */
+    public CLCommandQueue finish() {
+        int ret = cl.clFinish(ID);
+        checkForError(ret, "can not finish command queue");
+        return this;
+    }
+
+    /**
+     * Calls {@native clFlush}.
+     */
+    public CLCommandQueue flush() {
+        int ret = cl.clFlush(ID);
+        checkForError(ret, "can not flush command queue");
+        return this;
+    }
+
+    /**
+     * Returns true only when {@link Mode#PROFILING_MODE} has been enabled.
+     */
+    public boolean isProfilingEnabled() {
+        return (Mode.PROFILING_MODE.QUEUE_MODE & properties) != 0;
+    }
+
+    /**
+     * Returns true only when {@link Mode#OUT_OF_ORDER_MODE} mode has been enabled.
+     */
+    public boolean isOutOfOrderModeEnabled() {
+        return (Mode.OUT_OF_ORDER_MODE.QUEUE_MODE & properties) != 0;
+    }
+
+    public void release() {
+        int ret = cl.clReleaseCommandQueue(ID);
+        context.onCommandQueueReleased(device, this);
+        checkForError(ret, "can not release command queue");
+    }
+
+    public void close() {
+        release();
+    }
+
+    private static PointerBuffer copy2NIO(PointerBuffer buffer, long a) {
+        return (PointerBuffer) buffer.put(0, a);
+    }
+
+//    private static PointerBuffer copy2NIO(PointerBuffer buffer, long a, long b) {
+//        return buffer.position(1).put(a).put(b).position(1);
+//    }
+//
+//    private static PointerBuffer copy2NIO(PointerBuffer buffer, long a, long b, long c) {
+//        return buffer.rewind().put(a).put(b).put(c).rewind();
+//    }
+
+    private static Int64Buffer copy2NIO(Int64Buffer buffer, long a) {
+        return (Int64Buffer) buffer.put(2, a).position(2);
+    }
+
+    private static Int64Buffer copy2NIO(Int64Buffer buffer, long a, long b) {
+        return (Int64Buffer) ((Int64Buffer)buffer.position(1)).put(a).put(b).position(1);
+    }
+
+    private static Int64Buffer copy2NIO(Int64Buffer buffer, long a, long b, long c) {
+        return (Int64Buffer) ((Int64Buffer)buffer.rewind()).put(a).put(b).put(c).rewind();
+    }
+
+    private static String toStr(Int64Buffer buffer) {
+        if(buffer == null) {
+            return null;
+        }
+        StringBuilder sb = new StringBuilder();
+        sb.append('{');
+        for (int i = buffer.position(); i < buffer.capacity(); i++) {
+            sb.append(buffer.get(i));
+            if(i != buffer.capacity()-1) {
+                sb.append(", ");
+            }
+        }
+        return sb.append('}').toString();
+    }
+
+    private static String toStr(CLEventList condition, CLEventList events) {
+        return " cond.: " + condition +" events: "+events;
+    }
+
+    private String toStr(int... values) {
+        return Arrays.asList(values).toString();
+    }
+
+    /**
+     * Returns the device of this command queue.
+     */
+    public CLDevice getDevice() {
+        return device;
+    }
+
+    /**
+     * Returns the command queue properties as EnumSet.
+     */
+    public EnumSet<Mode> getProperties() {
+        return Mode.valuesOf(properties);
+    }
+
+    /**
+     * Setting properties after a command queue has been created can be implementation specific,
+     * please refer to the specification ({@native clSetCommandQueueProperty}) or vendor documentation.
+     */
+    public void setProperty(Mode property, boolean enabled) {
+        int ret = cl.clSetCommandQueueProperty(ID, property.QUEUE_MODE, clBoolean(enabled), null);
+        if(ret != CL_SUCCESS) {
+            checkForError(ret, "can not set command queue property: " + property);
+        }
+        if(enabled) {
+            properties |=  property.QUEUE_MODE;
+        }else{
+            properties &= ~property.QUEUE_MODE;
+        }
+    }
+
+    @Override
+    public boolean equals(Object obj) {
+        if (obj == null) {
+            return false;
+        }
+        if (getClass() != obj.getClass()) {
+            return false;
+        }
+        final CLCommandQueue other = (CLCommandQueue) obj;
+        if (this.ID != other.ID) {
+            return false;
+        }
+        if (this.context != other.context && (this.context == null || !this.context.equals(other.context))) {
+            return false;
+        }
+        if (this.device != other.device && (this.device == null || !this.device.equals(other.device))) {
+            return false;
+        }
+        return true;
+    }
+
+    @Override
+    public int hashCode() {
+        int hash = 3;
+        hash = 89 * hash + (int) (this.ID ^ (this.ID >>> 32));
+        hash = 89 * hash + (this.context != null ? this.context.hashCode() : 0);
+        hash = 89 * hash + (this.device != null ? this.device.hashCode() : 0);
+        return hash;
+    }
+
+    /**
+     * Enumeration for the command-queue settings.
+     */
+    public enum Mode {
+        /**
+         * If set, the commands in the command-queue are
+         * executed out-of-order. Otherwise, commands are executed in-order.
+         */
+        OUT_OF_ORDER_MODE(CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE),
+
+        /**
+         * Enables profiling of commands in the command-queue.
+         * If set, the profiling of commands is enabled. Otherwise profiling of
+         * commands is disabled. See {@link com.jogamp.opencl.CLEvent} for more information.
+         */
+        PROFILING_MODE(CL_QUEUE_PROFILING_ENABLE);
+
+        /**
+         * Value of wrapped OpenCL device type.
+         */
+        public final int QUEUE_MODE;
+
+        private Mode(int value) {
+            this.QUEUE_MODE = value;
+        }
+
+        public static Mode valueOf(int queueMode) {
+            switch(queueMode) {
+                case(CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE):
+                    return OUT_OF_ORDER_MODE;
+                case(CL_QUEUE_PROFILING_ENABLE):
+                    return PROFILING_MODE;
+            }
+            return null;
+        }
+
+        public static EnumSet<Mode> valuesOf(long bitfield) {
+            List<Mode> matching = new ArrayList<Mode>();
+            Mode[] values = Mode.values();
+            for (Mode value : values) {
+                if((value.QUEUE_MODE & bitfield) != 0)
+                    matching.add(value);
+            }
+            if(matching.isEmpty())
+                return EnumSet.noneOf(Mode.class);
+            else
+                return EnumSet.copyOf(matching);
+        }
+
+    }
+}
diff --git a/src/com/jogamp/opencl/CLContext.java b/src/com/jogamp/opencl/CLContext.java
new file mode 100644
index 00000000..a3157777
--- /dev/null
+++ b/src/com/jogamp/opencl/CLContext.java
@@ -0,0 +1,504 @@
+package com.jogamp.opencl;
+
+import com.jogamp.opencl.CLDevice.Type;
+import com.jogamp.opencl.CLMemory.Mem;
+import com.jogamp.opencl.CLSampler.AddressingMode;
+import com.jogamp.opencl.CLSampler.FilteringMode;
+import com.jogamp.common.nio.Int64Buffer;
+import com.jogamp.common.nio.PointerBuffer;
+import java.io.BufferedReader;
+import java.io.IOException;
+import java.io.InputStream;
+import java.io.InputStreamReader;
+import java.nio.Buffer;
+import java.nio.ByteBuffer;
+import java.nio.ByteOrder;
+import java.nio.DoubleBuffer;
+import java.nio.FloatBuffer;
+import java.nio.IntBuffer;
+import java.nio.LongBuffer;
+import java.nio.ShortBuffer;
+
+import java.util.ArrayList;
+import java.util.Collections;
+import java.util.HashMap;
+import java.util.List;
+import java.util.Map;
+import static com.jogamp.opencl.CLException.*;
+import static com.jogamp.common.nio.Buffers.*;
+import static com.jogamp.common.os.Platform.*;
+
+/**
+ * CLContext is responsible for managing objects such as command-queues, memory,
+ * program and kernel objects and for executing kernels on one or more devices
+ * specified in the context.
+ * @author Michael Bien
+ */
+public class CLContext extends CLObject implements CLResource {
+
+    protected CLDevice[] devices;
+
+    protected final List<CLProgram> programs;
+    protected final List<CLSampler> samplers;
+    protected final List<CLMemory<? extends Buffer>> memoryObjects;
+    
+    protected final Map<CLDevice, List<CLCommandQueue>> queuesMap;
+
+    protected final CLPlatform platform;
+
+
+    protected CLContext(CLPlatform platform, long contextID) {
+        super(CLPlatform.getLowLevelCLInterface(), contextID);
+        this.platform = platform;
+        this.programs = new ArrayList<CLProgram>();
+        this.samplers = new ArrayList<CLSampler>();
+        this.memoryObjects = new ArrayList<CLMemory<? extends Buffer>>();
+        this.queuesMap = new HashMap<CLDevice, List<CLCommandQueue>>();
+    }
+
+    private void initDevices() {
+        
+        if (devices == null) {
+
+            Int64Buffer deviceCount = Int64Buffer.allocateDirect(1);
+
+            int ret = cl.clGetContextInfo(ID, CL.CL_CONTEXT_DEVICES, 0, null, deviceCount);
+            checkForError(ret, "can not enumerate devices");
+
+            ByteBuffer deviceIDs = ByteBuffer.allocateDirect((int)deviceCount.get()).order(ByteOrder.nativeOrder());
+            ret = cl.clGetContextInfo(ID, CL.CL_CONTEXT_DEVICES, deviceIDs.capacity(), deviceIDs, null);
+            checkForError(ret, "can not enumerate devices");
+
+            devices = new CLDevice[deviceIDs.capacity() / (is32Bit() ? 4 : 8)];
+            for (int i = 0; i < devices.length; i++) {
+                devices[i] = new CLDevice(this, is32Bit() ? deviceIDs.getInt() : deviceIDs.getLong());
+            }
+        }
+    }
+
+    /**
+     * Creates a context on all available devices (CL_DEVICE_TYPE_ALL).
+     * The platform to be used is implementation dependent.
+     */
+    public static CLContext create() {
+        return create((CLPlatform)null, Type.ALL);
+    }
+
+    /**
+     * Creates a context on the specified device types.
+     * The platform to be used is implementation dependent.
+     */
+    public static CLContext create(CLDevice.Type... deviceTypes) {
+        return create(null, deviceTypes);
+    }
+
+    /**
+     * Creates a context on the specified devices.
+     * The platform to be used is implementation dependent.
+     */
+    public static CLContext create(CLDevice... devices) {
+        return create(null, devices);
+    }
+
+    /**
+     * Creates a context on the specified platform on all available devices (CL_DEVICE_TYPE_ALL).
+     */
+    public static CLContext create(CLPlatform platform) {
+        return create(platform, CLDevice.Type.ALL);
+    }
+
+    /**
+     * Creates a context on the specified platform and with the specified
+     * device types.
+     */
+    public static CLContext create(CLPlatform platform, CLDevice.Type... deviceTypes) {
+
+        if(platform == null) {
+            platform = CLPlatform.getDefault();
+        }
+
+        long type = toDeviceBitmap(deviceTypes);
+
+        PointerBuffer properties = setupContextProperties(platform);
+        return new CLContext(platform, createContextFromType(properties, type));
+    }
+
+    /**
+     * Creates a context on the specified platform and with the specified
+     * devices.
+     */
+    public static CLContext create(CLPlatform platform, CLDevice... devices) {
+
+        if(platform == null) {
+            platform = CLPlatform.getDefault();
+        }
+
+        PointerBuffer properties = setupContextProperties(platform);
+        CLContext context = new CLContext(platform, createContext(properties, devices));
+        if(devices != null) {
+            for (int i = 0; i < devices.length; i++) {
+                devices[i].setContext(context);
+            }
+        }
+        return context;
+    }
+
+    protected static long createContextFromType(PointerBuffer properties, long deviceType) {
+
+        IntBuffer status = IntBuffer.allocate(1);
+        long context = CLPlatform.getLowLevelCLInterface().clCreateContextFromType(properties, deviceType, null, null, status);
+
+        checkForError(status.get(), "can not create CL context");
+
+        return context;
+    }
+
+    protected static long createContext(PointerBuffer properties, CLDevice... devices) {
+
+        IntBuffer status = newDirectIntBuffer(1);
+        PointerBuffer pb = null;
+        if(devices != null && devices.length != 0) {
+            pb = PointerBuffer.allocateDirect(devices.length);
+            for (int i = 0; i < devices.length; i++) {
+                pb.put(i, devices[i].ID);
+            }
+        }
+        long context = CLPlatform.getLowLevelCLInterface().clCreateContext(properties, pb, null, null, status);
+
+        checkForError(status.get(), "can not create CL context");
+
+        return context;
+    }
+
+    private static PointerBuffer setupContextProperties(CLPlatform platform) {
+
+        if(platform == null) {
+            throw new RuntimeException("no OpenCL installation found");
+        }
+
+        return (PointerBuffer)PointerBuffer.allocateDirect(3).put(CL.CL_CONTEXT_PLATFORM)
+                                              .put(platform.ID).put(0) // 0 terminated array
+                                              .rewind();
+    }
+
+    /**
+     * Creates a program from the given sources, the program is not build yet.
+     */
+    public CLProgram createProgram(String src) {
+        CLProgram program = CLProgram.create(this, src);
+        programs.add(program);
+        return program;
+    }
+
+    /**
+     * Creates a program and reads the source from stream, the program is not build yet.
+     * @throws IOException when a IOException occurred while reading or closing the stream.
+     */
+    public CLProgram createProgram(InputStream source) throws IOException {
+
+        if(source == null)
+            throw new IllegalArgumentException("input stream for program source must not be null");
+
+        BufferedReader reader = new BufferedReader(new InputStreamReader(source));
+        StringBuilder sb = new StringBuilder();
+
+        String line;
+        try {
+            while ((line = reader.readLine()) != null)
+                sb.append(line).append("\n");
+        } finally {
+            source.close();
+        }
+
+        return createProgram(sb.toString());
+    }
+
+    /**
+     * Creates a program from the given binaries, the program is not build yet.
+     * <br/>Creating a program will fail if:<br/>
+     * <ul>
+     * <li>the submitted binaries are invalid or can not be loaded from the OpenCL driver</li>
+     * <li>the binaries do not fit to the CLDevices associated with this context</li>
+     * <li>binaries are missing for one or more CLDevices</li>
+     * </ul>
+     */
+    public CLProgram createProgram(Map<CLDevice, byte[]> binaries) {
+        CLProgram program = CLProgram.create(this, binaries);
+        programs.add(program);
+        return program;
+    }
+
+    /**
+     * Creates a CLBuffer with the specified flags and element count. No flags creates a MEM.READ_WRITE buffer.
+     */
+    public final CLBuffer<ShortBuffer> createShortBuffer(int size, Mem... flags) {
+        return createBuffer(newDirectShortBuffer(size), flags);
+    }
+
+    /**
+     * Creates a CLBuffer with the specified flags and element count. No flags creates a MEM.READ_WRITE buffer.
+     */
+    public final CLBuffer<IntBuffer> createIntBuffer(int size, Mem... flags) {
+        return createBuffer(newDirectIntBuffer(size), flags);
+    }
+
+    /**
+     * Creates a CLBuffer with the specified flags and element count. No flags creates a MEM.READ_WRITE buffer.
+     */
+    public final CLBuffer<LongBuffer> createLongBuffer(int size, Mem... flags) {
+        return createBuffer(newDirectLongBuffer(size), flags);
+    }
+
+    /**
+     * Creates a CLBuffer with the specified flags and element count. No flags creates a MEM.READ_WRITE buffer.
+     */
+    public final CLBuffer<FloatBuffer> createFloatBuffer(int size, Mem... flags) {
+        return createBuffer(newDirectFloatBuffer(size), flags);
+    }
+
+    /**
+     * Creates a CLBuffer with the specified flags and element count. No flags creates a MEM.READ_WRITE buffer.
+     */
+    public final CLBuffer<DoubleBuffer> createDoubleBuffer(int size, Mem... flags) {
+        return createBuffer(newDirectDoubleBuffer(size), flags);
+    }
+
+    /**
+     * Creates a CLBuffer with the specified flags and buffer size in bytes. No flags creates a MEM.READ_WRITE buffer.
+     */
+    public final CLBuffer<ByteBuffer> createByteBuffer(int size, Mem... flags) {
+        return createByteBuffer(size, Mem.flagsToInt(flags));
+    }
+
+    /**
+     * Creates a CLBuffer with the specified flags and buffer size in bytes.
+     */
+    public final CLBuffer<ByteBuffer> createByteBuffer(int size, int flags) {
+        return createBuffer(newDirectByteBuffer(size), flags);
+    }
+
+    /**
+     * Creates a CLBuffer with the specified flags. No flags creates a MEM.READ_WRITE buffer.
+     */
+    public final CLBuffer<?> createBuffer(int size, Mem... flags) {
+        return createBuffer(size, Mem.flagsToInt(flags));
+    }
+
+    /**
+     * Creates a CLBuffer with the specified flags.
+     */
+    public final CLBuffer<?> createBuffer(int size, int flags) {
+        CLBuffer<?> buffer = CLBuffer.create(this, size, flags);
+        memoryObjects.add(buffer);
+        return buffer;
+    }
+
+    /**
+     * Creates a CLBuffer with the specified flags. No flags creates a MEM.READ_WRITE buffer.
+     */
+    public final <B extends Buffer> CLBuffer<B> createBuffer(B directBuffer, Mem... flags) {
+        return createBuffer(directBuffer, Mem.flagsToInt(flags));
+    }
+
+    /**
+     * Creates a CLBuffer with the specified flags.
+     */
+    public final <B extends Buffer> CLBuffer<B> createBuffer(B directBuffer, int flags) {
+        CLBuffer<B> buffer = CLBuffer.create(this, directBuffer, flags);
+        memoryObjects.add(buffer);
+        return buffer;
+    }
+
+    CLCommandQueue createCommandQueue(CLDevice device, long properties) {
+
+        CLCommandQueue queue = CLCommandQueue.create(this, device, properties);
+
+        List<CLCommandQueue> list = queuesMap.get(device);
+        if(list == null) {
+            list = new ArrayList<CLCommandQueue>();
+            queuesMap.put(device, list);
+        }
+        list.add(queue);
+
+        return queue;
+    }
+
+    public CLSampler createSampler(AddressingMode addrMode, FilteringMode filtMode, boolean normalizedCoords) {
+        CLSampler sampler = CLSampler.create(this, addrMode, filtMode, normalizedCoords);
+        samplers.add(sampler);
+        return sampler;
+    }
+
+    void onProgramReleased(CLProgram program) {
+        programs.remove(program);
+    }
+
+    void onMemoryReleased(CLMemory<?> buffer) {
+        memoryObjects.remove(buffer);
+    }
+
+    void onCommandQueueReleased(CLDevice device, CLCommandQueue queue) {
+        List<CLCommandQueue> list = queuesMap.get(device);
+        list.remove(queue);
+        // remove empty lists from map
+        if(list.isEmpty())
+            queuesMap.remove(device);
+    }
+
+    void onSamplerReleased(CLSampler sampler) {
+        samplers.remove(sampler);
+    }
+
+    /**
+     * Releases the context and all resources.
+     */
+    public void release() {
+
+        //release all resources
+        while(!programs.isEmpty())
+            programs.get(0).release();
+
+        while(!memoryObjects.isEmpty())
+            memoryObjects.get(0).release();
+
+        while(!samplers.isEmpty())
+            samplers.get(0).release();
+
+        for (CLDevice device : devices) {
+            List<CLCommandQueue> list = queuesMap.get(device);
+            if(list != null) {
+                while(!list.isEmpty()) {
+                    list.get(0).release();
+                }
+            }
+        }
+
+        int ret = cl.clReleaseContext(ID);
+        checkForError(ret, "error releasing context");
+    }
+
+    public void close() {
+        release();
+    }
+
+    protected void overrideContext(CLDevice device) {
+        device.setContext(this);
+    }
+
+    /**
+     * Returns the CLPlatform this context is running on.
+     */
+    @Override
+    public CLPlatform getPlatform() {
+        return platform;
+    }
+
+    @Override
+    public CLContext getContext() {
+        return this;
+    }
+
+    /**
+     * Returns a read only view of all programs associated with this context.
+     */
+    public List<CLProgram> getPrograms() {
+        return Collections.unmodifiableList(programs);
+    }
+
+    /**
+     * Returns a read only view of all allocated memory objects associated with this context.
+     */
+    public List<CLMemory<? extends Buffer>> getMemoryObjects() {
+        return Collections.unmodifiableList(memoryObjects);
+    }
+
+    /**
+     * Returns a read only view of all samplers associated with this context.
+     */
+    public List<CLSampler> getSamplers() {
+        return Collections.unmodifiableList(samplers);
+    }
+
+    /**
+     * Returns the device with maximal FLOPS from this context.
+     * The device speed is estimated by calculating the product of
+     * MAX_COMPUTE_UNITS and MAX_CLOCK_FREQUENCY.
+     * @see #getMaxFlopsDevice(com.jogamp.opencl.CLDevice.Type)
+     */
+    public CLDevice getMaxFlopsDevice() {
+        return CLPlatform.findMaxFlopsDevice(getDevices());
+    }
+
+    /**
+     * Returns the device with maximal FLOPS of the specified device type from this context.
+     * The device speed is estimated by calculating the product of
+     * MAX_COMPUTE_UNITS and MAX_CLOCK_FREQUENCY.
+     */
+    public CLDevice getMaxFlopsDevice(CLDevice.Type type) {
+        return CLPlatform.findMaxFlopsDevice(getDevices(), type);
+    }
+
+    /**
+     * Returns all devices associated with this CLContext.
+     */
+    public CLDevice[] getDevices() {
+        initDevices();
+        return devices;
+    }
+
+    /**
+     * Return the low level OpenCL interface.
+     */
+    public CL getCL() {
+        return cl;
+    }
+
+    CLDevice getDevice(long dID) {
+        CLDevice[] deviceArray = getDevices();
+        for (int i = 0; i < deviceArray.length; i++) {
+            if(dID == deviceArray[i].ID)
+                return deviceArray[i];
+        }
+        return null;
+    }
+
+    protected static long toDeviceBitmap(Type[] deviceTypes) {
+        long type = 0;
+        if (deviceTypes != null) {
+            for (int i = 0; i < deviceTypes.length; i++) {
+                type |= deviceTypes[i].TYPE;
+            }
+        }
+        return type;
+    }
+
+    @Override
+    public String toString() {
+        return "CLContext [id: " + ID
+                      + " #devices: " + getDevices().length
+                      + "]";
+    }
+
+    @Override
+    public boolean equals(Object obj) {
+        if (obj == null) {
+            return false;
+        }
+        if (getClass() != obj.getClass()) {
+            return false;
+        }
+        final CLContext other = (CLContext) obj;
+        if (this.ID != other.ID) {
+            return false;
+        }
+        return true;
+    }
+
+    @Override
+    public int hashCode() {
+        int hash = 7;
+        hash = 23 * hash + (int) (this.ID ^ (this.ID >>> 32));
+        return hash;
+    }
+
+}
diff --git a/src/com/jogamp/opencl/CLDevice.java b/src/com/jogamp/opencl/CLDevice.java
new file mode 100644
index 00000000..3b26e507
--- /dev/null
+++ b/src/com/jogamp/opencl/CLDevice.java
@@ -0,0 +1,823 @@
+package com.jogamp.opencl;
+
+import com.jogamp.opencl.util.CLUtil;
+import com.jogamp.common.nio.Int64Buffer;
+import com.jogamp.common.os.Platform;
+import java.nio.Buffer;
+import java.nio.ByteBuffer;
+import java.util.ArrayList;
+import java.util.Collections;
+import java.util.EnumSet;
+import java.util.HashSet;
+import java.util.List;
+import java.util.Map;
+import java.util.Scanner;
+import java.util.Set;
+
+import static com.jogamp.opencl.CL.*;
+
+/**
+ * This object represents an OpenCL device.
+ * @see CLPlatform#listCLDevices(com.jogamp.opencl.CLDevice.Type...)
+ * @see CLPlatform#getMaxFlopsDevice(com.jogamp.opencl.CLDevice.Type...)
+ * @see CLContext#getDevices()
+ * @see CLContext#getMaxFlopsDevice(com.jogamp.opencl.CLDevice.Type)
+ * @author Michael Bien
+ */
+public final class CLDevice extends CLObject {
+
+    private Set<String> extensions;
+    
+    private final CLDeviceInfoAccessor deviceInfo;
+
+    CLDevice(CL cl, long id) {
+        super(cl, id);
+        this.deviceInfo = new CLDeviceInfoAccessor();
+    }
+
+    CLDevice(CLContext context, long id) {
+        super(context, id);
+        this.deviceInfo = new CLDeviceInfoAccessor();
+    }
+
+    public CLCommandQueue createCommandQueue() {
+        return createCommandQueue(0);
+    }
+
+    public CLCommandQueue createCommandQueue(CLCommandQueue.Mode property) {
+        return createCommandQueue(property.QUEUE_MODE);
+    }
+
+    public CLCommandQueue createCommandQueue(CLCommandQueue.Mode... properties) {
+        int flags = 0;
+        if(properties != null) {
+            for (int i = 0; i < properties.length; i++) {
+                flags |= properties[i].QUEUE_MODE;
+            }
+        }
+        return createCommandQueue(flags);
+    }
+    
+    public CLCommandQueue createCommandQueue(long properties) {
+        if(context == null)
+            throw new IllegalStateException("this device is not associated with a context");
+        return context.createCommandQueue(this, properties);
+    }
+    
+    /*keep this package private*/
+    void setContext(CLContext context) {
+        this.context = context;
+    }
+
+    /**
+     * Returns the name of this device.
+     */
+    public String getName() {
+        return deviceInfo.getString(CL_DEVICE_NAME);
+    }
+
+    /**
+     * Returns the OpenCL profile of this device.
+     */
+    public String getProfile() {
+        return deviceInfo.getString(CL_DEVICE_PROFILE);
+    }
+
+    /**
+     * Returns the vendor of this device.
+     */
+    public String getVendor() {
+        return deviceInfo.getString(CL_DEVICE_VENDOR);
+    }
+
+    /**
+     * Returns the vendor id of this device.
+     */
+    public long getVendorID() {
+        return deviceInfo.getLong(CL_DEVICE_VENDOR_ID);
+    }
+
+    /**
+     * Returns OpenCL version string. Returns the OpenCL version supported by the device.
+     * This version string has the following format:<br>
+     * OpenCL[space][major_version.minor_version][space][vendor-specific information]
+     */
+    public String getVersion() {
+        return deviceInfo.getString(CL_DEVICE_VERSION);
+    }
+
+    /**
+     * Returns OpenCL software driver version string in the form major_number.minor_number.
+     */
+    public String getDriverVersion() {
+        return deviceInfo.getString(CL_DRIVER_VERSION);
+    }
+
+    /**
+     * Returns the type of this device.
+     */
+    public Type getType() {
+        return Type.valueOf((int)deviceInfo.getLong(CL_DEVICE_TYPE));
+    }
+
+    /**
+     * The default compute device address space size specified in bits.
+     * Currently supported values are 32 or 64 bits.
+     */
+    public int getAddressBits() {
+        return (int)deviceInfo.getLong(CL_DEVICE_ADDRESS_BITS);
+    }
+
+    /**
+     * Preferred native vector width size for built-in short vectors.
+     * The vector width is defined as the number of scalar elements that can be stored in the vector.
+     */
+    public int getPreferredShortVectorWidth() {
+        return (int)deviceInfo.getLong(CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT);
+    }
+
+    /**
+     * Preferred native vector width size for built-in char vectors.
+     * The vector width is defined as the number of scalar elements that can be stored in the vector.
+     */
+    public int getPreferredCharVectorWidth() {
+        return (int)deviceInfo.getLong(CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR);
+    }
+
+    /**
+     * Preferred native vector width size for built-in int vectors.
+     * The vector width is defined as the number of scalar elements that can be stored in the vector.
+     */
+    public int getPreferredIntVectorWidth() {
+        return (int)deviceInfo.getLong(CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT);
+    }
+
+    /**
+     * Preferred native vector width size for built-in long vectors.
+     * The vector width is defined as the number of scalar elements that can be stored in the vector.
+     */
+    public int getPreferredLongVectorWidth() {
+        return (int)deviceInfo.getLong(CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG);
+    }
+
+    /**
+     * Preferred native vector width size for built-in float vectors.
+     * The vector width is defined as the number of scalar elements that can be stored in the vector.
+     */
+    public int getPreferredFloatVectorWidth() {
+        return (int)deviceInfo.getLong(CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT);
+    }
+
+    /**
+     * Preferred native vector width size for built-in double vectors.
+     * The vector width is defined as the number of scalar elements that can be stored in the vector.
+     */
+    public int getPreferredDoubleVectorWidth() {
+        return (int)deviceInfo.getLong(CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE);
+    }
+
+    /**
+     * Returns the number of parallel compute cores on the OpenCL device.
+     * The minimum value is 1.
+     */
+    public int getMaxComputeUnits() {
+        return (int) deviceInfo.getLong(CL_DEVICE_MAX_COMPUTE_UNITS);
+    }
+
+    /**
+     * Returns the maximum number of work-items in a work-group executing
+     * a kernel using the data parallel execution model.
+     * The minimum value is 1.
+     */
+    public int getMaxWorkGroupSize() {
+        return (int) deviceInfo.getLong(CL_DEVICE_MAX_WORK_GROUP_SIZE);
+    }
+
+    /**
+     * Returns the maximum configured clock frequency of the device in MHz.
+     */
+    public int getMaxClockFrequency() {
+        return (int) (deviceInfo.getLong(CL_DEVICE_MAX_CLOCK_FREQUENCY));
+    }
+
+    /**
+     * Returns the maximum dimensions that specify the global and local work-item
+     * IDs used by the data parallel execution model.
+     * The minimum value is 3.
+     */
+    public int getMaxWorkItemDimensions() {
+        return (int) deviceInfo.getLong(CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS);
+    }
+
+    /**
+     * Returns the maximum number of work-items that can be specified in each
+     * dimension of the work-group.
+     * The minimum value is (1, 1, 1).
+     */
+    public int[] getMaxWorkItemSizes() {
+        int n = (int) deviceInfo.getLong(CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS);
+        return deviceInfo.getInts(n, CL_DEVICE_MAX_WORK_ITEM_SIZES);
+    }
+
+    /**
+     * Returns the max size in bytes of the arguments that can be passed to a kernel.
+     * The minimum value is 256.
+     */
+    public long getMaxParameterSize() {
+        return deviceInfo.getLong(CL_DEVICE_MAX_PARAMETER_SIZE);
+    }
+
+    /**
+     * Returns the maximal allocatable memory on this device.
+     */
+    public long getMaxMemAllocSize() {
+        return deviceInfo.getLong(CL_DEVICE_MAX_MEM_ALLOC_SIZE);
+    }
+
+    /**
+     * Returns the global memory size in bytes.
+     */
+    public long getGlobalMemSize() {
+        return deviceInfo.getLong(CL_DEVICE_GLOBAL_MEM_SIZE);
+    }
+
+    /**
+     * Returns the local memory size in bytes.
+     */
+    public long getLocalMemSize() {
+        return deviceInfo.getLong(CL_DEVICE_LOCAL_MEM_SIZE);
+    }
+
+    /**
+     * Returns the max size in bytes of a constant buffer allocation.
+     * The minimum value is 64 KB.
+     */
+    public long getMaxConstantBufferSize() {
+        return deviceInfo.getLong(CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE);
+    }
+
+    /**
+     * Returns the size of global memory cache line in bytes.
+     */
+    public long getGlobalMemCachelineSize() {
+        return deviceInfo.getLong(CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE);
+    }
+
+    /**
+     * Returns the size of global memory cache in bytes.
+     */
+    public long getGlobalMemCacheSize() {
+        return deviceInfo.getLong(CL_DEVICE_GLOBAL_MEM_CACHE_SIZE);
+    }
+
+    /**
+     * Returns the max number of arguments declared with the <code>constant</code>
+     * qualifier in a kernel. The minimum value is 8.
+     */
+    public long getMaxConstantArgs() {
+        return deviceInfo.getLong(CL_DEVICE_MAX_CONSTANT_ARGS);
+    }
+
+    /**
+     * Returns true if images are supported by the OpenCL device and false otherwise.
+     */
+    public boolean isImageSupportAvailable() {
+        return deviceInfo.getLong(CL_DEVICE_IMAGE_SUPPORT) == CL_TRUE;
+    }
+
+    /**
+     * Returns the max number of simultaneous image objects that can be read by a kernel.
+     * The minimum value is 128 if image support is available.
+     */
+    public int getMaxReadImageArgs() {
+        return (int)deviceInfo.getLong(CL_DEVICE_MAX_READ_IMAGE_ARGS);
+    }
+
+    /**
+     * Returns the max number of simultaneous image objects that can be written by a kernel.
+     * The minimum value is 8 if image support is available.
+     */
+    public int getMaxWriteImageArgs() {
+        return (int)deviceInfo.getLong(CL_DEVICE_MAX_WRITE_IMAGE_ARGS);
+    }
+
+    /**
+     * Returns the max width of 2D image in pixels. The minimum value is 8192 if
+     * image support is available.
+     */
+    public int getMaxImage2dWidth() {
+        return (int)deviceInfo.getLong(CL_DEVICE_IMAGE2D_MAX_WIDTH);
+    }
+
+    /**
+     * Returns the max height of 2D image in pixels. The minimum value is 8192 if
+     * image support is available.
+     */
+    public int getMaxImage2dHeight() {
+        return (int)deviceInfo.getLong(CL_DEVICE_IMAGE2D_MAX_HEIGHT);
+    }
+
+    /**
+     * Returns the max width of 3D image in pixels. The minimum value is 2048 if
+     * image support is available.
+     */
+    public int getMaxImage3dWidth() {
+        return (int)deviceInfo.getLong(CL_DEVICE_IMAGE3D_MAX_WIDTH);
+    }
+
+    /**
+     * Returns the max height of 3D image in pixels. The minimum value is 2048 if
+     * image support is available.
+     */
+    public int getMaxImage3dHeight() {
+        return (int)deviceInfo.getLong(CL_DEVICE_IMAGE3D_MAX_HEIGHT);
+    }
+
+    /**
+     * Returns the max depth of 3D image in pixels. The minimum value is 2048 if
+     * image support is available.
+     */
+    public int getMaxImage3dDepth() {
+        return (int)deviceInfo.getLong(CL_DEVICE_IMAGE3D_MAX_DEPTH);
+    }
+
+    /**
+     * Returns the maximum number of samplers that can be used in a kernel. The
+     * minimum value is 16 if image support is available.
+     */
+    public int getMaxSamplers() {
+        return (int)deviceInfo.getLong(CL_DEVICE_MAX_SAMPLERS);
+    }
+
+    /**
+     * Returns the resolution of device timer. This is measured in nanoseconds.
+     */
+    public long getProfilingTimerResolution() {
+        return deviceInfo.getLong(CL_DEVICE_PROFILING_TIMER_RESOLUTION);
+    }
+
+    /**
+     * Returns the execution capabilities as EnumSet.
+     */
+    public EnumSet<Capabilities> getExecutionCapabilities() {
+        return Capabilities.valuesOf((int)deviceInfo.getLong(CL_DEVICE_EXECUTION_CAPABILITIES));
+    }
+
+    /**
+     * Returns the optional half precision floating-point capability of the device.
+     * The required minimum half precision floating-point capabilities as implemented by this
+     * extension are {@link FPConfig#ROUND_TO_ZERO}, {@link FPConfig#ROUND_TO_INF}
+     * and {@link FPConfig#INF_NAN}.
+     * @return An EnumSet containing the extensions, never null.
+     */
+    public EnumSet<FPConfig> getHalfFPConfig() {
+        if(isHalfFPAvailable())
+            return FPConfig.valuesOf((int)deviceInfo.getLong(CL_DEVICE_HALF_FP_CONFIG));
+        else
+            return EnumSet.noneOf(FPConfig.class);
+    }
+
+    /**
+     * Returns the single precision floating-point capability of the device.
+     * The mandated minimum floating-point capabilities are {@link FPConfig#ROUND_TO_NEAREST} and
+     * {@link FPConfig#INF_NAN}.
+     * @return An EnumSet containing the extensions, never null.
+     */
+    public EnumSet<FPConfig> getSingleFPConfig() {
+        return FPConfig.valuesOf((int)deviceInfo.getLong(CL_DEVICE_SINGLE_FP_CONFIG));
+    }
+
+    /**
+     * Returns the optional double precision floating-point capability of the device.
+     * The mandated minimum double precision floating-point capabilities are {@link FPConfig#FMA},
+     * {@link FPConfig#ROUND_TO_NEAREST}, {@link FPConfig#ROUND_TO_ZERO},
+     * {@link FPConfig#ROUND_TO_INF}, {@link FPConfig#INF_NAN}, and {@link FPConfig#DENORM}.
+     * @return An EnumSet containing the extensions, never null.
+     */
+    public EnumSet<FPConfig> getDoubleFPConfig() {
+        if(isDoubleFPAvailable())
+            return FPConfig.valuesOf((int)deviceInfo.getLong(CL_DEVICE_DOUBLE_FP_CONFIG));
+        else
+            return EnumSet.noneOf(FPConfig.class);
+    }
+
+    /**
+     * Returns the local memory type.
+     */
+    public LocalMemType getLocalMemType() {
+        return LocalMemType.valueOf((int)deviceInfo.getLong(CL_DEVICE_LOCAL_MEM_TYPE));
+    }
+
+    /**
+     * Returns the type of global memory cache supported.
+     */
+    public GlobalMemCacheType getGlobalMemCacheType() {
+        return GlobalMemCacheType.valueOf((int)deviceInfo.getLong(CL_DEVICE_GLOBAL_MEM_CACHE_TYPE));
+    }
+
+    /**
+     * Returns the command-queue properties supported by the device.
+     */
+    public EnumSet<CLCommandQueue.Mode> getQueueProperties() {
+        return CLCommandQueue.Mode.valuesOf((int)deviceInfo.getLong(CL_DEVICE_QUEUE_PROPERTIES));
+    }
+
+    /**
+     * Returns true if this device is available.
+     */
+    public boolean isAvailable() {
+        return deviceInfo.getLong(CL_DEVICE_AVAILABLE) == CL_TRUE;
+    }
+
+    /**
+     * Returns false if the implementation does not have a compiler available to
+     * compile the program source. Is true if the compiler is available.
+     * This can be false for the OpenCL ES profile only.
+     */
+    public boolean isCompilerAvailable() {
+        return deviceInfo.getLong(CL_DEVICE_COMPILER_AVAILABLE) == CL_TRUE;
+    }
+
+    /**
+     * Returns true if the OpenCL device is a little endian device and false otherwise.
+     */
+    public boolean isLittleEndian() {
+        return deviceInfo.getLong(CL_DEVICE_ENDIAN_LITTLE) == CL_TRUE;
+    }
+
+    /**
+     * Returns true if the device implements error correction for the memories,
+     * caches, registers etc. in the device. Is false if the device does not
+     * implement error correction.
+     */
+    public boolean isErrorCorrectionSupported() {
+        return deviceInfo.getLong(CL_DEVICE_ERROR_CORRECTION_SUPPORT) == CL_TRUE;
+    }
+
+    /**
+     * Returns {@link #isExtensionAvailable}("cl_khr_fp16").
+     * @see #getExtensions()
+     */
+    public boolean isHalfFPAvailable() {
+        return isExtensionAvailable("cl_khr_fp16");
+    }
+
+    /**
+     * Returns {@link #isExtensionAvailable}("cl_khr_fp64").
+     * @see #getExtensions()
+     */
+    public boolean isDoubleFPAvailable() {
+        return isExtensionAvailable("cl_khr_fp64");
+    }
+
+    /**
+     * Returns {@link #isExtensionAvailable}("cl_khr_gl_sharing") || {@link #isExtensionAvailable}("cl_apple_gl_sharing").
+     * @see #getExtensions()
+     */
+    public boolean isGLMemorySharingSupported() {
+        return isExtensionAvailable("cl_khr_gl_sharing") || isExtensionAvailable("cl_apple_gl_sharing");
+    }
+
+    /**
+     * Returns true if the extension is supported on this device.
+     * @see #getExtensions()
+     */
+    public boolean isExtensionAvailable(String extension) {
+        return getExtensions().contains(extension);
+    }
+
+    /**
+     * Returns all device extension names as unmodifiable Set.
+     */
+    public Set<String> getExtensions() {
+
+        if(extensions == null) {
+            extensions = new HashSet<String>();
+            String ext = deviceInfo.getString(CL_DEVICE_EXTENSIONS);
+            Scanner scanner = new Scanner(ext);
+
+            while(scanner.hasNext())
+                extensions.add(scanner.next());
+
+            extensions = Collections.unmodifiableSet(extensions);
+        }
+
+        return extensions;
+    }
+
+    /**
+     * Returns a Map of device properties with the enum names as keys.
+     * @see CLUtil#obtainDeviceProperties(com.jogamp.opencl.CLDevice)
+     */
+    public Map<String, String> getProperties() {
+        return CLUtil.obtainDeviceProperties(this);
+    }
+
+    private final class CLDeviceInfoAccessor extends CLInfoAccessor {
+
+        @Override
+        protected int getInfo(int name, long valueSize, Buffer value, Int64Buffer valueSizeRet) {
+            return cl.clGetDeviceInfo(ID, name, valueSize, value, valueSizeRet);
+        }
+
+        private int[] getInts(int n, int key) {
+
+            ByteBuffer buffer = localBB.get();
+            int ret = getInfo(key, buffer.capacity(), buffer, null);
+            CLException.checkForError(ret, "error while asking device for infos");
+
+            int[] array = new int[n];
+            for(int i = 0; i < array.length; i++) {
+                if(Platform.is32Bit()) {
+                    array[i] = buffer.getInt();
+                }else{
+                    array[i] = (int)buffer.getLong();
+                }
+            }
+            buffer.rewind();
+            
+            return array;
+        }
+
+    }
+
+
+    @Override
+    public String toString() {
+        return "CLDevice [id: " + ID
+                      + " name: " + getName()
+                      + " type: " + getType()
+                      + " profile: " + getProfile()+"]";
+    }
+
+    @Override
+    public boolean equals(Object obj) {
+        if (obj == null) {
+            return false;
+        }
+        if (getClass() != obj.getClass()) {
+            return false;
+        }
+        final CLDevice other = (CLDevice) obj;
+        if (this.ID != other.ID) {
+            return false;
+        }
+        return true;
+    }
+
+    @Override
+    public int hashCode() {
+        int hash = 3;
+        hash = 79 * hash + (int) (this.ID ^ (this.ID >>> 32));
+        return hash;
+    }
+
+    /**
+     * Enumeration for the execution capabilities of the device.
+     */
+    public enum Capabilities {
+
+        /**
+         * The OpenCL device can execute OpenCL kernels.
+         */
+        EXEC_KERNEL(CL_EXEC_KERNEL),
+
+        /**
+         * The OpenCL device can execute native kernels.
+         */
+        EXEC_NATIVE_KERNEL(CL_EXEC_NATIVE_KERNEL);
+
+        /**
+         * Value of wrapped OpenCL device type.
+         */
+        public final int CAPS;
+
+        private Capabilities(int type) {
+            this.CAPS = type;
+        }
+
+        public static Capabilities valueOf(int caps) {
+            switch(caps) {
+                case(CL_EXEC_KERNEL):
+                    return EXEC_KERNEL;
+                case(CL_EXEC_NATIVE_KERNEL):
+                    return EXEC_NATIVE_KERNEL;
+            }
+            return null;
+        }
+
+        public static EnumSet<Capabilities> valuesOf(int bitfield) {
+            if((EXEC_KERNEL.CAPS & bitfield) != 0) {
+                if((EXEC_NATIVE_KERNEL.CAPS & bitfield) != 0) {
+                    return EnumSet.of(EXEC_KERNEL, EXEC_NATIVE_KERNEL);
+                }else{
+                    return EnumSet.of(EXEC_KERNEL);
+                }
+            }else if((EXEC_NATIVE_KERNEL.CAPS & bitfield) != 0){
+                return EnumSet.of(EXEC_NATIVE_KERNEL);
+            }
+            return null;
+        }
+
+    }
+
+    /**
+     * Enumeration for the type of a device.
+     */
+    public enum Type {
+        /**
+         * CL_DEVICE_TYPE_CPU
+         */
+        CPU(CL_DEVICE_TYPE_CPU),
+        /**
+         * CL_DEVICE_TYPE_GPU
+         */
+        GPU(CL_DEVICE_TYPE_GPU),
+        /**
+         * CL_DEVICE_TYPE_ACCELERATOR
+         */
+        ACCELERATOR(CL_DEVICE_TYPE_ACCELERATOR),
+        /**
+         * CL_DEVICE_TYPE_DEFAULT. This type can be used for creating a context on
+         * the default device, a single device can never have this type.
+         */
+        DEFAULT(CL_DEVICE_TYPE_DEFAULT),
+        /**
+         * CL_DEVICE_TYPE_ALL. This type can be used for creating a context on
+         * all devices, a single device can never have this type.
+         */
+        ALL(CL_DEVICE_TYPE_ALL);
+
+        /**
+         * Value of wrapped OpenCL device type.
+         */
+        public final long TYPE;
+
+        private Type(long type) {
+            this.TYPE = type;
+        }
+
+        public static Type valueOf(long clDeviceType) {
+
+            if(clDeviceType == CL_DEVICE_TYPE_ALL)
+                return ALL;
+
+            switch((int)clDeviceType) {
+                case(CL_DEVICE_TYPE_DEFAULT):
+                    return DEFAULT;
+                case(CL_DEVICE_TYPE_CPU):
+                    return CPU;
+                case(CL_DEVICE_TYPE_GPU):
+                    return GPU;
+                case(CL_DEVICE_TYPE_ACCELERATOR):
+                    return ACCELERATOR;
+            }
+            return null;
+        }
+    }
+
+    /**
+     * Describes floating-point capability of the device.
+     * Zero or more values are possible.
+     */
+    public enum FPConfig {
+
+        /**
+         * denorms are supported.
+         */
+        DENORM(CL_FP_DENORM),
+
+        /**
+         * INF and quiet NaNs are supported.
+         */
+        INF_NAN(CL_FP_INF_NAN),
+
+        /**
+         * round to nearest rounding mode supported.
+         */
+        ROUND_TO_NEAREST(CL_FP_ROUND_TO_NEAREST),
+
+        /**
+         * round to +ve and –ve infinity rounding modes supported.
+         */
+        ROUND_TO_INF(CL_FP_ROUND_TO_INF),
+
+        /**
+         * round to zero rounding mode supported.
+         */
+        ROUND_TO_ZERO(CL_FP_ROUND_TO_ZERO),
+
+        /**
+         * IEEE754-2008 fused multiply-add is supported.
+         */
+        FMA(CL_FP_FMA);
+
+
+        /**
+         * Value of wrapped OpenCL bitfield.
+         */
+        public final int CONFIG;
+
+        private FPConfig(int config) {
+            this.CONFIG = config;
+        }
+
+        /**
+         * Returns a EnumSet for the given bitfield.
+         */
+        public static EnumSet<FPConfig> valuesOf(int bitfield) {
+            List<FPConfig> matching = new ArrayList<FPConfig>();
+            FPConfig[] values = FPConfig.values();
+            for (FPConfig value : values) {
+                if((value.CONFIG & bitfield) != 0)
+                    matching.add(value);
+            }
+            if(matching.isEmpty())
+                return EnumSet.noneOf(FPConfig.class);
+            else
+                return EnumSet.copyOf(matching);
+        }
+
+    }
+
+    /**
+     * Type of global memory cache supported.
+     */
+    public enum GlobalMemCacheType {
+
+        /**
+         * Global memory cache not supported.
+         */
+        NONE(CL_NONE),
+
+        /**
+         * Read only cache.
+         */
+        READ_ONLY(CL_READ_ONLY_CACHE),
+
+        /**
+         * Read-write cache.
+         */
+        READ_WRITE(CL_READ_WRITE_CACHE);
+        
+
+        /**
+         * Value of wrapped OpenCL value.
+         */
+        public final int TYPE;
+
+        private GlobalMemCacheType(int type) {
+            this.TYPE = type;
+        }
+
+        /**
+         * Returns the matching GlobalMemCacheType for the given cl type.
+         */
+        public static GlobalMemCacheType valueOf(int bitfield) {
+            GlobalMemCacheType[] values = GlobalMemCacheType.values();
+            for (GlobalMemCacheType value : values) {
+                if(value.TYPE == bitfield)
+                    return value;
+            }
+            return null;
+        }
+    }
+
+    /**
+     * Type of local memory cache supported.
+     */
+    public enum LocalMemType {
+
+        /**
+         * GLOBAL implies that no dedicated memory storage is available (global mem is used instead).
+         */
+        GLOBAL(CL_GLOBAL),
+
+        /**
+         * LOCAL implies dedicated local memory storage such as SRAM.
+         */
+        LOCAL(CL_LOCAL);
+
+        /**
+         * Value of wrapped OpenCL value.
+         */
+        public final int TYPE;
+
+        private LocalMemType(int type) {
+            this.TYPE = type;
+        }
+
+        /**
+         * Returns the matching LocalMemCacheType for the given cl type.
+         */
+        public static LocalMemType valueOf(int clLocalCacheType) {
+            if(clLocalCacheType == CL_GLOBAL)
+                return GLOBAL;
+            else if(clLocalCacheType == CL_LOCAL)
+                return LOCAL;
+            return null;
+        }
+
+    }
+
+}
diff --git a/src/com/jogamp/opencl/CLEvent.java b/src/com/jogamp/opencl/CLEvent.java
new file mode 100644
index 00000000..46b92a6d
--- /dev/null
+++ b/src/com/jogamp/opencl/CLEvent.java
@@ -0,0 +1,263 @@
+package com.jogamp.opencl;
+
+import com.jogamp.common.nio.Int64Buffer;
+import java.nio.Buffer;
+
+import static com.jogamp.opencl.CL.*;
+import static com.jogamp.opencl.CLException.*;
+
+/**
+ * Event objects can be used for synchronizing command queues, e.g you can wait until a
+ * event occurs or they can also be used to capture profiling information that
+ * measure execution time of a command.
+ * Profiling of OpenCL commands can be enabled by using a {@link com.jogamp.opencl.CLCommandQueue} created with
+ * {@link com.jogamp.opencl.CLCommandQueue.Mode#PROFILING_MODE}.
+ * @author Michael Bien
+ */
+public class CLEvent extends CLObject implements CLResource {
+
+    private final CLEventInfoAccessor eventInfo;
+    private final CLEventProfilingInfoAccessor eventProfilingInfo;
+
+    CLEvent(CLContext context, long id) {
+        super(context, id);
+        this.eventInfo = new CLEventInfoAccessor();
+        this.eventProfilingInfo = new CLEventProfilingInfoAccessor();
+    }
+
+    public void release() {
+        int ret = cl.clReleaseEvent(ID);
+        checkForError(ret, "can not release event");
+    }
+
+    public void close() {
+        release();
+    }
+
+    /**
+     * Returns the execution status of the command which triggers this event.
+     */
+    public ExecutionStatus getStatus() {
+        return ExecutionStatus.valueOf(getStatusCode());
+    }
+    
+    public int getStatusCode() {
+        return (int)eventInfo.getLong(CL_EVENT_COMMAND_EXECUTION_STATUS);
+    }
+
+    public CommandType getType() {
+        int status = (int)eventInfo.getLong(CL_EVENT_COMMAND_TYPE);
+        return CommandType.valueOf(status);
+    }
+
+    public long getProfilingInfo(ProfilingCommand command) {
+        return eventProfilingInfo.getLong(command.COMMAND);
+    }
+
+
+    @Override
+    public String toString() {
+        return "CLEvent [id: " + ID
+                      + " name: " + getType()
+                      + " status: " + getStatus()+"]";
+    }
+
+    @Override
+    public boolean equals(Object obj) {
+        if (obj == null) {
+            return false;
+        }
+        if (getClass() != obj.getClass()) {
+            return false;
+        }
+        final CLEvent other = (CLEvent) obj;
+        if (this.context != other.context && (this.context == null || !this.context.equals(other.context))) {
+            return false;
+        }
+        if (this.ID != other.ID) {
+            return false;
+        }
+        return true;
+    }
+
+    @Override
+    public int hashCode() {
+        int hash = 5;
+        hash = 13 * hash + (this.context != null ? this.context.hashCode() : 0);
+        hash = 13 * hash + (int) (this.ID ^ (this.ID >>> 32));
+        return hash;
+    }
+
+    
+
+    private class CLEventInfoAccessor extends CLInfoAccessor {
+
+        @Override
+        protected int getInfo(int name, long valueSize, Buffer value, Int64Buffer valueSizeRet) {
+            return cl.clGetEventInfo(ID, name, valueSize, value, valueSizeRet);
+        }
+
+    }
+
+    private class CLEventProfilingInfoAccessor extends CLInfoAccessor {
+
+        @Override
+        protected int getInfo(int name, long valueSize, Buffer value, Int64Buffer valueSizeRet) {
+            return cl.clGetEventProfilingInfo(ID, name, valueSize, value, valueSizeRet);
+        }
+
+    }
+
+    // TODO merge with ExecutionStatus?
+
+    public enum ProfilingCommand {
+
+        /**
+         * A 64-bit value that describes the current device time counter in nanoseconds
+         * when the command identified by event is enqueued in a command-queue by the host.
+         */
+        QUEUED(CL_PROFILING_COMMAND_QUEUED),
+
+        /**
+         * A 64-bit value that describes the current device time counter in nanoseconds when
+         * the command identified by event that has been enqueued is submitted by the host to
+         * the device associated with the commandqueue.
+         */
+        SUBMIT(CL_PROFILING_COMMAND_SUBMIT),
+
+        /**
+         * A 64-bit value that describes the current device time counter in nanoseconds when
+         * the command identified by event starts execution on the device.
+         */
+        START(CL_PROFILING_COMMAND_START),
+
+        /**
+         * A 64-bit value that describes the current device time counter in nanoseconds when
+         * the command identified by event has finished execution on the device.
+         */
+        END(CL_PROFILING_COMMAND_END);
+
+        /**
+         * Value of wrapped OpenCL profiling command.
+         */
+        public final int COMMAND;
+
+        private ProfilingCommand(int command) {
+            this.COMMAND = command;
+        }
+
+        public static ProfilingCommand valueOf(int status) {
+            switch(status) {
+                case(CL_PROFILING_COMMAND_QUEUED):
+                    return QUEUED;
+                case(CL_PROFILING_COMMAND_SUBMIT):
+                    return SUBMIT;
+                case(CL_PROFILING_COMMAND_START):
+                    return START;
+                case(CL_PROFILING_COMMAND_END):
+                    return END;
+            }
+            return null;
+        }
+
+    }
+
+
+
+    public enum ExecutionStatus {
+
+        /**
+         * Command has been enqueued in the command-queue.
+         */
+        QUEUED(CL_QUEUED),
+
+        /**
+         * Enqueued command has been submitted by the host to the device
+         * associated with the command-queue.
+         */
+        SUBMITTED(CL_SUBMITTED),
+        
+        /**
+         * Device is currently executing this command.
+         */
+        RUNNING(CL_RUNNING),
+
+        /**
+         * The command has completed.
+         */
+        COMPLETE(CL_COMPLETE),
+
+        /**
+         * The command did not complete because of an error.
+         */
+        ERROR(-1);
+
+
+        /**
+         * Value of wrapped OpenCL command execution status.
+         */
+        public final int STATUS;
+
+        private ExecutionStatus(int status) {
+            this.STATUS = status;
+        }
+
+        public static ExecutionStatus valueOf(int status) {
+            switch(status) {
+                case(CL_QUEUED):
+                    return QUEUED;
+                case(CL_SUBMITTED):
+                    return SUBMITTED;
+                case(CL_RUNNING):
+                    return RUNNING;
+                case(CL_COMPLETE):
+                    return COMPLETE;
+            }
+            if(status < 0) {
+                return ERROR;
+            }
+            return null;
+        }
+    }
+
+    public enum CommandType {
+
+        NDRANGE_KERNEL(CL_COMMAND_NDRANGE_KERNEL),
+        TASK(CL_COMMAND_TASK),
+        NATIVE_KERNEL(CL_COMMAND_NATIVE_KERNEL),
+        READ_BUFFER(CL_COMMAND_READ_BUFFER),
+        WRITE_BUFFER(CL_COMMAND_WRITE_BUFFER),
+        COPY_BUFFER(CL_COMMAND_COPY_BUFFER),
+        READ_IMAGE(CL_COMMAND_READ_IMAGE),
+        WRITE_IMAGE(CL_COMMAND_WRITE_IMAGE),
+        COPY_IMAGE(CL_COMMAND_COPY_IMAGE),
+        COPY_BUFFER_TO_IMAGE(CL_COMMAND_COPY_BUFFER_TO_IMAGE),
+        COPY_IMAGE_TO_BUFFER(CL_COMMAND_COPY_IMAGE_TO_BUFFER),
+        MAP_BUFFER(CL_COMMAND_MAP_BUFFER),
+        MAP_IMAGE(CL_COMMAND_MAP_IMAGE),
+        UNMAP_MEM_OBJECT(CL_COMMAND_UNMAP_MEM_OBJECT),
+        MARKER(CL_COMMAND_MARKER),
+        ACQUIRE_GL_OBJECTS(CL_COMMAND_ACQUIRE_GL_OBJECTS),
+        RELEASE_GL_OBJECTS(CL_COMMAND_RELEASE_GL_OBJECTS);
+
+        /**
+         * Value of wrapped OpenCL command type.
+         */
+        public final int TYPE;
+
+        private CommandType(int type) {
+            this.TYPE = type;
+        }
+
+        public static CommandType valueOf(int commandType) {
+            CommandType[] values = CommandType.values();
+            for (CommandType value : values) {
+                if(value.TYPE == commandType)
+                    return value;
+            }
+            return null;
+        }
+
+    }
+
+}
diff --git a/src/com/jogamp/opencl/CLEventList.java b/src/com/jogamp/opencl/CLEventList.java
new file mode 100644
index 00000000..b9b4cd4b
--- /dev/null
+++ b/src/com/jogamp/opencl/CLEventList.java
@@ -0,0 +1,99 @@
+package com.jogamp.opencl;
+
+import com.jogamp.common.nio.PointerBuffer;
+import java.util.Iterator;
+
+/**
+ * Fixed size list for storing CLEvents.
+ * @author Michael Bien
+ */
+public final class CLEventList implements CLResource, Iterable<CLEvent> {
+
+    private final CLEvent[] events;
+
+    final PointerBuffer IDs;
+    int size;
+
+    public CLEventList(int capacity) {
+        this.events = new CLEvent[capacity];
+        this.IDs = PointerBuffer.allocateDirect(capacity);
+    }
+
+    void createEvent(CLContext context) {
+
+        if(events[size] != null)
+            events[size].release();
+
+        events[size] = new CLEvent(context, IDs.get());
+        size++;
+    }
+
+    /**
+     * Releases all CLEvents in this list.
+     */
+    public void release() {
+        for (int i = 0; i < size; i++) {
+            events[i].release();
+            events[i] = null;
+        }
+        size = 0;
+        IDs.rewind();
+    }
+
+    public void close() {
+        release();
+    }
+
+    public CLEvent getEvent(int index) {
+        if(index >= size)
+            throw new IndexOutOfBoundsException("list contains "+size+" events, can not return event with index "+index);
+        return events[index];
+    }
+
+    /**
+     * Returns the current size of this list.
+     */
+    public int size() {
+        return size;
+    }
+
+    /**
+     * Returns the maximum size of this list.
+     */
+    public int capacity() {
+        return events.length;
+    }
+
+    public Iterator<CLEvent> iterator() {
+        return new EventIterator(events, size);
+    }
+
+    private static class EventIterator implements Iterator<CLEvent> {
+
+        private final CLEvent[] events;
+        private final int size;
+        private int index;
+
+        private EventIterator(CLEvent[] events, int size) {
+            this.events = events;
+            this.size = size;
+        }
+
+        public boolean hasNext() {
+            return index < size;
+        }
+
+        public CLEvent next() {
+            if(hasNext())
+                return events[index++];
+            else
+                return null;
+        }
+
+        public void remove() {
+            throw new UnsupportedOperationException("remove() not supported.");
+        }
+
+    }
+
+}
diff --git a/src/com/jogamp/opencl/CLException.java b/src/com/jogamp/opencl/CLException.java
new file mode 100644
index 00000000..29491bc7
--- /dev/null
+++ b/src/com/jogamp/opencl/CLException.java
@@ -0,0 +1,704 @@
+package com.jogamp.opencl;
+
+import static com.jogamp.opencl.CL.*;
+
+/**
+ * Main Exception type for runtime OpenCL errors and failed function calls (e.g. returning not CL_SUCCESS).
+ * @author Michael Bien
+ */
+public class CLException extends RuntimeException {
+
+    // must be positive
+    private static final long serialVersionUID = 6573520735486076436L;
+
+    public final int errorcode;
+    public final String error;
+
+    private final static String ERROR_CODE_DOC =
+            "http://www.khronos.org/opencl/sdk/1.0/docs/man/xhtml/errors.html";
+
+    public CLException(String message) {
+        super(message);
+        errorcode = 0;
+        error = "none";
+    }
+
+    private CLException(int errorcode, String errorStr, String message) {
+        super(message + "\nerror: " + errorStr + " (man page: "+ERROR_CODE_DOC+")");
+        this.error = errorStr;
+        this.errorcode = errorcode;
+    }
+
+    /**
+     * Throws a CLException when <code>status != CL_SUCCESS</code>.
+     */
+    public static void checkForError(int status, String message) {
+        if(status != CL_SUCCESS) {
+            CLException ex = newException(status, message);
+            ex.fillInStackTrace();
+            throw ex;
+        }
+    }
+
+    /**
+     * Returns a CLException specific to the error code.
+     */
+    public static CLException newException(int status, String message) {
+        CLException specificEx = createSpecificException(status, message);
+        if(specificEx != null) {
+            specificEx.fillInStackTrace();
+            return specificEx;
+        }
+        return new CLException(status, "unknown", "unknown cause: code " + status);
+    }
+
+    /**
+     * Returns a human readable String for the OpenCL error code.
+     */
+    public String getCLErrorString() {
+        return error;
+    }
+
+
+     // - - - generated code do not edit - - -
+
+    /**
+     * Returns a human readable String for the OpenCL error code or null if not known.
+     */
+    public static String resolveErrorCode(int error) {
+        switch(error) {
+            case CL_DEVICE_NOT_FOUND:                       return "CL_DEVICE_NOT_FOUND";
+            case CL_DEVICE_NOT_AVAILABLE:                   return "CL_DEVICE_NOT_AVAILABLE";
+            case CL_COMPILER_NOT_AVAILABLE:                 return "CL_COMPILER_NOT_AVAILABLE";
+            case CL_MEM_OBJECT_ALLOCATION_FAILURE:          return "CL_MEM_OBJECT_ALLOCATION_FAILURE";
+            case CL_OUT_OF_RESOURCES:                       return "CL_OUT_OF_RESOURCES";
+            case CL_OUT_OF_HOST_MEMORY:                     return "CL_OUT_OF_HOST_MEMORY";
+            case CL_PROFILING_INFO_NOT_AVAILABLE:           return "CL_PROFILING_INFO_NOT_AVAILABLE";
+            case CL_MEM_COPY_OVERLAP:                       return "CL_MEM_COPY_OVERLAP";
+            case CL_IMAGE_FORMAT_MISMATCH:                  return "CL_IMAGE_FORMAT_MISMATCH";
+            case CL_IMAGE_FORMAT_NOT_SUPPORTED:             return "CL_IMAGE_FORMAT_NOT_SUPPORTED";
+            case CL_BUILD_PROGRAM_FAILURE:                  return "CL_BUILD_PROGRAM_FAILURE";
+            case CL_MAP_FAILURE:                            return "CL_MAP_FAILURE";
+            case CL_INVALID_VALUE:                          return "CL_INVALID_VALUE";
+            case CL_INVALID_DEVICE_TYPE:                    return "CL_INVALID_DEVICE_TYPE";
+            case CL_INVALID_PLATFORM:                       return "CL_INVALID_PLATFORM";
+            case CL_INVALID_DEVICE:                         return "CL_INVALID_DEVICE";
+            case CL_INVALID_CONTEXT:                        return "CL_INVALID_CONTEXT";
+            case CL_INVALID_QUEUE_PROPERTIES:               return "CL_INVALID_QUEUE_PROPERTIES";
+            case CL_INVALID_COMMAND_QUEUE:                  return "CL_INVALID_COMMAND_QUEUE";
+            case CL_INVALID_HOST_PTR:                       return "CL_INVALID_HOST_PTR";
+            case CL_INVALID_MEM_OBJECT:                     return "CL_INVALID_MEM_OBJECT";
+            case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR:        return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR";
+            case CL_INVALID_IMAGE_SIZE:                     return "CL_INVALID_IMAGE_SIZE";
+            case CL_INVALID_SAMPLER:                        return "CL_INVALID_SAMPLER";
+            case CL_INVALID_BINARY:                         return "CL_INVALID_BINARY";
+            case CL_INVALID_BUILD_OPTIONS:                  return "CL_INVALID_BUILD_OPTIONS";
+            case CL_INVALID_PROGRAM:                        return "CL_INVALID_PROGRAM";
+            case CL_INVALID_PROGRAM_EXECUTABLE:             return "CL_INVALID_PROGRAM_EXECUTABLE";
+            case CL_INVALID_KERNEL_NAME:                    return "CL_INVALID_KERNEL_NAME";
+            case CL_INVALID_KERNEL_DEFINITION:              return "CL_INVALID_KERNEL_DEFINITION";
+            case CL_INVALID_KERNEL:                         return "CL_INVALID_KERNEL";
+            case CL_INVALID_ARG_INDEX:                      return "CL_INVALID_ARG_INDEX";
+            case CL_INVALID_ARG_VALUE:                      return "CL_INVALID_ARG_VALUE";
+            case CL_INVALID_ARG_SIZE:                       return "CL_INVALID_ARG_SIZE";
+            case CL_INVALID_KERNEL_ARGS:                    return "CL_INVALID_KERNEL_ARGS";
+            case CL_INVALID_WORK_DIMENSION:                 return "CL_INVALID_WORK_DIMENSION";
+            case CL_INVALID_WORK_GROUP_SIZE:                return "CL_INVALID_WORK_GROUP_SIZE";
+            case CL_INVALID_WORK_ITEM_SIZE:                 return "CL_INVALID_WORK_ITEM_SIZE";
+            case CL_INVALID_GLOBAL_OFFSET:                  return "CL_INVALID_GLOBAL_OFFSET";
+            case CL_INVALID_EVENT_WAIT_LIST:                return "CL_INVALID_EVENT_WAIT_LIST";
+            case CL_INVALID_EVENT:                          return "CL_INVALID_EVENT";
+            case CL_INVALID_OPERATION:                      return "CL_INVALID_OPERATION";
+            case CL_INVALID_GL_OBJECT:                      return "CL_INVALID_GL_OBJECT";
+            case CL_INVALID_BUFFER_SIZE:                    return "CL_INVALID_BUFFER_SIZE";
+            case CL_INVALID_MIP_LEVEL:                      return "CL_INVALID_MIP_LEVEL";
+            case CL_INVALID_GLOBAL_WORK_SIZE:               return "CL_INVALID_GLOBAL_WORK_SIZE";
+            case CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR:    return "CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR";
+            case CL_PLATFORM_NOT_FOUND_KHR:                 return "CL_PLATFORM_NOT_FOUND_KHR";
+            default: return null;
+        }
+    }
+
+    private static CLException createSpecificException(int error, String message) {
+        switch(error) {
+            case CL_DEVICE_NOT_FOUND:                       return new CLDeviceNotFoundException(message);
+            case CL_DEVICE_NOT_AVAILABLE:                   return new CLDeviceNotAvailableException(message);
+            case CL_COMPILER_NOT_AVAILABLE:                 return new CLCompilerNotAvailableException(message);
+            case CL_MEM_OBJECT_ALLOCATION_FAILURE:          return new CLMemObjectAllocationFailureException(message);
+            case CL_OUT_OF_RESOURCES:                       return new CLOutOfResourcesException(message);
+            case CL_OUT_OF_HOST_MEMORY:                     return new CLOutOfHostMemoryException(message);
+            case CL_PROFILING_INFO_NOT_AVAILABLE:           return new CLProfilingInfoNotAvailableException(message);
+            case CL_MEM_COPY_OVERLAP:                       return new CLMemCopyOverlapException(message);
+            case CL_IMAGE_FORMAT_MISMATCH:                  return new CLImageFormatMismatchException(message);
+            case CL_IMAGE_FORMAT_NOT_SUPPORTED:             return new CLImageFormatNotSupportedException(message);
+            case CL_BUILD_PROGRAM_FAILURE:                  return new CLBuildProgramFailureException(message);
+            case CL_MAP_FAILURE:                            return new CLMapFailureException(message);
+            case CL_INVALID_VALUE:                          return new CLInvalidValueException(message);
+            case CL_INVALID_DEVICE_TYPE:                    return new CLInvalidDeviceTypeException(message);
+            case CL_INVALID_PLATFORM:                       return new CLInvalidPlatformException(message);
+            case CL_INVALID_DEVICE:                         return new CLInvalidDeviceException(message);
+            case CL_INVALID_CONTEXT:                        return new CLInvalidContextException(message);
+            case CL_INVALID_QUEUE_PROPERTIES:               return new CLInvalidQueuePropertiesException(message);
+            case CL_INVALID_COMMAND_QUEUE:                  return new CLInvalidCommandQueueException(message);
+            case CL_INVALID_HOST_PTR:                       return new CLInvalidHostPtrException(message);
+            case CL_INVALID_MEM_OBJECT:                     return new CLInvalidMemObjectException(message);
+            case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR:        return new CLInvalidImageFormatDescriptorException(message);
+            case CL_INVALID_IMAGE_SIZE:                     return new CLInvalidImageSizeException(message);
+            case CL_INVALID_SAMPLER:                        return new CLInvalidSamplerException(message);
+            case CL_INVALID_BINARY:                         return new CLInvalidBinaryException(message);
+            case CL_INVALID_BUILD_OPTIONS:                  return new CLInvalidBuildOptionsException(message);
+            case CL_INVALID_PROGRAM:                        return new CLInvalidProgramException(message);
+            case CL_INVALID_PROGRAM_EXECUTABLE:             return new CLInvalidProgramExecutableException(message);
+            case CL_INVALID_KERNEL_NAME:                    return new CLInvalidKernelNameException(message);
+            case CL_INVALID_KERNEL_DEFINITION:              return new CLInvalidKernelDefinitionException(message);
+            case CL_INVALID_KERNEL:                         return new CLInvalidKernelException(message);
+            case CL_INVALID_ARG_INDEX:                      return new CLInvalidArgIndexException(message);
+            case CL_INVALID_ARG_VALUE:                      return new CLInvalidArgValueException(message);
+            case CL_INVALID_ARG_SIZE:                       return new CLInvalidArgSizeException(message);
+            case CL_INVALID_KERNEL_ARGS:                    return new CLInvalidKernelArgsException(message);
+            case CL_INVALID_WORK_DIMENSION:                 return new CLInvalidWorkDimensionException(message);
+            case CL_INVALID_WORK_GROUP_SIZE:                return new CLInvalidWorkGroupSizeException(message);
+            case CL_INVALID_WORK_ITEM_SIZE:                 return new CLInvalidWorkItemSizeException(message);
+            case CL_INVALID_GLOBAL_OFFSET:                  return new CLInvalidGlobalOffsetException(message);
+            case CL_INVALID_EVENT_WAIT_LIST:                return new CLInvalidEventWaitListException(message);
+            case CL_INVALID_EVENT:                          return new CLInvalidEventException(message);
+            case CL_INVALID_OPERATION:                      return new CLInvalidOperationException(message);
+            case CL_INVALID_GL_OBJECT:                      return new CLInvalidGLObjectException(message);
+            case CL_INVALID_BUFFER_SIZE:                    return new CLInvalidBufferSizeException(message);
+            case CL_INVALID_MIP_LEVEL:                      return new CLInvalidMipLevelException(message);
+            case CL_INVALID_GLOBAL_WORK_SIZE:               return new CLInvalidGlobalWorkSizeException(message);
+            case CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR:    return new CLInvalidGLSharegroupReferenceKhrException(message);
+            case CL_PLATFORM_NOT_FOUND_KHR:                 return new CLPlatformNotFoundKhrException(message);
+            default: return null;
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_DEVICE_NOT_FOUND errors.
+     * @author Michael Bien
+     */
+    public final static class CLDeviceNotFoundException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_DEVICE_NOT_FOUND;
+        public CLDeviceNotFoundException(String message) {
+            super(CL_DEVICE_NOT_FOUND, "CL_DEVICE_NOT_FOUND", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_DEVICE_NOT_AVAILABLE errors.
+     * @author Michael Bien
+     */
+    public final static class CLDeviceNotAvailableException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_DEVICE_NOT_AVAILABLE;
+        public CLDeviceNotAvailableException(String message) {
+            super(CL_DEVICE_NOT_AVAILABLE, "CL_DEVICE_NOT_AVAILABLE", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_COMPILER_NOT_AVAILABLE errors.
+     * @author Michael Bien
+     */
+    public final static class CLCompilerNotAvailableException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_COMPILER_NOT_AVAILABLE;
+        public CLCompilerNotAvailableException(String message) {
+            super(CL_COMPILER_NOT_AVAILABLE, "CL_COMPILER_NOT_AVAILABLE", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_MEM_OBJECT_ALLOCATION_FAILURE errors.
+     * @author Michael Bien
+     */
+    public final static class CLMemObjectAllocationFailureException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_MEM_OBJECT_ALLOCATION_FAILURE;
+        public CLMemObjectAllocationFailureException(String message) {
+            super(CL_MEM_OBJECT_ALLOCATION_FAILURE, "CL_MEM_OBJECT_ALLOCATION_FAILURE", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_OUT_OF_RESOURCES errors.
+     * @author Michael Bien
+     */
+    public final static class CLOutOfResourcesException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_OUT_OF_RESOURCES;
+        public CLOutOfResourcesException(String message) {
+            super(CL_OUT_OF_RESOURCES, "CL_OUT_OF_RESOURCES", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_OUT_OF_HOST_MEMORY errors.
+     * @author Michael Bien
+     */
+    public final static class CLOutOfHostMemoryException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_OUT_OF_HOST_MEMORY;
+        public CLOutOfHostMemoryException(String message) {
+            super(CL_OUT_OF_HOST_MEMORY, "CL_OUT_OF_HOST_MEMORY", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_PROFILING_INFO_NOT_AVAILABLE errors.
+     * @author Michael Bien
+     */
+    public final static class CLProfilingInfoNotAvailableException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_PROFILING_INFO_NOT_AVAILABLE;
+        public CLProfilingInfoNotAvailableException(String message) {
+            super(CL_PROFILING_INFO_NOT_AVAILABLE, "CL_PROFILING_INFO_NOT_AVAILABLE", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_MEM_COPY_OVERLAP errors.
+     * @author Michael Bien
+     */
+    public final static class CLMemCopyOverlapException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_MEM_COPY_OVERLAP;
+        public CLMemCopyOverlapException(String message) {
+            super(CL_MEM_COPY_OVERLAP, "CL_MEM_COPY_OVERLAP", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_IMAGE_FORMAT_MISMATCH errors.
+     * @author Michael Bien
+     */
+    public final static class CLImageFormatMismatchException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_IMAGE_FORMAT_MISMATCH;
+        public CLImageFormatMismatchException(String message) {
+            super(CL_IMAGE_FORMAT_MISMATCH, "CL_IMAGE_FORMAT_MISMATCH", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_IMAGE_FORMAT_NOT_SUPPORTED errors.
+     * @author Michael Bien
+     */
+    public final static class CLImageFormatNotSupportedException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_IMAGE_FORMAT_NOT_SUPPORTED;
+        public CLImageFormatNotSupportedException(String message) {
+            super(CL_IMAGE_FORMAT_NOT_SUPPORTED, "CL_IMAGE_FORMAT_NOT_SUPPORTED", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_BUILD_PROGRAM_FAILURE errors.
+     * @author Michael Bien
+     */
+    public final static class CLBuildProgramFailureException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_BUILD_PROGRAM_FAILURE;
+        public CLBuildProgramFailureException(String message) {
+            super(CL_BUILD_PROGRAM_FAILURE, "CL_BUILD_PROGRAM_FAILURE", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_MAP_FAILURE errors.
+     * @author Michael Bien
+     */
+    public final static class CLMapFailureException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_MAP_FAILURE;
+        public CLMapFailureException(String message) {
+            super(CL_MAP_FAILURE, "CL_MAP_FAILURE", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_VALUE errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidValueException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_VALUE;
+        public CLInvalidValueException(String message) {
+            super(CL_INVALID_VALUE, "CL_INVALID_VALUE", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_DEVICE_TYPE errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidDeviceTypeException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_DEVICE_TYPE;
+        public CLInvalidDeviceTypeException(String message) {
+            super(CL_INVALID_DEVICE_TYPE, "CL_INVALID_DEVICE_TYPE", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_PLATFORM errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidPlatformException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_PLATFORM;
+        public CLInvalidPlatformException(String message) {
+            super(CL_INVALID_PLATFORM, "CL_INVALID_PLATFORM", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_DEVICE errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidDeviceException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_DEVICE;
+        public CLInvalidDeviceException(String message) {
+            super(CL_INVALID_DEVICE, "CL_INVALID_DEVICE", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_CONTEXT errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidContextException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_CONTEXT;
+        public CLInvalidContextException(String message) {
+            super(CL_INVALID_CONTEXT, "CL_INVALID_CONTEXT", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_QUEUE_PROPERTIES errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidQueuePropertiesException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_QUEUE_PROPERTIES;
+        public CLInvalidQueuePropertiesException(String message) {
+            super(CL_INVALID_QUEUE_PROPERTIES, "CL_INVALID_QUEUE_PROPERTIES", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_COMMAND_QUEUE errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidCommandQueueException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_COMMAND_QUEUE;
+        public CLInvalidCommandQueueException(String message) {
+            super(CL_INVALID_COMMAND_QUEUE, "CL_INVALID_COMMAND_QUEUE", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_HOST_PTR errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidHostPtrException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_HOST_PTR;
+        public CLInvalidHostPtrException(String message) {
+            super(CL_INVALID_HOST_PTR, "CL_INVALID_HOST_PTR", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_MEM_OBJECT errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidMemObjectException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_MEM_OBJECT;
+        public CLInvalidMemObjectException(String message) {
+            super(CL_INVALID_MEM_OBJECT, "CL_INVALID_MEM_OBJECT", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_IMAGE_FORMAT_DESCRIPTOR errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidImageFormatDescriptorException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_IMAGE_FORMAT_DESCRIPTOR;
+        public CLInvalidImageFormatDescriptorException(String message) {
+            super(CL_INVALID_IMAGE_FORMAT_DESCRIPTOR, "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_IMAGE_SIZE errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidImageSizeException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_IMAGE_SIZE;
+        public CLInvalidImageSizeException(String message) {
+            super(CL_INVALID_IMAGE_SIZE, "CL_INVALID_IMAGE_SIZE", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_SAMPLER errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidSamplerException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_SAMPLER;
+        public CLInvalidSamplerException(String message) {
+            super(CL_INVALID_SAMPLER, "CL_INVALID_SAMPLER", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_BINARY errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidBinaryException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_BINARY;
+        public CLInvalidBinaryException(String message) {
+            super(CL_INVALID_BINARY, "CL_INVALID_BINARY", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_BUILD_OPTIONS errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidBuildOptionsException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_BUILD_OPTIONS;
+        public CLInvalidBuildOptionsException(String message) {
+            super(CL_INVALID_BUILD_OPTIONS, "CL_INVALID_BUILD_OPTIONS", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_PROGRAM errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidProgramException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_PROGRAM;
+        public CLInvalidProgramException(String message) {
+            super(CL_INVALID_PROGRAM, "CL_INVALID_PROGRAM", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_PROGRAM_EXECUTABLE errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidProgramExecutableException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_PROGRAM_EXECUTABLE;
+        public CLInvalidProgramExecutableException(String message) {
+            super(CL_INVALID_PROGRAM_EXECUTABLE, "CL_INVALID_PROGRAM_EXECUTABLE", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_KERNEL_NAME errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidKernelNameException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_KERNEL_NAME;
+        public CLInvalidKernelNameException(String message) {
+            super(CL_INVALID_KERNEL_NAME, "CL_INVALID_KERNEL_NAME", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_KERNEL_DEFINITION errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidKernelDefinitionException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_KERNEL_DEFINITION;
+        public CLInvalidKernelDefinitionException(String message) {
+            super(CL_INVALID_KERNEL_DEFINITION, "CL_INVALID_KERNEL_DEFINITION", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_KERNEL errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidKernelException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_KERNEL;
+        public CLInvalidKernelException(String message) {
+            super(CL_INVALID_KERNEL, "CL_INVALID_KERNEL", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_ARG_INDEX errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidArgIndexException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_ARG_INDEX;
+        public CLInvalidArgIndexException(String message) {
+            super(CL_INVALID_ARG_INDEX, "CL_INVALID_ARG_INDEX", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_ARG_VALUE errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidArgValueException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_ARG_VALUE;
+        public CLInvalidArgValueException(String message) {
+            super(CL_INVALID_ARG_VALUE, "CL_INVALID_ARG_VALUE", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_ARG_SIZE errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidArgSizeException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_ARG_SIZE;
+        public CLInvalidArgSizeException(String message) {
+            super(CL_INVALID_ARG_SIZE, "CL_INVALID_ARG_SIZE", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_KERNEL_ARGS errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidKernelArgsException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_KERNEL_ARGS;
+        public CLInvalidKernelArgsException(String message) {
+            super(CL_INVALID_KERNEL_ARGS, "CL_INVALID_KERNEL_ARGS", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_WORK_DIMENSION errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidWorkDimensionException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_WORK_DIMENSION;
+        public CLInvalidWorkDimensionException(String message) {
+            super(CL_INVALID_WORK_DIMENSION, "CL_INVALID_WORK_DIMENSION", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_WORK_GROUP_SIZE errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidWorkGroupSizeException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_WORK_GROUP_SIZE;
+        public CLInvalidWorkGroupSizeException(String message) {
+            super(CL_INVALID_WORK_GROUP_SIZE, "CL_INVALID_WORK_GROUP_SIZE", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_WORK_ITEM_SIZE errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidWorkItemSizeException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_WORK_ITEM_SIZE;
+        public CLInvalidWorkItemSizeException(String message) {
+            super(CL_INVALID_WORK_ITEM_SIZE, "CL_INVALID_WORK_ITEM_SIZE", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_GLOBAL_OFFSET errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidGlobalOffsetException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_GLOBAL_OFFSET;
+        public CLInvalidGlobalOffsetException(String message) {
+            super(CL_INVALID_GLOBAL_OFFSET, "CL_INVALID_GLOBAL_OFFSET", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_EVENT_WAIT_LIST errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidEventWaitListException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_EVENT_WAIT_LIST;
+        public CLInvalidEventWaitListException(String message) {
+            super(CL_INVALID_EVENT_WAIT_LIST, "CL_INVALID_EVENT_WAIT_LIST", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_EVENT errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidEventException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_EVENT;
+        public CLInvalidEventException(String message) {
+            super(CL_INVALID_EVENT, "CL_INVALID_EVENT", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_OPERATION errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidOperationException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_OPERATION;
+        public CLInvalidOperationException(String message) {
+            super(CL_INVALID_OPERATION, "CL_INVALID_OPERATION", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_GL_OBJECT errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidGLObjectException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_GL_OBJECT;
+        public CLInvalidGLObjectException(String message) {
+            super(CL_INVALID_GL_OBJECT, "CL_INVALID_GL_OBJECT", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_BUFFER_SIZE errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidBufferSizeException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_BUFFER_SIZE;
+        public CLInvalidBufferSizeException(String message) {
+            super(CL_INVALID_BUFFER_SIZE, "CL_INVALID_BUFFER_SIZE", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_MIP_LEVEL errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidMipLevelException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_MIP_LEVEL;
+        public CLInvalidMipLevelException(String message) {
+            super(CL_INVALID_MIP_LEVEL, "CL_INVALID_MIP_LEVEL", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_GLOBAL_WORK_SIZE errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidGlobalWorkSizeException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_GLOBAL_WORK_SIZE;
+        public CLInvalidGlobalWorkSizeException(String message) {
+            super(CL_INVALID_GLOBAL_WORK_SIZE, "CL_INVALID_GLOBAL_WORK_SIZE", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR errors.
+     * @author Michael Bien
+     */
+    public final static class CLInvalidGLSharegroupReferenceKhrException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR;
+        public CLInvalidGLSharegroupReferenceKhrException(String message) {
+            super(CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR, "CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR", message);
+        }
+    }
+
+    /**
+     * {@link CLException} thrown on CL.CL_PLATFORM_NOT_FOUND_KHR errors.
+     * @author Michael Bien
+     */
+    public final static class CLPlatformNotFoundKhrException extends CLException {
+        private static final long serialVersionUID = CLException.serialVersionUID+CL_PLATFORM_NOT_FOUND_KHR;
+        public CLPlatformNotFoundKhrException(String message) {
+            super(CL_PLATFORM_NOT_FOUND_KHR, "CL_PLATFORM_NOT_FOUND_KHR", message);
+        }
+    }
+
+}
\ No newline at end of file
diff --git a/src/com/jogamp/opencl/CLImage.java b/src/com/jogamp/opencl/CLImage.java
new file mode 100644
index 00000000..f8916fdb
--- /dev/null
+++ b/src/com/jogamp/opencl/CLImage.java
@@ -0,0 +1,90 @@
+package com.jogamp.opencl;
+
+import com.jogamp.common.nio.Int64Buffer;
+import java.nio.Buffer;
+
+import static com.jogamp.opencl.CL.*;
+
+/**
+ *
+ * @author Michael Bien
+ */
+public abstract class CLImage<B extends Buffer> extends CLMemory<B>  {
+
+    protected CLImageFormat format;
+
+    final CLInfoAccessor imageInfo;
+
+    public final int width;
+    public final int height;
+
+    protected CLImage(CLContext context, B directBuffer, CLImageFormat format, int width, int height, long id, int flags) {
+        this(context, directBuffer, format, new CLImageInfoAccessor(context.cl, id), width, height, id, flags);
+    }
+
+    protected CLImage(CLContext context, B directBuffer, CLImageFormat format, CLImageInfoAccessor accessor, int width, int height, long id, int flags) {
+        super(context, directBuffer, id, flags);
+        this.imageInfo = accessor;
+        this.format = format;
+        this.width = width;
+        this.height = height;
+    }
+
+    protected static CLImageFormat createUninitializedImageFormat() {
+        return new CLImageFormat();
+    }
+
+    /**
+     * Returns the image format descriptor specified when image was created.
+     */
+    public CLImageFormat getFormat() {
+        return format;
+    }
+
+    /**
+     * Returns the size of each element of the image memory object given by image.
+     * An element is made up of n channels. The value of n is given in {@link CLImageFormat} descriptor.
+     */
+    public int getElementSize() {
+        return (int)imageInfo.getLong(CL_IMAGE_ELEMENT_SIZE);
+    }
+
+    /**
+     * Returns the size in bytes of a row of elements of the image object given by image.
+     */
+    public int getRowPitch() {
+        return (int)imageInfo.getLong(CL_IMAGE_ROW_PITCH);
+    }
+
+    /**
+     * Returns width of this image in pixels.
+     */
+    public int getWidth() {
+        return width;
+    }
+
+    /**
+     * Returns the height of this image in pixels.
+     */
+    public int getHeight() {
+        return height;
+    }
+
+
+    protected final static class CLImageInfoAccessor extends CLInfoAccessor {
+
+        private final long id;
+        private final CL cl;
+
+        public CLImageInfoAccessor(CL cl, long id) {
+            this.cl = cl;
+            this.id = id;
+        }
+        @Override
+        public int getInfo(int name, long valueSize, Buffer value, Int64Buffer valueSizeRet) {
+            return cl.clGetImageInfo(id, name, valueSize, value, valueSizeRet);
+        }
+    }
+
+
+}
diff --git a/src/com/jogamp/opencl/CLImage2d.java b/src/com/jogamp/opencl/CLImage2d.java
new file mode 100644
index 00000000..7cf04cf1
--- /dev/null
+++ b/src/com/jogamp/opencl/CLImage2d.java
@@ -0,0 +1,46 @@
+package com.jogamp.opencl;
+
+import com.jogamp.common.nio.Buffers;
+import java.nio.Buffer;
+import java.nio.IntBuffer;
+
+import static com.jogamp.opencl.CLException.*;
+
+/**
+ *
+ * @author Michael Bien
+ */
+public class CLImage2d<B extends Buffer> extends CLImage<B> {
+
+    private CLImage2d(CLContext context, B directBuffer, CLImageFormat format, int width, int height, long id, int flags) {
+        super(context, directBuffer, format, width, height, id, flags);
+    }
+    
+    protected CLImage2d(CLContext context, B directBuffer, CLImageFormat format, CLImageInfoAccessor accessor, int width, int height, long id, int flags) {
+        super(context, directBuffer, format, accessor, width, height, id, flags);
+    }
+
+    static <B extends Buffer> CLImage2d<B> createImage(CLContext context, B directBuffer,
+            int width, int height, int rowPitch, CLImageFormat format, int flags) {
+
+        CL cl = context.cl;
+        IntBuffer err = Buffers.newDirectByteBuffer(4).asIntBuffer();
+
+        long id = cl.clCreateImage2D(context.ID, flags, format.getFormatImpl(), width, height, rowPitch, directBuffer, err);
+        checkForError(err.get(), "can not create 2d image");
+
+        return new CLImage2d<B>(context, directBuffer, format, width, height, id, flags);
+    }
+
+    @Override
+    public <T extends Buffer> CLImage2d<T> cloneWith(T directBuffer) {
+        return new CLImage2d<T>(context, directBuffer, format, width, height, ID, FLAGS);
+    }
+
+
+    @Override
+    public String toString() {
+        return "CLImage2d [id: " + ID+" width: "+width+" height: "+height+"]";
+    }
+
+}
diff --git a/src/com/jogamp/opencl/CLImage3d.java b/src/com/jogamp/opencl/CLImage3d.java
new file mode 100644
index 00000000..fd127864
--- /dev/null
+++ b/src/com/jogamp/opencl/CLImage3d.java
@@ -0,0 +1,64 @@
+package com.jogamp.opencl;
+
+import com.jogamp.common.nio.Buffers;
+import java.nio.Buffer;
+import java.nio.IntBuffer;
+
+import static com.jogamp.opencl.CL.*;
+import static com.jogamp.opencl.CLException.*;
+
+/**
+ *
+ * @author Michael Bien
+ */
+public class CLImage3d<B extends Buffer> extends CLImage<B> {
+
+    public final int depth;
+
+    private CLImage3d(CLContext context, B directBuffer, CLImageFormat format, int width, int height, int depth, long id, int flags) {
+        super(context, directBuffer, format, width, height, id, flags);
+        this.depth = depth;
+    }
+
+    protected CLImage3d(CLContext context, B directBuffer, CLImageFormat format, CLImageInfoAccessor accessor, int width, int height, int depth, long id, int flags) {
+        super(context, directBuffer, format, accessor, width, height, id, flags);
+        this.depth = depth;
+    }
+    
+
+    static <B extends Buffer> CLImage3d<B> createImage(CLContext context, B directBuffer,
+            int width, int height, int depth, int rowPitch, int slicePitch, CLImageFormat format, int flags) {
+
+        CL cl = context.cl;
+        IntBuffer err = Buffers.newDirectByteBuffer(4).asIntBuffer();
+
+        long id = cl.clCreateImage3D(context.ID, flags, format.getFormatImpl(), width, height, depth, rowPitch, slicePitch, directBuffer, err);
+        checkForError(err.get(), "can not create 2d image");
+
+        return new CLImage3d<B>(context, directBuffer, format, width, height, depth, id, flags);
+    }
+
+    @Override
+    public <T extends Buffer> CLImage3d<T> cloneWith(T directBuffer) {
+        return new CLImage3d<T>(context, directBuffer, format, width, height, depth, ID, FLAGS);
+    }
+
+    /**
+     * Returns the size in bytes of a 2D slice of this 3D image.
+     */
+    public int getSlicePitch() {
+        return (int)imageInfo.getLong(CL_IMAGE_SLICE_PITCH);
+    }
+
+    /**
+     * Returns the depth of this image in pixels.
+     */
+    public int getDepth() {
+        return depth;
+    }
+
+    @Override
+    public String toString() {
+        return "CLImage3d [id: " + ID+" width: "+width+" height: "+height+" depth: "+depth+"]";
+    }
+}
diff --git a/src/com/jogamp/opencl/CLImageFormat.java b/src/com/jogamp/opencl/CLImageFormat.java
new file mode 100644
index 00000000..20f2c4b1
--- /dev/null
+++ b/src/com/jogamp/opencl/CLImageFormat.java
@@ -0,0 +1,280 @@
+package com.jogamp.opencl;
+
+import com.jogamp.opencl.impl.CLImageFormatImpl;
+
+import static com.jogamp.opencl.CL.*;
+
+/**
+ *
+ * @author Michael Bien
+ */
+public final class CLImageFormat {
+
+    private final CLImageFormatImpl format = CLImageFormatImpl.create();
+
+    CLImageFormat() {
+    }
+
+    public CLImageFormat(ChannelOrder order, ChannelType type) {
+        setImageChannelOrder(order);
+        setImageChannelDataType(type);
+    }
+
+    public CLImageFormat setImageChannelOrder(ChannelOrder order) {
+        format.setImageChannelOrder(order.ORDER);
+        return this;
+    }
+
+    public CLImageFormat setImageChannelDataType(ChannelType type) {
+        format.setImageChannelDataType(type.TYPE);
+        return this;
+    }
+
+    public ChannelOrder getImageChannelOrder() {
+        return ChannelOrder.valueOf(format.getImageChannelOrder());
+    }
+
+    public ChannelType getImageChannelDataType() {
+        return ChannelType.valueOf(format.getImageChannelDataType());
+    }
+
+    /**
+     * Returns the struct accessor for the cl_image_format struct.
+     */
+    public CLImageFormatImpl getFormatImpl() {
+        return format;
+    }
+
+    /**
+     * Specifies the number of channels and the channel layout i.e. the memory
+     * layout in which channels are stored in the image.
+     */
+    public enum ChannelOrder {
+                
+        /**
+         *
+         */
+        R(CL_R),
+
+        /**
+         *
+         */
+        A(CL_A),
+
+        /**
+         *
+         */
+        RG(CL_RG),
+
+        /**
+         *
+         */
+        RA(CL_RA),
+
+        /**
+         * This format can only be used if channel data type is one of the following values:
+         * {@link ChannelType#UNORM_SHORT_565}, {@link ChannelType#UNORM_SHORT_555}
+         * or {@link ChannelType#UNORM_INT_101010}.
+         */
+        RGB(CL_RGB),
+
+        /**
+         *
+         */
+        RGBA(CL_RGBA),
+
+        /**
+         * This format can only be used if channel data type is one of the following values:
+         * {@link ChannelType#UNORM_INT8}, {@link ChannelType#SNORM_INT8}, {@link ChannelType#SIGNED_INT8}
+         * or {@link ChannelType#UNSIGNED_INT8}.
+         */
+        ARGB(CL_ARGB),
+
+        /**
+         * @see #ARGB
+         */
+        BGRA(CL_BGRA),
+
+        /**
+         * This format can only be used if channel data type is one of the following values:
+         * {@link ChannelType#UNORM_INT8}, {@link ChannelType#UNORM_INT16}, {@link ChannelType#SNORM_INT8},
+         * {@link ChannelType#SNORM_INT16}, {@link ChannelType#HALF_FLOAT}, or {@link ChannelType#FLOAT}.
+         */
+        INTENSITY(CL_INTENSITY),
+
+        /**
+         * This format can only be used if channel data type is one of the following values:
+         * {@link ChannelType#UNORM_INT8}, {@link ChannelType#UNORM_INT16}, {@link ChannelType#SNORM_INT8},
+         * {@link ChannelType#SNORM_INT16}, {@link ChannelType#HALF_FLOAT}, or {@link ChannelType#FLOAT}.
+         */
+        LUMINANCE(CL_LUMINANCE);
+        
+        
+        /**
+         * Value of wrapped OpenCL flag.
+         */
+        public final int ORDER;
+
+        private ChannelOrder(int order) {
+            this.ORDER = order;
+        }
+
+        public static ChannelOrder valueOf(int orderFlag) {
+            switch (orderFlag) {
+                case CL_R:
+                    return R;
+                case CL_A:
+                    return A;
+                case CL_INTENSITY:
+                    return INTENSITY;
+                case CL_LUMINANCE:
+                    return LUMINANCE;
+                case CL_RG:
+                    return RG;
+                case CL_RA:
+                    return RA;
+                case CL_RGB:
+                    return RGB;
+                case CL_RGBA:
+                    return RGBA;
+                case CL_ARGB:
+                    return ARGB;
+                case CL_BGRA:
+                    return BGRA;
+            }
+            return null;
+        }
+
+    }
+
+
+    /**
+     * Describes the size of the channel data type.
+     */
+    public enum ChannelType {
+
+        /**
+         * Each channel component is a normalized signed 8-bit integer value.
+         */
+        SNORM_INT8(CL_SNORM_INT8),
+        
+        /**
+         * Each channel component is a normalized signed 16-bit integer value.
+         */
+        SNORM_INT16(CL_SNORM_INT16),
+
+        /**
+         * Each channel component is a normalized unsigned 8-bit integer value.
+         */
+        UNORM_INT8(CL_UNORM_INT8),
+
+        /**
+         * Each channel component is a normalized unsigned 16-bit integer value.
+         */
+        UNORM_INT16(CL_UNORM_INT16),
+
+        /**
+         * Represents a normalized 5-6-5 3-channel RGB image. The channel order must
+         * be {@link ChannelOrder#RGB}.
+         */
+        UNORM_SHORT_565(CL_UNORM_SHORT_565),
+
+        /**
+         * Represents a normalized x-5-5-5 4-channel xRGB image. The channel order must
+         * be {@link ChannelOrder#RGB}.
+         */
+        UNORM_SHORT_555(CL_UNORM_SHORT_555),
+
+        /**
+         * Represents a normalized x-10-10-10 4-channel xRGB image. The channel order
+         * must be {@link ChannelOrder#RGB}.
+         */
+        UNORM_INT_101010(CL_UNORM_INT_101010),
+
+        /**
+         * Each channel component is an unnormalized signed 8-bit integer value.
+         */
+        SIGNED_INT8(CL_SIGNED_INT8),
+
+        /**
+         * Each channel component is an unnormalized signed 16-bit integer value.
+         */
+        SIGNED_INT16(CL_SIGNED_INT16),
+
+        /**
+         * Each channel component is an unnormalized signed 32-bit integer value.
+         */
+        SIGNED_INT32(CL_SIGNED_INT32),
+
+        /**
+         * Each channel component is an unnormalized unsigned 8-bit integer value.
+         */
+        UNSIGNED_INT8(CL_UNSIGNED_INT8),
+
+        /**
+         * Each channel component is an unnormalized unsigned 16-bit integer value.
+         */
+        UNSIGNED_INT16(CL_UNSIGNED_INT16),
+
+        /**
+         * Each channel component is an unnormalized unsigned 32-bit integer value.
+         */
+        UNSIGNED_INT32(CL_UNSIGNED_INT32),
+
+        /**
+         * Each channel component is a 16-bit half-float value.
+         */
+        HALF_FLOAT(CL_HALF_FLOAT),
+
+        /**
+         * Each channel component is a single precision floating-point value.
+         */
+        FLOAT(CL_FLOAT);
+
+        /**
+         * Value of wrapped OpenCL flag.
+         */
+        public final int TYPE;
+
+        private ChannelType(int channel) {
+            this.TYPE = channel;
+        }
+
+        public static ChannelType valueOf(int channelFlag) {
+            switch (channelFlag) {
+                case CL_SNORM_INT8:
+                    return SNORM_INT8;
+                case CL_SNORM_INT16:
+                    return SNORM_INT16;
+                case CL_UNORM_INT8:
+                    return UNORM_INT8;
+                case CL_UNORM_INT16:
+                    return UNORM_INT16;
+                case CL_UNORM_SHORT_565:
+                    return UNORM_SHORT_565;
+                case CL_UNORM_SHORT_555:
+                    return UNORM_SHORT_555;
+                case CL_UNORM_INT_101010:
+                    return UNORM_INT_101010;
+                case CL_SIGNED_INT8:
+                    return SIGNED_INT8;
+                case CL_SIGNED_INT16:
+                    return SIGNED_INT16;
+                case CL_SIGNED_INT32:
+                    return SIGNED_INT32;
+                case CL_UNSIGNED_INT8:
+                    return UNSIGNED_INT8;
+                case CL_UNSIGNED_INT16:
+                    return UNSIGNED_INT16;
+                case CL_UNSIGNED_INT32:
+                    return UNSIGNED_INT32;
+                case CL_HALF_FLOAT:
+                    return HALF_FLOAT;
+                case CL_FLOAT:
+                    return FLOAT;
+            }
+            return null;
+        }
+
+    }
+}
diff --git a/src/com/jogamp/opencl/CLInfoAccessor.java b/src/com/jogamp/opencl/CLInfoAccessor.java
new file mode 100644
index 00000000..56f2318d
--- /dev/null
+++ b/src/com/jogamp/opencl/CLInfoAccessor.java
@@ -0,0 +1,62 @@
+package com.jogamp.opencl;
+
+import com.jogamp.common.nio.Buffers;
+import com.jogamp.common.nio.Int64Buffer;
+import com.jogamp.opencl.util.CLUtil;
+import java.nio.Buffer;
+import java.nio.ByteBuffer;
+
+import static com.jogamp.opencl.CLException.*;
+
+/**
+ * Internal utility for common OpenCL clGetFooInfo calls.
+ * Threadsafe.
+ * @author Michael Bien
+ */
+abstract class CLInfoAccessor {
+
+    protected final static ThreadLocal<ByteBuffer> localBB = new ThreadLocal<ByteBuffer>() {
+
+        @Override
+        protected ByteBuffer initialValue() {
+            return Buffers.newDirectByteBuffer(512);
+        }
+
+    };
+    protected final static ThreadLocal<Int64Buffer> localPB = new ThreadLocal<Int64Buffer>() {
+
+        @Override
+        protected Int64Buffer initialValue() {
+            return Int64Buffer.allocateDirect(1);
+        }
+
+    };
+
+    public final long getLong(int key) {
+
+        ByteBuffer buffer = localBB.get();
+        int ret = getInfo(key, 8, buffer, null);
+        checkForError(ret, "error while asking for info value");
+
+        return buffer.getLong(0);
+    }
+
+    public final String getString(int key) {
+        
+        ByteBuffer buffer = localBB.get();
+        Int64Buffer sizeBuffer = localPB.get();
+        int ret = getInfo(key, buffer.capacity(), buffer, sizeBuffer);
+        checkForError(ret, "error while asking for info string");
+
+        int clSize = (int)sizeBuffer.get(0);
+        byte[] array = new byte[clSize-1]; // last char is always null
+        buffer.get(array).rewind();
+
+        return CLUtil.clString2JavaString(array, clSize);
+
+    }
+
+    protected abstract int getInfo(int name, long valueSize, Buffer value, Int64Buffer valueSizeRet);
+
+
+}
diff --git a/src/com/jogamp/opencl/CLKernel.java b/src/com/jogamp/opencl/CLKernel.java
new file mode 100644
index 00000000..ca7f6786
--- /dev/null
+++ b/src/com/jogamp/opencl/CLKernel.java
@@ -0,0 +1,306 @@
+package com.jogamp.opencl;
+
+import com.jogamp.opencl.util.CLUtil;
+import com.jogamp.common.nio.Buffers;
+import com.jogamp.common.os.Platform;
+import com.jogamp.common.nio.Int64Buffer;
+import java.nio.Buffer;
+import java.nio.ByteBuffer;
+import java.nio.ByteOrder;
+
+import static com.jogamp.opencl.CLException.*;
+import static com.jogamp.opencl.CL.*;
+
+/**
+ * High level abstraction for an OpenCL Kernel.
+ * A kernel is a function declared in a program. A kernel is identified by the <code>kernel</code> qualifier
+ * applied to any function in a program. A kernel object encapsulates the specific <code>kernel</code>
+ * function declared in a program and the argument values to be used when executing this
+ * <code>kernel</code> function.
+ * CLKernel is not threadsafe.
+ * @see CLProgram#createCLKernel(java.lang.String)
+ * @see CLProgram#createCLKernels()
+ * @author Michael Bien
+ */
+public class CLKernel extends CLObject implements CLResource, Cloneable {
+
+    public final String name;
+    public final int numArgs;
+
+    private final CLProgram program;
+
+    private final ByteBuffer buffer;
+
+    private int argIndex;
+    private boolean force32BitArgs;
+
+    CLKernel(CLProgram program, long id) {
+        super(program.getContext(), id);
+        this.program = program;
+        this.buffer = Buffers.newDirectByteBuffer(8);
+
+        Int64Buffer size = Int64Buffer.allocateDirect(1);
+
+        // get function name
+        int ret = cl.clGetKernelInfo(ID, CL_KERNEL_FUNCTION_NAME, 0, null, size);
+        checkForError(ret, "error while asking for kernel function name");
+
+        ByteBuffer bb = ByteBuffer.allocateDirect((int)size.get(0)).order(ByteOrder.nativeOrder());
+
+        ret = cl.clGetKernelInfo(ID, CL_KERNEL_FUNCTION_NAME, bb.capacity(), bb, null);
+        checkForError(ret, "error while asking for kernel function name");
+
+        this.name = CLUtil.clString2JavaString(bb, bb.capacity());
+
+        // get number of arguments
+        ret = cl.clGetKernelInfo(ID, CL_KERNEL_NUM_ARGS, bb.capacity(), bb, null);
+        checkForError(ret, "error while asking for number of function arguments.");
+
+        numArgs = bb.getInt(0);
+
+    }
+
+//    public CLKernel putArg(Buffer value) {
+//        setArg(argIndex++, value);
+//        return this;
+//    }
+    
+    public CLKernel putArg(CLMemory<?> value) {
+        setArg(argIndex++, value);
+        return this;
+    }
+
+    public CLKernel putArg(int value) {
+        setArg(argIndex++, value);
+        return this;
+    }
+
+    public CLKernel putArg(long value) {
+        setArg(argIndex++, value);
+        return this;
+    }
+
+    public CLKernel putArg(float value) {
+        setArg(argIndex++, value);
+        return this;
+    }
+
+    public CLKernel putArg(double value) {
+        setArg(argIndex++, value);
+        return this;
+    }
+
+    public CLKernel putNullArg(int size) {
+        setNullArg(argIndex++, size);
+        return this;
+    }
+
+    public CLKernel putArgs(CLMemory<?>... values) {
+        setArgs(argIndex, values);
+        argIndex += values.length;
+        return this;
+    }
+
+    public CLKernel rewind() {
+        argIndex = 0;
+        return this;
+    }
+
+//    public CLKernel setArg(int argumentIndex, Buffer value) {
+//        setArgument(argumentIndex, CLMemory.sizeOfBufferElem(value)*value.capacity(), value);
+//        return this;
+//    }
+
+    public CLKernel setArg(int argumentIndex, CLMemory<?> value) {
+        setArgument(argumentIndex, Platform.is32Bit()?4:8, wrap(value.ID));
+        return this;
+    }
+
+    public CLKernel setArg(int argumentIndex, int value) {
+        setArgument(argumentIndex, 4, wrap(value));
+        return this;
+    }
+
+    public CLKernel setArg(int argumentIndex, long value) {
+        if(force32BitArgs) {
+            setArgument(argumentIndex, 4, wrap((int)value));
+        }else{
+            setArgument(argumentIndex, 8, wrap(value));
+        }
+        return this;
+    }
+
+    public CLKernel setArg(int argumentIndex, float value) {
+        setArgument(argumentIndex, 4, wrap(value));
+        return this;
+    }
+
+    public CLKernel setArg(int argumentIndex, double value) {
+        if(force32BitArgs) {
+            setArgument(argumentIndex, 4, wrap((float)value));
+        }else{
+            setArgument(argumentIndex, 8, wrap(value));
+        }
+        return this;
+    }
+
+    public CLKernel setNullArg(int argumentIndex, int size) {
+        setArgument(argumentIndex, size, null);
+        return this;
+    }
+
+    public CLKernel setArgs(CLMemory<?>... values) {
+        setArgs(0, values);
+        return this;
+    }
+
+    private void setArgs(int startIndex, CLMemory<?>... values) {
+        for (int i = 0; i < values.length; i++) {
+            setArg(i+startIndex, values[i]);
+        }
+    }
+
+    private void setArgument(int argumentIndex, int size, Buffer value) {
+        if(argumentIndex >= numArgs || argumentIndex < 0) {
+            throw new IndexOutOfBoundsException("kernel "+ toString() +" has "+numArgs+
+                    " arguments, can not set argument with index "+argumentIndex);
+        }
+        if(!program.isExecutable()) {
+            throw new IllegalStateException("can not set program" +
+                    " arguments for a not executable program. "+program);
+        }
+
+        int ret = cl.clSetKernelArg(ID, argumentIndex, size, value);
+        checkForError(ret, "error on clSetKernelArg");
+    }
+
+    /**
+     * Forces double and long arguments to be passed as float and int to the OpenCL kernel.
+     * This can be used in applications which want to mix kernels with different floating point precision.
+     */
+    public CLKernel setForce32BitArgs(boolean force) {
+        this.force32BitArgs = force;
+        return this;
+    }
+    
+    public CLProgram getProgram() {
+        return program;
+    }
+
+    /**
+     * @see #setForce32BitArgs(boolean) 
+     */
+    public boolean isForce32BitArgsEnabled() {
+        return force32BitArgs;
+    }
+
+    private Buffer wrap(float value) {
+        return buffer.putFloat(value).rewind();
+    }
+
+    private Buffer wrap(double value) {
+        return buffer.putDouble(value).rewind();
+    }
+
+    private Buffer wrap(int value) {
+        return buffer.putInt(value).rewind();
+    }
+
+    private Buffer wrap(long value) {
+        return buffer.putLong(value).rewind();
+    }
+
+    /**
+     * Returns the amount of local memory in bytes being used by a kernel.
+     * This includes local memory that may be needed by an implementation to execute the kernel,
+     * variables declared inside the kernel with the <code>__local</code> address qualifier and local memory
+     * to be allocated for arguments to the kernel declared as pointers with the <code>__local</code> address
+     * qualifier and whose size is specified with clSetKernelArg.
+     * If the local memory size, for any pointer argument to the kernel declared with
+     * the <code>__local</code> address qualifier, is not specified, its size is assumed to be 0.
+     */
+    public long getLocalMemorySize(CLDevice device) {
+        return getWorkGroupInfo(device, CL_KERNEL_LOCAL_MEM_SIZE);
+    }
+
+    /**
+     * Returns the work group size for this kernel on the given device.
+     * This provides a mechanism for the application to query the work-group size
+     * that can be used to execute a kernel on a specific device given by device.
+     * The OpenCL implementation uses the resource requirements of the kernel
+     * (register usage etc.) to determine what this work-group size should be. 
+     */
+    public long getWorkGroupSize(CLDevice device) {
+        return getWorkGroupInfo(device, CL_KERNEL_WORK_GROUP_SIZE);
+    }
+
+    /**
+     * Returns the work-group size specified by the <code>__attribute__((reqd_work_gr oup_size(X, Y, Z)))</code> qualifier in kernel sources.
+     * If the work-group size is not specified using the above attribute qualifier <code>new long[]{(0, 0, 0)}</code> is returned.
+     */
+    public long[] getCompileWorkGroupSize(CLDevice device) {
+        int ret = cl.clGetKernelWorkGroupInfo(ID, device.ID, CL_KERNEL_COMPILE_WORK_GROUP_SIZE, 8*3, buffer, null);
+        checkForError(ret, "error while asking for clGetKernelWorkGroupInfo");
+        return new long[] { buffer.getLong(0), buffer.getLong(1), buffer.getLong(2) };
+    }
+
+    private long getWorkGroupInfo(CLDevice device, int flag) {
+        int ret = cl.clGetKernelWorkGroupInfo(ID, device.ID, flag, 8, buffer, null);
+        checkForError(ret, "error while asking for clGetKernelWorkGroupInfo");
+        return buffer.getLong(0);
+    }
+
+    /**
+     * Releases all resources of this kernel from its context.
+     */
+    public void release() {
+        int ret = cl.clReleaseKernel(ID);
+        program.onKernelReleased(this);
+        checkForError(ret, "can not release kernel");
+    }
+
+    public void close() {
+        release();
+    }
+
+    @Override
+    public String toString() {
+        return "CLKernel [id: " + ID
+                      + " name: " + name+"]";
+    }
+
+    @Override
+    public boolean equals(Object obj) {
+        if (obj == null) {
+            return false;
+        }
+        if (getClass() != obj.getClass()) {
+            return false;
+        }
+        final CLKernel other = (CLKernel) obj;
+        if (this.ID != other.ID) {
+            return false;
+        }
+        if (!this.program.equals(other.program)) {
+            return false;
+        }
+        return true;
+    }
+
+    @Override
+    public int hashCode() {
+        int hash = 7;
+        hash = 43 * hash + (int) (this.ID ^ (this.ID >>> 32));
+        hash = 43 * hash + (this.program != null ? this.program.hashCode() : 0);
+        return hash;
+    }
+
+    /**
+     * Returns a new instance of this kernel with uninitialized arguments.
+     */
+    @Override
+    public CLKernel clone() {
+        return program.createCLKernel(name).setForce32BitArgs(force32BitArgs);
+    }
+
+}
diff --git a/src/com/jogamp/opencl/CLMemory.java b/src/com/jogamp/opencl/CLMemory.java
new file mode 100644
index 00000000..55a04af9
--- /dev/null
+++ b/src/com/jogamp/opencl/CLMemory.java
@@ -0,0 +1,392 @@
+package com.jogamp.opencl;
+
+import com.jogamp.opencl.gl.CLGLI;
+import com.jogamp.common.nio.Buffers;
+import com.jogamp.common.nio.PointerBuffer;
+import java.nio.Buffer;
+import java.nio.ByteBuffer;
+import java.nio.DoubleBuffer;
+import java.nio.FloatBuffer;
+import java.nio.IntBuffer;
+import java.nio.ShortBuffer;
+import java.util.ArrayList;
+import java.util.EnumSet;
+import java.util.List;
+
+import static com.jogamp.opencl.CLException.*;
+import static com.jogamp.opencl.gl.CLGLI.*;
+
+/**
+ * Common superclass for all OpenCL memory types.
+ * @author Michael Bien
+ */
+public abstract class CLMemory <B extends Buffer> extends CLObject implements CLResource {
+    
+    B buffer;
+    protected final int FLAGS;
+    
+    protected <Buffer> CLMemory(CLContext context, long id, int flags) {
+        super(context, id);
+        this.FLAGS = flags;
+    }
+    
+    protected CLMemory(CLContext context, B directBuffer, long id, int flags) {
+        super(context, id);
+        this.buffer = directBuffer;
+        this.FLAGS = flags;
+    }
+
+    /**
+     * Returns true if a host pointer must be specified on mem object creation.
+     */
+    protected static boolean isHostPointerFlag(int flags) {
+        return (flags & CL_MEM_COPY_HOST_PTR) != 0
+            || (flags & CL_MEM_USE_HOST_PTR)  != 0;
+    }
+
+    protected static int sizeOfBufferElem(Buffer buffer) {
+        if (buffer instanceof ByteBuffer) {
+            return Buffers.SIZEOF_BYTE;
+        } else if (buffer instanceof IntBuffer) {
+            return Buffers.SIZEOF_INT;
+        } else if (buffer instanceof ShortBuffer) {
+            return Buffers.SIZEOF_SHORT;
+        } else if (buffer instanceof FloatBuffer) {
+            return Buffers.SIZEOF_FLOAT;
+        } else if (buffer instanceof DoubleBuffer) {
+            return Buffers.SIZEOF_DOUBLE;
+        }
+        throw new RuntimeException("Unexpected buffer type " + buffer.getClass().getName());
+    }
+
+    protected static CL getCL(CLContext context) {
+        return context.cl;
+    }
+
+    /**
+     * Returns a new instance of CLMemory pointing to the same CLResource but using a different Buffer.
+     */
+    public abstract <T extends Buffer> CLMemory<T> cloneWith(T directBuffer);
+
+
+    public CLMemory<B> use(B buffer) {
+        if(this.buffer != null && buffer != null && this.buffer.getClass() != buffer.getClass()) {
+            throw new IllegalArgumentException(
+                    "expected a Buffer of class " + this.buffer.getClass()
+                    +" but got " + buffer.getClass());
+        }
+        this.buffer = buffer;
+        return this;
+    }
+
+    /**
+     * Returns the optional NIO buffer for this memory object.
+     */
+    public B getBuffer() {
+        return buffer;
+    }
+
+    /**
+     * Returns the capacity of the wrapped direct buffer or 0 if no buffer available.
+     */
+    public int getCapacity() {
+        if(buffer == null) {
+            return 0;
+        }
+        return buffer.capacity();
+    }
+
+    /**
+     * Returns the size of the wrapped direct buffer in byte or 0 if no buffer available.
+     */
+    public int getSize() {
+        if(buffer == null) {
+            return 0;
+        }
+        return sizeOfBufferElem(buffer) * buffer.capacity();
+    }
+
+    /**
+     * Returns the size of the allocated OpenCL memory.
+     */
+    public long getCLSize() {
+        PointerBuffer pb = PointerBuffer.allocateDirect(1);
+        int ret = cl.clGetMemObjectInfo(ID, CL_MEM_SIZE, PointerBuffer.elementSize(), pb.getBuffer(), null);
+        checkForError(ret, "can not obtain buffer info");
+        return pb.get();
+    }
+
+    /**
+     * Returns the configuration of this memory object.
+     */
+    public EnumSet<Mem> getConfig() {
+        return Mem.valuesOf(FLAGS);
+    }
+
+    /**
+     * Returns true if this memory object was created with the {@link Mem#READ_ONLY} flag.
+     */
+    public boolean isReadOnly() {
+        return (Mem.READ_ONLY.CONFIG & FLAGS) != 0;
+    }
+
+    /**
+     * Returns true if this memory object was created with the {@link Mem#WRITE_ONLY} flag.
+     */
+    public boolean isWriteOnly() {
+        return (Mem.WRITE_ONLY.CONFIG & FLAGS) != 0;
+    }
+
+    /**
+     * Returns true if this memory object was created with the {@link Mem#READ_WRITE} flag.
+     */
+    public boolean isReadWrite() {
+        return (Mem.READ_WRITE.CONFIG & FLAGS) != 0;
+    }
+
+    public void release() {
+        int ret = cl.clReleaseMemObject(ID);
+        context.onMemoryReleased(this);
+        checkForError(ret, "can not release mem object");
+    }
+
+    public void close() {
+        release();
+    }
+
+    // kept only for debugging purposes
+    /**
+     * Returns the OpenGL buffer type of this shared buffer.
+     */
+    /*public*/ final GLObjectType _getGLObjectType() {
+        int[] array = new int[1];
+        int ret = ((CLGLI)cl).clGetGLObjectInfo(ID, array, 0, null, 0);
+        CLException.checkForError(ret, "error while asking for gl object info");
+        return GLObjectType.valueOf(array[0]);
+    }
+
+    /**
+     * Returns the OpenGL object id of this shared buffer.
+     */
+    /*public*/ final int _getGLObjectID() {
+        int[] array = new int[1];
+        int ret = ((CLGLI)cl).clGetGLObjectInfo(ID, null, 0, array, 0);
+        CLException.checkForError(ret, "error while asking for gl object info");
+        return array[0];
+    }
+
+    @Override
+    public boolean equals(Object obj) {
+        if (obj == null) {
+            return false;
+        }
+        if (getClass() != obj.getClass()) {
+            return false;
+        }
+        final CLMemory<?> other = (CLMemory<?>) obj;
+        if (this.ID != other.ID) {
+            return false;
+        }
+        if (this.context != other.context && (this.context == null || !this.context.equals(other.context))) {
+            return false;
+        }
+        return true;
+    }
+
+    @Override
+    public int hashCode() {
+        int hash = 7;
+        hash = 83 * hash + (int) (this.ID ^ (this.ID >>> 32));
+        hash = 83 * hash + (this.context != null ? this.context.hashCode() : 0);
+        return hash;
+    }
+
+    @Override
+    public String toString() {
+        return "CLMemory [id: " + ID+"]";
+    }
+
+    /**
+     * Memory settings for configuring CLMemory.
+     */
+    public enum Mem {
+
+        /**
+         * Enum representing CL_MEM_READ_WRITE.
+         * This flag specifies that the memory object will be read and
+         * written by a kernel.
+         */
+        READ_WRITE(CL_MEM_READ_WRITE),
+
+        /**
+         * Enum representing CL_MEM_WRITE_ONLY.
+         * This flags specifies that the memory object will be written
+         * but not read by a kernel.
+         * Reading from a buffer or image object created with WRITE_ONLY
+         * inside a kernel is undefined.
+         */
+        WRITE_ONLY(CL_MEM_WRITE_ONLY),
+
+        /**
+         * Enum representing CL_MEM_READ_ONLY.
+         * This flag specifies that the memory object is a read-only memory
+         * object when used inside a kernel. Writing to a buffer or image object
+         * created withREAD_ONLY inside a kernel is undefined.
+         */
+        READ_ONLY(CL_MEM_READ_ONLY),
+
+        /**
+         * Enum representing CL_MEM_USE_HOST_PTR.
+         * If specified, it indicates that the application wants the OpenCL
+         * implementation to use memory referenced by host_ptr as the storage
+         * bits for the memory object. OpenCL implementations are allowed
+         * to cache the buffer contents pointed to by host_ptr in device memory.
+         * This cached copy can be used when kernels are executed on a device.
+         */
+        USE_BUFFER(CL_MEM_USE_HOST_PTR),
+
+        /**
+         * Enum representing CL_MEM_ALLOC_HOST_PTR.
+         * This flag specifies that the application wants the OpenCL implementation
+         * to allocate memory from host accessible memory.
+         * {@link #ALLOCATE_BUFFER} and {@link #USE_BUFFER} are mutually exclusive.
+         */
+        ALLOCATE_BUFFER(CL_MEM_ALLOC_HOST_PTR),
+
+        /**
+         * Enum representing CL_MEM_COPY_HOST_PTR.
+         * If {@link #COPY_BUFFER} specified, it indicates that the application
+         * wants the OpenCL implementation to allocate memory for the memory object
+         * and copy the data from memory referenced by host_ptr.<br/>
+         * {@link #COPY_BUFFER} and {@link #USE_BUFFER} are mutually exclusive.
+         */
+        COPY_BUFFER(CL_MEM_COPY_HOST_PTR);
+
+        /**
+         * Value of wrapped OpenCL flag.
+         */
+        public final int CONFIG;
+
+        private Mem(int config) {
+            this.CONFIG = config;
+        }
+
+        public static Mem valueOf(int bufferFlag) {
+            switch (bufferFlag) {
+                case CL_MEM_READ_WRITE:
+                    return Mem.READ_WRITE;
+                case CL_MEM_READ_ONLY:
+                    return Mem.READ_ONLY;
+                case CL_MEM_WRITE_ONLY:
+                    return Mem.WRITE_ONLY;
+                case CL_MEM_USE_HOST_PTR:
+                    return Mem.USE_BUFFER;
+                case(CL_MEM_ALLOC_HOST_PTR):
+                    return ALLOCATE_BUFFER;
+                case CL_MEM_COPY_HOST_PTR:
+                    return Mem.COPY_BUFFER;
+            }
+            return null;
+        }
+
+        public static EnumSet<Mem> valuesOf(int bitfield) {
+            List<Mem> matching = new ArrayList<Mem>();
+            Mem[] values = Mem.values();
+            for (Mem value : values) {
+                if((value.CONFIG & bitfield) != 0)
+                    matching.add(value);
+            }
+            if(matching.isEmpty())
+                return EnumSet.noneOf(Mem.class);
+            else
+                return EnumSet.copyOf(matching);
+        }
+
+        public static int flagsToInt(Mem[] flags) {
+            int clFlags = 0;
+            if (flags != null) {
+                for (int i = 0; i < flags.length; i++) {
+                    clFlags |= flags[i].CONFIG;
+                }
+            }
+            if (clFlags == 0) {
+                clFlags = CL_MEM_READ_WRITE;
+            }
+            return clFlags;
+        }
+    }
+
+    /**
+     * Configures the mapping process.
+     * @see com.jogamp.opencl.CLCommandQueue#putMapBuffer(CLBuffer, com.jogamp.opencl.CLMemory.Map, boolean).
+     * @see com.jogamp.opencl.CLCommandQueue#putMapImage(CLImage2d, com.jogamp.opencl.CLMemory.Map, boolean)
+     * @see com.jogamp.opencl.CLCommandQueue#putMapImage(CLImage3d, com.jogamp.opencl.CLMemory.Map, boolean)
+     */
+    public enum Map {
+
+        /**
+         * Enum representing CL_MAP_READ | CL_MAP_WRITE.
+         * This flag specifies that the memory object will be mapped for read and write operation.
+         */
+        READ_WRITE(CL_MAP_READ | CL_MAP_WRITE),
+
+        /**
+         * Enum representing CL_MAP_WRITE.
+         * This flag specifies that the memory object will be mapped for write operation.
+         */
+        WRITE(CL_MAP_WRITE),
+
+        /**
+         * Enum representing CL_MAP_READ.
+         * This flag specifies that the memory object will be mapped for read operation.
+         */
+        READ(CL_MAP_READ);
+
+        /**
+         * Value of wrapped OpenCL flag.
+         */
+        public final int FLAGS;
+
+        private Map(int flags) {
+            this.FLAGS = flags;
+        }
+
+        public Map valueOf(int flag) {
+            if(flag == WRITE.FLAGS)
+                return WRITE;
+            else if(flag == READ.FLAGS)
+                return READ;
+            else if(flag == READ_WRITE.FLAGS)
+                return READ_WRITE;
+            return null;
+        }
+
+    }
+
+    public enum GLObjectType {
+
+        GL_OBJECT_BUFFER(CL_GL_OBJECT_BUFFER),
+        GL_OBJECT_TEXTURE2D(CL_GL_OBJECT_TEXTURE2D),
+        GL_OBJECT_TEXTURE3D(CL_GL_OBJECT_TEXTURE3D),
+        GL_OBJECT_RENDERBUFFER(CL_GL_OBJECT_RENDERBUFFER);
+
+        public final int TYPE;
+
+        private GLObjectType(int type) {
+            this.TYPE = type;
+        }
+
+        public static GLObjectType valueOf(int type) {
+            if(type == CL_GL_OBJECT_BUFFER)
+                return GL_OBJECT_BUFFER;
+            else if(type == CL_GL_OBJECT_TEXTURE2D)
+                return GL_OBJECT_TEXTURE2D;
+            else if(type == CL_GL_OBJECT_TEXTURE3D)
+                return GL_OBJECT_TEXTURE3D;
+            else if(type == CL_GL_OBJECT_RENDERBUFFER)
+                return GL_OBJECT_RENDERBUFFER;
+            return null;
+        }
+    }
+
+}
diff --git a/src/com/jogamp/opencl/CLObject.java b/src/com/jogamp/opencl/CLObject.java
new file mode 100644
index 00000000..58c8485c
--- /dev/null
+++ b/src/com/jogamp/opencl/CLObject.java
@@ -0,0 +1,57 @@
+package com.jogamp.opencl;
+
+/**
+ * Common superclass for all OpenCL objects.
+ * @author Michael Bien
+ */
+abstract class CLObject {
+
+    /**
+     * The OpenCL object handle.
+     */
+    public final long ID;
+    
+    protected CLContext context;
+
+    protected final CL cl;
+
+    CLObject(CL cl, long ID) {
+        this.cl = cl;
+        this.context = null;
+        this.ID = ID;
+    }
+
+    CLObject(CLContext context, long ID) {
+        this.cl = context.cl;
+        this.context = context;
+        this.ID = ID;
+    }
+
+    /**
+     * Returns the context for this OpenCL object.
+     */
+    public CLContext getContext() {
+        return context;
+    }
+
+    /**
+     * Returns the platform for this OpenCL object.
+     */
+    public CLPlatform getPlatform() {
+        return context.getPlatform();
+    }
+
+    /**
+     * Returns the OpenCL object handle
+     */
+    public long getID() {
+        return ID;
+    }
+
+    @Override
+    public String toString() {
+        return "CLObject [id: " + ID
+                      + " context: " + context+"]";
+    }
+
+}
diff --git a/src/com/jogamp/opencl/CLPlatform.java b/src/com/jogamp/opencl/CLPlatform.java
new file mode 100644
index 00000000..f3c00926
--- /dev/null
+++ b/src/com/jogamp/opencl/CLPlatform.java
@@ -0,0 +1,302 @@
+package com.jogamp.opencl;
+
+import com.jogamp.common.nio.Int64Buffer;
+import com.jogamp.opencl.util.CLUtil;
+import com.jogamp.opencl.impl.CLImpl;
+import com.jogamp.common.nio.PointerBuffer;
+import java.nio.ByteBuffer;
+import java.nio.ByteOrder;
+import java.nio.IntBuffer;
+import java.util.ArrayList;
+import java.util.Collections;
+import java.util.HashSet;
+import java.util.List;
+import java.util.Map;
+import java.util.Scanner;
+import java.util.Set;
+
+import static com.jogamp.opencl.CLException.*;
+import static com.jogamp.opencl.CL.*;
+
+/**
+ *
+ * @author Michael Bien
+ */
+public final class CLPlatform {
+
+    /**
+     * OpenCL platform id for this platform.
+     */
+    public final long ID;
+
+    private static final CL cl;
+
+    private Set<String> extensions;
+
+    static{
+        NativeLibLoader.loadJOCL();
+//        System.loadLibrary("gluegen-rt");
+//        ProcAddressHelper.resetProcAddressTable(table, null);
+        cl = new CLImpl();
+    }
+
+    private CLPlatform(long id) {
+        this.ID = id;
+    }
+
+    /**
+     * Returns the default OpenCL platform or null when no platform found.
+     */
+    public static CLPlatform getDefault() {
+        CLPlatform[] platforms = listCLPlatforms();
+        if(platforms.length > 0)
+            return platforms[0];
+        return null;
+    }
+
+    /**
+     * Lists all available OpenCL implementations.
+     * @throws CLException if something went wrong initializing OpenCL
+     */
+    public static CLPlatform[] listCLPlatforms() {
+
+        IntBuffer ib = ByteBuffer.allocateDirect(4).order(ByteOrder.nativeOrder()).asIntBuffer();
+        // find all available OpenCL platforms
+        int ret = cl.clGetPlatformIDs(0, null, ib);
+        checkForError(ret, "can not enumerate platforms");
+
+        // receive platform ids
+        PointerBuffer platformId = PointerBuffer.allocateDirect(ib.get(0));
+        ret = cl.clGetPlatformIDs(platformId.capacity(), platformId, null);
+        checkForError(ret, "can not enumerate platforms");
+
+        CLPlatform[] platforms = new CLPlatform[platformId.capacity()];
+
+        for (int i = 0; i < platformId.capacity(); i++)
+            platforms[i] = new CLPlatform(platformId.get(i));
+
+        return platforms;
+    }
+
+    /**
+     * Returns the low level binding interface to the OpenCL APIs.
+     */
+    public static CL getLowLevelCLInterface() {
+        return cl;
+    }
+
+    /**
+     * Hint to allow the implementation to release the resources allocated by the OpenCL compiler.
+     * Calls to {@link CLProgram#build()} after unloadCompiler will reload the compiler if necessary.
+     */
+    public static void unloadCompiler() {
+        int ret = cl.clUnloadCompiler();
+        checkForError(ret, "error while sending unload compiler hint");
+    }
+
+    /**
+     * Lists all physical devices available on this platform.
+     * @see #listCLDevices(com.jogamp.opencl.CLDevice.Type...)
+     */
+    public CLDevice[] listCLDevices() {
+        return this.listCLDevices(CLDevice.Type.ALL);
+    }
+
+    /**
+     * Lists all physical devices available on this platform matching the given {@link CLDevice.Type}.
+     */
+    public CLDevice[] listCLDevices(CLDevice.Type... types) {
+
+        IntBuffer ib = ByteBuffer.allocateDirect(4).order(ByteOrder.nativeOrder()).asIntBuffer();
+
+        List<CLDevice> list = new ArrayList<CLDevice>();
+        for(int t = 0; t < types.length; t++) {
+            CLDevice.Type type = types[t];
+
+            //find all devices
+            int ret = cl.clGetDeviceIDs(ID, type.TYPE, 0, null, ib);
+
+            // return an empty array rather than throwing an exception
+            if(ret == CL.CL_DEVICE_NOT_FOUND) {
+                continue;
+            }
+
+            checkForError(ret, "error while enumerating devices");
+
+            PointerBuffer deviceIDs = PointerBuffer.allocateDirect(ib.get(0));
+            ret = cl.clGetDeviceIDs(ID, type.TYPE, deviceIDs.capacity(), deviceIDs, null);
+            checkForError(ret, "error while enumerating devices");
+
+            //add device to list
+            for (int n = 0; n < deviceIDs.capacity(); n++)
+                list.add(new CLDevice(cl, deviceIDs.get(n)));
+        }
+
+        CLDevice[] devices = new CLDevice[list.size()];
+        for (int i = 0; i < list.size(); i++) {
+            devices[i] = list.get(i);
+        }
+
+        return devices;
+
+    }
+
+    static CLDevice findMaxFlopsDevice(CLDevice[] devices) {
+        return findMaxFlopsDevice(devices, null);
+    }
+    
+    static CLDevice findMaxFlopsDevice(CLDevice[] devices, CLDevice.Type type) {
+
+        CLDevice maxFLOPSDevice = null;
+
+        int maxflops = -1;
+
+        for (int i = 0; i < devices.length; i++) {
+
+            CLDevice device = devices[i];
+
+            if(type == null || type.equals(device.getType())) {
+
+                int maxComputeUnits     = device.getMaxComputeUnits();
+                int maxClockFrequency   = device.getMaxClockFrequency();
+                int flops = maxComputeUnits*maxClockFrequency;
+
+                if(flops > maxflops) {
+                    maxflops = flops;
+                    maxFLOPSDevice = device;
+                }
+            }
+
+        }
+
+        return maxFLOPSDevice;
+    }
+
+
+    /**
+     * Returns the device with maximal FLOPS from this platform.
+     * The device speed is estimated by calculating the product of
+     * MAX_COMPUTE_UNITS and MAX_CLOCK_FREQUENCY.
+     * @see #getMaxFlopsDevice(com.jogamp.opencl.CLDevice.Type...)
+     */
+    public CLDevice getMaxFlopsDevice() {
+        return findMaxFlopsDevice(listCLDevices());
+    }
+
+    /**
+     * Returns the device with maximal FLOPS and the specified type from this platform.
+     * The device speed is estimated by calculating the product of
+     * MAX_COMPUTE_UNITS and MAX_CLOCK_FREQUENCY.
+     */
+    public CLDevice getMaxFlopsDevice(CLDevice.Type... types) {
+        return findMaxFlopsDevice(listCLDevices(types));
+    }
+
+    /**
+     * Returns the platform name.
+     */
+    public String getName() {
+        return getInfoString(CL_PLATFORM_NAME);
+    }
+
+    /**
+     * Returns the platform version.
+     */
+    public String getVersion() {
+        return getInfoString(CL_PLATFORM_VERSION);
+    }
+
+    /**
+     * Returns the platform profile.
+     */
+    public String getProfile() {
+        return getInfoString(CL_PLATFORM_PROFILE);
+    }
+
+    /**
+     * Returns the platform vendor.
+     */
+    public String getVendor() {
+        return getInfoString(CL_PLATFORM_VENDOR);
+    }
+
+    /**
+     * Returns true if the extension is supported on this platform.
+     */
+    public boolean isExtensionAvailable(String extension) {
+        return getExtensions().contains(extension);
+    }
+
+    /**
+     * Returns all platform extension names as unmodifiable Set.
+     */
+    public Set<String> getExtensions() {
+
+        if(extensions == null) {
+            extensions = new HashSet<String>();
+            String ext = getInfoString(CL_PLATFORM_EXTENSIONS);
+            Scanner scanner = new Scanner(ext);
+
+            while(scanner.hasNext())
+                extensions.add(scanner.next());
+
+            extensions = Collections.unmodifiableSet(extensions);
+        }
+
+        return extensions;
+    }
+
+    /**
+     * Returns a Map of platform properties with the enum names as keys.
+     * @see CLUtil#obtainPlatformProperties(com.jogamp.opencl.CLPlatform)
+     */
+    public Map<String, String> getProperties() {
+        return CLUtil.obtainPlatformProperties(this);
+    }
+
+    /**
+     * Returns a info string in exchange for a key (CL_PLATFORM_*).
+     */
+    public String getInfoString(int key) {
+        Int64Buffer size = Int64Buffer.allocateDirect(1);
+        // TODO use cache/query size
+        ByteBuffer bb = ByteBuffer.allocateDirect(512);
+
+        int ret = cl.clGetPlatformInfo(ID, key, bb.capacity(), bb, size);
+        checkForError(ret, "can not receive info string");
+
+        return CLUtil.clString2JavaString(bb, (int)size.get(0));
+    }
+
+    @Override
+    public String toString() {
+        return "CLPlatform [name:" + getName()
+                         +" vendor:"+getVendor()
+                         +" profile:"+getProfile()
+                         +" version:"+getVersion()+"]";
+    }
+
+    @Override
+    public boolean equals(Object obj) {
+        if (obj == null) {
+            return false;
+        }
+        if (getClass() != obj.getClass()) {
+            return false;
+        }
+        final CLPlatform other = (CLPlatform) obj;
+        if (this.ID != other.ID) {
+            return false;
+        }
+        return true;
+    }
+
+    @Override
+    public int hashCode() {
+        int hash = 7;
+        hash = 71 * hash + (int) (this.ID ^ (this.ID >>> 32));
+        return hash;
+    }
+
+
+}
diff --git a/src/com/jogamp/opencl/CLProgram.java b/src/com/jogamp/opencl/CLProgram.java
new file mode 100644
index 00000000..bb5abccf
--- /dev/null
+++ b/src/com/jogamp/opencl/CLProgram.java
@@ -0,0 +1,682 @@
+package com.jogamp.opencl;
+
+import com.jogamp.opencl.util.CLProgramConfiguration;
+import com.jogamp.opencl.util.CLUtil;
+import com.jogamp.common.nio.Int64Buffer;
+import com.jogamp.common.os.Platform;
+import com.jogamp.common.nio.PointerBuffer;
+import java.nio.ByteBuffer;
+import java.nio.IntBuffer;
+import java.util.Collections;
+import java.util.HashMap;
+import java.util.HashSet;
+import java.util.LinkedHashMap;
+import java.util.Set;
+import java.util.Map;
+
+import static com.jogamp.opencl.CLException.*;
+import static com.jogamp.opencl.CL.*;
+import static com.jogamp.common.nio.Buffers.*;
+
+/**
+ * Represents a OpenCL program executed on one or more {@link CLDevice}s.
+ * A CLProgram must be build using one of the build methods before creating {@link CLKernel}s.
+ * @see CLContext#createProgram(java.io.InputStream)
+ * @see CLContext#createProgram(java.lang.String)
+ * @see CLContext#createProgram(java.util.Map)
+ * @author Michael Bien
+ */
+public class CLProgram extends CLObject implements CLResource {
+
+//    private final static Object buildLock = new Object();
+    
+    private final Set<CLKernel> kernels;
+    private Map<CLDevice, Status> buildStatusMap;
+
+    private boolean executable;
+    private boolean released;
+
+    private CLProgram(CLContext context, long id) {
+        super(context, id);
+        this.kernels = new HashSet<CLKernel>();
+    }
+    
+    static CLProgram create(CLContext context, String src) {
+
+        IntBuffer status = newDirectByteBuffer(4).asIntBuffer();
+        // Create the program
+        long id = context.cl.clCreateProgramWithSource(context.ID, 1, new String[] {src},
+                               (Int64Buffer)Int64Buffer.allocateDirect(1).put(src.length()), status);
+
+        checkForError(status.get(), "can not create program with source");
+        
+        return new CLProgram(context, id);
+    }
+
+    static CLProgram create(CLContext context, Map<CLDevice, byte[]> binaries) {
+
+        PointerBuffer devices = PointerBuffer.allocateDirect(binaries.size());
+        PointerBuffer codeBuffers = PointerBuffer.allocateDirect(binaries.size());
+        Int64Buffer lengths = Int64Buffer.allocateDirect(binaries.size());
+
+        int i = 0;
+        Set<CLDevice> keys = binaries.keySet();
+        for (CLDevice device : keys) {
+
+            byte[] bytes = binaries.get(device);
+
+            devices.put(device.ID);
+            lengths.put(bytes.length);
+
+            codeBuffers.referenceBuffer(i, newDirectByteBuffer(bytes));
+            i++;
+        }
+        devices.rewind();
+        lengths.rewind();
+
+        IntBuffer err = newDirectIntBuffer(1);
+//        IntBuffer status = newDirectByteBuffer(binaries.size()*4).asIntBuffer();
+        long id = context.cl.clCreateProgramWithBinary(context.ID, devices.capacity(), devices, lengths, codeBuffers, /*status*/null, err);
+
+//        while(status.remaining() != 0) {
+//            checkForError(status.get(), "unable to load binaries on all devices");
+//        }
+
+        checkForError(err.get(), "can not create program with binary");
+
+        return new CLProgram(context, id);
+    }
+
+    private void initBuildStatus() {
+
+        if(buildStatusMap == null) {
+//            synchronized(buildLock) {
+                Map<CLDevice, Status> map = new HashMap<CLDevice, Status>();
+                CLDevice[] devices = getCLDevices();
+                for (CLDevice device : devices) {
+                    Status status = getBuildStatus(device);
+                    if(status == Status.BUILD_SUCCESS) {
+                        executable = true;
+                    }
+                    map.put(device, status);
+                }
+                this.buildStatusMap = Collections.unmodifiableMap(map);
+//            }
+        }
+    }
+
+    private String getBuildInfoString(long device, int flag) {
+
+        if(released) {
+            return "";
+        }
+
+        Int64Buffer size = Int64Buffer.allocateDirect(1);
+
+        int ret = cl.clGetProgramBuildInfo(ID, device, flag, 0, null, size);
+        checkForError(ret, "on clGetProgramBuildInfo");
+
+        ByteBuffer buffer = newDirectByteBuffer((int)size.get(0));
+
+        ret = cl.clGetProgramBuildInfo(ID, device, flag, buffer.capacity(), buffer, null);
+        checkForError(ret, "on clGetProgramBuildInfo");
+
+        return CLUtil.clString2JavaString(buffer, (int)size.get(0));
+    }
+
+    private String getProgramInfoString(int flag) {
+
+        if(released) {
+            return "";
+        }
+
+        Int64Buffer size = Int64Buffer.allocateDirect(1);
+
+        int ret = cl.clGetProgramInfo(ID, flag, 0, null, size);
+        checkForError(ret, "on clGetProgramInfo");
+
+        ByteBuffer buffer = newDirectByteBuffer((int)size.get(0));
+
+        ret = cl.clGetProgramInfo(ID, flag, buffer.capacity(), buffer, null);
+        checkForError(ret, "on clGetProgramInfo");
+
+        return CLUtil.clString2JavaString(buffer, (int)size.get(0));
+    }
+
+//    private int getProgramInfoInt(int flag) {
+//
+//        ByteBuffer bb = ByteBuffer.allocateDirect(4).order(ByteOrder.nativeOrder());
+//
+//        int ret = cl.clGetProgramInfo(programID, flag, bb.capacity(), bb, null, 0);
+//        checkForError(ret, "");
+//
+//        return bb.getInt();
+//    }
+
+    private int getBuildInfoInt(long device, int flag) {
+
+        ByteBuffer buffer = newDirectByteBuffer(4);
+
+        int ret = cl.clGetProgramBuildInfo(ID, device, flag, buffer.capacity(), buffer, null);
+        checkForError(ret, "error on clGetProgramBuildInfo");
+
+        return buffer.getInt();
+    }
+
+
+    /**
+     * Builds this program for all devices associated with the context.
+     * @return this
+     */
+    public CLProgram build() {
+        build(null, (CLDevice[])null);
+        return this;
+    }
+    
+    /**
+     * Builds this program for the given devices.
+     * @return this
+     * @param devices A list of devices this program should be build on or null for all devices of its context.
+     */
+    public CLProgram build(CLDevice... devices) {
+        build(null, devices);
+        return this;
+    }
+
+    /**
+     * Builds this program for all devices associated with the context using the specified build options.
+     * @see CompilerOptions
+     * @return this
+     */
+    public CLProgram build(String options) {
+        build(options, (CLDevice[])null);
+        return this;
+    }
+
+    /**
+     * Builds this program for all devices associated with the context using the specified build options.
+     * @see CompilerOptions
+     * @return this
+     */
+    public CLProgram build(String... options) {
+        build(optionsOf(options), (CLDevice[])null);
+        return this;
+    }
+
+    /**
+     * Builds this program for the given devices and with the specified build options. In case this program was
+     * already built and there are kernels associated with this program they will be released first before rebuild.
+     * @see CompilerOptions
+     * @return this
+     * @param devices A list of devices this program should be build on or null for all devices of its context.
+     */
+    public CLProgram build(String options, CLDevice... devices) {
+
+        if(released) {
+            throw new CLException("can not build a released program");
+        }
+
+        if(!kernels.isEmpty()) {
+            //No changes to the program executable are allowed while there are
+            //kernel objects associated with a program object.
+            releaseKernels();
+        }
+
+        PointerBuffer deviceIDs = null;
+        int count = 0;
+        if(devices != null && devices.length != 0) {
+            deviceIDs = PointerBuffer.allocateDirect(devices.length);
+            for (int i = 0; i < devices.length; i++) {
+                deviceIDs.put(i, devices[i].ID);
+            }
+            deviceIDs.rewind();
+            count = devices.length;
+        }
+
+        // nvidia driver doesn't like empty strings
+        if(options != null && options.trim().isEmpty()) {
+            options = null;
+        }
+
+        // invalidate build status
+        buildStatusMap = null;
+        executable = false;
+
+        // Build the program
+        int ret = 0;
+        // building programs is not threadsafe
+//        synchronized(buildLock) {
+            ret = cl.clBuildProgram(ID, count, deviceIDs, options, null, null);
+//        }
+
+        if(ret != CL_SUCCESS) {
+            throw newException(ret, "\n"+getBuildLog());
+        }
+
+        return this;
+    }
+
+    /**
+     * Prepares the build for this program by returning a new {@link CLProgramConfiguration}.
+     */
+    public CLProgramConfiguration prepare() {
+        return CLProgramBuilder.createConfiguration(this);
+    }
+
+    /**
+     * Creates a kernel with the specified kernel name.
+     */
+    public CLKernel createCLKernel(String kernelName) {
+
+        if(released) {
+            return null;
+        }
+
+        int[] err = new int[1];
+        long id = cl.clCreateKernel(ID, kernelName, err, 0);
+        checkForError(err[0], "unable to create Kernel with name: "+kernelName);
+
+        CLKernel kernel = new CLKernel(this, id);
+        kernels.add(kernel);
+        return kernel;
+    }
+
+    /**
+     * Creates all kernels of this program and stores them a Map with the kernel name as key.
+     */
+    public Map<String, CLKernel> createCLKernels() {
+
+        if(released) {
+            return Collections.emptyMap();
+        }
+
+        HashMap<String, CLKernel> newKernels = new HashMap<String, CLKernel>();
+
+        IntBuffer numKernels = newDirectByteBuffer(4).asIntBuffer();
+        int ret = cl.clCreateKernelsInProgram(ID, 0, null, numKernels);
+        checkForError(ret, "can not create kernels for program");
+
+        if(numKernels.get(0) > 0) {
+
+            PointerBuffer kernelIDs = PointerBuffer.allocateDirect(numKernels.get(0));
+            ret = cl.clCreateKernelsInProgram(ID, kernelIDs.capacity(), kernelIDs, null);
+            checkForError(ret, "can not create kernels for program");
+
+            for (int i = 0; i < kernelIDs.capacity(); i++) {
+                CLKernel kernel = new CLKernel(this, kernelIDs.get(i));
+                kernels.add(kernel);
+                newKernels.put(kernel.name, kernel);
+            }
+        }else{
+            initBuildStatus();
+            if(!isExecutable()) {
+                // It is illegal to create kernels from a not executable program.
+                // For consistency between AMD and NVIDIA drivers throw an exception at this point.
+                throw newException(CL_INVALID_PROGRAM_EXECUTABLE,
+                        "can not initialize kernels, program is not executable. status: "+buildStatusMap);
+            }
+        }
+
+        return newKernels;
+    }
+
+    void onKernelReleased(CLKernel kernel) {
+        this.kernels.remove(kernel);
+    }
+
+    /**
+     * Releases this program with its kernels.
+     */
+    public void release() {
+
+        releaseKernels();
+
+        executable = false;
+        released = true;
+        buildStatusMap = null;
+
+        int ret = cl.clReleaseProgram(ID);
+        context.onProgramReleased(this);
+        checkForError(ret, "can not release program");
+    }
+
+    public void close() {
+        release();
+    }
+
+    private void releaseKernels() {
+        if(!kernels.isEmpty()) {
+            // copy to array to prevent concurrent modification exception
+            CLKernel[] array = kernels.toArray(new CLKernel[kernels.size()]);
+            for (CLKernel kernel : array) {
+                kernel.release();
+            }
+        }
+    }
+
+    /**
+     * Returns all devices associated with this program.
+     */
+    public CLDevice[] getCLDevices() {
+        if(released) {
+            return new CLDevice[0];
+        }
+        Int64Buffer size = Int64Buffer.allocateDirect(1);
+        int ret = cl.clGetProgramInfo(ID, CL_PROGRAM_DEVICES, 0, null, size);
+        checkForError(ret, "on clGetProgramInfo");
+
+        ByteBuffer bb = newDirectByteBuffer((int) size.get(0));
+        ret = cl.clGetProgramInfo(ID, CL_PROGRAM_DEVICES, bb.capacity(), bb, null);
+        checkForError(ret, "on clGetProgramInfo");
+
+        int count = bb.capacity() / (Platform.is32Bit()?4:8);
+        CLDevice[] devices = new CLDevice[count];
+        for (int i = 0; i < count; i++) {
+            devices[i] = context.getDevice(Platform.is32Bit()?bb.getInt():bb.getLong());
+        }
+
+        return devices;
+
+    }
+
+    /**
+     * Returns the build log of this program on all devices. The contents of the log are
+     * implementation dependent.
+     */
+    public String getBuildLog() {
+        if(released) {
+            return "";
+        }
+        StringBuilder sb = new StringBuilder();
+        CLDevice[] devices = getCLDevices();
+        for (int i = 0; i < devices.length; i++) {
+            CLDevice device = devices[i];
+            sb.append(device).append(" build log:\n");
+            String log = getBuildLog(device).trim();
+            sb.append(log.isEmpty()?"    <empty>":log);
+            if(i != devices.length-1)
+                sb.append("\n");
+        }
+        return sb.toString();
+    }
+
+    /**
+     * Returns the build status enum of this program for each device as Map.
+     */
+    public Map<CLDevice,Status> getBuildStatus() {
+        if(released) {
+            return Collections.emptyMap();
+        }
+        initBuildStatus();
+        return buildStatusMap;
+    }
+
+    /**
+     * Returns true if the build status 'BUILD_SUCCESS' for at least one device
+     * of this program exists.
+     */
+    public boolean isExecutable() {
+        if(released) {
+            return false;
+        }
+        initBuildStatus();
+        return executable;
+    }
+
+    /**
+     * Returns the build log for this program on the specified device. The contents
+     * of the log are implementation dependent log can be an empty String.
+     */
+    public String getBuildLog(CLDevice device) {
+        return getBuildInfoString(device.ID, CL_PROGRAM_BUILD_LOG);
+    }
+
+    /**
+     * Returns the build status enum for this program on the specified device.
+     */
+    public Status getBuildStatus(CLDevice device) {
+        if(released) {
+            return Status.BUILD_NONE;
+        }
+        int clStatus = getBuildInfoInt(device.ID, CL_PROGRAM_BUILD_STATUS);
+        return Status.valueOf(clStatus);
+    }
+
+    /**
+     * Returns the source code of this program. Note: sources are not cached,
+     * each call of this method calls into Open
+     */
+    public String getSource() {
+        return getProgramInfoString(CL_PROGRAM_SOURCE);
+    }
+
+    /**
+     * Returns the binaries for this program in an ordered Map containing the device as key
+     * and the program binaries as value.
+     */
+    public Map<CLDevice, byte[]> getBinaries() {
+
+        if(!isExecutable()) {
+            return Collections.emptyMap();
+        }
+        
+        CLDevice[] devices = getCLDevices();
+
+        ByteBuffer sizes = newDirectByteBuffer(8*devices.length);
+        int ret = cl.clGetProgramInfo(ID, CL_PROGRAM_BINARY_SIZES, sizes.capacity(), sizes, null);
+        checkForError(ret, "on clGetProgramInfo");
+
+        int binariesSize = 0;
+        while(sizes.remaining() != 0) {
+            int size = (int) sizes.getLong();
+            binariesSize += size;
+        }
+        ByteBuffer binaries = newDirectByteBuffer(binariesSize);
+
+        
+        long address = InternalBufferUtil.getDirectBufferAddress(binaries);
+        PointerBuffer addresses = PointerBuffer.allocateDirect(sizes.capacity());
+        sizes.rewind();
+        while(sizes.remaining() != 0) {
+            addresses.put(address);
+            address += sizes.getLong();
+        }
+        
+        ret = cl.clGetProgramInfo(ID, CL_PROGRAM_BINARIES, addresses.capacity(), addresses.getBuffer(), null);
+        checkForError(ret, "on clGetProgramInfo");
+
+        Map<CLDevice, byte[]> map = new LinkedHashMap<CLDevice, byte[]>();
+        sizes.rewind();
+        for (int i = 0; i < devices.length; i++) {
+            byte[] bytes = new byte[(int)sizes.getLong()];
+            binaries.get(bytes);
+            map.put(devices[i], bytes);
+        }
+
+        return map;
+    }
+
+    /**
+     * Utility method which builds a properly seperated option string.
+     */
+    public static String optionsOf(String... options) {
+        StringBuilder sb = new StringBuilder(options.length * 24);
+        for (int i = 0; i < options.length; i++) {
+            sb.append(options[i]);
+            if(i!= options.length-1)
+                sb.append(" ");
+        }
+        return sb.toString();
+    }
+
+    /**
+     * Utility method for defining macros as build options (Returns "-D name").
+     */
+    public static String define(String name) {
+        return "-D "+name;
+    }
+
+    /**
+     * Utility method for defining macros as build options (Returns "-D name=value").
+     */
+    public static String define(String name, Object value) {
+        return "-D "+name+"="+value;
+    }
+
+    @Override
+    public String toString() {
+        return "CLProgram [id: " + ID
+                       + " status: "+getBuildStatus()+"]";
+    }
+
+    @Override
+    public boolean equals(Object obj) {
+        if (obj == null) {
+            return false;
+        }
+        if (getClass() != obj.getClass()) {
+            return false;
+        }
+        final CLProgram other = (CLProgram) obj;
+        if (this.ID != other.ID) {
+            return false;
+        }
+        if (!this.context.equals(other.context)) {
+            return false;
+        }
+        return true;
+    }
+
+    @Override
+    public int hashCode() {
+        int hash = 7;
+        hash = 37 * hash + (this.context != null ? this.context.hashCode() : 0);
+        hash = 37 * hash + (int) (this.ID ^ (this.ID >>> 32));
+        return hash;
+    }
+    
+    public enum Status {
+
+        BUILD_SUCCESS(CL_BUILD_SUCCESS),
+        BUILD_NONE(CL_BUILD_NONE),
+        BUILD_IN_PROGRESS(CL_BUILD_IN_PROGRESS),
+        BUILD_ERROR(CL_BUILD_ERROR);
+
+        /**
+         * Value of wrapped OpenCL device type.
+         */
+        public final int STATUS;
+
+        private Status(int status) {
+            this.STATUS = status;
+        }
+
+        public static Status valueOf(int clBuildStatus) {
+            switch(clBuildStatus) {
+                case(CL_BUILD_SUCCESS):
+                    return BUILD_SUCCESS;
+                case(CL_BUILD_NONE):
+                    return BUILD_NONE;
+                case(CL_BUILD_IN_PROGRESS):
+                    return BUILD_IN_PROGRESS;
+                case(CL_BUILD_ERROR):
+                    return BUILD_ERROR;
+// is this a standard state?
+//              case (CL_BUILD_PROGRAM_FAILURE):
+//                    return BUILD_PROGRAM_FAILURE;
+            }
+            return null;
+        }
+    }
+
+    /**
+     * Common compiler options for the OpenCL compiler.
+     */
+    public interface CompilerOptions {
+
+        /**
+         * Treat double precision floating-point constant as single precision constant.
+         */
+        public final static String SINGLE_PRECISION_CONSTANTS = "-cl-single-precision-constant";
+
+        /**
+         * This option controls how single precision and double precision denormalized numbers are handled.
+         * If specified as a build option, the single precision denormalized numbers may be flushed to zero
+         * and if the optional extension for double precision is supported, double precision denormalized numbers
+         * may also be flushed to zero. This is intended to be a performance hint and the OpenCL compiler can choose
+         * not to flush denorms to zero if the device supports single precision (or double precision) denormalized numbers.<br>
+         * This option is ignored for single precision numbers if the device does not support single precision denormalized
+         * numbers i.e. {@link CLDevice.FPConfig#DENORM} is not present in the set returned by {@link CLDevice#getSingleFPConfig()}<br>
+         * This option is ignored for double precision numbers if the device does not support double precision or if it does support
+         * double precision but {@link CLDevice.FPConfig#DENORM} is not present in the set returned by {@link CLDevice#getDoubleFPConfig()}.<br>
+         * This flag only applies for scalar and vector single precision floating-point variables and computations on
+         * these floating-point variables inside a program. It does not apply to reading from or writing to image objects.
+         */
+        public final static String DENORMS_ARE_ZERO = "-cl-denorms-are-zero";
+
+        /**
+         * This option disables all optimizations. The default is optimizations are enabled.
+         */
+        public final static String DISABLE_OPT = "-cl-opt-disable";
+
+        /**
+         * This option allows the compiler to assume the strictest aliasing rules.
+         */
+        public final static String STRICT_ALIASING = "-cl-strict-aliasing";
+
+        /**
+         * Allow a * b + c to be replaced by a mad. The mad computes a * b + c with reduced accuracy.
+         * For example, some OpenCL devices implement mad as truncate the result of a * b before adding it to c.
+         */
+        public final static String ENABLE_MAD = "-cl-mad-enable";
+
+        /**
+         * Allow optimizations for floating-point arithmetic that ignore the signedness of zero.
+         * IEEE 754 arithmetic specifies the behavior of distinct +0.0 and -0.0 values, which then prohibits
+         * simplification of expressions such as x+0.0 or 0.0*x (even with -cl-finite-math-only ({@link #FINITE_MATH_ONLY})).
+         * This option implies that the sign of a zero result isn't significant.
+         */
+        public final static String NO_SIGNED_ZEROS = "-cl-no-signed-zeros";
+
+        /**
+         * Allow optimizations for floating-point arithmetic that<br>
+         * (a) assume that arguments and results are valid,<br>
+         * (b) may violate IEEE 754 standard and<br>
+         * (c) may violate the OpenCL numerical compliance requirements as defined in section
+         * 7.4 for single-precision floating-point, section 9.3.9 for double-precision floating-point,
+         * and edge case behavior in section 7.5.
+         * This option includes the -cl-no-signed-zeros ({@link #NO_SIGNED_ZEROS})
+         * and -cl-mad-enable ({@link #ENABLE_MAD}) options.
+         */
+        public final static String UNSAFE_MATH = "-cl-unsafe-math-optimizations";
+
+        /**
+         * Allow optimizations for floating-point arithmetic that assume that arguments and results are not NaNs or ±∞.
+         * This option may violate the OpenCL numerical compliance requirements defined in in section 7.4 for
+         * single-precision floating-point, section 9.3.9 for double-precision floating-point, and edge case behavior in section 7.5.
+         */
+        public final static String FINITE_MATH_ONLY = "-cl-finite-math-only";
+
+        /**
+         * Sets the optimization options -cl-finite-math-only ({@link #FINITE_MATH_ONLY}) and -cl-unsafe-math-optimizations ({@link #UNSAFE_MATH}).
+         * This allows optimizations for floating-point arithmetic that may violate the IEEE 754
+         * standard and the OpenCL numerical compliance requirements defined in the specification
+         * in section 7.4 for single-precision floating-point, section 9.3.9 for double-precision
+         * floating-point, and edge case behavior in section 7.5. This option causes the preprocessor
+         * macro __FAST_RELAXED_MATH__ to be defined in the OpenCL program.
+         */
+        public final static String FAST_RELAXED_MATH = "-cl-fast-relaxed-math";
+
+        /**
+         * Inhibit all warning messages.
+         */
+        public final static String DISABLE_WARNINGS = "-w";
+
+        /**
+         * Make all warnings into errors.
+         */
+        public final static String WARNINGS_ARE_ERRORS = "-Werror";
+
+    }
+
+}
diff --git a/src/com/jogamp/opencl/CLProgramBuilder.java b/src/com/jogamp/opencl/CLProgramBuilder.java
new file mode 100644
index 00000000..855d9837
--- /dev/null
+++ b/src/com/jogamp/opencl/CLProgramBuilder.java
@@ -0,0 +1,325 @@
+package com.jogamp.opencl;
+
+import com.jogamp.opencl.util.CLBuildConfiguration;
+import com.jogamp.opencl.util.CLProgramConfiguration;
+import java.io.IOException;
+import java.io.ObjectInputStream;
+import java.io.ObjectOutputStream;
+import java.io.Serializable;
+import java.util.ArrayList;
+import java.util.Arrays;
+import java.util.Iterator;
+import java.util.LinkedHashMap;
+import java.util.LinkedHashSet;
+import java.util.List;
+import java.util.Map;
+import java.util.Set;
+
+
+/**
+ * CLProgramBuilder is a helper for building programs with more complex configurations or
+ * building multiple programs with similar configurations.
+ * @see CLProgram#prepare()
+ * @see #createConfiguration()
+ * @see #createConfiguration(com.jogamp.opencl.CLProgram)
+ * @see #loadConfiguration(java.io.ObjectInputStream)
+ * @see #loadConfiguration(java.io.ObjectInputStream, com.jogamp.opencl.CLContext)
+ * @author Michael Bien
+ */
+public final class CLProgramBuilder implements CLProgramConfiguration, Serializable {
+
+    static final long serialVersionUID = 42;
+
+    private static final byte[] NO_BINARIES = new byte[0];
+
+    private transient CLProgram program;
+    private transient Map<CLDevice, byte[]> binariesMap = new LinkedHashMap<CLDevice, byte[]>();
+
+    private String source;
+
+    private final Set<String> optionSet = new LinkedHashSet<String>();
+    private final Set<String> defineSet = new LinkedHashSet<String>();
+
+
+    private CLProgramBuilder() {
+        this(null);
+    }
+
+    private CLProgramBuilder(CLProgram program) {
+        this(program, null, null);
+    }
+
+    private CLProgramBuilder(CLProgram program, String source, Map<CLDevice, byte[]> map) {
+        this.program = program;
+        this.source = source;
+        if(map != null) {
+            this.binariesMap.putAll(map);
+        }
+    }
+
+    /**
+     * Creates a new CLBuildConfiguration.
+     */
+    public static CLBuildConfiguration createConfiguration() {
+        return createConfiguration(null);
+    }
+
+    /**
+     * Creates a new CLProgramConfiguration for this program.
+     */
+    public static CLProgramConfiguration createConfiguration(CLProgram program) {
+        return new CLProgramBuilder(program);
+    }
+
+    /**
+     * Loads a CLBuildConfiguration.
+     * @param ois The ObjectInputStream for reading the object.
+     */
+    public static CLBuildConfiguration loadConfiguration(ObjectInputStream ois) throws IOException, ClassNotFoundException {
+        return (CLBuildConfiguration) ois.readObject();
+    }
+
+    /**
+     * Loads a CLProgramConfiguration containing a CLProgram.
+     * The CLProgram is initialized and ready to be build after this method call.
+     * This method preferes program initialization from binaries if this fails or if
+     * no binaries have been found, it will try to load the program from sources. If
+     * This also fails an appropriate exception will be thrown.
+     * @param ois The ObjectInputStream for reading the object.
+     * @param context The context used for program initialization.
+     */
+    public static CLProgramConfiguration loadConfiguration(ObjectInputStream ois, CLContext context) throws IOException, ClassNotFoundException {
+        CLProgramBuilder config = (CLProgramBuilder) ois.readObject();
+        if(config.binariesMap.size() > 0 && config.binariesMap.values().iterator().next().length > 0) {
+            try{
+                config.program = context.createProgram(config.binariesMap);
+            }catch(CLException.CLInvalidBinaryException ex) {
+                if(config.source != null) {
+                    config.program = context.createProgram(config.source);
+                }else{
+                    throw new IOException("Program configuration contained invalid program binaries and no source.", ex);
+                }
+            }
+        }else if(config.source != null) {
+            config.program = context.createProgram(config.source);
+        }else{
+            throw new IOException("Program configuration did not contain program sources or binaries");
+        }
+        return config;
+    }
+
+    @Override
+    public void save(ObjectOutputStream oos) throws IOException {
+        if(program != null) {
+            this.source = program.getSource();
+            if(program.isExecutable()) {
+                binariesMap = program.getBinaries();
+            }
+        }
+        oos.writeObject(this);
+    }
+
+
+    @Override
+    public CLProgramBuilder withOption(String option) {
+        optionSet.add(option);
+        return this;
+    }
+
+    @Override
+    public CLProgramBuilder withOptions(String... options) {
+        for (String option : options) {
+            optionSet.add(option);
+        }
+        return this;
+    }
+
+    @Override
+    public CLProgramBuilder withDefine(String name) {
+        defineSet.add(CLProgram.define(name));
+        return this;
+    }
+
+    @Override
+    public CLProgramBuilder withDefines(String... names) {
+        for (String name : names) {
+            defineSet.add(CLProgram.define(name));
+        }
+        return this;
+    }
+
+    @Override
+    public CLProgramBuilder withDefine(String name, Object value) {
+        defineSet.add(CLProgram.define(name, value.toString()));
+        return this;
+    }
+
+    @Override
+    public CLProgramBuilder withDefines(Map<String, ? extends Object> defines) {
+        for (String name : defines.keySet()) {
+            defineSet.add(CLProgram.define(name, defines.get(name)));
+        }
+        return this;
+    }
+
+    @Override
+    public CLProgramBuilder forDevice(CLDevice device) {
+        binariesMap.put(device, NO_BINARIES);
+        return this;
+    }
+
+    @Override
+    public CLProgramBuilder forDevices(CLDevice... devices) {
+        for (CLDevice device : devices) {
+            binariesMap.put(device, NO_BINARIES);
+        }
+        return this;
+    }
+
+    @Override
+    public CLProgram build() {
+        return build(program);
+    }
+
+    @Override
+    public CLProgram build(CLProgram program) {
+        if(program == null) {
+            throw new NullPointerException("no program has been set");
+        }
+        List<String> setup = new ArrayList<String>();
+        setup.addAll(optionSet);
+        setup.addAll(defineSet);
+        String options = CLProgram.optionsOf(setup.toArray(new String[setup.size()]));
+        CLDevice[] devices = binariesMap.keySet().toArray(new CLDevice[binariesMap.size()]);
+        return program.build(options, devices);
+    }
+
+    @Override
+    public CLProgramBuilder reset() {
+        resetOptions();
+        resetDefines();
+        resetDevices();
+        return this;
+    }
+
+    @Override
+    public CLProgramConfiguration resetDefines() {
+        defineSet.clear();
+        return this;
+    }
+
+    @Override
+    public CLProgramConfiguration resetDevices() {
+        binariesMap.clear();
+        return this;
+    }
+
+    @Override
+    public CLProgramConfiguration resetOptions() {
+        optionSet.clear();
+        return this;
+    }
+
+    private void writeObject(ObjectOutputStream out) throws IOException {
+        out.defaultWriteObject();
+        out.writeInt(binariesMap.size());
+
+        for (CLDevice device : binariesMap.keySet()) {
+            byte[] binaries = binariesMap.get(device);
+            out.writeLong(device.ID);
+            out.writeInt(binaries.length);
+            out.write(binaries);
+        }
+    }
+
+    private void readObject(ObjectInputStream in) throws IOException, ClassNotFoundException {
+        in.defaultReadObject();
+        this.binariesMap = new LinkedHashMap<CLDevice, byte[]>();
+        int mapSize = in.readInt();
+
+        for (int i = 0; i < mapSize; i++) {
+            long deviceID = in.readLong();
+            int length = in.readInt();
+            byte[] binaries = new byte[length];
+            in.readFully(binaries);
+
+            CLDevice device = new CLDevice(CLPlatform.getLowLevelCLInterface(), deviceID);
+            binariesMap.put(device, binaries);
+        }
+    }
+
+    @Override
+    public CLProgramBuilder asBuildConfiguration() {
+        CLProgramBuilder builder = new CLProgramBuilder();
+        builder.defineSet.addAll(defineSet);
+        builder.optionSet.addAll(optionSet);
+        return builder;
+    }
+    
+    @Override
+    public CLProgramBuilder clone() {
+        CLProgramBuilder builder = new CLProgramBuilder(program, source, binariesMap);
+        builder.defineSet.addAll(defineSet);
+        builder.optionSet.addAll(optionSet);
+        return builder;
+    }
+
+    public CLProgram getProgram() {
+        return program;
+    }
+
+    public CLProgramBuilder setProgram(CLProgram program) {
+        this.program = program;
+        return this;
+    }
+
+    @Override
+    public String toString() {
+        final StringBuilder sb = new StringBuilder();
+        sb.append("CLProgramBuilder");
+        sb.append("{options=").append(optionSet);
+        sb.append(", defines=").append(defineSet);
+        sb.append(", devices=").append(binariesMap);
+        sb.append('}');
+        return sb.toString();
+    }
+
+    @Override
+    public boolean equals(Object o) {
+        if (this == o) return true;
+        if (o == null || getClass() != o.getClass()) return false;
+
+        CLProgramBuilder that = (CLProgramBuilder) o;
+
+        if (source    != null ? !source.equals(that.source)       : that.source    != null) return false;
+        if (defineSet != null ? !defineSet.equals(that.defineSet) : that.defineSet != null) return false;
+        if (optionSet != null ? !optionSet.equals(that.optionSet) : that.optionSet != null) return false;
+
+        if(binariesMap != null && that.binariesMap != null) {
+            if(binariesMap.size() != that.binariesMap.size()) {
+                return false;
+            }
+            Iterator<CLDevice> iterator0 = binariesMap.keySet().iterator();
+            Iterator<CLDevice> iterator1 = that.binariesMap.keySet().iterator();
+            for (int i = 0; i < binariesMap.size(); i++) {
+                CLDevice device0 = iterator0.next();
+                CLDevice device1 = iterator1.next();
+                if(!device0.equals(device1) || !Arrays.equals(binariesMap.get(device0), that.binariesMap.get(device1)))
+                    return false;
+            }
+        }else if(binariesMap != null || that.binariesMap != null) {
+            return false;
+        }
+
+        return true;
+    }
+
+    @Override
+    public int hashCode() {
+        int result = optionSet != null ? optionSet.hashCode() : 0;
+        result = 31 * result + (defineSet != null ? defineSet.hashCode() : 0);
+        result = 31 * result + (binariesMap != null ? binariesMap.hashCode() : 0);
+        return result;
+    }
+
+}
diff --git a/src/com/jogamp/opencl/CLResource.java b/src/com/jogamp/opencl/CLResource.java
new file mode 100644
index 00000000..63b4749e
--- /dev/null
+++ b/src/com/jogamp/opencl/CLResource.java
@@ -0,0 +1,20 @@
+package com.jogamp.opencl;
+
+/**
+ * Releasable OpenCL resource.
+ * @author Michael Bien
+ */
+public interface CLResource extends Disposable<CLException> {
+
+    /**
+     * Releases the OpenCL resource.
+     */
+    public void release();
+
+    /**
+     * Calls {@link #release()};
+     * @see #release()
+     */
+    @Override public void close();
+
+}
diff --git a/src/com/jogamp/opencl/CLSampler.java b/src/com/jogamp/opencl/CLSampler.java
new file mode 100644
index 00000000..79ef6067
--- /dev/null
+++ b/src/com/jogamp/opencl/CLSampler.java
@@ -0,0 +1,123 @@
+package com.jogamp.opencl;
+
+import com.jogamp.common.nio.Int64Buffer;
+
+import java.nio.Buffer;
+
+import static com.jogamp.opencl.CLException.*;
+import static com.jogamp.opencl.CL.*;
+import static com.jogamp.opencl.util.CLUtil.*;
+
+/**
+ * Object representing an OpenCL sampler.
+ * @see CLContext#createSampler(com.jogamp.opencl.CLSampler.AddressingMode, com.jogamp.opencl.CLSampler.FilteringMode, boolean)
+ * @author Michael Bien
+ */
+public class CLSampler extends CLObject implements CLResource {
+
+    private final CLSamplerInfoAccessor samplerInfo;
+
+    private CLSampler(CLContext context, long id,  AddressingMode addrMode, FilteringMode filtMode, boolean normalizedCoords) {
+        super(context, id);
+        this.samplerInfo = new CLSamplerInfoAccessor();
+    }
+
+    static CLSampler create(CLContext context, AddressingMode addrMode, FilteringMode filtMode, boolean normalizedCoords) {
+        int[] error = new int[1];
+
+        long id = context.cl.clCreateSampler(context.ID, clBoolean(normalizedCoords), addrMode.MODE, filtMode.MODE, error, 0);
+
+        checkForError(error[0], "can not create sampler");
+        return new CLSampler(context, id, addrMode, filtMode, normalizedCoords);
+    }
+
+    public FilteringMode getFilteringMode() {
+        int info = (int)samplerInfo.getLong(CL_SAMPLER_FILTER_MODE);
+        return FilteringMode.valueOf(info);
+    }
+
+    public AddressingMode getAddressingMode() {
+        int info = (int)samplerInfo.getLong(CL_SAMPLER_ADDRESSING_MODE);
+        return AddressingMode.valueOf(info);
+    }
+
+    public boolean hasNormalizedCoords() {
+        return samplerInfo.getLong(CL_SAMPLER_NORMALIZED_COORDS) == CL_TRUE;
+    }
+
+    public void release() {
+        int ret = cl.clReleaseSampler(ID);
+        context.onSamplerReleased(this);
+        checkForError(ret, "can not release sampler");
+    }
+
+    public void close() {
+        release();
+    }
+
+    private class CLSamplerInfoAccessor extends CLInfoAccessor {
+
+        @Override
+        protected int getInfo(int name, long valueSize, Buffer value, Int64Buffer valueSizeRet) {
+            return cl.clGetSamplerInfo(ID, name, valueSize, value, valueSizeRet);
+        }
+
+    }
+
+    public enum FilteringMode {
+
+        NEAREST(CL_FILTER_NEAREST),
+        LINEAR(CL_FILTER_LINEAR);
+
+        /**
+         * Value of wrapped OpenCL sampler filtering mode type.
+         */
+        public final int MODE;
+
+        private FilteringMode(int mode) {
+            this.MODE = mode;
+        }
+
+        public static FilteringMode valueOf(int mode) {
+            switch(mode) {
+                case(CL_FILTER_NEAREST):
+                    return NEAREST;
+                case(CL_FILTER_LINEAR):
+                    return LINEAR;
+            }
+            return null;
+        }
+    }
+
+    public enum AddressingMode {
+
+        REPEAT(CL_ADDRESS_REPEAT),
+        CLAMP_TO_EDGE(CL_ADDRESS_CLAMP_TO_EDGE),
+        CLAMP(CL_ADDRESS_CLAMP),
+        NONE(CL_ADDRESS_NONE);
+
+        /**
+         * Value of wrapped OpenCL sampler addressing mode type.
+         */
+        public final int MODE;
+
+        private AddressingMode(int mode) {
+            this.MODE = mode;
+        }
+
+        public static AddressingMode valueOf(int mode) {
+            switch(mode) {
+                case(CL_ADDRESS_REPEAT):
+                    return REPEAT;
+                case(CL_ADDRESS_CLAMP_TO_EDGE):
+                    return CLAMP_TO_EDGE;
+                case(CL_ADDRESS_CLAMP):
+                    return CLAMP;
+                case(CL_ADDRESS_NONE):
+                    return NONE;
+            }
+            return null;
+        }
+    }
+
+}
diff --git a/src/com/jogamp/opencl/CreateContextCallback.java b/src/com/jogamp/opencl/CreateContextCallback.java
new file mode 100644
index 00000000..dc25860a
--- /dev/null
+++ b/src/com/jogamp/opencl/CreateContextCallback.java
@@ -0,0 +1,14 @@
+package com.jogamp.opencl;
+
+import java.nio.ByteBuffer;
+
+/**
+ *
+ * @author Michael Bien
+ */
+// TODO implement callbacks
+public interface CreateContextCallback {
+
+    public void createContextCallback(String errinfo, ByteBuffer private_info, long cb, Object user_data);
+
+}
diff --git a/src/com/jogamp/opencl/Disposable.java b/src/com/jogamp/opencl/Disposable.java
new file mode 100644
index 00000000..2cd1784b
--- /dev/null
+++ b/src/com/jogamp/opencl/Disposable.java
@@ -0,0 +1,21 @@
+package com.jogamp.opencl;
+
+/*
+ * JDK7 ARM proposal, cypied to be forward compatible with java 7 automatic resource managment blocks.
+ * @author Michael Bien
+ */
+
+//package java.lang;
+
+/**
+ * A resource that must be closed when it is no longer needed.
+ *
+ * @param X the type of exception thrown by the close method (or
+ *     {@link RuntimeException} if the close method is not permitted
+ *     to throw any checked exceptions).
+ */
+/*public*/ interface Disposable<X extends Throwable> {
+
+    void close() throws X;
+
+}
diff --git a/src/com/jogamp/opencl/InternalBufferUtil.java b/src/com/jogamp/opencl/InternalBufferUtil.java
new file mode 100644
index 00000000..9cd1fac9
--- /dev/null
+++ b/src/com/jogamp/opencl/InternalBufferUtil.java
@@ -0,0 +1,39 @@
+package com.jogamp.opencl;
+
+import java.lang.reflect.Field;
+import java.nio.Buffer;
+import sun.misc.Unsafe;
+
+/**
+ *
+ * @author Michael Bien
+ */
+class InternalBufferUtil {
+
+    private static final long addressFieldOffset;
+    private static Unsafe unsafe;
+
+    static {
+        try {
+            Field f = Buffer.class.getDeclaredField("address");
+
+            Field[] fields = Unsafe.class.getDeclaredFields();
+            for (int i = 0; i < fields.length; i++) {
+                if (fields[i].getName().equals("theUnsafe")) {
+                    fields[i].setAccessible(true);
+                    unsafe = (Unsafe)fields[i].get(Unsafe.class);
+                    break;
+                }
+            }
+
+            addressFieldOffset = unsafe.objectFieldOffset(f);
+        } catch (Exception e) {
+            throw new RuntimeException(e);
+        }
+    }
+
+    public static long getDirectBufferAddress(Buffer buffer) {
+        return ((buffer == null) ? 0 : unsafe.getLong(buffer, addressFieldOffset));
+    }
+
+}
\ No newline at end of file
diff --git a/src/com/jogamp/opencl/NativeLibLoader.java b/src/com/jogamp/opencl/NativeLibLoader.java
new file mode 100644
index 00000000..67824b95
--- /dev/null
+++ b/src/com/jogamp/opencl/NativeLibLoader.java
@@ -0,0 +1,21 @@
+package com.jogamp.opencl;
+
+import java.security.AccessController;
+import java.security.PrivilegedAction;
+import com.jogamp.nativewindow.impl.NativeLibLoaderBase;
+
+/**
+ *
+ * @author Michael Bien
+ */
+class NativeLibLoader extends NativeLibLoaderBase {
+
+    public static void loadJOCL() {
+        AccessController.doPrivileged(new PrivilegedAction<Object>() {
+            public Object run() {
+                loadLibrary("jocl", null, true);
+                return null;
+            }
+        });
+    }
+}
diff --git a/src/com/jogamp/opencl/gl/CLGLBuffer.java b/src/com/jogamp/opencl/gl/CLGLBuffer.java
new file mode 100644
index 00000000..c7f14c4c
--- /dev/null
+++ b/src/com/jogamp/opencl/gl/CLGLBuffer.java
@@ -0,0 +1,77 @@
+package com.jogamp.opencl.gl;
+
+import com.jogamp.opencl.CL;
+import com.jogamp.opencl.CLBuffer;
+import com.jogamp.opencl.CLContext;
+
+import java.nio.Buffer;
+import javax.media.opengl.GLContext;
+
+
+/**
+ * Shared buffer between OpenGL and OpenCL contexts.
+ * @author Michael Bien
+ */
+public final class CLGLBuffer<B extends Buffer> extends CLBuffer<B> implements CLGLObject {
+
+
+    /**
+     * The OpenGL object handle.
+     */
+    public final int GLID;
+
+    private CLGLBuffer(CLContext context, B directBuffer, long id, int glObject, int flags) {
+        super(context, directBuffer, id, flags);
+        this.GLID = glObject;
+    }
+
+
+    static <B extends Buffer> CLGLBuffer<B> create(CLContext context, B directBuffer, int flags, int glObject) {
+        checkBuffer(directBuffer, flags);
+        
+        CL cl = getCL(context);
+        int[] result = new int[1];
+        CLGLI clgli = (CLGLI)cl;
+        
+        long id = clgli.clCreateFromGLBuffer(context.ID, flags, glObject, result, 0);
+
+        return new CLGLBuffer<B>(context, directBuffer, id, glObject, flags);
+    }
+
+    static <B extends Buffer> void checkBuffer(B directBuffer, int flags) throws IllegalArgumentException {
+        if (directBuffer != null && !directBuffer.isDirect()) {
+            throw new IllegalArgumentException("buffer is not a direct buffer");
+        }
+        if (isHostPointerFlag(flags)) {
+            throw new IllegalArgumentException("CL_MEM_COPY_HOST_PTR or CL_MEM_USE_HOST_PTR can not be used with OpenGL Buffers.");
+        }
+    }
+
+    public int getGLObjectID() {
+        return GLID;
+    }
+
+    public GLObjectType getGLObjectType() {
+        return GLObjectType.GL_OBJECT_BUFFER;
+    }
+
+    @Override
+    public CLGLContext getContext() {
+        return (CLGLContext) super.getContext();
+    }
+
+    public GLContext getGLContext() {
+        return getContext().getGLContext();
+    }
+
+    @Override
+    public <T extends Buffer> CLGLBuffer<T> cloneWith(T directBuffer) {
+        return new CLGLBuffer<T>(context, directBuffer, ID, GLID, FLAGS);
+    }
+
+    @Override
+    public String toString() {
+        return "CLGLBuffer [id: " + ID+" glID: "+GLID+"]";
+    }
+
+}
diff --git a/src/com/jogamp/opencl/gl/CLGLContext.java b/src/com/jogamp/opencl/gl/CLGLContext.java
new file mode 100644
index 00000000..92618729
--- /dev/null
+++ b/src/com/jogamp/opencl/gl/CLGLContext.java
@@ -0,0 +1,250 @@
+package com.jogamp.opencl.gl;
+
+import com.jogamp.opencl.CLContext;
+import com.jogamp.opencl.CLDevice;
+import java.nio.Buffer;
+import com.jogamp.opencl.CLMemory.Mem;
+import com.jogamp.opencl.CLPlatform;
+import com.jogamp.common.nio.PointerBuffer;
+import com.jogamp.opengl.impl.GLContextImpl;
+import com.jogamp.opengl.impl.macosx.cgl.MacOSXCGLContext;
+import com.jogamp.opengl.impl.windows.wgl.WindowsWGLContext;
+import com.jogamp.opengl.impl.x11.glx.X11GLXContext;
+import javax.media.nativewindow.DefaultGraphicsConfiguration;
+import javax.media.opengl.GLContext;
+
+import static com.jogamp.opencl.gl.CLGLI.*;
+
+/**
+ * OpenCL Context supporting JOGL-JOCL interoperablity.
+ * @author Michael Bien
+ */
+public final class CLGLContext extends CLContext {
+
+    final long glID;
+    private final GLContext glContext;
+
+    private CLGLContext(CLPlatform platform, GLContext glContext, long clContextID, long glContextID) {
+        super(platform, clContextID);
+        this.glID = glContextID;
+        this.glContext = glContext;
+    }
+
+    /**
+     * Creates a shared context on all available devices (CL_DEVICE_TYPE_ALL).
+     */
+    public static CLGLContext create(GLContext glContext) {
+        return create(glContext, (CLPlatform)null, CLDevice.Type.ALL);
+    }
+
+    /**
+     * Creates a shared context on the specified platform on all available devices (CL_DEVICE_TYPE_ALL).
+     */
+    public static CLGLContext create(GLContext glContext, CLPlatform platform) {
+        return create(glContext, platform, CLDevice.Type.ALL);
+    }
+
+    /**
+     * Creates a shared context on the specified platform and with the specified
+     * device types.
+     */
+    public static CLGLContext create(GLContext glContext, CLDevice.Type... deviceTypes) {
+        return create(glContext, null, deviceTypes);
+    }
+
+    /**
+     * Creates a shared context on the specified devices.
+     * The platform to be used is implementation dependent.
+     */
+    public static CLGLContext create(GLContext glContext, CLDevice... devices) {
+        return create(glContext, null, devices);
+    }
+
+    /**
+     * Creates a shared context on the specified platform and with the specified
+     * device types.
+     */
+    public static CLGLContext create(GLContext glContext, CLPlatform platform, CLDevice.Type... deviceTypes) {
+
+        if(platform == null) {
+            platform = CLPlatform.getDefault();
+        }
+
+        long[] glID = new long[1];
+        PointerBuffer properties = setupContextProperties(platform, glContext, glID);
+        long clID = createContextFromType(properties, toDeviceBitmap(deviceTypes));
+
+        return new CLGLContext(platform, glContext, clID, glID[0]);
+
+    }
+
+    /**
+     * Creates a shared context on the specified platform and with the specified
+     * devices.
+     */
+    public static CLGLContext create(GLContext glContext, CLPlatform platform, CLDevice... devices) {
+
+        if(platform == null) {
+            platform = CLPlatform.getDefault();
+        }
+
+        long[] glID = new long[1];
+        PointerBuffer properties = setupContextProperties(platform, glContext, glID);
+        long clID = createContext(properties, devices);
+
+        CLGLContext context = new CLGLContext(platform, glContext, clID, glID[0]);
+        if(devices != null) {
+            for (int i = 0; i < devices.length; i++) {
+                context.overrideContext(devices[i]);
+            }
+        }
+        return context;
+    }
+
+
+    private static PointerBuffer setupContextProperties(CLPlatform platform, GLContext glContext, long[] glID) {
+
+        if(platform == null) {
+            throw new RuntimeException("no OpenCL installation found");
+        }
+
+        GLContextImpl ctxImpl = (GLContextImpl)glContext;
+
+        DefaultGraphicsConfiguration config = (DefaultGraphicsConfiguration)ctxImpl.getDrawableImpl()
+             .getNativeWindow().getGraphicsConfiguration().getNativeGraphicsConfiguration();
+
+        PointerBuffer properties;
+        if(glContext instanceof X11GLXContext) {
+            properties = PointerBuffer.allocateDirect(7);
+            long handle = config.getScreen().getDevice().getHandle();
+            glID[0] = ((X11GLXContext)glContext).getContext();
+            properties.put(CL_GL_CONTEXT_KHR).put(glID[0])
+                      .put(CL_GLX_DISPLAY_KHR).put(handle)
+                      .put(CL_CONTEXT_PLATFORM).put(platform.ID);
+        }else if(glContext instanceof WindowsWGLContext) {
+            // TODO test on windows
+            //WIN32
+            //cl_context_properties props[] = {
+            //         CL_GL_CONTEXT_KHR, (cl_context_properties)0,
+            //         CL_WGL_HDC_KHR, (cl_context_properties)0,
+            //         CL_CONTEXT_PLATFORM, (cl_context_properties)cpPlatform, 0};
+            properties = PointerBuffer.allocateDirect(7);
+            long handle = config.getScreen().getDevice().getHandle();
+            glID[0] = ((WindowsWGLContext)glContext).getHGLRC();
+            properties.put(CL_GL_CONTEXT_KHR).put(glID[0])
+                      .put(CL_WGL_HDC_KHR).put(handle)
+                      .put(CL_CONTEXT_PLATFORM).put(platform.ID);
+        }else if(glContext instanceof MacOSXCGLContext) {
+            // TODO test on mac
+            //MACOSX
+            //cl_context_properties props[] = {
+            //         CL_CGL_SHAREGROUP_KHR, (cl_context_properties)0,
+            //         CL_CONTEXT_PLATFORM, (cl_context_properties)cpPlatform, 0};
+            properties = PointerBuffer.allocateDirect(5);
+            glID[0] = ((MacOSXCGLContext)glContext).getCGLContext();
+            properties.put(CL_CGL_SHAREGROUP_KHR).put(glID[0])
+                      .put(CL_CONTEXT_PLATFORM).put(platform.ID);
+        }else{
+            throw new RuntimeException("unsupported GLContext: "+glContext);
+        }
+
+        return (PointerBuffer)properties.put(0).rewind(); // 0 terminated array
+    }
+
+    // Buffers
+    public final CLGLBuffer<?> createFromGLBuffer(int glBuffer, Mem... flags) {
+        return createFromGLBuffer(null, glBuffer, Mem.flagsToInt(flags));
+    }
+
+    public final CLGLBuffer<?> createFromGLBuffer(int glBuffer, int flags) {
+        return createFromGLBuffer(null, glBuffer, flags);
+    }
+
+    public final <B extends Buffer> CLGLBuffer<B> createFromGLBuffer(B directBuffer, int glBuffer, Mem... flags) {
+        return createFromGLBuffer(directBuffer, glBuffer, Mem.flagsToInt(flags));
+    }
+
+    public final <B extends Buffer> CLGLBuffer<B> createFromGLBuffer(B directBuffer, int glBuffer, int flags) {
+        CLGLBuffer<B> buffer = CLGLBuffer.create(this, directBuffer, flags, glBuffer);
+        memoryObjects.add(buffer);
+        return buffer;
+    }
+
+    // Renderbuffers
+    public final CLGLImage2d<?> createFromGLRenderbuffer(int glBuffer, Mem... flags) {
+        return createFromGLRenderbuffer(null, glBuffer, Mem.flagsToInt(flags));
+    }
+
+    public final CLGLImage2d<?> createFromGLRenderbuffer(int glBuffer, int flags) {
+        return createFromGLRenderbuffer(null, glBuffer, flags);
+    }
+
+    public final <B extends Buffer> CLGLImage2d<B> createFromGLRenderbuffer(B directBuffer, int glBuffer, Mem... flags) {
+        return createFromGLRenderbuffer(directBuffer, glBuffer, Mem.flagsToInt(flags));
+    }
+    
+    public final <B extends Buffer> CLGLImage2d<B> createFromGLRenderbuffer(B directBuffer, int glBuffer, int flags) {
+        CLGLImage2d<B> buffer = CLGLImage2d.createFromGLRenderbuffer(this, directBuffer, flags, glBuffer);
+        memoryObjects.add(buffer);
+        return buffer;
+    }
+    
+    //2d Textures
+    public final CLGLTexture2d<?> createFromGLTexture2d(int target, int texture, int mipmap, Mem... flags) {
+        return createFromGLTexture2d(null, target, texture, mipmap, Mem.flagsToInt(flags));
+    }
+
+    public final CLGLTexture2d<?> createFromGLTexture2d(int target, int texture, int mipmap, int flags) {
+        return createFromGLTexture2d(null, target, texture, mipmap, flags);
+    }
+
+    public final <B extends Buffer> CLGLTexture2d<B> createFromGLTexture2d(B directBuffer, int target, int texture, int mipmap, Mem... flags) {
+        return createFromGLTexture2d(directBuffer, target, texture, mipmap, Mem.flagsToInt(flags));
+    }
+    
+    public final <B extends Buffer> CLGLTexture2d<B> createFromGLTexture2d(B directBuffer, int target, int texture, int mipmap, int flags) {
+        CLGLTexture2d<B> buffer = CLGLTexture2d.createFromGLTexture2d(this, directBuffer, target, texture, mipmap, flags);
+        memoryObjects.add(buffer);
+        return buffer;
+    }
+    
+    //3d Textures
+    public final CLGLTexture3d<?> createFromGLTexture3d(int target, int texture, int mipmap, Mem... flags) {
+        return createFromGLTexture3d(null, target, texture, mipmap, Mem.flagsToInt(flags));
+    }
+
+    public final CLGLTexture3d<?> createFromGLTexture3d(int target, int texture, int mipmap, int flags) {
+        return createFromGLTexture3d(null, target, texture, mipmap, flags);
+    }
+
+    public final <B extends Buffer> CLGLTexture3d<B> createFromGLTexture3d(B directBuffer, int target, int texture, int mipmap, Mem... flags) {
+        return createFromGLTexture3d(directBuffer, target, texture, mipmap, Mem.flagsToInt(flags));
+    }
+
+    public final <B extends Buffer> CLGLTexture3d<B> createFromGLTexture3d(B directBuffer, int target, int texture, int mipmap, int flags) {
+        CLGLTexture3d<B> buffer = CLGLTexture3d.createFromGLTexture3d(this, directBuffer, target, texture, mipmap, flags);
+        memoryObjects.add(buffer);
+        return buffer;
+    }
+
+    /**
+     * Return the low level OpenCL interface with OpenGL interoperability.
+     */
+    @Override
+    public CLGLI getCL() {
+        return (CLGLI)super.getCL();
+    }
+
+    /**
+     * Returns the OpenGL context this context was shared with.
+     */
+    public GLContext getGLContext() {
+        return glContext;
+    }
+
+    @Override
+    public CLGLContext getContext() {
+        return this;
+    }
+
+}
diff --git a/src/com/jogamp/opencl/gl/CLGLImage2d.java b/src/com/jogamp/opencl/gl/CLGLImage2d.java
new file mode 100644
index 00000000..82021816
--- /dev/null
+++ b/src/com/jogamp/opencl/gl/CLGLImage2d.java
@@ -0,0 +1,73 @@
+package com.jogamp.opencl.gl;
+
+import com.jogamp.opencl.CL;
+import com.jogamp.opencl.CLContext;
+import com.jogamp.opencl.CLImage2d;
+import com.jogamp.opencl.CLImageFormat;
+import com.jogamp.opencl.impl.CLImageFormatImpl;
+import java.nio.Buffer;
+import javax.media.opengl.GLContext;
+
+import static com.jogamp.opencl.CL.*;
+
+/**
+ * 2D OpenCL image representing an OpenGL renderbuffer.
+ * @author Michael Bien
+ */
+public class CLGLImage2d<B extends Buffer> extends CLImage2d<B> implements CLGLObject {
+
+    /**
+     * The OpenGL object handle.
+     */
+    public final int GLID;
+
+    protected CLGLImage2d(CLContext context, B directBuffer, CLImageFormat format, CLImageInfoAccessor accessor, int width, int height, long id, int glid, int flags) {
+        super(context, directBuffer, format, accessor, width, height, id, flags);
+        this.GLID = glid;
+    }
+
+    static <B extends Buffer> CLGLImage2d<B> createFromGLRenderbuffer(CLContext context, B directBuffer, int flags, int glObject) {
+
+        CLGLBuffer.checkBuffer(directBuffer, flags);
+
+        CL cl = getCL(context);
+        int[] result = new int[1];
+        CLGLI clgli = (CLGLI)cl;
+
+        long id = clgli.clCreateFromGLRenderbuffer(context.ID, flags, glObject, result, 0);
+
+        return createImage(context, id, directBuffer, glObject, flags);
+    }
+
+    static <B extends Buffer> CLGLImage2d<B> createImage(CLContext context, long id, B directBuffer, int glObject, int flags) {
+        CLImageInfoAccessor accessor = new CLImageInfoAccessor(getCL(context), id);
+
+        CLImageFormat format = createUninitializedImageFormat();
+        accessor.getInfo(CL_IMAGE_FORMAT, CLImageFormatImpl.size(), format.getFormatImpl().getBuffer(), null);
+
+        int width = (int)accessor.getLong(CL_IMAGE_WIDTH);
+        int height = (int)accessor.getLong(CL_IMAGE_HEIGHT);
+
+        return new CLGLImage2d<B>(context, directBuffer, format, accessor, width, height, id, glObject, flags);
+    }
+
+    @Override
+    public GLObjectType getGLObjectType() {
+        return GLObjectType.GL_OBJECT_RENDERBUFFER;
+    }
+
+    @Override
+    public int getGLObjectID() {
+        return GLID;
+    }
+
+    @Override
+    public CLGLContext getContext() {
+        return (CLGLContext) super.getContext();
+    }
+
+    public GLContext getGLContext() {
+        return getContext().getGLContext();
+    }
+
+}
diff --git a/src/com/jogamp/opencl/gl/CLGLObject.java b/src/com/jogamp/opencl/gl/CLGLObject.java
new file mode 100644
index 00000000..6e2aa9d2
--- /dev/null
+++ b/src/com/jogamp/opencl/gl/CLGLObject.java
@@ -0,0 +1,35 @@
+/*
+ * Created on Friday, February 26 2010
+ */
+package com.jogamp.opencl.gl;
+
+import com.jogamp.opencl.CLMemory.GLObjectType;
+import javax.media.opengl.GLContext;
+
+/**
+ *
+ * @author Michael Bien
+ */
+interface CLGLObject {
+
+    /**
+     * Returns the OpenGL object id of this shared object.
+     */
+    public int getGLObjectID();
+
+    /**
+     * Returns the OpenGL buffer type of this shared object.
+     */
+    public GLObjectType getGLObjectType();
+    
+    /**
+     * Returns the OpenCL context of this shared object.
+     */
+    public CLGLContext getContext();
+
+    /**
+     * Returns the OpenGL context of this shared object.
+     */
+    public GLContext getGLContext();
+
+}
diff --git a/src/com/jogamp/opencl/gl/CLGLTexture.java b/src/com/jogamp/opencl/gl/CLGLTexture.java
new file mode 100644
index 00000000..312df80b
--- /dev/null
+++ b/src/com/jogamp/opencl/gl/CLGLTexture.java
@@ -0,0 +1,23 @@
+/*
+ * Created on Friday, February 26 2010
+ */
+
+package com.jogamp.opencl.gl;
+
+/**
+ *
+ * @author Michael Bien
+ */
+interface CLGLTexture extends CLGLObject {
+
+    /**
+     * Returns the OpenGL texture target of this texture.
+     */
+    public int getTextureTarget();
+
+    /**
+     * Returns the OpenGL mipmap level of this texture.
+     */
+    public int getMipMapLevel();
+
+}
diff --git a/src/com/jogamp/opencl/gl/CLGLTexture2d.java b/src/com/jogamp/opencl/gl/CLGLTexture2d.java
new file mode 100644
index 00000000..88b32286
--- /dev/null
+++ b/src/com/jogamp/opencl/gl/CLGLTexture2d.java
@@ -0,0 +1,63 @@
+package com.jogamp.opencl.gl;
+
+import com.jogamp.opencl.CL;
+import com.jogamp.opencl.CLContext;
+import com.jogamp.opencl.CLImageFormat;
+import com.jogamp.opencl.impl.CLImageFormatImpl;
+import java.nio.Buffer;
+
+import static com.jogamp.opencl.CL.*;
+
+/**
+ * 2D OpenCL image representing an 2D OpenGL texture.
+ * @author Michael Bien
+ */
+public class CLGLTexture2d<B extends Buffer> extends CLGLImage2d<B> implements CLGLTexture {
+
+    public final int target;
+    
+    public final int mipMapLevel;
+
+    public CLGLTexture2d(CLContext context, B directBuffer, CLImageFormat format, CLImageInfoAccessor accessor, int target, int mipLevel, int width, int height, long id, int glid, int flags) {
+        super(context, directBuffer, format, accessor, width, height, id, glid, flags);
+        this.target = target;
+        this.mipMapLevel = mipLevel;
+    }
+
+    static <B extends Buffer> CLGLTexture2d<B> createFromGLTexture2d(CLContext context, B directBuffer, int target, int texture, int mipLevel, int flags) {
+
+        CLGLBuffer.checkBuffer(directBuffer, flags);
+
+        CL cl = getCL(context);
+        int[] result = new int[1];
+        CLGLI clgli = (CLGLI)cl;
+
+        long id = clgli.clCreateFromGLTexture2D(context.ID, flags, target, mipLevel, texture, result, 0);
+
+        CLImageInfoAccessor accessor = new CLImageInfoAccessor(cl, id);
+
+        CLImageFormat format = createUninitializedImageFormat();
+        accessor.getInfo(CL_IMAGE_FORMAT, CLImageFormatImpl.size(), format.getFormatImpl().getBuffer(), null);
+
+        int width = (int)accessor.getLong(CL_IMAGE_WIDTH);
+        int height = (int)accessor.getLong(CL_IMAGE_HEIGHT);
+
+        return new CLGLTexture2d<B>(context, directBuffer, format, accessor, target, mipLevel, width, height, id, width, flags);
+
+    }
+
+    public int getTextureTarget() {
+        return target;
+    }
+
+    public int getMipMapLevel() {
+        return mipMapLevel;
+    }
+
+    @Override
+    public GLObjectType getGLObjectType() {
+        return GLObjectType.GL_OBJECT_TEXTURE2D;
+    }
+
+
+}
diff --git a/src/com/jogamp/opencl/gl/CLGLTexture3d.java b/src/com/jogamp/opencl/gl/CLGLTexture3d.java
new file mode 100644
index 00000000..8731b5e7
--- /dev/null
+++ b/src/com/jogamp/opencl/gl/CLGLTexture3d.java
@@ -0,0 +1,86 @@
+package com.jogamp.opencl.gl;
+
+import com.jogamp.opencl.CL;
+import com.jogamp.opencl.CLContext;
+import com.jogamp.opencl.CLImage3d;
+import com.jogamp.opencl.CLImageFormat;
+import com.jogamp.opencl.impl.CLImageFormatImpl;
+import java.nio.Buffer;
+import javax.media.opengl.GLContext;
+
+import static com.jogamp.opencl.CL.*;
+
+/**
+ * 3D OpenCL image representing an 3D OpenGL texture.
+ * @author Michael Bien
+ */
+public class CLGLTexture3d<B extends Buffer> extends CLImage3d<B> implements CLGLObject, CLGLTexture {
+
+    /**
+     * The OpenGL object handle.
+     */
+    public final int GLID;
+
+    public final int target;
+    
+    public final int mipMapLevel;
+
+    private CLGLTexture3d(CLContext context, B directBuffer, CLImageFormat format, CLImageInfoAccessor accessor, int target, int mipLevel, int width, int height, int depth, long id, int glid, int flags) {
+        super(context, directBuffer, format, accessor, width, height, depth, id, flags);
+        this.GLID = glid;
+        this.target = target;
+        this.mipMapLevel = mipLevel;
+    }
+
+    static <B extends Buffer> CLGLTexture3d<B> createFromGLTexture3d(CLContext context, B directBuffer, int flags, int target, int mipLevel, int texture) {
+
+        CLGLBuffer.checkBuffer(directBuffer, flags);
+
+        CL cl = getCL(context);
+        int[] result = new int[1];
+        CLGLI clgli = (CLGLI)cl;
+
+        long id = clgli.clCreateFromGLTexture3D(context.ID, flags, target, mipLevel, texture, result, 0);
+
+        CLImageInfoAccessor accessor = new CLImageInfoAccessor(cl, id);
+
+        CLImageFormat format = createUninitializedImageFormat();
+        accessor.getInfo(CL_IMAGE_FORMAT, CLImageFormatImpl.size(), format.getFormatImpl().getBuffer(), null);
+
+        int width = (int)accessor.getLong(CL_IMAGE_WIDTH);
+        int height = (int)accessor.getLong(CL_IMAGE_HEIGHT);
+        int depth = (int)accessor.getLong(CL_IMAGE_DEPTH);
+
+        return new CLGLTexture3d<B>(context, directBuffer, format, accessor, target, mipLevel, width, height, depth, id, texture, flags);
+    }
+
+    public int getGLObjectID() {
+        return GLID;
+    }
+
+    public int getTextureTarget() {
+        return target;
+    }
+
+    public int getMipMapLevel() {
+        return mipMapLevel;
+    }
+
+    public GLObjectType getGLObjectType() {
+        return GLObjectType.GL_OBJECT_TEXTURE3D;
+    }
+
+    /**
+     * Returns the shared CLGLContext.
+     */
+    @Override
+    public CLGLContext getContext() {
+        return (CLGLContext) super.getContext();
+    }
+
+    @Override
+    public GLContext getGLContext() {
+        return getContext().getGLContext();
+    }
+
+}
diff --git a/src/com/jogamp/opencl/util/CLBuildConfiguration.java b/src/com/jogamp/opencl/util/CLBuildConfiguration.java
new file mode 100644
index 00000000..1de62637
--- /dev/null
+++ b/src/com/jogamp/opencl/util/CLBuildConfiguration.java
@@ -0,0 +1,105 @@
+package com.jogamp.opencl.util;
+
+import com.jogamp.opencl.CLDevice;
+import com.jogamp.opencl.CLProgram;
+import java.io.IOException;
+import java.io.ObjectOutputStream;
+import java.util.Map;
+
+/**
+ * Configuration representing everything needed to build an OpenCL program.
+ * @author Michael Bien
+ * @see com.jogamp.opencl.CLProgramBuilder#createConfiguration()
+ * @see com.jogamp.opencl.CLProgramBuilder#loadConfiguration(java.io.ObjectInputStream)
+ */
+public interface CLBuildConfiguration extends Cloneable {
+
+    /**
+     * Builds or rebuilds the program.
+     * @param program The program which should be build.
+     */
+    public CLProgram build(CLProgram program);
+
+    /**
+     * Sets the program which should be build.
+     */
+    public CLProgramConfiguration setProgram(CLProgram program);
+
+    /**
+     * Adds the device as build target.
+     */
+    public CLBuildConfiguration forDevice(CLDevice device);
+
+    /**
+     * Adds the devices as build target.
+     */
+    public CLBuildConfiguration forDevices(CLDevice... devices);
+
+    /**
+     * Resets this builder's configuration like options, devices and definitions.
+     */
+    public CLBuildConfiguration reset();
+
+    /**
+     * Resets this builder's configuration options.
+     */
+    public CLBuildConfiguration resetOptions();
+
+    /**
+     * Resets this builder's macro definitions.
+     */
+    public CLBuildConfiguration resetDefines();
+
+    /**
+     * Resets this builder's device list.
+     */
+    public CLBuildConfiguration resetDevices();
+
+    /**
+     * Adds the definition to the build configuration.
+     * @see CLProgram#define(java.lang.String)
+     */
+    public CLBuildConfiguration withDefine(String name);
+
+    /**
+     * Adds the definition to the build configuration.
+     * @see CLProgram#define(java.lang.String, java.lang.Object)
+     */
+    public CLBuildConfiguration withDefine(String name, Object value);
+
+    /**
+     * Adds the definitions to the build configuration.
+     * @see com.jogamp.opencl.CLProgram#define(java.lang.String)
+     */
+    public CLBuildConfiguration withDefines(String... names);
+
+    /**
+     * Adds the definitions to the build configuration.
+     * @see com.jogamp.opencl.CLProgram#define(java.lang.String, java.lang.Object)
+     */
+    public CLBuildConfiguration withDefines(Map<String, ? extends Object> defines);
+
+    /**
+     * Adds the compiler option to the build configuration.
+     * @see com.jogamp.opencl.CLProgram.CompilerOptions
+     */
+    public CLBuildConfiguration withOption(String option);
+
+    /**
+     * Adds the compiler options to the build configuration.
+     * @see com.jogamp.opencl.CLProgram.CompilerOptions
+     */
+    public CLBuildConfiguration withOptions(String... options);
+
+    /**
+     * Clones this configuration.
+     */
+    public CLBuildConfiguration clone();
+
+    /**
+     * Saves this configuration to the ObjectOutputStream.
+     * The caller is responsible for closing the stream.
+     */
+    public void save(ObjectOutputStream oos) throws IOException;
+
+}
diff --git a/src/com/jogamp/opencl/util/CLProgramConfiguration.java b/src/com/jogamp/opencl/util/CLProgramConfiguration.java
new file mode 100644
index 00000000..901e28ce
--- /dev/null
+++ b/src/com/jogamp/opencl/util/CLProgramConfiguration.java
@@ -0,0 +1,50 @@
+package com.jogamp.opencl.util;
+
+import com.jogamp.opencl.CLDevice;
+import com.jogamp.opencl.CLProgram;
+import java.util.Map;
+
+/**
+ * Configuration representing everything needed to build an OpenCL program (program included).
+ * CLProgramConfiguration is a helper for building programs with more complex configurations or
+ * building multiple programs with the similar configuration.
+ * @see CLProgram#prepare()
+ * @see com.jogamp.opencl.CLProgramBuilder#createConfiguration(com.jogamp.opencl.CLProgram)
+ * @see com.jogamp.opencl.CLProgramBuilder#loadConfiguration(java.io.ObjectInputStream, com.jogamp.opencl.CLContext)
+ * @author Michael Bien
+ */
+public interface CLProgramConfiguration extends CLBuildConfiguration {
+
+    /**
+     * Builds or rebuilds a program.
+     */
+    public CLProgram build();
+
+    /**
+     * Returns the program.
+     */
+    public CLProgram getProgram();
+
+    /**
+     * Returns a new instance of of this configuration without a {@link CLProgram},
+     * program binaries or sources associated with it.
+     */
+    public CLBuildConfiguration asBuildConfiguration();
+
+
+    // overwrite with CLProgramConfiguration as return type
+    @Override public CLProgramConfiguration forDevice(CLDevice device);
+    @Override public CLProgramConfiguration forDevices(CLDevice... devices);
+    @Override public CLProgramConfiguration withDefine(String name);
+    @Override public CLProgramConfiguration withDefine(String name, Object value);
+    @Override public CLProgramConfiguration withDefines(String... names);
+    @Override public CLProgramConfiguration withDefines(Map<String, ? extends Object> defines);
+    @Override public CLProgramConfiguration withOption(String option);
+    @Override public CLProgramConfiguration withOptions(String... options);
+    @Override public CLProgramConfiguration reset();
+    @Override public CLProgramConfiguration resetOptions();
+    @Override public CLProgramConfiguration resetDefines();
+    @Override public CLProgramConfiguration resetDevices();
+    @Override public CLProgramConfiguration clone();
+
+}
diff --git a/src/com/jogamp/opencl/util/CLUtil.java b/src/com/jogamp/opencl/util/CLUtil.java
new file mode 100644
index 00000000..6649698c
--- /dev/null
+++ b/src/com/jogamp/opencl/util/CLUtil.java
@@ -0,0 +1,119 @@
+package com.jogamp.opencl.util;
+
+import com.jogamp.opencl.CL;
+import com.jogamp.opencl.CLDevice;
+import com.jogamp.opencl.CLPlatform;
+import java.nio.ByteBuffer;
+import java.util.Arrays;
+import java.util.LinkedHashMap;
+import java.util.Map;
+
+/**
+ *
+ * @author Michael Bien
+ */
+public class CLUtil {
+
+    public static String clString2JavaString(byte[] chars, int clLength) {
+        return clLength==0 ? "" : new String(chars, 0, clLength-1);
+    }
+
+    public static String clString2JavaString(ByteBuffer chars, int clLength) {
+        if (clLength==0) {
+            return "";
+        }else{
+            byte[] array = new byte[clLength-1]; // last char is always null
+            chars.get(array).rewind();
+            return new String(array, 0, clLength-1);
+        }
+    }
+
+    /**
+     * Returns true if clBoolean == CL.CL_TRUE.
+     */
+    public static boolean clBoolean(int clBoolean) {
+        return clBoolean == CL.CL_TRUE;
+    }
+
+    /**
+     * Returns b ? CL.CL_TRUE : CL.CL_FALSE
+     */
+    public static int clBoolean(boolean b) {
+        return b ? CL.CL_TRUE : CL.CL_FALSE;
+    }
+
+    public static Map<String, String> obtainPlatformProperties(CLPlatform platform) {
+
+        Map<String, String> map = new LinkedHashMap<String, String>();
+        map.put("CL_PLATFORM_NAME",     platform.getName());
+        map.put("CL_PLATFORM_PROFILE",  platform.getProfile());
+        map.put("CL_PLATFORM_VERSION",  platform.getVersion());
+        map.put("CL_PLATFORM_VENDOR",   platform.getVendor());
+        map.put("CL_PLATFORM_EXTENSIONS",   platform.getExtensions().toString());
+//        map.put("fastest device (estimated)", platform.getMaxFlopsDevice().toString());
+
+        return map;
+    }
+
+    public static Map<String, String> obtainDeviceProperties(CLDevice dev) {
+        
+        Map<String, String> map = new LinkedHashMap<String, String>();
+        map.put("CL_DEVICE_NAME",       dev.getName());
+        map.put("CL_DEVICE_PROFILE",    dev.getProfile());
+        map.put("CL_DEVICE_VENDOR",     dev.getVendor());
+        map.put("CL_DEVICE_VENDOR_ID",  dev.getVendorID()+"");
+        map.put("CL_DEVICE_VERSION",    dev.getVersion());
+        map.put("CL_DRIVER_VERSION",    dev.getDriverVersion());
+        map.put("CL_DEVICE_TYPE",       dev.getType().toString());
+
+        map.put("CL_DEVICE_GLOBAL_MEM_SIZE",    dev.getGlobalMemSize()/(1024*1024)+" MB");
+        map.put("CL_DEVICE_MAX_MEM_ALLOC_SIZE", dev.getMaxMemAllocSize()/(1024*1024)+" MB");
+        map.put("CL_DEVICE_MAX_PARAMETER_SIZE", dev.getMaxParameterSize()+" Byte");
+        map.put("CL_DEVICE_LOCAL_MEM_SIZE",     dev.getLocalMemSize()/1024+" KB");
+        map.put("CL_DEVICE_LOCAL_MEM_TYPE",     dev.getLocalMemType()+"");
+        map.put("CL_DEVICE_GLOBAL_MEM_CACHE_SIZE", dev.getGlobalMemCacheSize()+"");
+        map.put("CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE", dev.getGlobalMemCachelineSize()+"");
+        map.put("CL_DEVICE_GLOBAL_MEM_CACHE_TYPE", dev.getGlobalMemCacheType()+"");
+        map.put("CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE", dev.getMaxConstantBufferSize()+"");
+        map.put("CL_DEVICE_MAX_CONSTANT_ARGS",        dev.getMaxConstantArgs()+"");
+        map.put("CL_DEVICE_ERROR_CORRECTION_SUPPORT", dev.isErrorCorrectionSupported()+"");
+
+        map.put("CL_DEVICE_MAX_CLOCK_FREQUENCY",        dev.getMaxClockFrequency()+" MHz");
+        map.put("CL_DEVICE_PROFILING_TIMER_RESOLUTION", dev.getProfilingTimerResolution()+" ns");
+        map.put("CL_DEVICE_QUEUE_PROPERTIES",           dev.getQueueProperties()+"");
+        map.put("CL_DEVICE_MAX_WORK_GROUP_SIZE",    dev.getMaxWorkGroupSize()+"");
+        map.put("CL_DEVICE_MAX_COMPUTE_UNITS",      dev.getMaxComputeUnits()+"");
+        map.put("CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS", dev.getMaxWorkItemDimensions()+"");
+        map.put("CL_DEVICE_MAX_WORK_ITEM_SIZES",    Arrays.toString(dev.getMaxWorkItemSizes()));
+        map.put("CL_DEVICE_COMPILER_AVAILABLE",     dev.isCompilerAvailable()+"");
+
+        map.put("CL_DEVICE_IMAGE_SUPPORT",          dev.isImageSupportAvailable()+"");
+        map.put("CL_DEVICE_MAX_READ_IMAGE_ARGS",    dev.getMaxReadImageArgs()+"");
+        map.put("CL_DEVICE_MAX_WRITE_IMAGE_ARGS",   dev.getMaxWriteImageArgs()+"");
+        map.put("CL_DEVICE_IMAGE2D_MAX_DIMENSIONS", Arrays.asList(dev.getMaxImage2dWidth(), dev.getMaxImage2dHeight()).toString());
+        map.put("CL_DEVICE_IMAGE3D_MAX_DIMENSIONS", Arrays.asList(dev.getMaxImage2dWidth(), dev.getMaxImage2dHeight(), dev.getMaxImage3dDepth()).toString());
+        map.put("CL_DEVICE_MAX_SAMPLERS",           dev.getMaxSamplers()+"");
+        map.put("CL_DEVICE_EXECUTION_CAPABILITIES", dev.getExecutionCapabilities()+"");
+
+        map.put("CL_DEVICE_ADDRESS_BITS",   dev.getAddressBits()+"");
+        map.put("cl_khr_fp16",              dev.isHalfFPAvailable()+"");
+        map.put("cl_khr_fp64",              dev.isDoubleFPAvailable()+"");
+        map.put("CL_DEVICE_ENDIAN_LITTLE",  dev.isLittleEndian()+"");
+        map.put("CL_DEVICE_HALF_FP_CONFIG", dev.getHalfFPConfig()+"");
+        map.put("CL_DEVICE_SINGLE_FP_CONFIG", dev.getSingleFPConfig()+"");
+        map.put("CL_DEVICE_DOUBLE_FP_CONFIG", dev.getDoubleFPConfig()+"");
+        map.put("CL_DEVICE_EXTENSIONS",     dev.getExtensions()+"");
+
+        map.put("CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT",   dev.getPreferredShortVectorWidth()+"");
+        map.put("CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR",    dev.getPreferredCharVectorWidth()+"");
+        map.put("CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT",     dev.getPreferredIntVectorWidth()+"");
+        map.put("CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG",    dev.getPreferredLongVectorWidth()+"");
+        map.put("CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT",   dev.getPreferredFloatVectorWidth()+"");
+        map.put("CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE",  dev.getPreferredDoubleVectorWidth()+"");
+
+        //TODO device extensions -> properties
+
+        return map;
+    }
+
+}
diff --git a/src/com/jogamp/opencl/util/MultiQueueBarrier.java b/src/com/jogamp/opencl/util/MultiQueueBarrier.java
new file mode 100644
index 00000000..59398b5e
--- /dev/null
+++ b/src/com/jogamp/opencl/util/MultiQueueBarrier.java
@@ -0,0 +1,141 @@
+package com.jogamp.opencl.util;
+
+import com.jogamp.opencl.CLCommandQueue;
+import com.jogamp.opencl.CLEventList;
+import java.util.Collections;
+import java.util.HashSet;
+import java.util.Set;
+import java.util.concurrent.CountDownLatch;
+import java.util.concurrent.TimeUnit;
+
+/**
+ * An utility for synchronizing multiple concurrent {@link CLCommandQueue}s.
+ * This Barrier can be reused after it has been broken.
+ * @author Michael Bien
+ */
+public class MultiQueueBarrier {
+
+    private CountDownLatch latch;
+    private final Set<CLCommandQueue> queues;
+    private final int count;
+
+    /**
+     * Creates a new MultiQueueBarrier with the given queueCount.
+     * It is recommented to use {@link #MultiQueueBarrier(CLCommandQueue... allowedQueues)} if possible
+     * which restricts the set of allowed queues for the barrier.
+     */
+    public MultiQueueBarrier(int queueCount) {
+        if(queueCount == 0) {
+            throw new IllegalArgumentException("queueCount was 0");
+        }
+        this.latch = new CountDownLatch(queueCount);
+        this.queues = null;
+        this.count = queueCount;
+    }
+
+    /**
+     * Creates a new MultiQueueBarrier for the given queues.
+     */
+    public MultiQueueBarrier(CLCommandQueue... allowedQueues) {
+        if(allowedQueues.length == 0) {
+            throw new IllegalArgumentException("allowedQueues was empty");
+        }
+        this.latch = new CountDownLatch(allowedQueues.length);
+        this.count = allowedQueues.length;
+
+        HashSet<CLCommandQueue> set = new HashSet<CLCommandQueue>(allowedQueues.length);
+        for (CLCommandQueue queue : allowedQueues) {
+            set.add(queue);
+        }
+        this.queues = Collections.unmodifiableSet(set);
+    }
+
+    /**
+     * Blocks the current Thread until all commands on the {@link CLCommandQueue} finished excecution.
+     * This method may be invoked concurrently without synchronization on the MultiQueueBarrier object
+     * as long each Thread passes a distinct CLCommandQueue as parameter to this method.
+     */
+    public MultiQueueBarrier waitFor(CLCommandQueue queue) {
+        checkQueue(queue);
+
+        queue.putBarrier();
+        synchronized(this) {
+            latch.countDown();
+        }
+        return this;
+    }
+
+    /**
+     * Blocks the current Thread until the given events on the {@link CLCommandQueue} occurred.
+     * This method may be invoked concurrently without synchronization on the MultiQueueBarrier object
+     * as long each Thread passes a distinct CLCommandQueue as parameter to this method.
+     */
+    public MultiQueueBarrier waitFor(CLCommandQueue queue, CLEventList events) {
+        checkQueue(queue);
+
+        queue.putWaitForEvents(events, true);
+        synchronized(this) {
+            latch.countDown();
+        }
+        return this;
+    }
+
+    /**
+     * Blocks until all Threads which called {@link #waitFor}
+     * continue execution.
+     * This method blocks only once, all subsequent calls are ignored.
+     */
+    public MultiQueueBarrier await() throws InterruptedException {
+        latch.await();
+        rebuildBarrierIfBroken();
+        return this;
+    }
+    
+    /**
+     * @see #await()
+     * @param timeout the maximum time to wait
+     * @param unit the time unit of the {@code timeout} argument
+     */
+    public MultiQueueBarrier await(long timeout, TimeUnit unit) throws InterruptedException {
+        latch.await(timeout, unit);
+        rebuildBarrierIfBroken();
+        return this;
+    }
+    
+    /**
+     * Resets this barrier and unblocks all waiting threads.
+     */
+    public void resetBarrier() {
+        synchronized(this) {
+            while(latch.getCount() > 0) {
+                latch.countDown();
+            }
+            // thats OK. Another Thread can not rebuild the barrier since we have the lock.
+            // we have to rebuid it here in case there was no thread waiting.
+            latch = new CountDownLatch(count);
+        }
+    }
+
+    private void rebuildBarrierIfBroken() {
+        synchronized (this) {
+            if (latch.getCount() == 0) {
+                latch = new CountDownLatch(count);
+            }
+        }
+    }
+
+    /**
+     * Returns the current number of events which must occure before this barrier unblocks the waiting threads.
+     * This method is typically used for debugging and testing purposes.
+     */
+    public long getCount() {
+        return latch.getCount();
+    }
+
+    private void checkQueue(CLCommandQueue queue) throws IllegalArgumentException {
+        if (queues != null && !queues.contains(queue)) {
+            throw new IllegalArgumentException(queue + " is not in the allowedQueues Set: " + queues);
+        }
+    }
+
+}