1

我正在编写一个 cuda 程序，用于匹配每个分辨率 ~180X180 的输入图像，大约 10,000 个分辨率 ~128*128 的模板图像。目标是实现实时性能，即在 1 秒内对 25~30 个输入图像（每个都有 10,000 个模板）进行模板匹配。

目前我正在使用以下方法

  1. 在 GPU 全局内存上预加载所有模板以节省运行时 I/O 操作。
  2. 创建单个内核以将一个源图像与所有模板图像匹配，并为所有正匹配返回一个数组。
  3. 在时域中执行所有操作（不使用 FFT）。原因是，我尝试了 Radix-4 fft 实现，但它需要大量的中间全局读取和写入，最终会花费更多时间。

到目前为止，1 个输入图像到 10,000 个模板，大约需要 2 秒。

我的问题是：

  1. 有没有办法确定这个任务是否可以实时完成？我的意思是借助最大 FLOPS 和 I/O 带宽限制等
  2. 如何计算 GPU 是否被充分利用？
  3. 提高性能的可能方法？

机器规格：[i7-4770、8GB、GTX-680]

当前内核代码说明：

  1. 所有模板图像 [大小约为 128X128 RGB] 都加载到 GPU 内存上。想法是在运行时操作期间节省 I/O。
  2. 每个输入图像都加载到纹理内存中，原因是纹理是 2D 寻址的好选择。
  3. 每个“块”有 1024 个线程。
  4. 每个线程计算每个输出像素的值，输出大小为 [31X31 = 961 像素]。
  5. 启动的块数等于匹配的模板图像数。

内核代码：

__global__ void cudaMatchTemplate(TemplateArray *templates, uchar *Match)
{
    int global = blockIdx.x*blockDim.x + threadIdx.x;

    __shared__ int idx[TEMPLATE_MATCH_DIM];
    __shared__ float out_shared[TEMPLATE_MATCH_DIM];

    //halving the template size....
    int rows = (templates[blockIdx.x].nHeight)/2;
    int cols = (templates[blockIdx.x].nWidth)/2;

    int fullCol = templates[blockIdx.x].nWidth;

    int x = templates[blockIdx.x].nMatchLeft;
    int y = templates[blockIdx.x].nMatchTop;

    int offset_y =  (threadIdx.x/TEMPLATE_MATCH_SIZE);
    int offset_x =  (threadIdx.x - offset_y*TEMPLATE_MATCH_SIZE);

    // *************** Performing match in time domain *****************************//
    int sum = 0;
    float temp;
    int idxXFactor = 3*(2*(offset_x) + x);
    int idxYFactor = 2*(offset_y) + y ;
    
    for (int i = 0; i < rows; i++)
    {
        int I=3*i*fullCol;
        int sourceIdxY = idxYFactor + 2*i;
        for (int j = 0; j < cols; j++)
        {
            int J=3*j;
            int sourceIdxX = idxXFactor + 2*J;          
            int templateIdx = 2*I+2*J;
            //**** R *****//
            temp = float(tex2D(SourceImgColorTex,sourceIdxX,sourceIdxY)) - float(templates[blockIdx.x].pRAWPixels_gpu[templateIdx]);
            sum = sum + temp*temp;
            //**** G *****//
            temp = float(tex2D(SourceImgColorTex,sourceIdxX+1,sourceIdxY)) - float(templates[blockIdx.x].pRAWPixels_gpu[templateIdx +1]);
            sum = sum + temp*temp;
            //**** B *****//
            temp = float(tex2D(SourceImgColorTex,sourceIdxX+2,sourceIdxY)) - float(templates[blockIdx.x].pRAWPixels_gpu[templateIdx +2]);
            sum = sum + temp*temp;
        }
    }

    __syncthreads();
    
//placing all values in shared memory for comparison.
    if(threadIdx.x < TEMPLATE_MATCH_DIM)
    {
        idx[threadIdx.x] = threadIdx.x;
        out_shared[threadIdx.x] = sum;
    }
    __syncthreads();


// //computing the Min location.....//

#pragma unroll
    for(int s=512; s>0; s>>=1) 
    {
        if ((threadIdx.x < s) &&((threadIdx.x + s)<TEMPLATE_MATCH_DIM))
        {
            idx[threadIdx.x] = out_shared[threadIdx.x] < out_shared[threadIdx.x + s] ? idx[threadIdx.x] : idx[threadIdx.x + s];
            out_shared[threadIdx.x]  = out_shared[threadIdx.x] < out_shared[threadIdx.x + s] ? out_shared[threadIdx.x] : out_shared[threadIdx.x + s];           
        }
        
    }

    __syncthreads();

    if(threadIdx.x <1)
    {
        int half_Margin = MARGIN_FOR_TEMPLATE_MATCH/2;
        int matchY = idx[0]/TEMPLATE_MATCH_SIZE ;
        int matchX = idx[0] - matchY * TEMPLATE_MATCH_SIZE;

        int diff = absolute(half_Margin - matchX) + absolute(half_Margin - matchY);
        if(diff < THRESHOLD)
        {
            Match[blockIdx.x] = 1;
        }
        else
            Match[blockIdx.x] = 0;

    }
}
4

1 回答 1

1

我会尽力回答你的大部分问题。

有没有办法确定这个任务是否可以实时完成？我的意思是借助最大 FLOPS 和 I/O 带宽限制等。

我不知道如何确定内核是否可以实时实现，您可以使用CUDA 占用计算器最大化您的 CUDA 内核。您可以考虑使用纹理、表面内存、常量内存、固定主机内存等。这些取决于您的算法实现。

如何计算 GPU 是否被充分利用？

您可以使用 CUDA 占用计算器和 CUDA 视觉分析器。我强烈推荐使用视觉分析器，它会引导你理解 CUDA。

提高性能的可能方法？

有几种有趣的方法可以这样做。第一，您可以使用上述方法最大化您的内核调用。如果这还不够，请尝试使用流对象实现管道，以便同时复制数据和计算作业。

如果还是不行，尝试处理延迟，同时操作多个线程访问 GPU，因为 CC 3.5 CUDA 推出了 HyperQ，这可能会帮助您并行完成多个调用。

如果这不起作用，请考虑使用多个 GPU 设备。

于 2014-01-11T17:26:24.377 回答