performance - 提高二维图像 'tracing' CUDA 内核性能的技巧？

Question

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

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

现在必须在所有方向和角度跟踪每个图像，并且必须对从跟踪中提取的值进行一些计算。因此，例如，对于 500x500我需要的图片229080 threads计算低于它的内核是 SUM 的值(这就是为什么我检查线程 ID idHilo 没有超过它)。我将几个数组一个接一个地复制到设备的全局内存中，因为我需要访问它们来计算所有长度 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;
        }
    }
}

}

编辑:

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

Answer 1

对于非结构化的答案，我很抱歉，我只想抛出一些通常有用的注释，并引用您的代码，使其对其他人更有用。

算法更改始终是优化的第一位。是否有另一种方法可以解决需要较少数学/迭代/内存等的问题？

如果精度不是一个大问题，请使用浮点(或使用较新架构的半精度浮点)。当您短暂尝试时它没有对您的性能产​​生太大影响的部分原因是因为您仍在对浮点数据使用 double 计算(fabs 需要 double ，因此如果您使用 float，它会将您的浮点数转换为 double 数, 进行 double 运算，返回 double 并转换为浮点数，使用 fabsf)。

如果您不需要使用 float 的绝对全精度，请使用快速数学(编译器选项)。

乘法比除法快得多(特别是对于全精度/非快速数学)。在内核外计算 1/var，然后在内核内进行乘法而不是除法。

不知道它是否得到优化，但您应该使用增量和减量运算符。 v=v-1；可能是 v--;等等。

转换为 int 将向零截断。 floor() 将向负无穷大截断。您可能不需要显式的 floor()，也不需要上述的用于 float 的 floorf()。当您将它用于整数类型的中间计算时，它们已经被截断了。因此，您无缘无故地转换为 double 并返回。使用适当类型的函数(abs、fabs、fabsf 等)

if(fabs((double)neighborhood[v]) < M && fabs((double)ji)<N)
change to
if(abs(neighborhood[v]) < M && abs(ji)<N)

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)])]);
change to 
vector_trazo_cuda[20+idHilo*31] +=
    fabsf(imagen_cuda[ji*M+(neighborhood[v-(15+k)])]-
        imagen_cuda[ji*M+(neighborhood[v-(15-k)])]);

.

xi = xb_cuda[idHilo] + floor((double)t*xinc_cuda[idHilo]);
change to
xi = xb_cuda[idHilo] + t*xinc_cuda[idHilo];

上面的行是不必要的复杂。从本质上讲，您正在这样做，
将 t 转换为双倍，
将 xinc_cuda 转换为加倍和乘法，
地板它(返回双倍)，
将 xb_cuda 转换为 double 并添加，
转换为长。

新行将在更短的时间内存储相同的结果(也更好，因为如果您在前一种情况下超过 double 的精度，您将四舍五入到最接近的 2 次方)。 此外，这四行应该在 for 循环之外 ...如果它们不依赖于 t，则不需要重新计算它们。总之，如果这将您的运行时间减少了 10-30 倍，我不会感到惊讶。

您的结构会导致大量全局内存读取，尝试从全局读取一次，处理本地内存上的计算，并写入一次到全局(如果可能的话)。

始终使用 -lineinfo 编译。使分析更容易，而且我无法评估任何开销(使用 0.1 到 10 毫秒执行时间范围内的内核)。

如果您受计算或内存限制，请使用分析器确定并相应地投入时间。

尽量允许编译器在可能的情况下使用寄存器，这是一个很大的话题。

与往常一样，不要一下子改变一切。我通过编译/测试输入了所有这些，所以我可能有错误。