如何正确使用__shfl_up_sync在少数通道之间进行数据传输

Question

假设我需要将一些值从29泳道洗牌到30泳道，也从30泳道洗牌到31泳道；我该用什么面膜？

当然mask应该覆盖lane 30, 31AND29（cuda指南说源lane也需要参与）；但随后必须在通道 29 上调用该命令 - 那么我是否需要不断扩展掩码（在本例中最终扩展为完整掩码），或者我是否只使用 0xE0000000 并让未定义的结果返回到通道 29 （无论如何它不需要结果）？

Answer 1

好的，让我们列出您的要求

29->29  //dest 29, offset = 0
29->30  //dest 30, offset = 1
30->31  //dest 31, offset = 1

代码将如下所示：

#include <stdio.h>
#include <cuda.h>


__global__ void shfl_29() {
    auto dummydata = threadIdx.x % 32;
    const auto activemask = 0b111u << 29;
    const auto laneid = threadIdx.x % 32;
    const auto offset = int(laneid > 29); //29 = 0, else 1
    if (laneid >= 29) {
        dummydata = __shfl_up_sync(activemask, dummydata, offset);
        printf("old: %i, new: %i\n", laneid, dummydata);
    }
    
}

int main() {
    shfl_29<<<1,256>>>();
    cudaDeviceSynchronize();   
}

您可以在 Godbolt 上看到它的运行情况：https://cuda.godbolt.org/z/6E9bGbx1f