OpenCL 内核问题

Question

您好，我创建了两个内核来执行一个简单的匹配 deshredder 程序运行与 OpenCL 和定时。这两个内核做它们应该做的事情，但是一个运行s 比另一个慢得多，原因是我无法破译：/唯一真正的区别是我如何存储正在发送的数据以及如何进行匹配.

__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 =像素数，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 = number碎片数，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 字节的二次方大小的块）。让我们假设以下示例的大小为 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 的读取。这不是顺序访问，因此不会合并。

OpenCL 内核问题

OpenCL Kernel Troubles

c#

c++

opencl