提高 2d 图像“跟踪”CUDA 内核性能的技巧？

Question

你能给我一些优化这个 CUDA 代码的技巧吗？

我在计算能力为 1.3 的设备上运行它（我需要它用于 Tesla C1060，尽管我现在正在具有相同计算能力的 GTX 260 上测试它），并且我有几个如下所示的内核。执行该内核所需的线程数量由 给出long SUM并取决于size_t M以及size_t N作为参数接收的矩形图像的尺寸，它可以以像素或更多为单位变化50x50很大10000x10000。尽管我最感兴趣的是使用 Cuda 处理更大的图像。

现在，必须在所有方向和角度上跟踪每个图像，并且必须对从跟踪中提取的值进行一些计算。因此，例如，对于一个500x500图像，我需要229080 threads计算该内核，其下方是 SUM 的值（这就是为什么我检查线程 ididHilo是否不超过它）。我将几个数组一个接一个地复制到设备的全局内存中，因为我需要访问它们以进行所有长度的计算SUM。像这样

cudaMemcpy(xb_cuda,xb_host,(SUM*sizeof(long)),cudaMemcpyHostToDevice);

cudaMemcpy(yb_cuda,yb_host,(SUM*sizeof(long)),cudaMemcpyHostToDevice);

...etc

因此每个数组的每个值都可以由一个线程访问。所有这些都在内核调用之前完成。根据 Nsight 上的 Cuda Profiler，最长的内存复制持续时间是246.016 us针对500x500图像的，因此不会花费那么长时间。

但是，像我在下面复制的那样的内核对于任何实际用途来说都花费了太长的时间（根据下面的内核的 Cuda 分析器，对于 500x500 图像，需要 3.25 秒，对于持续时间最长的内核，需要 5.052 秒），所以我需要看看是否我可以优化它们。

我这样安排数据

首先是块尺寸

dim3 dimBlock(256,1,1);

然后是每个网格的块数

dim3 dimGrid((SUM+255)/256);

895 blocks对于一个图像的多个500x500。

我不确定在我的情况下如何使用合并和共享内存，或者即使使用数据的不同部分多次调用内核是个好主意。数据彼此独立，因此理论上我可以多次调用该内核，而不是在需要时一次性使用 229080 个线程。

现在考虑外for循环

for(t=15;t<=tendbegin_cuda[idHilo]-15;t++){

依赖于取决于

tendbegin_cuda[idHilo]

它的值取决于每个线程，但大多数线程都有相似的值。

根据 Cuda Profiler，全局存储效率是针对该内核的0.619，全局加载效率是0.951针对该内核的。其他内核也有类似的值。

这样好吗？坏的？我该如何解释这些价值观？遗憾的是，计算能力 1.3 的设备不提供其他有用的信息来评估代码，例如多处理器和内核内存或指令分析。分析后得到的唯一结果是“全局内存存储效率低”和“全局内存加载效率低”，但我不确定如何优化它们。

void __global__ t21_trazo(long SUM,int cT, double Bn, size_t M, size_t N, float* imagen_cuda, double* vector_trazo_cuda, long* xb_cuda, long* yb_cuda, long* xinc_cuda, long* yinc_cuda, long* tbegin_cuda, long* tendbegin_cuda){

long xi;
long yi;
int t;
int k;
int a;
int ji;
long idHilo=blockIdx.x*blockDim.x+threadIdx.x;

int neighborhood[31];
int v=0;

if(idHilo<SUM){

    for(t=15;t<=tendbegin_cuda[idHilo]-15;t++){

        xi = xb_cuda[idHilo] + floor((double)t*xinc_cuda[idHilo]);
        yi = yb_cuda[idHilo] + floor((double)t*yinc_cuda[idHilo]);
        neighborhood[v]=floor(xi/Bn);
        ji=floor(yi/Bn);

        if(fabs((double)neighborhood[v]) < M && fabs((double)ji)<N)
        {
            if(tendbegin_cuda[idHilo]>30 && v==30){

                if(t==0)
                vector_trazo_cuda[20+idHilo*31]=0;

                for(k=1;k<=15;k++)
                vector_trazo_cuda[20+idHilo*31]=vector_trazo_cuda[20+idHilo*31]+fabs(imagen_cuda[ji*M+(neighborhood[v-(15+k)])]-
                            imagen_cuda[ji*M+(neighborhood[v-(15-k)])]);


                for(a=0;a<30;a++)
                neighborhood[a]=neighborhood[a+1];

                v=v-1;
            }

            v=v+1;
        }
    }
}

}

编辑：

将 DP 触发器更改为 SP 触发器仅略微改善了持续时间。循环展开内部循环实际上没有帮助。

Answer 1

您可能同时运行太多线程。当您运行正确数量的线程时，似乎就会出现最佳性能：足够的线程保持忙碌，但数量又不会过多，以免导致每个并发线程可用的本地内存过度碎片化。

去年秋天，我编写了一个教程来研究使用 CUDA 和 CUDAFY 优化旅行商问题 (TSP)。尽管问题领域不同，但我在从已发布的算法中实现数倍加速时所经历的步骤可能对指导您的工作很有用。该教程和代码可在CUDA Tuning with CUDAFY中找到。