低性能 - 补丁匹配。 GPU上的图像处理（CUDA）

Question

我遇到了性能问题：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代码可以做任何改进？

Answer 1

一些东西。尽可能避免使用if。你可以写在这里：

sumZero += (1 - D_SUB[i])
success += (1 - D_SUB[i]) * (1 - D_SMALL[j])

但是我不认为你会在这里看到巨大的差异。我看到两个原因。

一个是调用cuda的开销很大。需要将数据复制到图形卡并返回。这吃了你获得的一些加速。不知道它有多少，但由于运行时间太短，它可以发挥作用。我希望你没有时间编译内核和其他一次性的东西（通过在循环中运行代码并忽略前几次迭代来取出它们）。

其次你的大矩阵太小而你的小矩阵太大了。因为小矩阵是如此之大（1000列），我猜它在CPU缓存行中运行得非常好。如果小矩阵较小，则必须更频繁地转到下一行，这会增加破坏缓存行的机会。 gpu使用矩形进行缓存，因此不会出现问题。如果大矩阵要大一些，你也会增加所需的计算量，这样GPU就会开始领先。