c# – OpenCL内核故障

2023年7月12日 133次阅读

嗨我创建了两个内核来做一个简单的匹配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);
    }
}

这个内核水平获取碎片边缘,因此一个碎片的像素占据阵列allShreds中的pos 0到n,然后下一个碎片的像素从pos n 1到m存储(其中n =像素数,m =添加的像素数). GPU的每个线程都可以使用一个像素,并将其与所有其他Shred(包括其自身)的相应像素进行匹配

__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)的碎片数量.该过程类似于前一个过程,其中每个线程获得一个像素并将其与每个其他Shred的相应像素相匹配.

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

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

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

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

任何帮助将不胜感激

最佳答案

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

点赞