Create a gist now

Instantly share code, notes, and snippets.

Embed
Using a __device__ lambda in CUDA 8.0
/*
* The following example shows the use of lambdas on the device with CUDA 8.0.
*
* needs the nvcc compiler option --expt-extended-lambda to compile.
*
* (c) 2016 Joern Dinkla, www.dinkla.net
*/
#include <thrust/device_vector.h>
#include <thrust/sequence.h>
#include <thrust/transform.h>
#include <iostream>
using namespace std;
using namespace thrust;
/*
* With CUDA 7.0 you have to use a functor class that implements operator().
*/
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);
transform(d.begin(), d.end(), d.begin(), functor); // pass the instance of the 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;
}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment