Visible to Intel only — GUID: GUID-233C06C0-F462-4E64-AB26-27E6F5E9AF4D
Visible to Intel only — GUID: GUID-233C06C0-F462-4E64-AB26-27E6F5E9AF4D
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