title: OpenCL - how to effectively distribute work items to different devices
body: <p>I'm writing an openCL application where I have N work items that I want to distribute to D devices where N > D and in turn each device can process the elements of its own work item in parallel and thus achieve a sort of "double" parallelism. </p>
  
  <p>Here is the code I have written already to try and achieve this.</p>
  
  <p>First I create a an event for each of my devices and set them all to complete:</p>
  
  <pre class="lang-c prettyprint-override"><code>cl_int err;
  cl_event *events = new cl_event[deviceCount];
  for(int i = 0; i &lt; deviceCount; i++)
  {
      events[i] = clCreateUserEvent(context, &amp;err);
      events[i] = clSetUserEventStatus(events[i], CL_COMPLETE);
  
  }
  </code></pre>
  
  <p>Each device also has its own command queue and its own "instance" of a kernel. </p>
  
  <p>Then I enter into my "main loop" for distributing work items. The code finds the first available device and enqueues it with the work item. </p>
  
  <pre class="lang-c prettyprint-override"><code>/*---Loop over all available jobs---*/
  for(int i = 0; i &lt; numWorkItems; i++)
  {   
      WorkItem item = workItems[i];
  
      bool found = false; //Check for device availability
      int index = -1;     //Index of found device
      while(!found)       //Continuously loop until free device is found.
      {
          for(int j = 0; j &lt; deviceCount; j++) //Total number of CPUs + GPUs
          {
              cl_int status;
              err = clGetEventInfo(events[j], CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(cl_int), &amp;status, NULL);
              if(status == CL_COMPLETE) /*Current device has completed all of its tasks*/
              {
                  found = true; //Exit infinite loop
                  index = j;    //Choose current device
                  break;        //Break out of inner loop
              }
          }
      }
  
      //Enqueue my kernel
      clSetKernelArg(kernels[index], 0, sizeof(cl_mem), &amp;item);
      clEnqueueNDRangeKernel(queues[index], kernels[index], 1, NULL, &amp;glob, &amp;loc, 0, NULL, &amp;events[index]);
  
      clFlush(commandQueues[index]);
  }
  </code></pre>
  
  <p>And then finally I wrap up by calling clFinish on all my devices:</p>
  
  <pre class="lang-c prettyprint-override"><code>/*---Wait For Completion---*/
  for(int i = 0; i &lt; deviceCount; i++)
  {
      clFinish(queues[i]);
  }
  </code></pre>
  
  <p><strong>This approach has a few problems however:</strong></p>
  
  <p>1) It doesn't distribute the work to all my devices. On my current computer I have 3 devices. My algorithm above only distributes the work to devices 1 and 2. Device 3 always gets left out because devices 1 and 2 finish so quickly that they can snatch up more work items before 3 gets a chance. </p>
  
  <p>2) Even with devices 1 and 2 running together, I only see a very, very mild speed increase. For instance if i were to assign all work items to device 1 it might take 10 seconds to complete, and if I assign all work items to device 2 it might take 11 seconds to complete, but if I try to split the work between them, combined it might take 8-9 seconds when what I would hope for might be between 4-5 seconds. I get the feeling that they might not really be running in parallel with each other the way I want. </p>
  
  <p>How do I fix these issues?</p>