7 files changed, 37 insertions, 6 deletions
diff --git a/src/com/mbien/opencl/CLBuffer.java b/src/com/mbien/opencl/CLBuffer.java
index 74df71e6..f9f3a239 100644
--- a/src/com/mbien/opencl/CLBuffer.java
+++ b/src/com/mbien/opencl/CLBuffer.java
@@ -33,8 +33,9 @@ public class CLBuffer {
     }
 
     public void release() {
-        cl.clReleaseMemObject(ID);
+        int ret = cl.clReleaseMemObject(ID);
         context.bufferReleased(this);
+        checkForError(ret, "can not release mem object");
     }
 
     @Override
diff --git a/src/com/mbien/opencl/CLCommandQueue.java b/src/com/mbien/opencl/CLCommandQueue.java
index ff91131b..b0a6980f 100644
--- a/src/com/mbien/opencl/CLCommandQueue.java
+++ b/src/com/mbien/opencl/CLCommandQueue.java
@@ -10,11 +10,13 @@ public class CLCommandQueue {
 
     public final long ID;
     private final CLContext context;
+    private final CLDevice device;
     private final CL cl;
 
     CLCommandQueue(CLContext context, CLDevice device, long properties) {
         this.context = context;
         this.cl = context.cl;
+        this.device = device;
 
         int[] status = new int[1];
         this.ID = cl.clCreateCommandQueue(context.ID, device.ID, properties, status, 0);
@@ -55,7 +57,7 @@ public class CLCommandQueue {
 
        int ret = cl.clEnqueueNDRangeKernel(
                 ID, kernel.ID, 1,
-                null, 0,
+                globalWorkOffset, 0,
                 globalWorkSize, 0,
                 localWorkSize, 0,
                 0,
@@ -70,6 +72,7 @@ public class CLCommandQueue {
 
     public void release() {
         int ret = cl.clReleaseCommandQueue(ID);
+        context.commandQueueReleased(device, this);
         checkForError(ret, "can not release command queue");
     }
 
diff --git a/src/com/mbien/opencl/CLContext.java b/src/com/mbien/opencl/CLContext.java
index fcbb8041..dda8eb05 100644
--- a/src/com/mbien/opencl/CLContext.java
+++ b/src/com/mbien/opencl/CLContext.java
@@ -7,7 +7,9 @@ import java.nio.IntBuffer;
 
 import java.util.ArrayList;
 import java.util.Collections;
+import java.util.HashMap;
 import java.util.List;
+import java.util.Map;
 import static com.mbien.opencl.CLException.*;
 
 /**
@@ -23,6 +25,7 @@ public final class CLContext {
 
     private final List<CLProgram> programs;
     private final List<CLBuffer> buffers;
+    private final Map<CLDevice, List<CLCommandQueue>> queuesMap;
 
     static{
         System.loadLibrary("gluegen-rt");
@@ -34,6 +37,7 @@ public final class CLContext {
         this.ID = contextID;
         this.programs = new ArrayList<CLProgram>();
         this.buffers = new ArrayList<CLBuffer>();
+        this.queuesMap = new HashMap<CLDevice, List<CLCommandQueue>>();
     }
 
     /**
@@ -78,6 +82,20 @@ public final class CLContext {
         return buffer;
     }
 
+    CLCommandQueue createCommandQueue(CLDevice device, long properties) {
+
+        CLCommandQueue queue = new CLCommandQueue(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;
+    }
+
     void programReleased(CLProgram program) {
         programs.remove(program);
     }
@@ -86,6 +104,14 @@ public final class CLContext {
         buffers.remove(buffer);
     }
 
+    void commandQueueReleased(CLDevice device, CLCommandQueue queue) {
+        List<CLCommandQueue> list = queuesMap.get(device);
+        list.remove(queue);
+        // remove empty lists from map
+        if(list.isEmpty())
+            queuesMap.remove(device);
+    }
+
     /**
      * Releases the context and all resources.
      */
diff --git a/src/com/mbien/opencl/CLDevice.java b/src/com/mbien/opencl/CLDevice.java
index c69afef5..87616516 100644
--- a/src/com/mbien/opencl/CLDevice.java
+++ b/src/com/mbien/opencl/CLDevice.java
@@ -86,7 +86,7 @@ public final class CLDevice {
     public CLCommandQueue createCommandQueue(long properties) {
         if(context == null)
             throw new IllegalStateException("this device is not associated with a context");
-        return new CLCommandQueue(context, this, properties);
+        return context.createCommandQueue(this, properties);
     }
 
     /*keep this package private for now, may be null*/
diff --git a/src/com/mbien/opencl/CLKernel.java b/src/com/mbien/opencl/CLKernel.java
index 8f470719..1db3da38 100644
--- a/src/com/mbien/opencl/CLKernel.java
+++ b/src/com/mbien/opencl/CLKernel.java
@@ -53,8 +53,9 @@ public class CLKernel {
     }
 
     public void release() {
-        cl.clReleaseKernel(ID);
+        int ret = cl.clReleaseKernel(ID);
         program.kernelReleased(this);
+        checkForError(ret, "can not release kernel");
     }
 
     @Override
diff --git a/src/com/mbien/opencl/CLProgram.java b/src/com/mbien/opencl/CLProgram.java
index 0404e070..020de17c 100644
--- a/src/com/mbien/opencl/CLProgram.java
+++ b/src/com/mbien/opencl/CLProgram.java
@@ -140,8 +140,8 @@ public class CLProgram {
         }
 
         int ret = cl.clReleaseProgram(ID);
-        checkForError(ret, "can not release program");
         context.programReleased(this);
+        checkForError(ret, "can not release program");
         
     }
 
diff --git a/test/com/mbien/opencl/JOCLTest.java b/test/com/mbien/opencl/JOCLTest.java
index e0dfdee8..224a6768 100644
--- a/test/com/mbien/opencl/JOCLTest.java
+++ b/test/com/mbien/opencl/JOCLTest.java
@@ -419,7 +419,7 @@ public class JOCLTest {
         // Asynchronous write of data to GPU device, blocking read later
         queue.putWriteBuffer(clBufferA, false)
              .putWriteBuffer(clBufferB, false)
-             .putNDRangeKernel(vectorAddKernel, 1, new long[]{0}, new long[]{ globalWorkSize }, new long[]{ localWorkSize })
+             .putNDRangeKernel(vectorAddKernel, 1, null, new long[]{ globalWorkSize }, new long[]{ localWorkSize })
              .putReadBuffer(clBufferC, true).release();
 
         out.println("a+b=c result snapshot: ");