Repository navigation
Add dpnp.random.Generator on the oneMKL device RNG API - #3157
ndgrigorian wants to merge 2 commits into
Conversation
dbdb2c8 to
3963d20
Compare
3963d20 to
6c3ff46
Compare
|
View rendered docs @ https://intelpython.github.io/dpnp/pull/3157/index.html |
|
Array API standard conformance tests for dpnp=0.22.0dev1=np2py314h8d9cdd5_3 ran successfully. |
|
coverage: 84.226% (-0.2%) from 84.397% — feature/device-rng-generator into master |
There was a problem hiding this comment.
🟡 Changes recommended
Unresolved RNG correctness issues can deadlock state access, duplicate random streams through races or invalid sizes, and produce incorrect results for large draws or overflowing float32 ranges.
5 open issues
What changed in this PR
Introduces the foundation for device-side random generation through NumPy’s Generator API, the first step in the RNG migration series. Legacy RandomState and module-level functions remain unchanged.
Changes:
- Adds
Generator,default_rng, and Philox bit generators with state management and spawning. - Implements base distributions in a header-only oneMKL device RNG extension, with strided and GPU-optimized blocked sampling.
- Adds API documentation and tests for distributions, reproducibility, queues, and allocation types.
| File | Description |
|---|---|
dpnp/tests/test_usm_type.py |
Tests RNG allocation types. |
dpnp/tests/test_sycl_queue.py |
Tests queue placement and mismatch errors. |
dpnp/tests/test_random_generator.py |
Tests generators, distributions, state, and reproducibility. |
dpnp/tests/AGENTS.md |
Adds RNG testing guidance. |
dpnp/random/_generator.py |
Implements Generator and base distributions. |
dpnp/random/_bit_generator.py |
Implements Philox state, raw draws, and spawning. |
dpnp/random/__init__.py |
Exports the new random APIs. |
dpnp/libtensor/extensions/rng/uniform.cpp |
Registers uniform and raw-bit samplers. |
dpnp/libtensor/extensions/rng/sample.cpp |
Validates inputs and manages sampling lifetimes. |
dpnp/libtensor/extensions/rng/rng_py.cpp |
Defines the Python extension module. |
dpnp/libtensor/extensions/rng/normal.cpp |
Registers the normal sampler. |
dpnp/libtensor/extensions/rng/kernels/samplers.hpp |
Defines distribution sampler adapters. |
dpnp/libtensor/extensions/rng/kernels/sample.hpp |
Implements strided and blocked sampling layouts. |
dpnp/libtensor/extensions/rng/kernels/engines.hpp |
Wraps Philox engines and uniform-bit conversion. |
dpnp/libtensor/extensions/rng/kernels/distributions.hpp |
Implements parameters and Box–Muller sampling. |
dpnp/libtensor/extensions/rng/distributions.hpp |
Declares sampler registration functions. |
dpnp/libtensor/extensions/rng/common.hpp |
Provides dispatch and kernel submission helpers. |
dpnp/libtensor/extensions/rng/CMakeLists.txt |
Configures the header-only RNG extension. |
dpnp/libtensor/AGENTS.md |
Documents RNG architecture and stream invariants. |
dpnp/CMakeLists.txt |
Includes the RNG extension in the build. |
dpnp/AGENTS.md |
Adds Python RNG implementation guidance. |
doc/reference/random.rst |
Documents modern RNG APIs alongside legacy functions. |
CHANGELOG.md |
Records the new device RNG API. |
Review effort: Balanced
This comment was marked as off-topic.
This comment was marked as off-topic.
antonwolfy
left a comment
antonwolfy
left a comment
There was a problem hiding this comment.
The sampling core is right: the stream matches Philox4x32-10 bit for bit on GPU and CPU, and it passes PractRand and TestU01. The problems are at the edges, in the Python layer, the docs and the tests. Copilot's five points all hold; of the inline comments below, I'd fix the overlap check, the device claim in the docs and the skipped CuPy tests before merge.
Open bot points I confirm:
- Copilot,
lockdeadlock (_bit_generator.py:110). New evidence: a thread runningwith bg.lock: bg.statewas still alive after a 3 s join. - Copilot,
random_raw(output=False)rewind (_bit_generator.py:242). New evidence:random_raw(16), thenrandom_raw(-16, output=False), thenrandom_raw(16)returns the same 16 values, while NumPy'sPhilox.random_raw(-16, output=False)leaves the counter unchanged.random_raw((2, -8), output=False)rewinds too, andrandom_raw((2**40, 2**40), output=False)wrapsnumpy.prodto 0 and leaves the state alone. - Copilot, key read after the lock is released (
_bit_generator.py:391). - Copilot,
LARGE_RANGES(rng/CMakeLists.txt:40). It takes more than 1.7e10 elements to reach. - Copilot, float32
uniformrange (_generator.py:129). New evidence:uniform(-3e38, 3e38, size=1000, dtype=numpy.float32)returns 1000infs.
What I ran, on a Panther Lake U iGPU (Level Zero and OpenCL) and the OpenCL CPU device, with the PR built in place:
test_random_generator.py: 102 passed on level_zero:gpu; 100 passed and 2 skipped (the two-device tests) on opencl:cpu. The newtest_sycl_queue.pyandtest_usm_type.pytests,third_party/cupy/random_tests/as shipped,test_random.pyandtest_random_state.pypass on both, with no pytest warnings. The build had no compiler warnings.- A NumPy implementation of Philox4x32-10, checked against the Random123 known-answer vectors, reproduces
random_rawbit for bit in 274 checks on both devices: sizes from 1 to 220+3, a counter carrying into the high word, the wrap at 2128, consecutive calls andoutput=False.randomanduniformmatch its float conversion bit for bit. - PractRand 0.94: 256 GB of the GPU stream with no anomaly. 4-16 GB each, also with none: the CPU stream; keys 0-3 with the same counter, seeds 0-3 and four
spawnchildren, each interleaved by word and by chunk; and streams of 1-, 9- and 17-value calls (built with the reference, which matches the device). - TestU01 1.2.3: SmallCrush passes on the GPU stream, the CPU stream and the interleaved related keys. Crush passes all 144 tests on two seeds; a third seed had p = 6.4e-4 (CouponCollector) and 0.9992 (CollisionPermut), which the other two seeds didn't repeat. BigCrush is still running.
- Distributions, 1e8 draws per case on both devices: KS and chi-square for
random,uniform(ranges up to ±1e30 in float32 and ±1e300 in float64),standard_normalandnormal, in float32 and float64, all in range. Over 6.7e9 normals there is no correlation between the two values of a Box-Muller pair or between neighbouring work-items.
Not inline:
uniformskips the checks of its host parameters when any parameter is on the device.uniform(dpnp.zeros(1), numpy.array([numpy.inf]))returns[inf], where NumPy raisesOverflowError: Range exceeds valid bounds._check_arrays(_generator.py:142-152) returnsNonefor every parameter as soon as one is a device array, whilenormalchecks each parameter on its own (normal(dpnp.zeros(1), -1.0)raises). Suggestion: check that the host parameters are finite even when the range can't be computed on the host.- The output ignores the parameters' USM type, and the new
dpnp/AGENTS.mdline says something else about temporaries.normal(dpnp.zeros(4, usm_type="shared"), 1.0).usm_typeis"device". The line says "Temporaries take the output'susm_type", but_to_kernel_param(_generator.py:360) callsdpnp.asarraywithoutusm_type, so a parameter given as a list becomes a"device"array even withusm_type="shared". Suggestion: passusm_type=usm_typethere, and either follow the inputs' USM type as other dpnp functions do or list the keyword-only rule under "Differences from NumPy". - The method docstrings don't describe
dtype,usm_typeandout. They appear only as one-lineLimitations, and the methods have noParameters,ReturnsorExamples, unlike the legacydpnp.random.uniform.random.rstdoesn't listGenerator.spawn,Generator.bit_generator,BitGenerator.random_raworBitGenerator.state, and thedevicetext at_bit_generator.py:87refers todpctl.tensor.Device, which isdpnp.tensor.Deviceelsewhere in dpnp now. - Nits:
- Nit: the
_random_rawdoc string (uniform.cpp:46-47) says "low half first", butBitGen::next_uint64returns(first_word << 32) | second_word, so the first 32-bit word is the high half. - Nit:
NormalSampler::operator()(kernels/samplers.hpp) lacks the#pragma clang fp contract(off)thatUniformSamplerhas, soloc + scale * zmay become an FMA on one device and not on another (icpx defaults to-ffp-contract=fast-honor-pragmas). - Nit: the Sphinx spelling step reports 41 new words from these docstrings (Philox, BitGenerator, SeedSequence, ziggurat, ...); they belong in
doc/known_words.txt. - Non-blocking:
pickleandcopy.deepcopyof aGeneratoror aBitGeneratorfail withTypeError(from theSyclQueueand theLock), andadvance/jumpedare absent; NumPy has all of them. Worth a line under "Differences from NumPy" until a follow-up adds them.
- Nit: the
Checked, no issues
- Engine setup:
{key, counter_lo, counter_hi + w}maps to oneMKL's key k0/k1 and counter c0-c3 the way_launchassumes, and onegenerateonphilox4x32x10<8>or<16>yields the words of 8 or 16 calls on<1>. Raw and uniform values are bit-identical between the GPU (blocked) and CPU (strided) layouts. - Stream separation: each work-item owns 264 counters, and the carry from the low word keeps the blocks disjoint up to the wrap at 2128.
- Lifetimes and validation:
keep_args_aliveholds the result and the parameters until the kernel ends; the queue, writability, contiguity, size, dtype, engine and seed-length checks are in place. - fp64: the double kernels are guarded and split per kernel, and the float paths use only float literals.
- NumPy parity: size normalisation, dtype handling, the
scale < 0,-0.0, NaN andNonechecks, broadcasting and theExecutionPlacementErrorcases match NumPy. - Statistical tests: the mean and variance checks of
test_momentsgave no failure over 200 other seeds. - CI: GitHub Actions and the required Jenkins lanes are green. The failed Jenkins Windows py3.14 wheel build is an icx out-of-memory crash in
copy_and_cast_usm_to_usm.cpp, which this PR doesn't touch.
Add the dpnp.libtensor._rng_impl extension, which samples on the device from oneMKL's device philox4x32x10 engine, and the Generator API of NumPy on top of it: BitGenerator, Philox4x3210, SeedSequence (NumPy's), default_rng, and the random, uniform, standard_normal and normal methods. A bit generator keeps its state on the host and every launch rebuilds its engines from it, so the values depend only on the seed and the output size, not on the device. The legacy RandomState and module functions are unchanged.
9fef498 to
1bd3cee
Compare
- Make the bit generator lock reentrant and read the Philox key under it - Validate `size` in `random_raw(output=False)` and count values without overflow - Drop the seed sequence of a Philox4x3210 given its key, as NumPy does - Check the `uniform` range in the sampling dtype, and host bounds next to device ones - Raise TypeError for parameters that don't cast to the sampling dtype - Follow the USM type of the array parameters when `usm_type` is None - Accept a usm_ndarray and, for `random` and `standard_normal`, an F-contiguous `out` - Reject a parameter array overlapping the result, and build the module with LARGE_RANGES - Take the float32 Box-Muller u1 from a whole 32-bit word, and keep FMA contraction off in the normal sampler - Correct the device-independence claims, document the output-size dependence, and expand the method docstrings - Enable the CuPy Generator and Philox tests, skipping the classes of later distributions and bit generators
|
Thanks for the thorough review. All of the inline points are addressed in cab62a2, with replies on each thread. On the points from the review body:
The six PRs above are rebased onto this one, carrying the changes where they apply: |
|
BigCrush (TestU01 1.2.3) has finished on the GPU stream: all 160 statistics of its 106 tests passed, over 3.57e11 32-bit words. That completes the randomness testing from my review: no battery found a problem with the generator. |
antonwolfy
left a comment
antonwolfy
left a comment
There was a problem hiding this comment.
Re-checked at this head on level_zero:gpu and opencl:cpu: all the earlier points are fixed, and the tests pass on both devices. A few new points inline: the first two are worth fixing, the rest are nits.
| // no fused multiply-add, to round as NumPy does | ||
| #pragma clang fp contract(off) | ||
| const T low = low_[i]; | ||
| return low + (high_[i] - low) * bg.template next_uniform<T>(); |
There was a problem hiding this comment.
On opencl:cpu, uniform over a tiny range returns values above high.
uniform(0, 1.5 * finfo(float32).tiny, 4099, dtype=float32) on opencl:cpu returns values up to 1.333 × high, while level_zero:gpu stays below high. The CPU output equals k * (range * 2**-24) bit for bit, with k = w >> 8, instead of (k * 2**-24) * range: the two products are reassociated, so range * 2**-24 underflows into a denormal and is rounded. float64 does the same with 2**-53.
- Any range below 2**-102 (about 2e-31) in float32 or 2**-969 in float64 can give other values than the GPU, which contradicts "the uniform floats are the same on every device" (
random.rst:31-33). - A range below about 4 ×
tinygives values outside[low, high); at 7 ×tinythe largest is still 1.14 ×high.
_rng_impl is built with icpx's default fast floating-point model: its compile lines add only -fno-approx-func -fno-finite-math-only. fp contract(off) turns off contraction but still allows reassociation. libtensor builds its precision-sensitive sources with -fno-fast-math (dpnp/libtensor/source/CMakeLists.txt:324-329). This sampler is unchanged since the first commit.
Suggestion: build the rng sources with -fno-fast-math (or add #pragma clang fp reassociate(off) here), and add a test that compares uniform(0, 1.5 * tiny) with NumPy's low + (high - low) * u, where u is random() of the same seed.
|
|
||
| def test_float32_normal_reference(self): | ||
| # float32 normals take u1 from a whole 32-bit word | ||
| key, counter, n = 0x0123456789ABCDEF, 3, 32 |
There was a problem hiding this comment.
This test also passes with the old 24-bit u1.
At counter=3 all 16 u1 words are far from 0, where the mappings differ by less than 1e-6 relative in the output, well inside rtol=1e-5. Against the device output on level_zero:gpu and opencl:cpu, the test still passes when the reference uses the old ((w >> 8) + 0.5) * 2**-24, or float32(w) * 2**-32 without the + 0.5.
With counter=14589945 (same key and n), the second u1 word of work-item 0 is 6, so elements 8 and 12 come from u1 = 6.5 * 2**-32, where the Box-Muller radius is 6.37. There the current mapping passes on both devices, and the two mutants fail, with relative differences of 8.3e-2 and 2.0e-3.
Suggestion: use that counter, so the test covers the tail it is meant to pin.
| Python integer when `size` is ``None``. With ``output=False``, the | ||
| state advances by the number of elements of `size`, and an invalid | ||
| `size` raises as it does with ``output=True``. NumPy advances by | ||
| the sum of the dimensions of `size` instead, and ignores a negative |
There was a problem hiding this comment.
NumPy doesn't ignore a negative dimension here, only a negative total.
With output=False NumPy advances by np.asarray(size).sum() (numpy/random/_common.pyx:88), so a negative dimension lowers the count: with Philox(1), random_raw((5, -2), output=False) advances by 3 values, while (2, -8) and (-16,) don't advance it.
Suggestion: "NumPy advances by the sum of the dimensions of size instead, and not at all when that sum is negative."
| checked as in NumPy. Parameters in :class:`dpnp.ndarray` are | ||
| not, since that would wait for the device. | ||
| Invalid ones give undefined values, such as NaN. | ||
| * Generators and bit generators cannot be pickled or copied, and the bit |
There was a problem hiding this comment.
copy.copy works; only pickle and copy.deepcopy fail.
copy.copy of a Generator returns one that shares the bit generator, and copy.copy of a Philox4x3210 an independent one in the same state, as in NumPy (it also shares the original's lock, which NumPy's copy doesn't). pickle.dumps and copy.deepcopy of either raise TypeError. _generator.py:202 says the same as this line.
Suggestion: "cannot be pickled or deep-copied".
| def_sampler<rng_ns::RawSampler, RawTypes>( | ||
| m, "_random_raw", | ||
| "Fill the result with the raw 64-bit output of the engine, or with " | ||
| "its 32-bit halves, high half first", |
There was a problem hiding this comment.
The 32-bit output isn't the halves of the 64-bit one.
Element i comes from work-item i % ceil(n / 8) for both dtypes, so a uint32 draw of 2 * n elements equals the high-first halves of a uint64 draw of n elements only for n <= 4. Each matches the Philox reference on its own.
Suggestion: "or with one 32-bit word per element".
| oneAPI | ||
| ord | ||
| orthonormal | ||
| Philox |
There was a problem hiding this comment.
Nit: Philox belongs after Penrose.
On the 41 words from my review: that was the number of new warning lines from the spelling step, not of words. At this head the step reports 10 lines more than master. The only new word is rng, from the dpnp.random.default_rng page; oneMKL, sycl and reproducibility are flagged on master already. So rng is the only word left to add.
| with bg.lock: | ||
| bg.state = bg.state | ||
|
|
||
| t = threading.Thread(target=_use_state) |
There was a problem hiding this comment.
Nit: make this thread a daemon.
If the lock stops being reentrant, the assert fails but t stays blocked, and as a non-daemon thread it keeps the pytest process from exiting. threading.Thread(target=_use_state, daemon=True) avoids that.
| @bit_generator.setter | ||
| def bit_generator(self, value): | ||
| if not isinstance(value, BitGenerator): | ||
| raise TypeError( |
There was a problem hiding this comment.
Minor: a few lines of the Python layer have no test.
The coverage job at this head shows _generator.py at 96.6% and _bit_generator.py at 95.3%. The lines no test reaches:
- here,
_generator.py:238: theTypeErrorof thebit_generatorsetter, e.g.rng.bit_generator = 5(the same check in__init__is covered); _bit_generator.py:154: a seed given as a dpnp array, which goes throughdpnp.asnumpy;_generator.py:251and_bit_generator.py:192: thesycl_deviceproperties, used only in a docstring example;_generator.py:318-319: the C-onlyoutcheck. No method of this PR reaches it, sincerandomandstandard_normaltake the path that allows an F-contiguousout;standard_gammareaches it in Add continuous distributions todpnp.random.Generator#3158, so its test can land there;__repr__and__str__(_generator.py:220, 223,_bit_generator.py:158, 161), and theNotImplementedErrorof the baseBitGenerator.stateand_launch(_bit_generator.py:197, 212).
Suggestion: tests for the setter error, the dpnp-array seed and sycl_device; the rest is optional.
| { | ||
| } | ||
|
|
||
| void operator()(sycl::id<1> id) const |
There was a problem hiding this comment.
Minor: the kernel headers in rng/kernels/ count in the coverage report.
scripts/gen_coverage.py:325 drops the SYCL kernel headers from the C++ report with -ignore-filename-regex=.*/libtensor/include/kernels/.*\.hpp$, since host coverage can't see code that runs in a kernel (#2623, #2816). These headers are in extensions/rng/kernels/, which the pattern doesn't match, so Coveralls lists all four:
- only their host lines count as covered: the functor constructors and the
n_wihelper, 13 lines; - the 105 device lines, this
operator()body and the others, count as missed:engines.hppanddistributions.hppare at 0%, this file at 18% andsamplers.hppat 27%.
That is most of this PR's -0.2%: without these headers, line coverage would be about 0.26 points higher.
Suggestion: widen the pattern to .*/libtensor/(include|extensions/[^/]+)/kernels/.*\.hpp$, or move the headers to libtensor/include/kernels/rng/. The latter ships them with dpnp.get_include(), where they would need oneMKL's device RNG headers.
|
|
||
| #include <sycl/sycl.hpp> | ||
|
|
||
| #include <oneapi/mkl/rng/device.hpp> |
There was a problem hiding this comment.
Minor: with DPNP_USE_ONEMATH, this could include oneMath's device RNG header.
_rng_impl always includes oneMKL's device RNG API, while blas, fft, lapack and sparse switch to oneMath with DPNP_USE_ONEMATH. oneMath's device API (oneapi/math/rng/device.hpp) is header-only, generic SYCL code. At the v0.9 tag that dpnp pins (CMakeLists.txt:274) and at current develop, it has the same engines and distributions as oneMKL 2026.2's. Across the series this extension uses only philox4x32x10<N>, mrg32k3a (#3161), bits<std::uint32_t> and generate, and oneMath has all of them. Its distributions don't matter here, since the samplers are dpnp's own. The only part missing from oneMath is oneMKL's engine_helpers.hpp (engine_descriptor), which isn't used.
This matters mainly for the CUDA and HIP builds, which go through oneMath. I haven't checked whether oneMKL's header compiles for those targets.
Suggestion: on USE_ONEMATH, include <oneapi/math/rng/device.hpp>, alias oneapi::math::rng::device instead of oneapi::mkl::rng::device, take the include directories from ${ONEMATH_LIB}, and run the Philox reference tests on a oneMath build to confirm the streams are the same. This is fine for a follow-up PR.
| a :class:`dpnp.random.Generator` is returned unaltered. | ||
|
|
||
| Default: ``None``. | ||
| device : {None, string, SyclDevice, SyclQueue, Device}, optional |
There was a problem hiding this comment.
The device and sycl_queue texts here aren't the standard ones.
dpnp/AGENTS.md:35 asks to copy them from dpnp.empty, as BitGenerator's device (_bit_generator.py:119) and RandomState do. This one doesn't list what device accepts (a filter selector string, a SyclDevice, a SyclQueue or a dpnp.tensor.Device), and sycl_queue doesn't say that None takes the queue from device or uses a default one. The same holds for Philox4x3210 (_bit_generator.py:353, 357), whose device and sycl_queue only say "See :class:dpnp.random.BitGenerator.", and for the sycl_queue of BitGenerator (_bit_generator.py:128).
Suggestion: use the dpnp.empty text in all three, keeping the extra sentence here ("Ignored when seed is a bit generator or a generator.").
| BitGenerator | ||
| BitGenerator.random_raw | ||
| BitGenerator.spawn | ||
| BitGenerator.state |
There was a problem hiding this comment.
Some documented properties have no page.
dpnp's class pages don't list their members (the RandomState page doesn't either), so a member is rendered only when this file lists it. Generator.sycl_queue and Generator.sycl_device, and BitGenerator.lock, BitGenerator.seed_seq, BitGenerator.sycl_queue and BitGenerator.sycl_device have docstrings but aren't listed, so they appear nowhere in the rendered docs. lock matters most: its docstring is the only place that says the lock is reentrant and how to use it. ndarray.rst:73-75 lists ndarray.sycl_device, ndarray.sycl_queue and ndarray.usm_type.
Suggestion: add them to the Generator and BitGenerator autosummary lists.
|
|
||
| class BitGenerator: | ||
| """ | ||
| Base class for the bit generators of :class:`dpnp.random.Generator`. |
There was a problem hiding this comment.
Nit: the "For full documentation" line is missing.
Generator and its methods have it, and NumPy has this class. dpnp/AGENTS.md:31 puts it after the summary: For full documentation refer to :obj:`numpy.random.BitGenerator`.
|
|
||
| def spawn(self, n_children): | ||
| """ | ||
| Create new independent child generators on the same SYCL queue. |
There was a problem hiding this comment.
n_children isn't described.
This docstring has only the summary and the reference line, while BitGenerator.spawn (_bit_generator.py:301) has Parameters and Returns. So the rendered Generator.spawn page doesn't say what n_children is or what is returned.
Suggestion: add the same Parameters and Returns, with out : list of Generator.
| * Generators and bit generators cannot be pickled or copied, and the bit | ||
| generators have no ``advance`` or ``jumped`` method. | ||
|
|
||
| Random Generator |
There was a problem hiding this comment.
Nit: this heading and the legacy "Random generator" (line 181) differ only by one capital letter.
The page's table of contents shows "Random Generator" for the new API and "Random generator" for RandomState under "Legacy random generation". The other headings on this page are in sentence case ("Accessing the bit generator and spawning", "Bit generators", "Legacy random generation"); "Random Generator" follows NumPy's page title.
Suggestion: rename one of them, e.g. the legacy one to "RandomState", so the two sections are easy to tell apart.
| // Most of the sampling algorithms below are adapted from NumPy's | ||
| // numpy/random/src/distributions/distributions.c, which carries this notice: | ||
| // | ||
| // Copyright (c) 2005-2025, NumPy Developers. |
There was a problem hiding this comment.
This NumPy notice isn't needed in this PR.
The code below it is Param and the Box-Muller StandardNormal, which isn't NumPy's algorithm: NumPy's Generator.standard_normal uses the ziggurat, and the legacy gauss the polar method. So "Most of the sampling algorithms below are adapted from ..." isn't true yet, and the notice isn't needed here.
And I believe we need to follow the procedure and get approval if we would like to add the notice.
| bg = Philox4x3210(key=key, counter=counter) | ||
| res = bg.random_raw(n) | ||
| assert res.dtype == dpnp.uint64 | ||
| assert_array_equal(dpnp.asnumpy(res), _reference_raw(key, counter, n)) |
There was a problem hiding this comment.
dpnp.asnumpy() is not needed here and below:
| assert_array_equal(dpnp.asnumpy(res), _reference_raw(key, counter, n)) | |
| assert_array_equal(res, _reference_raw(key, counter, n)) |
This PR is the first of a series splitting #3147 into smaller pieces. It adds the new framework for device-side random number generation and NumPy's
GeneratorAPI, with only the base distributionsThe new
dpnp.libtensor._rng_implextension samples on the device using the oneMKL device RNG API (header-only, so no oneMKL RNG library is linked). A bit generator keeps its state on the host, and each kernel launch rebuilds its engines from that state, so sampling never waits on the device. For a given seed, output size and dtype, the raw bits and the uniform floats are the same on every device, and the normal samples agree to a few ULPAdds
dpnp.random.Generatorwithrandom,uniform,standard_normal,normalandspawn, anddefault_rng; withusm_type=None(the default) the samples take the coerced USM type of the array parametersdpnp.random.BitGenerator(withrandom_raw,spawn,stateandlock),SeedSequence(NumPy's) and thePhilox4x3210bit generator on oneMKL's devicephilox4x32x10The legacy
RandomStateand the module-level functions are untouched until later in the seriesSeries:
dpnp.random.Generatoron the oneMKL device RNG API #3157 (this PR): framework,Philox4x3210and the base distributionsdpnp.random.Generator#3158: Continuous distributionsdpnp.random.Generator#3159: Discrete distributionsdpnp.random.Generator#3160: Multivariate distributions and permutationsPCG64DXSMandMRG32k3abit generators todpnp.random#3161:PCG64DXSMandMRG32k3abit generatorsdpnp.random.RandomStateon the device RNG #3162:RandomStateand the module functions on the device RNG