c++ - 在 Thrust 仿函数中使用 CURAND

Question

是否可以在设备仿函数中将 CURAND 与 Thrust 一起使用？最小代码示例可以是:

#include <thrust/device_vector.h>

struct Move
{
    Move() {}

    using Position = thrust::tuple<double, double>;

    __host__ __device__
    Position operator()(Position p)
    {
        thrust::get<0>(p) += 1.0; // use CURAND to add a random N(0,1)
        thrust::get<1>(p) += 1.0; // use CURAND to add a random N(0,1)
        return p;
    }
};

int main()
{
    // Create vectors on device
    thrust::device_vector<double> p1(10, 0.0);
    thrust::device_vector<double> p2(10, 0.0);

    // Create zip iterators
    auto pBeg = thrust::make_zip_iterator(thrust::make_tuple(p1.begin(), p2.begin()));
    auto pEnd = thrust::make_zip_iterator(thrust::make_tuple(p1.end(),   p2.end()  ));

    // Move points in the vectors
    thrust::transform(pBeg, pEnd, pBeg, Move());

    // Print result (just for debug)
    thrust::copy(p1.begin(), p1.end(), std::ostream_iterator<double>(std::cout, "\n"));
    thrust::copy(p2.begin(), p2.end(), std::ostream_iterator<double>(std::cout, "\n"));

    return 0;
}

在运算符函数中创建随机数的正确方法是什么？

Answer 1

Is it possible to use CURAND together with Thrust inside a device functor?

是的，这是可能的。如@m.s 所示。您需要从 curand 获得的大部分内容都可以从 curand device api example 获得在 curand 文档中。 (事实上​​ ，文档中甚至有完整的 thrust/curand 示例代码 here )

我们可以通过推力算法调用来模拟那里指示的设置内核的行为，例如。 thrust::for_each_n 为每个设备 vector 元素设置初始 curand 状态变量。

之后，只需通过 zip 迭代器中的附加迭代器将初始化的 curand 状态传递给 Move 仿函数，然后调用 curand_uniform(例如) 在仿函数中。

这是一个基于您的代码的完整示例:

$ cat t20.cu
#include <thrust/device_vector.h>
#include <curand_kernel.h>
#include <iostream>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/transform.h>
#include <thrust/for_each.h>

const int seed = 1234;
const int ds = 10;
const int offset = 0;

struct Move
{
    Move() {}

    using Position = thrust::tuple<double, double, curandState>;

    __device__
    Position operator()(Position p)
    {
        curandState s = thrust::get<2>(p);
        thrust::get<0>(p) += curand_uniform(&s); // use CURAND to add a random N(0,1)
        thrust::get<1>(p) += curand_uniform(&s); // use CURAND to add a random N(0,1)
        thrust::get<2>(p) = s;
        return p;
    }
};

struct curand_setup
{
    using init_tuple = thrust::tuple<int, curandState &>;
    __device__
    void operator()(init_tuple t){
      curandState s;
      int id = thrust::get<0>(t);
      curand_init(seed, id, offset, &s);
      thrust::get<1>(t) = s;
      }
};

int main()
{
    // Create vectors on device
    thrust::device_vector<double> p1(ds, 0.0);
    thrust::device_vector<double> p2(ds, 0.0);
    thrust::device_vector<curandState> s1(ds);

    // Create zip iterators
    auto pBeg = thrust::make_zip_iterator(thrust::make_tuple(p1.begin(), p2.begin(), s1.begin()));
    auto pEnd = thrust::make_zip_iterator(thrust::make_tuple(p1.end(),   p2.end(), s1.end()  ));
    auto pInit = thrust::make_zip_iterator(thrust::make_tuple(thrust::counting_iterator<int>(0), s1.begin()));
    // initialize random generator
    thrust::for_each_n(pInit, ds, curand_setup());
    // Move points in the vectors
    thrust::transform(pBeg, pEnd, pBeg, Move());

    // Print result (just for debug)
    thrust::copy(p1.begin(), p1.end(), std::ostream_iterator<double>(std::cout, "\n"));
    thrust::copy(p2.begin(), p2.end(), std::ostream_iterator<double>(std::cout, "\n"));

    return 0;
}
$ nvcc -arch=sm_61 -std=c++11 t20.cu -o t20 -lcurand
$ ./t20
0.145468
0.820181
0.550399
0.29483
0.914733
0.868979
0.321921
0.782857
0.0113023
0.28545
0.434899
0.926417
0.811845
0.308556
0.557235
0.501246
0.206681
0.123377
0.539587
0.198575
$

关于这个问题:

What is the right way to create random numbers inside the operator function?

在 thrust 中使用 curand 没有问题，但您可能还想知道 thrust 有一个内置的 RNG facility并且有一个完整的用法示例 here .