c++ - 如何确定 CUDA gpu 性能？

Question

我正在编写一个 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. 所有模板图像 [RGB 大小约为 128X128] 都预先加载到 GPU 内存中。想法是在运行时操作期间节省 I/O。
2. 每个输入图像都加载到纹理内存中，原因是纹理是二维寻址的不错选择。
3. 每个“ block ”有 1024 个线程。
4. 每个线程计算每个输出像素的值，输出大小为 [31X31 = 961 像素]。
5. 启动的 block 数等于匹配的模板图像数。

内核代码:

__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;

    }
}

Answer 1

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

Is there is way to determine if this task is achievable in realtime or not? I mean with the help of maximum FLOPS and I/O bandwidth limitations etc.

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

How to compute if the GPU is being fully utilitzed at its maximum?

您可以使用 CUDA Occupancy Calculator 和 CUDA visual profiler。 我强烈建议使用可视化分析器，它会指导您了解 CUDA。

Possible ways to improve the performance?

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

如果这行不通，请尝试使用延迟，操作多个线程同时访问 GPU，因为 CC 3.5 CUDA 启动了 HyperQ，这可能会帮助您并行完成多个调用。

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