11

我为康威的生活游戏编写了这个 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。输出从全局内存到主机，然后进行处理。

如您所见，线程处理单个像素。我不确定这是否是最快的方式，因为一些消息来源建议每个线程执行一行或更多。如果我理解正确，NVidia 自己说线程越多越好。我希望有实践经验的人对此提出建议。

4

3 回答 3

11

我的两分钱。

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

使用纹理内存是解决这个问题的好方法，但不一定是最佳方法。

至少，尝试每个线程一次（水平）做 4 个像素。全局内存一次可以访问 128 个字节（只要你有一个 warp 试图访问 128 字节间隔内的任何字节，你还不如在几乎没有额外成本的情况下拉入整个高速缓存行）。由于扭曲是 32 个线程，因此让每个线程在 4 个像素上工作应该是有效的。

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

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

在 2.0 之前的计算兼容性设备上，您需要使用共享内存。在计算兼容性 2.0 和 2.1 的设备上，缓存能力大大提高，因此全局内存可能没问题。

通过确保每个 warp 只访问输入像素的每个水平行中的两个缓存行而不是三个缓存行，可以节省一些重要的节省，就像在每个线程 4 个像素、每个 warp 32 个线程上工作的幼稚实现中会发生的那样。

没有充分的理由使用 float 作为缓冲区类型。您不仅最终获得了四倍的内存带宽，而且代码变得不可靠且容易出错。（例如，您确定它if(neighbors == 3)可以正常工作，因为您正在比较浮点数和整数？）使用无符号字符。更好的是，如果未定义，请使用 uint8_t 和 typedef 来表示 unsigned char。

最后，不要低估实验的价值。很多时候 CUDA 代码的性能不能用逻辑来简单地解释，你必须求助于调整参数并看看会发生什么。

于 2011-01-03T08:52:24.470 回答
4

TL；DR：见： http: //golly.sourceforge.net

问题是大多数 CUDA 实现都遵循手动计算邻居的脑死亡想法。这太慢了，以至于任何智能串行 CPU 实现都将胜过它。

进行 GoL 计算的唯一明智方法是使用查找表。
CPU 上当前最快的实现使用查找方形 4x4 = 16 位块来查看未来的 2x2 单元。

在此设置中，单元格的布局如下：

 01234567
0xxxxxxxx //byte0
1xxxxxxxx //byte1 
2  etc
3
4
5
6
7

使用一些位移来使 4x4 块适合单词，并使用查找表查找该单词。查找表也包含单词，这样可以将 4 个不同版本的结果存储在查找表中，因此您可以最大限度地减少需要对输入和/或输出进行的位移量。

此外，不同代是交错的，因此您只需查看 4 个相邻的平板，而不是 9 个。像这样：

AAAAAAAA 
AAAAAAAA   BBBBBBBB
AAAAAAAA   BBBBBBBB
AAAAAAAA   BBBBBBBB
AAAAAAAA   BBBBBBBB
AAAAAAAA   BBBBBBBB
AAAAAAAA   BBBBBBBB
AAAAAAAA   BBBBBBBB
           BBBBBBBB
//odd generations (A) are 1 pixel above and to the right of B,
//even generations (B) are 1 pixels below and to the left of A.

与愚蠢的计数实现相比，仅此一项就可以提高 1000 倍以上的速度。

然后是不计算静态或周期性为 2 的板的优化。

然后是HashLife，但那是完全不同的野兽。
HashLife 可以在 O(log n) 时间内生成生命模式，而不是在 O(n) 时间内正常实现。这使您可以在几秒钟内计算生成：6,366,548,773,467,669,985,195,496,000（6 octillion）。
不幸的是，Hashlife 需要递归，因此在 CUDA 上很难。

于 2016-02-06T19:31:08.623 回答
3

看看这个线程，我们在那里做了很多改进......

http://forums.nvidia.com/index.php?showtopic=152757&st=60

于 2011-01-09T21:46:18.643 回答