嗨,我创建了两个内核来执行一个简单的匹配的粉碎程序,该程序与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个字节(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
远离的读取。这不是顺序访问,因此无法合并。
本文收集自互联网,转载请注明来源。
如有侵权,请联系 [email protected] 删除。
我来说两句