diff --git a/dpnp/CMakeLists.txt b/dpnp/CMakeLists.txt index d9c95b62c0b..c4a3889fda3 100644 --- a/dpnp/CMakeLists.txt +++ b/dpnp/CMakeLists.txt @@ -58,6 +58,8 @@ 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/device) +add_subdirectory(backend/extensions/rng/host) add_subdirectory(backend/extensions/vm) add_subdirectory(backend/extensions/sycl_ext) add_subdirectory(backend/extensions/ufunc) diff --git a/dpnp/backend/extensions/rng/device/CMakeLists.txt b/dpnp/backend/extensions/rng/device/CMakeLists.txt new file mode 100644 index 00000000000..37df343f40b --- /dev/null +++ b/dpnp/backend/extensions/rng/device/CMakeLists.txt @@ -0,0 +1,78 @@ +# ***************************************************************************** +# 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_dev_impl) +pybind11_add_module(${python_module_name} MODULE + rng_py.cpp + gaussian.cpp + uniform.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}) +target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/engine) +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/device" +) diff --git a/dpnp/backend/extensions/rng/device/common_impl.hpp b/dpnp/backend/extensions/rng/device/common_impl.hpp new file mode 100644 index 00000000000..2b4744c251f --- /dev/null +++ b/dpnp/backend/extensions/rng/device/common_impl.hpp @@ -0,0 +1,122 @@ +//***************************************************************************** +// Copyright (c) 2024, 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. +//***************************************************************************** + +#pragma once + +#include + +#include +#include + +namespace dpnp::backend::ext::rng::device::details +{ +namespace py = pybind11; + +namespace mkl_rng_dev = oneapi::mkl::rng::device; + +template +struct RngContigFunctor +{ +private: + using DataT = typename DistributorBuilderT::result_type; + + EngineBuilderT engine_; + DistributorBuilderT distr_; + DataT *const res_ = nullptr; + const std::size_t nelems_; + +public: + RngContigFunctor(EngineBuilderT &engine, + DistributorBuilderT &distr, + DataT *res, + const std::size_t n_elems) + : engine_(engine), distr_(distr), res_(res), nelems_(n_elems) + { + } + + void operator()(sycl::nd_item<1> nd_it) const + { + auto sg = nd_it.get_sub_group(); + const std::uint8_t sg_size = sg.get_local_range()[0]; + const std::uint8_t max_sg_size = sg.get_max_local_range()[0]; + + using EngineT = typename EngineBuilderT::EngineType; + using DistrT = typename DistributorBuilderT::distr_type; + + constexpr std::size_t vec_sz = EngineT::vec_size; + constexpr std::size_t vi_per_wi = vec_sz * items_per_wi; + + EngineT engine = engine_(nd_it.get_global_id() * vi_per_wi); + 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); + + 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< + 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); + } + } + } + else { + 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) + { + res_[offset] = mkl_rng_dev::generate_single( + distr, engine); + } + } + } +}; +} // namespace dpnp::backend::ext::rng::device::details diff --git a/dpnp/backend/extensions/rng/device/dispatch/matrix.hpp b/dpnp/backend/extensions/rng/device/dispatch/matrix.hpp new file mode 100644 index 00000000000..6cbd4253858 --- /dev/null +++ b/dpnp/backend/extensions/rng/device/dispatch/matrix.hpp @@ -0,0 +1,84 @@ +//***************************************************************************** +// Copyright (c) 2024, 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. +//***************************************************************************** + +#pragma once + +#include + +#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> +{ + static constexpr bool is_defined = true; +}; + +template +struct GaussianTypePairSupportFactory +{ + static constexpr bool is_defined = std::disjunction< + TypePairDefinedEntry, + TypePairDefinedEntry, + // fall-through + dpctl_td_ns::NotDefinedEntry>::is_defined; +}; + +template +struct UniformTypePairSupportFactory +{ + static constexpr bool is_defined = std::disjunction< + TypePairDefinedEntry, + TypePairDefinedEntry, + TypePairDefinedEntry, + TypePairDefinedEntry, + // fall-through + dpctl_td_ns::NotDefinedEntry>::is_defined; +}; +} // 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 new file mode 100644 index 00000000000..44044b36af2 --- /dev/null +++ b/dpnp/backend/extensions/rng/device/dispatch/table_builder.hpp @@ -0,0 +1,105 @@ +//***************************************************************************** +// Copyright (c) 2024, 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. +//***************************************************************************** + +#pragma once + +#include + +namespace dpnp::backend::ext::rng::device::dispatch +{ +namespace mkl_rng_dev = oneapi::mkl::rng::device; + +template + typename factory, + int _no_of_engines, + int _no_of_types, + int _no_of_methods> +class Dispatch3DTableBuilder +{ +private: + template + const std::vector row_per_method() const + { + std::vector per_method = { + factory{}.get()..., + }; + assert(per_method.size() == _no_of_methods); + return per_method; + } + + 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, Methods...>(), + row_per_method, Methods...>()}; + assert(table_by_type.size() == _no_of_types); + return table_by_type; + } + +public: + Dispatch3DTableBuilder() = default; + ~Dispatch3DTableBuilder() = default; + + template + void populate(funcPtrT table[][_no_of_types][_no_of_methods], + std::integer_sequence) const + { + const auto map_by_engine = { + table_per_type_and_method, Methods...>()..., + table_per_type_and_method, Methods...>()..., + table_per_type_and_method, Methods...>()..., + table_per_type_and_method, Methods...>()...}; + assert(map_by_engine.size() == _no_of_engines); + + std::uint16_t engine_id = 0; + for (auto &table_by_type : map_by_engine) { + std::uint16_t type_id = 0; + for (auto &row_by_method : table_by_type) { + std::uint16_t method_id = 0; + for (auto &fn_ptr : row_by_method) { + table[engine_id][type_id][method_id] = fn_ptr; + ++method_id; + } + ++type_id; + } + ++engine_id; + } + } +}; +} // 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 new file mode 100644 index 00000000000..d6f49595c06 --- /dev/null +++ b/dpnp/backend/extensions/rng/device/engine/base_engine.hpp @@ -0,0 +1,159 @@ +//***************************************************************************** +// Copyright (c) 2024, 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. +//***************************************************************************** + +#pragma once + +#include + +namespace dpnp::backend::ext::rng::device::engine +{ +class EngineType +{ +public: + enum Type : std::uint8_t + { + MRG32k3a = 0, + PHILOX4x32x10, + MCG31M1, + MCG59, + Base, // must be the last always + }; + + EngineType() = default; + constexpr EngineType(Type type) : type_(type) {} + + constexpr std::uint8_t id() const + { + return static_cast(type_); + } + + static constexpr std::uint8_t base_id() + { + return EngineType(Base).id(); + } + +private: + Type type_; +}; + +// A total number of supported engines == EngineType::Base +constexpr std::uint8_t no_of_engines = EngineType::base_id(); + +class EngineBase +{ +private: + sycl::queue q_{}; + std::vector seed_vec{}; + std::vector offset_vec{}; + + void validate_vec_size(const std::size_t size) + { + if (size > max_vec_n) { + throw std::runtime_error("TODO: add text"); + } + } + +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::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) + { + 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::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) + { + validate_vec_size(seeds.size()); + validate_vec_size(offsets.size()); + + seed_vec.reserve(seeds.size()); + seed_vec.assign(seeds.begin(), seeds.end()); + } + + virtual ~EngineBase() {} + + virtual EngineType get_type() const noexcept + { + return EngineType::Base; + } + + sycl::queue &get_queue() noexcept + { + return q_; + } + + std::vector &get_seeds() noexcept + { + return seed_vec; + } + + std::vector &get_offsets() noexcept + { + return offset_vec; + } + + // + static constexpr std::uint8_t max_vec_n = 1; +}; +} // 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 new file mode 100644 index 00000000000..b6f0fea3ffd --- /dev/null +++ b/dpnp/backend/extensions/rng/device/engine/builder/base_builder.hpp @@ -0,0 +1,130 @@ +//***************************************************************************** +// Copyright (c) 2024, 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. +//***************************************************************************** + +#pragma once + +#include "base_engine.hpp" + +// TODO: remove the include once issue in MKL is resolved +#include +namespace mkl_rng_dev = oneapi::mkl::rng::device; + +namespace dpnp::backend::ext::rng::device::engine::builder +{ +template +class BaseBuilder +{ +private: + static constexpr std::uint8_t max_n = EngineBase::max_vec_n; + + std::uint8_t no_of_seeds; + std::uint8_t no_of_offsets; + + std::array seeds{}; + std::array offsets{}; + +public: + BaseBuilder(EngineBase *engine) + { + auto seed_values = engine->get_seeds(); + no_of_seeds = seed_values.size(); + if (no_of_seeds > max_n) { + throw std::runtime_error(""); + } + + // TODO: implement a caster + for (std::uint16_t i = 0; i < no_of_seeds; i++) { + seeds[i] = static_cast(seed_values[i]); + } + + auto offset_values = engine->get_offsets(); + no_of_offsets = offset_values.size(); + if (no_of_offsets > max_n) { + throw std::runtime_error(""); + } + + // TODO: implement a caster + for (std::uint16_t i = 0; i < no_of_seeds; i++) { + offsets[i] = static_cast(offset_values[i]); + } + } + + 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]); + } + } + // TODO: implement full switch + default: + break; + } + return EngineT(); + } + + 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}); + } + } + // TODO: implement full switch + default: + break; + } + return EngineT(); + } + + // TODO: remove + void print() + { + std::cout << "vector size = " << std::to_string(EngineT::vec_size) + << std::endl; + std::cout << "list_of_seeds: "; + for (auto &val : seeds) { + std::cout << std::to_string(val) << ", "; + } + std::cout << std::endl; + } +}; +} // 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 new file mode 100644 index 00000000000..1fd9c9c9e89 --- /dev/null +++ b/dpnp/backend/extensions/rng/device/engine/builder/builder.hpp @@ -0,0 +1,39 @@ +//***************************************************************************** +// Copyright (c) 2024, 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. +//***************************************************************************** + +#pragma once + +namespace dpnp::backend::ext::rng::device::engine::builder +{ +template +class Builder +{ +}; +} // namespace dpnp::backend::ext::rng::device::engine::builder + +#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 new file mode 100644 index 00000000000..b94b93b35b3 --- /dev/null +++ b/dpnp/backend/extensions/rng/device/engine/builder/mcg31m1.hpp @@ -0,0 +1,51 @@ +//***************************************************************************** +// Copyright (c) 2024, 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. +//***************************************************************************** + +#pragma once + +#include + +#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> +{ +public: + using EngineType = mkl_rng_dev::mcg31m1; + + Builder(EngineBase *engine) + : BaseBuilder(engine) + { + } +}; +} // 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 new file mode 100644 index 00000000000..7f8e003b87c --- /dev/null +++ b/dpnp/backend/extensions/rng/device/engine/builder/mcg59.hpp @@ -0,0 +1,51 @@ +//***************************************************************************** +// Copyright (c) 2024, 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. +//***************************************************************************** + +#pragma once + +#include + +#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> +{ +public: + using EngineType = mkl_rng_dev::mcg59; + + Builder(EngineBase *engine) + : BaseBuilder(engine) + { + } +}; +} // 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 new file mode 100644 index 00000000000..a5ab3470b03 --- /dev/null +++ b/dpnp/backend/extensions/rng/device/engine/builder/mrg32k3a.hpp @@ -0,0 +1,51 @@ +//***************************************************************************** +// Copyright (c) 2024, 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. +//***************************************************************************** + +#pragma once + +#include + +#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> +{ +public: + using EngineType = mkl_rng_dev::mrg32k3a; + + Builder(EngineBase *engine) + : BaseBuilder(engine) + { + } +}; +} // 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 new file mode 100644 index 00000000000..932c5ee00d7 --- /dev/null +++ b/dpnp/backend/extensions/rng/device/engine/builder/philox4x32x10.hpp @@ -0,0 +1,51 @@ +//***************************************************************************** +// Copyright (c) 2024, 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. +//***************************************************************************** + +#pragma once + +#include + +#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> +{ +public: + using EngineType = mkl_rng_dev::philox4x32x10; + + Builder(EngineBase *engine) + : BaseBuilder(engine) + { + } +}; +} // 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 new file mode 100644 index 00000000000..4022bd33c18 --- /dev/null +++ b/dpnp/backend/extensions/rng/device/engine/mcg31m1_engine.hpp @@ -0,0 +1,52 @@ +//***************************************************************************** +// Copyright (c) 2024, 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. +//***************************************************************************** + +#pragma once + +#include "base_engine.hpp" + +namespace dpnp::backend::ext::rng::device::engine +{ +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::vector &seeds, + std::uint64_t offset = 0) + : EngineBase(q, seeds, offset) + { + } + + virtual EngineType get_type() const noexcept override + { + return EngineType::MCG31M1; + } +}; +} // 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 new file mode 100644 index 00000000000..e0fbf1741a3 --- /dev/null +++ b/dpnp/backend/extensions/rng/device/engine/mcg59_engine.hpp @@ -0,0 +1,52 @@ +//***************************************************************************** +// Copyright (c) 2024, 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. +//***************************************************************************** + +#pragma once + +#include "base_engine.hpp" + +namespace dpnp::backend::ext::rng::device::engine +{ +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::vector &seeds, + std::uint64_t offset = 0) + : EngineBase(q, seeds, offset) + { + } + + virtual EngineType get_type() const noexcept override + { + return EngineType::MCG59; + } +}; +} // 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 new file mode 100644 index 00000000000..d6657d6d8d8 --- /dev/null +++ b/dpnp/backend/extensions/rng/device/engine/mrg32k3a_engine.hpp @@ -0,0 +1,66 @@ +//***************************************************************************** +// Copyright (c) 2024, 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. +//***************************************************************************** + +#pragma once + +#include "base_engine.hpp" + +namespace dpnp::backend::ext::rng::device::engine +{ +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::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::vector &seeds, + std::vector &offsets) + : EngineBase(q, seeds, offsets) + { + } + + virtual EngineType get_type() const noexcept override + { + return EngineType::MRG32k3a; + } +}; +} // 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 new file mode 100644 index 00000000000..bef1e7d2119 --- /dev/null +++ b/dpnp/backend/extensions/rng/device/engine/philox4x32x10_engine.hpp @@ -0,0 +1,66 @@ +//***************************************************************************** +// Copyright (c) 2024, 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. +//***************************************************************************** + +#pragma once + +#include "base_engine.hpp" + +namespace dpnp::backend::ext::rng::device::engine +{ +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::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::vector &seeds, + std::vector &offsets) + : EngineBase(q, seeds, offsets) + { + } + + virtual EngineType get_type() const noexcept override + { + return EngineType::PHILOX4x32x10; + } +}; +} // 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 new file mode 100644 index 00000000000..643c4091fee --- /dev/null +++ b/dpnp/backend/extensions/rng/device/gaussian.cpp @@ -0,0 +1,296 @@ +//***************************************************************************** +// 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 + +// dpctl tensor headers +#include "kernels/alignment.hpp" +#include "utils/output_validation.hpp" +#include "utils/type_dispatch.hpp" +#include "utils/type_utils.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; +namespace dpctl_td_ns = dpctl::tensor::type_dispatch; +namespace dpctl_tu_ns = dpctl::tensor::type_utils; +namespace mkl_rng_dev = oneapi::mkl::rng::device; +namespace py = pybind11; + +using dpctl_krn_ns::disabled_sg_loadstore_wrapper_krn; +using dpctl_krn_ns::is_aligned; +using dpctl_krn_ns::required_alignment; + +constexpr auto no_of_methods = 1; // number of methods of gaussian distribution + +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 +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; +} + +template +struct DistributorBuilder +{ +private: + const DataT mean_; + const DataT stddev_; + +public: + using result_type = DataT; + using method_type = Method; + using distr_type = typename mkl_rng_dev::gaussian; + + DistributorBuilder(const DataT mean, const DataT stddev) + : mean_(mean), stddev_(stddev) + { + } + + inline auto operator()(void) const + { + return distr_type(mean_, stddev_); + } +}; + +typedef sycl::event (*gaussian_impl_fn_ptr_t)(engine::EngineBase *engine, + const double, + const double, + const std::uint64_t, + char *, + const std::vector &); + +static gaussian_impl_fn_ptr_t gaussian_dispatch_table[no_of_engines] + [dpctl_td_ns::num_types] + [no_of_methods]; + +template +class gaussian_kernel; + +template +static sycl::event gaussian_impl(engine::EngineBase *engine, + const double mean_val, + const double stddev_val, + const std::uint64_t n, + char *out_ptr, + const std::vector &depends) +{ + auto &exec_q = engine->get_queue(); + dpctl_tu_ns::validate_type_for_device(exec_q); + + DataT *out = reinterpret_cast(out_ptr); + DataT mean = static_cast(mean_val); + DataT stddev = static_cast(stddev_val); + + constexpr std::size_t vec_sz = EngineT::vec_size; + 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; + + sycl::event distr_event; + + try { + distr_event = exec_q.submit([&](sycl::handler &cgh) { + cgh.depends_on(depends); + + using EngineBuilderT = engine::builder::Builder; + EngineBuilderT eng_builder(engine); + // eng_builder.print(); // TODO: remove + + using DistributorBuilderT = DistributorBuilder; + DistributorBuilderT dist_builder(mean, stddev); + + 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)); + } + 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)); + } + }); + } catch (oneapi::mkl::exception const &e) { + std::stringstream error_msg; + + 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(); + 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) +{ + auto &exec_q = engine->get_queue(); + + const int res_nd = res.get_ndim(); + const py::ssize_t *res_shape = res.get_shape_raw(); + + size_t res_nelems(1); + for (int i = 0; i < res_nd; ++i) { + res_nelems *= static_cast(res_shape[i]); + } + + if (res_nelems == 0) { + // nothing to do + return std::make_pair(sycl::event(), sycl::event()); + } + + // ensure that output is ample enough to accommodate all elements + 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"); + } + + 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."); + } + + 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."); + } + + if (method_id >= no_of_methods) { + 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"); + } + 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]; + if (gaussian_fn == nullptr) { + 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 ht_ev = + dpctl::utils::keep_args_alive(exec_q, {res}, {gaussian_ev}); + return std::make_pair(ht_ev, gaussian_ev); +} + +template +struct GaussianContigFactory +{ + fnT get() + { + if constexpr (dispatch::GaussianTypePairSupportFactory::is_defined) { + return gaussian_impl; + } + else { + return nullptr; + } + } +}; + +void init_gaussian_dispatch_3d_table(void) +{ + dispatch::Dispatch3DTableBuilder + contig; + contig.populate(gaussian_dispatch_table, seq_of_vec_sizes); +} +} // namespace dpnp::backend::ext::rng::device diff --git a/dpnp/backend/extensions/rng/device/gaussian.hpp b/dpnp/backend/extensions/rng/device/gaussian.hpp new file mode 100644 index 00000000000..00973a5d4e5 --- /dev/null +++ b/dpnp/backend/extensions/rng/device/gaussian.hpp @@ -0,0 +1,45 @@ +//***************************************************************************** +// Copyright (c) 2024, 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. +//***************************************************************************** + +#pragma once + +#include + +#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 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 new file mode 100644 index 00000000000..27cbdb80d8d --- /dev/null +++ b/dpnp/backend/extensions/rng/device/rng_py.cpp @@ -0,0 +1,131 @@ +//***************************************************************************** +// 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 "gaussian.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 +{ +public: + // inherit the constructor + using EngineBase::EngineBase; + + // trampoline (need one for each virtual function) + // sycl::queue &get_queue() { + // PYBIND11_OVERRIDE_PURE( + // sycl::queue&, /* Return type */ + // EngineBase, /* Parent class */ + // 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") + .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, "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"), + py::arg("depends") = py::list()); +} diff --git a/dpnp/backend/extensions/rng/device/uniform.cpp b/dpnp/backend/extensions/rng/device/uniform.cpp new file mode 100644 index 00000000000..62ef69a29ea --- /dev/null +++ b/dpnp/backend/extensions/rng/device/uniform.cpp @@ -0,0 +1,295 @@ +//***************************************************************************** +// 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 + +// dpctl tensor headers +#include "kernels/alignment.hpp" +#include "utils/output_validation.hpp" +#include "utils/type_dispatch.hpp" +#include "utils/type_utils.hpp" + +#include "common_impl.hpp" +#include "uniform.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; +namespace dpctl_td_ns = dpctl::tensor::type_dispatch; +namespace dpctl_tu_ns = dpctl::tensor::type_utils; +namespace mkl_rng_dev = oneapi::mkl::rng::device; +namespace py = pybind11; + +using dpctl_krn_ns::disabled_sg_loadstore_wrapper_krn; +using dpctl_krn_ns::is_aligned; +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 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 +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; +} + +template +struct DistributorBuilder +{ +private: + const DataT mean_; + const DataT stddev_; + +public: + using result_type = DataT; + using method_type = Method; + using distr_type = typename mkl_rng_dev::uniform; + + DistributorBuilder(const DataT mean, const DataT stddev) + : mean_(mean), stddev_(stddev) + { + } + + inline auto operator()(void) const + { + return distr_type(mean_, stddev_); + } +}; + +typedef sycl::event (*uniform_impl_fn_ptr_t)(engine::EngineBase *engine, + const double, + const double, + const std::uint64_t, + char *, + const std::vector &); + +static uniform_impl_fn_ptr_t uniform_dispatch_table[no_of_engines] + [dpctl_td_ns::num_types] + [no_of_methods]; + +template +class uniform_kernel; + +template +static sycl::event uniform_impl(engine::EngineBase *engine, + const double a_val, + const double b_val, + const std::uint64_t n, + char *out_ptr, + const std::vector &depends) +{ + auto &exec_q = engine->get_queue(); + dpctl_tu_ns::validate_type_for_device(exec_q); + + DataT *out = reinterpret_cast(out_ptr); + DataT a = static_cast(a_val); + DataT b = static_cast(b_val); + + constexpr std::size_t vec_sz = EngineT::vec_size; + 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; + + sycl::event distr_event; + + try { + distr_event = exec_q.submit([&](sycl::handler &cgh) { + cgh.depends_on(depends); + + using EngineBuilderT = engine::builder::Builder; + EngineBuilderT eng_builder(engine); + // eng_builder.print(); // TODO: remove + + using DistributorBuilderT = DistributorBuilder; + DistributorBuilderT dist_builder(a, b); + + if (is_aligned(out_ptr)) { + constexpr bool enable_sg_load = true; + using KernelName = + uniform_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 = + uniform_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(); + 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(); + throw std::runtime_error(error_msg.str()); + } + return distr_event; +} + +std::pair + uniform(engine::EngineBase *engine, + const std::uint8_t method_id, + const std::uint8_t vec_size, + const double a, + const double b, + const std::uint64_t n, + dpctl::tensor::usm_ndarray res, + const std::vector &depends) +{ + auto &exec_q = engine->get_queue(); + + const int res_nd = res.get_ndim(); + const py::ssize_t *res_shape = res.get_shape_raw(); + + size_t res_nelems(1); + for (int i = 0; i < res_nd; ++i) { + res_nelems *= static_cast(res_shape[i]); + } + + if (res_nelems == 0) { + // nothing to do + return std::make_pair(sycl::event(), sycl::event()); + } + + // ensure that output is ample enough to accommodate all elements + 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"); + } + + 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."); + } + + 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."); + } + + if (method_id >= no_of_methods) { + 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"); + } + 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 uniform_fn = + uniform_dispatch_table[enginge_id][res_type_id][method_id]; + if (uniform_fn == nullptr) { + throw py::value_error( + "No gaussian implementation defined for a required type"); + } + + char *res_data = res.get_data(); + sycl::event uniform_ev = + uniform_fn(engine, a, b, n, res_data, depends); + + sycl::event ht_ev = + dpctl::utils::keep_args_alive(exec_q, {res}, {uniform_ev}); + return std::make_pair(ht_ev, uniform_ev); +} + +template +struct UniformContigFactory +{ + fnT get() + { + if constexpr (dispatch::UniformTypePairSupportFactory::is_defined) { + return uniform_impl; + } + else { + return nullptr; + } + } +}; + +void init_uniform_dispatch_3d_table(void) +{ + dispatch::Dispatch3DTableBuilder + contig; + contig.populate(uniform_dispatch_table, seq_of_vec_sizes); +} +} // namespace dpnp::backend::ext::rng::device diff --git a/dpnp/backend/extensions/rng/device/uniform.hpp b/dpnp/backend/extensions/rng/device/uniform.hpp new file mode 100644 index 00000000000..e45bc45f31d --- /dev/null +++ b/dpnp/backend/extensions/rng/device/uniform.hpp @@ -0,0 +1,45 @@ +//***************************************************************************** +// Copyright (c) 2024, 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. +//***************************************************************************** + +#pragma once + +#include + +#include "engine/base_engine.hpp" + +namespace dpnp::backend::ext::rng::device +{ +extern std::pair + uniform(engine::EngineBase *engine, + const std::uint8_t method_id, + const std::uint8_t vec_size, + const double a, + const double b, + const std::uint64_t n, + dpctl::tensor::usm_ndarray res, + const std::vector &depends = {}); + +extern void init_uniform_dispatch_3d_table(void); +} // namespace dpnp::backend::ext::rng::device diff --git a/dpnp/backend/extensions/rng/host/CMakeLists.txt b/dpnp/backend/extensions/rng/host/CMakeLists.txt new file mode 100644 index 00000000000..f3db23c2a23 --- /dev/null +++ b/dpnp/backend/extensions/rng/host/CMakeLists.txt @@ -0,0 +1,77 @@ +# ***************************************************************************** +# 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_host_impl) +pybind11_add_module(${python_module_name} MODULE + rng_py.cpp + gaussian.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}) +target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/engine) +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/host" +) diff --git a/dpnp/backend/extensions/rng/host/dispatch/matrix.hpp b/dpnp/backend/extensions/rng/host/dispatch/matrix.hpp new file mode 100644 index 00000000000..eb7aa43450f --- /dev/null +++ b/dpnp/backend/extensions/rng/host/dispatch/matrix.hpp @@ -0,0 +1,65 @@ +//***************************************************************************** +// Copyright (c) 2024, 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. +//***************************************************************************** + +#pragma once + +#include + +#include "utils/type_dispatch.hpp" + +namespace dpnp::backend::ext::rng::host::dispatch +{ +namespace dpctl_td_ns = dpctl::tensor::type_dispatch; +namespace mkl_rng = oneapi::mkl::rng; + +template +struct TypePairDefinedEntry + : std::bool_constant && + std::is_same_v> +{ + static constexpr bool is_defined = true; +}; + +template +struct GaussianTypePairSupportFactory +{ + static constexpr bool is_defined = std::disjunction< + TypePairDefinedEntry, + TypePairDefinedEntry, + TypePairDefinedEntry, + TypePairDefinedEntry, + // fall-through + dpctl_td_ns::NotDefinedEntry>::is_defined; +}; +} // namespace dpnp::backend::ext::rng::host::dispatch diff --git a/dpnp/backend/extensions/rng/host/dispatch/table_builder.hpp b/dpnp/backend/extensions/rng/host/dispatch/table_builder.hpp new file mode 100644 index 00000000000..772fa8fc9f3 --- /dev/null +++ b/dpnp/backend/extensions/rng/host/dispatch/table_builder.hpp @@ -0,0 +1,106 @@ +//***************************************************************************** +// Copyright (c) 2024, 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. +//***************************************************************************** + +#pragma once + +#include + +namespace dpnp::backend::ext::rng::host::dispatch +{ +namespace mkl_rng = oneapi::mkl::rng; + +template + typename factory, + int _no_of_engines, + int _no_of_types, + int _no_of_methods> +class Dispatch3DTableBuilder +{ +private: + template + const std::vector row_per_method() const + { + std::vector per_method = { + factory{} + .get(), + factory{} + .get(), + }; + assert(per_method.size() == _no_of_methods); + return per_method; + } + + 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>()}; + assert(table_by_type.size() == _no_of_types); + return table_by_type; + } + +public: + Dispatch3DTableBuilder() = default; + ~Dispatch3DTableBuilder() = default; + + void populate(funcPtrT table[][_no_of_types][_no_of_methods]) 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()}; + assert(map_by_engine.size() == _no_of_engines); + + std::uint16_t engine_id = 0; + for (auto &table_by_type : map_by_engine) { + std::uint16_t type_id = 0; + for (auto &row_by_method : table_by_type) { + std::uint16_t method_id = 0; + for (auto &fn_ptr : row_by_method) { + table[engine_id][type_id][method_id] = fn_ptr; + ++method_id; + } + ++type_id; + } + ++engine_id; + } + } +}; +} // namespace dpnp::backend::ext::rng::host::dispatch diff --git a/dpnp/backend/extensions/rng/host/gaussian.cpp b/dpnp/backend/extensions/rng/host/gaussian.cpp new file mode 100644 index 00000000000..d6b4a238702 --- /dev/null +++ b/dpnp/backend/extensions/rng/host/gaussian.cpp @@ -0,0 +1,181 @@ +//***************************************************************************** +// 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 + +// dpctl tensor headers +#include "utils/output_validation.hpp" +#include "utils/type_dispatch.hpp" +#include "utils/type_utils.hpp" + +#include "gaussian.hpp" + +#include "dispatch/matrix.hpp" +#include "dispatch/table_builder.hpp" + +namespace dpnp::backend::ext::rng::host +{ +namespace dpctl_td_ns = dpctl::tensor::type_dispatch; +namespace dpctl_tu_ns = dpctl::tensor::type_utils; +namespace mkl_rng = oneapi::mkl::rng; +namespace py = pybind11; + +constexpr auto no_of_methods = 2; // number of methods of gaussian distribution +constexpr auto no_of_engines = device::engine::no_of_engines; + +typedef sycl::event (*gaussian_impl_fn_ptr_t)( + device::engine::EngineBase *engine, + const double, + const double, + const std::uint64_t, + char *, + const std::vector &); + +static gaussian_impl_fn_ptr_t gaussian_dispatch_table[no_of_engines] + [dpctl_td_ns::num_types] + [no_of_methods]; + +template +static sycl::event gaussian_impl(device::engine::EngineBase *engine, + const double mean_val, + const double stddev_val, + const std::uint64_t n, + char *out_ptr, + const std::vector &depends) +{ + auto &exec_q = engine->get_queue(); + dpctl_tu_ns::validate_type_for_device(exec_q); + + DataT *out = reinterpret_cast(out_ptr); + DataT mean = static_cast(mean_val); + DataT stddev = static_cast(stddev_val); + + auto seed_values = engine->get_seeds(); + auto no_of_seeds = seed_values.size(); + if (no_of_seeds > 1) { + throw std::runtime_error(""); + } + + mkl_rng::gaussian distribution(mean, stddev); + mkl_rng::mcg59 eng(exec_q, seed_values[0]); + + return mkl_rng::generate(distribution, eng, n, out, depends); +} + +std::pair + gaussian(device::engine::EngineBase *engine, + const std::uint8_t method_id, + 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(); + + const int res_nd = res.get_ndim(); + const py::ssize_t *res_shape = res.get_shape_raw(); + + size_t res_nelems(1); + for (int i = 0; i < res_nd; ++i) { + res_nelems *= static_cast(res_shape[i]); + } + + if (res_nelems == 0) { + // nothing to do + return std::make_pair(sycl::event(), sycl::event()); + } + + // ensure that output is ample enough to accommodate all elements + 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"); + } + + 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."); + } + + auto enginge_id = engine->get_type().id(); + if (enginge_id >= device::engine::no_of_engines) { + 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."); + } + + 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]; + if (gaussian_fn == nullptr) { + 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 ht_ev = + dpctl::utils::keep_args_alive(exec_q, {res}, {gaussian_ev}); + return std::make_pair(ht_ev, gaussian_ev); +} + +template +struct GaussianContigFactory +{ + fnT get() + { + if constexpr (dispatch::GaussianTypePairSupportFactory::is_defined) { + return gaussian_impl; + } + else { + return nullptr; + } + } +}; + +void init_gaussian_dispatch_3d_table(void) +{ + dispatch::Dispatch3DTableBuilder + contig; + contig.populate(gaussian_dispatch_table); +} +} // namespace dpnp::backend::ext::rng::host diff --git a/dpnp/backend/extensions/rng/host/gaussian.hpp b/dpnp/backend/extensions/rng/host/gaussian.hpp new file mode 100644 index 00000000000..2ebf5f976e2 --- /dev/null +++ b/dpnp/backend/extensions/rng/host/gaussian.hpp @@ -0,0 +1,44 @@ +//***************************************************************************** +// Copyright (c) 2024, 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. +//***************************************************************************** + +#pragma once + +#include + +#include "../device/engine/base_engine.hpp" + +namespace dpnp::backend::ext::rng::host +{ +extern std::pair + gaussian(device::engine::EngineBase *engine, + const std::uint8_t method_id, + 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::host diff --git a/dpnp/backend/extensions/rng/host/rng_py.cpp b/dpnp/backend/extensions/rng/host/rng_py.cpp new file mode 100644 index 00000000000..d76b07242e7 --- /dev/null +++ b/dpnp/backend/extensions/rng/host/rng_py.cpp @@ -0,0 +1,142 @@ +//***************************************************************************** +// 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 "gaussian.hpp" + +// #include "../device/engine/mcg31m1_engine.hpp" +// #include "../device/engine/mcg59_engine.hpp" +// #include "../device/engine/mrg32k3a_engine.hpp" +// #include "../device/engine/philox4x32x10_engine.hpp" + +namespace mkl_rng = oneapi::mkl::rng; +namespace rng_host_ext = dpnp::backend::ext::rng::host; +// 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_host_ext::init_gaussian_dispatch_3d_table(); +} + +// class PyEngineBase : public rng_dev_engine::EngineBase +// { +// public: +// // inherit the constructor +// using EngineBase::EngineBase; + +// // trampoline (need one for each virtual function) +// // sycl::queue &get_queue() { +// // PYBIND11_OVERRIDE_PURE( +// // sycl::queue&, /* Return type */ +// // EngineBase, /* Parent class */ +// // get_queue, /* Name of function in C++ (must match Python +// name) +// // */ +// // ); +// // } +// }; + +PYBIND11_MODULE(_rng_host_impl, m) +{ + init_dispatch_3d_tables(); + + // 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, + // "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_host_ext::gaussian, "", py::arg("engine"), + py::arg("method_id"), py::arg("mean"), py::arg("stddev"), + py::arg("n"), py::arg("res"), py::arg("depends") = py::list()); +}