时间戳

首页 > 解决方案 > 支持 `__shfl()` 和 `__shfl_sync()` 指令的正确方法是什么？

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

问题描述

据我了解，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指令。我不确定我做错了什么。

标签： cudaptxptxas

解决方案

我做对了，结果发现另一个依赖项不支持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;
}

推荐阅读