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.