initial import of CLBuffer and CLKernel.

added hashCode(), equals() and toString() methods.
updated JUnit test to test new classes.

7 files changed, 402 insertions, 45 deletions

diff --git a/src/com/mbien/opencl/CLBuffer.java b/src/com/mbien/opencl/CLBuffer.java
new file mode 100644
index 00000000..0f6e34a4
--- /dev/null
+++ b/src/com/mbien/opencl/CLBuffer.java
@@ -0,0 +1,70 @@
+package com.mbien.opencl;
+
+import java.nio.ByteBuffer;
+import static com.mbien.opencl.CLException.*;
+
+/**
+ *
+ * @author Michael Bien
+ */
+public class CLBuffer {
+
+    public final ByteBuffer buffer;
+    public final long bufferID;
+    
+    private final CLContext context;
+    private final CL cl;
+
+    CLBuffer(CLContext context, int flags, ByteBuffer directBuffer) {
+        
+        if(!directBuffer.isDirect())
+            throw new IllegalArgumentException("buffer is not a direct buffer");
+
+        this.buffer = directBuffer;
+        this.context = context;
+        this.cl = context.cl;
+
+        int[] intArray = new int[1];
+
+        this.bufferID = cl.clCreateBuffer(context.contextID, flags, directBuffer.capacity(), null, intArray, 0);
+
+        checkForError(intArray[0], "can not create cl buffer");
+        
+    }
+
+    public CLBuffer release() {
+        cl.clReleaseMemObject(bufferID);
+        context.bufferReleased(this);
+        return this;
+    }
+
+    @Override
+    public boolean equals(Object obj) {
+        if (obj == null) {
+            return false;
+        }
+        if (getClass() != obj.getClass()) {
+            return false;
+        }
+        final CLBuffer other = (CLBuffer) obj;
+        if (this.buffer != other.buffer && (this.buffer == null || !this.buffer.equals(other.buffer))) {
+            return false;
+        }
+        if (this.context.contextID != other.context.contextID) {
+            return false;
+        }
+        return true;
+    }
+
+    @Override
+    public int hashCode() {
+        int hash = 3;
+        hash = 29 * hash + (this.buffer != null ? this.buffer.hashCode() : 0);
+        hash = 29 * hash + (int) (this.context.contextID ^ (this.context.contextID >>> 32));
+        return hash;
+    }
+
+
+
+
+}
diff --git a/src/com/mbien/opencl/CLContext.java b/src/com/mbien/opencl/CLContext.java
index 601a2b79..9cb649d7 100644
--- a/src/com/mbien/opencl/CLContext.java
+++ b/src/com/mbien/opencl/CLContext.java
@@ -22,6 +22,7 @@ public final class CLContext {
     private CLDevice[] devices;
 
     private final List<CLProgram> programs;
+    private final List<CLBuffer> buffers;
 
     static{
         System.loadLibrary("gluegen-rt");
@@ -32,6 +33,7 @@ public final class CLContext {
     private CLContext(long contextID) {
         this.contextID = contextID;
         this.programs = new ArrayList<CLProgram>();
+        this.buffers = new ArrayList<CLBuffer>();
     }
 
     /**
@@ -70,14 +72,32 @@ public final class CLContext {
         return program;
     }
 
+    public CLBuffer createBuffer(int flags, ByteBuffer directBuffer) {
+        CLBuffer buffer = new CLBuffer(this, flags, directBuffer);
+        buffers.add(buffer);
+        return buffer;
+    }
+
     void programReleased(CLProgram program) {
         programs.remove(program);
     }
 
+    void bufferReleased(CLBuffer buffer) {
+        buffers.remove(buffer);
+    }
+
     /**
      * Releases the context and all resources.
      */
     public CLContext release() {
+
+        //release all resources
+        while(!programs.isEmpty())
+            programs.get(0).release();
+
+        while(!buffers.isEmpty())
+            buffers.get(0).release();
+
         int ret = cl.clReleaseContext(contextID);
         checkForError(ret, "error releasing context");
         return this;
@@ -86,10 +106,17 @@ public final class CLContext {
     /**
      * Returns a read only view of all programs associated with this context.
      */
-    public List<CLProgram> getPrograms() {
+    public List<CLProgram> getCLPrograms() {
         return Collections.unmodifiableList(programs);
     }
 
+    /**
+     * Returns a read only view of all buffers associated with this context.
+     */
+    public List<CLBuffer> getCLBuffers() {
+        return Collections.unmodifiableList(buffers);
+    }
+
 
     /**
      * Gets the device with maximal FLOPS from this context.
@@ -157,7 +184,7 @@ public final class CLContext {
         return devices;
     }
 
-    CLDevice getCLDevices(long dID) {
+    CLDevice getCLDevice(long dID) {
         CLDevice[] deviceArray = getCLDevices();
         for (int i = 0; i < deviceArray.length; i++) {
             if(dID == deviceArray[i].deviceID)
@@ -198,4 +225,34 @@ public final class CLContext {
     }
 
 
+    @Override
+    public String toString() {
+        return "CLContext [id: " + contextID
+                      + " #devices: " + getCLDevices().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.contextID != other.contextID) {
+            return false;
+        }
+        return true;
+    }
+
+    @Override
+    public int hashCode() {
+        int hash = 7;
+        hash = 23 * hash + (int) (this.contextID ^ (this.contextID >>> 32));
+        return hash;
+    }
+
+
 }
diff --git a/src/com/mbien/opencl/CLDevice.java b/src/com/mbien/opencl/CLDevice.java
index d9f643ce..dcc9ee97 100644
--- a/src/com/mbien/opencl/CLDevice.java
+++ b/src/com/mbien/opencl/CLDevice.java
@@ -191,9 +191,17 @@ public final class CLDevice {
 
     @Override
     public boolean equals(Object obj) {
-        if(obj != null && obj instanceof CLDevice)
-            return ((CLDevice)obj).deviceID == deviceID;
-        return false;
+        if (obj == null) {
+            return false;
+        }
+        if (getClass() != obj.getClass()) {
+            return false;
+        }
+        final CLDevice other = (CLDevice) obj;
+        if (this.deviceID != other.deviceID) {
+            return false;
+        }
+        return true;
     }
 
     @Override
diff --git a/src/com/mbien/opencl/CLKernel.java b/src/com/mbien/opencl/CLKernel.java
new file mode 100644
index 00000000..be5e03b6
--- /dev/null
+++ b/src/com/mbien/opencl/CLKernel.java
@@ -0,0 +1,93 @@
+package com.mbien.opencl;
+
+import com.sun.gluegen.runtime.BufferFactory;
+import java.nio.ByteBuffer;
+import java.nio.ByteOrder;
+import static com.mbien.opencl.CLException.*;
+
+/**
+ *
+ * @author Michael Bien
+ */
+public class CLKernel {
+
+    public final long kernelID;
+    public final String name;
+
+    private final CLProgram program;
+    private final CL cl;
+
+    CLKernel(CLProgram program, long id) {
+        this.kernelID = id;
+        this.program = program;
+        this.cl = program.context.cl;
+
+        long[] longArray = new long[1];
+
+        int ret = cl.clGetKernelInfo(kernelID, CL.CL_KERNEL_FUNCTION_NAME, 0, null, longArray, 0);
+        checkForError(ret, "error while asking for kernel function name");
+
+        ByteBuffer bb = ByteBuffer.allocate((int)longArray[0]).order(ByteOrder.nativeOrder());
+
+        ret = cl.clGetKernelInfo(kernelID, CL.CL_KERNEL_FUNCTION_NAME, bb.capacity(), bb, null, 0);
+        checkForError(ret, "error while asking for kernel function name");
+
+        this.name = new String(bb.array(), 0, (int)longArray[0]).trim();
+
+    }
+
+    public CLKernel setArg(int argumentIndex, int argumentSize, CLBuffer value) {
+        int ret = cl.clSetKernelArg(kernelID, argumentIndex, argumentSize, wrapLong(value.bufferID));
+        checkForError(ret, "error on clSetKernelArg");
+        return this;
+    }
+
+    public CLKernel setArg(int argumentIndex, int argumentSize, long value) {
+        int ret = cl.clSetKernelArg(kernelID, argumentIndex, argumentSize, wrapLong(value));
+        checkForError(ret, "error on clSetKernelArg");
+        return this;
+    }
+
+    private final ByteBuffer wrapLong(long value) {
+        return (ByteBuffer) BufferFactory.newDirectByteBuffer(8).putLong(value).rewind();
+    }
+
+    public CLKernel release() {
+        cl.clReleaseKernel(kernelID);
+        program.kernelReleased(this);
+        return this;
+    }
+
+    @Override
+    public String toString() {
+        return "CLKernel [id: " + kernelID
+                      + " 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.kernelID != other.kernelID) {
+            return false;
+        }
+        if (!this.program.equals(other.program)) {
+            return false;
+        }
+        return true;
+    }
+
+    @Override
+    public int hashCode() {
+        int hash = 7;
+        hash = 43 * hash + (int) (this.kernelID ^ (this.kernelID >>> 32));
+        hash = 43 * hash + (this.program != null ? this.program.hashCode() : 0);
+        return hash;
+    }
+
+}
diff --git a/src/com/mbien/opencl/CLPlatform.java b/src/com/mbien/opencl/CLPlatform.java
index a717564a..d9f8dd25 100644
--- a/src/com/mbien/opencl/CLPlatform.java
+++ b/src/com/mbien/opencl/CLPlatform.java
@@ -11,7 +11,7 @@ public final class CLPlatform {
     /**
      * OpenCL platform id for this platform.
      */
-    public  final long platformID;
+    public final long platformID;
 
     private final CL cl;
 
@@ -94,7 +94,27 @@ public final class CLPlatform {
                          +" 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.platformID != other.platformID) {
+            return false;
+        }
+        return true;
+    }
 
+    @Override
+    public int hashCode() {
+        int hash = 7;
+        hash = 71 * hash + (int) (this.platformID ^ (this.platformID >>> 32));
+        return hash;
+    }
 
 
 }
diff --git a/src/com/mbien/opencl/CLProgram.java b/src/com/mbien/opencl/CLProgram.java
index e93e7246..73fe8cac 100644
--- a/src/com/mbien/opencl/CLProgram.java
+++ b/src/com/mbien/opencl/CLProgram.java
@@ -2,6 +2,10 @@ package com.mbien.opencl;
 
 import java.nio.ByteBuffer;
 import java.nio.ByteOrder;
+import java.util.Collections;
+import java.util.HashMap;
+import java.util.Map;
+import java.util.Set;
 import static com.mbien.opencl.CLException.*;
 
 /**
@@ -10,9 +14,12 @@ import static com.mbien.opencl.CLException.*;
  */
 public class CLProgram {
     
-    private final CLContext context;
-    private final long programID;
+    public final CLContext context;
+    public final long programID;
+    
     private final CL cl;
+
+    private final Map<String, CLKernel> kernels;
     
     public enum Status {
         
@@ -53,6 +60,8 @@ public class CLProgram {
         this.cl = context.cl;
         this.context = context;
 
+        this.kernels = new HashMap<String, CLKernel>();
+
         int[] intArray = new int[1];
         // Create the program
         programID = cl.clCreateProgramWithSource(contextID, 1, new String[] {src}, new long[]{src.length()}, 0, intArray, 0);
@@ -92,10 +101,46 @@ public class CLProgram {
     }
 
     /**
+     * Returns all kernels of this program in a unmodifiable view of a map with the kernel function names as keys.
+     */
+    public Map<String, CLKernel> getCLKernels() {
+
+        if(kernels.isEmpty()) {
+            
+            int[] intArray = new int[1];
+            int ret = cl.clCreateKernelsInProgram(programID, 0, null, 0, intArray, 0);
+            checkForError(ret, "can not create kernels for program");
+
+            long[] kernelIDs = new long[intArray[0]];
+            ret = cl.clCreateKernelsInProgram(programID, kernelIDs.length, kernelIDs, 0, null, 0);
+            checkForError(ret, "can not create kernels for program");
+
+            for (int i = 0; i < intArray[0]; i++) {
+                CLKernel kernel = new CLKernel(this, kernelIDs[i]);
+                kernels.put(kernel.name, kernel);
+            }
+        }
+
+        return Collections.unmodifiableMap(kernels);
+    }
+
+    void kernelReleased(CLKernel kernel) {
+        this.kernels.remove(kernel.name);
+    }
+
+    /**
      * Releases this program.
      * @return this
      */
     public CLProgram release() {
+
+        if(!kernels.isEmpty()) {
+            String[] names = kernels.keySet().toArray(new String[kernels.size()]);
+            for (String name : names) {
+                kernels.get(name).release();
+            }
+        }
+
         int ret = cl.clReleaseProgram(programID);
         checkForError(ret, "can not release program");
         context.programReleased(this);
@@ -119,7 +164,7 @@ public class CLProgram {
         int count = bb.capacity() / 8; // TODO sizeof cl_device
         CLDevice[] devices = new CLDevice[count];
         for (int i = 0; i < count; i++) {
-            devices[i] = context.getCLDevices(bb.getLong());
+            devices[i] = context.getCLDevice(bb.getLong());
         }
 
         return devices;
@@ -194,4 +239,30 @@ public class CLProgram {
         return bb.getInt();
     }
 
+    @Override
+    public boolean equals(Object obj) {
+        if (obj == null) {
+            return false;
+        }
+        if (getClass() != obj.getClass()) {
+            return false;
+        }
+        final CLProgram other = (CLProgram) obj;
+        if (this.programID != other.programID) {
+            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.programID ^ (this.programID >>> 32));
+        return hash;
+    }
+
 }
diff --git a/test/com/mbien/opencl/JOCLTest.java b/test/com/mbien/opencl/JOCLTest.java
index 05b1d28d..5a5b7a27 100644
--- a/test/com/mbien/opencl/JOCLTest.java
+++ b/test/com/mbien/opencl/JOCLTest.java
@@ -4,6 +4,7 @@ import com.sun.gluegen.runtime.BufferFactory;
 import java.nio.ByteBuffer;
 import java.nio.ByteOrder;
 import java.util.Arrays;
+import java.util.Map;
 import java.util.Random;
 import org.junit.BeforeClass;
 import org.junit.Test;
@@ -27,7 +28,15 @@ public class JOCLTest {
             + "    }                                                                                                            \n"
             + "    // add the vector elements                                                                                   \n"
             + "    c[iGID] = a[iGID] + b[iGID];                                                                                 \n"
-            + "    //c[iGID] = iGID;                                                                                            \n"
+            + "}                                                                                                                \n"
+            + "__kernel void Test(__global const int* a, __global const int* b, __global int* c, int iNumElements) {       \n"
+            + "    // get index into global data array                                                                          \n"
+            + "    int iGID = get_global_id(0);                                                                                 \n"
+            + "    // bound check (equivalent to the limit on a 'for' loop for standard/serial C code                           \n"
+            + "    if (iGID >= iNumElements)  {                                                                                 \n"
+            + "        return;                                                                                                  \n"
+            + "    }                                                                                                            \n"
+            + "    c[iGID] = iGID;                                                                                            \n"
             + "}                                                                                                                \n";
 
     public JOCLTest() {
@@ -156,15 +165,19 @@ public class JOCLTest {
         int localWorkSize = 256;      // set and log Global and Local work size dimensions
         int globalWorkSize = roundUp(localWorkSize, elementCount);  // rounded up to the nearest multiple of the LocalWorkSize
 
-        out.println(globalWorkSize);
+        out.println("allocateing buffers of size: "+globalWorkSize);
+
+        ByteBuffer srcA = BufferFactory.newDirectByteBuffer(globalWorkSize*BufferFactory.SIZEOF_INT);
+        ByteBuffer srcB = BufferFactory.newDirectByteBuffer(globalWorkSize*BufferFactory.SIZEOF_INT);
+        ByteBuffer dest = BufferFactory.newDirectByteBuffer(globalWorkSize*BufferFactory.SIZEOF_INT);
 
         // TODO sizeof int ...
         // Allocate the OpenCL buffer memory objects for source and result on the device GMEM
-        long devSrcA = cl.clCreateBuffer(context, CL.CL_MEM_READ_ONLY, BufferFactory.SIZEOF_INT * globalWorkSize, null, intArray, 0);
+        long devSrcA = cl.clCreateBuffer(context, CL.CL_MEM_READ_ONLY, srcA.capacity(), null, intArray, 0);
         checkError("on clCreateBuffer", intArray[0]);
-        long devSrcB = cl.clCreateBuffer(context, CL.CL_MEM_READ_ONLY, BufferFactory.SIZEOF_INT * globalWorkSize, null, intArray, 0);
+        long devSrcB = cl.clCreateBuffer(context, CL.CL_MEM_READ_ONLY, srcB.capacity(), null, intArray, 0);
         checkError("on clCreateBuffer", intArray[0]);
-        long devDst  = cl.clCreateBuffer(context, CL.CL_MEM_WRITE_ONLY, BufferFactory.SIZEOF_INT * globalWorkSize, null, intArray, 0);
+        long devDst  = cl.clCreateBuffer(context, CL.CL_MEM_WRITE_ONLY, dest.capacity(), null, intArray, 0);
         checkError("on clCreateBuffer", intArray[0]);
 
 
@@ -198,7 +211,7 @@ public class JOCLTest {
         ret = cl.clGetProgramBuildInfo(program, firstDeviceID, CL.CL_PROGRAM_BUILD_STATUS, bb.capacity(), bb, null, 0);
         checkError("on clGetProgramBuildInfo1", ret);
 
-        out.println("program build status: " + getBuildStatus(bb.getInt(0)));
+        out.println("program build status: " + CLProgram.Status.valueOf(bb.getInt(0)));
         assertEquals("build status", CL.CL_BUILD_SUCCESS, bb.getInt(0));
 
         // Read build log
@@ -216,11 +229,6 @@ public class JOCLTest {
         long kernel = cl.clCreateKernel(program, "VectorAdd", intArray, 0);
         checkError("on clCreateKernel", intArray[0]);
 
-
-        ByteBuffer srcA = BufferFactory.newDirectByteBuffer(globalWorkSize*BufferFactory.SIZEOF_INT);
-        ByteBuffer srcB = BufferFactory.newDirectByteBuffer(globalWorkSize*BufferFactory.SIZEOF_INT);
-        ByteBuffer dst  = BufferFactory.newDirectByteBuffer(globalWorkSize*BufferFactory.SIZEOF_INT);
-
 //        srcA.limit(elementCount*BufferFactory.SIZEOF_FLOAT);
 //        srcB.limit(elementCount*BufferFactory.SIZEOF_FLOAT);
 
@@ -233,7 +241,7 @@ public class JOCLTest {
         ret = cl.clSetKernelArg(kernel, 2, BufferFactory.SIZEOF_LONG, wrap(devDst));   checkError("on clSetKernelArg2", ret);
         ret = cl.clSetKernelArg(kernel, 3, BufferFactory.SIZEOF_INT,  wrap(elementCount));  checkError("on clSetKernelArg3", ret);
 
-        out.println("used device memory: "+ (srcA.capacity()+srcB.capacity()+dst.capacity())/1000000 +"MB");
+        out.println("used device memory: "+ (srcA.capacity()+srcB.capacity()+dest.capacity())/1000000 +"MB");
 
         // Asynchronous write of data to GPU device
         ret = cl.clEnqueueWriteBuffer(commandQueue, devSrcA, CL.CL_FALSE, 0, srcA.capacity(), srcA, 0, null, 0, null, 0);
@@ -250,13 +258,13 @@ public class JOCLTest {
         checkError("on clEnqueueNDRangeKernel", ret);
 
         // Synchronous/blocking read of results
-        ret = cl.clEnqueueReadBuffer(commandQueue, devDst, CL.CL_TRUE, 0, BufferFactory.SIZEOF_INT * globalWorkSize, dst, 0, null, 0, null, 0);
+        ret = cl.clEnqueueReadBuffer(commandQueue, devDst, CL.CL_TRUE, 0, dest.capacity(), dest, 0, null, 0, null, 0);
         checkError("on clEnqueueReadBuffer", ret);
 
         out.println("a+b=c result snapshot: ");
         for(int i = 0; i < 10; i++)
-            out.print(dst.getInt()+", ");
-        out.println();
+            out.print(dest.getInt()+", ");
+        out.println("...; "+dest.remaining()/BufferFactory.SIZEOF_INT + " more");
 
 
         // cleanup
@@ -284,8 +292,9 @@ public class JOCLTest {
 
     }
 
-    @Test
+//    @Test
     public void loadTest() {
+        //for memory leak detection; e.g watch out for "out of host memory" errors
         out.println(" - - - loadTest - - - ");
         for(int i = 0; i < 100; i++) {
             out.println("###iteration "+i);
@@ -307,24 +316,6 @@ public class JOCLTest {
         return (ByteBuffer) BufferFactory.newDirectByteBuffer(8).putLong(value).rewind();
     }
 
-    private String getBuildStatus(int status) {
-        switch(status) {
-            case CL.CL_BUILD_SUCCESS:
-                return "CL_BUILD_SUCCESS";
-            case CL.CL_BUILD_NONE:
-                return "CL_BUILD_NONE";
-            case CL.CL_BUILD_IN_PROGRESS:
-                return "CL_BUILD_IN_PROGRESS";
-            case CL.CL_BUILD_ERROR:
-                return "CL_BUILD_ERROR";
-// can't find this flag in spec...
-//            case CL.CL_BUILD_PROGRAM_FAILURE:
-//                return "CL_BUILD_PROGRAM_FAILURE";
-            default:
-                return "unknown status: " + status;
-        }
-    }
-
     @Test
     public void highLevelTest1() {
         
@@ -389,9 +380,56 @@ public class JOCLTest {
 
         out.println("source:\n"+program.getSource());
 
-        assertTrue(1 == context.getPrograms().size());
+        int elementCount = 11444777;	// Length of float arrays to process (odd # for illustration)
+        int localWorkSize = 256;      // set and log Global and Local work size dimensions
+        int globalWorkSize = roundUp(localWorkSize, elementCount);  // rounded up to the nearest multiple of the LocalWorkSize
+
+        out.println("allocateing buffers of size: "+globalWorkSize);
+
+        ByteBuffer srcA = BufferFactory.newDirectByteBuffer(globalWorkSize*BufferFactory.SIZEOF_INT);
+        ByteBuffer srcB = BufferFactory.newDirectByteBuffer(globalWorkSize*BufferFactory.SIZEOF_INT);
+        ByteBuffer dest = BufferFactory.newDirectByteBuffer(globalWorkSize*BufferFactory.SIZEOF_INT);
+
+
+        CLBuffer clBufferA = context.createBuffer(CL.CL_MEM_READ_ONLY, srcA);
+        CLBuffer clBufferB = context.createBuffer(CL.CL_MEM_READ_ONLY, srcB);
+        CLBuffer clBufferC = context.createBuffer(CL.CL_MEM_WRITE_ONLY, dest);
+        
+        Map<String, CLKernel> kernels = program.getCLKernels();
+        for (CLKernel kernel : kernels.values()) {
+            out.println("kernel: "+kernel.toString());
+        }
+
+        assertNotNull(kernels.get("VectorAdd"));
+        assertNotNull(kernels.get("Test"));
+        
+        CLKernel vectorAddKernel = kernels.get("VectorAdd");
+
+        vectorAddKernel
+            .setArg(0, BufferFactory.SIZEOF_LONG, clBufferA)
+            .setArg(1, BufferFactory.SIZEOF_LONG, clBufferB)
+            .setArg(2, BufferFactory.SIZEOF_LONG, clBufferC)
+            .setArg(3, BufferFactory.SIZEOF_LONG, elementCount);
+
+        //TODO CLComandQueue...
+
+
+        assertTrue(3 == context.getCLBuffers().size());
+        clBufferA.release();
+        assertTrue(2 == context.getCLBuffers().size());
+
+        assertTrue(2 == context.getCLBuffers().size());
+        clBufferB.release();
+        assertTrue(1 == context.getCLBuffers().size());
+
+        assertTrue(1 == context.getCLBuffers().size());
+        clBufferC.release();
+        assertTrue(0 == context.getCLBuffers().size());
+
+
+        assertTrue(1 == context.getCLPrograms().size());
         program.release();
-        assertTrue(0 == context.getPrograms().size());
+        assertTrue(0 == context.getCLPrograms().size());
 
 //        CLDevice device = ctx.getMaxFlopsDevice();
 //        out.println("max FLOPS device: " + device);