CUDA：车道之间的 __shfl delta 可以不同吗？答案

【问题标题】：CUDA: can __shfl delta be different between lanes?CUDA：车道之间的 __shfl delta 可以不同吗？
【发布时间】：2016-01-03 03:11:32
【问题描述】：

根据CUDA documentation，__shfl() 内在函数允许在线程之间交换变量。更具体地说，__shfl_up() 和__shfl_down()（也称为__shfl_xor()）允许交换不同通道的变量。 unsigned int delta 是这些的第二个参数，并指定（正或负）要交换变量的通道的 offset 通道 ID。我一直认为delta 对于经线中的所有车道必须保持不变。因此，例如，必须像这样调用

val += __shfl_down(val, 3);

或

for (i = 1; i < warpSize / 2; i ++)
  val += __shfl_down(val, i);

但是，我刚刚意识到没有明确说明 delta 对于所有车道必须相同（只要所有车道都参与）。所以，是下面undefined还是罚款

val += __shfl_down(val, threadIdx.x % warpSize);

这只是一个例子，许多通道不会添加任何内容，因为 delta "will not wrap around ... so effectively the upper delta lanes will remain unchanged.". 也没有指定 delta 必须是 >0，0 必须只返回与 out 相同的值-of-range delta?

除此之外，这里的文档是不是自相矛盾？

线程只能从另一个积极参与__shfl() 命令的线程读取数据。如果目标线程处于非活动状态，则检索到的值未定义。

所以它对于非活动线程是未定义的，但是...

所有__shfl() 内部函数都将源通道 ID 中 var 引用的 4 字节字作为无符号整数返回。如果源通道 ID 超出范围或源线程已退出，则返回调用线程自己的 var。

... 这意味着对于非活动线程，它返回调用线程自己的 var。这进一步意味着，在我上面的示例中，如果delta 越界，它会自行添加，而不是什么？

【问题讨论】：

标签： c cuda

【解决方案1】：

车道之间的 __shfl delta 可以不同吗？

是的，delta 在通道之间可能不同。例如，这可以用于将 32 位数量从单个通道广播到其他通道（尽管这不是执行此类broadcast 的唯一方法）：

#include <stdio.h>
__global__ void bcast() {
  int value = threadIdx.x;
  value = __shfl_up(value, threadIdx.x); // Get "value" from lane 0
  if (value) printf("Thread %u failed: %d\n", threadIdx.x, value);
}

int main() {
  bcast<<< 1, 32 >>>();
  cudaDeviceSynchronize();
  return 0;
}

关于您关于delta 值的问题，（请注意，__shfl_up 和__shfl_down 的function prototypes 坚持认为delta 是unsigned int 数量，因此不能为负数），问题关于delta 值0 对我来说不是很清楚。这将引用原始通道，根据定义，该通道不能超出范围或处于非活动状态。所以，是的，零索引值将返回起始车道值，就像超出范围的索引一样。

关于您的最后一个问题，越界计算的源通道 ID 与碰巧引用非活动线程的界内计算的源通道 ID 不同。在前一种情况下，返回线程自己的 shuffle 变量。在后一种情况下，返回的值是未定义的。

计算大于 31 的 delta 值的源通道 ID 也可能具有一些额外的复杂性。随意地，这样的计算似乎总是越界（但实际行为可能更复杂一些）。我不确定文档是否明确涉及这一点，我不确定这是您任何问题的症结所在。

【讨论】：