在opencl中防止計算兩次

Question

我是ocl編程的新手。

我有200萬個多邊形（每個4行 - 對於這個例子 - 但它是變種），我需要找到1000個橢圓的交集。

我需要知道哪個橢圓與至少一個多邊形相交。

爲此，我創建了一個包含所有點和橢圓緩衝區的poylgons緩衝區。

我的輸出緩衝區是在所有項目中設置爲0的1000項目int。內核會根據橢圓索引在右邊的索引中設置1（當他找到交集）。我運行的內核有一個全局的-2 dim，{2million，1000}。

__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 (<<isIntersects>>) { 
     output[ellipseIdx] = 1; 
    } 
} 

問題是，一旦其中一個多邊形與橢圓相交，我不需要計算剩餘的多邊形。

我試過在路口測試前檢查output[ellipseIdx] != 0，但是性能沒有太大變化。

我試圖做單暗全局 - 給1000（橢圓）和內核運行在數百萬個多邊形，並停止當我找到一個，但仍然沒有太多的變化。

我做錯了什麼？我可以加快這個操作嗎？ 任何提示？

編輯

使用從@Moises尖端和做研究配發我改變我的代碼運行200萬次，單維。使用小組工作項目。將我所有的結構改爲原生類型，跳過模數運算。基本上，我可以將數據從全局複製到私有/本地內存，我做到了。

我的本地尺寸是我的設備CL_DEVICE_MAX_WORK_GROUP_SIZE，在我的cpu & gpu它是1024，所以在一次運行中我覆蓋了我所有的1000個橢圓。

主機端

size_t global = 1999872; // 2 million divided by 1024 - for the test 
size_t local = 1024; 

我的代碼看起來現在這個樣子

__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 < 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 < 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 < ellipseCount && output[ellipseIdx] == 0; ++ellipseIdx) { 
     if (<<isIntersects Using localEllipses array and private_pts>>) { 
      output[ellipseIdx] = 1; 
     } 
    } 
} 

結果

CPU就沒有那麼多的改進 - 更改後1.1快。

GPU - 快6.5倍（我很高興）

有沒有我可以更好地改善任何地方嗎？ 提醒，一旦其中一個多邊形與橢圓相交，我們就不需要檢查其餘的多邊形。我怎麼做 ？我的要求輸出值的技巧並不真正起作用 - 性能是相同或不相似的

Answer 1

我知道您所有的2百萬x1000線程都讀取了它自己的多邊形數據和橢圓嗎？因此，對於每個多邊形，每個線程讀取1000次相同的內存位置（與多邊形數據），不是嗎？爲了避免這種內存綁定行爲，您只能創建2百萬個線程，並使用1000次迭代的循環遍歷橢圓的數量。或者是一箇中間解決方案，具有2百萬x 64個線程的網格，其中每個線程爲每個多邊形計算16個橢圓。我不知道這些是否比你的解決方案好，但他們避免了多餘的內存訪問。

問候， 莫伊塞斯

Answer 2

優化：

• 使用的內存最小量。 __global int* output舉行一個布爾值太多，請使用char來代替。或者甚至更好，使用二進制數組。 （與全局讀取相比，二進制操作速度很快）

• 不應從每個線程的全局內存中讀取output[ellipseIdx] == 0。這是非常緩慢的，相反，將它保存到本地存儲器中，並在開始時使用elipses數據。注意：只有在發現一場比賽的小組之後發起的本地小組才能從加速中受益。但是，這將節省大量的全局讀取，這比保存一些計算要好得多。此外，本地組無法獲益，因爲當工作項目找到匹配項時，所有本地工作項目都已處理該橢圓。

  __kernel無效polygonsIntersectsEllipses（

  __global const float4* pts, 
  __global const float4* ellipses, 
  int ellipseCount, 
  __local float4* localEllipses, 
  __local char* localOuts, 
  __global char* output){ 
  
  // Saving the ellipses to local memory 
  int localId = get_local_id(0); 
  if (localId < eCount) 
      localOuts[localId] = output[localId]; 
  barrier(CLK_LOCAL_MEM_FENCE); 
  if (localId < eCount && localOuts[localId]) // Do not copy elipses if we are not going to check them anyway 
      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 < 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 < ellipseCount; ++ellipseIdx) { 
      if (localOuts[ellipseIdx] == 0){ 
       if (<<isIntersects Using localEllipses array and private_pts>>) { 
        localOuts[ellipseIdx] = 1; 
       } 
       barrier(CLK_LOCAL_MEM_FENCE); 
       if(localOuts[ellipseIdx] && localId == 0){ 
        output[ellipseIdx] = 1; 
       } 
      } 
  } 
  

  }