Как объявить функцию с аргументом, который является замыканием в C ++ 11 в функции устройства Cuda? - PullRequest
Новая
       40

Как объявить функцию с аргументом, который является замыканием в C ++ 11 в функции устройства Cuda?

0 голосов
/ 30 марта 2019

Я работаю над Cuda с C ++ 11 (я не думаю, что Cuda пока поддерживает более поздние версии C ++).У меня есть объект замыкания, который передается в функцию Process(), которая вызывает замыкание для каждой итерации.

Я понимаю, что функциональность std:: обычно недоступна в Cuda.Например, когда я пытаюсь использовать std::function< float(uint32_t) >, я получаю эту ошибку:

ошибка: вызов функции host ("std :: function :: function <::, void, void> ") из global функции (" _ NV_ANON_NAMESPACE :: LargeKernel ") не разрешено

Чем я могу заменить lookupFunc, чтобы она компилировалась безstd::function в наличии?Я смог обойти это, создав шаблон функции для определения типа лямбда-функции.

Этот код работает и показывает работу, которую я использовал: 

//using lookupFunc = std::function< float(uint32_t) >;

template< typename Lambda > // Work around with function template
__device__
void Process(float       * const outData,
             const  int32_t      locationX,
             const Lambda /* lookupFunc */ lambda)
{
    float answer = 0.f;

    for( int32_t offset = -1 ; ++offset < 1024 ; )
    {
        const float value = lambda( offset );

        answer += value;
    }

    outData[ locationX ] = answer;
}

__global__
void LargeKernel(const float * const inData,
                 float       * const outData)
{
    constexpr uint32_t cellStride = 1;
    const     int32_t  locationX  = threadIdx.x + blockDim.x * blockIdx.x;
    const auto lambda
        = [locationX, inData, cellStride](const int32_t offset)
          {
              return inData[ locationX + offset + cellStride ];
          };

    Process( outData, locationX, lambda );
}

Iтакже пробовал: 

using lookupFunc = float(* const)(uint32_t);

Но это дает ошибку:

Ошибка: нет подходящей функции преобразования из "const lambda -> float" в "float (*) (uint32_t)"Существует

Как я могу объявить тип третьего аргумента для Process() без использования шаблона?

Ответы [ 2 ]

3 голосов
/ 30 марта 2019

CUDA, эквивалентный std::function, равен nvstd::function. Начиная с CUDA 8.0 nvstd::function может использоваться в коде хоста и устройства - но «экземпляры не могут быть переданы из кода хоста в код устройства (и наоборот) во время выполнения». Его использование объясняется в руководстве по программированию CUDA .

0 голосов
/ 30 марта 2019

Нужно ли использовать лямбду?

В противном случае вы можете смоделировать его, объявив struct

struct noLambda
 {
   std::int32_t const   locationX;
   float const * const  inData;
   std::uint32_t const  cellStride;

   noLambda (std::int32_t l0, float const * const i0, std::uint32_t c0)
      : locationX{l0}, inData{i0}, cellStride{c0}
    { }

   float operator() (std::int32_t const offset) const
    { return inData[ locationX + offset + cellStride ]; }
 };

Таким образом, подпись процесса становится 

void Process(float       * const outData,
             const  int32_t      locationX,
             const noLambda lambda)

и вызывается следующим образом 

Process( outData, locationX, noLambda{locationX, inData, cellStride} );

(осторожно: код не проверен)

...