我有一个CUDA代码，我想对其进行优化。我的内核正在使用dim3 grid=(35,48)和dim3 threads=(18,18)。首先，每个块执行独立的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的线程将跟随另一个分支。只要第一组线程在经线＃0,1，..，10中，并且线程＃324在经线＃11中，就不会有分支分支。但是，我读到，通常最好避免在此类内核中使用任何
if语句，并使用strided index或类似的索引替换它们。因此，我可以以某种方式改进此代码吗？我的GPU是带有cc 5.2的GTX 980，我使用VS2013进行编码。
谢谢，米哈伊尔
最佳答案
让我们考虑一个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
请记住，经线由32个线程组成。因此，线程0..31、32..63等以锁定步骤执行。您可能在上面的架构中注意到，每18个线程中只有一个不 Activity 线程。换句话说，这意味着所有您的经线会发散。
但是，这可能不会对性能造成巨大的影响(如果有的话)，因为分支之一始终是“不执行任何操作”。话虽这么说，我绝对会鼓励您修复您的设计，并且我相信您会注意到性能的提高(尽管更多的是由于内存访问模式而不是分歧本身)。
一个明显的解决方案是仅启动290个线程而不是324个线程，并自己映射到x和y，但是最后一次扭曲将以明显的方式发散。
另一种解决方案是启动足够的经线以覆盖前289个线程(这意味着10个经线，最后一个浪费31个线程)，并运行补充经线，其中将线程用于第二个分支(例如，最后一个) 。这样一来，将浪费11个经纱，352个线程，浪费62个。就效率而言，这似乎更糟，但实际上由于内存访问模式而比它更复杂，因此请尝试一下。
还要注意，如果
if/else语句的主体实际上在代码上并没有在数据上有所不同(如您所暗示的那样...)，那么使用分支是没有意义的。只是玩指针。可能会出现其他问题(与内存访问合并有关)，但是不会出现代码流差异。我建议您进行更多改进，但是如果您看不到您的代码或不知道数据的布局方式，那简直就是一片黑暗。您在评论中说，您不能让NSIGHT起作用:我强烈建议您将其作为优先事项。
关于c++ - 用分支优化CUDA代码，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/28100671/