2 files changed, 29 insertions, 3 deletions

diff --git a/src/com/jogamp/opencl/CLCommandQueue.java b/src/com/jogamp/opencl/CLCommandQueue.java
index b5ea41b2..970ef0ce 100644
--- a/src/com/jogamp/opencl/CLCommandQueue.java
+++ b/src/com/jogamp/opencl/CLCommandQueue.java
@@ -1631,7 +1631,7 @@ public class CLCommandQueue extends CLObjectResource {
      * Calls {@native clEnqueueNDRangeKernel}.
      */
     public CLCommandQueue putWork(CLWork work) {
-        this.putNDRangeKernel(work.getKernel(), work.getDimension(), work.getWorkOffset(), work.getWorkSize(), work.getGroupSize(), null, null);
+        this.putWork(work, null, null);
         return this;
     }
 
@@ -1639,7 +1639,7 @@ public class CLCommandQueue extends CLObjectResource {
      * Calls {@native clEnqueueNDRangeKernel}.
      */
     public CLCommandQueue putWork(CLWork work, CLEventList events) {
-        this.putNDRangeKernel(work.getKernel(), work.getDimension(), work.getWorkOffset(), work.getWorkSize(), work.getGroupSize(), null, events);
+        this.putWork(work, null, events);
         return this;
     }
 
@@ -1647,7 +1647,14 @@ public class CLCommandQueue extends CLObjectResource {
      * Calls {@native clEnqueueNDRangeKernel}.
      */
     public CLCommandQueue putWork(CLWork work, CLEventList condition, CLEventList events) {
-        this.putNDRangeKernel(work.getKernel(), work.getDimension(), work.getWorkOffset(), work.getWorkSize(), work.getGroupSize(), condition, events);
+        NativeSizeBuffer groupSize = null;
+        if(                        work.getGroupSize().get(0) != 0
+         || work.dimension >= 2 && work.getGroupSize().get(1) != 0
+         || work.dimension == 3 && work.getGroupSize().get(2) != 0) {
+            groupSize = work.getGroupSize();
+        }
+
+        this.putNDRangeKernel(work.getKernel(), work.dimension, work.getWorkOffset(), work.getWorkSize(), groupSize, condition, events);
         return this;
     }
 
diff --git a/test/com/jogamp/opencl/CLProgramTest.java b/test/com/jogamp/opencl/CLProgramTest.java
index 47eb42e0..cf7f45ff 100644
--- a/test/com/jogamp/opencl/CLProgramTest.java
+++ b/test/com/jogamp/opencl/CLProgramTest.java
@@ -436,7 +436,11 @@ public class CLProgramTest {
 
             CLWork1D work = CLWork.create1D(program.createCLKernel("add"));
             work.getKernel().setArgs(buffer, 5, buffer.getNIOCapacity());
+
+            //optimal values
             work.setWorkSize(20, 1).optimizeFor(device);
+            assertEquals(work.workSize.get(0), 20);
+            assertNotSame(work.groupSize.get(0), 0);
 
             queue.putWriteBuffer(buffer, false)
                  .putWork(work)
@@ -445,6 +449,21 @@ public class CLProgramTest {
             while(buffer.getBuffer().hasRemaining()) {
                 assertEquals(5, buffer.getBuffer().get());
             }
+            buffer.getBuffer().rewind();
+
+            // driver choice
+            work.setWorkSize(20);
+            assertEquals(work.workSize.get(0), 20);
+            assertEquals(work.groupSize.get(0), 0);
+
+            queue.putWriteBuffer(buffer, false)
+                 .putWork(work)
+                 .putReadBuffer(buffer, true);
+
+            while(buffer.getBuffer().hasRemaining()) {
+                assertEquals(10, buffer.getBuffer().get());
+            }
+            buffer.getBuffer().rewind();
 
         }finally{
             context.release();