CUDA版本比CPU版本慢吗？

Question

我正在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内核。

Answer 1

这可能是因为内核执行得不好，或者可能是将数据移入或移出GPU卡的开销正在减少任何计算上的好处。尝试单独对内核进行基准测试（无需CPU <-> GPU内存传输），查看内核占用您的总时间以及内存传输的总时间。然后，您可以根据这些测量结果来决定是否需要在内核上做更多的工作。

Answer 2

虽然我不知道你的硬件正在运行这一个，你应该能够让这个内核接近1000 fps的执行，而不是1000毫秒/帧:)

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

首先，我假设你在每个方向上对1280x1024图像进行二次采样，得到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个字节。如果32字节索引与warp中的线索索引相匹配，则所有32个访问都被放入一个128个事务中。这就是你如何获得GPU的100GB/s带宽。

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

我希望这会有所帮助 - 如果您愿意，我可以在某些部分回复更多细节。