使用分支优化 CUDA 代码

Question

我有一个 CUDA 代码，我想对其进行优化。我的内核正在使用 dim3 grid=(35,48) 和 dim3 threads=(18,18)。首先，每个块执行独立的290次向量计算，其中每个线程执行1次向量计算（即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 个线程的块，编号从 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) 不采用分支（见下图）。然而，在一个块的跨度内，确实有 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 个线程就有一个非活动线程。换句话说，这意味着所有你的扭曲发散。

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

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

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

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

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

Answer 2

按照我的理解，如果要对分支进行优化，必须提前对数据进行处理（即那些位于18号的数据要提前聚簇在一起，在原位置删除）。

Answer 3

This 是一篇简短的博客，非常清楚地解释了 分支分歧 问题。

正常情况下，仅存在经内发散，无经间发散。

使用分支优化 CUDA 代码

optimizing CUDA code with branches

c++

cuda

branch