低性能 - 补丁匹配。 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 时性能不佳

1063gpu频率低是啥故障

深度学习框架只为GPU? 答案在这里

markdown TVM使用autotvm调优NVIDIA GPU上的高性能卷积

gpu占用低和内存频率有关系吗