如何优化Conway对CUDA的生活游戏？

Question

我为Conway的生活游戏编写了这个CUDA内核:

__global__ void gameOfLife(float* returnBuffer, int width, int height) {  
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;  
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;  
    float p = tex2D(inputTex, x, y);  
    float neighbors = 0;  
    neighbors += tex2D(inputTex, x+1, y);  
    neighbors += tex2D(inputTex, x-1, y);  
    neighbors += tex2D(inputTex, x, y+1);  
    neighbors += tex2D(inputTex, x, y-1);  
    neighbors += tex2D(inputTex, x+1, y+1);  
    neighbors += tex2D(inputTex, x-1, y-1);  
    neighbors += tex2D(inputTex, x-1, y+1);  
    neighbors += tex2D(inputTex, x+1, y-1);  
    __syncthreads();  
    float final = 0;  
    if(neighbors < 2) final = 0;  
    else if(neighbors > 3) final = 0;  
    else if(p != 0) final = 1;  
    else if(neighbors == 3) final = 1;  
    __syncthreads();  
    returnBuffer[x + y*width] = final;  
}

我正在寻找错误/优化.并行编程对我来说很新,我不确定我是否能够正确地完成它.

其余的是从输入数组到绑定到CUDA数组的2D纹理inputTex的memcpy.输出从全局内存到主机进行memcpy-ed然后处理.

正如您所看到的,线程处理单个像素.我不确定这是否是最快的方式,因为一些消息来源建议每个线程执行一行或更多.如果我理解正确NVidia自己说越多线程越好.我很乐意从有实际经验的人那里得到建议.

Answer 1

我的两分钱.

整个事情看起来很可能受到多处理器和GPU内存之间通信延迟的限制.您的代码应该采用30-50个时钟周期来自行执行,并且它生成至少3个内存访问,如果必需数据不在缓存中,则每个内存访问需要200多个时钟周期.

使用纹理内存是解决这个问题的好方法,但它不一定是最佳方式.

至少,尝试每个线程一次(水平)做4个像素.全局内存一次可以访问128个字节(只要你有一个warp试图访问128字节间隔内的任何字节,你也可以在几乎没有额外成本的情况下拉入整个缓存行).由于warp是32个线程,每个线程在4个像素上工作应该是有效的.

此外,您希望由同一个多处理器处理垂直相邻的像素.原因是相邻行共享相同的输入数据.如果由一个MP处理像素(x = 0,y = 0)并且像素(x = 0,y = 1)由不同的MP处理,则两个MP必须分别发出三个全局存储器请求.如果它们都由同一个MP处理并且结果被正确缓存(隐式或显式),则您只需要总共四个.这可以通过让每个线程在几个垂直像素上工作,或者让blockDim.y> 1来完成.

更一般地说,您可能希望每个32线程warp加载与MP上可用的内存(16-48 kb,或至少128x128块)一样多,然后处理该窗口内的所有像素.

在2.0之前的计算兼容性设备上,您将需要使用共享内存.在计算兼容性2.0和2.1的设备上,缓存功能得到了很大改善,因此全局内存可能很好.

通过确保每个warp只访问输入像素的每个水平行中的两个缓存行而不是三个缓存行,可以获得一些重要的节省,就像在每个线程4个像素,每个warp 32个线程的朴素实现中一样.

使用float作为缓冲区类型是没有充分理由的.您不仅最终获得了四倍的内存带宽,而且代码变得不可靠并且容易出错.(例如,你确定它if(neighbors == 3)正常工作,因为你要比较一个浮点数和一个整数吗？)使用unsigned char.更好的是,使用uint8_t并将其定义为无符号字符(如果未定义).

最后,不要低估实验的价值.很多时候,CUDA代码的性能很难通过逻辑来解释,你不得不求助于调整参数并看看会发生什么.