Visible to Intel only — GUID: GUID-59FA97CE-853B-4459-944C-39FB71F380EC
Visible to Intel only — GUID: GUID-59FA97CE-853B-4459-944C-39FB71F380EC
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
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; }
Additionally, examples that demonstrate usage of random number generators device routines are available in:
${MKL}/examples/dpcpp_device/rng/source