title: Prevent calculation twice in opencl
body: <p>I'm pretty new in ocl programming.</p>
  
  <p>I have 2 million polygons (4 lines each - for this examples - but it is variant) that i need to find intersection with 1000 ellipses.</p>
  
  <p>I need to know which if each ellipse intersected at least with one of the polygons.</p>
  
  <p>For that i've created a poylgons buffer with all the points, and ellipses buffer.</p>
  
  <p>My output buffer is a 1000 items int set to 0 in all items. and the kernel will set 1 (when he found intersection) in the right index according to the ellipse index.</p>
  
  <p>I ran the kernel with a global - 2 dim, {2million, 1000}.</p>
  
  <pre><code>__kernel void polygonsIntersectsEllipses(   __global const Point* pts, 
                                              __global const Ellipse* ellipses, 
                                              __global int* output)
  {
      int polygonIdx = get_global_id(0);
      int ellipseIdx = get_global_id(1);
  
      if (&lt;&lt;isIntersects&gt;&gt;) {
          output[ellipseIdx] = 1;
      }
  }
  </code></pre>
  
  <p>The problem is that once one of the polygons intersects the ellipse i don't need to calculate the rest of the polygons.</p>
  
  <p>I've tried to check <code>output[ellipseIdx] != 0</code> before intersection test, but the performance didn't change so much.</p>
  
  <p>I've tried to do single dim global - giving the 1000 (ellipses) and in the kernel run over the millions polygons and stop when i've found one, but still not so much change.</p>
  
  <p>Am i doing something wrong ? can i speed this operation up ?
  any tips ?</p>
  
  <p><strong>EDIT</strong></p>
  
  <p>Using the tip from @Moises and doing allot of research I've change my code to run 2 million times, single dimension. using group work items. changed all my structs to native types, skipped modulus operation. And basically where i could copy data from the global to the private/local memory, i did it.</p>
  
  <p>My local size is my device <code>CL_DEVICE_MAX_WORK_GROUP_SIZE</code>, in my cpu&amp;gpu it is 1024, so in a single run i'm covering all my 1000 ellipses.</p>
  
  <p>Host side</p>
  
  <pre><code>size_t global = 1999872; // 2 million divided by 1024 - for the test
  size_t local = 1024;
  </code></pre>
  
  <p>My code look like this now</p>
  
  <pre><code> __kernel void polygonsIntersectsEllipses(  __global const float4* pts, 
                                              __global const float4* ellipses, 
                                              int ellipseCount,
                                              __local float4* localEllipses,
                                              __global int* output)
  {
      // Saving the ellipses to local memory
      int localId = get_local_id(0);
      if (localId &lt; eCount)
          localEllipses[localId] = ellipses[localId];
  
      barrier(CLK_LOCAL_MEM_FENCE);
  
      // Saving the current polygon into private memory
      int polygonIdx = get_global_id(0);
      float2 private_pts[5];
      for (int currPtsIdx = 0; currPtsIdx &lt; 4; currPtsIdx++)
      {
          private_pts[currPtsIdx] = pts[polygonIdx * 4 + currPtsIdx];
      }
  
      // saving the last point as first so i will not use modulus for cycling, in the intersection algorithm
      private_pts[4] = private_pts[0];
  
      // Run over all the ellipse in the local memory including checking if already there is an output
      for (int ellipseIdx = 0; ellipseIdx &lt; ellipseCount &amp;&amp; output[ellipseIdx] == 0; ++ellipseIdx) {
         if (&lt;&lt;isIntersects Using localEllipses array and private_pts&gt;&gt;) {
             output[ellipseIdx] = 1;
         }
      }
  }
  </code></pre>
  
  <p>The results </p>
  
  <p>CPU was not so much improved - 1.1 faster after change.</p>
  
  <p><strong>GPU - 6.5</strong> times faster (i'm thrilled)</p>
  
  <p>Is there any place i can improve even more ?
  Reminder, once one of the polygons intersects the ellipse, we don't need to check with the rest of the polygons. How do i do that ? my trick with asking the output value is not really working - the performances are the same with or without it</p>