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

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

0 голосов
/ 30 апреля 2019

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

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

Поддержка shfl и голосование инструкций без.sync квалификатор был удален для .targetsm_70 и выше.Эта поддержка устарела с версии 6.0 PTX ISA, как описано в версии 6.2 ISA PTX.

Как правильно поддерживать 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 инструкций.Я не уверен, что я делаю неправильно.

1 Ответ

2 голосов
/ 30 апреля 2019

Я поступил правильно, оказалось, что другая зависимость не поддерживала 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;
}
...