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

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

Re: 线程分化-以warp为分支单位
« 回复 #1 于: 五月 09, 2022, 11:18:22 am »
或者怎么看作用于if的和作用于else的不同warp是否同时执行了呢

Re: 线程分化-以warp为分支单位
« 回复 #2 于: 五月 19, 2022, 07:08:25 pm »
对开发者而言,应该尽量避免在同一个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的你还得等)。

Re: 线程分化-以warp为分支单位
« 回复 #3 于: 五月 25, 2022, 09:09:00 am »
感谢您的解答。
请问您说的异步载入是指这个吧?
程序代码: [选择]
__pipeline_memcpy_async(&shared[blockDim.x * i + threadIdx.x],
                            &global[blockDim.x * i + threadIdx.x], sizeof(T));
这个代码的效果是不是和我上述使用if~else效果相同,只是异步载入使用起来会更方便?

Re: 线程分化-以warp为分支单位
« 回复 #4 于: 五月 25, 2022, 09:35:56 am »
感谢您的解答。
请问您说的异步载入是指这个吧?
程序代码: [选择]
__pipeline_memcpy_async(&shared[blockDim.x * i + threadIdx.x],
                            &global[blockDim.x * i + threadIdx.x], sizeof(T));
这个代码的效果是不是和我上述使用if~else效果相同,只是异步载入使用起来会更方便?
但是我又有一个疑问,如果使用异步载入的话,执行计算是一个核函数,进行异步载入又是一个核函数,可是在单流中,两个核函数只能串行进行,那就达不到计算和搬运数据同时进行的效果了。

Re: 线程分化-以warp为分支单位
« 回复 #5 于: 六月 06, 2022, 04:54:32 pm »
感谢您的解答。
请问您说的异步载入是指这个吧?
程序代码: [选择]
__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的,不会的。这点手册上已经说的很明白了,就不重复说了。你过度担忧了。