c-CUDA:车道之间的__shfl delta是否不同

Question

根据CUDA文档，__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"不会环绕……因此有效地，上部delta通道将保持不变。"。也没有任何内容指定delta必须是>0，0必须只返回与超出范围的delta相同的内容？

除此之外，文件在这里是否自相矛盾？

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

因此，它对于非活动线程是未定义的，然而。。。

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

这意味着对于非活动线程，它返回调用线程自己的变量。这进一步意味着，在我上面的例子中，如果delta超出了界限，它会添加自己，而不是什么都不添加？

Answer 1

不同车道之间的__shfl delta会不同吗？

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

#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: %dn", threadIdx.x, value);
}
int main() {
  bcast<<< 1, 32 >>>();
  cudaDeviceSynchronize();
  return 0;
}

关于你关于delta值的问题（注意，__shfl_up和__shfl_down的函数原型坚持认为delta是unsigned int的量，所以它不能是负的），关于0的delta值的问题我不太清楚。这将引用源通道，根据定义，它不能超出范围或不活动。因此，是的，零索引值将返回起始车道值，就像超出范围的索引一样。

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

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