Using a __device__ lambda in CUDA 8.0

In CUDA 8.0 C++ 11 lambdas can be used on the device. See the following code.

#include <thrust/device_vector.h>
#include <thrust/sequence.h>
#include <thrust/transform.h>
#include <iostream>

using namespace std;
using namespace thrust;

struct double_functor {
__device__ int operator()(const int value) {
return value * 2;
}
};

void show(device_vector<int>& d) {
for (int i = 0; i < d.size(); i++) {
cout << "d[" << i << "] = " << d[i] << endl;
}
}

void thrust_map_example_cuda7() {
device_vector<int> d(16);
double_functor functor;
sequence(d.begin(), d.end(), 1);
// pass the instance of the functor
transform(d.begin(), d.end(), d.begin(), functor);
show(d);
}

/*
* with CUDA 8.0 you can use __device__ lambdas
*/

void thrust_map_example_cuda8() {
device_vector<int> d(16);
sequence(d.begin(), d.end(), 1);
// here comes the device lambda !
transform(d.begin(), d.end(), d.begin(), [=] __device__ (const int value) {
return value * 2;
});
show(d);
}

int main(int argc, char** argv) {
thrust_map_example_cuda7();
thrust_map_example_cuda8();
return 0;
}

See also on gist.

Remark: This post was adapted to the new blog format in November 2016.

 "Guter Roman über IT und DevOps" "Web-App für die Visualisierung der Verbreitung von Familiennamen"