CUDA版本比CPU版本慢？

Question

我正在 CUDA 中编写一个图像子采样器，并使用线程执行平均操作。但是，如果我在不调用内核的情况下执行此操作，那么与我实际调用 CUDA 内核时相比，它的运行速度要快得多。现在的图像大小是 1280x1024。 内核调用通常需要花费大量时间还是我的实现有问题？

PS 我尝试仅调用内核（删除了代码），并且它与带有代码的内核几乎相同。此外，我的代码（不带内核调用）运行大约 350 毫秒，而带内核调用运行接近 1000 毫秒。

__global__ void subsampler(int *r_d,int *g_d,int *b_d, int height,int width,int *f_r,int*f_g,int*f_b){ 
        int id=blockIdx.x * blockDim.x*blockDim.y+ threadIdx.y*blockDim.x+threadIdx.x+blockIdx.y*gridDim.x*blockDim.x*blockDim.y;
        if (id<height*width/4){
        f_r[id]=(r_d[4*id]+r_d[4*id+1]+r_d[4*id+2]+r_d[4*id+3])/4;
        f_g[id]=(g_d[4*id]+g_d[4*id+1]+g_d[4*id+2]+g_d[4*id+3])/4;
        f_b[id]=(b_d[4*id]+b_d[4*id+1]+b_d[4*id+2]+b_d[4*id+3])/4;
        }
        }

我将 blockSizeX 和 blockSizeY 定义为 1 和 1（我尝试将它们设为 4,16），但不知怎的，这是最快的

 dim3 blockSize(blocksizeX,blocksizeY);
  int new_width=img_width/2;
  int new_height=img_height/2;

  int n_blocks_x=new_width/blocksizeX+(new_width/blocksizeY == 0 ?0:1);
  int n_blocks_y=new_height/blocksizeX+(new_height/blocksizeY == 0 ?0:1);
  dim3 gridSize(n_blocks_x,n_blocks_y);

，然后我用 gridSize,BlockSize 调用内核。

Answer 1

可能是内核实现得不是很好，或者可能是从 GPU 卡移入和移出数据的开销淹没了任何计算优势。尝试单独对内核进行基准测试（没有 CPU <-> GPU 内存传输），看看内核占用了总时间，内存传输占用了多少时间。然后，您可以根据这些测量结果决定是否需要在内核上做更多工作。

Answer 2

虽然我不确定您运行的是什么硬件，但您应该能够使该内核的执行速度接近 1000 fps，而不是 1000 毫秒/帧:)

建议 1：如果此处理通过 OpenGL 与可视化有任何交互/DirectX 或类似的，只需将其作为着色器执行即可 - 网格/块大小、内存布局等所有细节都会为您处理。如果您确实需要自己在 CUDA 中实现此功能，请继续阅读：

首先，我假设您在每个方向上对 1280x1024 图像进行 2 倍子采样，从而生成 640x512 图像。生成图像中的每个像素都是源图像中四个像素的平均值。图像具有三个通道：RGB。

问题 1：您真的想要每通道 32 位还是想要 RGB888（每通道 8 位）？ RGB888 相当常见——我假设这就是你的意思。

问题 2：您的数据实际上是平面的，还是从交错格式中提取的？ RGB888 是一种交错格式，其中像素在内存中显示为 RGBRGBRGB。我会编写你的内核来以其本机格式处理图像。我假设您的数据实际上是平面的，因此您有三个平面：R8、G8 和 B8。

首先要做的是考虑内存布局。您需要为目标图像中的每个像素分配一个线程。鉴于子采样的内存访问模式未合并，您将需要将像素数据读入共享内存。考虑 32x8 线程的块大小。这允许每个块以 3bpp 读取 40*8*4 像素，即 3072 字节。实际上，您将读入稍微多一点的内容，以保持负载合并，每个块总共 4096 字节。现在给你：

dim3 block(32, 8);
dim3 grid(1280 / 2 / 32, 1024 / 2 / 8); // 20x64 blocks of 256 threads

现在有趣的部分来了：共享内存。你的内核可能看起来像这样：

__global__ void subsample(uchar* r, uchar* g, uchar* b,    // in
                          uchar* ro, uchar* go, uchar* bo) // out
{
    /* Global offset into output pixel arrays */
    int gid = blockIdx.y * gridDim.x * blockDim.x + blockIdx.x * blockDim.x;

    /* Global offset into input pixel arrays */
    int gidin = gid * 2;

    __shared__ uchar* rc[1024];
    __shared__ uchar* gc[1024];
    __shared__ uchar* bc[1024];

    /* Read r, g, and b, into shmem cache */
    ((int*)rc)[threadIdx.x] = ((int*)r)[gidin + threadIdx.x];
    ((int*)gc)[threadIdx.x] = ((int*)g)[gidin + threadIdx.x];
    ((int*)bc)[threadIdx.x] = ((int*)b)[gidin + threadIdx.x];

    __syncthreads();

    /* Shared memory for output */
    __shared__ uchar* roc[256];
    __shared__ uchar* goc[256];
    __shared__ uchar* boc[256];

    /* Do the subsampling, one pixel per thread. Store into the output shared memory */

    ...

    __syncthreads();

    /* Finally, write the result to global memory with coalesced stores */
    if (threadIdx.x < 64) {
        ((int*)ro)[gid + threadIdx.x] =  ((int*)roc)[threadIdx.x];
    } else if (threadIdx.x < 128) {
        ((int*)go)[gid + threadIdx.x-64] =  ((int*)goc)[threadIdx.x-64];
    } else if (threadIdx.x < 192) {
        ((int*)bo)[gid + threadIdx.x-128] =  ((int*)boc)[threadIdx.x-128];
    }
}

哇！那里有很多东西，抱歉代码转储。需要记住的一些原则：

1) 当您使用合并加载/存储时，内存速度会很快。这意味着对于 32 个 warp 中的每个线程，每个线程访问 32 个字节。如果 32 字节索引与 warp 中的线程索引匹配，则所有 32 次访问都会放入一个 128 事务中。这就是获得 GPU 100GB/s 带宽的方法。

2) 进行子采样时的内存访问模式不会合并，因为它依赖于原始内存所不具备的二维空间局部性。 （也可以为此使用纹理内存...）通过将输入存储在共享内存中，然后进行处理，可以最大限度地减少对计算性能的影响。

我希望这会有所帮助——如果您愿意，我可以回复有关某些部分的更多详细信息。