为什么这个 CUDA 内核很慢？答案

【问题标题】：Why is this CUDA kernel slow?为什么这个 CUDA 内核很慢？
【发布时间】：2016-10-13 01:11:40
【问题描述】：

我需要帮助让我的 cuda 程序运行得更快。 NVIDIA 视觉分析器显示性能不佳，显示“Low Compute Utilization 1.4%”：

代码如下。第一个内核准备工作：

void laskeSyvyydet(int& tiilet0, int& tiilet1, int& tiilet2, int& tiilet3) {

cudaArray *tekstuuriSisaan, *tekstuuriUlos;

//take care of synchronazion
cudaEvent_t cEvent;
cudaEventCreate(&cEvent);

//let's take control of OpenGL textures
cudaGraphicsMapResources(1, &cuda.cMaxSyvyys);
cudaEventRecord(cEvent, 0);
cudaGraphicsMapResources(1, &cuda.cDepthTex);
cudaEventRecord(cEvent, 0);

//need to create CUDA pointers
cudaGraphicsSubResourceGetMappedArray(&tekstuuriSisaan, cuda.cDepthTex, 0, 0);
cudaGraphicsSubResourceGetMappedArray(&tekstuuriUlos, cuda.cMaxSyvyys, 0, 0);

cudaProfilerStart();

//launch kernel
cLaskeSyvyydet(tiilet0, tiilet1, tiilet2, tiilet3, tekstuuriSisaan, tekstuuriUlos);
cudaEventRecord(cEvent, 0);

cudaProfilerStop();

//release textures back to OpenGL
cudaGraphicsUnmapResources(1, &cuda.cMaxSyvyys, 0);
cudaEventRecord(cEvent, 0);
cudaGraphicsUnmapResources(1, &cuda.cDepthTex, 0);
cudaEventRecord(cEvent, 0);

//final synchronazion
cudaEventSynchronize(cEvent);
cudaEventDestroy(cEvent);
}

内核启动：

void cLaskeSyvyydet(int& tiilet0, int& tiilet1, int& tiilet2, int& tiilet3, cudaArray* tekstuuriSisaan, cudaArray* tekstuuriUlos) {

cudaBindTextureToArray(surfRefSisaan, tekstuuriSisaan);
cudaBindSurfaceToArray(surfRefUlos, tekstuuriUlos);

    int blocksW = (int)ceilf( tiilet0 / 32.0f );
    int blocksH = (int)ceilf( tiilet1 / 32.0f );
    dim3 gridDim( blocksW, blocksH, 1 );
    dim3 blockDim(32, 32, 1 );

kLaskeSyvyydet<<<gridDim, blockDim>>>(tiilet0, tiilet1, tiilet2, tiilet3);

}

还有内核：

__global__ void kLaskeSyvyydet(const int tiilet0, const int tiilet1, const int tiilet2, const int tiilet3) {

//first define indexes
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i >= tiilet0 || j >= tiilet1) return;

//if we are inside boundaries, let's find the greatest depth value
    unsigned int takana=0;
    unsigned int ddd;
    uchar4 syvyys;
    uchar4 dd;

//there's possibly four different tile sizes to choose between
    if (j!=tiilet1-1 && i!=tiilet0-1) {

    for (int y=j*BLOCK_SIZE; y<(j+1)*BLOCK_SIZE; y++) {
        for (int x=i*BLOCK_SIZE; x<(i+1)*BLOCK_SIZE; x++) {
            dd=tex2D(surfRefSisaan, x, y);
            ddd=(dd.x << 24) | (dd.y << 16) | (dd.z << 8) | (dd.w);
            takana=max(takana, ddd);
        }
    }

    } else if (j==tiilet1-1 && i!=tiilet0-1) {

    for (int y=j*BLOCK_SIZE; y<j*BLOCK_SIZE+tiilet3; y++) {
        for (int x=i*BLOCK_SIZE; x<(i+1)*BLOCK_SIZE; x++) {
            dd=tex2D(surfRefSisaan, x, y);
            ddd=(dd.x << 24) | (dd.y << 16) | (dd.z << 8) | (dd.w);
            takana=max(takana, ddd);
            }
        }

    } else if (j!=tiilet1-1 && i==tiilet0-1) {

    for (int y=j*BLOCK_SIZE; y<(j+1)*BLOCK_SIZE; y++) {
        for (int x=i*BLOCK_SIZE; x<i*BLOCK_SIZE+tiilet2; x++) {
            dd=tex2D(surfRefSisaan, x, y);
            ddd=(dd.x << 24) | (dd.y << 16) | (dd.z << 8) | (dd.w);
            takana=max(takana, ddd);
        }
    }

    } else if (j==tiilet1-1 && i==tiilet0-1) {

    for (int y=j*BLOCK_SIZE; y<j*BLOCK_SIZE+tiilet3; y++) {
        for (int x=i*BLOCK_SIZE; x<i*BLOCK_SIZE+tiilet2; x++) {
            dd=tex2D(surfRefSisaan, x, y);
            ddd=(dd.x << 24) | (dd.y << 16) | (dd.z << 8) | (dd.w);
            takana=max(takana, ddd);
        }
    }
    }

//if there's empty texture, then we choose the maximum possible value
    if (takana==0) {
    takana=1000000000;
    }

//after slicing the greatest 32bit depth value into four 8bit pieces we write the value into another texture
    syvyys.x=(takana & 0xFF000000) >> 24;
    syvyys.y=(takana & 0x00FF0000) >> 16;
    syvyys.z=(takana & 0x0000FF00) >> 8;
    syvyys.w=(takana & 0x000000FF) >> 0;

    surf2Dwrite(syvyys, surfRefUlos, i*sizeof(syvyys), j, cudaBoundaryModeZero);

}

请帮我更快地完成这项工作，我不知道......

【问题讨论】：

提供其他人可以编译和运行的完整程序（minimal reproducible example）。还包括您的时间或性能测量，以及您的平台（GPU、操作系统、CUDA 版本）。
你的程序的目的是什么？
Robert Crovella：我会尽量给出一个完整的，但它很难与不必要的代码分开。这包含一个单独的 dll，我试图摆脱它。 Eric：它简化了 Tiled Forward Engine 的深度缓冲。 talonmies：请告诉更多...
@mamannon：这里没什么好说的——你用谷歌搜索“cuda reduction”，你得到了 390,000 次点击。其中前十个都是优秀的参考。但是看看你的分析器输出，你的内核需要 5ms 才能运行。这真的慢吗？
让你的内核运行得更快不会改善“低计算利用率”的结果。它只会使情况变得更糟。你似乎不明白分析器告诉你什么。您的程序被分析了 0.3 秒，而此时，内核中只花费了 5 毫秒。是什么让您认为让这个内核运行得更快会改善测量结果？它只会让它“更糟”。

标签： c++ cuda bit-shift texture2d

【解决方案1】：

看起来你有一个大小为 2D int 的输入数组

((tiilet0-1)*BLOCK_SIZE+tiilet2, ((tiilet1-1)*BLOCK_SIZE)+tiilet3)

您的每个线程将顺序读取大小输入块中的所有元素

(BLOCK_SIZE, BLOCK_SIZE)

并将每个输入块的最大值写入大小为 2D 的结果数组

(tiilet0, tiilet1)

与合并的内存访问相比，这可能是访问全局内存的最糟糕的方式，即使使用 2D 纹理也是如此。你们很多人都想了解合并的内存访问。

https://devblogs.nvidia.com/parallelforall/how-access-global-memory-efficiently-cuda-c-kernels/

通常你把太多的工作放在一个线程上。鉴于您将 CUDA 线程块映射到输入数组的方式，我想除非您有非常大的输入，否则您的 gridDim 将太小而无法充分利用 GPU。

为了获得更好的性能，您可能希望将每个输入块一个 CUDA 线程更改为每个输入块一个 CUDA 线程块 (int[BLOCK_SIZE][BLOCK_SIZE])，并使用 parallel reduction 来查找块级最大值。

【讨论】：