Как правильно поддерживать инструкции `__shfl()` и `__shfl_sync()`?

Question

Насколько я понимаю, CUDA 10.1 удалил shfl инструкции:

PTX ISA версии 6.4 удаляет следующие функции:

Поддержка для shfl и голосовать за инструкции без .sync квалификатор был удален для.targetsm_70 и выше. Эта поддержка устарела с версии 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 инструкции. Я не уверен, что я делаю неправильно.

Answer 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;
}