Лямбда-выражения с CUDA

Если я использую thrust::transform в thrust::host, использование лямбда в порядке

thrust::transform(thrust::host, a, a+arraySize,b,d,[](int a, int b)->int
{
    return a + b;
});

Однако, если я изменяю thrust::host на thrust::device, код не будет передавать компилятор. Вот ошибка на VS2013:

Тип закрытия для лямбда ( "lambda [] (int, int) → int" ) не может использоваться в типе аргумента шаблона экземпляра шаблона функции __global__, если только лямбда не определена в пределах __device__ или __global__ функция

Таким образом, проблема заключается в том, как использовать __device__ или __global__ в связи с устройством lambdas.

Ответ 1

В CUDA 7 это невозможно. Цитата из Марк Харрис:

Это не поддерживается сегодня в CUDA, потому что лямбда - это хост-код. Передача lambdas от хоста к устройству является сложной проблемой, но это то, что мы будем исследовать для будущего выпуска CUDA.

Что вы можете сделать в CUDA 7, это алгоритмы выталкивания вызовов из вашего кода устройства, и в этом случае вы можете передать им лямбда...

С помощью CUDA 7 алгоритмы тяги можно вызывать из кода устройства (например, ядра CUDA или функторов __device__). В таких ситуациях вы можете использовать (устройство) лямбда с тягой. Пример приведен в сообщении в блоге parallelforall здесь.

Однако CUDA 7.5 представляет экспериментальную функцию лямбда-устройства. Эта функция описана здесь:

CUDA 7.5 представляет экспериментальную функцию: GPU lambdas. GPU lambdas - это анонимные объекты функций устройства, которые вы можете определить в коде хоста, аннотируя их спецификатором __device__.

Чтобы включить компиляцию для этой функции (в настоящее время с CUDA 7.5), необходимо указать --expt-extended-lambda в командной строке nvcc.

Ответ 2

Этот простой код с использованием устройства lambdas работает под CUDA 8.0 RC, хотя устройство lambdas для этой версии CUDA все еще находится на экспериментальной стадии:

#include <thrust/device_vector.h>
#include <thrust/functional.h>
#include <thrust/transform.h>

using namespace thrust::placeholders;

int main(void)
{
    // --- Input data 
    float a = 2.0f;
    float x[4] = { 1, 2, 3, 4 };
    float y[4] = { 1, 1, 1, 1 };

    thrust::device_vector<float> X(x, x + 4);
    thrust::device_vector<float> Y(y, y + 4);

    thrust::transform(X.begin(), 
                      X.end(),  
                      Y.begin(), 
                      Y.begin(),
                      [=] __host__ __device__ (float x, float y) { return a * x + y; }      // --- Lambda expression 
                     );        

    for (size_t i = 0; i < 4; i++) std::cout << a << " * " << x[i] << " + " << y[i] << " = " << Y[i] << std::endl;

    return 0;
}

Не забудьте использовать

--expt-extended-lambda

для компиляции.