在推力函子中使用 CURAND

Posted 2023-02-21

【中文标题】在推力函子中使用 CURAND

【英文标题】：Using CURAND inside a Thrust functor

【发布时间】：2016-10-19 22:23:26

【问题描述】：

是否可以在设备函子内将 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;

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

【问题讨论】：

您必须使用此处记录的 cuRAND 设备 API：docs.nvidia.com/cuda/curand/…

【参考方案1】：

是否可以在设备函子内将 CURAND 与 Thrust 一起使用？

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

我们可以通过推力算法调用来模仿那里指示的设置内核的行为，例如。 thrust::for_each_n 为每个设备向量元素设置初始 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
$

关于这个问题：

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

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

【讨论】：

这非常有效。我更喜欢 CURAND，因为它使用序列而不是不同的种子来确保随机数的质量。但是，我尝试使用 ds = 1000000 并且程序崩溃了。应该不是内存问题（每个curandState都是48b）。

ds=1000000 对我来说不会崩溃。您正在为什么操作系统、GPU、CUDA 版本和 GPU 架构进行编译？如果您使用的是 Windows Visual Studio，请确保您构建的是 64 位 (x64) 发布项目，而不是 win32 项目，也不是调试项目。顺便说一句，CURAND 文档本身有一个 thrust + CURAND example code

我使用的是 Ubuntu 16.04，GeForce GTX 745/PCIe/SSE2，CUDA 8.0，arch=sm_50，我遇到的错误是：terminate called after throwing an instance of 'thrust::system::system_error' what(): the launch timed out and was terminated

感谢您的链接，尽管我认为最好使用 thread_id 作为序列 ID 而不是种子（就像您在此解决方案中所做的那样）。

您的内核在托管显示器的 GPU 上运行时间过长。阅读this。这不是由于代码中的任何缺陷造成的。