OpenCL内核问题

4

你好,我创建了两个内核来运行一个简单的匹配拼接程序,并使用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的位置0到n,然后下一个碎片的像素从位置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);
    }
}

这个内核垂直地获取碎片边缘,因此所有碎片的第一个像素都存储在位置0到n中,然后所有碎片的第二个像素存储在位置n + 1到m中(其中n =碎片数量,m =添加到n的碎片数量)。该过程类似于先前的过程,每个线程获取一个像素,并将其与其他碎片的相应像素进行匹配。
两种方法都可以正确返回结果,并已针对纯顺序程序进行过测试。理论上,它们应该以大致相同的时间运行,可能是垂直方法更快,因为原子添加不应该对其产生太多影响...但它运行得更慢...有什么想法吗?
这是我用来启动它的代码(我正在使用C#包装器):
theContext.EnqueueNDRangeKernel(1, null, new int[] { minRows * shredcount }, null, out clEvent);

全球总工作量等于每个像素的总数(# Shreds X #每个像素)。

任何帮助将不胜感激。

1个回答

2
这两个内核都可以完成它们应该完成的工作,但其中一个由于我无法解密的原因运行得比另一个慢得多 :/ 唯一的真正区别是我如何存储发送的数据以及匹配的方式。
这就产生了很大的差异。这是一个经典的合并问题。您在问题中没有指定GPU型号或供应商,因此我必须保持模糊,因为实际数字和行为完全取决于硬件,但一般思路是相当可移植的。
GPU中的工作项一起发出内存请求(读取和写入)到内存引擎(通过“warp”/“wavefront”/“子组”)。该引擎以事务方式提供内存(大小为2的幂次方的16到128字节的块)。让我们假设以下示例的大小为128。
进入内存访问合并:如果一个warp中32个工作项读取4个连续的字节(int或float),内存引擎将发出单个事务来服务所有32个请求。但对于每个与另一个超过128个字节的读取,需要发出另一个事务。在最坏的情况下,这是128字节的32个事务,其成本要高得多。
您的水平内核执行以下访问:
allShreds[(i * pixelCount) + thisPixel]

"

我可以帮忙翻译关于IT技术的内容,下面是您需要翻译的内容:

(i * pixelCount) 在所有工作项中都是常量,只有thisPixel 不同。根据您的代码并假设工作项0具有 thisPixel = 0,则工作项1具有thisPixel = 1, 以此类推。这意味着您的工作项正在请求相邻的读取,因此您会得到完美共享访问。对于调用 atomic_add 也是如此。

另一方面,您的垂直核心执行以下访问:

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

matchShrednumShreds在所有线程中都是常量,只有thisRowmyShred会变化。这意味着您正在请求相互间隔numShreds的读取。这不是顺序访问,因此不是协同的。


你一直在哪里啊 :) 所以基本上,垂直内核需要比水平内核更多的内存调用,因此会降低性能? - Simon Naude
这对我来说最有意义了 :) 非常感谢,你真的帮了我大忙 :) - Simon Naude

网页内容由stack overflow 提供, 点击上面的
可以查看英文原文,
原文链接