__shfl_up_sync 调用中的掩码是否自适应？答案

【问题标题】：Is mask adaptive in __shfl_up_sync call?__shfl_up_sync 调用中的掩码是否自适应？
【发布时间】：2020-03-11 03:45:38
【问题描述】：

基本上，它是this post 的物化版本。假设一个 warp 需要处理 4 个对象（例如，图像中的像素），每 8 个通道组合在一起以处理一个对象：现在我需要在处理一个对象期间（即在该对象的 8 个通道中）进行内部 shuffle 操作，它适用于每个对象，只需将 mask 设置为 0xff：

uint32_t mask = 0xff;
__shfl_up_sync(mask,val,1);

但是，据我了解，将mask 设置为0xff 将强制对象0（或对象3？也停留在这一点上）的lane0:lane7 参与，但我确保上述用法适用于每个对象之后大量的试验。那么，我的问题是__shfl_up_sync 调用是否可以调整参数mask 来强制相应的车道参与？

更新
实际上，这个问题来自我试图解析的libSGM 的代码。特别是，它以相当并行的方式解决了dynamic programming 的最小成本路径。一旦程序在使用执行配置启动内核aggregate_vertical_path_kernel 后到达此line：

//MAX_DISPARITY is 128 and BLOCK_SIZE is 256
//Basically, each block serves to process 32 pixels in which each warp serves to process 4.
const int gdim = (width + PATHS_PER_BLOCK - 1) / PATHS_PER_BLOCK;
const int bdim = BLOCK_SIZE;
aggregate_vertical_path_kernel<1, MAX_DISPARITY><<<gdim, bdim, 0, stream>>>(...)

一个对象dp从DynamicProgramming<DP_BLOCK_SIZE, SUBGROUP_SIZE>实例化：

static constexpr unsigned int DP_BLOCK_SIZE = 16u;
...
//MAX_DISPARITY is 128
static const unsigned int SUBGROUP_SIZE = MAX_DISPARITY / DP_BLOCK_SIZE;
...
DynamicProgramming<DP_BLOCK_SIZE, SUBGROUP_SIZE> dp;

继续跟随程序，dp.updata() 将被调用，其中__shfl_up_sync 用于访问前面DP_BLOCK 的最后一个元素，__shfl_down_sync 用于访问后面DP_BLOCK 的第一个元素。此外，一个warp中的每8个lane被分组在一起：

//So each 8 threads are grouped together to process one pixel in which each lane is contributed to one DP_BLOCK for corresponding pixel.
const unsigned int lane_id = threadIdx.x % SUBGROUP_SIZE;

它来了，一旦程序到达这个line：

//mask is specified as 0xff(255)
const uint32_t prev =__shfl_up_sync(mask, dp[DP_BLOCK_SIZE - 1], 1);

一个经线中的每个通道确实使用相同的掩码 0xff 进行洗牌，这导致了我的上述问题。

【问题讨论】：

我不相信“一个经线中的每个通道都使用相同的掩码 0xff 进行洗牌”。是一个真实的陈述。整个扭曲的蒙版似乎不是相同的值。见here。 group_id 在整个经线上发生变化。
好的，我明白了...

标签： cuda shuffle intrinsics

【解决方案1】：

当你这样做时它会令人困惑：

lane0:lane7 | lane0:lane7 | lane0:lane7 | lane0:lane7

因为 warp 没有 4 组车道，编号为 0 到 7 号车道。它有一组车道，编号为 0 到 31 号车道。

lane 31 | lane 30 | ... | lane 0

请注意，我以这种方式对通道进行了排序，因为这对应于mask 中的位顺序。哪个位对应哪个通道应该很明显。 mask参数中的bit 0对应lane 0，以此类推。

由于您在 mask 中仅指定 8 位，即 8 个通道，这一事实更加复杂：

uint32_t mask = 0xff;

如果您希望扭曲有正确的可能性使用所有 32 个通道来处理所有 4 个对象，则必须指定 32 位 mask：

uint32_t mask = 0xffffffff;

没有对 8 位 mask 的“适配”以应用于 warp 中的每组 8 个通道。您必须为 32 个通道中的每一个显式指定 mask。即使使用了width 参数也是如此（见下文）。

如果您想让 shuffle 操作仅在 8 位组（具有 4 个逻辑 shuffle）中工作，这就是 width parameter 的用途：

T __shfl_up_sync(unsigned mask, T var, unsigned int delta, int width=warpSize);
                                                               ^^^^^

【讨论】：