线程分化-以warp为分支单位

对开发者而言，应该尽量避免在同一个Warp中有不同的执行路径。当必须写分支语句时，尽量让分支宽度大于WarpSize，即不以Thread为分支单位，而是以Warp为分支单位，这样就能保证同一个Warp内部不会出现分支。
请问如果分配线程块 block(256)
核函数中是
int warpIdx = threadIdx.x/32;
if(warpIdx < 2){
    执行语句1；
}else{
    执行语句2；
}代码是以warp为分支单位，那这种情况下if和else可以同时运行吗？
即现在有一个需求，想让线程块中一部分线程来执行计算，另一部分线程来把数据从全局内存搬运到共享内存，请问我可以通过上述if-else达到这两种需求同时运行的效果吗？

或者怎么看作用于if的和作用于else的不同warp是否同时执行了呢

对开发者而言，应该尽量避免在同一个Warp中有不同的执行路径。当必须写分支语句时，尽量让分支宽度大于WarpSize，即不以Thread为分支单位，而是以Warp为分支单位，这样就能保证同一个Warp内部不会出现分支。
请问如果分配线程块 block(256)
核函数中是

程序代码: [选择]

int warpIdx = threadIdx.x/32;
if(warpIdx < 2){
    执行语句1；
}else{
    执行语句2；
}代码是以warp为分支单位，那这种情况下if和else可以同时运行吗？
即现在有一个需求，想让线程块中一部分线程来执行计算，另一部分线程来把数据从全局内存搬运到共享内存，请问我可以通过上述if-else达到这两种需求同时运行的效果吗？

关于你的几个问题：
（1）“当必须写分支语句时，尽量让分支宽度大于WarpSize”---我猜测你的正确意思应该是分支宽度是warpSize的倍数。
（2）“那这种情况下if和else可以同时运行吗”--如果编译器生成了常规的跳转指令，同时2个warp执行到了不同分支的位置，则在SM有多个scheduler的情况下，的确2个warp可能在同时执行这2个不同的代码路径。
（3）“即现在有一个需求，想让线程块中一部分线程来执行计算，另一部分线程来把数据从全局内存搬运到共享内存，请问我可以通过上述if-else达到这两种需求同时运行的效果吗？”---这个的确可以，前人有一篇老文章，讲述了通过"warp specilization来执行特殊的载入任务，一个warp为其他warp服务“（大致意思，原文记不清了，你可以搜搜对应的英文原文”。

不过考虑到现在已经有了异步载入指令（例如你的8.6计算能力的RTX30卡），和最新9.0计算能力的内嵌在SM内部的微型DMA引擎，都可以完成你设想的载入global->shared的任务，你的独立使用1个warp来进行的主意，可能没有太大的用途。我建议你使用异步载入，手册有介绍（针对8.6的，9.X的你还得等）。

感谢您的解答。
请问您说的异步载入是指这个吧？
__pipeline_memcpy_async(&shared[blockDim.x * i + threadIdx.x],
                            &global[blockDim.x * i + threadIdx.x], sizeof(T));这个代码的效果是不是和我上述使用if~else效果相同，只是异步载入使用起来会更方便？

感谢您的解答。
请问您说的异步载入是指这个吧？

程序代码: [选择]

__pipeline_memcpy_async(&shared[blockDim.x * i + threadIdx.x],
                            &global[blockDim.x * i + threadIdx.x], sizeof(T));这个代码的效果是不是和我上述使用if~else效果相同，只是异步载入使用起来会更方便？

但是我又有一个疑问，如果使用异步载入的话，执行计算是一个核函数，进行异步载入又是一个核函数，可是在单流中，两个核函数只能串行进行，那就达不到计算和搬运数据同时进行的效果了。

感谢您的解答。
请问您说的异步载入是指这个吧？

程序代码: [选择]

__pipeline_memcpy_async(&shared[blockDim.x * i + threadIdx.x],
                            &global[blockDim.x * i + threadIdx.x], sizeof(T));这个代码的效果是不是和我上述使用if~else效果相同，只是异步载入使用起来会更方便？

但是我又有一个疑问，如果使用异步载入的话，执行计算是一个核函数，进行异步载入又是一个核函数，可是在单流中，两个核函数只能串行进行，那就达不到计算和搬运数据同时进行的效果了。

首先说：
（1）和你用if...else, 在warp的边界上，不同的warp干不同的活（例如一个warp计算，一个warp载入）不同的。异步载入shared memory在载入命令发出后，warp可以继续往下做其他事情，而无需卡住。这样所有的warps都可以发出大量的进行中的load请求，也都可以计算一些不依赖的其他事情。和单独选出1个warp单独干载入（以前有篇文章专门讲这个），可能效果并不太一样。

其次：
（2）异步载入并不会产生一个新的kernel的，不会的。这点手册上已经说的很明白了，就不重复说了。你过度担忧了。