OpenCL内核故障

西蒙·瑙德

嗨,我创建了两个内核来执行一个简单的匹配的粉碎程序,该程序与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中,然后所有碎片的第二个像素存储在pos n + 1 ot m中(其中n =碎片数量,m =添加到n中的碎片数量)。该过程类似于上一个过程,在该过程中,每个线程都获得一个像素,并将其与其他每个碎片的对应像素匹配。

两者都给出相同的结果,而对于纯顺序程序而言,这些结果是正确的。从理论上讲,它们都应该在大致相同的时间内运行,并且垂直运行的速度可能会更快,因为原子加法不应对其产生太大的影响……但是运行速度要慢得多……有什么想法吗?

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

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

全局总工作负载等于像素总数(每个像素中的#Shreds X #Pixels)。

任何帮助将不胜感激

用户名

这两个内核完成了他们应该做的事情,但是一个内核的运行速度比另一个内核慢得多,原因是我无法解读:/唯一真正的区别是我如何存储正在发送的数据以及匹配如何发生。

这一切都与众不同。这是经典的合并问题。您没有在问题中指定GPU型号或供应商,因此由于实际数字和行为完全取决于硬件,因此我必须保持模棱两可,但总体思路是可移植的。

GPU中的工作项一起(通过“扭曲” /“波前” /“子组”)向存储引擎发出存储请求(读取和写入)。该引擎为事务中的内存提供服务(16到128字节的2幂大小的块)。下面的示例假定大小为128。

输入内存访问合并:如果warp的32个工作项在内存中连续读取4个字节(intfloat),则内存引擎将发出一个事务来满足所有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远离的读取这不是顺序访问,因此无法合并。

本文收集自互联网,转载请注明来源。

如有侵权,请联系 [email protected] 删除。

编辑于
0

我来说两句

0 条评论
登录 后参与评论

相关文章