OpenCL 内核问题

Question

您好，我创建了两个内核来执行一个简单的匹配碎纸机程序，以使用 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);
    }
}

这个内核水平获取碎片边缘，因此一个碎片的像素占据数组allShreds中的pos 0到n，然后下一个碎片的像素从pos n+1存储到m（其中n =像素，m = 添加的像素数）。 GPU 的每个线程都获得一个像素来处理，并将其与所有其他碎片（包括其自身）的相应像素进行匹配

__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 中，然后所有碎片的第 2 个像素存储在 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。

进入内存访问合并：如果一个 warp 的 32 个工作项读取内存中连续的 4 个字节（int 或 float），则内存引擎将发出单个事务来服务所有 32 个请求。但是对于每一次超过 128 字节的读取，都需要发出另一个事务。在最坏的情况下，这是 32 个事务，每个事务 128 字节，这会更加昂贵。

---

您的水平内核执行以下访问：

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 的读取。这不是顺序访问，因此不会合并。