Thrust::transform custom function

Question

I am trying to do a very basic example within CUDA. I would like to do a simple calculation on a list of floats.

vh[x] * k1 + k2

Currently I am trying this and it is not working:

Code 1

#include <vector>
#include <iostream>
#include <thrust/transform.h>
#include <thrust/functional.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>

using namespace std;
using namespace thrust;

float k1 = 42, k2 = 7;

int main(void)
{
    vector<float> vh = { 0, 1, 2, 3, 4, 5, 6, 7 };
    device_vector<float> v = vh;
    device_vector<float> v_out(v.size());

    thrust::transform(v.begin(), v.end(), v_out.begin(), [=] __device__(float x) {
        return x*k1 + k2;
    });

    for (size_t i = 0; i < v_out.size(); i++)
        std::cout << v_out[i] << std::endl;
}

I am getting a very annoying lambda function error with the above code so I have tried to use a custom function as the code below shows:

Code 2

#include <vector>
#include <iostream>
#include <thrust/transform.h>
#include <thrust/functional.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>

using namespace std;
using namespace thrust;

float k1 = 42, k2 = 7;

float multiply(float x)
{
    return x * k1 + k2;
}

int main(void) {
    vector<float> vh = { 0, 1, 2, 3, 4, 5, 6, 7 };
    device_vector<float> v = vh;
    device_vector<float> v_out(v.size());
    thrust::negate<float> op;

    thrust::transform(v.begin(), v.end(), v_out.begin(), multiply __device__(float x) );

    for (size_t i = 0; i < v_out.size(); i++)
        std::cout << v_out[i] << std::endl;

    std::getwchar();
}

Can anyone tell my why Code 1 and/or Code 2 is not working?

Answer 1

For Code 2, you must wrap your function in an object to create a functor.

For Code 1, you have to use --expt-extended-lambda nvcc option to enable full lambda support.

You also must declare k1, k2 as const, or not make them static by (for instance) declaring it inside the main.

Use functor for production code, unless your lambda is very simple.

Check out the following code for a working example:

#include <vector>
#include <iostream>
#include <thrust/transform.h>
#include <thrust/functional.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>

using namespace std;
using namespace thrust;



template<class T>
struct saxpi{
    T k1;
    T k2;
    saxpi(T _k1, T _k2){
        k1=_k1;
        k2=_k2;
    }
    __host__ __device__ T operator()(T &x) const{
        return x*k1+k2;
    }
};


int main(void)
{
    float kk1=1, kk2=5;
    vector<float> vh = { 0, 1, 2, 3, 4, 5, 6, 7 };
    device_vector<float> v = vh;
    device_vector<float> v_out(v.size());
    cout<<"Lambda:"<<endl;
    auto ff = [=]  __device__ (float x) {return kk1*x +kk2;};

    thrust::transform(v.begin(),v.end(),v_out.begin(),ff);

    for (size_t i = 0; i < v_out.size(); i++)
        std::cout << v_out[i] << std::endl;

    cout<<"Functor:"<<endl;
    saxpi<float> f(kk1,kk2);
    v_out.clear();
    v_out.resize(v.size());
    thrust::transform(v.begin(),v.end(),v_out.begin(),f);


    for (size_t i = 0; i < v_out.size(); i++)
            std::cout << v_out[i] << std::endl;

}

Compile it using the following options: --expt-extended-lambda -std=c++11

Lambda:
5
6
7
8
9
10
11
12
Functor:
5
6
7
8
9
10
11
12