OpenCL核心问题

Question

嗨我创建了两个内核来做一个简单的匹配deshredder程序与OpenCL和定时运行。这两个内核做他们应该做的事情，但是由于我无法破译的原因，其中一个运行速度比另一个慢得多：/唯一真正的区别是我如何存储发送的数据以及匹配是如何发生的。

__kernel void Horizontal_Match_Orig( 
__global int* allShreds, 
__global int* matchOut, 
const unsigned int shredCount, 
const unsigned int pixelCount) 

{ 
    int match = 0; 
    int GlobalID = get_global_id(0); 
    int currShred = GlobalID/pixelCount; 
    int thisPixel = GlobalID - (currShred * pixelCount); 
    int matchPixel = allShreds[GlobalID];//currShred*pixelCount+thisPixel]; 
    for (int i = 0; i < shredCount; i++) 
    { 

     match = 0; 
     if (matchPixel == allShreds[(i * pixelCount) + thisPixel]) 
     { 
      if (matchPixel == 0) 
      { 
       match = match + 150; 
      } 
      else match = match + 1; 
     } 
     else match = match - 50; 
     atomic_add(&matchOut[(currShred * shredCount) + i], match); 
    } 
} 

此内核水平得到切丝边缘，所以一个切丝的像素占用POS 0到n在阵列allShreds再下切丝的像素被从POS n + 1个存储到M（其中n =像素数量，m =增加的像素数量）。 GPU的每个线程都有一个像素与合作，并与它匹配反对所有其他Shred的对应像素（包括自身）

__kernel void Vertical(
    __global int* allShreds, 
    __global int* matchOut, 
    const int numShreds, 
    const int pixelsPerEdge) 
{ 
    int GlobalID = get_global_id(0); 
    int myMatch = allShreds[GlobalID]; 
    int myShred = GlobalID % numShreds; 
    int thisRow = GlobalID/numShreds; 
    for (int matchShred = 0; matchShred < numShreds; matchShred++) 
    { 
     int match = 0; 
     int matchPixel = allShreds[(thisRow * numShreds) + matchShred]; 
     if (myMatch == matchPixel) 
     { 
      if (myMatch == 0) 
       match = 150; 
      else 
       match = 1; 
     } 
     else match = -50; 
      atomic_add(&matchOut[(myShred * numShreds) + matchShred], match); 
    } 
} 

这个内核垂直得到一丝一毫的边缘，因此，所有的碎片的第一像素被存储在pos 0到n中，则所有碎片的第二个像素被存储在pos n + 1 ot m（其中n =碎片的数量，并且m =被添加到n的碎片的数量）。该过程类似于前一个过程，每个线程获取一个像素并将其与其他每个细丝的相应像素进行匹配。

两者都给出了相同的结果针对纯序列程序测试的正确结果。从理论上讲，它们应该运行在大致相同的时间内，垂直运行的可能性会更快，因为原子添加对它的影响不应该太大......但是运行速度会更慢......任何想法？

这是我使用启动它的代码（我使用它为C＃包装）：

theContext.EnqueueNDRangeKernel(1, null, new int[] { minRows * shredcount }, null, out clEvent); 

与全球总工作量等于像素中的每个的总数量（#Shreds X #Pixels一）。

任何帮助，将不胜感激

Answer 1

两个内核做他们应该做的事情，但一跑远远超过了其他的原因，我无法破译的要慢/ 唯一的区别是如何我存储正在发送的数据以及匹配是如何发生的。

这让一切变得不同。这是一个典型的聚结问题。你没有在你的问题中指定你的GPU模型或供应商，所以我必须保持模糊，因为实际的数字和行为完全依赖于硬件，但总的想法是合理的便携。

GPU中的工作项发出内存请求（读取和写入）在一起（通过“warp”/“wavefront”/“sub-group”）到内存引擎。该引擎在事务中提供内存（两个大小的功耗为16到128字节的块）。下面的例子假设大小为128。

输入存储器存取凝聚：如果经的32项工作，阅读4个字节（int或float）是连续的内存，存储引擎将问题单笔交易为所有32个请求。但是对于每个超过128字节的读取，需要发出另一个事务。在最坏的情况下，这是每个128字节的32个事务，这样会更昂贵。

---

你的水平内核做以下访问：

allShreds[(i * pixelCount) + thisPixel] 

的(i * pixelCount)是整个工作项目不变，只是thisPixel变化。给定您的代码并假设工作项目0具有thisPixel = 0，那么工作项目1具有thisPixel = 1等等。这意味着您的工作项目正在请求相邻的读取，因此您可以获得完美的合并访问权限。同样的电话atomic_add。

在另一方面，你的垂直内核执行以下操作访问：

allShreds[(thisRow * numShreds) + matchShred] 
// ... 
matchOut[(myShred * numShreds) + matchShred] 

matchShred和numShreds是跨线程不变，仅thisRow和myShred变化。这意味着您请求的读取距离彼此为numShreds。这不是顺序访问，因此不能合并。