防止在opencl中计算两次

Question

我是 ocl 编程的新手。

我有 200 万个多边形（每个 4 行 - 对于这个例子 - 但它是变体），我需要找到与 1000 个椭圆的交点。

我需要知道每个椭圆是否至少与其中一个多边形相交。

为此，我创建了一个包含所有点的多边形缓冲区和椭圆缓冲区。

我的输出缓冲区是 1000 个项目 int 在所有项目中设置为 0。并且内核将根据椭圆索引在正确的索引中设置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

我了解您所有的 200 万x1000 线程，读取其自己的多边形数据和椭圆，对吗？所以，对于每个多边形，每个线程读取 1000 次相同的内存位置（使用多边形数据），不是吗？为了避免这种内存绑定行为，您可以只创建 200 万个线程并使用 1000 次迭代的 for 循环来迭代椭圆的数量。或者一个中间解决方案，有一个由 2millionx64 个线程组成的网格，每个线程为每个多边形计算 16 个椭圆。我不知道这些是否比您的解决方案更好，但它们避免了多余的内存访问。

问候， 莫伊塞斯

Answer 2

优化：

• 使用最少的内存量。 __global int* output 保存布尔值太多，请改用char。或者更好的是，使用二进制数组。 （与全局读取相比，二进制操作速度更快）

• 您不应从每个线程output[ellipseIdx] == 0 再次从全局内存中读取。这非常慢，相反，将其与开头的省略号数据一起保存到本地内存中。注意：只有在找到匹配的组之后启动的本地组才会受益于加速。然而，这将节省大量的全局读取，这比节省一些计算要好得多。此外，本地组无法受益，因为当工作项找到匹配项时，所有本地工作项都已处理该省略号。

  __kernel void 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;
         }
     }
  }
  

  }