fixed a bug which used a wrong eventlist offset under certain conditions and added a regression test.

4 files changed, 128 insertions, 41 deletions

diff --git a/src/com/jogamp/opencl/CLCommandQueue.java b/src/com/jogamp/opencl/CLCommandQueue.java
index d24fb115..2c64b1f7 100644
--- a/src/com/jogamp/opencl/CLCommandQueue.java
+++ b/src/com/jogamp/opencl/CLCommandQueue.java
@@ -118,7 +118,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
         
@@ -162,7 +162,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
         
@@ -211,7 +211,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -281,7 +281,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -359,7 +359,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -440,7 +440,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -508,7 +508,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -570,7 +570,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -631,7 +631,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -693,7 +693,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -757,7 +757,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -826,7 +826,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -890,7 +890,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -957,7 +957,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -1021,7 +1021,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -1088,7 +1088,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -1146,7 +1146,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -1206,7 +1206,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -1272,7 +1272,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -1316,7 +1316,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -1348,11 +1348,11 @@ public class CLCommandQueue extends CLObject implements CLResource {
      * 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);
+
+        PointerBuffer ids = PointerBuffer.wrap(list.IDs.getBuffer()).position(index);
+        
+        int ret = blockingWait ? cl.clWaitForEvents(1, ids)
+                               : cl.clEnqueueWaitForEvents(ID, 1, ids);
         if(ret != CL_SUCCESS) {
             throw newException(ret, "can not "+ (blockingWait?"blocking": "") +" wait for event #" + index+ " in "+list);
         }
@@ -1363,9 +1363,8 @@ public class CLCommandQueue extends CLObject implements CLResource {
      * 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);
+        int ret = blockingWait ? cl.clWaitForEvents(list.size, list.IDsView)
+                               : cl.clEnqueueWaitForEvents(ID, list.size, list.IDsView);
         if(ret != CL_SUCCESS) {
             throw newException(ret, "can not "+ (blockingWait?"blocking": "") +" wait for events " + list);
         }
@@ -1410,7 +1409,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -1537,7 +1536,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -1588,7 +1587,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
@@ -1635,7 +1634,7 @@ public class CLCommandQueue extends CLObject implements CLResource {
         PointerBuffer conditionIDs = null;
         int conditions = 0;
         if(condition != null) {
-            conditionIDs = condition.IDs;
+            conditionIDs = condition.IDsView;
             conditions   = condition.size;
         }
 
diff --git a/src/com/jogamp/opencl/CLEventList.java b/src/com/jogamp/opencl/CLEventList.java
index 03a6f838..f2b98adf 100644
--- a/src/com/jogamp/opencl/CLEventList.java
+++ b/src/com/jogamp/opencl/CLEventList.java
@@ -40,17 +40,29 @@ public final class CLEventList implements CLResource, AutoCloseable, Iterable<CL
 
     private final CLEvent[] events;
 
+    /**
+     * stores event ids for fast access.
+     */
     final PointerBuffer IDs;
+    
+    /**
+     * Points always to the first element of the id buffer.
+     */
+    final PointerBuffer IDsView;
+    
     int size;
 
     public CLEventList(int capacity) {
         this.events = new CLEvent[capacity];
         this.IDs = PointerBuffer.allocateDirect(capacity);
+        this.IDsView = PointerBuffer.wrap(IDs.getBuffer());
     }
 
     public CLEventList(CLEvent... events) {
         this.events = events;
         this.IDs = PointerBuffer.allocateDirect(events.length);
+        this.IDsView = PointerBuffer.wrap(IDs.getBuffer());
+        
         for (CLEvent event : events) {
             IDs.put(event.ID);
         }
diff --git a/test/com/jogamp/opencl/CLCommandQueueTest.java b/test/com/jogamp/opencl/CLCommandQueueTest.java
index e40d07e4..1d47ced6 100644
--- a/test/com/jogamp/opencl/CLCommandQueueTest.java
+++ b/test/com/jogamp/opencl/CLCommandQueueTest.java
@@ -35,8 +35,11 @@ import java.util.concurrent.CountDownLatch;
 import com.jogamp.opencl.util.MultiQueueBarrier;
 import com.jogamp.opencl.CLCommandQueue.Mode;
 import com.jogamp.opencl.CLMemory.Mem;
+import com.jogamp.opencl.util.CLDeviceFilters;
+import com.jogamp.opencl.util.CLPlatformFilters;
 import java.io.IOException;
 import java.nio.ByteBuffer;
+import java.nio.IntBuffer;
 import java.util.EnumSet;
 import java.util.concurrent.TimeUnit;
 import org.junit.Test;
@@ -47,6 +50,7 @@ import static com.jogamp.opencl.TestUtils.*;
 import static com.jogamp.opencl.CLEvent.*;
 import static com.jogamp.opencl.CLVersion.*;
 import static com.jogamp.common.nio.Buffers.*;
+import static com.jogamp.opencl.CLCommandQueue.Mode.*;
 
 /**
  *
@@ -62,8 +66,8 @@ public class CLCommandQueueTest {
 
         //CLCommandQueueEnums
         EnumSet<Mode> queueMode = Mode.valuesOf(CL.CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL.CL_QUEUE_PROFILING_ENABLE);
-        assertTrue(queueMode.contains(Mode.OUT_OF_ORDER_MODE));
-        assertTrue(queueMode.contains(Mode.PROFILING_MODE));
+        assertTrue(queueMode.contains(OUT_OF_ORDER_MODE));
+        assertTrue(queueMode.contains(PROFILING_MODE));
 
         assertNotNull(Mode.valuesOf(0));
         assertEquals(0, Mode.valuesOf(0).size());
@@ -151,6 +155,67 @@ public class CLCommandQueueTest {
             context.release();
         }
     }
+    
+    @Test
+    public void eventConditionsTest() throws IOException {
+        
+        out.println(" - - - event conditions test - - - ");
+
+        CLPlatform platform = CLPlatform.getDefault(CLPlatformFilters.queueMode(OUT_OF_ORDER_MODE));
+        
+        CLDevice device = null;
+        // we can still test this with in-order queues
+        if(platform == null) {
+            device = CLPlatform.getDefault().getMaxFlopsDevice();
+        }else{
+            device = platform.getMaxFlopsDevice(CLDeviceFilters.queueMode(OUT_OF_ORDER_MODE));
+        }
+        
+        CLContext context = CLContext.create(device);
+        
+        try{
+            
+            CLProgram program = context.createProgram(getClass().getResourceAsStream("testkernels.cl")).build();
+            
+            CLBuffer<IntBuffer> buffer = context.createBuffer(newDirectIntBuffer(new int[]{ 1,1,1, 1,1,1, 1,1,1 }));
+            
+            int elements = buffer.getNIOCapacity();
+            
+            CLCommandQueue queue;
+            if(device.getQueueProperties().contains(OUT_OF_ORDER_MODE)) {
+                queue = device.createCommandQueue(OUT_OF_ORDER_MODE);
+            }else{
+                queue = device.createCommandQueue();
+            }
+            
+            CLEventList writeEvent   = new CLEventList(1);
+            CLEventList kernelEvents = new CLEventList(2);
+            
+            // (1+1)*2 = 4; conditions enforce propper order
+            CLKernel addKernel = program.createCLKernel("add").putArg(buffer).putArg(1).putArg(elements);
+            CLKernel mulKernel = program.createCLKernel("mul").putArg(buffer).putArg(2).putArg(elements);
+            
+            queue.putWriteBuffer(buffer, false, writeEvent);
+            
+            queue.put1DRangeKernel(addKernel, 0, elements, 1, writeEvent, kernelEvents);
+            queue.put1DRangeKernel(mulKernel, 0, elements, 1, writeEvent, kernelEvents);
+            
+            queue.putReadBuffer(buffer, false, kernelEvents, null);
+            
+            queue.finish();
+            
+            writeEvent.release();
+            kernelEvents.release();
+            
+            for (int i = 0; i < elements; i++) {
+                assertEquals(4, buffer.getBuffer().get(i));
+            }
+            
+        }finally{
+            context.release();
+        }
+        
+    }
 
     @Test
     public void profilingEventsTest() throws IOException {
@@ -174,7 +239,7 @@ public class CLCommandQueueTest {
 
             CLProgram program = context.createProgram(getClass().getResourceAsStream("testkernels.cl")).build();
             CLKernel vectorAddKernel = program.createCLKernel("VectorAddGM").setArg(3, elements);
-            CLCommandQueue queue = device.createCommandQueue(Mode.PROFILING_MODE);
+            CLCommandQueue queue = device.createCommandQueue(PROFILING_MODE);
 
             out.println(queue);
 
diff --git a/test/com/jogamp/opencl/testkernels.cl b/test/com/jogamp/opencl/testkernels.cl
index ec7e8bf6..2b8c097d 100644
--- a/test/com/jogamp/opencl/testkernels.cl
+++ b/test/com/jogamp/opencl/testkernels.cl
@@ -1,22 +1,33 @@
 
-    // OpenCL Kernel Function for element by element vector addition
     kernel void VectorAddGM(global const int* a, global const int* b, global int* c, int iNumElements) {
-        // get index into global data array
         int iGID = get_global_id(0);
-        // bound check (equivalent to the limit on a 'for' loop for standard/serial C code
         if (iGID >= iNumElements)  {
             return;
         }
-        // add the vector elements
         c[iGID] = a[iGID] + b[iGID];
     }
 
     kernel void Test(global const int* a, global const int* b, global int* c, int iNumElements) {
-        // get index into global data array
         int iGID = get_global_id(0);
-        // bound check (equivalent to the limit on a 'for' loop for standard/serial C code
         if (iGID >= iNumElements)  {
             return;
         }
         c[iGID] = iGID;
     }
+
+    kernel void add(global int* a, int value, int iNumElements) {
+        int iGID = get_global_id(0);
+        if (iGID >= iNumElements)  {
+            return;
+        }
+        a[iGID] += value;
+    }
+
+    kernel void mul(global int* a, int value, int iNumElements) {
+
+        int iGID = get_global_id(0);
+        if (iGID >= iNumElements)  {
+            return;
+        }
+        a[iGID] *= value;
+    }