c# - 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);
    }
}

这个kernel是水平获取shred的边缘，所以一个shred的像素在数组allShreds中占据pos 0到n，然后下一个shred的像素从pos n+1到m存储(其中n = number of of像素，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 中，然后所有碎片的第二个像素存储在 pos n+1 ot m 中(其中 n = 碎片数, m = 添加到 n) 的碎片数。该过程类似于之前的过程，其中每个线程获取一个像素并将其与每个其他碎片的相应像素相匹配。

两者都给出了相同的结果，是针对纯顺序程序测试的正确结果。从理论上讲，它们都应该在大致相同的时间内运行，垂直的可能运行得更快，因为原子添加不应该影响它那么多......但是它运行得慢得多......有什么想法吗？

这是我用来启动它的代码(我正在为它使用 C# 包装器):

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

全局总工作负载等于像素总数(#Shreds X #Pixels in each one)。

任何帮助将不胜感激

Answer 1

The two kernels do what they are supposed to do, but one runs far slower than the other for a reason i cannot decipher :/ The only real difference is how i store the data being sent up and how the matching happens.

这一切都不同了。这是一个经典的聚结问题。您没有在您的问题中指定您的 GPU 型号或供应商，因此我必须保持含糊不清，因为实际数字和行为完全取决于硬件，但总体思路是可移植的。

GPU 中的工作项一起(通过“warp”/“wavefront”/“sub-group”)向内存引擎发出内存请求(读取和写入)。该引擎在事务中提供内存(16 到 128 字节的二次方大小的 block )。对于以下示例，我们假设大小为 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 的读取。这不是顺序访问，因此不会合并。