如何确定CUDA GPU性能？

StackOverflow https://stackoverflow.com//questions/21062607

题

我正在编写一个CUDA程序，用于匹配分辨率〜180x180的每个输入图像，大约10,000个分辨率〜128 * 128的模板图像。目标是实现实时性能，即25〜30个输入图像的模板匹配（每个包含所有10,000个模板）在1秒内。

目前我正在使用以下方法

  1. 预加载GPU全局内存的所有模板，以保存运行时I / O操作。
  2. 创建了一个内核，将一个源图像与所有模板图像匹配，并返回所有正匹配的数组。
  3. 在时间域中进行所有操作（不使用FFT）。原因是，我尝试了RADIX-4 FFT恳求，但它需要大量的中间全球读取和结束时花费更多的时间。
    4. 到目前为止，对于1个输入图像到10,000个模板，它需要大约2秒钟。

    我的问题是：

    1. 是否有方法可以确定是否可以实时实现这项任务吗？我的意思是在最大的闪光波和I / O带宽限制的帮助下。
    2. 如何计算GPU在其最大值最大限度地被充分利用？
    3. 提高性能的可能方法？

      4. 机器规格：[I7-4770，8GB，GTX-680]

      当前内核代码的解释：

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

        6. 内核代码：

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

    }
}
        .

有帮助吗？

解决方案

我会尝试回答你的大多数问题

是否有方法可以确定是否可以实时实现这项任务吗？我的意思是借助MaxMimumlumlow和I / O带宽限制e.t.c。

我不知道如何确定内核是否是实时可实现的，您可以使用 CUDA占用计算器， 您可以考虑使用纹理，曲面内存，常量内存，固定主机和更多，这些算法实现。

如何计算GPU在其最大值最大限度地被充分利用？

可以使用CUDA占用计算器和CUDA VICESS PROFILER。 我强烈推荐使用视觉探查器，它将通过CUDA的理解来指导您。

改善性能的可能方法？

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

如果这不起作用，请尝试使用延迟，操作多个线程同时访问GPU，自CC 3.5 CUDA推出HyperQ，这可能会帮助您并行完成多个呼叫。

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

请让我们发布

许可以下： CC-BY-SA 和 归因
不隶属于 StackOverflow
scroll top