CUDA版本比CPU版本慢？答案

【问题标题】：CUDA version slower than CPU version?CUDA版本比CPU版本慢？
【发布时间】：2012-04-15 07:09:47
【问题描述】：

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

P.S 我尝试只调用内核（删除了代码），它与带代码的内核几乎相同。此外，我的代码不带内核调用运行大约 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 调用内核。

【问题讨论】：

多少线程/块？为什么不指定线程数以便摆脱 if() ？
我在上面编辑了线程/块。我不确定如何摆脱“if”以及它是否会损害性能。（因为我测量了移除该部分并调用空内核的性能，并且几乎需要相同的时间）
好吧，我还运行了另一个简单的程序来添加 2 个数组，因此 CPU 版本的线程执行速度似乎比 GPU 版本快。

标签： cuda

【解决方案1】：

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

【讨论】：

至于内存复制，我使用手册中提到的零复制。所以我不确定是否有任何转移发生。我还编辑了上面的帖子以包含我的代码。跨度>
即使是零拷贝，GPU 卡和主板之间仍然存在 I/O - 您需要确定是 I/O 受限还是计算受限。

【解决方案2】：

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

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

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

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

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

首先要做的是考虑内存布局。您需要为目标图像中的每个像素创建一个线程。鉴于二次采样的内存访问模式未合并，您需要将像素数据读入共享内存。考虑 32x8 线程的块大小。这允许每个块以 40*8*4 像素或 3bpp 读取 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 个字节。如果 32byte 索引与 warp 中的线程索引匹配，则所有 32 次访问都被放入一个 128 事务中。这就是您获得 GPU 100GB/s 带宽的方式。

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

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

【讨论】：