diff --git a/CHANGELOG.md b/CHANGELOG.md index 8acc4098d6ad..cc4393ed723e 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -10,6 +10,10 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 * Added support for boolean input in `dpnp.positive` and the `+` operator of `dpnp.ndarray` [#3128](https://github.com/IntelPython/dpnp/pull/3128) * Added `dpnp.top_k` to return the `k` largest or smallest elements and their indices along an axis, with a `sorted` keyword to return them in sorted order [#3133](https://github.com/IntelPython/dpnp/pull/3133) +* Added the `dpnp.random.Generator` API with `default_rng`, `BitGenerator`, `SeedSequence` and the `Philox4x3210` bit generator, on the oneMKL device RNG API [#3147](https://github.com/IntelPython/dpnp/pull/3147) +* Added the `dpnp.random.PCG64DXSM` bit generator, which matches `numpy.random.PCG64DXSM` for raw values, `random`, `uniform` and `bytes` [#3147](https://github.com/IntelPython/dpnp/pull/3147) +* Added `dpnp.random.get_bit_generator` and `dpnp.random.set_bit_generator`, and support for a bit generator as the seed of `dpnp.random.RandomState` [#3147](https://github.com/IntelPython/dpnp/pull/3147) +* Added the `dpnp.random.MRG32k3a` bit generator [#3147](https://github.com/IntelPython/dpnp/pull/3147) ### Changed @@ -18,6 +22,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 * Merged the `ufunc` and `indexing` extensions into `dpnp.libtensor._tensor_elementwise_impl` and `dpnp.libtensor._tensor_impl`, while `statistics` and `window` extensions became part of libtensor [#3131](https://github.com/IntelPython/dpnp/pull/3131) * Moved `blas`, `fft`, `lapack`, `sparse` and `vm` extensions into `dpnp/libtensor/extensions` [#3131](https://github.com/IntelPython/dpnp/pull/3131) * Moved the libtensor headers shipped with dpnp from `dpnp/tensor/libtensor/include` to `dpnp/libtensor/include`, and changed `dpnp.get_include` to return that directory [#3131](https://github.com/IntelPython/dpnp/pull/3131) +* Reimplemented `dpnp.random.RandomState` and the `dpnp.random` functions on device kernels, giving the same values on every device. `device`, `usm_type` and `sycl_queue` are now keyword-only, `randint(dtype=int)` returns `int64` and `shuffle` requires a `dpnp.ndarray` [#3147](https://github.com/IntelPython/dpnp/pull/3147) ### Deprecated @@ -29,6 +34,8 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 ### Removed * Removed the unused `FourZeroOffsets_Indexer`, `FixedDimStridedIndexer` and `ThreeOffsets_FixedDimStridedIndexer` from the libtensor header `utils/offset_utils.hpp` [#3131](https://github.com/IntelPython/dpnp/pull/3131) +* Removed the legacy backend implementation of `dpnp.random`, which no longer needs `onemkl-sycl-rng` [#3147](https://github.com/IntelPython/dpnp/pull/3147) +* Removed the unused parts of the legacy `dpnp_backend_c` library, including its empty GTest suite, the `DPNP_BACKEND_TESTS` CMake option and the NumPy fallback helper `call_origin` [#3147](https://github.com/IntelPython/dpnp/pull/3147) ### Fixed diff --git a/CMakeLists.txt b/CMakeLists.txt index 252698dc1fea..68beb804774e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -42,7 +42,6 @@ option( "Build dpnp tensor pybind11 offloading extensions with coverage instrumentation" OFF ) -option(DPNP_BACKEND_TESTS "Enable building of DPNP backend test suite" OFF) option( DPNP_WITH_REDIST "Build DPNP assuming DPC++ redistributable is installed into Python prefix" diff --git a/conda-recipe-cf/meta.yaml b/conda-recipe-cf/meta.yaml index 1fa8ab3fecb7..9524ddb0a60c 100644 --- a/conda-recipe-cf/meta.yaml +++ b/conda-recipe-cf/meta.yaml @@ -56,7 +56,6 @@ requirements: - {{ pin_compatible('onemkl-sycl-blas', min_pin='x.x', max_pin='x') }} - {{ pin_compatible('onemkl-sycl-dft', min_pin='x.x', max_pin='x') }} - {{ pin_compatible('onemkl-sycl-lapack', min_pin='x.x', max_pin='x') }} - - {{ pin_compatible('onemkl-sycl-rng', min_pin='x.x', max_pin='x') }} - {{ pin_compatible('onemkl-sycl-vm', min_pin='x.x', max_pin='x') }} - numpy - ocl-icd-system # [linux] diff --git a/conda-recipe/meta.yaml b/conda-recipe/meta.yaml index 18d1c0d11629..318a7ff78102 100644 --- a/conda-recipe/meta.yaml +++ b/conda-recipe/meta.yaml @@ -69,7 +69,6 @@ requirements: - {{ pin_compatible('onemkl-sycl-blas', min_pin='x.x', max_pin='x') }} - {{ pin_compatible('onemkl-sycl-dft', min_pin='x.x', max_pin='x') }} - {{ pin_compatible('onemkl-sycl-lapack', min_pin='x.x', max_pin='x') }} - - {{ pin_compatible('onemkl-sycl-rng', min_pin='x.x', max_pin='x') }} - {{ pin_compatible('onemkl-sycl-vm', min_pin='x.x', max_pin='x') }} - {{ pin_compatible('onemkl-sycl-sparse', min_pin='x.x', max_pin='x') }} - numpy diff --git a/conda-recipe/rattler_recipe.yaml b/conda-recipe/rattler_recipe.yaml index eeb25ca31b4e..0a0cca5f862b 100644 --- a/conda-recipe/rattler_recipe.yaml +++ b/conda-recipe/rattler_recipe.yaml @@ -70,7 +70,6 @@ requirements: - ${{ pin_compatible('onemkl-sycl-blas', lower_bound='x.x', upper_bound='x') }} - ${{ pin_compatible('onemkl-sycl-dft', lower_bound='x.x', upper_bound='x') }} - ${{ pin_compatible('onemkl-sycl-lapack', lower_bound='x.x', upper_bound='x') }} - - ${{ pin_compatible('onemkl-sycl-rng', lower_bound='x.x', upper_bound='x') }} - ${{ pin_compatible('onemkl-sycl-vm', lower_bound='x.x', upper_bound='x') }} - ${{ pin_compatible('onemkl-sycl-sparse', lower_bound='x.x', upper_bound='x') }} - numpy diff --git a/doc/known_words.txt b/doc/known_words.txt index b8f0ec87e530..64e2bb159d18 100644 --- a/doc/known_words.txt +++ b/doc/known_words.txt @@ -42,9 +42,11 @@ Hanning histogrammed Hypergeometric kwargs +Kelton iaxis iinfo Infs +L'Ecuyer intp ints iterable @@ -86,15 +88,19 @@ prepended prepending representable resampling +RngStreams runtimes scikit se signbit signum +Simard sinc subarray subarrays subclasses +substream +substreams subtype SyclDevice SyclQueue diff --git a/doc/reference/random.rst b/doc/reference/random.rst index 55ac0271307a..d528920f5267 100644 --- a/doc/reference/random.rst +++ b/doc/reference/random.rst @@ -3,10 +3,185 @@ Random sampling (:mod:`dpnp.random`) ==================================== +.. hint:: `NumPy API Reference: Random sampling `_ + +DPNP draws random numbers on the SYCL device, by default with the +counter-based Philox4x32x10 generator of the oneMKL device RNG API, or with +NumPy's PCG64DXSM generator or the MRG32k3a generator of the oneMKL device +RNG API. It provides both NumPy random APIs: the :class:`Generator` API recommended for new code, +and the legacy :class:`RandomState` API with the module-level functions +built on it. + +.. code-block:: python + + import dpnp + + rng = dpnp.random.default_rng(seed=42) + x = rng.standard_normal((3, 4)) # on the default device + y = rng.integers(0, 10, size=5, usm_type="shared") + + rng = dpnp.random.default_rng(42, device="cpu") + +A generator draws all its samples on one SYCL queue, chosen with the +``device`` or ``sycl_queue`` keyword when it is created, and the samples +depend only on the seed, not on the device, except that the discrete +distributions also depend on its support of double precision. All samples are +:class:`dpnp.ndarray`, zero-dimensional when ``size`` is ``None``. + +Differences from NumPy +---------------------- + +* The same seed gives the same stream on every device, but for the discrete + distributions noted below. The streams differ from NumPy's, except that + :class:`PCG64DXSM` gives NumPy's raw values, :meth:`Generator.random`, + :meth:`Generator.uniform` and :meth:`Generator.bytes`, and can exchange + states with :class:`numpy.random.PCG64DXSM`. +* The floating-point distributions take a ``dtype`` keyword, as in CuPy, + which defaults to the default floating-point type of the device. All + methods take a ``usm_type`` keyword. +* Parameters on the host (Python scalars, lists and NumPy arrays) are + checked as in NumPy. In :class:`Generator`, parameters in + :class:`dpnp.ndarray` are not, since that would wait for the device. + Invalid ones give undefined values, such as NaN or the smallest ``int64``, + but the kernels handle them safely: they never loop without end or access + memory out of bounds. Floating-point arrays given for integer parameters + are converted as by :meth:`dpnp.ndarray.astype`, whose result for NaN and + out-of-range values is undefined. +* The discrete distributions return ``int64`` values. They compute in double + precision on devices supporting it, and in single precision otherwise. + There, the streams differ from the ones of devices with double precision, + and some samples are only multiples of a power of two: the ones above + ``2**24`` of ``geometric``, ``logseries`` and ``zipf``, and the ones of + ``poisson`` and ``binomial`` with a standard deviation above about + ``2**22``. +* ``zipf`` follows the Zipf distribution truncated to the largest ``int64``. + NumPy's rejection test rounds to accepting every sample above ``2**53``, + which makes its tail too heavy for ``a`` close to 1. +* :meth:`RandomState.randint` and :func:`randint` return ``int64`` values by + default, as do :meth:`RandomState.random_integers`, + :func:`random_integers` and :meth:`RandomState.tomaxint`, whose values + reach the largest ``int64``. NumPy returns a C ``long`` instead, which is + ``int32`` on Windows. +* ``hypergeometric`` requires ``ngood`` and ``nbad`` below ``10**9``, also in + :class:`RandomState`. +* :meth:`Generator.shuffle` and :func:`shuffle` only accept a + :class:`dpnp.ndarray`, as in CuPy. + +Random Generator +---------------- + +.. autosummary:: + :toctree: generated/ + :nosignatures: + + default_rng + Generator + +Simple random data +^^^^^^^^^^^^^^^^^^ + +.. autosummary:: + :toctree: generated/ + :nosignatures: + + Generator.integers + Generator.random + Generator.choice + Generator.bytes + +Permutations +^^^^^^^^^^^^ + +.. autosummary:: + :toctree: generated/ + :nosignatures: + + Generator.shuffle + Generator.permutation + Generator.permuted + +Distributions +^^^^^^^^^^^^^ + +.. autosummary:: + :toctree: generated/ + :nosignatures: + + Generator.beta + Generator.binomial + Generator.chisquare + Generator.dirichlet + Generator.exponential + Generator.f + Generator.gamma + Generator.geometric + Generator.gumbel + Generator.hypergeometric + Generator.laplace + Generator.logistic + Generator.lognormal + Generator.logseries + Generator.multinomial + Generator.multivariate_hypergeometric + Generator.multivariate_normal + Generator.negative_binomial + Generator.noncentral_chisquare + Generator.noncentral_f + Generator.normal + Generator.pareto + Generator.poisson + Generator.power + Generator.rayleigh + Generator.standard_cauchy + Generator.standard_exponential + Generator.standard_gamma + Generator.standard_normal + Generator.standard_t + Generator.triangular + Generator.uniform + Generator.vonmises + Generator.wald + Generator.weibull + Generator.zipf + +Bit generators +-------------- + +.. autosummary:: + :toctree: generated/ + :nosignatures: + + BitGenerator + Philox4x3210 + PCG64DXSM + MRG32k3a + SeedSequence + +Legacy random generation +------------------------ + .. hint:: `NumPy API Reference: Legacy random generation `_ +:class:`RandomState` keeps the parameter checks of NumPy's legacy one, also +for parameters in :class:`dpnp.ndarray`, as CuPy's does, so that a call with +such parameters waits for the device to check them. The module-level +functions use one :class:`RandomState` per SYCL queue: the +queue given by the ``device`` or ``sycl_queue`` keyword, or else the queue of +a :class:`dpnp.ndarray` parameter, or else the default queue. + +.. autosummary:: + :toctree: generated/ + :nosignatures: + + RandomState + seed + get_state + set_state + get_bit_generator + set_bit_generator + Simple random data ------------------- +^^^^^^^^^^^^^^^^^^ .. autosummary:: :toctree: generated/ @@ -23,9 +198,8 @@ Simple random data choice bytes - Permutations ------------- +^^^^^^^^^^^^ .. autosummary:: :toctree: generated/ @@ -34,9 +208,8 @@ Permutations shuffle permutation - Distributions -------------- +^^^^^^^^^^^^^ .. autosummary:: :toctree: generated/ @@ -77,16 +250,3 @@ Distributions wald weibull zipf - - -Random generator ----------------- - -.. autosummary:: - :toctree: generated/ - :nosignatures: - - RandomState - seed - get_random_state - set_random_state diff --git a/dpnp/AGENTS.md b/dpnp/AGENTS.md index b7d3956d4134..35eb219e84af 100644 --- a/dpnp/AGENTS.md +++ b/dpnp/AGENTS.md @@ -6,7 +6,8 @@ - `dpnp_algo/dpnp_elementwise_common.py`: ufunc wrapper classes (`DPNPUnaryFunc`, `DPNPBinaryFunc`, two-output and specialised variants) - `dpnp_utils/`: shared helpers (`get_usm_allocations`, `map_dtype_to_device`, reduction wrappers) - `linalg/`, `fft/`, `random/`, `scipy/`, `lib/`: subpackages, each exporting through its own `__init__.py` -- Legacy, don't extend: the Cython `dpnp_algo` path and `dpnp_backend_c` (still used by `dpnp.partition` and `dpnp.random`), and the NumPy fallback `call_origin`. New kernels go into `dpnp/libtensor` (see `dpnp/libtensor/AGENTS.md`). +- `random/`: `_bit_generator.py` (`BitGenerator`, `Philox4x3210`, `PCG64DXSM`, `MRG32k3a`; `SeedSequence` is NumPy's), `_generator.py` (`Generator`, which implements every distribution), `dpnp_random_state.py` (`RandomState`, the same samplers with NumPy's legacy checks) and `dpnp_iface_random.py` (module functions, one `RandomState` per queue). The kernels are in `dpnp.libtensor._rng_impl`. +- Legacy, don't extend: the Cython `dpnp_algo` path and `dpnp_backend_c` (still used by `dpnp.partition`). New kernels go into `dpnp/libtensor` (see `dpnp/libtensor/AGENTS.md`). ## Adding a Public Function 1. Implement it in the matching `dpnp_iface_*.py` or subpackage module. Iface modules have no `__all__`. @@ -20,6 +21,7 @@ - Extension calls are ordered with `_manager = dpctl.utils.SequentialOrderManager[exec_q]`: pass `depends=_manager.submitted_events`, then `_manager.add_event_pair(ht_ev, ev)`. Don't add blocking waits. - Use `dpnp.default_float_type(...)` or `map_dtype_to_device(...)` for the default float type. - Deprecated ufuncs use `DPNPDeprecatedUnaryFunc(..., deprecated_msg=...)`. +- A `dpnp.random` distribution is a `Generator` method: it calls `self._check_queues(...)` on its array parameters (raises `ExecutionPlacementError`), checks host parameters as NumPy does, and draws through `_cont`, `_float_sample` or `_discrete_sample` into `BitGenerator._fill`. `RandomState` and the module function delegate to it. Parameters on the device are checked only by `RandomState` (`_check_device`), as the check waits for the device. Temporaries take the output's `usm_type`. - Unsupported NumPy keywords raise `NotImplementedError` (`dpnp.check_limitations`) and are listed under `Limitations` in the docstring. - The Array API spec version is set in three places that change together: `__array_api_version__` in `dpnp/tensor/_array_api.py`, `ARRAY_API_TESTS_VERSION` in `.github/workflows/conda-package.yml` and `doc/reference/array_api.rst` (plus the version test in `dpnp/tests/test_ndarray.py`). Inspection utilities come from `__array_namespace_info__` in `dpnp_array_api_info.py`. - `dpnp.tensor` replaces `dpctl.tensor`, which no longer exists here; `SyclQueue`, `SyclDevice`, `dpctl.memory` and `SequentialOrderManager` still come from `dpctl`. diff --git a/dpnp/CMakeLists.txt b/dpnp/CMakeLists.txt index e96cf5bbc5a9..edd7fdd628f1 100644 --- a/dpnp/CMakeLists.txt +++ b/dpnp/CMakeLists.txt @@ -229,9 +229,9 @@ add_subdirectory(libtensor/source) add_subdirectory(libtensor/extensions/blas) add_subdirectory(libtensor/extensions/fft) add_subdirectory(libtensor/extensions/lapack) +add_subdirectory(libtensor/extensions/rng) add_subdirectory(libtensor/extensions/vm) add_subdirectory(libtensor/extensions/sparse) add_subdirectory(dpnp_algo) add_subdirectory(dpnp_utils) -add_subdirectory(random) diff --git a/dpnp/backend/AGENTS.md b/dpnp/backend/AGENTS.md index 9066c8c933a6..827834d66de8 100644 --- a/dpnp/backend/AGENTS.md +++ b/dpnp/backend/AGENTS.md @@ -3,5 +3,5 @@ Legacy, don't extend: new kernels and bindings go into `dpnp/libtensor` (see `dpnp/libtensor/AGENTS.md`, which also holds the C++ rules and build options). ## Layout -- `src/`, `include/` and `kernels/dpnp_krnl_*.cpp` build `dpnp_backend_c` (still used by `dpnp.partition` and `dpnp.random`); `tests/` (GTest, `-DDPNP_BACKEND_TESTS=ON`) and `examples/` are not built in CI +- `src/`, `include/` and `kernels/dpnp_krnl_*.cpp` build `dpnp_backend_c` (still used by `dpnp.partition`) - `cmake/Modules/`: vendored `IntelSYCL`, `MKL`, `TBB` and `oneDPL` config modules, used by the whole build diff --git a/dpnp/backend/CMakeLists.txt b/dpnp/backend/CMakeLists.txt index 7a474b29170c..ef0276fea695 100644 --- a/dpnp/backend/CMakeLists.txt +++ b/dpnp/backend/CMakeLists.txt @@ -27,17 +27,7 @@ # THE POSSIBILITY OF SUCH DAMAGE. # ***************************************************************************** -set(DPNP_SRC - kernels/dpnp_krnl_arraycreation.cpp - kernels/dpnp_krnl_common.cpp - kernels/dpnp_krnl_random.cpp - kernels/dpnp_krnl_sorting.cpp - src/dpnp_iface_fptr.cpp - src/memory_sycl.cpp - src/queue_sycl.cpp - src/verbose.cpp - src/dpnp_random_state.cpp -) +set(DPNP_SRC kernels/dpnp_krnl_sorting.cpp src/dpnp_iface_fptr.cpp) set(_trgt "dpnp_backend_c") @@ -69,12 +59,6 @@ if(DPNP_GENERATE_COVERAGE) target_link_options(${_trgt} PRIVATE -fprofile-instr-generate -fcoverage-mapping) endif() -# Link only the MKL SYCL domains actually used by the backend sources -# (BLAS: mkl_blas::scal; RNG: engines/distributions; VM: vector math ops) -target_link_libraries( - ${_trgt} - PUBLIC MKL::MKL_SYCL::BLAS MKL::MKL_SYCL::RNG MKL::MKL_SYCL::VM -) target_link_libraries(${_trgt} PUBLIC oneDPL) if(UNIX) @@ -103,10 +87,6 @@ if(DPNP_WITH_REDIST) set_target_properties(${_trgt} PROPERTIES INSTALL_RPATH "$ORIGIN/../../../") endif() -if(DPNP_BACKEND_TESTS) - add_subdirectory(tests) -endif() - set(_lib_destination dpnp) set(_include_destination dpnp/backend/include) install( diff --git a/dpnp/backend/examples/example10.cpp b/dpnp/backend/examples/example10.cpp deleted file mode 100644 index 86284c2bfe94..000000000000 --- a/dpnp/backend/examples/example10.cpp +++ /dev/null @@ -1,137 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// - Neither the name of the copyright holder nor the names of its contributors -// may be used to endorse or promote products derived from this software -// without specific prior written permission. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -/** - * Example 10. - * - * Possible compile line: - * clang++ -fsycl dpnp/backend/examples/example10.cpp -Idpnp - * -Idpnp/backend/include -Ldpnp -Wl,-rpath='$ORIGIN'/dpnp -ldpnp_backend_c -o - * example10 -lmkl_sycl -lmkl_intel_ilp64 -lmkl_core - */ - -#include -#include - -#include -#include - -#include - -void test_dpnp_random_normal(const size_t size, - const size_t iters, - const size_t seed, - const double loc, - const double scale) -{ - clock_t start, end; - double dev_time_used = 0.0; - double sum_dev_time_used = 0.0; - - double *result = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - - dpnp_rng_srand_c(seed); // TODO: will move - - for (size_t i = 0; i < iters; ++i) { - start = clock(); - dpnp_rng_normal_c(result, loc, scale, size); - end = clock(); - sum_dev_time_used += ((double)(end - start)) / CLOCKS_PER_SEC; - // TODO: cumulative addition error - // div error - } - - dpnp_memory_free_c(result); - - dev_time_used = sum_dev_time_used / iters; - std::cout << "DPNPC random normal: " << dev_time_used << std::endl; - - return; -} - -// TODO: name check -void test_mkl_random_normal(const size_t size, - const size_t iters, - const size_t seed, - const double loc, - const double scale) -{ - clock_t start, end; - double dev_time_used = 0.0; - double sum_dev_time_used = 0.0; - - sycl::queue queue{sycl::gpu_selector()}; - - double *result = - reinterpret_cast(malloc_shared(size * sizeof(double), queue)); - if (result == nullptr) { - throw std::runtime_error("Error: out of memory."); - } - - oneapi::mkl::rng::mt19937 rng_engine(queue, seed); - oneapi::mkl::rng::gaussian distribution(loc, scale); - - for (size_t i = 0; i < iters; ++i) { - start = clock(); - - auto event_out = - oneapi::mkl::rng::generate(distribution, rng_engine, size, result); - event_out.wait(); - - end = clock(); - - sum_dev_time_used += ((double)(end - start)) / CLOCKS_PER_SEC; - // TODO: cumulative addition error - // div error - } - free(result, queue); - - std::cout << "MKL time: "; - dev_time_used = sum_dev_time_used / iters; - std::cout << dev_time_used << std::endl; - - return; -} - -int main(int, char **) -{ - const size_t size = 100000000; - const size_t iters = 30; - - const size_t seed = 10; - const double loc = 0.0; - const double scale = 1.0; - - std::cout << "Normal distr. params:\nloc is " << loc << ", scale is " - << scale << std::endl; - - test_dpnp_random_normal(size, iters, seed, loc, scale); - test_mkl_random_normal(size, iters, seed, loc, scale); - - return 0; -} diff --git a/dpnp/backend/include/dpnp_iface.hpp b/dpnp/backend/include/dpnp_iface.hpp index b7627101aab4..cd75f6d1cfac 100644 --- a/dpnp/backend/include/dpnp_iface.hpp +++ b/dpnp/backend/include/dpnp_iface.hpp @@ -61,8 +61,6 @@ typedef ssize_t shape_elem_type; #include -#include "dpnp_iface_random.hpp" - /** * @defgroup BACKEND_API Backend C++ library interface API * @{ @@ -70,39 +68,6 @@ typedef ssize_t shape_elem_type; * @} */ -/** - * @ingroup BACKEND_API - * @brief SYCL queue device status. - * - * Return 1 if current @ref queue is related to cpu device. return 0 otherwise. - */ -INP_DLLEXPORT size_t dpnp_queue_is_cpu_c(); - -/** - * @ingroup BACKEND_API - * @brief SYCL queue memory allocation. - * - * Memory allocation on the SYCL backend. - * - * @param [in] size_in_bytes Number of bytes for requested memory allocation. - * @param [in] q_ref Reference to SYCL queue. - * - * @return A pointer to newly created memory on SYCL device. - */ -INP_DLLEXPORT char *dpnp_memory_alloc_c(DPCTLSyclQueueRef q_ref, - size_t size_in_bytes); -INP_DLLEXPORT char *dpnp_memory_alloc_c(size_t size_in_bytes); - -INP_DLLEXPORT void dpnp_memory_free_c(DPCTLSyclQueueRef q_ref, void *ptr); -INP_DLLEXPORT void dpnp_memory_free_c(void *ptr); - -INP_DLLEXPORT void dpnp_memory_memcpy_c(DPCTLSyclQueueRef q_ref, - void *dst, - const void *src, - size_t size_in_bytes); -INP_DLLEXPORT void - dpnp_memory_memcpy_c(void *dst, const void *src, size_t size_in_bytes); - /** * @ingroup BACKEND_API * @brief Return a partitioned copy of an array. @@ -127,109 +92,4 @@ INP_DLLEXPORT DPCTLSyclEventRef const size_t ndim, const DPCTLEventVectorRef dep_event_vec_ref); -template -INP_DLLEXPORT void dpnp_partition_c(void *array, - void *array2, - void *result, - const size_t kth, - const shape_elem_type *shape, - const size_t ndim); - -/** - * @ingroup BACKEND_API - * @brief implementation of creating filled with value array function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result1 Output array. - * @param [in] value Value in array. - * @param [in] size Number of elements in array. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_initval_c(DPCTLSyclQueueRef q_ref, - void *result1, - void *value, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_initval_c(void *result1, void *value, size_t size); - -/** - * @ingroup BACKEND_API - * @brief Implementation of ones function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] size Number of elements in the output array. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_ones_c(DPCTLSyclQueueRef q_ref, - void *result, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_ones_c(void *result, size_t size); - -/** - * @ingroup BACKEND_API - * @brief Implementation of ones_like function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] size Number of elements in the output array. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_ones_like_c(DPCTLSyclQueueRef q_ref, - void *result, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_ones_like_c(void *result, size_t size); - -/** - * @ingroup BACKEND_API - * @brief Implementation of zeros function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] size Number of elements in the output array. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_zeros_c(DPCTLSyclQueueRef q_ref, - void *result, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_zeros_c(void *result, size_t size); - -/** - * @ingroup BACKEND_API - * @brief Implementation of zeros_like function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] size Number of elements in the output array. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_zeros_like_c(DPCTLSyclQueueRef q_ref, - void *result, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_zeros_like_c(void *result, size_t size); - #endif // BACKEND_IFACE_H diff --git a/dpnp/backend/include/dpnp_iface_fptr.hpp b/dpnp/backend/include/dpnp_iface_fptr.hpp index 868a31bdaac5..74f8de6b2c6c 100644 --- a/dpnp/backend/include/dpnp_iface_fptr.hpp +++ b/dpnp/backend/include/dpnp_iface_fptr.hpp @@ -61,136 +61,10 @@ */ enum class DPNPFuncName : size_t { - DPNP_FN_NONE, /**< Very first element of the enumeration */ - DPNP_FN_INITVAL, /**< Used in numpy ones, ones_like, zeros, zeros_like impls - */ - DPNP_FN_INITVAL_EXT, /**< Used in numpy ones, ones_like, zeros, zeros_like - impls */ - DPNP_FN_ONES, /**< Used in numpy.ones() impl */ - DPNP_FN_ONES_LIKE, /**< Used in numpy.ones_like() impl */ - DPNP_FN_PARTITION, /**< Used in numpy.partition() impl */ + DPNP_FN_NONE, /**< Very first element of the enumeration */ DPNP_FN_PARTITION_EXT, /**< Used in numpy.partition() impl, requires extra parameters */ - DPNP_FN_RNG_BETA, /**< Used in numpy.random.beta() impl */ - DPNP_FN_RNG_BETA_EXT, /**< Used in numpy.random.beta() impl, requires extra - parameters */ - DPNP_FN_RNG_BINOMIAL, /**< Used in numpy.random.binomial() impl */ - DPNP_FN_RNG_BINOMIAL_EXT, /**< Used in numpy.random.binomial() impl, - requires extra parameters */ - DPNP_FN_RNG_CHISQUARE, /**< Used in numpy.random.chisquare() impl */ - DPNP_FN_RNG_CHISQUARE_EXT, /**< Used in numpy.random.chisquare() impl, - requires extra parameters */ - DPNP_FN_RNG_EXPONENTIAL, /**< Used in numpy.random.exponential() impl */ - DPNP_FN_RNG_EXPONENTIAL_EXT, /**< Used in numpy.random.exponential() impl, - requires extra parameters */ - DPNP_FN_RNG_F, /**< Used in numpy.random.f() impl */ - DPNP_FN_RNG_F_EXT, /**< Used in numpy.random.f() impl, requires extra - parameters */ - DPNP_FN_RNG_GAMMA, /**< Used in numpy.random.gamma() impl */ - DPNP_FN_RNG_GAMMA_EXT, /**< Used in numpy.random.gamma() impl, requires - extra parameters */ - DPNP_FN_RNG_GAUSSIAN, /**< Used in numpy.random.randn() impl */ - DPNP_FN_RNG_GAUSSIAN_EXT, /**< Used in numpy.random.randn() impl, requires - extra parameters */ - DPNP_FN_RNG_GEOMETRIC, /**< Used in numpy.random.geometric() impl */ - DPNP_FN_RNG_GEOMETRIC_EXT, /**< Used in numpy.random.geometric() impl, - requires extra parameters */ - DPNP_FN_RNG_GUMBEL, /**< Used in numpy.random.gumbel() impl */ - DPNP_FN_RNG_GUMBEL_EXT, /**< Used in numpy.random.gumbel() impl, requires - extra parameters */ - DPNP_FN_RNG_HYPERGEOMETRIC, /**< Used in numpy.random.hypergeometric() impl - */ - DPNP_FN_RNG_HYPERGEOMETRIC_EXT, /**< Used in numpy.random.hypergeometric() - impl, requires extra parameters */ - DPNP_FN_RNG_LAPLACE, /**< Used in numpy.random.laplace() impl */ - DPNP_FN_RNG_LAPLACE_EXT, /**< Used in numpy.random.laplace() impl */ - DPNP_FN_RNG_LOGISTIC, /**< Used in numpy.random.logistic() impl */ - DPNP_FN_RNG_LOGISTIC_EXT, /**< Used in numpy.random.logistic() impl, - requires extra parameters */ - DPNP_FN_RNG_LOGNORMAL, /**< Used in numpy.random.lognormal() impl */ - DPNP_FN_RNG_LOGNORMAL_EXT, /**< Used in numpy.random.lognormal() impl, - requires extra parameters */ - DPNP_FN_RNG_MULTINOMIAL, /**< Used in numpy.random.multinomial() impl */ - DPNP_FN_RNG_MULTINOMIAL_EXT, /**< Used in numpy.random.multinomial() impl, - requires extra parameters */ - DPNP_FN_RNG_MULTIVARIATE_NORMAL, /**< Used in - numpy.random.multivariate_normal() impl - */ - DPNP_FN_RNG_MULTIVARIATE_NORMAL_EXT, /**< Used in - numpy.random.multivariate_normal() - impl */ - DPNP_FN_RNG_NEGATIVE_BINOMIAL, /**< Used in numpy.random.negative_binomial() - impl */ - DPNP_FN_RNG_NEGATIVE_BINOMIAL_EXT, /**< Used in - numpy.random.negative_binomial() impl - */ - DPNP_FN_RNG_NONCENTRAL_CHISQUARE, /**< Used in - numpy.random.noncentral_chisquare() - impl */ - DPNP_FN_RNG_NONCENTRAL_CHISQUARE_EXT, /**< Used in - numpy.random.noncentral_chisquare() - impl */ - DPNP_FN_RNG_NORMAL, /**< Used in numpy.random.normal() impl */ - DPNP_FN_RNG_NORMAL_EXT, /**< Used in numpy.random.normal() impl, requires - extra parameters */ - DPNP_FN_RNG_PARETO, /**< Used in numpy.random.pareto() impl */ - DPNP_FN_RNG_PARETO_EXT, /**< Used in numpy.random.pareto() impl, requires - extra parameters */ - DPNP_FN_RNG_POISSON, /**< Used in numpy.random.poisson() impl */ - DPNP_FN_RNG_POISSON_EXT, /**< Used in numpy.random.poisson() impl, requires - extra parameters */ - DPNP_FN_RNG_POWER, /**< Used in numpy.random.power() impl */ - DPNP_FN_RNG_POWER_EXT, /**< Used in numpy.random.power() impl, requires - extra parameters */ - DPNP_FN_RNG_RAYLEIGH, /**< Used in numpy.random.rayleigh() impl */ - DPNP_FN_RNG_RAYLEIGH_EXT, /**< Used in numpy.random.rayleigh() impl, - requires extra parameters */ - DPNP_FN_RNG_SRAND, /**< Used in numpy.random.seed() impl */ - DPNP_FN_RNG_SRAND_EXT, /**< Used in numpy.random.seed() impl, requires extra - parameters */ - DPNP_FN_RNG_SHUFFLE, /**< Used in numpy.random.shuffle() impl */ - DPNP_FN_RNG_SHUFFLE_EXT, /**< Used in numpy.random.shuffle() impl, requires - extra parameters */ - DPNP_FN_RNG_STANDARD_CAUCHY, /**< Used in numpy.random.standard_cauchy() - impl */ - DPNP_FN_RNG_STANDARD_CAUCHY_EXT, /**< Used in numpy.random.standard_cauchy() - impl */ - DPNP_FN_RNG_STANDARD_EXPONENTIAL, /**< Used in - numpy.random.standard_exponential() - impl */ - DPNP_FN_RNG_STANDARD_EXPONENTIAL_EXT, /**< Used in - numpy.random.standard_exponential() - impl */ - DPNP_FN_RNG_STANDARD_GAMMA, /**< Used in numpy.random.standard_gamma() impl - */ - DPNP_FN_RNG_STANDARD_GAMMA_EXT, /**< Used in numpy.random.standard_gamma() - impl, requires extra parameters */ - DPNP_FN_RNG_STANDARD_NORMAL, /**< Used in numpy.random.standard_normal() - impl */ - DPNP_FN_RNG_STANDARD_T, /**< Used in numpy.random.standard_t() impl */ - DPNP_FN_RNG_STANDARD_T_EXT, /**< Used in numpy.random.standard_t() impl, - requires extra parameters */ - DPNP_FN_RNG_TRIANGULAR, /**< Used in numpy.random.triangular() impl */ - DPNP_FN_RNG_TRIANGULAR_EXT, /**< Used in numpy.random.triangular() impl, - requires extra parameters */ - DPNP_FN_RNG_UNIFORM, /**< Used in numpy.random.uniform() impl */ - DPNP_FN_RNG_UNIFORM_EXT, /**< Used in numpy.random.uniform() impl, requires - extra parameters */ - DPNP_FN_RNG_VONMISES, /**< Used in numpy.random.vonmises() impl */ - DPNP_FN_RNG_VONMISES_EXT, /**< Used in numpy.random.vonmises() impl, - requires extra parameters */ - DPNP_FN_RNG_WALD, /**< Used in numpy.random.wald() impl */ - DPNP_FN_RNG_WALD_EXT, /**< Used in numpy.random.wald() impl, requires extra - parameters */ - DPNP_FN_RNG_WEIBULL, /**< Used in numpy.random.weibull() impl */ - DPNP_FN_RNG_WEIBULL_EXT, /**< Used in numpy.random.weibull() impl, requires - extra parameters */ - DPNP_FN_RNG_ZIPF, /**< Used in numpy.random.zipf() impl */ - DPNP_FN_RNG_ZIPF_EXT, /**< Used in numpy.random.zipf() impl, requires extra - parameters */ - DPNP_FN_ZEROS, /**< Used in numpy.zeros() impl */ - DPNP_FN_ZEROS_LIKE, /**< Used in numpy.zeros_like() impl */ - DPNP_FN_LAST, /**< The latest element of the enumeration */ + DPNP_FN_LAST, /**< The latest element of the enumeration */ }; /** diff --git a/dpnp/backend/include/dpnp_iface_random.hpp b/dpnp/backend/include/dpnp_iface_random.hpp deleted file mode 100644 index 182e8caaf2eb..000000000000 --- a/dpnp/backend/include/dpnp_iface_random.hpp +++ /dev/null @@ -1,942 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// - Neither the name of the copyright holder nor the names of its contributors -// may be used to endorse or promote products derived from this software -// without specific prior written permission. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -/* - * This header file is for interface Cython with C++. - * It should not contains any backend specific headers (like SYCL or math - * library) because all included headers will be exposed in Cython compilation - * procedure - * - * We would like to avoid backend specific things in higher level Cython - * modules. Any backend interface functions and types should be defined here. - * - * Also, this file should contains documentation on functions and types - * which are used in the interface - */ - -#pragma once -#ifndef BACKEND_IFACE_RANDOM_H // Cython compatibility -#define BACKEND_IFACE_RANDOM_H - -/** - * @defgroup BACKEND_RANDOM_API Backend C++ library interface RANDOM API - * @{ - * This section describes Backend API for RANDOM NUMBER GENERATION (RNG) part. - * @} - */ - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (beta - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] a Alpha, shape param. - * @param [in] b Beta, scalefactor. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_beta_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType a, - const _DataType b, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_beta_c(void *result, - const _DataType a, - const _DataType b, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (binomial - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] ntrial Number of independent trials. - * @param [in] p Success probability p of a single trial. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_binomial_c(DPCTLSyclQueueRef q_ref, - void *result, - const int ntrial, - const double p, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_binomial_c(void *result, - const int ntrial, - const double p, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (chi-square - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] df Degrees of freedom. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_chisquare_c(DPCTLSyclQueueRef q_ref, - void *result, - const int df, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_rng_chisquare_c(void *result, const int df, const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (exponential - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] beta Beta, scalefactor. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_exponential_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType beta, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_exponential_c(void *result, - const _DataType beta, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (F - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] df_num Degrees of freedom in numerator. - * @param [in] df_den Degrees of freedom in denominator. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_f_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType df_num, - const _DataType df_den, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_f_c(void *result, - const _DataType df_num, - const _DataType df_den, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (gamma - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] shape The shape of the gamma distribution. - * @param [in] scale The scale of the gamma distribution. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_gamma_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType shape, - const _DataType scale, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_gamma_c(void *result, - const _DataType shape, - const _DataType scale, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (gaussian - * continuous distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] mean Mean value. - * @param [in] stddev Standard deviation. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_gaussian_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType mean, - const _DataType stddev, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_gaussian_c(void *result, - const _DataType mean, - const _DataType stddev, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (Geometric - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] p The probability of success of an individual - * trial. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_geometric_c(DPCTLSyclQueueRef q_ref, - void *result, - const float p, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_rng_geometric_c(void *result, const float p, const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (Gumbel - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] loc The location of the mode of the - * distribution. - * @param [in] scale The scale parameter of the distribution. - * Default is 1. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_gumbel_c(DPCTLSyclQueueRef q_ref, - void *result, - const double loc, - const double scale, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_gumbel_c(void *result, - const double loc, - const double scale, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (hypergeometric - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] l Lot size of l. - * @param [in] s Size of sampling without replacement. - * @param [in] m Number of marked elements m. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_hypergeometric_c(DPCTLSyclQueueRef q_ref, - void *result, - const int l, - const int s, - const int m, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_hypergeometric_c(void *result, - const int l, - const int s, - const int m, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (laplace - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] loc The position of the distribution peak. - * @param [in] scale The exponential decay. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_laplace_c(DPCTLSyclQueueRef q_ref, - void *result, - const double loc, - const double scale, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_laplace_c(void *result, - const double loc, - const double scale, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (logistic - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] loc The position of the distribution peak. - * @param [in] scale The exponential decay. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_logistic_c(DPCTLSyclQueueRef q_ref, - void *result, - const double loc, - double const scale, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_logistic_c(void *result, - const double loc, - double const scale, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (lognormal - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] mean Mean value. - * @param [in] stddev Standard deviation. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_lognormal_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType mean, - const _DataType stddev, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_lognormal_c(void *result, - const _DataType mean, - const _DataType stddev, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (multinomial - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] ntrial Number of independent trials. - * @param [in] p_in Probability of possible outcomes (k length). - * @param [in] p_size Length of `p_in`. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_multinomial_c(DPCTLSyclQueueRef q_ref, - void *result, - const int ntrial, - const double *p_in, - const size_t p_size, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_multinomial_c(void *result, - const int ntrial, - const double *p_in, - const size_t p_size, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (multivariate - * normal distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] dimen Dimension of output random vectors. - * @param [in] mean_in Mean array of dimension. - * @param [in] mean_size Length of `mean_in`. - * @param [in] cov Variance-covariance matrix. - * @param [in] cov_size Length of `cov_in`. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_multivariate_normal_c(DPCTLSyclQueueRef q_ref, - void *result, - const int dimen, - const double *mean_in, - const size_t mean_size, - const double *cov_in, - const size_t cov_size, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_multivariate_normal_c(void *result, - const int dimen, - const double *mean_in, - const size_t mean_size, - const double *cov_in, - const size_t cov_size, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (negative - * binomial distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] a The first distribution parameter a, > 0. - * @param [in] p The second distribution parameter p, >= 0 - * and <=1. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - * - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_negative_binomial_c(DPCTLSyclQueueRef q_ref, - void *result, - const double a, - const double p, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_negative_binomial_c(void *result, - const double a, - const double p, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (noncentral - * chisquare distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] df Degrees of freedom. - * @param [in] nonc Non-centrality param. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef dpnp_rng_noncentral_chisquare_c( - DPCTLSyclQueueRef q_ref, - void *result, - const _DataType df, - const _DataType nonc, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_noncentral_chisquare_c(void *result, - const _DataType df, - const _DataType nonc, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (normal - * continuous distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result_out Output array. - * @param [in] mean Mean value. - * @param [in] stddev Standard deviation. - * @param [in] size Number of elements in `result` arrays. - * @param [in] random_state_in Pointer on random state. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_normal_c(DPCTLSyclQueueRef q_ref, - void *result_out, - const double mean, - const double stddev, - const int64_t size, - void *random_state_in, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_normal_c(void *result, - const _DataType mean, - const _DataType stddev, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (Pareto - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] alpha Shape of the distribution, alpha. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_pareto_c(DPCTLSyclQueueRef q_ref, - void *result, - const double alpha, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_rng_pareto_c(void *result, const double alpha, const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (poisson - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] lambda Distribution parameter lambda. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_poisson_c(DPCTLSyclQueueRef q_ref, - void *result, - const double lambda, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_rng_poisson_c(void *result, const double lambda, const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (power - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] alpha Shape of the distribution, alpha. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_power_c(DPCTLSyclQueueRef q_ref, - void *result, - const double alpha, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_rng_power_c(void *result, const double alpha, const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (rayleigh - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] scale Distribution parameter, scalefactor. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_rayleigh_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType scale, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_rng_rayleigh_c(void *result, const _DataType scale, const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random in-place shuffle. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Input/output array. - * @param [in] itemsize Length of `result` array element in bytes. - * @param [in] ndim Number of array dimensions in `result` - * arrays. - * @param [in] high_dim_size Number of elements in `result` arrays higher - * dimension, or len(result). - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_shuffle_c(DPCTLSyclQueueRef q_ref, - void *result, - const size_t itemsize, - const size_t ndim, - const size_t high_dim_size, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_shuffle_c(void *result, - const size_t itemsize, - const size_t ndim, - const size_t high_dim_size, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief initializer for basic random number generator. - * - * @param [in] seed The seed value. - */ -INP_DLLEXPORT void dpnp_rng_srand_c(size_t seed = 1); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (standard - * cauchy distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_standard_cauchy_c(DPCTLSyclQueueRef q_ref, - void *result, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_standard_cauchy_c(void *result, const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (standard - * exponential distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef dpnp_rng_standard_exponential_c( - DPCTLSyclQueueRef q_ref, - void *result, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_standard_exponential_c(void *result, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (standard gamma - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] shape Shape value. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_standard_gamma_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType shape, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_standard_gamma_c(void *result, - const _DataType shape, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (standard - * normal distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result_out Output array. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT void dpnp_rng_standard_normal_c(void *result, const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (standard - * Student's t distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] df Degrees of freedom. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_standard_t_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType df, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_rng_standard_t_c(void *result, const _DataType df, const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (Triangular - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] x_min Lower limit. - * @param [in] x_mode The value where the peak of the distribution - * occurs. - * @param [in] x_max Upper limit. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_triangular_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType x_min, - const _DataType x_mode, - const _DataType x_max, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_triangular_c(void *result, - const _DataType x_min, - const _DataType x_mode, - const _DataType x_max, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (uniform - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result_out Output array. - * @param [in] low Left bound of array values. - * @param [in] high Right bound of array values. - * @param [in] size Number of elements in `result` array. - * @param [in] random_state_in Pointer on random state. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_uniform_c(DPCTLSyclQueueRef q_ref, - void *result_out, - const double low, - const double high, - const int64_t size, - void *random_state_in, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_uniform_c(void *result, - const long low, - const long high, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (Vonmises - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] mu Mode of the distribution. - * @param [in] kappa Dispersion of the distribution. - * @param [in] size Number of elements in `result` array. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_vonmises_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType mu, - const _DataType kappa, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_vonmises_c(void *result, - const _DataType mu, - const _DataType kappa, - const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (Wald's - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] mean Mean value. - * @param [in] scale The scale of the distribution. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_wald_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType mean, - const _DataType scale, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_rng_wald_c(void *result, - const _DataType mean, - const _DataType scale, - size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (weibull - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] alpha Shape parameter of the distribution, alpha. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_weibull_c(DPCTLSyclQueueRef q_ref, - void *result, - const double alpha, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_rng_weibull_c(void *result, const double alpha, const size_t size); - -/** - * @ingroup BACKEND_RANDOM_API - * @brief math library implementation of random number generator (Zipf - * distribution) - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] a Distribution parameter. - * @param [in] size Number of elements in `result` arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_rng_zipf_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType a, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_rng_zipf_c(void *result, const _DataType a, const size_t size); - -#endif // BACKEND_IFACE_RANDOM_H diff --git a/dpnp/backend/kernels/dpnp_krnl_arraycreation.cpp b/dpnp/backend/kernels/dpnp_krnl_arraycreation.cpp deleted file mode 100644 index 308c311b2ba7..000000000000 --- a/dpnp/backend/kernels/dpnp_krnl_arraycreation.cpp +++ /dev/null @@ -1,217 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// - Neither the name of the copyright holder nor the names of its contributors -// may be used to endorse or promote products derived from this software -// without specific prior written permission. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -#include - -#include "dpnp_fptr.hpp" -#include "dpnp_iface.hpp" -#include "dpnp_utils.hpp" -#include "dpnpc_memory_adapter.hpp" -#include "queue_sycl.hpp" - -template -class dpnp_ones_c_kernel; - -template -DPCTLSyclEventRef dpnp_ones_c(DPCTLSyclQueueRef q_ref, - void *result, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - sycl::queue q = *(reinterpret_cast(q_ref)); - - _DataType *fill_value = reinterpret_cast<_DataType *>( - sycl::malloc_shared(sizeof(_DataType), q)); - fill_value[0] = 1; - - DPCTLSyclEventRef event_ref = dpnp_initval_c<_DataType>( - q_ref, result, fill_value, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); - - sycl::free(fill_value, q); - - return nullptr; -} - -template -void dpnp_ones_c(void *result, size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_ones_c<_DataType>(q_ref, result, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_ones_default_c)(void *, size_t) = dpnp_ones_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_ones_like_c(DPCTLSyclQueueRef q_ref, - void *result, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - return dpnp_ones_c<_DataType>(q_ref, result, size, dep_event_vec_ref); -} - -template -void dpnp_ones_like_c(void *result, size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_ones_like_c<_DataType>(q_ref, result, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_ones_like_default_c)(void *, size_t) = dpnp_ones_like_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_zeros_c(DPCTLSyclQueueRef q_ref, - void *result, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - sycl::queue q = *(reinterpret_cast(q_ref)); - - _DataType *fill_value = reinterpret_cast<_DataType *>( - sycl::malloc_shared(sizeof(_DataType), q)); - fill_value[0] = 0; - - DPCTLSyclEventRef event_ref = dpnp_initval_c<_DataType>( - q_ref, result, fill_value, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); - - sycl::free(fill_value, q); - - return nullptr; -} - -template -void dpnp_zeros_c(void *result, size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_zeros_c<_DataType>(q_ref, result, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_zeros_default_c)(void *, size_t) = dpnp_zeros_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_zeros_like_c(DPCTLSyclQueueRef q_ref, - void *result, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - return dpnp_zeros_c<_DataType>(q_ref, result, size, dep_event_vec_ref); -} - -template -void dpnp_zeros_like_c(void *result, size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_zeros_like_c<_DataType>(q_ref, result, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_zeros_like_default_c)(void *, - size_t) = dpnp_zeros_like_c<_DataType>; - -void func_map_init_arraycreation(func_map_t &fmap) -{ - // Used in dpnp_rng_geometric_c - fmap[DPNPFuncName::DPNP_FN_ONES][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_ones_default_c}; - fmap[DPNPFuncName::DPNP_FN_ONES][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_ones_default_c}; - fmap[DPNPFuncName::DPNP_FN_ONES][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_ones_default_c}; - fmap[DPNPFuncName::DPNP_FN_ONES][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_ones_default_c}; - fmap[DPNPFuncName::DPNP_FN_ONES][eft_BLN][eft_BLN] = { - eft_BLN, (void *)dpnp_ones_default_c}; - fmap[DPNPFuncName::DPNP_FN_ONES][eft_C128][eft_C128] = { - eft_C128, (void *)dpnp_ones_default_c>}; - - fmap[DPNPFuncName::DPNP_FN_ONES_LIKE][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_ones_like_default_c}; - fmap[DPNPFuncName::DPNP_FN_ONES_LIKE][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_ones_like_default_c}; - fmap[DPNPFuncName::DPNP_FN_ONES_LIKE][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_ones_like_default_c}; - fmap[DPNPFuncName::DPNP_FN_ONES_LIKE][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_ones_like_default_c}; - fmap[DPNPFuncName::DPNP_FN_ONES_LIKE][eft_BLN][eft_BLN] = { - eft_BLN, (void *)dpnp_ones_like_default_c}; - fmap[DPNPFuncName::DPNP_FN_ONES_LIKE][eft_C128][eft_C128] = { - eft_C128, (void *)dpnp_ones_like_default_c>}; - - // Used in dpnp_rng_binomial_c, dpnp_rng_gamma_c, dpnp_rng_hypergeometric_c - // dpnp_rng_laplace_c, dpnp_rng_multinomial_c, dpnp_rng_weibull_c - fmap[DPNPFuncName::DPNP_FN_ZEROS][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_zeros_default_c}; - fmap[DPNPFuncName::DPNP_FN_ZEROS][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_zeros_default_c}; - fmap[DPNPFuncName::DPNP_FN_ZEROS][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_zeros_default_c}; - fmap[DPNPFuncName::DPNP_FN_ZEROS][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_zeros_default_c}; - fmap[DPNPFuncName::DPNP_FN_ZEROS][eft_BLN][eft_BLN] = { - eft_BLN, (void *)dpnp_zeros_default_c}; - fmap[DPNPFuncName::DPNP_FN_ZEROS][eft_C128][eft_C128] = { - eft_C128, (void *)dpnp_zeros_default_c>}; - - fmap[DPNPFuncName::DPNP_FN_ZEROS_LIKE][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_zeros_like_default_c}; - fmap[DPNPFuncName::DPNP_FN_ZEROS_LIKE][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_zeros_like_default_c}; - fmap[DPNPFuncName::DPNP_FN_ZEROS_LIKE][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_zeros_like_default_c}; - fmap[DPNPFuncName::DPNP_FN_ZEROS_LIKE][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_zeros_like_default_c}; - fmap[DPNPFuncName::DPNP_FN_ZEROS_LIKE][eft_BLN][eft_BLN] = { - eft_BLN, (void *)dpnp_zeros_like_default_c}; - fmap[DPNPFuncName::DPNP_FN_ZEROS_LIKE][eft_C128][eft_C128] = { - eft_C128, (void *)dpnp_zeros_like_default_c>}; - - return; -} diff --git a/dpnp/backend/kernels/dpnp_krnl_common.cpp b/dpnp/backend/kernels/dpnp_krnl_common.cpp deleted file mode 100644 index ae5f4fb76b6c..000000000000 --- a/dpnp/backend/kernels/dpnp_krnl_common.cpp +++ /dev/null @@ -1,126 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// - Neither the name of the copyright holder nor the names of its contributors -// may be used to endorse or promote products derived from this software -// without specific prior written permission. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -#include -#include -#include -#include -#include - -#include "dpnp_fptr.hpp" -#include "dpnp_utils.hpp" -#include "dpnpc_memory_adapter.hpp" -#include "queue_sycl.hpp" -#include - -template -class dpnp_initval_c_kernel; - -template -DPCTLSyclEventRef dpnp_initval_c(DPCTLSyclQueueRef q_ref, - void *result, - void *value, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - _DataType val = *(static_cast<_DataType *>(value)); - - validate_type_for_device<_DataType>(q); - - auto event = q.fill<_DataType>(result, val, size); - event_ref = reinterpret_cast(&event); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_initval_c(void *result1, void *value, size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_initval_c<_DataType>( - q_ref, result1, value, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_initval_default_c)(void *, - void *, - size_t) = dpnp_initval_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_initval_ext_c)(DPCTLSyclQueueRef, - void *, - void *, - size_t, - const DPCTLEventVectorRef) = - dpnp_initval_c<_DataType>; - -void func_map_init_linalg(func_map_t &fmap) -{ - fmap[DPNPFuncName::DPNP_FN_INITVAL][eft_BLN][eft_BLN] = { - eft_BLN, (void *)dpnp_initval_default_c}; - fmap[DPNPFuncName::DPNP_FN_INITVAL][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_initval_default_c}; - fmap[DPNPFuncName::DPNP_FN_INITVAL][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_initval_default_c}; - fmap[DPNPFuncName::DPNP_FN_INITVAL][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_initval_default_c}; - fmap[DPNPFuncName::DPNP_FN_INITVAL][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_initval_default_c}; - fmap[DPNPFuncName::DPNP_FN_INITVAL][eft_C128][eft_C128] = { - eft_C128, (void *)dpnp_initval_default_c>}; - - fmap[DPNPFuncName::DPNP_FN_INITVAL_EXT][eft_BLN][eft_BLN] = { - eft_BLN, (void *)dpnp_initval_ext_c}; - fmap[DPNPFuncName::DPNP_FN_INITVAL_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_initval_ext_c}; - fmap[DPNPFuncName::DPNP_FN_INITVAL_EXT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_initval_ext_c}; - fmap[DPNPFuncName::DPNP_FN_INITVAL_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_initval_ext_c}; - fmap[DPNPFuncName::DPNP_FN_INITVAL_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_initval_ext_c}; - fmap[DPNPFuncName::DPNP_FN_INITVAL_EXT][eft_C64][eft_C64] = { - eft_C64, (void *)dpnp_initval_ext_c>}; - fmap[DPNPFuncName::DPNP_FN_INITVAL_EXT][eft_C128][eft_C128] = { - eft_C128, (void *)dpnp_initval_ext_c>}; - - return; -} diff --git a/dpnp/backend/kernels/dpnp_krnl_random.cpp b/dpnp/backend/kernels/dpnp_krnl_random.cpp deleted file mode 100644 index d1219d286407..000000000000 --- a/dpnp/backend/kernels/dpnp_krnl_random.cpp +++ /dev/null @@ -1,3241 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// - Neither the name of the copyright holder nor the names of its contributors -// may be used to endorse or promote products derived from this software -// without specific prior written permission. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -#include -#include -#include -#include -#include -#include -#include - -#include - -#include "dpnp_fptr.hpp" -#include "dpnp_random_state.hpp" -#include "dpnp_utils.hpp" -#include "dpnpc_memory_adapter.hpp" -#include "queue_sycl.hpp" - -namespace mkl_blas = oneapi::mkl::blas; -namespace mkl_rng = oneapi::mkl::rng; -namespace mkl_vm = oneapi::mkl::vm; - -static std::mutex rng_stream_mutex; -/** - * Use get/set functions to access/modify this variable. Both require - * `rng_stream_mutex` to be held, and it must stay held for as long as the - * stream returned by the getter is in use, since the setter destroys the - * stream. - */ -static VSLStreamStatePtr rng_stream = nullptr; - -// only called with rng_stream_mutex held -static void set_rng_stream_nolock(size_t seed = 1) -{ - if (rng_stream) { - vslDeleteStream(&rng_stream); - rng_stream = nullptr; - } - - vslNewStream(&rng_stream, VSL_BRNG_MT19937, seed); -} - -// only called with rng_stream_mutex held -static VSLStreamStatePtr get_rng_stream_nolock() -{ - if (!rng_stream) { - set_rng_stream_nolock(); - } - - return rng_stream; -} - -void dpnp_rng_srand_c(size_t seed) -{ - auto &be = backend_sycl::get(); - be.set_rng_engines_seed(seed); - - std::lock_guard lock(rng_stream_mutex); - set_rng_stream_nolock(seed); -} - -template -static inline DPCTLSyclEventRef dpnp_rng_generate(const _DistrType &distr, - _EngineType &engine, - const int64_t size, - _DataType *result) -{ - DPCTLSyclEventRef event_ref = nullptr; - sycl::event event; - - // perform rng generation - try { - event = mkl_rng::generate<_DistrType, _EngineType>(distr, engine, size, - result); - event_ref = reinterpret_cast(&event); - } catch (const std::exception &e) { - // TODO: add error reporting - return event_ref; - } - return DPCTLEvent_Copy(event_ref); -} - -template -static inline DPCTLSyclEventRef dpnp_rng_generate_uniform(_EngineType &engine, - sycl::queue *q, - const _DataType a, - const _DataType b, - const int64_t size, - _DataType *result) -{ - DPCTLSyclEventRef event_ref = nullptr; - - if constexpr (std::is_same<_DataType, int32_t>::value) { - if (q->get_device().has(sycl::aspect::fp64)) { - /** - * A note from oneMKL for oneapi::mkl::rng::uniform (Discrete): - * The oneapi::mkl::rng::uniform_method::standard uses the s BRNG - * type on GPU devices. This might cause the produced numbers to - * have incorrect statistics (due to rounding error) - * when abs(b-a) > 2^23 || abs(b) > 2^23 || abs(a) > 2^23. - * To get proper statistics for this case, use the - * oneapi::mkl::rng::uniform_method::accurate method instead. - */ - using method_type = mkl_rng::uniform_method::accurate; - mkl_rng::uniform<_DataType, method_type> distribution(a, b); - - // perform generation - try { - sycl::event event = - mkl_rng::generate(distribution, engine, size, result); - - event_ref = reinterpret_cast(&event); - return DPCTLEvent_Copy(event_ref); - } catch (const oneapi::mkl::unsupported_device &) { - // fall through to try with uniform_method::standard - } catch (const oneapi::mkl::unimplemented &) { - // fall through to try with uniform_method::standard - } catch (const std::exception &e) { - // TODO: add error reporting - return event_ref; - } - } - } - - // uniform_method::standard is a method used by default - using method_type = mkl_rng::uniform_method::standard; - mkl_rng::uniform<_DataType, method_type> distribution(a, b); - - // perform generation - return dpnp_rng_generate(distribution, engine, size, result); -} - -template -DPCTLSyclEventRef dpnp_rng_beta_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType a, - const _DataType b, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - _DataType displacement = _DataType(0.0); - - _DataType scalefactor = _DataType(1.0); - - _DataType *result1 = reinterpret_cast<_DataType *>(result); - - mkl_rng::beta<_DataType> distribution(a, b, displacement, scalefactor); - // perform generation - auto event_out = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - - event_ref = reinterpret_cast(&event_out); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_beta_c(void *result, - const _DataType a, - const _DataType b, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_beta_c<_DataType>( - q_ref, result, a, b, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_beta_default_c)(void *, - const _DataType, - const _DataType, - const size_t) = dpnp_rng_beta_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_beta_ext_c)(DPCTLSyclQueueRef, - void *, - const _DataType, - const _DataType, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_beta_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_binomial_c(DPCTLSyclQueueRef q_ref, - void *result, - const int ntrial, - const double p, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (result == nullptr) { - return event_ref; - } - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - if (ntrial == 0 || p == 0) { - event_ref = - dpnp_zeros_c<_DataType>(q_ref, result, size, dep_event_vec_ref); - } - else if (p == 1) { - _DataType *fill_value = reinterpret_cast<_DataType *>( - sycl::malloc_shared(sizeof(_DataType), q)); - fill_value[0] = static_cast<_DataType>(ntrial); - - event_ref = dpnp_initval_c<_DataType>(q_ref, result, fill_value, size, - dep_event_vec_ref); - DPCTLEvent_Wait(event_ref); - dpnp_memory_free_c(q_ref, fill_value); - } - else { - _DataType *result1 = reinterpret_cast<_DataType *>(result); - mkl_rng::binomial<_DataType> distribution(ntrial, p); - auto event_out = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_ref = reinterpret_cast(&event_out); - return DPCTLEvent_Copy(event_ref); - } - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_binomial_c(void *result, - const int ntrial, - const double p, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_binomial_c<_DataType>( - q_ref, result, ntrial, p, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_binomial_default_c)(void *, - const int, - const double, - const size_t) = - dpnp_rng_binomial_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_binomial_ext_c)(DPCTLSyclQueueRef, - void *, - const int, - const double, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_binomial_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_chisquare_c(DPCTLSyclQueueRef q_ref, - void *result, - const int df, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - _DataType *result1 = reinterpret_cast<_DataType *>(result); - - mkl_rng::chi_square<_DataType> distribution(df); - auto event_out = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - - event_ref = reinterpret_cast(&event_out); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_chisquare_c(void *result, const int df, const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_chisquare_c<_DataType>( - q_ref, result, df, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_chisquare_default_c)(void *, const int, const size_t) = - dpnp_rng_chisquare_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_chisquare_ext_c)(DPCTLSyclQueueRef, - void *, - const int, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_chisquare_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_exponential_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType beta, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - // set displacement a - const _DataType a = (_DataType(0.0)); - - _DataType *result1 = reinterpret_cast<_DataType *>(result); - - mkl_rng::exponential<_DataType> distribution(a, beta); - // perform generation - auto event_out = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - - event_ref = reinterpret_cast(&event_out); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_exponential_c(void *result, - const _DataType beta, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_exponential_c<_DataType>( - q_ref, result, beta, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_exponential_default_c)(void *, const _DataType, const size_t) = - dpnp_rng_exponential_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_exponential_ext_c)(DPCTLSyclQueueRef, - void *, - const _DataType, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_exponential_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_rng_f_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType df_num, - const _DataType df_den, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - std::vector no_deps; - - const _DataType d_zero = (_DataType(0.0)); - - _DataType shape = 0.5 * df_num; - _DataType scale = 2.0 / df_num; - _DataType *den = nullptr; - - DPNPC_ptr_adapter<_DataType> result1_ptr(q_ref, result, size, true, true); - _DataType *result1 = result1_ptr.get_ptr(); - - mkl_rng::gamma<_DataType> gamma_distribution1(shape, d_zero, scale); - auto event_gamma_distribution1 = - mkl_rng::generate(gamma_distribution1, DPNP_RNG_ENGINE, size, result1); - - den = reinterpret_cast<_DataType *>( - sycl::malloc_shared(size * sizeof(_DataType), q)); - shape = 0.5 * df_den; - scale = 2.0 / df_den; - mkl_rng::gamma<_DataType> gamma_distribution2(shape, d_zero, scale); - auto event_gamma_distribution2 = - mkl_rng::generate(gamma_distribution2, DPNP_RNG_ENGINE, size, den); - - auto event_out = - mkl_vm::div(q, size, result1, den, result1, - {event_gamma_distribution1, event_gamma_distribution2}, - mkl_vm::mode::ha); - event_out.wait(); - - sycl::free(den, q); - - return event_ref; -} - -template -void dpnp_rng_f_c(void *result, - const _DataType df_num, - const _DataType df_den, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_f_c<_DataType>( - q_ref, result, df_num, df_den, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_f_default_c)(void *, - const _DataType, - const _DataType, - const size_t) = dpnp_rng_f_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_f_ext_c)(DPCTLSyclQueueRef, - void *, - const _DataType, - const _DataType, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_f_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_rng_gamma_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType shape, - const _DataType scale, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - sycl::event event_out; - - if (!size || result == nullptr) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - if (shape == 0.0 || scale == 0.0) { - event_ref = - dpnp_zeros_c<_DataType>(q_ref, result, size, dep_event_vec_ref); - } - else { - _DataType *result1 = reinterpret_cast<_DataType *>(result); - const _DataType a = (_DataType(0.0)); - - mkl_rng::gamma<_DataType> distribution(shape, a, scale); - event_out = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_ref = reinterpret_cast(&event_out); - } - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_gamma_c(void *result, - const _DataType shape, - const _DataType scale, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_gamma_c<_DataType>( - q_ref, result, shape, scale, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_gamma_default_c)(void *, - const _DataType, - const _DataType, - const size_t) = dpnp_rng_gamma_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_gamma_ext_c)(DPCTLSyclQueueRef, - void *, - const _DataType, - const _DataType, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_gamma_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_gaussian_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType mean, - const _DataType stddev, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - _DataType *result1 = reinterpret_cast<_DataType *>(result); - - mkl_rng::gaussian<_DataType> distribution(mean, stddev); - // perform generation - auto event_out = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_ref = reinterpret_cast(&event_out); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_gaussian_c(void *result, - const _DataType mean, - const _DataType stddev, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_gaussian_c<_DataType>( - q_ref, result, mean, stddev, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_gaussian_default_c)(void *, - const _DataType, - const _DataType, - const size_t) = - dpnp_rng_gaussian_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_gaussian_ext_c)(DPCTLSyclQueueRef, - void *, - const _DataType, - const _DataType, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_gaussian_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_geometric_c(DPCTLSyclQueueRef q_ref, - void *result, - const float p, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - sycl::event event_out; - - if (!size || !result) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - if (p == 1.0) { - event_ref = - dpnp_ones_c<_DataType>(q_ref, result, size, dep_event_vec_ref); - } - else { - _DataType *result1 = reinterpret_cast<_DataType *>(result); - mkl_rng::geometric<_DataType> distribution(p); - // perform generation - event_out = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_ref = reinterpret_cast(&event_out); - } - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_geometric_c(void *result, const float p, const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_geometric_c<_DataType>( - q_ref, result, p, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_geometric_default_c)(void *, const float, const size_t) = - dpnp_rng_geometric_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_geometric_ext_c)(DPCTLSyclQueueRef, - void *, - const float, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_geometric_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_rng_gumbel_c(DPCTLSyclQueueRef q_ref, - void *result, - const double loc, - const double scale, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - sycl::event event_out; - - if (!size || !result) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - if (scale == 0.0) { - _DataType *fill_value = reinterpret_cast<_DataType *>( - sycl::malloc_shared(sizeof(_DataType), q)); - fill_value[0] = static_cast<_DataType>(loc); - - event_ref = dpnp_initval_c<_DataType>(q_ref, result, fill_value, size, - dep_event_vec_ref); - DPCTLEvent_Wait(event_ref); - dpnp_memory_free_c(q_ref, fill_value); - } - else { - const _DataType alpha = (_DataType(-1.0)); - std::int64_t incx = 1; - _DataType *result1 = reinterpret_cast<_DataType *>(result); - double negloc = loc * (double(-1.0)); - - mkl_rng::gumbel<_DataType> distribution(negloc, scale); - auto event_distribution = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - - event_out = - mkl_blas::scal(q, size, alpha, result1, incx, {event_distribution}); - event_ref = reinterpret_cast(&event_out); - } - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_gumbel_c(void *result, - const double loc, - const double scale, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_gumbel_c<_DataType>( - q_ref, result, loc, scale, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_gumbel_default_c)(void *, - const double, - const double, - const size_t) = dpnp_rng_gumbel_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_gumbel_ext_c)(DPCTLSyclQueueRef, - void *, - const double, - const double, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_gumbel_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_hypergeometric_c(DPCTLSyclQueueRef q_ref, - void *result, - const int l, - const int s, - const int m, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - sycl::event event_out; - - if (!size || !result) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - if (m == 0) { - event_ref = - dpnp_zeros_c<_DataType>(q_ref, result, size, dep_event_vec_ref); - } - else if (l == m) { - _DataType *fill_value = reinterpret_cast<_DataType *>( - sycl::malloc_shared(sizeof(_DataType), q)); - fill_value[0] = static_cast<_DataType>(s); - - event_ref = dpnp_initval_c<_DataType>(q_ref, result, fill_value, size, - dep_event_vec_ref); - DPCTLEvent_Wait(event_ref); - dpnp_memory_free_c(q_ref, fill_value); - } - else { - _DataType *result1 = reinterpret_cast<_DataType *>(result); - mkl_rng::hypergeometric<_DataType> distribution(l, s, m); - event_out = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_ref = reinterpret_cast(&event_out); - } - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_hypergeometric_c(void *result, - const int l, - const int s, - const int m, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_hypergeometric_c<_DataType>( - q_ref, result, l, s, m, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_hypergeometric_default_c)(void *, - const int, - const int, - const int, - const size_t) = - dpnp_rng_hypergeometric_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_hypergeometric_ext_c)(DPCTLSyclQueueRef, - void *, - const int, - const int, - const int, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_hypergeometric_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_laplace_c(DPCTLSyclQueueRef q_ref, - void *result, - const double loc, - const double scale, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - sycl::event event_out; - - if (!size || !result) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - if (scale == 0.0) { - event_ref = - dpnp_zeros_c<_DataType>(q_ref, result, size, dep_event_vec_ref); - } - else { - _DataType *result1 = reinterpret_cast<_DataType *>(result); - mkl_rng::laplace<_DataType> distribution(loc, scale); - // perform generation - event_out = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_ref = reinterpret_cast(&event_out); - } - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_laplace_c(void *result, - const double loc, - const double scale, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_laplace_c<_DataType>( - q_ref, result, loc, scale, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_laplace_default_c)(void *, - const double, - const double, - const size_t) = - dpnp_rng_laplace_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_laplace_ext_c)(DPCTLSyclQueueRef, - void *, - const double, - const double, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_laplace_c<_DataType>; - -template -class dpnp_rng_logistic_c_kernel; - -/* Logistic(loc, scale) ~ loc + scale * log(u/(1.0 - u)) */ -template -DPCTLSyclEventRef - dpnp_rng_logistic_c(DPCTLSyclQueueRef q_ref, - void *result, - const double loc, - const double scale, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size || !result) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - const _DataType d_zero = _DataType(0.0); - const _DataType d_one = _DataType(1.0); - - _DataType *result1 = reinterpret_cast<_DataType *>(result); - - mkl_rng::uniform<_DataType> distribution(d_zero, d_one); - auto event_distribution = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - - sycl::range<1> gws(size); - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { - size_t i = global_id[0]; - result1[i] = sycl::log(result1[i] / (1.0 - result1[i])); - result1[i] = loc + scale * result1[i]; - }; - auto kernel_func = [&](sycl::handler &cgh) { - cgh.depends_on({event_distribution}); - cgh.parallel_for>( - gws, kernel_parallel_for_func); - }; - auto event = q.submit(kernel_func); - event_ref = reinterpret_cast(&event); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_logistic_c(void *result, - const double loc, - const double scale, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_logistic_c<_DataType>( - q_ref, result, loc, scale, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_logistic_default_c)(void *, - const double, - const double, - const size_t) = - dpnp_rng_logistic_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_logistic_ext_c)(DPCTLSyclQueueRef, - void *, - const double, - const double, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_logistic_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_lognormal_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType mean, - const _DataType stddev, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - sycl::event event_out; - - if (!size || !result) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - _DataType *result1 = reinterpret_cast<_DataType *>(result); - - if (stddev == 0.0) { - _DataType *fill_value = reinterpret_cast<_DataType *>( - sycl::malloc_shared(sizeof(_DataType), q)); - fill_value[0] = - static_cast<_DataType>(std::exp(mean + (stddev * stddev) / 2)); - - event_ref = dpnp_initval_c<_DataType>(q_ref, result, fill_value, size, - dep_event_vec_ref); - DPCTLEvent_Wait(event_ref); - dpnp_memory_free_c(q_ref, fill_value); - } - else { - const _DataType displacement = _DataType(0.0); - const _DataType scalefactor = _DataType(1.0); - - mkl_rng::lognormal<_DataType> distribution(mean, stddev, displacement, - scalefactor); - event_out = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_ref = reinterpret_cast(&event_out); - } - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_lognormal_c(void *result, - const _DataType mean, - const _DataType stddev, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_lognormal_c<_DataType>( - q_ref, result, mean, stddev, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_lognormal_default_c)(void *, - const _DataType, - const _DataType, - const size_t) = - dpnp_rng_lognormal_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_lognormal_ext_c)(DPCTLSyclQueueRef, - void *, - const _DataType, - const _DataType, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_lognormal_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_multinomial_c(DPCTLSyclQueueRef q_ref, - void *result, - const int ntrial, - void *p_in, - const size_t p_size, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - sycl::event event_out; - - if (!size || !result || (ntrial < 0)) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - if (ntrial == 0) { - event_ref = - dpnp_zeros_c<_DataType>(q_ref, result, size, dep_event_vec_ref); - } - else { - DPNPC_ptr_adapter p_ptr(q_ref, p_in, p_size, true); - double *p_data = p_ptr.get_ptr(); - - // size = size - // `result` is a array for random numbers - // `size` is a `result`'s len. `size = n * p_size` - // `n` is a number of random values to be generated. - size_t n = size / p_size; - - size_t is_cpu_queue = dpnp_queue_is_cpu_c(); - - // math library supports the distribution generation on GPU device with - // input parameters which follow the condition - if (is_cpu_queue || - (!is_cpu_queue && (p_size >= ((size_t)ntrial * 16)) && - (ntrial <= 16))) { - DPNPC_ptr_adapter<_DataType> result_ptr(q_ref, result, size, true, - true); - _DataType *result1 = result_ptr.get_ptr(); - - auto p = sycl::span{p_data, p_size}; - mkl_rng::multinomial<_DataType> distribution(ntrial, p); - - // perform generation - event_out = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, n, result1); - event_ref = reinterpret_cast(&event_out); - - p_ptr.depends_on(event_out); - result_ptr.depends_on(event_out); - } - else { - DPNPC_ptr_adapter<_DataType> result_ptr(q_ref, result, size, true, - true); - _DataType *result1 = result_ptr.get_ptr(); - - int errcode; - { - // the lock is held across the call, as the stream is mutated - // by the generation and may be destroyed by a concurrent - // reseeding - std::lock_guard lock(rng_stream_mutex); - errcode = - viRngMultinomial(VSL_RNG_METHOD_MULTINOMIAL_MULTPOISSON, - get_rng_stream_nolock(), n, result1, - ntrial, p_size, p_data); - } - if (errcode != VSL_STATUS_OK) { - throw std::runtime_error( - "DPNP RNG Error: dpnp_rng_multinomial_c() failed."); - } - } - } - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_multinomial_c(void *result, - const int ntrial, - void *p_in, - const size_t p_size, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_multinomial_c<_DataType>( - q_ref, result, ntrial, p_in, p_size, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_rng_multinomial_default_c)(void *, - const int, - void *, - const size_t, - const size_t) = - dpnp_rng_multinomial_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_multinomial_ext_c)(DPCTLSyclQueueRef, - void *, - const int, - void *, - const size_t, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_multinomial_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_multivariate_normal_c(DPCTLSyclQueueRef q_ref, - void *result, - const int dimen, - void *mean_in, - const size_t mean_size, - void *cov_in, - const size_t cov_size, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter mean_ptr(q_ref, mean_in, mean_size, true); - double *mean_data = mean_ptr.get_ptr(); - DPNPC_ptr_adapter cov_ptr(q_ref, cov_in, cov_size, true); - double *cov_data = cov_ptr.get_ptr(); - - _DataType *result1 = static_cast<_DataType *>(result); - - auto mean = sycl::span{mean_data, mean_size}; - auto cov = sycl::span{cov_data, cov_size}; - - // `result` is a array for random numbers - // `size` is a `result`'s len. - // `size1` is a number of random values to be generated for each dimension. - size_t size1 = size / dimen; - - mkl_rng::gaussian_mv<_DataType> distribution(dimen, mean, cov); - auto event_out = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size1, result1); - event_ref = reinterpret_cast(&event_out); - - mean_ptr.depends_on(event_out); - cov_ptr.depends_on(event_out); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_multivariate_normal_c(void *result, - const int dimen, - void *mean_in, - const size_t mean_size, - void *cov_in, - const size_t cov_size, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_multivariate_normal_c<_DataType>( - q_ref, result, dimen, mean_in, mean_size, cov_in, cov_size, size, - dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_rng_multivariate_normal_default_c)(void *, - const int, - void *, - const size_t, - void *, - const size_t, - const size_t) = - dpnp_rng_multivariate_normal_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_multivariate_normal_ext_c)( - DPCTLSyclQueueRef, - void *, - const int, - void *, - const size_t, - void *, - const size_t, - const size_t, - const DPCTLEventVectorRef) = dpnp_rng_multivariate_normal_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_negative_binomial_c(DPCTLSyclQueueRef q_ref, - void *result, - const double a, - const double p, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - _DataType *result1 = reinterpret_cast<_DataType *>(result); - mkl_rng::negative_binomial<_DataType> distribution(a, p); - auto event_out = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_ref = reinterpret_cast(&event_out); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_negative_binomial_c(void *result, - const double a, - const double p, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_negative_binomial_c<_DataType>( - q_ref, result, a, p, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_negative_binomial_default_c)(void *, - const double, - const double, - const size_t) = - dpnp_rng_negative_binomial_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_negative_binomial_ext_c)( - DPCTLSyclQueueRef, - void *, - const double, - const double, - const size_t, - const DPCTLEventVectorRef) = dpnp_rng_negative_binomial_c<_DataType>; - -template -class dpnp_rng_noncentral_chisquare_c_kernel1; -template -class dpnp_rng_noncentral_chisquare_c_kernel2; - -template -DPCTLSyclEventRef - dpnp_rng_noncentral_chisquare_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType df, - const _DataType nonc, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - sycl::event event_out; - - if (!size || !result) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter<_DataType> result1_ptr(q_ref, result, size, false, true); - _DataType *result1 = result1_ptr.get_ptr(); - - const _DataType d_zero = _DataType(0.0); - const _DataType d_one = _DataType(1.0); - const _DataType d_two = _DataType(2.0); - - _DataType shape, loc; - size_t i; - - if (df > 1) { - _DataType *nvec = nullptr; - - shape = 0.5 * (df - 1.0); - /* res has chi^2 with (df - 1) */ - mkl_rng::gamma<_DataType> gamma_distribution(shape, d_zero, d_two); - auto event_gamma_distr = mkl_rng::generate( - gamma_distribution, DPNP_RNG_ENGINE, size, result1); - - nvec = reinterpret_cast<_DataType *>( - sycl::malloc_shared(size * sizeof(_DataType), q)); - - loc = sqrt(nonc); - - mkl_rng::gaussian<_DataType> gaussian_distribution(loc, d_one); - auto event_gaussian_distr = mkl_rng::generate( - gaussian_distribution, DPNP_RNG_ENGINE, size, nvec); - - /* squaring could result in an overflow */ - auto event_sqr_out = mkl_vm::sqr( - q, size, nvec, nvec, {event_gamma_distr, event_gaussian_distr}, - mkl_vm::mode::ha); - auto event_add_out = mkl_vm::add(q, size, result1, nvec, result1, - {event_sqr_out}, mkl_vm::mode::ha); - event_add_out.wait(); - sycl::free(nvec, q); - } - else if (df < 1) { - /* noncentral_chisquare(df, nonc) ~ G( df/2 + Poisson(nonc/2), 2) */ - double lambda; - int *pvec = nullptr; - pvec = - reinterpret_cast(sycl::malloc_shared(size * sizeof(int), q)); - lambda = 0.5 * nonc; - - mkl_rng::poisson poisson_distribution(lambda); - event_out = mkl_rng::generate(poisson_distribution, DPNP_RNG_ENGINE, - size, pvec); - event_out.wait(); - - shape = 0.5 * df; - if (0.125 * size > sqrt(lambda)) { - size_t *idx = nullptr; - _DataType *tmp = nullptr; - idx = reinterpret_cast( - sycl::malloc_shared(size * sizeof(size_t), q)); - - sycl::range<1> gws1(size); - auto kernel_parallel_for_func1 = [=](sycl::id<1> global_id) { - size_t i = global_id[0]; - idx[i] = i; - }; - auto kernel_func1 = [&](sycl::handler &cgh) { - cgh.parallel_for< - class dpnp_rng_noncentral_chisquare_c_kernel1<_DataType>>( - gws1, kernel_parallel_for_func1); - }; - event_out = q.submit(kernel_func1); - event_out.wait(); - - std::sort(idx, idx + size, [pvec](size_t i1, size_t i2) { - return pvec[i1] < pvec[i2]; - }); - /* idx now contains original indexes of ordered Poisson outputs */ - - /* allocate workspace to store samples of gamma, enough to hold - * entire output */ - tmp = reinterpret_cast<_DataType *>( - sycl::malloc_shared(size * sizeof(_DataType), q)); - for (i = 0; i < size;) { - size_t j; - int cv = pvec[idx[i]]; - // TODO vectorize - for (j = i + 1; (j < size) && (pvec[idx[j]] == cv); j++) {} - - if (j <= i) { - throw std::runtime_error( - "DPNP RNG Error: dpnp_rng_noncentral_chisquare_c() " - "failed."); - } - mkl_rng::gamma<_DataType> gamma_distribution(shape + cv, d_zero, - d_two); - event_out = mkl_rng::generate(gamma_distribution, - DPNP_RNG_ENGINE, j - i, tmp); - event_out.wait(); - - sycl::range<1> gws2(j - i); - auto kernel_parallel_for_func2 = [=](sycl::id<1> global_id) { - size_t index = global_id[0]; - result1[idx[index + i]] = tmp[index]; - }; - auto kernel_func2 = [&](sycl::handler &cgh) { - cgh.parallel_for< - class dpnp_rng_noncentral_chisquare_c_kernel2< - _DataType>>(gws2, kernel_parallel_for_func2); - }; - event_out = q.submit(kernel_func2); - event_out.wait(); - - i = j; - } - sycl::free(tmp, q); - sycl::free(idx, q); - } - else { - for (i = 0; i < size; i++) { - mkl_rng::gamma<_DataType> gamma_distribution(shape + pvec[i], - d_zero, d_two); - event_out = mkl_rng::generate(gamma_distribution, - DPNP_RNG_ENGINE, 1, result1 + 1); - event_out.wait(); - } - } - sycl::free(pvec, q); - } - else { - /* noncentral_chisquare(1, nonc) ~ (Z + sqrt(nonc))**2 for df == 1 */ - loc = sqrt(nonc); - mkl_rng::gaussian<_DataType> gaussian_distribution(loc, d_one); - auto event_gaussian_distr = mkl_rng::generate( - gaussian_distribution, DPNP_RNG_ENGINE, size, result1); - event_out = mkl_vm::sqr(q, size, result1, result1, - {event_gaussian_distr}, mkl_vm::mode::ha); - event_ref = reinterpret_cast(&event_out); - } - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_noncentral_chisquare_c(void *result, - const _DataType df, - const _DataType nonc, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_noncentral_chisquare_c<_DataType>( - q_ref, result, df, nonc, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_noncentral_chisquare_default_c)(void *, - const _DataType, - const _DataType, - const size_t) = - dpnp_rng_noncentral_chisquare_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_noncentral_chisquare_ext_c)( - DPCTLSyclQueueRef, - void *, - const _DataType, - const _DataType, - const size_t, - const DPCTLEventVectorRef) = dpnp_rng_noncentral_chisquare_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_rng_normal_c(DPCTLSyclQueueRef q_ref, - void *result_out, - const double mean_in, - const double stddev_in, - const int64_t size, - void *random_state_in, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - sycl::queue *q = reinterpret_cast(q_ref); - - if (!size) { - return event_ref; - } - assert(q != nullptr); - - _DataType *result = static_cast<_DataType *>(result_out); - - // set mean of distribution - const _DataType mean = static_cast<_DataType>(mean_in); - // set standard deviation of distribution - const _DataType stddev = static_cast<_DataType>(stddev_in); - - mkl_rng::gaussian<_DataType> distribution(mean, stddev); - - if (q->get_device().is_cpu()) { - mt19937_struct *random_state = - static_cast(random_state_in); - mkl_rng::mt19937 *engine = - static_cast(random_state->engine); - - // perform generation with MT19937 engine - event_ref = dpnp_rng_generate(distribution, *engine, size, result); - } - else { - mcg59_struct *random_state = - static_cast(random_state_in); - mkl_rng::mcg59 *engine = - static_cast(random_state->engine); - - // perform generation with MCG59 engine - event_ref = dpnp_rng_generate(distribution, *engine, size, result); - } - return event_ref; -} - -template -void dpnp_rng_normal_c(void *result, - const _DataType mean, - const _DataType stddev, - const size_t size) -{ - sycl::queue *q = &DPNP_QUEUE; - DPCTLSyclQueueRef q_ref = reinterpret_cast(q); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = nullptr; - - if (q->get_device().is_cpu()) { - mt19937_struct *mt19937 = new mt19937_struct(); - mt19937->engine = &DPNP_RNG_ENGINE; - - event_ref = dpnp_rng_normal_c<_DataType>(q_ref, result, mean, stddev, - static_cast(size), - mt19937, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); - delete mt19937; - } - else { - // MCG59 engine is assumed to provide a better performance on GPU than - // MT19937 - mcg59_struct *mcg59 = new mcg59_struct(); - mcg59->engine = &DPNP_RNG_MCG59_ENGINE; - - event_ref = dpnp_rng_normal_c<_DataType>(q_ref, result, mean, stddev, - static_cast(size), - mcg59, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); - delete mcg59; - } -} - -template -void (*dpnp_rng_normal_default_c)(void *, - const _DataType, - const _DataType, - const size_t) = dpnp_rng_normal_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_normal_ext_c)(DPCTLSyclQueueRef, - void *, - const double, - const double, - const int64_t, - void *, - const DPCTLEventVectorRef) = - dpnp_rng_normal_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_rng_pareto_c(DPCTLSyclQueueRef q_ref, - void *result, - const double alpha, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - std::vector no_deps; - - const _DataType d_zero = _DataType(0.0); - const _DataType d_one = _DataType(1.0); - _DataType neg_rec_alp = -1.0 / alpha; - - _DataType *result1 = reinterpret_cast<_DataType *>(result); - - mkl_rng::uniform<_DataType> distribution(d_zero, d_one); - auto event_out = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_out.wait(); - - event_out = mkl_vm::powx(q, size, result1, neg_rec_alp, result1, no_deps, - mkl_vm::mode::ha); - event_ref = reinterpret_cast(&event_out); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_pareto_c(void *result, const double alpha, const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_pareto_c<_DataType>( - q_ref, result, alpha, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_pareto_default_c)(void *, - const double, - const size_t) = dpnp_rng_pareto_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_pareto_ext_c)(DPCTLSyclQueueRef, - void *, - const double, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_pareto_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_poisson_c(DPCTLSyclQueueRef q_ref, - void *result, - const double lambda, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - _DataType *result1 = reinterpret_cast<_DataType *>(result); - - mkl_rng::poisson<_DataType> distribution(lambda); - // perform generation - auto event_out = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_ref = reinterpret_cast(&event_out); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_poisson_c(void *result, const double lambda, const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_poisson_c<_DataType>( - q_ref, result, lambda, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_poisson_default_c)(void *, const double, const size_t) = - dpnp_rng_poisson_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_poisson_ext_c)(DPCTLSyclQueueRef, - void *, - const double, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_poisson_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_rng_power_c(DPCTLSyclQueueRef q_ref, - void *result, - const double alpha, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - std::vector no_deps; - - const _DataType d_zero = _DataType(0.0); - const _DataType d_one = _DataType(1.0); - _DataType neg_rec_alp = 1.0 / alpha; - - _DataType *result1 = reinterpret_cast<_DataType *>(result); - - mkl_rng::uniform<_DataType> distribution(d_zero, d_one); - auto event_out = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_out.wait(); - - event_out = mkl_vm::powx(q, size, result1, neg_rec_alp, result1, no_deps, - mkl_vm::mode::ha); - event_ref = reinterpret_cast(&event_out); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_power_c(void *result, const double alpha, const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_power_c<_DataType>( - q_ref, result, alpha, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_power_default_c)(void *, - const double, - const size_t) = dpnp_rng_power_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_power_ext_c)(DPCTLSyclQueueRef, - void *, - const double, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_power_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_rayleigh_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType scale, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - std::vector no_deps; - - const _DataType a = 0.0; - const _DataType beta = 2.0; - - DPNPC_ptr_adapter<_DataType> result1_ptr(q_ref, result, size); - _DataType *result1 = result1_ptr.get_ptr(); - - mkl_rng::exponential<_DataType> distribution(a, beta); - - auto exponential_rng_event = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - auto sqrt_event = mkl_vm::sqrt(q, size, result1, result1, - {exponential_rng_event}, mkl_vm::mode::ha); - auto scal_event = mkl_blas::scal(q, size, scale, result1, 1, {sqrt_event}); - event_ref = reinterpret_cast(&scal_event); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_rayleigh_c(void *result, const _DataType scale, const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_rayleigh_c<_DataType>( - q_ref, result, scale, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_rayleigh_default_c)(void *, const _DataType, const size_t) = - dpnp_rng_rayleigh_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_rayleigh_ext_c)(DPCTLSyclQueueRef, - void *, - const _DataType, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_rayleigh_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_shuffle_c(DPCTLSyclQueueRef q_ref, - void *result, - const size_t itemsize, - const size_t ndim, - const size_t high_dim_size, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!result) { - return event_ref; - } - - if (!size || !ndim || !(high_dim_size > 1)) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter result1_ptr(q_ref, result, size * itemsize, true, - true); - char *result1 = result1_ptr.get_ptr(); - - size_t uvec_size = high_dim_size - 1; - double *Uvec = reinterpret_cast( - sycl::malloc_shared(uvec_size * sizeof(double), q)); - mkl_rng::uniform uniform_distribution(0.0, 1.0); - auto uniform_event = mkl_rng::generate(uniform_distribution, - DPNP_RNG_ENGINE, uvec_size, Uvec); - uniform_event.wait(); - - if (ndim == 1) { - // Fast, statically typed path: shuffle the underlying buffer. - // Only for non-empty, 1d objects of class ndarray (subclasses such - // as MaskedArrays may not support this approach). - void *buf = sycl::malloc_device(itemsize, q); - for (size_t i = uvec_size; i > 0; i--) { - size_t j = (size_t)(floor((i + 1) * Uvec[i - 1])); - if (i != j) { - auto memcpy1 = q.memcpy(buf, result1 + j * itemsize, itemsize); - auto memcpy2 = - q.memcpy(result1 + j * itemsize, result1 + i * itemsize, - itemsize, memcpy1); - q.memcpy(result1 + i * itemsize, buf, itemsize, memcpy2).wait(); - } - } - sycl::free(buf, q); - } - else { - // Multidimensional ndarrays require a bounce buffer. - size_t step_size = - (size / high_dim_size) * itemsize; // size in bytes for x[i] element - void *buf = sycl::malloc_device(step_size, q); - for (size_t i = uvec_size; i > 0; i--) { - size_t j = (size_t)(floor((i + 1) * Uvec[i - 1])); - if (j < i) { - auto memcpy1 = - q.memcpy(buf, result1 + j * step_size, step_size); - auto memcpy2 = - q.memcpy(result1 + j * step_size, result1 + i * step_size, - step_size, memcpy1); - q.memcpy(result1 + i * step_size, buf, step_size, memcpy2) - .wait(); - } - } - sycl::free(buf, q); - } - - sycl::free(Uvec, q); - - return event_ref; -} - -template -void dpnp_rng_shuffle_c(void *result, - const size_t itemsize, - const size_t ndim, - const size_t high_dim_size, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_shuffle_c<_DataType>( - q_ref, result, itemsize, ndim, high_dim_size, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_shuffle_default_c)(void *, - const size_t, - const size_t, - const size_t, - const size_t) = - dpnp_rng_shuffle_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_shuffle_ext_c)(DPCTLSyclQueueRef, - void *, - const size_t, - const size_t, - const size_t, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_shuffle_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_standard_cauchy_c(DPCTLSyclQueueRef q_ref, - void *result, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - _DataType *result1 = reinterpret_cast<_DataType *>(result); - - const _DataType displacement = _DataType(0.0); - - const _DataType scalefactor = _DataType(1.0); - - mkl_rng::cauchy<_DataType> distribution(displacement, scalefactor); - // perform generation - auto event_out = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_ref = reinterpret_cast(&event_out); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_standard_cauchy_c(void *result, const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_standard_cauchy_c<_DataType>( - q_ref, result, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_standard_cauchy_default_c)(void *, const size_t) = - dpnp_rng_standard_cauchy_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_standard_cauchy_ext_c)(DPCTLSyclQueueRef, - void *, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_standard_cauchy_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_standard_exponential_c(DPCTLSyclQueueRef q_ref, - void *result, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - // set displacement a - const _DataType beta = (_DataType(1.0)); - - event_ref = - dpnp_rng_exponential_c(q_ref, result, beta, size, dep_event_vec_ref); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_standard_exponential_c(void *result, const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_standard_exponential_c<_DataType>( - q_ref, result, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_standard_exponential_default_c)(void *, const size_t) = - dpnp_rng_standard_exponential_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_standard_exponential_ext_c)( - DPCTLSyclQueueRef, - void *, - const size_t, - const DPCTLEventVectorRef) = dpnp_rng_standard_exponential_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_standard_gamma_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType shape, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - const _DataType scale = _DataType(1.0); - - event_ref = - dpnp_rng_gamma_c(q_ref, result, shape, scale, size, dep_event_vec_ref); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_standard_gamma_c(void *result, - const _DataType shape, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_standard_gamma_c<_DataType>( - q_ref, result, shape, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_standard_gamma_default_c)(void *, - const _DataType, - const size_t) = - dpnp_rng_standard_gamma_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_standard_gamma_ext_c)(DPCTLSyclQueueRef, - void *, - const _DataType, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_standard_gamma_c<_DataType>; - -template -void dpnp_rng_standard_normal_c(void *result, size_t size) -{ - dpnp_rng_normal_c(result, _DataType(0.0), _DataType(1.0), size); -} - -template -void (*dpnp_rng_standard_normal_default_c)(void *, const size_t) = - dpnp_rng_standard_normal_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_standard_t_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType df, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size || !result) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - _DataType *result1 = reinterpret_cast<_DataType *>(result); - const _DataType d_zero = 0.0, d_one = 1.0; - _DataType shape = df / 2; - _DataType *sn = nullptr; - - mkl_rng::gamma<_DataType> gamma_distribution(shape, d_zero, 1.0 / shape); - auto gamma_distr_event = - mkl_rng::generate(gamma_distribution, DPNP_RNG_ENGINE, size, result1); - - auto invsqrt_event = mkl_vm::invsqrt(q, size, result1, result1, - {gamma_distr_event}, mkl_vm::mode::ha); - - sn = reinterpret_cast<_DataType *>( - sycl::malloc_shared(size * sizeof(_DataType), q)); - - mkl_rng::gaussian<_DataType> gaussian_distribution(d_zero, d_one); - auto gaussian_distr_event = - mkl_rng::generate(gaussian_distribution, DPNP_RNG_ENGINE, size, sn); - - auto event_out = - mkl_vm::mul(q, size, result1, sn, result1, - {invsqrt_event, gaussian_distr_event}, mkl_vm::mode::ha); - event_out.wait(); - - sycl::free(sn, q); - - return event_ref; -} - -template -void dpnp_rng_standard_t_c(void *result, const _DataType df, const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_standard_t_c<_DataType>( - q_ref, result, df, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_standard_t_default_c)(void *, const _DataType, const size_t) = - dpnp_rng_standard_t_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_standard_t_ext_c)(DPCTLSyclQueueRef, - void *, - const _DataType, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_standard_t_c<_DataType>; - -template -class dpnp_rng_triangular_ration_acceptance_c_kernel; - -template -DPCTLSyclEventRef - dpnp_rng_triangular_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType x_min, - const _DataType x_mode, - const _DataType x_max, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - _DataType *result1 = reinterpret_cast<_DataType *>(result); - const _DataType d_zero = (_DataType(0)); - const _DataType d_one = (_DataType(1)); - - _DataType ratio, lpr, rpr; - - mkl_rng::uniform<_DataType> uniform_distribution(d_zero, d_one); - auto event_uniform = - mkl_rng::generate(uniform_distribution, DPNP_RNG_ENGINE, size, result1); - - { - _DataType wtot, wl, wr; - - wtot = x_max - x_min; - wl = x_mode - x_min; - wr = x_max - x_mode; - - ratio = wl / wtot; - lpr = wl * wtot; - rpr = wr * wtot; - } - - if (!(0 <= ratio && ratio <= 1)) { - throw std::runtime_error( - "DPNP RNG Error: dpnp_rng_triangular_c() failed."); - } - - sycl::range<1> gws(size); - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { - size_t i = global_id[0]; - if (ratio <= 0) { - result1[i] = x_max - sycl::sqrt(result1[i] * rpr); - } - else if (ratio >= 1) { - result1[i] = x_min + sycl::sqrt(result1[i] * lpr); - } - else { - _DataType ui = result1[i]; - if (ui > ratio) { - result1[i] = x_max - sycl::sqrt((1.0 - ui) * rpr); - } - else { - result1[i] = x_min + sycl::sqrt(ui * lpr); - } - } - }; - auto kernel_func = [&](sycl::handler &cgh) { - cgh.depends_on({event_uniform}); - cgh.parallel_for< - class dpnp_rng_triangular_ration_acceptance_c_kernel<_DataType>>( - gws, kernel_parallel_for_func); - }; - auto event_ration_acceptance = q.submit(kernel_func); - event_ref = reinterpret_cast(&event_ration_acceptance); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_triangular_c(void *result, - const _DataType x_min, - const _DataType x_mode, - const _DataType x_max, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_triangular_c<_DataType>( - q_ref, result, x_min, x_mode, x_max, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_triangular_default_c)(void *, - const _DataType, - const _DataType, - const _DataType, - const size_t) = - dpnp_rng_triangular_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_triangular_ext_c)(DPCTLSyclQueueRef, - void *, - const _DataType, - const _DataType, - const _DataType, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_triangular_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_uniform_c(DPCTLSyclQueueRef q_ref, - void *result_out, - const double low, - const double high, - const int64_t size, - void *random_state_in, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue *q = reinterpret_cast(q_ref); - - _DataType *result = static_cast<_DataType *>(result_out); - - // set left bound of distribution - const _DataType a = static_cast<_DataType>(low); - // set right bound of distribution - const _DataType b = static_cast<_DataType>(high); - - if (q->get_device().is_cpu()) { - mt19937_struct *random_state = - static_cast(random_state_in); - mkl_rng::mt19937 *engine = - static_cast(random_state->engine); - - // perform generation with MT19937 engine - event_ref = dpnp_rng_generate_uniform(*engine, q, a, b, size, result); - } - else { - mcg59_struct *random_state = - static_cast(random_state_in); - mkl_rng::mcg59 *engine = - static_cast(random_state->engine); - - // perform generation with MCG59 engine - event_ref = dpnp_rng_generate_uniform(*engine, q, a, b, size, result); - } - return event_ref; -} - -template -void dpnp_rng_uniform_c(void *result, - const long low, - const long high, - const size_t size) -{ - sycl::queue *q = &DPNP_QUEUE; - DPCTLSyclQueueRef q_ref = reinterpret_cast(q); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = nullptr; - - if (q->get_device().is_cpu()) { - mt19937_struct *mt19937 = new mt19937_struct(); - mt19937->engine = &DPNP_RNG_ENGINE; - - event_ref = dpnp_rng_uniform_c<_DataType>( - q_ref, result, static_cast(low), static_cast(high), - static_cast(size), mt19937, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); - delete mt19937; - } - else { - // MCG59 engine is assumed to provide a better performance on GPU than - // MT19937 - mcg59_struct *mcg59 = new mcg59_struct(); - mcg59->engine = &DPNP_RNG_MCG59_ENGINE; - - event_ref = dpnp_rng_uniform_c<_DataType>( - q_ref, result, static_cast(low), static_cast(high), - static_cast(size), mcg59, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); - delete mcg59; - } -} - -template -void (*dpnp_rng_uniform_default_c)(void *, - const long, - const long, - const size_t) = - dpnp_rng_uniform_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_uniform_ext_c)(DPCTLSyclQueueRef, - void *, - const double, - const double, - const int64_t, - void *, - const DPCTLEventVectorRef) = - dpnp_rng_uniform_c<_DataType>; - -#ifndef M_PI -/* 128-bits worth of pi */ -#define M_PI 3.141592653589793238462643383279502884197 -#endif - -template -DPCTLSyclEventRef - dpnp_rng_vonmises_large_kappa_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType mu, - const _DataType kappa, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size || !result) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter<_DataType> result1_ptr(q_ref, result, size, true, true); - _DataType *result1 = result1_ptr.get_ptr(); - - _DataType r_over_two_kappa, recip_two_kappa; - _DataType s_minus_one, hpt, r_over_two_kappa_minus_one, rho_minus_one; - _DataType *Uvec = nullptr; - _DataType *Vvec = nullptr; - const _DataType d_zero = 0.0, d_one = 1.0; - - assert(kappa > 1.0); - - recip_two_kappa = 1 / (2 * kappa); - - /* variables here are dwindling to zero as kappa grows */ - hpt = sqrt(1 + recip_two_kappa * recip_two_kappa); - r_over_two_kappa_minus_one = - recip_two_kappa * (1 + recip_two_kappa / (1 + hpt)); - r_over_two_kappa = 1 + r_over_two_kappa_minus_one; - rho_minus_one = r_over_two_kappa_minus_one - - sqrt(2 * r_over_two_kappa * recip_two_kappa); - s_minus_one = rho_minus_one * (0.5 * rho_minus_one / (1 + rho_minus_one)); - - Uvec = reinterpret_cast<_DataType *>( - sycl::malloc_shared(size * sizeof(_DataType), q)); - Vvec = reinterpret_cast<_DataType *>( - sycl::malloc_shared(size * sizeof(_DataType), q)); - - for (size_t n = 0; n < size;) { - size_t diff_size = size - n; - mkl_rng::uniform<_DataType> uniform_distribution_u(d_zero, 0.5 * M_PI); - auto event_out = mkl_rng::generate(uniform_distribution_u, - DPNP_RNG_ENGINE, diff_size, Uvec); - event_out.wait(); - // TODO - // use deps case - mkl_rng::uniform<_DataType> uniform_distribution_v(d_zero, d_one); - event_out = mkl_rng::generate(uniform_distribution_v, DPNP_RNG_ENGINE, - diff_size, Vvec); - event_out.wait(); - - // TODO - // kernel - for (size_t i = 0; i < diff_size; i++) { - _DataType sn, cn, sn2, cn2; - _DataType neg_W_minus_one, V, Y; - - sn = sin(Uvec[i]); - cn = cos(Uvec[i]); - V = Vvec[i]; - sn2 = sn * sn; - cn2 = cn * cn; - - neg_W_minus_one = s_minus_one * sn2 / (0.5 * s_minus_one + cn2); - Y = kappa * (s_minus_one + neg_W_minus_one); - - if ((Y * (2 - Y) >= V) || (log(Y / V) + 1 >= Y)) { - Y = neg_W_minus_one * (2 - neg_W_minus_one); - if (Y < 0) - Y = 0.0; - else if (Y > 1.0) - Y = 1.0; - - result1[n++] = asin(sqrt(Y)); - } - } - } - - sycl::free(Uvec, q); - - mkl_rng::uniform<_DataType> uniform_distribution(d_zero, d_one); - auto uniform_distr_event = - mkl_rng::generate(uniform_distribution, DPNP_RNG_ENGINE, size, Vvec); - - sycl::range<1> gws(size); - - auto paral_kernel_acceptance = [&](sycl::handler &cgh) { - cgh.depends_on({uniform_distr_event}); - cgh.parallel_for(gws, [=](sycl::id<1> global_id) { - size_t i = global_id[0]; - double mod, resi; - resi = (Vvec[i] < 0.5) ? mu - result1[i] : mu + result1[i]; - mod = sycl::fabs(resi); - mod = (sycl::fmod(mod + M_PI, 2 * M_PI) - M_PI); - result1[i] = (resi < 0) ? -mod : mod; - }); - }; - auto acceptance_event = q.submit(paral_kernel_acceptance); - acceptance_event.wait(); - - sycl::free(Vvec, q); - return event_ref; -} - -template -void dpnp_rng_vonmises_large_kappa_c(void *result, - const _DataType mu, - const _DataType kappa, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_vonmises_large_kappa_c<_DataType>( - q_ref, result, mu, kappa, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_vonmises_large_kappa_default_c)(void *, - const _DataType, - const _DataType, - const size_t) = - dpnp_rng_vonmises_large_kappa_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_vonmises_large_kappa_ext_c)( - DPCTLSyclQueueRef, - void *, - const _DataType, - const _DataType, - const size_t, - const DPCTLEventVectorRef) = dpnp_rng_vonmises_large_kappa_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_vonmises_small_kappa_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType mu, - const _DataType kappa, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size || !result) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter<_DataType> result1_ptr(q_ref, result, size, true, true); - _DataType *result1 = result1_ptr.get_ptr(); - - _DataType rho_over_kappa, rho, r, s_kappa; - _DataType *Uvec = nullptr; - _DataType *Vvec = nullptr; - - const _DataType d_zero = 0.0, d_one = 1.0; - - assert(0. < kappa); - assert(kappa <= 1.0); - - r = 1 + sqrt(1 + 4 * kappa * kappa); - rho_over_kappa = (2) / (r + sqrt(2 * r)); - rho = rho_over_kappa * kappa; - - /* s times kappa */ - s_kappa = (1 + rho * rho) / (2 * rho_over_kappa); - - Uvec = reinterpret_cast<_DataType *>( - sycl::malloc_shared(size * sizeof(_DataType), q)); - Vvec = reinterpret_cast<_DataType *>( - sycl::malloc_shared(size * sizeof(_DataType), q)); - - for (size_t n = 0; n < size;) { - size_t diff_size = size - n; - mkl_rng::uniform<_DataType> uniform_distribution_u(d_zero, M_PI); - auto event_out = mkl_rng::generate(uniform_distribution_u, - DPNP_RNG_ENGINE, diff_size, Uvec); - event_out.wait(); - // TODO - // use deps case - mkl_rng::uniform<_DataType> uniform_distribution_v(d_zero, d_one); - event_out = mkl_rng::generate(uniform_distribution_v, DPNP_RNG_ENGINE, - diff_size, Vvec); - event_out.wait(); - - // TODO - // kernel - for (size_t i = 0; i < diff_size; i++) { - _DataType Z, W, Y, V; - Z = cos(Uvec[i]); - V = Vvec[i]; - W = (kappa + s_kappa * Z) / (s_kappa + kappa * Z); - Y = s_kappa - kappa * W; - if ((Y * (2 - Y) >= V) || (log(Y / V) + 1 >= Y)) { - result1[n++] = acos(W); - } - } - } - - sycl::free(Uvec, q); - - mkl_rng::uniform<_DataType> uniform_distribution(d_zero, d_one); - auto uniform_distr_event = - mkl_rng::generate(uniform_distribution, DPNP_RNG_ENGINE, size, Vvec); - - sycl::range<1> gws(size); - auto paral_kernel_acceptance = [&](sycl::handler &cgh) { - cgh.depends_on({uniform_distr_event}); - cgh.parallel_for(gws, [=](sycl::id<1> global_id) { - size_t i = global_id[0]; - double mod, resi; - resi = (Vvec[i] < 0.5) ? mu - result1[i] : mu + result1[i]; - mod = sycl::fabs(resi); - mod = (sycl::fmod(mod + M_PI, 2 * M_PI) - M_PI); - result1[i] = (resi < 0) ? -mod : mod; - }); - }; - auto acceptance_event = q.submit(paral_kernel_acceptance); - acceptance_event.wait(); - - sycl::free(Vvec, q); - return event_ref; -} - -template -void dpnp_rng_vonmises_small_kappa_c(void *result, - const _DataType mu, - const _DataType kappa, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_vonmises_small_kappa_c<_DataType>( - q_ref, result, mu, kappa, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_vonmises_small_kappa_default_c)(void *, - const _DataType, - const _DataType, - const size_t) = - dpnp_rng_vonmises_small_kappa_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_vonmises_small_kappa_ext_c)( - DPCTLSyclQueueRef, - void *, - const _DataType, - const _DataType, - const size_t, - const DPCTLEventVectorRef) = dpnp_rng_vonmises_small_kappa_c<_DataType>; - -/* Vonmisses uses the rejection algorithm compared against the wrapped - Cauchy distribution suggested by Best and Fisher and documented in - Chapter 9 of Luc's Non-Uniform Random Variate Generation. - http://cg.scs.carleton.ca/~luc/rnbookindex.html - (but corrected to match the algorithm in R and Python) -*/ -template -DPCTLSyclEventRef - dpnp_rng_vonmises_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType mu, - const _DataType kappa, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - sycl::queue q = *(reinterpret_cast(q_ref)); - - if (kappa > 1.0) - dpnp_rng_vonmises_large_kappa_c<_DataType>(result, mu, kappa, size); - else - dpnp_rng_vonmises_small_kappa_c<_DataType>(result, mu, kappa, size); - - return event_ref; -} - -template -void dpnp_rng_vonmises_c(void *result, - const _DataType mu, - const _DataType kappa, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_vonmises_c<_DataType>( - q_ref, result, mu, kappa, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_vonmises_default_c)(void *, - const _DataType, - const _DataType, - const size_t) = - dpnp_rng_vonmises_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_vonmises_ext_c)(DPCTLSyclQueueRef, - void *, - const _DataType, - const _DataType, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_vonmises_c<_DataType>; - -template -class dpnp_rng_wald_acceptance_kernel1; - -template -class dpnp_rng_wald_acceptance_kernel2; - -template -DPCTLSyclEventRef dpnp_rng_wald_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType mean, - const _DataType scale, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - _DataType *result1 = reinterpret_cast<_DataType *>(result); - _DataType *uvec = nullptr; - - const _DataType d_zero = _DataType(0.0); - const _DataType d_one = _DataType(1.0); - _DataType gsc = sqrt(0.5 * mean / scale); - - mkl_rng::gaussian<_DataType> gaussian_distribution(d_zero, gsc); - - auto gaussian_dstr_event = mkl_rng::generate( - gaussian_distribution, DPNP_RNG_ENGINE, size, result1); - - /* Y = mean/(2 scale) * Z^2 */ - auto sqr_event = mkl_vm::sqr(q, size, result1, result1, - {gaussian_dstr_event}, mkl_vm::mode::ha); - - sycl::range<1> gws(size); - auto acceptance_kernel1 = [=](sycl::id<1> global_id) { - size_t i = global_id[0]; - if (result1[i] <= 2.0) { - result1[i] = - 1.0 + result1[i] + sycl::sqrt(result1[i] * (result1[i] + 2.0)); - } - else { - result1[i] = - 1.0 + result1[i] * (1.0 + sycl::sqrt(1.0 + 2.0 / result1[i])); - } - }; - auto parallel_for_acceptance1 = [&](sycl::handler &cgh) { - cgh.depends_on({sqr_event}); - cgh.parallel_for>( - gws, acceptance_kernel1); - }; - auto event_ration_acceptance = q.submit(parallel_for_acceptance1); - - uvec = reinterpret_cast<_DataType *>( - sycl::malloc_shared(size * sizeof(_DataType), q)); - - mkl_rng::uniform<_DataType> uniform_distribution(d_zero, d_one); - auto uniform_distr_event = - mkl_rng::generate(uniform_distribution, DPNP_RNG_ENGINE, size, uvec); - - auto acceptance_kernel2 = [=](sycl::id<1> global_id) { - size_t i = global_id[0]; - if (uvec[i] * (1.0 + result1[i]) <= result1[i]) - result1[i] = mean / result1[i]; - else - result1[i] = mean * result1[i]; - }; - auto parallel_for_acceptance2 = [&](sycl::handler &cgh) { - cgh.depends_on({event_ration_acceptance, uniform_distr_event}); - cgh.parallel_for>( - gws, acceptance_kernel2); - }; - auto event_out = q.submit(parallel_for_acceptance2); - event_out.wait(); - - sycl::free(uvec, q); - - return event_ref; -} - -template -void dpnp_rng_wald_c(void *result, - const _DataType mean, - const _DataType scale, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_wald_c<_DataType>( - q_ref, result, mean, scale, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_wald_default_c)(void *, - const _DataType, - const _DataType, - const size_t) = dpnp_rng_wald_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_wald_ext_c)(DPCTLSyclQueueRef, - void *, - const _DataType, - const _DataType, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_wald_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_rng_weibull_c(DPCTLSyclQueueRef q_ref, - void *result, - const double alpha, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - sycl::event event_out; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - if (alpha == 0) { - event_ref = - dpnp_zeros_c<_DataType>(q_ref, result, size, dep_event_vec_ref); - } - else { - _DataType *result1 = reinterpret_cast<_DataType *>(result); - const _DataType a = (_DataType(0.0)); - const _DataType beta = (_DataType(1.0)); - - mkl_rng::weibull<_DataType> distribution(alpha, a, beta); - event_out = - mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_ref = reinterpret_cast(&event_out); - } - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_rng_weibull_c(void *result, const double alpha, const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_rng_weibull_c<_DataType>( - q_ref, result, alpha, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_weibull_default_c)(void *, const double, const size_t) = - dpnp_rng_weibull_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_weibull_ext_c)(DPCTLSyclQueueRef, - void *, - const double, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_weibull_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_rng_zipf_c(DPCTLSyclQueueRef q_ref, - void *result, - const _DataType a, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - sycl::event event_out; - - size_t i, n_accepted, batch_size; - _DataType T, U, V, am1, b; - _DataType *Uvec = nullptr, *Vvec = nullptr; - long X; - const _DataType d_zero = 0.0; - const _DataType d_one = 1.0; - - DPNPC_ptr_adapter<_DataType> result1_ptr(q_ref, result, size, true, true); - _DataType *result1 = result1_ptr.get_ptr(); - - am1 = a - d_one; - b = pow(2.0, am1); - - Uvec = reinterpret_cast<_DataType *>( - sycl::malloc_shared(size * 2 * sizeof(_DataType), q)); - Vvec = Uvec + size; - - // TODO - // kernel for acceptance - for (n_accepted = 0; n_accepted < size;) { - batch_size = size - n_accepted; - - mkl_rng::uniform<_DataType> uniform_distribution(d_zero, d_one); - event_out = mkl_rng::generate(uniform_distribution, DPNP_RNG_ENGINE, - batch_size, Uvec); - event_out.wait(); - event_out = mkl_rng::generate(uniform_distribution, DPNP_RNG_ENGINE, - batch_size, Vvec); - event_out.wait(); - for (i = 0; i < batch_size; i++) { - U = d_one - Uvec[i]; - V = Vvec[i]; - X = (long)floor(pow(U, (-1.0) / am1)); - /* The real result may be above what can be represented in a signed - * long. It will get casted to -sys.maxint-1. Since this is - * a straightforward rejection algorithm, we can just reject this - * value in the rejection condition below. This function then models - * a Zipf distribution truncated to sys.maxint. - */ - T = pow(d_one + d_one / X, am1); - if ((X > 0) && ((V * X) * (T - d_one) / (b - d_one) <= T / b)) { - result1[n_accepted++] = X; - } - } - } - - sycl::free(Uvec, q); - - return event_ref; -} - -template -void dpnp_rng_zipf_c(void *result, const _DataType a, const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_rng_zipf_c<_DataType>(q_ref, result, a, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_rng_zipf_default_c)(void *, - const _DataType, - const size_t) = dpnp_rng_zipf_c<_DataType>; - -template -DPCTLSyclEventRef (*dpnp_rng_zipf_ext_c)(DPCTLSyclQueueRef, - void *, - const _DataType, - const size_t, - const DPCTLEventVectorRef) = - dpnp_rng_zipf_c<_DataType>; - -void func_map_init_random(func_map_t &fmap) -{ - fmap[DPNPFuncName::DPNP_FN_RNG_BETA][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_beta_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_BETA_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_beta_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_BINOMIAL][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_rng_binomial_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_BINOMIAL_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_rng_binomial_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_CHISQUARE][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_chisquare_default_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_CHISQUARE][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_rng_chisquare_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_CHISQUARE_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_chisquare_ext_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_CHISQUARE_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_rng_chisquare_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_EXPONENTIAL][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_exponential_default_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_EXPONENTIAL][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_rng_exponential_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_EXPONENTIAL_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_exponential_ext_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_EXPONENTIAL_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_rng_exponential_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_F][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_f_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_F_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_f_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_GAMMA][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_gamma_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_GAMMA_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_gamma_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_GAUSSIAN][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_gaussian_default_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_GAUSSIAN][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_rng_gaussian_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_GAUSSIAN_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_gaussian_ext_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_GAUSSIAN_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_rng_gaussian_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_GEOMETRIC][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_rng_geometric_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_GEOMETRIC_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_rng_geometric_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_GUMBEL][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_gumbel_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_GUMBEL_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_gumbel_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_HYPERGEOMETRIC][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_rng_hypergeometric_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_HYPERGEOMETRIC_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_rng_hypergeometric_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_LAPLACE][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_laplace_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_LAPLACE_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_laplace_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_LOGISTIC][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_logistic_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_LOGISTIC_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_logistic_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_LOGNORMAL][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_lognormal_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_LOGNORMAL_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_lognormal_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_MULTINOMIAL][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_rng_multinomial_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_MULTINOMIAL_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_rng_multinomial_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_MULTIVARIATE_NORMAL][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_multivariate_normal_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_NEGATIVE_BINOMIAL][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_rng_negative_binomial_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_NEGATIVE_BINOMIAL_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_rng_negative_binomial_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_NONCENTRAL_CHISQUARE][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_noncentral_chisquare_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_NONCENTRAL_CHISQUARE_EXT][eft_DBL][eft_DBL] = - {eft_DBL, (void *)dpnp_rng_noncentral_chisquare_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_NORMAL][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_normal_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_NORMAL_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_normal_ext_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_NORMAL_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_rng_normal_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_PARETO][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_pareto_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_PARETO_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_pareto_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_POISSON][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_rng_poisson_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_POISSON_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_rng_poisson_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_POWER][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_power_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_POWER_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_power_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_RAYLEIGH][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_rayleigh_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_RAYLEIGH_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_rayleigh_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_SHUFFLE][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_shuffle_default_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_SHUFFLE][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_rng_shuffle_default_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_SHUFFLE][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_rng_shuffle_default_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_SHUFFLE][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_rng_shuffle_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_SHUFFLE_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_shuffle_ext_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_SHUFFLE_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_rng_shuffle_ext_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_SHUFFLE_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_rng_shuffle_ext_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_SHUFFLE_EXT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_rng_shuffle_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_SRAND][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_srand_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_CAUCHY][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_standard_cauchy_default_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_CAUCHY][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_rng_standard_cauchy_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_CAUCHY_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_standard_cauchy_ext_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_CAUCHY_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_rng_standard_cauchy_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_EXPONENTIAL][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_standard_exponential_default_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_EXPONENTIAL][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_rng_standard_exponential_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_EXPONENTIAL_EXT][eft_DBL][eft_DBL] = - {eft_DBL, (void *)dpnp_rng_standard_exponential_ext_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_EXPONENTIAL_EXT][eft_FLT][eft_FLT] = - {eft_FLT, (void *)dpnp_rng_standard_exponential_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_GAMMA][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_standard_gamma_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_GAMMA_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_standard_gamma_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_NORMAL][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_standard_normal_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_T][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_standard_t_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_T_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_standard_t_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_TRIANGULAR][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_triangular_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_TRIANGULAR_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_triangular_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_UNIFORM][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_uniform_default_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_UNIFORM][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_rng_uniform_default_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_UNIFORM][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_rng_uniform_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_UNIFORM_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_uniform_ext_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_UNIFORM_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_rng_uniform_ext_c}; - fmap[DPNPFuncName::DPNP_FN_RNG_UNIFORM_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_rng_uniform_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_VONMISES][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_vonmises_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_VONMISES_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_vonmises_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_WALD][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_wald_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_WALD_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_wald_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_WEIBULL][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_weibull_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_WEIBULL_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_weibull_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_ZIPF][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_zipf_default_c}; - - fmap[DPNPFuncName::DPNP_FN_RNG_ZIPF_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_rng_zipf_ext_c}; - - return; -} diff --git a/dpnp/backend/kernels/dpnp_krnl_sorting.cpp b/dpnp/backend/kernels/dpnp_krnl_sorting.cpp index a1495dfb0279..d01abcd324ea 100644 --- a/dpnp/backend/kernels/dpnp_krnl_sorting.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_sorting.cpp @@ -26,11 +26,15 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** -#include +#include +#include +#include +#include + +#include #include "dpnp_fptr.hpp" -#include "dpnpc_memory_adapter.hpp" -#include "queue_sycl.hpp" +#include "dpnp_pstl.hpp" #include template @@ -94,31 +98,6 @@ DPCTLSyclEventRef dpnp_partition_c(DPCTLSyclQueueRef q_ref, return event_ref; } -template -void dpnp_partition_c(void *array1_in, - void *array2_in, - void *result1, - const size_t kth, - const shape_elem_type *shape_, - const size_t ndim) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_partition_c<_DataType>(q_ref, array1_in, array2_in, result1, kth, - shape_, ndim, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_partition_default_c)(void *, - void *, - void *, - const size_t, - const shape_elem_type *, - const size_t) = dpnp_partition_c<_DataType>; - template DPCTLSyclEventRef (*dpnp_partition_ext_c)(DPCTLSyclQueueRef, void *, @@ -132,15 +111,6 @@ DPCTLSyclEventRef (*dpnp_partition_ext_c)(DPCTLSyclQueueRef, void func_map_init_sorting(func_map_t &fmap) { - fmap[DPNPFuncName::DPNP_FN_PARTITION][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_partition_default_c}; - fmap[DPNPFuncName::DPNP_FN_PARTITION][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_partition_default_c}; - fmap[DPNPFuncName::DPNP_FN_PARTITION][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_partition_default_c}; - fmap[DPNPFuncName::DPNP_FN_PARTITION][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_partition_default_c}; - fmap[DPNPFuncName::DPNP_FN_PARTITION_EXT][eft_BLN][eft_BLN] = { eft_BLN, (void *)dpnp_partition_ext_c}; fmap[DPNPFuncName::DPNP_FN_PARTITION_EXT][eft_INT][eft_INT] = { diff --git a/dpnp/backend/src/dpnp_fptr.hpp b/dpnp/backend/src/dpnp_fptr.hpp index 15d6b7ab20ee..9291daaf2d5e 100644 --- a/dpnp/backend/src/dpnp_fptr.hpp +++ b/dpnp/backend/src/dpnp_fptr.hpp @@ -158,9 +158,6 @@ class dpnp_less_comp /** * FPTR interface initialization functions */ -void func_map_init_arraycreation(func_map_t &fmap); -void func_map_init_linalg(func_map_t &fmap); -void func_map_init_random(func_map_t &fmap); void func_map_init_sorting(func_map_t &fmap); #endif // BACKEND_FPTR_H diff --git a/dpnp/backend/src/dpnp_iface_fptr.cpp b/dpnp/backend/src/dpnp_iface_fptr.cpp index ab1769c7e56c..78c3b4d645ff 100644 --- a/dpnp/backend/src/dpnp_iface_fptr.cpp +++ b/dpnp/backend/src/dpnp_iface_fptr.cpp @@ -98,9 +98,6 @@ static func_map_t func_map_init() { func_map_t fmap; - func_map_init_arraycreation(fmap); - func_map_init_linalg(fmap); - func_map_init_random(fmap); func_map_init_sorting(fmap); return fmap; diff --git a/dpnp/backend/src/dpnp_random_state.cpp b/dpnp/backend/src/dpnp_random_state.cpp deleted file mode 100644 index 3abfbad14403..000000000000 --- a/dpnp/backend/src/dpnp_random_state.cpp +++ /dev/null @@ -1,87 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2022, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// - Neither the name of the copyright holder nor the names of its contributors -// may be used to endorse or promote products derived from this software -// without specific prior written permission. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -#include - -#include "dpnp_random_state.hpp" -#include - -namespace mkl_rng = oneapi::mkl::rng; - -void MT19937_InitScalarSeed(mt19937_struct *mt19937, - DPCTLSyclQueueRef q_ref, - uint32_t seed) -{ - sycl::queue *q = reinterpret_cast(q_ref); - mt19937->engine = new mkl_rng::mt19937(*q, seed); -} - -void MT19937_InitVectorSeed(mt19937_struct *mt19937, - DPCTLSyclQueueRef q_ref, - uint32_t *seed, - unsigned int n) -{ - sycl::queue *q = reinterpret_cast(q_ref); - - switch (n) { - case 1: - mt19937->engine = new mkl_rng::mt19937(*q, {seed[0]}); - break; - case 2: - mt19937->engine = new mkl_rng::mt19937(*q, {seed[0], seed[1]}); - break; - case 3: - mt19937->engine = new mkl_rng::mt19937(*q, {seed[0], seed[1], seed[2]}); - break; - default: - // TODO need to get rid of the limitation for seed vector length - throw std::runtime_error("Too long seed vector"); - } -} - -void MT19937_Delete(mt19937_struct *mt19937) -{ - mkl_rng::mt19937 *engine = static_cast(mt19937->engine); - mt19937->engine = nullptr; - delete engine; -} - -void MCG59_InitScalarSeed(mcg59_struct *mcg59, - DPCTLSyclQueueRef q_ref, - uint64_t seed) -{ - sycl::queue *q = reinterpret_cast(q_ref); - mcg59->engine = new mkl_rng::mcg59(*q, seed); -} - -void MCG59_Delete(mcg59_struct *mcg59) -{ - mkl_rng::mcg59 *engine = static_cast(mcg59->engine); - mcg59->engine = nullptr; - delete engine; -} diff --git a/dpnp/backend/src/dpnp_random_state.hpp b/dpnp/backend/src/dpnp_random_state.hpp deleted file mode 100644 index bfd46031a14b..000000000000 --- a/dpnp/backend/src/dpnp_random_state.hpp +++ /dev/null @@ -1,143 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2022, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// - Neither the name of the copyright holder nor the names of its contributors -// may be used to endorse or promote products derived from this software -// without specific prior written permission. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -/* - * This header file is for interface Cython with C++. - * It should not contains any backend specific headers (like SYCL or math - * library) because all included headers will be exposed in Cython compilation - * procedure - * - * We would like to avoid backend specific things in higher level Cython - * modules. Any backend interface functions and types should be defined here. - * - * Also, this file should contains documentation on functions and types - * which are used in the interface - */ - -#pragma once -#ifndef BACKEND_RANDOM_STATE_H // Cython compatibility -#define BACKEND_RANDOM_STATE_H - -#ifdef _WIN32 -#define INP_DLLEXPORT __declspec(dllexport) -#else -#define INP_DLLEXPORT -#endif - -#include - -// Structure storing a base MKL engine -struct engine_struct -{ - void *engine; -}; - -// Structure storing MKL engine for MT199374x32x10 generator -struct mt19937_struct : engine_struct -{ -}; - -// Structure storing MKL engine for MCG59 generator -struct mcg59_struct : engine_struct -{ -}; - -/** - * @ingroup BACKEND_API - * @brief Create a MKL engine from scalar seed. - * - * Invoke a common seed initialization of the engine for MT199374x32x10 - * algorithm. - * - * @param [in] mt19937 A structure with MKL engine which will be filled - * with generated value by MKL. - * @param [in] q_ref A reference on SYCL queue which will be used to - * obtain random numbers. - * @param [in] seed An initial condition of the generator state. - */ -INP_DLLEXPORT void MT19937_InitScalarSeed(mt19937_struct *mt19937, - DPCTLSyclQueueRef q_ref, - uint32_t seed = 1); - -/** - * @ingroup BACKEND_API - * @brief Create a MKL engine from seed vector. - * - * Invoke an extended seed initialization of the engine for MT199374x32x10 - * algorithm.. - * - * @param [in] mt19937 A structure with MKL engine which will be filled - * with generated value by MKL. - * @param [in] q_ref A reference on SYCL queue which will be used to - * obtain random numbers. - * @param [in] seed A vector with the initial conditions of the - * generator state. - * @param [in] n Length of the vector. - */ -INP_DLLEXPORT void MT19937_InitVectorSeed(mt19937_struct *mt19937, - DPCTLSyclQueueRef q_ref, - uint32_t *seed, - unsigned int n); - -/** - * @ingroup BACKEND_API - * @brief Release a MKL engine. - * - * Release all resource required for storing of the MKL engine. - * - * @param [in] mt19937 A structure with the MKL engine. - */ -INP_DLLEXPORT void MT19937_Delete(mt19937_struct *mt19937); - -/** - * @ingroup BACKEND_API - * @brief Create a MKL engine from scalar seed. - * - * Invoke a common seed initialization of the engine for MCG59 algorithm. - * - * @param [in] mcg59 A structure with MKL engine which will be filled - * with generated value by MKL. - * @param [in] q_ref A reference on SYCL queue which will be used to - * obtain random numbers. - * @param [in] seed An initial condition of the generator state. - */ -INP_DLLEXPORT void MCG59_InitScalarSeed(mcg59_struct *mcg59, - DPCTLSyclQueueRef q_ref, - uint64_t seed); - -/** - * @ingroup BACKEND_API - * @brief Release a MKL engine. - * - * Release all resource required for storing of the MKL engine. - * - * @param [in] mcg59 A structure with the MKL engine. - */ -INP_DLLEXPORT void MCG59_Delete(mcg59_struct *mcg59); - -#endif // BACKEND_RANDOM_STATE_H diff --git a/dpnp/backend/src/dpnp_utils.hpp b/dpnp/backend/src/dpnp_utils.hpp deleted file mode 100644 index 999347b58d08..000000000000 --- a/dpnp/backend/src/dpnp_utils.hpp +++ /dev/null @@ -1,147 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// - Neither the name of the copyright holder nor the names of its contributors -// may be used to endorse or promote products derived from this software -// without specific prior written permission. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -#pragma once -#ifndef BACKEND_UTILS_H // Cython compatibility -#define BACKEND_UTILS_H - -#include -#include -#include -#include -#include -#include - -#include - -#include - -/** - * @defgroup BACKEND_UTILS Backend C++ library utilities - * @{ - * This section describes utilities used in Backend API. - * @} - */ - -/** - * @ingroup BACKEND_UTILS - * @brief check support of type T by SYCL device. - * - * To check if sycl::device may use templated type T. - * - * @param [in] q sycl::device which is examined for type support. - * - * @exception std::runtime_error type T is out of support by the queue. - */ -template -static inline void validate_type_for_device(const sycl::device &d) -{ - if constexpr (std::is_same_v) { - if (!d.has(sycl::aspect::fp64)) { - throw std::runtime_error("Device " + - d.get_info() + - " does not support type 'double'"); - } - } - else if constexpr (std::is_same_v>) { - if (!d.has(sycl::aspect::fp64)) { - throw std::runtime_error( - "Device " + d.get_info() + - " does not support type 'complex'"); - } - } - else if constexpr (std::is_same_v) { - if (!d.has(sycl::aspect::fp16)) { - throw std::runtime_error("Device " + - d.get_info() + - " does not support type 'half'"); - } - } -} - -/** - * @ingroup BACKEND_UTILS - * @brief check support of type T by SYCL queue. - * - * To check if sycl::queue assigned to a device may use templated type T. - * - * @param [in] q sycl::queue which is examined for type support. - * - * @exception std::runtime_error type T is out of support by the queue. - */ -template -static inline void validate_type_for_device(const sycl::queue &q) -{ - validate_type_for_device(q.get_device()); -} - -/** - * @ingroup BACKEND_UTILS - * @brief print std::vector to std::ostream. - * - * To print std::vector with POD types to std::out. - * - * @param [in] vec std::vector with types with ability to be printed. - */ -template -std::ostream &operator<<(std::ostream &out, const std::vector &vec) -{ - std::string delimiter; - out << "{"; - // std::copy(vec.begin(), vec.end(), std::ostream_iterator(out, ", ")); - // out << "\b\b}"; // last two 'backspaces' needs to eliminate last - // delimiter. ex: {2, 3, 4, } - for (auto &elem : vec) { - out << delimiter << elem; - if (delimiter.empty()) { - delimiter.assign(", "); - } - } - out << "}"; - - return out; -} - -/** - * @ingroup BACKEND_UTILS - * @brief print @ref DPNPFuncType to std::ostream. - * - * To print DPNPFuncType type to std::out. - * TODO implement string representation of the enum - * - * @param [in] elem DPNPFuncType value to be printed. - */ -template -std::ostream &operator<<(std::ostream &out, DPNPFuncType elem) -{ - out << static_cast(elem); - - return out; -} - -#endif // BACKEND_UTILS_H diff --git a/dpnp/backend/src/dpnpc_memory_adapter.hpp b/dpnp/backend/src/dpnpc_memory_adapter.hpp deleted file mode 100644 index cae39cfe7f08..000000000000 --- a/dpnp/backend/src/dpnpc_memory_adapter.hpp +++ /dev/null @@ -1,180 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// - Neither the name of the copyright holder nor the names of its contributors -// may be used to endorse or promote products derived from this software -// without specific prior written permission. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -#pragma once -#ifndef DPNP_MEMORY_ADAPTER_H // Cython compatibility -#define DPNP_MEMORY_ADAPTER_H - -#include "dpnp_utils.hpp" -#include "queue_sycl.hpp" - -/** - * @ingroup BACKEND_UTILS - * @brief Adapter for the memory given by parameters in the DPNPC functions - * - * This type should be used to accommodate memory in the function. For example, - * if the kernel must be executed in "queue_1" which is host based, but - * input arrays are located on other "queue_2" or unknown place. - * - * Also, some functions completely host based and has no SYCL environment. - * - */ -template -class DPNPC_ptr_adapter final -{ - DPCTLSyclQueueRef queue_ref; /**< reference to SYCL queue */ - sycl::queue queue; /**< SYCL queue */ - void *aux_ptr = nullptr; /**< pointer to allocated memory by this adapter */ - void *orig_ptr = - nullptr; /**< original pointer to memory given by parameters */ - size_t size_in_bytes = 0; /**< size of bytes of the memory */ - bool allocated = false; /**< True if the memory allocated by this procedure - and needs to be free */ - bool target_no_queue = false; /**< Indicates that original memory will be - accessed from non SYCL environment */ - bool copy_back = false; /**< If the memory is 'result' it needs to be copied - back to original */ - const bool verbose = false; - std::vector deps; - -public: - DPNPC_ptr_adapter() = delete; - - DPNPC_ptr_adapter(DPCTLSyclQueueRef q_ref, - const void *src_ptr, - const size_t size, - bool target_no_sycl = false, - bool copy_back_request = false) - { - queue_ref = q_ref; - queue = *(reinterpret_cast(queue_ref)); - target_no_queue = target_no_sycl; - copy_back = copy_back_request; - orig_ptr = const_cast(src_ptr); - size_in_bytes = size * sizeof(_DataType); - deps = std::vector{}; - - // enum class alloc { host = 0, device = 1, shared = 2, unknown = 3 }; - sycl::usm::alloc src_ptr_type = sycl::usm::alloc::unknown; - src_ptr_type = sycl::get_pointer_type(src_ptr, queue.get_context()); - if (verbose) { - std::cerr << "DPNPC_ptr_converter:"; - std::cerr << "\n\t target_no_queue=" << target_no_queue; - std::cerr << "\n\t copy_back=" << copy_back; - std::cerr << "\n\t pointer=" << src_ptr; - std::cerr << "\n\t size=" << size; - std::cerr << "\n\t size_in_bytes=" << size_in_bytes; - std::cerr << "\n\t pointer type=" << (long)src_ptr_type; - std::cerr << "\n\t queue inorder=" << queue.is_in_order(); - std::cerr << "\n\t queue device is_cpu=" - << queue.get_device().is_cpu(); - std::cerr << "\n\t queue device is_gpu=" - << queue.get_device().is_gpu(); - std::cerr << "\n\t queue device is_accelerator=" - << queue.get_device().is_accelerator(); - std::cerr << std::endl; - } - - if (is_memcpy_required(src_ptr_type)) { - aux_ptr = dpnp_memory_alloc_c(queue_ref, size_in_bytes); - dpnp_memory_memcpy_c(queue_ref, aux_ptr, src_ptr, size_in_bytes); - allocated = true; - if (verbose) { - std::cerr << "DPNPC_ptr_converter::alloc and copy memory" - << " from=" << src_ptr << " to=" << aux_ptr - << " size_in_bytes=" << size_in_bytes << std::endl; - } - } - else { - aux_ptr = const_cast(src_ptr); - } - } - - ~DPNPC_ptr_adapter() - { - if (allocated) { - if (verbose) { - std::cerr << "DPNPC_ptr_converter::free_memory at=" << aux_ptr - << std::endl; - } - - sycl::event::wait(deps); - - if (copy_back) { - copy_data_back(); - } - - dpnp_memory_free_c(queue_ref, aux_ptr); - } - } - - bool is_memcpy_required(sycl::usm::alloc src_ptr_type) - { - if (target_no_queue || queue.get_device().is_gpu()) { - if (src_ptr_type == sycl::usm::alloc::unknown) { - return true; - } - else if (target_no_queue && - src_ptr_type == sycl::usm::alloc::device) { - return true; - } - } - - return false; - } - - _DataType *get_ptr() const - { - return reinterpret_cast<_DataType *>(aux_ptr); - } - - void copy_data_back() const - { - if (verbose) { - std::cerr << "DPNPC_ptr_converter::copy_data_back:" - << " from=" << aux_ptr << " to=" << orig_ptr - << " size_in_bytes=" << size_in_bytes << std::endl; - } - - dpnp_memory_memcpy_c(queue_ref, orig_ptr, aux_ptr, size_in_bytes); - } - - void depends_on(const std::vector &new_deps) - { - assert(allocated); - deps.insert(std::end(deps), std::begin(new_deps), std::end(new_deps)); - } - - void depends_on(const sycl::event &new_dep) - { - assert(allocated); - deps.push_back(new_dep); - } -}; - -#endif // DPNP_MEMORY_ADAPTER_H diff --git a/dpnp/backend/src/memory_sycl.cpp b/dpnp/backend/src/memory_sycl.cpp deleted file mode 100644 index 9ca71485f35d..000000000000 --- a/dpnp/backend/src/memory_sycl.cpp +++ /dev/null @@ -1,134 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// - Neither the name of the copyright holder nor the names of its contributors -// may be used to endorse or promote products derived from this software -// without specific prior written permission. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -#include -#include -#include - -#include "dpnp_iface.hpp" -#include "queue_sycl.hpp" - -static bool use_sycl_device_memory() -{ - // TODO need to move all getenv() into common dpnpc place - char *dpnpc_memtype_device = nullptr; - -#ifdef _WIN32 - size_t dpnpc_memtype_device_size = 0; - _dupenv_s(&dpnpc_memtype_device, &dpnpc_memtype_device_size, - "DPNPC_OUTPUT_DPARRAY_USE_MEMORY_DEVICE"); -#else - dpnpc_memtype_device = - std::getenv("DPNPC_OUTPUT_DPARRAY_USE_MEMORY_DEVICE"); -#endif - - if (dpnpc_memtype_device != nullptr) { -#ifdef _WIN32 - free(dpnpc_memtype_device); -#endif - return true; - } - - return false; -} - -// This variable is needed for the NumPy corner case -// if we have zero memory array (ex. shape=(0,10)) we must keep the pointer to -// somewhere memory of this variable must not be used -const char *numpy_stub = "...the NumPy Stub..."; - -char *dpnp_memory_alloc_c(DPCTLSyclQueueRef q_ref, size_t size_in_bytes) -{ - char *array = const_cast(numpy_stub); - - // std::cout << "dpnp_memory_alloc_c(size=" << size_in_bytes << std::flush; - if (size_in_bytes > 0) { - sycl::queue q = *(reinterpret_cast(q_ref)); - sycl::usm::alloc memory_type = sycl::usm::alloc::shared; - if (use_sycl_device_memory()) { - memory_type = sycl::usm::alloc::device; - } - array = reinterpret_cast( - sycl::malloc(size_in_bytes, q, memory_type)); - if (array == nullptr) { - // TODO add information about number of allocated bytes - throw std::runtime_error( - "DPNP Error: dpnp_memory_alloc_c() out of memory."); - } - -#if not defined(NDEBUG) - if (memory_type != sycl::usm::alloc::device) { - for (size_t i = 0; i < size_in_bytes / sizeof(char); ++i) { - array[i] = 0; // type dependent is better. set double(42.42) - // instead zero - } - } - // std::cout << ") -> ptr=" << (void*)array << std::endl; -#endif - } - - return array; -} - -char *dpnp_memory_alloc_c(size_t size_in_bytes) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - return dpnp_memory_alloc_c(q_ref, size_in_bytes); -} - -void dpnp_memory_free_c(DPCTLSyclQueueRef q_ref, void *ptr) -{ - // std::cout << "dpnp_memory_free_c(ptr=" << (void*)ptr << ")" << std::endl; - sycl::queue q = *(reinterpret_cast(q_ref)); - if (ptr != numpy_stub) { - sycl::free(ptr, q); - } -} - -void dpnp_memory_free_c(void *ptr) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - dpnp_memory_free_c(q_ref, ptr); -} - -void dpnp_memory_memcpy_c(DPCTLSyclQueueRef q_ref, - void *dst, - const void *src, - size_t size_in_bytes) -{ - // std::cout << "dpnp_memory_memcpy_c(dst=" << dst << ", src=" << src << ")" - // << std::endl; - sycl::queue q = *(reinterpret_cast(q_ref)); - q.memcpy(dst, src, size_in_bytes).wait(); -} - -void dpnp_memory_memcpy_c(void *dst, const void *src, size_t size_in_bytes) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - dpnp_memory_memcpy_c(q_ref, dst, src, size_in_bytes); -} diff --git a/dpnp/backend/src/queue_sycl.cpp b/dpnp/backend/src/queue_sycl.cpp deleted file mode 100644 index 715193959cf1..000000000000 --- a/dpnp/backend/src/queue_sycl.cpp +++ /dev/null @@ -1,111 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// - Neither the name of the copyright holder nor the names of its contributors -// may be used to endorse or promote products derived from this software -// without specific prior written permission. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -#include -#include - -#include "dpnp_iface.hpp" -#include "dpnp_utils.hpp" -#include "queue_sycl.hpp" - -[[maybe_unused]] static void dpnpc_show_mathlib_version() -{ -#if 1 - const int len = 256; - std::string mathlib_version_str(len, 0x0); - - char *buf = const_cast( - mathlib_version_str.c_str()); // TODO avoid write into the container - - mkl_get_version_string(buf, len); - - std::cout << "Math backend version: " << mathlib_version_str << std::endl; -#else - // Failed to load library under Python environment die to unresolved symbol - MKLVersion version; - - mkl_get_version(&version); - - std::cout << "Math backend version: " << version.MajorVersion << "." - << version.UpdateVersion << "." << version.MinorVersion - << std::endl; -#endif -} - -#if (not defined(NDEBUG)) -[[maybe_unused]] static std::string - device_type_to_str(sycl::info::device_type devTy) -{ - std::stringstream ss; - switch (devTy) { - case sycl::info::device_type::cpu: - ss << "cpu"; - break; - case sycl::info::device_type::gpu: - ss << "gpu"; - break; - case sycl::info::device_type::accelerator: - ss << "accelerator"; - break; - case sycl::info::device_type::custom: - ss << "custom"; - break; - default: - ss << "unknown"; - } - return ss.str(); -} - -[[maybe_unused]] static void show_available_sycl_devices() -{ - const std::vector devices = sycl::device::get_devices(); - - std::cout << "Available SYCL devices:" << std::endl; - for (std::vector::const_iterator it = devices.cbegin(); - it != devices.cend(); ++it) { - std::cout - // not yet implemented error << " " << - // it->has(sycl::aspect::usm_shared_allocations) << " " - << " - id=" << it->get_info() - << ", type=" - << device_type_to_str( - it->get_info()) - << ", gws=" - << it->get_info() - << ", cu=" << it->get_info() - << ", name=" << it->get_info() - << std::endl; - } -} -#endif - -size_t dpnp_queue_is_cpu_c() -{ - const auto &be = backend_sycl::get(); - return be.backend_sycl_is_cpu(); -} diff --git a/dpnp/backend/src/queue_sycl.hpp b/dpnp/backend/src/queue_sycl.hpp deleted file mode 100644 index 3c41401ad62f..000000000000 --- a/dpnp/backend/src/queue_sycl.hpp +++ /dev/null @@ -1,141 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// - Neither the name of the copyright holder nor the names of its contributors -// may be used to endorse or promote products derived from this software -// without specific prior written permission. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -#pragma once -#ifndef QUEUE_SYCL_H // Cython compatibility -#define QUEUE_SYCL_H - -// #pragma clang diagnostic push -// #pragma clang diagnostic ignored "-Wpass-failed" -#include -// #pragma clang diagnostic pop - -#pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wunused-parameter" -#pragma clang diagnostic ignored "-Wreorder-ctor" -#include -#pragma clang diagnostic pop - -#include - -#include "dpnp_pstl.hpp" // this header must be included after - -#include "verbose.hpp" - -namespace mkl_rng = oneapi::mkl::rng; - -#define DPNP_QUEUE backend_sycl::get_queue() -#define DPNP_RNG_ENGINE backend_sycl::get_rng_engine() -#define DPNP_RNG_MCG59_ENGINE backend_sycl::get_rng_mcg59_engine() - -/** - * This is container for the SYCL queue, random number generation engine and - * related functions like queue and engine initialization and maintenance. The - * queue could not be initialized as a global object. Global object - * initialization order is undefined. This class postpone initialization of the - * SYCL queue and mt19937 random number generation engine. - */ -class backend_sycl -{ -public: - ~backend_sycl() {} - - static backend_sycl &get() - { -#if defined(_WIN32) && INTEL_MKL_VERSION == 20260000 - // TODO: remove once MKLD-19835 is resolved - // mt19937 (oneMKL 2026.0) destructor crashes during DLL_PROCESS_DETACH - // on Windows (Battlemage/Level Zero). Use a heap-allocated - // process-lifetime singleton to skip destructor; OS reclaims memory. - static backend_sycl *backend = new backend_sycl{}; - return *backend; -#else - static backend_sycl backend{}; - return backend; -#endif - } - - static sycl::queue &get_queue() - { - auto &be = backend_sycl::get(); - return be.queue_; - } - - static mkl_rng::mt19937 &get_rng_engine() - { - auto &be = backend_sycl::get(); - return be.rng_mt19937_engine_; - } - - static mkl_rng::mcg59 &get_rng_mcg59_engine() - { - auto &be = backend_sycl::get(); - return be.rng_mcg59_engine_; - } - - template - void set_rng_engines_seed(const SeedT &seed) - { - mkl_rng::mt19937 rng_eng_mt19937(queue_, seed); - mkl_rng::mcg59 rng_eng_mcg59(queue_, seed); - - // now that instances are created, let's move them - rng_mt19937_engine_ = std::move(rng_eng_mt19937); - rng_mcg59_engine_ = std::move(rng_eng_mcg59); - } - - bool backend_sycl_is_cpu() const - { - const auto &dev = queue_.get_device(); - return dev.is_cpu(); - } - -private: - static constexpr std::size_t default_seed = 1; - - backend_sycl() - : queue_{sycl::default_selector_v, - (is_verbose_mode()) - ? sycl::property_list{sycl::property::queue:: - enable_profiling()} - : sycl::property_list{}}, - rng_mt19937_engine_{queue_, default_seed}, - rng_mcg59_engine_{queue_, default_seed} - { - } - - backend_sycl(backend_sycl const &) = default; - backend_sycl &operator=(backend_sycl const &) = default; - backend_sycl &operator=(backend_sycl &&) = default; - - sycl::queue queue_; - mkl_rng::mt19937 rng_mt19937_engine_; - mkl_rng::mcg59 rng_mcg59_engine_; -}; - -#endif // QUEUE_SYCL_H diff --git a/dpnp/backend/tests/CMakeLists.txt b/dpnp/backend/tests/CMakeLists.txt deleted file mode 100644 index 2f5641d1ed92..000000000000 --- a/dpnp/backend/tests/CMakeLists.txt +++ /dev/null @@ -1,61 +0,0 @@ -# -*- coding: utf-8 -*- -# ***************************************************************************** -# Copyright (c) 2016, Intel Corporation -# All rights reserved. -# -# Redistribution and use in source and binary forms, with or without -# modification, are permitted provided that the following conditions are met: -# - Redistributions of source code must retain the above copyright notice, -# this list of conditions and the following disclaimer. -# - Redistributions in binary form must reproduce the above copyright notice, -# this list of conditions and the following disclaimer in the documentation -# and/or other materials provided with the distribution. -# - Neither the name of the copyright holder nor the names of its contributors -# may be used to endorse or promote products derived from this software -# without specific prior written permission. -# -# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -# ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -# LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -# CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -# SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -# INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -# CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -# ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -# THE POSSIBILITY OF SUCH DAMAGE. -# ***************************************************************************** - -find_package(GTest REQUIRED) -find_package(Threads REQUIRED) - -if(GTEST_FOUND) - message( - STATUS - "Found GTest: (include: ${GTEST_INCLUDE_DIRS}, library: ${GTEST_LIBRARIES})" - ) -else() - message(FATAL_ERROR "Cannot find Google Test Framework!") -endif() - -if(NOT Threads_FOUND) - message(FATAL_ERROR "Cannot find Threads library!") -endif() - -# Emulate autotools like make check target to build tests -set(CMAKE_CTEST_COMMAND ctest --progress --output-on-failure) -add_custom_target(check COMMAND ${CMAKE_CTEST_COMMAND}) -enable_testing() - -include_directories(${GTEST_INCLUDE_DIR}) -link_directories(${GTEST_LIB_DIR}) - -# TODO split -add_executable(dpnpc_tests test_main.cpp test_random.cpp) -target_link_libraries(dpnpc_tests GTest::GTest GTest::Main pthread dpnp_backend_library) - -# TODO split -add_test(dpnpc_tests dpnpc_tests) -gtest_discover_tests(dpnpc_tests) -add_dependencies(check dpnpc_tests) diff --git a/dpnp/backend/tests/test_random.cpp b/dpnp/backend/tests/test_random.cpp deleted file mode 100644 index bda2658e1de4..000000000000 --- a/dpnp/backend/tests/test_random.cpp +++ /dev/null @@ -1,221 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// - Neither the name of the copyright holder nor the names of its contributors -// may be used to endorse or promote products derived from this software -// without specific prior written permission. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -#include -#include - -#include - -#include -#include - -#include "gtest/gtest.h" - -// TODO -// * data management will be redesigned: allocation as chars (and casting on -// teste suits) -// * this class will be generalized -class RandomTestCase : public ::testing::Test -{ -public: - static void SetUpTestCase() { _get_device_mem(); } - - static void TearDownTestCase() - { - dpnp_memory_free_c(result1); - result1 = result2 = nullptr; - } - - void SetUp() override - { - for (size_t i = 0; i < size; ++i) { - result1[i] = 1; - result2[i] = 0; - } - } - - void TearDown() override {} - -private: - static void _get_device_mem() - { - result1 = (double *)dpnp_memory_alloc_c(size * 2 * sizeof(double)); - result2 = result1 + size; - } - -protected: - static size_t size; - static double *result1; - static double *result2; -}; - -double *RandomTestCase::result1 = nullptr; -double *RandomTestCase::result2 = nullptr; -size_t RandomTestCase::size = 10; - -template -bool check_statistics(_DataType *r, - double tM, - double tD, - double tQ, - size_t size) -{ - double tD2; - double sM, sD; - double sum, sum2; - double n, s; - double DeltaM, DeltaD; - - /***** Sample moments *****/ - sum = 0.0; - sum2 = 0.0; - for (size_t i = 0; i < size; i++) { - sum += (double)r[i]; - sum2 += (double)r[i] * (double)r[i]; - } - sM = sum / ((double)size); - sD = sum2 / (double)size - (sM * sM); - - /***** Comparison of theoretical and sample moments *****/ - n = (double)size; - tD2 = tD * tD; - s = ((tQ - tD2) / n) - (2 * (tQ - 2 * tD2) / (n * n)) + - ((tQ - 3 * tD2) / (n * n * n)); - - DeltaM = (tM - sM) / sqrt(tD / n); - DeltaD = (tD - sD) / sqrt(s); - if (fabs(DeltaM) > 3.0 || fabs(DeltaD) > 3.0) - return false; - else - return true; -} - -TEST_F(RandomTestCase, rng_beta_test_seed) -{ - const size_t seed = 10; - const double a = 0.4; - const double b = 0.5; - - dpnp_rng_srand_c(seed); - dpnp_rng_beta_c(result1, a, b, size); - - dpnp_rng_srand_c(seed); - dpnp_rng_beta_c(result2, a, b, size); - - for (size_t i = 0; i < size; ++i) { - EXPECT_NEAR(result1[i], result2[i], 0.004); - } -} - -TEST_F(RandomTestCase, rng_f_test_seed) -{ - const size_t seed = 10; - const double dfnum = 10.4; - const double dfden = 12.5; - - dpnp_rng_srand_c(seed); - dpnp_rng_f_c(result1, dfnum, dfden, size); - - dpnp_rng_srand_c(seed); - dpnp_rng_f_c(result2, dfnum, dfden, size); - - for (size_t i = 0; i < size; ++i) { - EXPECT_NEAR(result1[i], result2[i], 0.004); - } -} - -TEST_F(RandomTestCase, rng_normal_test_seed) -{ - const size_t seed = 10; - const double loc = 2.56; - const double scale = 0.8; - - dpnp_rng_srand_c(seed); - dpnp_rng_normal_c(result1, loc, scale, size); - - dpnp_rng_srand_c(seed); - dpnp_rng_normal_c(result2, loc, scale, size); - - for (size_t i = 0; i < size; ++i) { - EXPECT_NEAR(result1[i], result2[i], 0.004); - } -} - -TEST_F(RandomTestCase, rng_uniform_test_seed) -{ - const size_t seed = 10; - const long low = 1; - const long high = 120; - - dpnp_rng_srand_c(seed); - dpnp_rng_uniform_c(result1, low, high, size); - - dpnp_rng_srand_c(seed); - dpnp_rng_uniform_c(result2, low, high, size); - - for (size_t i = 0; i < size; ++i) { - EXPECT_NEAR(result1[i], result2[i], 0.004); - } -} - -TEST(TestBackendRandomUniform, test_statistics) -{ - const size_t size = 256; - size_t seed = 10; - long a = 1; - long b = 120; - bool check_statistics_res = false; - - /***** Theoretical moments *****/ - double tM = (b + a) / 2.0; - double tD = ((b - a) * (b - a)) / 12.0; - double tQ = ((b - a) * (b - a) * (b - a) * (b - a)) / 80.0; - - double *result = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - dpnp_rng_srand_c(seed); - dpnp_rng_uniform_c(result, a, b, size); - check_statistics_res = check_statistics(result, tM, tD, tQ, size); - ASSERT_TRUE(check_statistics_res); - dpnp_memory_free_c(result); -} - -// TODO: -// Generalization for all DPNPFuncName -TEST(TestBackendRandomSrand, test_func_ptr) -{ - void *fptr = nullptr; - DPNPFuncData kernel_data = get_dpnp_function_ptr( - DPNPFuncName::DPNP_FN_RNG_SRAND, DPNPFuncType::DPNP_FT_DOUBLE, - DPNPFuncType::DPNP_FT_DOUBLE); - - fptr = get_dpnp_function_ptr1( - kernel_data.return_type, DPNPFuncName::DPNP_FN_RNG_SRAND, - DPNPFuncType::DPNP_FT_DOUBLE, DPNPFuncType::DPNP_FT_DOUBLE); - - EXPECT_TRUE(fptr != nullptr); -} diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index 2ac2a0a04531..97b5311042f3 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -33,43 +33,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": cdef enum DPNPFuncName "DPNPFuncName": DPNP_FN_PARTITION_EXT - DPNP_FN_RNG_BETA_EXT - DPNP_FN_RNG_BINOMIAL_EXT - DPNP_FN_RNG_CHISQUARE_EXT - DPNP_FN_RNG_EXPONENTIAL_EXT - DPNP_FN_RNG_F_EXT - DPNP_FN_RNG_GAMMA_EXT - DPNP_FN_RNG_GAUSSIAN_EXT - DPNP_FN_RNG_GEOMETRIC_EXT - DPNP_FN_RNG_GUMBEL_EXT - DPNP_FN_RNG_HYPERGEOMETRIC_EXT - DPNP_FN_RNG_LAPLACE_EXT - DPNP_FN_RNG_LOGISTIC_EXT - DPNP_FN_RNG_LOGNORMAL_EXT - DPNP_FN_RNG_MULTINOMIAL_EXT - DPNP_FN_RNG_MULTIVARIATE_NORMAL - DPNP_FN_RNG_MULTIVARIATE_NORMAL_EXT - DPNP_FN_RNG_NEGATIVE_BINOMIAL_EXT - DPNP_FN_RNG_NONCENTRAL_CHISQUARE_EXT - DPNP_FN_RNG_NORMAL_EXT - DPNP_FN_RNG_PARETO_EXT - DPNP_FN_RNG_POISSON_EXT - DPNP_FN_RNG_POWER_EXT - DPNP_FN_RNG_RAYLEIGH_EXT - DPNP_FN_RNG_SHUFFLE_EXT - DPNP_FN_RNG_SRAND - DPNP_FN_RNG_SRAND_EXT - DPNP_FN_RNG_STANDARD_CAUCHY_EXT - DPNP_FN_RNG_STANDARD_EXPONENTIAL_EXT - DPNP_FN_RNG_STANDARD_GAMMA_EXT - DPNP_FN_RNG_STANDARD_NORMAL - DPNP_FN_RNG_STANDARD_T_EXT - DPNP_FN_RNG_TRIANGULAR_EXT - DPNP_FN_RNG_UNIFORM_EXT - DPNP_FN_RNG_VONMISES_EXT - DPNP_FN_RNG_WALD_EXT - DPNP_FN_RNG_WEIBULL_EXT - DPNP_FN_RNG_ZIPF_EXT # need this namespace for Enum import cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncType": diff --git a/dpnp/dpnp_iface.py b/dpnp/dpnp_iface.py index e3f36711c905..37d2e657bc9c 100644 --- a/dpnp/dpnp_iface.py +++ b/dpnp/dpnp_iface.py @@ -405,11 +405,9 @@ def get_dpnp_descriptor( if ext_obj.strides != shape_offsets or ext_obj_offset != 0: ext_obj = dpnp.array(ext_obj, order="C") - # while dpnp functions are based on DPNP_QUEUE - # we need to create a copy on device associated with DPNP_QUEUE - # if function get implementation for different queue - # then this behavior can be disabled with setting - # "copy_when_nondefault_queue" + # legacy backend functions ran on the default SYCL queue, so the input + # is copied to that queue; a function implemented for any queue can + # disable this by setting "copy_when_nondefault_queue" queue = getattr(ext_obj, "sycl_queue", None) if queue is not None and copy_when_nondefault_queue: default_queue = dpctl.SyclQueue() diff --git a/dpnp/dpnp_utils/dpnp_algo_utils.pxd b/dpnp/dpnp_utils/dpnp_algo_utils.pxd index 7b99aec3dd32..efec5ade978c 100644 --- a/dpnp/dpnp_utils/dpnp_algo_utils.pxd +++ b/dpnp/dpnp_utils/dpnp_algo_utils.pxd @@ -54,9 +54,6 @@ cpdef checker_throw_type_error(function_name, given_type) cpdef cpp_bool use_origin_backend(input1=*, size_t compute_size=*) -cpdef tuple _object_to_tuple(object obj) - - cdef class dpnp_descriptor: """array DPNP descriptor""" @@ -75,6 +72,3 @@ cdef dpnp_descriptor create_output_descriptor(shape_type_c output_shape, device=*, usm_type=*, sycl_queue=*) - -# Get common USM allocation in the form of (sycl_device, usm_type, sycl_queue) -cdef tuple get_common_usm_allocation(dpnp_descriptor x1, dpnp_descriptor x2) diff --git a/dpnp/dpnp_utils/dpnp_algo_utils.pyx b/dpnp/dpnp_utils/dpnp_algo_utils.pyx index 5a3fdd397501..390365bb2e3d 100644 --- a/dpnp/dpnp_utils/dpnp_algo_utils.pyx +++ b/dpnp/dpnp_utils/dpnp_algo_utils.pyx @@ -37,7 +37,6 @@ This module contains different helpers and utilities """ import dpctl -import numpy import dpnp import dpnp.config as config @@ -45,9 +44,6 @@ import dpnp.dpnp_container as dpnp_container from dpnp.dpnp_array import dpnp_array from dpnp.tensor import get_coerced_usm_type, get_execution_queue -cimport cpython -cimport cython -cimport numpy from libcpp cimport bool as cpp_bool from dpnp.dpnp_algo.dpnp_algo cimport ( @@ -56,15 +52,12 @@ from dpnp.dpnp_algo.dpnp_algo cimport ( # Python import functions __all__ = [ - "call_origin", "checker_throw_type_error", "checker_throw_value_error", "convert_item", "dpnp_descriptor", "get_usm_allocations", "map_dtype_to_device", - "_object_to_tuple", - "unwrap_array", "use_origin_backend" ] @@ -100,159 +93,6 @@ def convert_list_args(input_list): return result_list -def copy_from_origin(dst, src): - """Copy origin result to output result.""" - if hasattr(dst, "__sycl_usm_array_interface__"): - if src.size: - dst_dpt = unwrap_array(dst) - dst_dpt[...] = src - else: - for i in range(dst.size): - dst.flat[i] = src.item(i) - - -def call_origin(function, *args, **kwargs): - """ - Call fallback function for unsupported cases - """ - - allow_fallback = kwargs.pop("allow_fallback", False) - - if ( - not allow_fallback - and config.__DPNP_RAISE_EXCEPION_ON_NUMPY_FALLBACK__ == 1 - ): - raise NotImplementedError( - f"Requested function={function.__name__} " - f"with args={args} and kwargs={kwargs} " - "isn't currently supported and would fall " - "back on NumPy implementation. " - "Define environment variable " - "`DPNP_RAISE_EXCEPION_ON_NUMPY_FALLBACK` " - "to `0` if the fall back is required to be " - "supported without raising an exception." - ) - - dpnp_inplace = kwargs.pop("dpnp_inplace", False) - sycl_queue = kwargs.pop("sycl_queue", None) - # print(f"DPNP call_origin(): Fallback called. " - # f"\n\t function={function}, " - # f"\n\t args={args}, " - # f"\n\t kwargs={kwargs}, " - # f"\n\t dpnp_inplace={dpnp_inplace}") - - kwargs_out = kwargs.get("out", None) - alloc_queues = [sycl_queue] if sycl_queue else [] - if (kwargs_out is not None): - if isinstance(kwargs_out, numpy.ndarray): - kwargs["out"] = kwargs_out - else: - if hasattr(kwargs_out, "sycl_queue"): - alloc_queues.append(kwargs_out.sycl_queue) - kwargs["out"] = dpnp.asnumpy(kwargs_out) - - args_new_list = [] - for arg in args: - if hasattr(arg, "sycl_queue"): - alloc_queues.append(arg.sycl_queue) - argx = convert_item(arg) - args_new_list.append(argx) - args_new = tuple(args_new_list) - - kwargs_new = {} - for key, kwarg in kwargs.items(): - if hasattr(kwarg, "sycl_queue"): - alloc_queues.append(kwarg.sycl_queue) - kwargx = convert_item(kwarg) - kwargs_new[key] = kwargx - - exec_q = get_execution_queue(alloc_queues) - if exec_q is None: - exec_q = dpnp.get_normalized_queue_device( - sycl_queue=sycl_queue - ) - # print(f"DPNP call_origin(): backend called. " - # f"\n\t function={function}, " - # f"\n\t args_new={args_new}, " - # f"\n\t kwargs_new={kwargs_new}, " - # f"\n\t dpnp_inplace={dpnp_inplace}") - # TODO need to put array memory into NumPy call - result_origin = function(*args_new, **kwargs_new) - # print(f"DPNP call_origin(): result from backend. " - # f"\n\t result_origin={result_origin}, " - # f"\n\t args_new={args_new}, " - # f"\n\t kwargs_new={kwargs_new}, " - # f"\n\t dpnp_inplace={dpnp_inplace}") - result = result_origin - if dpnp_inplace: - # enough to modify only first argument in place - if args and args_new: - arg, arg_new = args[0], args_new[0] - if isinstance(arg_new, numpy.ndarray): - copy_from_origin(arg, arg_new) - elif isinstance(arg_new, list): - for i, val in enumerate(arg_new): - arg[i] = val - - elif isinstance(result, numpy.ndarray): - if kwargs_out is None: - # use dtype from input arguments if present - # or from the result otherwise - result_dtype = ( - kwargs.get("dtype", None) or result_origin.dtype - ) - - if exec_q is not None: - result_dtype = map_dtype_to_device( - result_origin.dtype, exec_q.sycl_device - ) - - result = dpnp_container.empty( - result_origin.shape, - dtype=result_dtype, - sycl_queue=exec_q, - ) - else: - result = kwargs_out - - copy_from_origin(result, result_origin) - - elif isinstance(result, tuple): - # convert tuple(fallback_array) to tuple(result_array) - result_list = [] - for res_origin in result: - res = res_origin - if isinstance(res_origin, numpy.ndarray): - if exec_q is not None: - result_dtype = map_dtype_to_device( - res_origin.dtype, - exec_q.sycl_device, - ) - else: - result_dtype = res_origin.d_type - res = dpnp_container.empty( - res_origin.shape, - dtype=result_dtype, - sycl_queue=exec_q, - ) - copy_from_origin(res, res_origin) - result_list.append(res) - - result = tuple(result_list) - - return result - - -def unwrap_array(x1): - """ - Get array from input object. - """ - if isinstance(x1, dpnp_array): - return x1.get_array() - - return x1 - - def _get_coerced_usm_type(objects): types_in_use = [ obj.usm_type for obj in objects @@ -428,53 +268,6 @@ cdef dpnp_descriptor create_output_descriptor( return result_desc -@cython.profile(False) -cpdef inline tuple _object_to_tuple(object obj): - """ Converts Python object into tuple - - """ - - if obj is None: - return () - - # dpnp.ndarray unconditionally succeeds in - # PySequence_Check as it implements __getitem__ - if ( - cpython.PySequence_Check(obj) - and not dpnp.is_supported_array_type(obj) - ): - if isinstance(obj, numpy.ndarray): - obj = numpy.atleast_1d(obj) - - nd = len(obj) - shape = [] - - for i in range(0, nd): - if cpython.PyBool_Check(obj[i]): - raise TypeError( - "DPNP object_to_tuple(): " - "no item in size can be bool" - ) - - # Assumes each item is castable to Py_ssize_t, - # otherwise TypeError will be raised - shape.append( obj[i]) - return tuple(shape) - - if dpnp.isscalar(obj): - if cpython.PyBool_Check(obj): - raise TypeError( - "DPNP object_to_tuple(): " - "'obj' can't be bool" - ) - return (obj, ) - - raise ValueError( - "DPNP object_to_tuple(): 'obj' should be " - "'None', collections.abc.Sequence, or 'int'" - ) - - cpdef cpp_bool use_origin_backend( input1=None, size_t compute_size=0 ): @@ -499,43 +292,6 @@ cpdef cpp_bool use_origin_backend( return False -cdef tuple get_common_usm_allocation( - dpnp_descriptor x1, dpnp_descriptor x2 -): - """ - Get common USM allocation in the form of - (sycl_device, usm_type, sycl_queue). - """ - array1_obj = x1.get_array() - array2_obj = x2.get_array() - - common_usm_type = get_coerced_usm_type( - (array1_obj.usm_type, array2_obj.usm_type) - ) - if common_usm_type is None: - raise ValueError( - "could not recognize common USM type " - f"for inputs of USM types {array1_obj.usm_type} " - f"and {array2_obj.usm_type}" - ) - - common_sycl_queue = get_execution_queue( - (array1_obj.sycl_queue, array2_obj.sycl_queue) - ) - if common_sycl_queue is None: - raise ValueError( - "could not recognize common SYCL queue " - f"for inputs in SYCL queues {array1_obj.sycl_queue} " - f"and {array2_obj.sycl_queue}" - ) - - return ( - common_sycl_queue.sycl_device, - common_usm_type, - common_sycl_queue, - ) - - cdef class dpnp_descriptor: def __init__(self, obj): """ Initialize variables """ diff --git a/dpnp/libtensor/AGENTS.md b/dpnp/libtensor/AGENTS.md index c6cec65c2044..a0e98b23381b 100644 --- a/dpnp/libtensor/AGENTS.md +++ b/dpnp/libtensor/AGENTS.md @@ -3,7 +3,7 @@ ## Layout - `include/kernels/`: header-only SYCL kernels in namespaces under `dpnp::tensor::kernels`; `include/utils/`: type dispatch, type, offset and memory-overlap helpers (`type_dispatch.hpp`, `type_dispatch_building.hpp`, `type_utils.hpp`, `memory_overlap.hpp`, ...); `include/ext/`: dispatch helpers (`init_dispatch_vector`, `init_dispatch_table`, `DispatchTable`) and validation helpers (`validation_utils.hpp`). The headers ship with dpnp; `dpnp.get_include()` and `python -m dpnp --tensor-include-dir` return `include/`. - `source/`: bindings built by `source/CMakeLists.txt` into `_tensor_impl`, `_tensor_elementwise_impl`, `_tensor_reductions_impl`, `_tensor_sorting_impl`, `_tensor_accumulation_impl`, `_tensor_linalg_impl`, `_statistics_impl` and `_window_impl`. The functions dpnp adds beyond the Array API are listed in the `_dpnp_*_sources` variables there and compiled into the same modules. -- `extensions/{blas,fft,lapack,sparse,vm}/`: oneMKL/oneMath extensions, one module `__impl` per directory. +- `extensions/{blas,fft,lapack,rng,sparse,vm}/`: oneMKL/oneMath extensions, one module `__impl` per directory. - Every module installs into `dpnp/libtensor` and is imported as `dpnp.libtensor.__impl`. A new kernel goes into `source/` and an existing module; only code that needs oneMKL gets a new extension. ## Adding a ufunc Kernel (model: `fabs`) @@ -21,8 +21,9 @@ A new extension's `CMakeLists.txt` lists its sources and calls `build_dpnp_pybin - Validate inputs before launching: queue compatibility, contiguity, writability, memory overlap (including the same logical tensor), dtype and ndim. Reuse the helpers neighbouring files use: `check_queue`, `check_c_contig`, `check_no_overlap`, `check_writable` from `include/ext/validation_utils.hpp` (`putmask`, `statistics`, `isclose`, `interpolate`), or `dpnp::utils::queues_are_compatible`, `dpnp::tensor::overlap::MemoryOverlap` and `dpnp::tensor::validation::CheckWritable` (extensions and the Array API functions). Don't silently accept strided input. - Throw `py::value_error` for bad arguments and `std::runtime_error` for internal failures; catch oneMKL and SYCL exceptions and rethrow them. - Guard double, complex and half with `dpnp::tensor::type_utils::validate_type_for_device(q)`; don't use double literals or casts in kernels that must run without fp64. -- `blas`, `fft`, `lapack` and `sparse` build against either oneMKL or oneMath (`USE_ONEMATH`); `vm` has kernels only with oneMKL: with oneMath it builds empty and `_is_available()` returns `False`. +- `blas`, `fft`, `lapack` and `sparse` build against either oneMKL or oneMath (`USE_ONEMATH`); `vm` has kernels only with oneMKL: with oneMath it builds empty and `_is_available()` returns `False`; `rng` uses the header-only oneMKL device RNG API and links no oneMKL library. - Release the GIL around blocking oneMKL calls (`py::gil_scoped_release`). +- `rng` values depend only on the engine state and the number of elements, never on the device or the layout chosen for it (`extensions/rng/kernels/sample.hpp`); each engine's `items_per_work_item` (`kernels/engines.hpp`) is part of its stream definition. Parameters on the device are unchecked, so a rejection sampler returns NaN or the smallest `int64` for the inputs that would make it loop forever, including values that are valid but round badly in fp32. Use `umul_hi` from `kernels/umul_hi.hpp` instead of `sycl::mul_hi`, which DPC++ 2026.2 miscompiles on Intel GPUs. - An extension whose kernels may be launched over more than `INT_MAX` work-items passes `LARGE_RANGES` to `build_dpnp_pybind11_ext`, which adds `-fno-sycl-id-queries-fit-in-int` at compile and link time (DPC++ 2026.2 and later need it at link time). `source/CMakeLists.txt` and `dpnp_backend_c` pass it to both `target_compile_options` and `target_link_options`. - clang-format 23 formats C++ (`.clang-format`: LLVM base, 4-space indent, 80 columns); gersemi formats CMake at 88 columns, except `dpnp/backend/cmake/Modules/`. diff --git a/dpnp/random/CMakeLists.txt b/dpnp/libtensor/extensions/rng/CMakeLists.txt similarity index 65% rename from dpnp/random/CMakeLists.txt rename to dpnp/libtensor/extensions/rng/CMakeLists.txt index 5a4689d99a27..e2289b19fd9d 100644 --- a/dpnp/random/CMakeLists.txt +++ b/dpnp/libtensor/extensions/rng/CMakeLists.txt @@ -1,6 +1,6 @@ # -*- coding: utf-8 -*- # ***************************************************************************** -# Copyright (c) 2023, Intel Corporation +# Copyright (c) 2026, Intel Corporation # All rights reserved. # # Redistribution and use in source and binary forms, with or without @@ -27,10 +27,29 @@ # THE POSSIBILITY OF SUCH DAMAGE. # ***************************************************************************** -# Building dpnp_algo_random Cython extension +set(python_module_name _rng_impl) +set(_module_src + ${CMAKE_CURRENT_SOURCE_DIR}/binomial.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/discrete.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/exponential.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/gamma.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/integers.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/noncentral.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/normal.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/poisson.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/rng_py.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/sample.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/uniform.cpp +) + +build_dpnp_pybind11_ext( + ${python_module_name} + SOURCES ${_module_src} +) -build_dpnp_cython_ext_with_backend( - dpnp_algo_random - ${CMAKE_CURRENT_SOURCE_DIR}/dpnp_algo_random.pyx - dpnp/random +# oneMKL device RNG API is header-only, so no oneMKL library is linked +target_include_directories( + ${python_module_name} + SYSTEM + PRIVATE $ ) diff --git a/dpnp/backend/src/verbose.hpp b/dpnp/libtensor/extensions/rng/binomial.cpp similarity index 77% rename from dpnp/backend/src/verbose.hpp rename to dpnp/libtensor/extensions/rng/binomial.cpp index 4b3fa077cb73..58c8bce687ef 100644 --- a/dpnp/backend/src/verbose.hpp +++ b/dpnp/libtensor/extensions/rng/binomial.cpp @@ -1,5 +1,5 @@ //***************************************************************************** -// Copyright (c) 2016, Intel Corporation +// Copyright (c) 2026, Intel Corporation // All rights reserved. // // Redistribution and use in source and binary forms, with or without @@ -25,17 +25,20 @@ // ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include -#pragma once -#ifndef VERBOSE_H // Cython compatibility -#define VERBOSE_H +#include "common.hpp" +#include "distributions.hpp" -#include +namespace dpnp::extensions::rng +{ +static discrete_dispatch_table_t binomial_dispatch_table; -bool is_verbose_mode(); -void set_barrier_event(sycl::queue queue, std::vector &depends); -void verbose_print(std::string header, - sycl::event first_event, - sycl::event last_event); - -#endif // VERBOSE_H +void init_binomial(py::module_ m) +{ + def_discrete_sampler( + m, "_binomial", + "Fill the result with binomial samples of n trials of probability p", + binomial_dispatch_table); +} +} // namespace dpnp::extensions::rng diff --git a/dpnp/libtensor/extensions/rng/common.hpp b/dpnp/libtensor/extensions/rng/common.hpp new file mode 100644 index 000000000000..4c4396807480 --- /dev/null +++ b/dpnp/libtensor/extensions/rng/common.hpp @@ -0,0 +1,371 @@ +//***************************************************************************** +// Copyright (c) 2026, Intel Corporation +// All rights reserved. +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// - Redistributions of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// - Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// - Neither the name of the copyright holder nor the names of its contributors +// may be used to endorse or promote products derived from this software +// without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE +// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF +// THE POSSIBILITY OF SUCH DAMAGE. +//***************************************************************************** + +#pragma once + +#include +#include +#include +#include +#include +#include +#include + +#include + +#include "dpnp4pybind11.hpp" +#include +#include + +#include "kernels/engines.hpp" +#include "kernels/sample.hpp" +#include "kernels/samplers.hpp" + +// utils extension header +#include "ext/common.hpp" + +// dpnp tensor headers +#include "utils/type_dispatch.hpp" +#include "utils/type_utils.hpp" + +namespace dpnp::extensions::rng +{ +namespace py = pybind11; +namespace td_ns = dpnp::tensor::type_dispatch; +namespace rng_ns = dpnp::kernels::rng; + +// A distribution parameter as passed from Python: either a C-contiguous +// usm_ndarray with the size of the result, whose data pointer is `data` and +// whose type lookup id is `type_id`, or a Python scalar. +struct ParamArg +{ + const char *data = nullptr; + int type_id = -1; + py::object scalar; +}; + +typedef sycl::event (*sample_fn_ptr_t)(sycl::queue &, + const rng_ns::EngineState &, + const std::size_t, + const int, + char *, + const std::vector &, + const std::vector &); + +// Samplers of a result type, indexed by the engine and the result type +using sample_dispatch_table_t = + sample_fn_ptr_t[rng_ns::num_engines][td_ns::num_types]; + +// Discrete samplers, whose result is always int64, indexed by the engine and +// by whether they compute in double (1) or in float (0) +using discrete_dispatch_table_t = sample_fn_ptr_t[rng_ns::num_engines][2]; + +namespace detail +{ +template +rng_ns::Param to_param(const ParamArg &arg, const std::size_t pos) +{ + if (arg.data != nullptr) { + if (arg.type_id != td_ns::GetTypeid{}.get()) { + throw py::value_error("Parameter array " + std::to_string(pos) + + " has an unexpected data type."); + } + return {reinterpret_cast(arg.data), T{}}; + } + return {nullptr, py::cast(arg.scalar)}; +} + +template +SamplerT make_sampler(const std::vector ¶ms, + std::index_sequence) +{ + using param_types = typename SamplerT::param_types; + return SamplerT( + to_param>(params[Is], Is)...); +} + +template +void validate_type_for_device(sycl::queue &exec_q) +{ + dpnp::tensor::type_utils::validate_type_for_device(exec_q); +} + +template +void validate_types_for_device(sycl::queue &exec_q, std::index_sequence) +{ + (validate_type_for_device>(exec_q), ...); +} +template +sycl::event submit_sample(sycl::queue &exec_q, + const rng_ns::EngineState &state, + const SamplerT &sampler, + typename SamplerT::result_type *res, + const std::size_t nelems, + const std::vector &depends) +{ + const std::size_t n_wi = rng_ns::num_work_items(nelems, L); + + return exec_q.submit([&](sycl::handler &cgh) { + cgh.depends_on(depends); + + using SampleKernel = rng_ns::SampleFunctor; + cgh.parallel_for( + sycl::range<1>(n_wi), + SampleKernel(state, sampler, res, nelems, n_wi)); + }); +} +} // namespace detail + +// Checks the parameters and the result against SamplerT, then submits the +// kernel, unless there is nothing to sample. ComputeT is the type the +// sampler computes in, which the device must support. +template +sycl::event sample_impl(sycl::queue &exec_q, + const rng_ns::EngineState &state, + const std::size_t nelems, + const int result_type_id, + char *result, + const std::vector ¶ms, + const std::vector &depends) +{ + using T = typename SamplerT::result_type; + using param_types = typename SamplerT::param_types; + constexpr std::size_t n_params = std::tuple_size_v; + + if (result_type_id != td_ns::GetTypeid{}.get()) { + throw py::value_error("The result array has an unexpected data type."); + } + if (params.size() != n_params) { + throw py::value_error("Expected " + std::to_string(n_params) + + " parameters, got " + + std::to_string(params.size()) + "."); + } + + detail::validate_type_for_device(exec_q); + detail::validate_type_for_device(exec_q); + detail::validate_types_for_device( + exec_q, std::make_index_sequence{}); + + const SamplerT sampler = detail::make_sampler( + params, std::make_index_sequence{}); + + if (nelems == 0) { + return sycl::event{}; + } + + T *res = reinterpret_cast(result); + + if constexpr (rng_ns::can_draw_sequential) { + return detail::submit_sample( + exec_q, state, sampler, res, nelems, depends); + } + else if constexpr (rng_ns::can_draw_blocked) { + // the blocked layout gives the same values, faster on GPUs but + // slower on CPUs + if (exec_q.get_device().is_gpu()) { + return detail::submit_sample( + exec_q, state, sampler, res, nelems, depends); + } + } + return detail::submit_sample( + exec_q, state, sampler, res, nelems, depends); +} + +// Types the floating-point samplers support +template +struct RealTypes + : std::bool_constant || std::is_same_v> +{ +}; + +// Types the integer samplers support +template +struct IntegerTypes : std::bool_constant> +{ +}; + +// The raw output of a bit generator, whole or as 32-bit halves +template +struct RawTypes : std::bool_constant || + std::is_same_v> +{ +}; + +template class SamplerT, + template class SupportedT> +struct SampleFactoryFor +{ + template + struct type + { + fnT get() + { + if constexpr (SupportedT::value) { + return sample_impl>; + } + else { + return nullptr; + } + } + }; +}; + +template