From 1f89cc556225e98ca793797de3dd2429d3bd194d Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Mon, 25 Mar 2024 15:28:51 +0100 Subject: [PATCH] Applied pre-commit formatting rules --- dpnp/CMakeLists.txt | 1 - dpnp/backend/extensions/rng/CMakeLists.txt | 74 ---------- .../extensions/rng/device/common_impl.hpp | 46 +++++-- .../extensions/rng/device/dispatch/matrix.hpp | 28 ++-- .../rng/device/dispatch/table_builder.hpp | 55 ++++---- .../rng/device/engine/base_engine.hpp | 88 ++++++++---- .../device/engine/builder/base_builder.hpp | 64 +++++---- .../rng/device/engine/builder/builder.hpp | 11 +- .../rng/device/engine/builder/mcg31m1.hpp | 15 +- .../rng/device/engine/builder/mcg59.hpp | 15 +- .../rng/device/engine/builder/mrg32k3a.hpp | 15 +- .../device/engine/builder/philox4x32x10.hpp | 15 +- .../rng/device/engine/mcg31m1_engine.hpp | 23 ++-- .../rng/device/engine/mcg59_engine.hpp | 23 ++-- .../rng/device/engine/mrg32k3a_engine.hpp | 39 ++++-- .../device/engine/philox4x32x10_engine.hpp | 39 ++++-- .../extensions/rng/device/gaussian.cpp | 130 ++++++++++++------ .../extensions/rng/device/gaussian.hpp | 18 +-- dpnp/backend/extensions/rng/device/rng_py.cpp | 88 +++++++----- dpnp/backend/extensions/rng/distr_proxy.hpp | 115 ---------------- dpnp/backend/extensions/rng/engine_proxy.hpp | 85 ------------ dpnp/backend/extensions/rng/rng_py.cpp | 89 ------------ 22 files changed, 460 insertions(+), 616 deletions(-) delete mode 100644 dpnp/backend/extensions/rng/CMakeLists.txt delete mode 100644 dpnp/backend/extensions/rng/distr_proxy.hpp delete mode 100644 dpnp/backend/extensions/rng/engine_proxy.hpp delete mode 100644 dpnp/backend/extensions/rng/rng_py.cpp diff --git a/dpnp/CMakeLists.txt b/dpnp/CMakeLists.txt index 2e531b5fafc..b4bdf13abbd 100644 --- a/dpnp/CMakeLists.txt +++ b/dpnp/CMakeLists.txt @@ -58,7 +58,6 @@ build_dpnp_cython_ext_with_backend(dparray ${CMAKE_CURRENT_SOURCE_DIR}/dparray.p add_subdirectory(backend) add_subdirectory(backend/extensions/blas) add_subdirectory(backend/extensions/lapack) -add_subdirectory(backend/extensions/rng) add_subdirectory(backend/extensions/rng/device) add_subdirectory(backend/extensions/vm) add_subdirectory(backend/extensions/sycl_ext) diff --git a/dpnp/backend/extensions/rng/CMakeLists.txt b/dpnp/backend/extensions/rng/CMakeLists.txt deleted file mode 100644 index a3821e96dcc..00000000000 --- a/dpnp/backend/extensions/rng/CMakeLists.txt +++ /dev/null @@ -1,74 +0,0 @@ -# ***************************************************************************** -# Copyright (c) 2023, 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. -# -# 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. -# ***************************************************************************** - - -set(python_module_name _rng_impl) -pybind11_add_module(${python_module_name} MODULE - rng_py.cpp -) - -if (WIN32) - if (${CMAKE_VERSION} VERSION_LESS "3.27") - # this is a work-around for target_link_options inserting option after -link option, cause - # linker to ignore it. - set(CMAKE_CXX_LINK_FLAGS "${CMAKE_CXX_LINK_FLAGS} -fsycl-device-code-split=per_kernel") - endif() -endif() - -set_target_properties(${python_module_name} PROPERTIES CMAKE_POSITION_INDEPENDENT_CODE ON) - -target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../include) -target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../src) - -target_include_directories(${python_module_name} PUBLIC ${Dpctl_INCLUDE_DIRS}) -target_include_directories(${python_module_name} PUBLIC ${Dpctl_TENSOR_INCLUDE_DIR}) - -if (WIN32) - target_compile_options(${python_module_name} PRIVATE - /clang:-fno-approx-func - /clang:-fno-finite-math-only - ) -else() - target_compile_options(${python_module_name} PRIVATE - -fno-approx-func - -fno-finite-math-only - ) -endif() - -target_link_options(${python_module_name} PUBLIC -fsycl-device-code-split=per_kernel) -if (UNIX) - # this option is support on Linux only - target_link_options(${python_module_name} PUBLIC -fsycl-link-huge-device-code) -endif() - -if (DPNP_GENERATE_COVERAGE) - target_link_options(${python_module_name} PRIVATE -fprofile-instr-generate -fcoverage-mapping) -endif() - -target_link_libraries(${python_module_name} PUBLIC MKL::MKL_DPCPP) - -install(TARGETS ${python_module_name} - DESTINATION "dpnp/backend/extensions/rng" -) diff --git a/dpnp/backend/extensions/rng/device/common_impl.hpp b/dpnp/backend/extensions/rng/device/common_impl.hpp index 405d7082023..e7e32baedb2 100644 --- a/dpnp/backend/extensions/rng/device/common_impl.hpp +++ b/dpnp/backend/extensions/rng/device/common_impl.hpp @@ -27,8 +27,8 @@ #include -#include #include +#include namespace dpnp { @@ -57,11 +57,14 @@ struct RngContigFunctor EngineBuilderT engine_; DistributorBuilderT distr_; - DataT * const res_ = nullptr; + DataT *const res_ = nullptr; const std::size_t nelems_; public: - RngContigFunctor(EngineBuilderT& engine, DistributorBuilderT& distr, DataT *res, const std::size_t n_elems) + RngContigFunctor(EngineBuilderT &engine, + DistributorBuilderT &distr, + DataT *res, + const std::size_t n_elems) : engine_(engine), distr_(distr), res_(res), nelems_(n_elems) { } @@ -82,21 +85,33 @@ struct RngContigFunctor DistrT distr = distr_(); if constexpr (enable_sg_load) { - const std::size_t base = vi_per_wi * (nd_it.get_group(0) * nd_it.get_local_range(0) + sg.get_group_id()[0] * max_sg_size); + const std::size_t base = + vi_per_wi * (nd_it.get_group(0) * nd_it.get_local_range(0) + + sg.get_group_id()[0] * max_sg_size); - if ((sg_size == max_sg_size) && (base + vi_per_wi * sg_size < nelems_)) { + if ((sg_size == max_sg_size) && + (base + vi_per_wi * sg_size < nelems_)) { #pragma unroll for (std::uint16_t it = 0; it < vi_per_wi; it += vec_sz) { - std::size_t offset = base + static_cast(it) * static_cast(sg_size); - auto out_multi_ptr = sycl::address_space_cast(&res_[offset]); - - sycl::vec rng_val_vec = mkl_rng_dev::generate(distr, engine); + std::size_t offset = + base + static_cast(it) * + static_cast(sg_size); + auto out_multi_ptr = sycl::address_space_cast< + sycl::access::address_space::global_space, + sycl::access::decorated::yes>(&res_[offset]); + + sycl::vec rng_val_vec = + mkl_rng_dev::generate(distr, engine); sg.store(out_multi_ptr, rng_val_vec); } } else { - for (std::size_t offset = base + sg.get_local_id()[0]; offset < nelems_; offset += sg_size) { - res_[offset] = mkl_rng_dev::generate_single(distr, engine); + for (std::size_t offset = base + sg.get_local_id()[0]; + offset < nelems_; offset += sg_size) + { + res_[offset] = + mkl_rng_dev::generate_single(distr, + engine); } } } @@ -104,9 +119,12 @@ struct RngContigFunctor std::size_t base = nd_it.get_global_linear_id(); base = (base / sg_size) * sg_size * vi_per_wi + (base % sg_size); - for (std::size_t offset = base; offset < std::min(nelems_, base + sg_size * vi_per_wi); offset += sg_size) + for (std::size_t offset = base; + offset < std::min(nelems_, base + sg_size * vi_per_wi); + offset += sg_size) { - res_[offset] = mkl_rng_dev::generate_single(distr, engine); + res_[offset] = mkl_rng_dev::generate_single( + distr, engine); } } } @@ -116,4 +134,4 @@ struct RngContigFunctor } // namespace rng } // namespace ext } // namespace backend -} // namespace dpnp \ No newline at end of file +} // namespace dpnp diff --git a/dpnp/backend/extensions/rng/device/dispatch/matrix.hpp b/dpnp/backend/extensions/rng/device/dispatch/matrix.hpp index d2e5b42fbc2..5b1f87ed66c 100644 --- a/dpnp/backend/extensions/rng/device/dispatch/matrix.hpp +++ b/dpnp/backend/extensions/rng/device/dispatch/matrix.hpp @@ -29,15 +29,15 @@ #include "utils/type_dispatch.hpp" - namespace dpnp::backend::ext::rng::device::dispatch { namespace dpctl_td_ns = dpctl::tensor::type_dispatch; namespace mkl_rng_dev = oneapi::mkl::rng::device; template -struct TypePairDefinedEntry : std::bool_constant && - std::is_same_v> +struct TypePairDefinedEntry + : std::bool_constant && + std::is_same_v> { static constexpr bool is_defined = true; }; @@ -46,11 +46,23 @@ template struct GaussianTypePairSupportFactory { static constexpr bool is_defined = std::disjunction< - TypePairDefinedEntry, - TypePairDefinedEntry, - TypePairDefinedEntry, - TypePairDefinedEntry, + TypePairDefinedEntry, + TypePairDefinedEntry, + TypePairDefinedEntry, + TypePairDefinedEntry, // fall-through dpctl_td_ns::NotDefinedEntry>::is_defined; }; -} // dpnp::backend::ext::rng::device::dispatch +} // namespace dpnp::backend::ext::rng::device::dispatch diff --git a/dpnp/backend/extensions/rng/device/dispatch/table_builder.hpp b/dpnp/backend/extensions/rng/device/dispatch/table_builder.hpp index 5c3b932a2af..e84f1753494 100644 --- a/dpnp/backend/extensions/rng/device/dispatch/table_builder.hpp +++ b/dpnp/backend/extensions/rng/device/dispatch/table_builder.hpp @@ -27,13 +27,13 @@ #include - namespace dpnp::backend::ext::rng::device::dispatch { namespace mkl_rng_dev = oneapi::mkl::rng::device; template typename factory, + template + typename factory, int _no_of_engines, int _no_of_types, int _no_of_methods> @@ -44,8 +44,10 @@ class Dispatch3DTableBuilder const std::vector row_per_method() const { std::vector per_method = { - factory{}.get(), - factory{}.get(), + factory{} + .get(), + factory{} + .get(), }; assert(per_method.size() == _no_of_methods); return per_method; @@ -54,21 +56,21 @@ class Dispatch3DTableBuilder template auto table_per_type_and_method() const { - std::vector> - table_by_type = {row_per_method(), - row_per_method(), - row_per_method(), - row_per_method(), - row_per_method(), - row_per_method(), - row_per_method(), - row_per_method(), - row_per_method(), - row_per_method(), - row_per_method(), - row_per_method(), - row_per_method>(), - row_per_method>()}; + std::vector> table_by_type = { + row_per_method(), + row_per_method(), + row_per_method(), + row_per_method(), + row_per_method(), + row_per_method(), + row_per_method(), + row_per_method(), + row_per_method(), + row_per_method(), + row_per_method(), + row_per_method(), + row_per_method>(), + row_per_method>()}; assert(table_by_type.size() == _no_of_types); return table_by_type; } @@ -78,12 +80,15 @@ class Dispatch3DTableBuilder ~Dispatch3DTableBuilder() = default; template - void populate(funcPtrT table[][_no_of_types][_no_of_methods], std::integer_sequence) const + void populate(funcPtrT table[][_no_of_types][_no_of_methods], + std::integer_sequence) const { - const auto map_by_engine = {table_per_type_and_method>()..., - table_per_type_and_method>()..., - table_per_type_and_method>()..., - table_per_type_and_method>()...}; + const auto map_by_engine = { + table_per_type_and_method>()..., + table_per_type_and_method< + mkl_rng_dev::philox4x32x10>()..., + table_per_type_and_method>()..., + table_per_type_and_method>()...}; assert(map_by_engine.size() == _no_of_engines); std::uint16_t engine_id = 0; @@ -101,4 +106,4 @@ class Dispatch3DTableBuilder } } }; -} // dpnp::backend::ext::rng::device::dispatch +} // namespace dpnp::backend::ext::rng::device::dispatch diff --git a/dpnp/backend/extensions/rng/device/engine/base_engine.hpp b/dpnp/backend/extensions/rng/device/engine/base_engine.hpp index a14a52ebe46..d6f49595c06 100644 --- a/dpnp/backend/extensions/rng/device/engine/base_engine.hpp +++ b/dpnp/backend/extensions/rng/device/engine/base_engine.hpp @@ -27,12 +27,13 @@ #include - namespace dpnp::backend::ext::rng::device::engine { -class EngineType { +class EngineType +{ public: - enum Type : std::uint8_t { + enum Type : std::uint8_t + { MRG32k3a = 0, PHILOX4x32x10, MCG31M1, @@ -43,28 +44,32 @@ class EngineType { EngineType() = default; constexpr EngineType(Type type) : type_(type) {} - constexpr std::uint8_t id() const { + constexpr std::uint8_t id() const + { return static_cast(type_); } - static constexpr std::uint8_t base_id() { + static constexpr std::uint8_t base_id() + { return EngineType(Base).id(); } private: - Type type_; + Type type_; }; // A total number of supported engines == EngineType::Base constexpr std::uint8_t no_of_engines = EngineType::base_id(); -class EngineBase { +class EngineBase +{ private: sycl::queue q_{}; std::vector seed_vec{}; std::vector offset_vec{}; - void validate_vec_size(const std::size_t size) { + void validate_vec_size(const std::size_t size) + { if (size > max_vec_n) { throw std::runtime_error("TODO: add text"); } @@ -73,33 +78,52 @@ class EngineBase { public: EngineBase() {} - EngineBase(sycl::queue &q, std::uint64_t seed, std::uint64_t offset) : - q_(q), seed_vec(1, seed), offset_vec(1, offset) {} + EngineBase(sycl::queue &q, std::uint64_t seed, std::uint64_t offset) + : q_(q), seed_vec(1, seed), offset_vec(1, offset) + { + } - EngineBase(sycl::queue &q, std::vector &seeds, std::uint64_t offset) : - q_(q), seed_vec(seeds), offset_vec(1, offset) { - validate_vec_size(seeds.size()); - } + EngineBase(sycl::queue &q, + std::vector &seeds, + std::uint64_t offset) + : q_(q), seed_vec(seeds), offset_vec(1, offset) + { + validate_vec_size(seeds.size()); + } - EngineBase(sycl::queue &q, std::vector &seeds, std::uint64_t offset) : q_(q), offset_vec(1, offset) { + EngineBase(sycl::queue &q, + std::vector &seeds, + std::uint64_t offset) + : q_(q), offset_vec(1, offset) + { validate_vec_size(seeds.size()); seed_vec.reserve(seeds.size()); seed_vec.assign(seeds.begin(), seeds.end()); } - EngineBase(sycl::queue &q, std::uint64_t seed, std::vector &offsets) : - q_(q), seed_vec(1, seed), offset_vec(offsets) { - validate_vec_size(offsets.size()); - } + EngineBase(sycl::queue &q, + std::uint64_t seed, + std::vector &offsets) + : q_(q), seed_vec(1, seed), offset_vec(offsets) + { + validate_vec_size(offsets.size()); + } - EngineBase(sycl::queue &q, std::vector &seeds, std::vector &offsets) : - q_(q), seed_vec(seeds), offset_vec(offsets) { - validate_vec_size(seeds.size()); - validate_vec_size(offsets.size()); - } + EngineBase(sycl::queue &q, + std::vector &seeds, + std::vector &offsets) + : q_(q), seed_vec(seeds), offset_vec(offsets) + { + validate_vec_size(seeds.size()); + validate_vec_size(offsets.size()); + } - EngineBase(sycl::queue &q, std::vector &seeds, std::vector &offsets) : q_(q), offset_vec(offsets) { + EngineBase(sycl::queue &q, + std::vector &seeds, + std::vector &offsets) + : q_(q), offset_vec(offsets) + { validate_vec_size(seeds.size()); validate_vec_size(offsets.size()); @@ -109,23 +133,27 @@ class EngineBase { virtual ~EngineBase() {} - virtual EngineType get_type() const noexcept { + virtual EngineType get_type() const noexcept + { return EngineType::Base; } - sycl::queue &get_queue() noexcept { + sycl::queue &get_queue() noexcept + { return q_; } - std::vector& get_seeds() noexcept { + std::vector &get_seeds() noexcept + { return seed_vec; } - std::vector& get_offsets() noexcept { + std::vector &get_offsets() noexcept + { return offset_vec; } // static constexpr std::uint8_t max_vec_n = 1; }; -} // dpnp::backend::ext::rng::device::engine +} // namespace dpnp::backend::ext::rng::device::engine diff --git a/dpnp/backend/extensions/rng/device/engine/builder/base_builder.hpp b/dpnp/backend/extensions/rng/device/engine/builder/base_builder.hpp index af515ced15b..b6f0fea3ffd 100644 --- a/dpnp/backend/extensions/rng/device/engine/builder/base_builder.hpp +++ b/dpnp/backend/extensions/rng/device/engine/builder/base_builder.hpp @@ -31,11 +31,11 @@ #include namespace mkl_rng_dev = oneapi::mkl::rng::device; - namespace dpnp::backend::ext::rng::device::engine::builder { template -class BaseBuilder { +class BaseBuilder +{ private: static constexpr std::uint8_t max_n = EngineBase::max_vec_n; @@ -74,18 +74,21 @@ class BaseBuilder { inline auto operator()(void) const { switch (no_of_seeds) { - case 1: { - if constexpr (std::is_same_v>) { - // issue with mcg59<>() constructor which breaks compilation - return EngineT(seeds[0], offsets[0]); - } - else { - return EngineT({seeds[0]}, offsets[0]); - } + case 1: + { + if constexpr (std::is_same_v>) + { + // issue with mcg59<>() constructor which breaks compilation + return EngineT(seeds[0], offsets[0]); } - // TODO: implement full switch - default: - break; + else { + return EngineT({seeds[0]}, offsets[0]); + } + } + // TODO: implement full switch + default: + break; } return EngineT(); } @@ -93,30 +96,35 @@ class BaseBuilder { inline auto operator()(const OffsetT offset) const { switch (no_of_seeds) { - case 1: { - if constexpr (std::is_same_v>) { - // issue with mcg59<>() constructor which breaks compilation - return EngineT(seeds[0], offsets[0] + offset); - } - else { - return EngineT({seeds[0]}, {offsets[0] + offset}); - } + case 1: + { + if constexpr (std::is_same_v>) + { + // issue with mcg59<>() constructor which breaks compilation + return EngineT(seeds[0], offsets[0] + offset); } - // TODO: implement full switch - default: - break; + else { + return EngineT({seeds[0]}, {offsets[0] + offset}); + } + } + // TODO: implement full switch + default: + break; } return EngineT(); } // TODO: remove - void print() { - std::cout << "vector size = " << std::to_string(EngineT::vec_size) << std::endl; + void print() + { + std::cout << "vector size = " << std::to_string(EngineT::vec_size) + << std::endl; std::cout << "list_of_seeds: "; - for (auto &val: seeds) { + for (auto &val : seeds) { std::cout << std::to_string(val) << ", "; } std::cout << std::endl; } }; -} // dpnp::backend::ext::rng::device::engine::builder +} // namespace dpnp::backend::ext::rng::device::engine::builder diff --git a/dpnp/backend/extensions/rng/device/engine/builder/builder.hpp b/dpnp/backend/extensions/rng/device/engine/builder/builder.hpp index d1bc4643ddc..1fd9c9c9e89 100644 --- a/dpnp/backend/extensions/rng/device/engine/builder/builder.hpp +++ b/dpnp/backend/extensions/rng/device/engine/builder/builder.hpp @@ -25,14 +25,15 @@ #pragma once - namespace dpnp::backend::ext::rng::device::engine::builder { template -class Builder {}; -} // dpnp::backend::ext::rng::device::engine::builder +class Builder +{ +}; +} // namespace dpnp::backend::ext::rng::device::engine::builder -#include "mrg32k3a.hpp" -#include "philox4x32x10.hpp" #include "mcg31m1.hpp" #include "mcg59.hpp" +#include "mrg32k3a.hpp" +#include "philox4x32x10.hpp" diff --git a/dpnp/backend/extensions/rng/device/engine/builder/mcg31m1.hpp b/dpnp/backend/extensions/rng/device/engine/builder/mcg31m1.hpp index 8e99afde372..b94b93b35b3 100644 --- a/dpnp/backend/extensions/rng/device/engine/builder/mcg31m1.hpp +++ b/dpnp/backend/extensions/rng/device/engine/builder/mcg31m1.hpp @@ -27,18 +27,25 @@ #include -#include "base_engine.hpp" #include "base_builder.hpp" +#include "base_engine.hpp" namespace dpnp::backend::ext::rng::device::engine::builder { namespace mkl_rng_dev = oneapi::mkl::rng::device; template -class Builder> : public BaseBuilder, std::uint32_t, std::uint64_t> { +class Builder> + : public BaseBuilder, + std::uint32_t, + std::uint64_t> +{ public: using EngineType = mkl_rng_dev::mcg31m1; - Builder(EngineBase *engine) : BaseBuilder(engine) {} + Builder(EngineBase *engine) + : BaseBuilder(engine) + { + } }; -} // dpnp::backend::ext::rng::device::engine::builder +} // namespace dpnp::backend::ext::rng::device::engine::builder diff --git a/dpnp/backend/extensions/rng/device/engine/builder/mcg59.hpp b/dpnp/backend/extensions/rng/device/engine/builder/mcg59.hpp index d4bc754d8bd..7f8e003b87c 100644 --- a/dpnp/backend/extensions/rng/device/engine/builder/mcg59.hpp +++ b/dpnp/backend/extensions/rng/device/engine/builder/mcg59.hpp @@ -27,18 +27,25 @@ #include -#include "base_engine.hpp" #include "base_builder.hpp" +#include "base_engine.hpp" namespace dpnp::backend::ext::rng::device::engine::builder { namespace mkl_rng_dev = oneapi::mkl::rng::device; template -class Builder> : public BaseBuilder, std::uint32_t, std::uint64_t> { +class Builder> + : public BaseBuilder, + std::uint32_t, + std::uint64_t> +{ public: using EngineType = mkl_rng_dev::mcg59; - Builder(EngineBase *engine) : BaseBuilder(engine) {} + Builder(EngineBase *engine) + : BaseBuilder(engine) + { + } }; -} // dpnp::backend::ext::rng::device::engine::builder +} // namespace dpnp::backend::ext::rng::device::engine::builder diff --git a/dpnp/backend/extensions/rng/device/engine/builder/mrg32k3a.hpp b/dpnp/backend/extensions/rng/device/engine/builder/mrg32k3a.hpp index 940df64f3a0..a5ab3470b03 100644 --- a/dpnp/backend/extensions/rng/device/engine/builder/mrg32k3a.hpp +++ b/dpnp/backend/extensions/rng/device/engine/builder/mrg32k3a.hpp @@ -27,18 +27,25 @@ #include -#include "base_engine.hpp" #include "base_builder.hpp" +#include "base_engine.hpp" namespace dpnp::backend::ext::rng::device::engine::builder { namespace mkl_rng_dev = oneapi::mkl::rng::device; template -class Builder> : public BaseBuilder, std::uint32_t, std::uint64_t> { +class Builder> + : public BaseBuilder, + std::uint32_t, + std::uint64_t> +{ public: using EngineType = mkl_rng_dev::mrg32k3a; - Builder(EngineBase *engine) : BaseBuilder(engine) {} + Builder(EngineBase *engine) + : BaseBuilder(engine) + { + } }; -} // dpnp::backend::ext::rng::device::engine::builder +} // namespace dpnp::backend::ext::rng::device::engine::builder diff --git a/dpnp/backend/extensions/rng/device/engine/builder/philox4x32x10.hpp b/dpnp/backend/extensions/rng/device/engine/builder/philox4x32x10.hpp index b0abc5c86b8..932c5ee00d7 100644 --- a/dpnp/backend/extensions/rng/device/engine/builder/philox4x32x10.hpp +++ b/dpnp/backend/extensions/rng/device/engine/builder/philox4x32x10.hpp @@ -27,18 +27,25 @@ #include -#include "base_engine.hpp" #include "base_builder.hpp" +#include "base_engine.hpp" namespace dpnp::backend::ext::rng::device::engine::builder { namespace mkl_rng_dev = oneapi::mkl::rng::device; template -class Builder> : public BaseBuilder, std::uint64_t, std::uint64_t> { +class Builder> + : public BaseBuilder, + std::uint64_t, + std::uint64_t> +{ public: using EngineType = mkl_rng_dev::philox4x32x10; - Builder(EngineBase *engine) : BaseBuilder(engine) {} + Builder(EngineBase *engine) + : BaseBuilder(engine) + { + } }; -} // dpnp::backend::ext::rng::device::engine::builder +} // namespace dpnp::backend::ext::rng::device::engine::builder diff --git a/dpnp/backend/extensions/rng/device/engine/mcg31m1_engine.hpp b/dpnp/backend/extensions/rng/device/engine/mcg31m1_engine.hpp index 2a8cdee310c..4022bd33c18 100644 --- a/dpnp/backend/extensions/rng/device/engine/mcg31m1_engine.hpp +++ b/dpnp/backend/extensions/rng/device/engine/mcg31m1_engine.hpp @@ -27,19 +27,26 @@ #include "base_engine.hpp" - namespace dpnp::backend::ext::rng::device::engine { -class MCG31M1 : public EngineBase { +class MCG31M1 : public EngineBase +{ public: - MCG31M1(sycl::queue &q, std::uint32_t seed, std::uint64_t offset = 0) : - EngineBase(q, seed, offset) {} + MCG31M1(sycl::queue &q, std::uint32_t seed, std::uint64_t offset = 0) + : EngineBase(q, seed, offset) + { + } - MCG31M1(sycl::queue &q, std::vector &seeds, std::uint64_t offset = 0) : - EngineBase(q, seeds, offset) {} + MCG31M1(sycl::queue &q, + std::vector &seeds, + std::uint64_t offset = 0) + : EngineBase(q, seeds, offset) + { + } - virtual EngineType get_type() const noexcept override { + virtual EngineType get_type() const noexcept override + { return EngineType::MCG31M1; } }; -} // dpnp::backend::ext::rng::device::engine +} // namespace dpnp::backend::ext::rng::device::engine diff --git a/dpnp/backend/extensions/rng/device/engine/mcg59_engine.hpp b/dpnp/backend/extensions/rng/device/engine/mcg59_engine.hpp index 313a351426b..e0fbf1741a3 100644 --- a/dpnp/backend/extensions/rng/device/engine/mcg59_engine.hpp +++ b/dpnp/backend/extensions/rng/device/engine/mcg59_engine.hpp @@ -27,19 +27,26 @@ #include "base_engine.hpp" - namespace dpnp::backend::ext::rng::device::engine { -class MCG59 : public EngineBase { +class MCG59 : public EngineBase +{ public: - MCG59(sycl::queue &q, std::uint32_t seed, std::uint64_t offset = 0) : - EngineBase(q, seed, offset) {} + MCG59(sycl::queue &q, std::uint32_t seed, std::uint64_t offset = 0) + : EngineBase(q, seed, offset) + { + } - MCG59(sycl::queue &q, std::vector &seeds, std::uint64_t offset = 0) : - EngineBase(q, seeds, offset) {} + MCG59(sycl::queue &q, + std::vector &seeds, + std::uint64_t offset = 0) + : EngineBase(q, seeds, offset) + { + } - virtual EngineType get_type() const noexcept override { + virtual EngineType get_type() const noexcept override + { return EngineType::MCG59; } }; -} // dpnp::backend::ext::rng::device::engine +} // namespace dpnp::backend::ext::rng::device::engine diff --git a/dpnp/backend/extensions/rng/device/engine/mrg32k3a_engine.hpp b/dpnp/backend/extensions/rng/device/engine/mrg32k3a_engine.hpp index a0fcad167af..d6657d6d8d8 100644 --- a/dpnp/backend/extensions/rng/device/engine/mrg32k3a_engine.hpp +++ b/dpnp/backend/extensions/rng/device/engine/mrg32k3a_engine.hpp @@ -27,25 +27,40 @@ #include "base_engine.hpp" - namespace dpnp::backend::ext::rng::device::engine { -class MRG32k3a : public EngineBase { +class MRG32k3a : public EngineBase +{ public: - MRG32k3a(sycl::queue &q, std::uint32_t seed, std::uint64_t offset = 0) : - EngineBase(q, seed, offset) {} + MRG32k3a(sycl::queue &q, std::uint32_t seed, std::uint64_t offset = 0) + : EngineBase(q, seed, offset) + { + } - MRG32k3a(sycl::queue &q, std::vector &seeds, std::uint64_t offset = 0) : - EngineBase(q, seeds, offset) {} + MRG32k3a(sycl::queue &q, + std::vector &seeds, + std::uint64_t offset = 0) + : EngineBase(q, seeds, offset) + { + } - MRG32k3a(sycl::queue &q, std::uint32_t seed, std::vector &offsets) : - EngineBase(q, seed, offsets) {} + MRG32k3a(sycl::queue &q, + std::uint32_t seed, + std::vector &offsets) + : EngineBase(q, seed, offsets) + { + } - MRG32k3a(sycl::queue &q, std::vector &seeds, std::vector &offsets) : - EngineBase(q, seeds, offsets) {} + MRG32k3a(sycl::queue &q, + std::vector &seeds, + std::vector &offsets) + : EngineBase(q, seeds, offsets) + { + } - virtual EngineType get_type() const noexcept override { + virtual EngineType get_type() const noexcept override + { return EngineType::MRG32k3a; } }; -} // dpnp::backend::ext::rng::device::engine +} // namespace dpnp::backend::ext::rng::device::engine diff --git a/dpnp/backend/extensions/rng/device/engine/philox4x32x10_engine.hpp b/dpnp/backend/extensions/rng/device/engine/philox4x32x10_engine.hpp index bba8a5c63eb..bef1e7d2119 100644 --- a/dpnp/backend/extensions/rng/device/engine/philox4x32x10_engine.hpp +++ b/dpnp/backend/extensions/rng/device/engine/philox4x32x10_engine.hpp @@ -27,25 +27,40 @@ #include "base_engine.hpp" - namespace dpnp::backend::ext::rng::device::engine { -class PHILOX4x32x10 : public EngineBase { +class PHILOX4x32x10 : public EngineBase +{ public: - PHILOX4x32x10(sycl::queue &q, std::uint64_t seed, std::uint64_t offset = 0) : - EngineBase(q, seed, offset) {} + PHILOX4x32x10(sycl::queue &q, std::uint64_t seed, std::uint64_t offset = 0) + : EngineBase(q, seed, offset) + { + } - PHILOX4x32x10(sycl::queue &q, std::vector &seeds, std::uint64_t offset = 0) : - EngineBase(q, seeds, offset) {} + PHILOX4x32x10(sycl::queue &q, + std::vector &seeds, + std::uint64_t offset = 0) + : EngineBase(q, seeds, offset) + { + } - PHILOX4x32x10(sycl::queue &q, std::uint64_t seed, std::vector &offsets) : - EngineBase(q, seed, offsets) {} + PHILOX4x32x10(sycl::queue &q, + std::uint64_t seed, + std::vector &offsets) + : EngineBase(q, seed, offsets) + { + } - PHILOX4x32x10(sycl::queue &q, std::vector &seeds, std::vector &offsets) : - EngineBase(q, seeds, offsets) {} + PHILOX4x32x10(sycl::queue &q, + std::vector &seeds, + std::vector &offsets) + : EngineBase(q, seeds, offsets) + { + } - virtual EngineType get_type() const noexcept override { + virtual EngineType get_type() const noexcept override + { return EngineType::PHILOX4x32x10; } }; -} // dpnp::backend::ext::rng::device::engine +} // namespace dpnp::backend::ext::rng::device::engine diff --git a/dpnp/backend/extensions/rng/device/gaussian.cpp b/dpnp/backend/extensions/rng/device/gaussian.cpp index 0d7529b2077..74866c628aa 100644 --- a/dpnp/backend/extensions/rng/device/gaussian.cpp +++ b/dpnp/backend/extensions/rng/device/gaussian.cpp @@ -26,20 +26,19 @@ #include // dpctl tensor headers +#include "kernels/alignment.hpp" #include "utils/output_validation.hpp" #include "utils/type_dispatch.hpp" #include "utils/type_utils.hpp" -#include "kernels/alignment.hpp" -#include "gaussian.hpp" #include "common_impl.hpp" +#include "gaussian.hpp" #include "engine/builder/builder.hpp" #include "dispatch/matrix.hpp" #include "dispatch/table_builder.hpp" - namespace dpnp::backend::ext::rng::device { namespace dpctl_krn_ns = dpctl::tensor::kernels::alignment_utils; @@ -54,18 +53,24 @@ using dpctl_krn_ns::required_alignment; constexpr auto no_of_methods = 2; // number of methods of gaussian distribution -constexpr auto seq_of_vec_sizes = std::integer_sequence{}; +constexpr auto seq_of_vec_sizes = + std::integer_sequence{}; constexpr auto vec_sizes_len = seq_of_vec_sizes.size(); constexpr auto no_of_engines = engine::no_of_engines * vec_sizes_len; -template -inline auto find_vec_size_impl(const VecSizeT vec_size, std::index_sequence) { - return std::min({ ((Ints == vec_size) ? Indices : sizeof...(Indices))... }); +template +inline auto find_vec_size_impl(const VecSizeT vec_size, + std::index_sequence) +{ + return std::min({((Ints == vec_size) ? Indices : sizeof...(Indices))...}); } -template -int find_vec_size(const VecSizeT vec_size, std::integer_sequence) { - auto res = find_vec_size_impl(vec_size, std::make_index_sequence{}); +template +int find_vec_size(const VecSizeT vec_size, + std::integer_sequence) +{ + auto res = find_vec_size_impl( + vec_size, std::make_index_sequence{}); return (res == sizeof...(Ints)) ? -1 : res; } @@ -99,9 +104,14 @@ typedef sycl::event (*gaussian_impl_fn_ptr_t)(engine::EngineBase *engine, char *, const std::vector &); -static gaussian_impl_fn_ptr_t gaussian_dispatch_table[no_of_engines][dpctl_td_ns::num_types][no_of_methods]; +static gaussian_impl_fn_ptr_t gaussian_dispatch_table[no_of_engines] + [dpctl_td_ns::num_types] + [no_of_methods]; -template +template class gaussian_kernel; template @@ -123,7 +133,8 @@ static sycl::event gaussian_impl(engine::EngineBase *engine, constexpr std::size_t items_per_wi = 4; constexpr std::size_t local_size = 256; const std::size_t wg_items = local_size * vec_sz * items_per_wi; - const std::size_t global_size = ((n + wg_items - 1) / (wg_items)) * local_size; + const std::size_t global_size = + ((n + wg_items - 1) / (wg_items)) * local_size; sycl::event distr_event; @@ -140,42 +151,57 @@ static sycl::event gaussian_impl(engine::EngineBase *engine, if (is_aligned(out_ptr)) { constexpr bool enable_sg_load = true; - using KernelName = gaussian_kernel; - - cgh.parallel_for(sycl::nd_range<1>({global_size}, {local_size}), - details::RngContigFunctor(eng_builder, dist_builder, out, n)); + using KernelName = + gaussian_kernel; + + cgh.parallel_for( + sycl::nd_range<1>({global_size}, {local_size}), + details::RngContigFunctor( + eng_builder, dist_builder, out, n)); } else { constexpr bool disable_sg_load = false; - using InnerKernelName = gaussian_kernel; - using KernelName = disabled_sg_loadstore_wrapper_krn; - - cgh.parallel_for(sycl::nd_range<1>({global_size}, {local_size}), - details::RngContigFunctor(eng_builder, dist_builder, out, n)); + using InnerKernelName = + gaussian_kernel; + using KernelName = + disabled_sg_loadstore_wrapper_krn; + + cgh.parallel_for( + sycl::nd_range<1>({global_size}, {local_size}), + details::RngContigFunctor( + eng_builder, dist_builder, out, n)); } }); } catch (oneapi::mkl::exception const &e) { std::stringstream error_msg; - error_msg << "Unexpected MKL exception caught during gaussian call:\nreason: " << e.what(); + error_msg + << "Unexpected MKL exception caught during gaussian call:\nreason: " + << e.what(); throw std::runtime_error(error_msg.str()); } catch (sycl::exception const &e) { std::stringstream error_msg; - error_msg << "Unexpected SYCL exception caught during gaussian call:\n" << e.what(); + error_msg << "Unexpected SYCL exception caught during gaussian call:\n" + << e.what(); throw std::runtime_error(error_msg.str()); } return distr_event; } -std::pair gaussian(engine::EngineBase *engine, - const std::uint8_t method_id, - const std::uint8_t vec_size, - const double mean, - const double stddev, - const std::uint64_t n, - dpctl::tensor::usm_ndarray res, - const std::vector &depends) +std::pair + gaussian(engine::EngineBase *engine, + const std::uint8_t method_id, + const std::uint8_t vec_size, + const double mean, + const double stddev, + const std::uint64_t n, + dpctl::tensor::usm_ndarray res, + const std::vector &depends) { auto &exec_q = engine->get_queue(); @@ -196,42 +222,52 @@ std::pair gaussian(engine::EngineBase *engine, dpctl::tensor::validation::AmpleMemory::throw_if_not_ample(res, res_nelems); if (!dpctl::utils::queues_are_compatible(exec_q, {res})) { - throw py::value_error("Execution queue is not compatible with the allocation queue"); + throw py::value_error( + "Execution queue is not compatible with the allocation queue"); } bool is_res_c_contig = res.is_c_contiguous(); if (!is_res_c_contig) { - throw std::runtime_error("Only population of contiguous array is supported."); + throw std::runtime_error( + "Only population of contiguous array is supported."); } auto enginge_id = engine->get_type().id(); if (enginge_id >= engine::no_of_engines) { - throw std::runtime_error("Unknown engine type=" + std::to_string(enginge_id) + " for gaussian distribution."); + throw std::runtime_error( + "Unknown engine type=" + std::to_string(enginge_id) + + " for gaussian distribution."); } if (method_id >= no_of_methods) { - throw std::runtime_error("Unknown method=" + std::to_string(method_id) + " for gaussian distribution."); + throw std::runtime_error("Unknown method=" + std::to_string(method_id) + + " for gaussian distribution."); } int vec_size_id = find_vec_size(vec_size, seq_of_vec_sizes); if (vec_size_id < 0) { - throw std::runtime_error("Vector size=" + std::to_string(vec_size) + " is out of supported range"); + throw std::runtime_error("Vector size=" + std::to_string(vec_size) + + " is out of supported range"); } enginge_id = enginge_id * vec_sizes_len + vec_size_id; auto array_types = dpctl_td_ns::usm_ndarray_types(); int res_type_id = array_types.typenum_to_lookup_id(res.get_typenum()); - auto gaussian_fn = gaussian_dispatch_table[enginge_id][res_type_id][method_id]; + auto gaussian_fn = + gaussian_dispatch_table[enginge_id][res_type_id][method_id]; if (gaussian_fn == nullptr) { - throw py::value_error("No gaussian implementation defined for a required type"); + throw py::value_error( + "No gaussian implementation defined for a required type"); } char *res_data = res.get_data(); - sycl::event gaussian_ev = gaussian_fn(engine, mean, stddev, n, res_data, depends); + sycl::event gaussian_ev = + gaussian_fn(engine, mean, stddev, n, res_data, depends); - sycl::event ht_ev = dpctl::utils::keep_args_alive(exec_q, {res}, {gaussian_ev}); - return std::make_pair(ht_ev, gaussian_ev); + sycl::event ht_ev = + dpctl::utils::keep_args_alive(exec_q, {res}, {gaussian_ev}); + return std::make_pair(ht_ev, gaussian_ev); } template @@ -239,7 +275,8 @@ struct GaussianContigFactory { fnT get() { - if constexpr (dispatch::GaussianTypePairSupportFactory::is_defined) { + if constexpr (dispatch::GaussianTypePairSupportFactory::is_defined) { return gaussian_impl; } else { @@ -250,7 +287,10 @@ struct GaussianContigFactory void init_gaussian_dispatch_3d_table(void) { - dispatch::Dispatch3DTableBuilder contig; + dispatch::Dispatch3DTableBuilder + contig; contig.populate(gaussian_dispatch_table, seq_of_vec_sizes); } -} // dpnp::backend::ext::rng::device +} // namespace dpnp::backend::ext::rng::device diff --git a/dpnp/backend/extensions/rng/device/gaussian.hpp b/dpnp/backend/extensions/rng/device/gaussian.hpp index 7b8b36cd98b..00973a5d4e5 100644 --- a/dpnp/backend/extensions/rng/device/gaussian.hpp +++ b/dpnp/backend/extensions/rng/device/gaussian.hpp @@ -29,17 +29,17 @@ #include "engine/base_engine.hpp" - namespace dpnp::backend::ext::rng::device { -extern std::pair gaussian(engine::EngineBase *engine, - const std::uint8_t method_id, - const std::uint8_t vec_size, - const double mean, - const double stddev, - const std::uint64_t n, - dpctl::tensor::usm_ndarray res, - const std::vector &depends = {}); +extern std::pair + gaussian(engine::EngineBase *engine, + const std::uint8_t method_id, + const std::uint8_t vec_size, + const double mean, + const double stddev, + const std::uint64_t n, + dpctl::tensor::usm_ndarray res, + const std::vector &depends = {}); extern void init_gaussian_dispatch_3d_table(void); } // namespace dpnp::backend::ext::rng::device diff --git a/dpnp/backend/extensions/rng/device/rng_py.cpp b/dpnp/backend/extensions/rng/device/rng_py.cpp index 354061597b6..27cbdb80d8d 100644 --- a/dpnp/backend/extensions/rng/device/rng_py.cpp +++ b/dpnp/backend/extensions/rng/device/rng_py.cpp @@ -35,25 +35,24 @@ #include "gaussian.hpp" -#include "engine/mrg32k3a_engine.hpp" -#include "engine/philox4x32x10_engine.hpp" #include "engine/mcg31m1_engine.hpp" #include "engine/mcg59_engine.hpp" - +#include "engine/mrg32k3a_engine.hpp" +#include "engine/philox4x32x10_engine.hpp" namespace mkl_rng = oneapi::mkl::rng; namespace rng_dev_ext = dpnp::backend::ext::rng::device; namespace rng_dev_engine = dpnp::backend::ext::rng::device::engine; namespace py = pybind11; - // populate dispatch 3-D tables void init_dispatch_3d_tables(void) { rng_dev_ext::init_gaussian_dispatch_3d_table(); } -class PyEngineBase : public rng_dev_engine::EngineBase { +class PyEngineBase : public rng_dev_engine::EngineBase +{ public: // inherit the constructor using EngineBase::EngineBase; @@ -63,45 +62,70 @@ class PyEngineBase : public rng_dev_engine::EngineBase { // PYBIND11_OVERRIDE_PURE( // sycl::queue&, /* Return type */ // EngineBase, /* Parent class */ - // get_queue, /* Name of function in C++ (must match Python name) */ + // get_queue, /* Name of function in C++ (must match Python name) + // */ // ); // } }; - PYBIND11_MODULE(_rng_dev_impl, m) { init_dispatch_3d_tables(); - py::class_(m, "EngineBase") + py::class_( + m, "EngineBase") .def(py::init<>()) .def("get_queue", &rng_dev_engine::EngineBase::get_queue); - py::class_(m, "MRG32k3a") - .def(py::init(), py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0) - .def(py::init &, std::uint64_t>(), py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0) - .def(py::init &>(), py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = py::list()) - .def(py::init &, std::vector &>(), py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = py::list()); - - py::class_(m, "PHILOX4x32x10") - .def(py::init(), py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0) - .def(py::init &, std::uint64_t>(), py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0) - .def(py::init &>(), py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = py::list()) - .def(py::init &, std::vector &>(), py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = py::list()); - - py::class_(m, "MCG31M1") - .def(py::init(), py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0) - .def(py::init &, std::uint64_t>(), py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0); + py::class_(m, + "MRG32k3a") + .def(py::init(), + py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0) + .def(py::init &, + std::uint64_t>(), + py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0) + .def(py::init &>(), + py::arg("sycl_queue"), py::arg("seed"), + py::arg("offset") = py::list()) + .def(py::init &, + std::vector &>(), + py::arg("sycl_queue"), py::arg("seed"), + py::arg("offset") = py::list()); + + py::class_( + m, "PHILOX4x32x10") + .def(py::init(), + py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0) + .def(py::init &, + std::uint64_t>(), + py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0) + .def(py::init &>(), + py::arg("sycl_queue"), py::arg("seed"), + py::arg("offset") = py::list()) + .def(py::init &, + std::vector &>(), + py::arg("sycl_queue"), py::arg("seed"), + py::arg("offset") = py::list()); + + py::class_(m, + "MCG31M1") + .def(py::init(), + py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0) + .def(py::init &, + std::uint64_t>(), + py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0); py::class_(m, "MCG59") - .def(py::init(), py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0) - .def(py::init &, std::uint64_t>(), py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0); - - m.def("_gaussian", &rng_dev_ext::gaussian, - "", - py::arg("engine"), - py::arg("method_id"), py::arg("vec_size"), - py::arg("mean"), py::arg("stddev"), - py::arg("n"), py::arg("res"), + .def(py::init(), + py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0) + .def(py::init &, + std::uint64_t>(), + py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0); + + m.def("_gaussian", &rng_dev_ext::gaussian, "", py::arg("engine"), + py::arg("method_id"), py::arg("vec_size"), py::arg("mean"), + py::arg("stddev"), py::arg("n"), py::arg("res"), py::arg("depends") = py::list()); } diff --git a/dpnp/backend/extensions/rng/distr_proxy.hpp b/dpnp/backend/extensions/rng/distr_proxy.hpp deleted file mode 100644 index 69c895f0608..00000000000 --- a/dpnp/backend/extensions/rng/distr_proxy.hpp +++ /dev/null @@ -1,115 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2023, 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. -// -// 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 - -namespace dpnp -{ -namespace backend -{ -namespace ext -{ -namespace rng -{ -namespace mkl_rng = oneapi::mkl::rng; - -template -class UnifromBase { -public: - using method_type = Method; - using result_type = DataType; - using distr_type = mkl_rng::uniform; - - - UnifromBase(DataType a, DataType b) { - distr = std::make_unique(a, b); - } - - distr_type& get_distr() const { return *distr; } - -private: - std::unique_ptr distr; -} - -class UniformB { -public: - - Uniform(sycl::queue queue, const int type_num, double a, double b): q(queue) { - if () - } - - void generate() { - // create distribution - } - -private: - int typenum; - std::unique_ptr engine; - - - - using method_type = Method; - // using result_type = RealType; - - DistrProxy(sycl::queue queue, std::uint32_t seed): q(queue) { - engine = std::make_unique(queue, seed); - }; - - // template - DistrProxy(sycl::queue queue, std::vector vec_seed): q(queue) { - switch (vec_seed.size()) { - case 1: - engine = std::make_unique(queue, std::initializer_list({vec_seed[0]})); - break; - case 2: - engine = std::make_unique(queue, std::initializer_list({vec_seed[0], vec_seed[1]})); - break; - case 3: - engine = std::make_unique(queue, std::initializer_list({vec_seed[0], vec_seed[1], vec_seed[2]})); - break; - default: - // TODO need to get rid of the limitation for seed vector length - throw std::runtime_error("Too long seed vector"); - } - }; - - ~DistrProxy() = default; - - sycl::queue& get_queue() { return q;} - EngineT& get_engine() const { return *engine;} - -private: - sycl::queue q; - std::unique_ptr engine; - // engine_t engine; - // sycl::queue q; - -}; -} // namespace lapack -} // namespace ext -} // namespace backend -} // namespace rng diff --git a/dpnp/backend/extensions/rng/engine_proxy.hpp b/dpnp/backend/extensions/rng/engine_proxy.hpp deleted file mode 100644 index 89346718fef..00000000000 --- a/dpnp/backend/extensions/rng/engine_proxy.hpp +++ /dev/null @@ -1,85 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2023, 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. -// -// 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 - -namespace dpnp -{ -namespace backend -{ -namespace ext -{ -namespace rng -{ -class EngineBase { -public: - EngineBase(sycl::queue queue) { - q = std::make_unique(queue); - }; - - sycl::queue& get_queue() { return *q; } - -private: - std::unique_ptr q; -}; - - -template -class EngineProxy: public EngineBase { -public: - using engine_t = EngineT; - - EngineProxy(sycl::queue queue, SeedT seed): EngineBase(queue) { - engine = std::make_unique(queue, seed); - }; - - // template - EngineProxy(sycl::queue queue, std::vector vec_seed): EngineBase(queue) { - switch (vec_seed.size()) { - case 1: - engine = std::make_unique(queue, std::initializer_list({vec_seed[0]})); - break; - case 2: - engine = std::make_unique(queue, std::initializer_list({vec_seed[0], vec_seed[1]})); - break; - case 3: - engine = std::make_unique(queue, std::initializer_list({vec_seed[0], vec_seed[1], vec_seed[2]})); - break; - default: - // TODO need to get rid of the limitation for seed vector length - throw std::runtime_error("Too long seed vector"); - } - }; - - // ~EngineProxy() = default; - -private: - std::unique_ptr engine; - -}; -} // namespace lapack -} // namespace ext -} // namespace backend -} // namespace rng diff --git a/dpnp/backend/extensions/rng/rng_py.cpp b/dpnp/backend/extensions/rng/rng_py.cpp deleted file mode 100644 index a360f3c8140..00000000000 --- a/dpnp/backend/extensions/rng/rng_py.cpp +++ /dev/null @@ -1,89 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2023, 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. -// -// 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 file defines functions of dpnp.backend._rng_impl extensions -// -//***************************************************************************** - -#include -#include - -#include - -#include - -#include "engine_proxy.hpp" -// #include "syevd.hpp" - -namespace mkl_rng = oneapi::mkl::rng; -namespace rng_ext = dpnp::backend::ext::rng; -namespace py = pybind11; - -// populate dispatch vectors -void init_dispatch_vectors(void) -{ - // lapack_ext::init_syevd_dispatch_vector(); -} - -// populate dispatch tables -void init_dispatch_tables(void) -{ - // lapack_ext::init_heevd_dispatch_table(); -} - - -PYBIND11_MODULE(_rng_impl, m) -{ - using engine_base_t = rng_ext::EngineBase; - py::class_ engine_base(m, "EngineBase"); - engine_base.def(py::init()) - .def("get_queue", &engine_base_t::get_queue); - - using mt19937_engine_t = rng_ext::EngineProxy; - py::class_(m, "mt19937", engine_base) - .def(py::init()) - .def(py::init>()); - - using mcg59_engine_t = rng_ext::EngineProxy; - py::class_(m, "mcg59", engine_base) - .def(py::init()); - - // init_dispatch_vectors(); - // init_dispatch_tables(); - - // m.def("_heevd", &lapack_ext::heevd, - // "Call `heevd` from OneMKL LAPACK library to return " - // "the eigenvalues and eigenvectors of a complex Hermitian matrix", - // py::arg("sycl_queue"), py::arg("jobz"), py::arg("upper_lower"), - // py::arg("eig_vecs"), py::arg("eig_vals"), - // py::arg("depends") = py::list()); - - // m.def("_syevd", &lapack_ext::syevd, - // "Call `syevd` from OneMKL LAPACK library to return " - // "the eigenvalues and eigenvectors of a real symmetric matrix", - // py::arg("sycl_queue"), py::arg("jobz"), py::arg("upper_lower"), - // py::arg("eig_vecs"), py::arg("eig_vals"), - // py::arg("depends") = py::list()); -}