增加块大小会降低性能

在我的cuda代码中，如果我增加blocksizeX，blocksizeY它实际上需要更多时间[因此，我以1x1运行它]另外还有一大块执行时间（例如，9个中的7个s）只是通过调用内核来实现的。事实上，我很惊讶即使我将整个内核注释掉，时间也几乎是一样的。有什么建议在哪里以及如何进行优化？增加块大小会降低性能

P.S.我已经用我的实际代码编辑了这篇文章。我对图像进行了下采样，因此每4个相邻像素（例如，对于例如1,2行和1,2行）给出一个输出像素。我得到一个有效的bw。 5GB/s，理论最大值为86.4 GB/s。我使用的时间是在调用内核和指令并调用一个空内核方面的差异。我现在看起来很糟糕，但我无法弄清楚我做错了什么。

__global__ void streamkernel(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; 
    int number=2*(id%(width/2))+(id/(width/2))*width*2; 

    if (id<height*width/4) 
    { 

     f_r[id]=(r_d[number]+r_d[number+1];+r_d[number+width];+r_d[number+width+1];)/4;        
     f_g[id]=(g_d[number]+g_d[number+1]+g_d[number+width]+g_d[number+width+1])/4;    
     f_b[id]=(g_d[number]+g_d[number+1]+g_d[number+width]+g_d[number+width+1];)/4; 
    } 


    }

来源

2011-02-19 Manish

谢谢！但我做了谷歌，并做了一些功课，然后发布在堆栈溢出。 – Manish 2011-02-19 06:41:07

@Nick：[LMGTFY网址不允许出于某种原因]（http://meta.stackexchange.com/questions/15650/ban-lmgtfy-let-me-google-that-for-you-links）。你会知道，如果你没有试图用tinyurl来混淆它，这也是*强烈不鼓励。我喜欢知道我被链接到哪里。 – 2011-02-19 07:02:53

@Manish - 我想帮助你，但是你对前面提到的关于cuda的两个问题没有选择正确的答案。如果你给我们更多的激励，这将有所帮助。 – jmilloy 2011-02-19 07:38:11

您忘记了一个多处理器可以同时执行多达8个块的事实，并且当时达到最高性能。然而，有许多因素限制块中的可以同时存在（不完全列表）的数量：

每个多处理器共享存储器的最大量限制如果#blocks *每块共享存储器将是块的数量>总共享内存。
如果#blocks * #threads/block>> max total #threads，则每个多处理器的最大线程数限制块的数量。
...

你应该尝试找到引起正好是8块要在一个多处理器上运行内核执行配置。即使入住率=/= 1.0，这几乎总是会产生最高的性能！从这一点开始，您可以尝试迭代地进行更改，以减少每个MP执行块的数量，但是因此会增加内核的占用量并查看性能是否提高。

nvidia occupancy calculator(excel sheet)会有很大的帮助。

来源

2011-02-20 03:29:35

尝试在CUDA SDK示例中查找矩阵乘法示例以了解如何使用共享内存。

当前内核的问题是，它为每3个添加和1个分区写入4个全局内存读取和1个全局内存写入。每次全局内存访问大约需要400个周期。这意味着你花费绝大多数时间进行内存访问（GPU不好），而不是计算（GPU优秀）。

共享内存有效地允许您缓存这个以便分期付款，您可以在每个像素上进行大约1次读取和1次写入，以进行3次添加和1次分割。对于CGMA比率（计算到全球存储器访问比率，GPU计算的圣杯），这仍然没那么好。总的来说，我认为对于这样一个简单的内核来说，考虑到通过PCI-E总线传输数据的开销，CPU实现可能会更快。

来源

2011-02-24 14:09:36 peakxu

增加块大小会降低性能

回答

相关问题