你好,我创建了两个内核来运行一个简单的匹配拼接程序,并使用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 #每个像素)。
任何帮助将不胜感激。