2016-10-19 4 views
0

Можно ли использовать CURAND вместе с Thrust внутри функтора устройства? Минимальный пример кода может быть:Использование 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; 
} 

Каков правильный способ создания случайных чисел внутри операторской функции?

+0

вы должны использовать API-интерфейс устройства cuRAND, который описан здесь: http://docs.nvidia.com/cuda/curand/device-api-overview.html#device-api-overview –

ответ

3

Можно ли использовать KURAND вместе с Thrust внутри функтора устройства?

Да, это возможно. Как указано @ m.s. большинство из того, что вам нужно от curand, можно получить от curand device api example в документации на curand. (На самом деле в документации имеется даже полный код примера тяги/куранда here)

Мы можем имитировать поведение ядра установки, указанное там, с вызовом алгоритма тяги, например. thrust::for_each_n, чтобы установить начальные переменные состояния curand для каждого элемента вектора устройства.

После этого необходимо только передать инициализированное состояние 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 
$ 

Что касается этого вопроса:

Что такое правильный способ для создания случайных чисел в функции оператора?

Там нет никаких проблем с использованием curand в тяге, но вы также можете быть в курсе, что тяга имеет встроенный RNG facility и есть полностью работал пример использования here.

+0

Это работает отлично. Я предпочитаю КУРРАН, потому что он обеспечивает хорошее качество случайных чисел, используя последовательности вместо разных семян. Однако я попытался использовать 'ds = 1000000', и программа выйдет из строя. Это не должно быть проблемой памяти (каждый curandState - 48b). – GerryR

+0

Это не сбой для меня с ds = 1000000. Какую версию ОС, GPU, CUDA и GPU вы компилируете? Если вы работаете в визуальной студии Windows, убедитесь, что вы создаете 64-разрядный (x64) проект выпуска, а не проект win32, а не проект отладки. Кстати, у самих документов CURAND есть код примера [thrust + CURAND] (http://docs.nvidia.com/cuda/curand/device-api-overview.html#thrust-and-curand-example) –

+0

I я использую Ubuntu 16.04, GeForce GTX 745/PCIe/SSE2, CUDA 8.0, arch = sm_50, у меня есть ошибка: «завершение вызова после вызова экземпляра« thrust :: system :: system_error »what(): время запуска и был прерван « – GerryR

Смежные вопросы