Using a __device__ lambda in CUDA 8.0
September 30, 2016
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.