支持 `__shfl()` 和 `__shfl_sync()` 指令的正确方法是什么？答案

【问题标题】：What is the correct way to support `__shfl()` and `__shfl_sync()` instructions?支持 `__shfl()` 和 `__shfl_sync()` 指令的正确方法是什么？
【发布时间】：2019-09-18 13:09:56
【问题描述】：

据我了解，CUDA 10.1 删除了shfl 指令：

PTX ISA 6.4 版删除了以下功能：

.targetsm_70 及更高版本已删除对 shfl 和不带 .sync 限定符的投票指令的支持。自 PTX ISA 6.0 版以来，此支持已被弃用，如 PTX ISA 6.2 版中所述。

支持shfl未来和过去的CUDA版本的正确方法是什么？

我当前的方法（在下面共享）导致使用 CUDA 10.1 时出错：

ptxas ... line 466727; error   : Instruction 'shfl' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4

template <typename T>
__device__ static __forceinline__
T _shfl_up(T var, unsigned int delta, int width=WARPSIZE, unsigned mask=MEMBERMASK)
{
#if (__CUDACC_VER_MAJOR__ >= 9)
  var = __shfl_up_sync(mask, var, delta, width);
#else
  var = __shfl_up(var, delta, width);
#endif
  return var;
}

另外，我想补充一点，我的项目的依赖项之一是 CUB，我相信他们使用相同的方法来拆分 _sync() 和旧的 shfl 指令。我不确定我做错了什么。

【问题讨论】：

标签： cuda ptx ptxas

【解决方案1】：

我做对了，结果发现另一个依赖项不支持sync，为此创建了一个拉取请求：https://github.com/moderngpu/moderngpu/pull/32

template <typename T>
__device__ static __forceinline__
T _shfl_up(T var, unsigned int delta, int width=WARPSIZE, unsigned mask=MEMBERMASK)
{
#if ( __CUDA_ARCH__ >= 300)
#if (__CUDACC_VER_MAJOR__ >= 9)
  var = __shfl_up_sync(mask, var, delta, width);
#else
  var = __shfl_up(var, delta, width);
#endif
#endif
  return var;
}

【讨论】：