低性能 - 补丁匹配。 GPU上的图像处理(CUDA)
Posted
tags:
篇首语:本文由小常识网(cha138.com)小编为大家整理,主要介绍了低性能 - 补丁匹配。 GPU上的图像处理(CUDA)相关的知识,希望对你有一定的参考价值。
我遇到了性能问题:CPU和GPU的性能几乎相同。
我处理的问题是PATCH MATCH。我有2个矩阵。我想找到大矩阵和小矩阵之间最大相似性的位置。
矩阵的二进制值为0/1(黑色和白色)。
当我使用i5 CPU检查小矩阵与大矩阵之间的匹配时,需要30ms(使用多线程)。
当我在Ge-force GT 730中检查小矩阵与大矩阵之间的匹配时,它也需要33ms。
我希望GPU在至少1个数量级的订单中工作得更快。我对目前的结果感到非常失望。
我有两个矩阵:
1)大 - 300000(300行,1000列)
2)小50000(50行,1000列)
比较过程是通过将大矩阵分成250个子矩阵,然后将每个矩阵与小矩阵进行比较,然后找到最高相似度来完成的。
相似性标准是两个矩阵(小和子大)上的相应黑色像素的总和除以子大的黑色像素的总和。
我使用以下CUDA代码完成了最后一项任务:
__global__ void matCompare_cuda (uint8_t *D_SUB , uint8_t *D_SMALL , float *D_RSLTS , unsigned int step, int numOfIndentations ,int SUB_size, int SMALL_size)
{
int i = 0 , j = 0 , success = 0, sumZero = 0;
int tid = threadIdx.x + blockIdx.x * blockDim.x;
int LoopIndex = ( tid * step );
if (tid < numOfIndentations)
{
for ( j = 0 ; j < (SMALL_size) ; j++)
{
i = j + LoopIndex;
if ( D_SUB[i] == 0 )
{
{
sumZero++;
if ( D_SMALL[j] == 0 )
success++;
}
}
}
if ( success > 0 && sumZero > 500)
D_RSLTS[tid] = 100*((float)success / sumZero) ;
}
}
Kernal发射:
int numOfIndentations = 300-50 //[ (big.row) - (small.row)]
int numBlock = 16;
int threadNumber = numOfIndentations/numBlock;
matCompare_cuda<<< numBlock , threadNumber >>> ( D_SUB , D_SMALL , D_RSLTS , step, numOfIndentations, SUB_size, SMALL_size );
Cpu代码:
for (i=0; i < (pixelNum) ; i++)
{
if (SUB[i]==0)
{
sumDots = sumDots +1;
if (SMALL->Image[i]==0)
{
success = success + 1;
}
}
}
if (success>0)
if (sumDots>500)
RSLT=((float)success/sumDots)*100;
你是否看到GPU代码可以做任何改进?
一些东西。尽可能避免使用if。你可以写在这里:
sumZero += (1 - D_SUB[i])
success += (1 - D_SUB[i]) * (1 - D_SMALL[j])
但是我不认为你会在这里看到巨大的差异。我看到两个原因。
一个是调用cuda的开销很大。需要将数据复制到图形卡并返回。这吃了你获得的一些加速。不知道它有多少,但由于运行时间太短,它可以发挥作用。我希望你没有时间编译内核和其他一次性的东西(通过在循环中运行代码并忽略前几次迭代来取出它们)。
其次你的大矩阵太小而你的小矩阵太大了。因为小矩阵是如此之大(1000列),我猜它在CPU缓存行中运行得非常好。如果小矩阵较小,则必须更频繁地转到下一行,这会增加破坏缓存行的机会。 gpu使用矩形进行缓存,因此不会出现问题。如果大矩阵要大一些,你也会增加所需的计算量,这样GPU就会开始领先。
以上是关于低性能 - 补丁匹配。 GPU上的图像处理(CUDA)的主要内容,如果未能解决你的问题,请参考以下文章
将 glDrawElements 性能与预期 GPU 性能进行比较?
同时使用 2 个 GPU 调用 cudaMalloc 时性能不佳