使用分支优化 CUDA 代码答案

【问题标题】：optimizing CUDA code with branches使用分支优化 CUDA 代码
【发布时间】：2015-03-21 22:43:50
【问题描述】：

我有一个 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 进行编码。

谢谢，米哈伊尔

【问题讨论】：

您是否分析了您的代码以测量经纱发散的实际数量？
你是说NSight？不，由于某种原因，我不能强迫它工作，但这是另一回事。
如果你想提高你的代码的性能，你应该首先衡量你的代码的性能，确定瓶颈，如果确实是瓶颈，才考虑发散。
根据你的问题，18 x 18 线程块，else if 语句将永远不会被执行。这段代码没有分歧。
共享内存访问对访问哪些银行非常敏感。最好是每个命令的每个扭曲只访问每个银行一次。（除非它是相同的内存位置）。但是如果不知道向量计算的内存位置是如何安排的，就无法说出更多。

标签： c++ cuda branch

【解决方案1】：

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

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

如果“前 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 个线程就有一个非活动线程。换句话说，这意味着所有你的经线都发散了。

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

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

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

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

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

【讨论】：

谢谢，这很有帮助。但是你能解释一下，我怎样才能避免分支。所以，让我们考虑一个简单的任务：我们有 2 个线程，两个输入数组（a，b），每个线程计算每个数组元素的总和。如果没有 if 语句，我该如何实现？
我也不明白，二维网格中的扭曲是如何组织的？在我的问题中，让我们通过 2D inder a_{ij} 枚举线程，其中 i,j=0,1,2,...,17。那么哪些线程属于，比如说，第一个经纱？
我假设您的意思是“扭曲如何映射到 2D 线程块的线程”？如果是这样，请查看 CUDA 文档的this section。
在 if 块之外进行计算之前，您需要一个简短的 if 语句来分配指针。

【解决方案2】：

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

【讨论】：

【解决方案3】：

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

通常，只有经线内发散，没有经线间发散。

【讨论】：