c++ - 使用分支优化 CUDA 代码

Question

我有一个 CUDA 代码，我想对其进行优化。我的内核正在使用 dim3 grid=(35,48) 和 dim3 threads=(18,18)。首先，每个 block 执行独立的290次 vector 计算，其中每个线程执行1次 vector 计算(即1024次加法-乘法)。

但是，此计算的前 17*17=289 个输入数据存储在共享数组 im1 中，最后一个数据存储在 im2 中(输出数组也不同)。之后，我使用所有获得的数据进行进一步的计算。

我是这样实现的:

if ((threadIdx.x < 17) && (threadIdx.y < 17)){
    **instructions for 289s vector calculations**
}
else if ((threadIdx.x == 17) && (threadIdx.y == 17)){
    **instruction for 290 vector calculation**
}
__syncthreads();
***further calculations***

所以，如果我理解正确的话，我的第一个 289 跟随 1 个分支，线程 #324 跟随另一个。只要第一组线程在 warp #0,1,..,10 中，并且线程 #324 在 warp #11 中，就没有不同的分支。然而，我读到，通常最好避免在此类内核中使用任何 if 语句，并将它们替换为跨步索引或类似的东西。那么，我能以某种方式改进这段代码吗？

我的 GPU 是 GTX 980 和 cc 5.2，我使用 VS2013 进行编码。

谢谢，米哈伊尔

Answer 1

让我们考虑一个包含 18 * 18 个线程的 block ，编号从 0 (0, 0) 到 323 (17, 17)。

So, if I understand correct, my first 289 follow 1 branch [...]

如果“first 289”指的是从 0 (0, 0) 到 288 (16, 16) 编号的线程，那么不，并非所有线程都采用第一个分支。例如，线程 17 (0, 17) 不采用分支(见下图)。然而，在一个 block 的跨度内，确实有 289 个线程采用该分支。

[...] and thread #324 follows another

没错，线程323(17, 17)走的是第二个分支。

线程 17 (0, 17), 35 (1, 17) ... 305 (16, 17) 和 306 (17, 0), 307 (17, 1) ... 322 (17, 16) (总共 35 个线程)不采取任何分支并被浪费。从性能的角度来看，这很糟糕，但也不是真正的灾难。

但是考虑一下您正在做的事情的以下模式:

    0  1  2  … 15 16 17     
0   *  *  *  *  *  *  -      * represents a thread that takes branch 1
1   *  *  *  *  *  *  -      X represents a thread that takes branch 2
2   *  *  *  *  *  *  -      - represents a thread that takes no branch
…   *  *  *  *  *  *  -
15  *  *  *  *  *  *  -
16  *  *  *  *  *  *  -
17  -  -  -  -  -  -  X

请记住，warp 由 32 个线程组成。因此线程 0..31、32..63 等以锁步方式执行。正如您在上面的模式中注意到的那样，每 18 个线程就有一个非事件线程。换句话说，这意味着你的所有经线发散。

不过，这可能不会对性能造成巨大影响(如果有的话)，因为其中一个分支总是“什么都不做”。话虽这么说，我肯定会鼓励您修复您的设计，并且我相信您会注意到性能改进(不过更多是由于内存访问模式而不是分歧本身)。

一个明显的解决方案是只启动 290 个线程而不是 324 个线程，然后自己映射到 x 和 y 坐标，但是最后的扭曲会以明显的方式发散。

另一种解决方案是启动足够多的 warp 来覆盖前 289 个线程(这意味着 10 个 warp 而最后一个浪费 31 个线程)并运行一个补充 warp，其中您为第二个分支(最后一个，例如)。所以这将是 11 个扭曲，352 个线程，62 个浪费。就效率而言，这可能看起来更差，但由于内存访问模式，它实际上比那更复杂，所以试一试吧。

另请注意，如果 if/else 语句的主体实际上在代码上没有不同，但在数据上有所不同(正如您似乎暗示的那样......)，那么使用分支毫无意义。只是玩指针。可能会出现其他问题(与内存访问合并相关)，但不会出现代码流分歧。

我会建议更多的改进，但没有看到您的代码或不知道您的数据是如何布局的，这有点像在黑暗中拍摄。你在评论中说你无法让 NSIGHT 工作:我强烈建议你将其作为优先事项。