Visible to Intel only — GUID: GUID-F203F08E-B8C3-4F43-A57C-9561BF7B53E4
Visible to Intel only — GUID: GUID-F203F08E-B8C3-4F43-A57C-9561BF7B53E4
oneMKL RNG Device Usage Model
Example of Random Numbers Generation by Engines Stored in sycl::buffer
Example of Random Numbers Generation with Host-side Helpers Usage
Example of Random Numbers Generation Using Gamma Distribution
A typical usage model for device routines is the same as described in oneMKL RNG Usage Model:
Create and initialize the object for basic random number generator.
Create and initialize the object for distribution generator.
Call the generate routine to get random numbers with appropriate statistical distribution.
Example of Scalar Random Numbers Generation
#include <iostream> #include <sycl/sycl.hpp> #include "oneapi/mkl/rng/device.hpp" int main() { sycl::queue queue; const int n = 1000; const int seed = 1; // Prepare an array for random numbers std::vector<float> r(n); sycl::buffer<float, 1> r_buf(r.data(), r.size()); // Submit a kernel to generate on device queue.submit([&](sycl::handler& cgh) { auto r_acc = r_buf.template get_access<sycl::access::mode::write>(cgh); cgh.parallel_for(sycl::range<1>(n), [=](sycl::item<1> item) { // Create an engine object oneapi::mkl::rng::device::philox4x32x10<> engine(seed, item.get_id(0)); // Create a distribution object oneapi::mkl::rng::device::uniform<> distr; // Call generate function to obtain scalar random number float res = oneapi::mkl::rng::device::generate(distr, engine); r_acc[item.get_id(0)] = res; }); }); auto r_acc = r_buf.template get_access<sycl::access::mode::read>(); std::cout << "Samples of uniform distribution" << std::endl; for(int i = 0; i < 10; i++) { std::cout << r_acc[i] << std::endl; } return 0; }
Example of Vector Random Numbers Generation
#include <iostream> #include <sycl/sycl.hpp> #include "oneapi/mkl/rng/device.hpp" int main() { sycl::queue queue; const int n = 1000; const int seed = 1; const int vec_size = 4; // Prepare an array for random numbers std::vector<float> r(n); sycl::buffer<float, 1> r_buf(r.data(), r.size()); // Submit a kernel to generate on device sycl::queue{}.submit([&](sycl::handler& cgh) { auto r_acc = r_buf.template get_access<sycl::access::mode::write>(cgh); cgh.parallel_for(sycl::range<1>(n / vec_size), [=](sycl::item<1> item) { // Create an engine object oneapi::mkl::rng::device::philox4x32x10<vec_size> engine(seed, item.get_id(0) * vec_size); // Create a distribution object oneapi::mkl::rng::device::uniform<> distr; // Call generate function to obtain sycl::vec<float, 4> with random numbers auto res = oneapi::mkl::rng::device::generate(distr, engine); res.store(ite.get_id(0), r_acc); }); }); auto r_acc = r_buf.template get_access<sycl::access::mode::read>(); std::cout << "Samples of uniform distribution" << std::endl; for(int i = 0; i < 10; i++) { std::cout << r_acc[i] << std::endl; } return 0; }
There is an opportunity to store engines between kernels manually via sycl::buffer / USM pointers or by using a specific host-side helper class called, engine descriptor.
Example of Random Numbers Generation by Engines Stored in sycl::buffer
Engines are initialized in the first kernel. Random number generation is performed in the second kernel.
#include <iostream> #include <sycl/sycl.hpp> #include "oneapi/mkl/rng/device.hpp" int main() { sycl::queue queue; const int n = 1000; const int seed = 1; const int vec_size = 4; // Prepare an array for random numbers std::vector<float> r(n); sycl::buffer<float, 1> r_buf(r.data(), r.size()); sycl::range<1> range(n / vec_size); sycl::buffer<oneapi::mkl::rng::device::mrg32k3a<vec_size>, 1> engine_buf(range); sycl::queue queue; // Kernel with initialization of engines queue.submit([&](sycl::handler& cgh) { // Create an accessor to sycl::buffer with engines to write initialized states auto engine_acc = engine_buf.template get_access<sycl::access::mode::write>(cgh); cgh.parallel_for(range, [=](sycl::item<1> item) { size_t id = item.get_id(0); // Create an engine object with offset id * 2^64 oneapi::mkl::rng::device::mrg32k3a<vec_size> engine(seed, {0, id}); engine_acc[id] = engine; }); }); // Kernel for random numbers generation queue.submit([&](sycl::handler& cgh) { auto r_acc = r_buf.template get_access<sycl::access::mode::write>(cgh); // Create an accessor to sycl::buffer with engines to read initialized states auto engine_acc = engine_buf.template get_access<sycl::access::mode::read>(cgh); cgh.parallel_for(range, [=](sycl::item<1> item) { size_t id = item.get_id(0); auto engine = engine_acc[id]; oneapi::mkl::rng::device::uniform distr; auto res = oneapi::mkl::rng::device::generate(distr, engine); res.store(id, r_acc); }); }); auto r_acc = r_buf.template get_access<sycl::access::mode::read>(); std::cout << "Samples of uniform distribution" << std::endl; for(int i = 0; i < 10; i++) { std::cout << r_acc[i] << std::endl; } return 0; }
Example of Random Numbers Generation with Host-side Helpers Usage
#include <iostream> #include <sycl/sycl.hpp> #include "oneapi/mkl/rng/device.hpp" int main() { sycl::queue queue; const int n = 1000; const int seed = 1; const int vec_size = 4; // prepare array for random numbers std::vector<float> r(n); sycl::buffer<float, 1> r_buf(r.data(), r.size()); sycl::range<1> range(n / vec_size); // offset of each engine in engine_descriptor int offset = vec_size; // each engine would be created in enqueued task as of specified range // as oneapi::mkl::rng::device::mrg32k3a<vec_size>(seed, id * offset); oneapi::mkl::rng::device::engine_descriptor<oneapi::mkl::rng::device::mrg32k3a<vec_size>> descr(queue, range, seed, offset); queue.submit([&](sycl::handler& cgh) { auto r_acc = r_buf.template get_access<sycl::access::mode::write>(cgh); // create engine_accessor auto engine_acc = descr.get_access(cgh); cgh.parallel_for(range, [=](sycl::item<1> item) { size_t id = item.get_id(0); // load engine from engine_accessor auto engine = engine_acc.load(id); oneapi::mkl::rng::device::uniform<Type> distr; auto res = oneapi::mkl::rng::device::generate(distr, engine); res.store(id, r_acc); // store engine for furter calculations if needed engine_acc.store(engine, id); }); }); auto r_acc = r_buf.template get_access<sycl::access::mode::read>(); std::cout << "Samples of uniform distribution" << std::endl; for(int i = 0; i < 10; i++) { std::cout << r_acc[i] << std::endl; } return 0; }
Example of Random Numbers Generation Using Gamma Distribution
#include <iostream> #include <vector> #include <sycl/sycl.hpp> #include "oneapi/mkl/rng/device.hpp" constexpr size_t seed = 42; int main() { sycl::queue queue; const size_t n = 10'000; // create USM allocator sycl::usm_allocator<double, sycl::usm::alloc::shared> allocator(queue); // create vector with USM allocator std::vector<double, decltype(allocator)> r_vec(n, allocator); double* r = r_vec.data(); size_t reduced_val = 0; // create a reductor to calculate the sum of rejected numbers across all work-items auto reductor = sycl::reduction(&reduced_val, size_t(0), std::plus<size_t>{}); queue.parallel_for(n, reductor, [=](size_t id, auto& sum) { // create basic random number generator object // 2^76 numbers for each work-item are skipped to allow acceptance rejection scheme ro reject as much numbers as it needs oneapi::mkl::rng::device::mrg32k3a<1> engine({ seed, seed, seed, seed, seed, seed }, { 0, (4096 * id) }); // create distribution object oneapi::mkl::rng::device::gamma<double> distr(0.7, 0.0, 1.0); // perform generation auto res = oneapi::mkl::rng::device::generate(distr, engine); r[id] = res; // calculate how much random numbers are rejected in this work-item size_t rej = distr.count_rejected_numbers(); sum += rej; }).wait(); // synchronization can be also done by queue.wait() std::cout << "First 10 gamma distributed numbers are following:"<< std::endl; for(int i = 0; i < 10; i++) { std::cout << r[i] << " "; } std::cout << std::endl; std::cout << "rejected " << reduced_val <<" random numbers" << std::endl; return 0; }
Additionally, examples that demonstrate usage of random number generators device routines are available in:
${MKL}/examples/dpcpp_device/rng/source