CUDA Thrust快捷数学函数

Question

有没有一种自动将CUDA数学函数包装在一个函数对象中的方法，以便可以应用thrust :: transform而无需手动编写函数对象？就像（我所了解的）std :: function提供的功能一样？ thrust :: placeholders 似乎不喜欢数学函数。 std :: function 似乎不可用。

示例代码：

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

struct myfunc{
    __device__ 
    double operator()(double x,double y){
    return hypot(x,y);
    }
};

int main(){

    double x0[10] = {3.,0.,1.,2.,3.,4.,5.,6.,7.,8.};
    double y0[10] = {4.,0.,1.,2.,3.,4.,5.,6.,7.,8.};

    thrust::device_vector<double> x(x0,x0+10);
    thrust::device_vector<double> y(y0,y0+10);
    thrust::device_vector<double> r(10);

    for (int i=0;i<10;i++) std::cout << x0[i] <<" ";    std::cout<<std::endl;
    for (int i=0;i<10;i++) std::cout << y0[i] <<" ";    std::cout<<std::endl;

    // this works:
    thrust::transform(x.begin(),x.end(),y.begin(),r.begin(), myfunc());

    // this doesn't compile:
    using namespace thrust::placeholders;
    thrust::transform(x.begin(),x.end(),y.begin(),r.begin(), hypot(_1,_2));

    // nor does this:
    thrust::transform(x.begin(),x.end(),y.begin(),r.begin(), std::function<double(double,double)>(hypot));


    for (int i=0;i<10;i++) std::cout << r[i] <<" ";    std::cout<<std::endl;
}

- Sean Dettrick

---

没有自动完成它的方法。实现类似功能的一种方法可能是创建一个与CUDA互操作的类似于std::bind的东西。然后，您需要使用bind来定义所有感兴趣的数学函数（例如hypot）的重载。 - Jared Hoberock

在CUDA 7.5中，您可以使用实验性的--expt-extended-lambda功能，并编写auto h = [] __device__(double x, double y){return hypot(x,y);}; thrust::transform(x.begin(),x.end(),y.begin(),r.begin(), h);。 - m.s.

如果你想提供答案，我会点赞。我认为Jared不会反对。 - Robert Crovella

Answer 1

将我的评论转换为此答案：

正如@JaredHoberock已经说明的那样，没有自动实现你想要的功能的方式。总是有一些语法/输入开销。

减少编写单独函数（就像你使用my_func一样）的开销的一种方法是使用lambda表达式。自CUDA 7.5以来，有一个实验性设备lambda特性，允许你做以下事情:

auto h = []__device__(double x, double y){return hypot(x,y);};
thrust::transform(x.begin(),x.end(),y.begin(),r.begin(), h);

您需要添加以下nvcc编译器开关来进行编译：

nvcc --expt-extended-lambda ...

---

另一种方法是使用以下Wrapper将函数转换为函数对象：

template<typename Sig, Sig& S>
struct Wrapper;

template<typename R, typename... T, R(&function)(T...)>
struct Wrapper<R(T...), function>
{
    __device__
    R operator() (T&... a)
    {
        return function(a...);
    }
};

然后您可以像这样使用它:

 thrust::transform(x.begin(),x.end(),y.begin(),r.begin(), Wrapper<double(double,double), hypot>());

- m.s.

Answer 2

正如m.s.所提到的，减少编写functor开销的一种可能方法是使用lambda表达式。请注意，CUDA 8.0 RC支持GPU lambda（虽然仍处于实验阶段）。另一种可能性是使用占位符技术。下面有两个工作示例，分别针对这两种情况：

LAMBDA表达式

#include <thrust/device_vector.h>
#include <thrust/functional.h>
#include <thrust/transform.h>

// Available for device operations only from CUDA 8.0 (experimental stage)
// Compile with the flag --expt-extended-lambda

using namespace thrust::placeholders;

int main(void)
{
    // --- Input data 
    float a = 2.0f;
    float x[4] = { 1, 2, 3, 4 };
    float y[4] = { 1, 1, 1, 1 };

    thrust::device_vector<float> X(x, x + 4);
    thrust::device_vector<float> Y(y, y + 4);

    thrust::transform(X.begin(), 
                      X.end(),  
                      Y.begin(), 
                      Y.begin(),
                      [=] __host__ __device__ (float x, float y) { return a * x + y; }      // --- Lambda expression 
                     );        

    for (size_t i = 0; i < 4; i++) std::cout << a << " * " << x[i] << " + " << y[i] << " = " << Y[i] << std::endl;

    return 0;
}

PLACEHOLDERS

#include <thrust/device_vector.h>
#include <thrust/functional.h>
#include <thrust/transform.h>

using namespace thrust::placeholders;

int main(void)
{
    // --- Input data 
    float a = 2.0f;
    float x[4] = { 1, 2, 3, 4 };
    float y[4] = { 1, 1, 1, 1 };

    thrust::host_vector<float> X(x, x + 4);
    thrust::host_vector<float> Y(y, y + 4);

    thrust::transform(X.begin(), X.end(),
                      Y.begin(),         
                      Y.begin(),
                      a * _1 + _2);

    for (size_t i = 0; i < 4; i++) std::cout << a << " * " << x[i] << " + " << y[i] << " = " << Y[i] << std::endl;

    return 0;
}

- Vitality