diff --git a/dpnp/CMakeLists.txt b/dpnp/CMakeLists.txt
index d9c95b62c0be..c4a3889fda3a 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 000000000000..37df343f40b9
--- /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 000000000000..2b4744c251fd
--- /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 <pybind11/pybind11.h>
+
+#include <oneapi/mkl/rng/device.hpp>
+#include <sycl/sycl.hpp>
+
+namespace dpnp::backend::ext::rng::device::details
+{
+namespace py = pybind11;
+
+namespace mkl_rng_dev = oneapi::mkl::rng::device;
+
+template <typename EngineBuilderT,
+          typename DistributorBuilderT,
+          unsigned int items_per_wi = 4,
+          bool enable_sg_load = true>
+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<std::size_t>(it) *
+                                   static_cast<std::size_t>(sg_size);
+                    auto out_multi_ptr = sycl::address_space_cast<
+                        sycl::access::address_space::global_space,
+                        sycl::access::decorated::yes>(&res_[offset]);
+
+                    sycl::vec<DataT, vec_sz> rng_val_vec =
+                        mkl_rng_dev::generate<DistrT, EngineT>(distr, engine);
+                    sg.store<vec_sz>(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<DistrT, EngineT>(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<DistrT, EngineT>(
+                    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 000000000000..6cbd42538584
--- /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 <oneapi/mkl/rng/device.hpp>
+
+#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 <typename Ty, typename ArgTy, typename Method, typename argMethod>
+struct TypePairDefinedEntry
+    : std::bool_constant<std::is_same_v<Ty, ArgTy> &&
+                         std::is_same_v<Method, argMethod>>
+{
+    static constexpr bool is_defined = true;
+};
+
+template <typename T, typename M>
+struct GaussianTypePairSupportFactory
+{
+    static constexpr bool is_defined = std::disjunction<
+        TypePairDefinedEntry<T,
+                             double,
+                             M,
+                             mkl_rng_dev::gaussian_method::box_muller2>,
+        TypePairDefinedEntry<T,
+                             float,
+                             M,
+                             mkl_rng_dev::gaussian_method::box_muller2>,
+        // fall-through
+        dpctl_td_ns::NotDefinedEntry>::is_defined;
+};
+
+template <typename T, typename M>
+struct UniformTypePairSupportFactory
+{
+    static constexpr bool is_defined = std::disjunction<
+        TypePairDefinedEntry<T,
+                             double,
+                             M,
+                             mkl_rng_dev::uniform_method::standard>,
+        TypePairDefinedEntry<T,
+                             double,
+                             M,
+                             mkl_rng_dev::uniform_method::accurate>,
+        TypePairDefinedEntry<T,
+                             float,
+                             M,
+                             mkl_rng_dev::uniform_method::standard>,
+        TypePairDefinedEntry<T,
+                             float,
+                             M,
+                             mkl_rng_dev::uniform_method::accurate>,
+        // 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 000000000000..44044b36af24
--- /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 <oneapi/mkl/rng/device.hpp>
+
+namespace dpnp::backend::ext::rng::device::dispatch
+{
+namespace mkl_rng_dev = oneapi::mkl::rng::device;
+
+template <typename funcPtrT,
+          template <typename fnT, typename E, typename T, typename M>
+          typename factory,
+          int _no_of_engines,
+          int _no_of_types,
+          int _no_of_methods>
+class Dispatch3DTableBuilder
+{
+private:
+    template <typename E, typename T, typename... Methods>
+    const std::vector<funcPtrT> row_per_method() const
+    {
+        std::vector<funcPtrT> per_method = {
+            factory<funcPtrT, E, T, Methods>{}.get()...,
+        };
+        assert(per_method.size() == _no_of_methods);
+        return per_method;
+    }
+
+    template <typename E, typename... Methods>
+    auto table_per_type_and_method() const
+    {
+        std::vector<std::vector<funcPtrT>> table_by_type = {
+            row_per_method<E, bool, Methods...>(),
+            row_per_method<E, int8_t, Methods...>(),
+            row_per_method<E, uint8_t, Methods...>(),
+            row_per_method<E, int16_t, Methods...>(),
+            row_per_method<E, uint16_t, Methods...>(),
+            row_per_method<E, int32_t, Methods...>(),
+            row_per_method<E, uint32_t, Methods...>(),
+            row_per_method<E, int64_t, Methods...>(),
+            row_per_method<E, uint64_t, Methods...>(),
+            row_per_method<E, sycl::half, Methods...>(),
+            row_per_method<E, float, Methods...>(),
+            row_per_method<E, double, Methods...>(),
+            row_per_method<E, std::complex<float>, Methods...>(),
+            row_per_method<E, std::complex<double>, Methods...>()};
+        assert(table_by_type.size() == _no_of_types);
+        return table_by_type;
+    }
+
+public:
+    Dispatch3DTableBuilder() = default;
+    ~Dispatch3DTableBuilder() = default;
+
+    template <typename... Methods, std::uint8_t... VecSizes>
+    void populate(funcPtrT table[][_no_of_types][_no_of_methods],
+                  std::integer_sequence<std::uint8_t, VecSizes...>) const
+    {
+        const auto map_by_engine = {
+            table_per_type_and_method<mkl_rng_dev::mrg32k3a<VecSizes>, Methods...>()...,
+            table_per_type_and_method<mkl_rng_dev::philox4x32x10<VecSizes>, Methods...>()...,
+            table_per_type_and_method<mkl_rng_dev::mcg31m1<VecSizes>, Methods...>()...,
+            table_per_type_and_method<mkl_rng_dev::mcg59<VecSizes>, 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 000000000000..d6f49595c06b
--- /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 <sycl/sycl.hpp>
+
+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<std::uint8_t>(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<std::uint64_t> seed_vec{};
+    std::vector<std::uint64_t> 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<std::uint64_t> &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<std::uint32_t> &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<std::uint64_t> &offsets)
+        : q_(q), seed_vec(1, seed), offset_vec(offsets)
+    {
+        validate_vec_size(offsets.size());
+    }
+
+    EngineBase(sycl::queue &q,
+               std::vector<std::uint64_t> &seeds,
+               std::vector<std::uint64_t> &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<std::uint32_t> &seeds,
+               std::vector<std::uint64_t> &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<std::uint64_t> &get_seeds() noexcept
+    {
+        return seed_vec;
+    }
+
+    std::vector<std::uint64_t> &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 000000000000..b6f0fea3ffd4
--- /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 <oneapi/mkl/rng/device.hpp>
+namespace mkl_rng_dev = oneapi::mkl::rng::device;
+
+namespace dpnp::backend::ext::rng::device::engine::builder
+{
+template <typename EngineT, typename SeedT, typename OffsetT>
+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<SeedT, max_n> seeds{};
+    std::array<OffsetT, max_n> 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<SeedT>(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<OffsetT>(offset_values[i]);
+        }
+    }
+
+    inline auto operator()(void) const
+    {
+        switch (no_of_seeds) {
+        case 1:
+        {
+            if constexpr (std::is_same_v<EngineT,
+                                         mkl_rng_dev::mcg59<EngineT::vec_size>>)
+            {
+                // 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<EngineT,
+                                         mkl_rng_dev::mcg59<EngineT::vec_size>>)
+            {
+                // 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 000000000000..1fd9c9c9e89b
--- /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 <typename Type>
+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 000000000000..b94b93b35b38
--- /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 <oneapi/mkl/rng/device.hpp>
+
+#include "base_builder.hpp"
+#include "base_engine.hpp"
+
+namespace dpnp::backend::ext::rng::device::engine::builder
+{
+namespace mkl_rng_dev = oneapi::mkl::rng::device;
+
+template <std::int32_t VecSize>
+class Builder<mkl_rng_dev::mcg31m1<VecSize>>
+    : public BaseBuilder<mkl_rng_dev::mcg31m1<VecSize>,
+                         std::uint32_t,
+                         std::uint64_t>
+{
+public:
+    using EngineType = mkl_rng_dev::mcg31m1<VecSize>;
+
+    Builder(EngineBase *engine)
+        : BaseBuilder<EngineType, std::uint32_t, std::uint64_t>(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 000000000000..7f8e003b87cf
--- /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 <oneapi/mkl/rng/device.hpp>
+
+#include "base_builder.hpp"
+#include "base_engine.hpp"
+
+namespace dpnp::backend::ext::rng::device::engine::builder
+{
+namespace mkl_rng_dev = oneapi::mkl::rng::device;
+
+template <std::int32_t VecSize>
+class Builder<mkl_rng_dev::mcg59<VecSize>>
+    : public BaseBuilder<mkl_rng_dev::mcg59<VecSize>,
+                         std::uint32_t,
+                         std::uint64_t>
+{
+public:
+    using EngineType = mkl_rng_dev::mcg59<VecSize>;
+
+    Builder(EngineBase *engine)
+        : BaseBuilder<EngineType, std::uint32_t, std::uint64_t>(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 000000000000..a5ab3470b03a
--- /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 <oneapi/mkl/rng/device.hpp>
+
+#include "base_builder.hpp"
+#include "base_engine.hpp"
+
+namespace dpnp::backend::ext::rng::device::engine::builder
+{
+namespace mkl_rng_dev = oneapi::mkl::rng::device;
+
+template <std::int32_t VecSize>
+class Builder<mkl_rng_dev::mrg32k3a<VecSize>>
+    : public BaseBuilder<mkl_rng_dev::mrg32k3a<VecSize>,
+                         std::uint32_t,
+                         std::uint64_t>
+{
+public:
+    using EngineType = mkl_rng_dev::mrg32k3a<VecSize>;
+
+    Builder(EngineBase *engine)
+        : BaseBuilder<EngineType, std::uint32_t, std::uint64_t>(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 000000000000..932c5ee00d76
--- /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 <oneapi/mkl/rng/device.hpp>
+
+#include "base_builder.hpp"
+#include "base_engine.hpp"
+
+namespace dpnp::backend::ext::rng::device::engine::builder
+{
+namespace mkl_rng_dev = oneapi::mkl::rng::device;
+
+template <std::int32_t VecSize>
+class Builder<mkl_rng_dev::philox4x32x10<VecSize>>
+    : public BaseBuilder<mkl_rng_dev::philox4x32x10<VecSize>,
+                         std::uint64_t,
+                         std::uint64_t>
+{
+public:
+    using EngineType = mkl_rng_dev::philox4x32x10<VecSize>;
+
+    Builder(EngineBase *engine)
+        : BaseBuilder<EngineType, std::uint64_t, std::uint64_t>(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 000000000000..4022bd33c18a
--- /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<std::uint32_t> &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 000000000000..e0fbf1741a3f
--- /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<std::uint32_t> &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 000000000000..d6657d6d8d8f
--- /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<std::uint32_t> &seeds,
+             std::uint64_t offset = 0)
+        : EngineBase(q, seeds, offset)
+    {
+    }
+
+    MRG32k3a(sycl::queue &q,
+             std::uint32_t seed,
+             std::vector<std::uint64_t> &offsets)
+        : EngineBase(q, seed, offsets)
+    {
+    }
+
+    MRG32k3a(sycl::queue &q,
+             std::vector<std::uint32_t> &seeds,
+             std::vector<std::uint64_t> &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 000000000000..bef1e7d2119b
--- /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<std::uint64_t> &seeds,
+                  std::uint64_t offset = 0)
+        : EngineBase(q, seeds, offset)
+    {
+    }
+
+    PHILOX4x32x10(sycl::queue &q,
+                  std::uint64_t seed,
+                  std::vector<std::uint64_t> &offsets)
+        : EngineBase(q, seed, offsets)
+    {
+    }
+
+    PHILOX4x32x10(sycl::queue &q,
+                  std::vector<std::uint64_t> &seeds,
+                  std::vector<std::uint64_t> &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 000000000000..643c4091fee9
--- /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 <pybind11/pybind11.h>
+
+// 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<std::uint8_t, 2, 4, 8, 16>{};
+constexpr auto vec_sizes_len = seq_of_vec_sizes.size();
+constexpr auto no_of_engines = engine::no_of_engines * vec_sizes_len;
+
+template <typename VecSizeT, VecSizeT... Ints, auto... Indices>
+inline auto find_vec_size_impl(const VecSizeT vec_size,
+                               std::index_sequence<Indices...>)
+{
+    return std::min({((Ints == vec_size) ? Indices : sizeof...(Indices))...});
+}
+
+template <typename VecSizeT, VecSizeT... Ints>
+int find_vec_size(const VecSizeT vec_size,
+                  std::integer_sequence<VecSizeT, Ints...>)
+{
+    auto res = find_vec_size_impl<VecSizeT, Ints...>(
+        vec_size, std::make_index_sequence<sizeof...(Ints)>{});
+    return (res == sizeof...(Ints)) ? -1 : res;
+}
+
+template <typename DataT, typename Method>
+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<DataT, Method>;
+
+    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<sycl::event> &);
+
+static gaussian_impl_fn_ptr_t gaussian_dispatch_table[no_of_engines]
+                                                     [dpctl_td_ns::num_types]
+                                                     [no_of_methods];
+
+template <typename EngineT,
+          typename DataT,
+          typename Method,
+          unsigned int items_per_wi>
+class gaussian_kernel;
+
+template <typename EngineT, typename DataT, typename Method>
+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<sycl::event> &depends)
+{
+    auto &exec_q = engine->get_queue();
+    dpctl_tu_ns::validate_type_for_device<DataT>(exec_q);
+
+    DataT *out = reinterpret_cast<DataT *>(out_ptr);
+    DataT mean = static_cast<DataT>(mean_val);
+    DataT stddev = static_cast<DataT>(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<EngineT>;
+            EngineBuilderT eng_builder(engine);
+            // eng_builder.print(); // TODO: remove
+
+            using DistributorBuilderT = DistributorBuilder<DataT, Method>;
+            DistributorBuilderT dist_builder(mean, stddev);
+
+            if (is_aligned<required_alignment>(out_ptr)) {
+                constexpr bool enable_sg_load = true;
+                using KernelName =
+                    gaussian_kernel<EngineT, DataT, Method, items_per_wi>;
+
+                cgh.parallel_for<KernelName>(
+                    sycl::nd_range<1>({global_size}, {local_size}),
+                    details::RngContigFunctor<EngineBuilderT,
+                                              DistributorBuilderT, items_per_wi,
+                                              enable_sg_load>(
+                        eng_builder, dist_builder, out, n));
+            }
+            else {
+                constexpr bool disable_sg_load = false;
+                using InnerKernelName =
+                    gaussian_kernel<EngineT, DataT, Method, items_per_wi>;
+                using KernelName =
+                    disabled_sg_loadstore_wrapper_krn<InnerKernelName>;
+
+                cgh.parallel_for<KernelName>(
+                    sycl::nd_range<1>({global_size}, {local_size}),
+                    details::RngContigFunctor<EngineBuilderT,
+                                              DistributorBuilderT, items_per_wi,
+                                              disable_sg_load>(
+                        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<sycl::event, sycl::event>
+    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<sycl::event> &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<size_t>(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 <typename fnT, typename E, typename T, typename M>
+struct GaussianContigFactory
+{
+    fnT get()
+    {
+        if constexpr (dispatch::GaussianTypePairSupportFactory<T,
+                                                               M>::is_defined) {
+            return gaussian_impl<E, T, M>;
+        }
+        else {
+            return nullptr;
+        }
+    }
+};
+
+void init_gaussian_dispatch_3d_table(void)
+{
+    dispatch::Dispatch3DTableBuilder<gaussian_impl_fn_ptr_t,
+                                     GaussianContigFactory, no_of_engines,
+                                     dpctl_td_ns::num_types, no_of_methods>
+        contig;
+    contig.populate<mkl_rng_dev::gaussian_method::box_muller2>(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 000000000000..00973a5d4e58
--- /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 <dpctl4pybind11.hpp>
+
+#include "engine/base_engine.hpp"
+
+namespace dpnp::backend::ext::rng::device
+{
+extern std::pair<sycl::event, sycl::event>
+    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<sycl::event> &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 000000000000..27cbdb80d8d7
--- /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 <pybind11/pybind11.h>
+#include <pybind11/stl.h>
+
+#include <dpctl4pybind11.hpp>
+#include <oneapi/mkl/rng.hpp>
+
+#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_<rng_dev_engine::EngineBase, PyEngineBase /* <--- trampoline */>(
+        m, "EngineBase")
+        .def(py::init<>())
+        .def("get_queue", &rng_dev_engine::EngineBase::get_queue);
+
+    py::class_<rng_dev_engine::MRG32k3a, rng_dev_engine::EngineBase>(m,
+                                                                     "MRG32k3a")
+        .def(py::init<sycl::queue &, std::uint32_t, std::uint64_t>(),
+             py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0)
+        .def(py::init<sycl::queue &, std::vector<std::uint32_t> &,
+                      std::uint64_t>(),
+             py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0)
+        .def(py::init<sycl::queue &, std::uint32_t,
+                      std::vector<std::uint64_t> &>(),
+             py::arg("sycl_queue"), py::arg("seed"),
+             py::arg("offset") = py::list())
+        .def(py::init<sycl::queue &, std::vector<std::uint32_t> &,
+                      std::vector<std::uint64_t> &>(),
+             py::arg("sycl_queue"), py::arg("seed"),
+             py::arg("offset") = py::list());
+
+    py::class_<rng_dev_engine::PHILOX4x32x10, rng_dev_engine::EngineBase>(
+        m, "PHILOX4x32x10")
+        .def(py::init<sycl::queue &, std::uint64_t, std::uint64_t>(),
+             py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0)
+        .def(py::init<sycl::queue &, std::vector<std::uint64_t> &,
+                      std::uint64_t>(),
+             py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0)
+        .def(py::init<sycl::queue &, std::uint64_t,
+                      std::vector<std::uint64_t> &>(),
+             py::arg("sycl_queue"), py::arg("seed"),
+             py::arg("offset") = py::list())
+        .def(py::init<sycl::queue &, std::vector<std::uint64_t> &,
+                      std::vector<std::uint64_t> &>(),
+             py::arg("sycl_queue"), py::arg("seed"),
+             py::arg("offset") = py::list());
+
+    py::class_<rng_dev_engine::MCG31M1, rng_dev_engine::EngineBase>(m,
+                                                                    "MCG31M1")
+        .def(py::init<sycl::queue &, std::uint32_t, std::uint64_t>(),
+             py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0)
+        .def(py::init<sycl::queue &, std::vector<std::uint32_t> &,
+                      std::uint64_t>(),
+             py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0);
+
+    py::class_<rng_dev_engine::MCG59, rng_dev_engine::EngineBase>(m, "MCG59")
+        .def(py::init<sycl::queue &, std::uint32_t, std::uint64_t>(),
+             py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") = 0)
+        .def(py::init<sycl::queue &, std::vector<std::uint32_t> &,
+                      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 000000000000..62ef69a29ea4
--- /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 <pybind11/pybind11.h>
+
+// 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<std::uint8_t, 2, 4, 8, 16>{};
+constexpr auto vec_sizes_len = seq_of_vec_sizes.size();
+constexpr auto no_of_engines = engine::no_of_engines * vec_sizes_len;
+
+template <typename VecSizeT, VecSizeT... Ints, auto... Indices>
+inline auto find_vec_size_impl(const VecSizeT vec_size,
+                               std::index_sequence<Indices...>)
+{
+    return std::min({((Ints == vec_size) ? Indices : sizeof...(Indices))...});
+}
+
+template <typename VecSizeT, VecSizeT... Ints>
+int find_vec_size(const VecSizeT vec_size,
+                  std::integer_sequence<VecSizeT, Ints...>)
+{
+    auto res = find_vec_size_impl<VecSizeT, Ints...>(
+        vec_size, std::make_index_sequence<sizeof...(Ints)>{});
+    return (res == sizeof...(Ints)) ? -1 : res;
+}
+
+template <typename DataT, typename Method>
+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<DataT, Method>;
+
+    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<sycl::event> &);
+
+static uniform_impl_fn_ptr_t uniform_dispatch_table[no_of_engines]
+                                                   [dpctl_td_ns::num_types]
+                                                   [no_of_methods];
+
+template <typename EngineT,
+          typename DataT,
+          typename Method,
+          unsigned int items_per_wi>
+class uniform_kernel;
+
+template <typename EngineT, typename DataT, typename Method>
+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<sycl::event> &depends)
+{
+    auto &exec_q = engine->get_queue();
+    dpctl_tu_ns::validate_type_for_device<DataT>(exec_q);
+
+    DataT *out = reinterpret_cast<DataT *>(out_ptr);
+    DataT a = static_cast<DataT>(a_val);
+    DataT b = static_cast<DataT>(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<EngineT>;
+            EngineBuilderT eng_builder(engine);
+            // eng_builder.print(); // TODO: remove
+
+            using DistributorBuilderT = DistributorBuilder<DataT, Method>;
+            DistributorBuilderT dist_builder(a, b);
+
+            if (is_aligned<required_alignment>(out_ptr)) {
+                constexpr bool enable_sg_load = true;
+                using KernelName =
+                    uniform_kernel<EngineT, DataT, Method, items_per_wi>;
+
+                cgh.parallel_for<KernelName>(
+                    sycl::nd_range<1>({global_size}, {local_size}),
+                    details::RngContigFunctor<EngineBuilderT,
+                                              DistributorBuilderT, items_per_wi,
+                                              enable_sg_load>(
+                        eng_builder, dist_builder, out, n));
+            }
+            else {
+                constexpr bool disable_sg_load = false;
+                using InnerKernelName =
+                    uniform_kernel<EngineT, DataT, Method, items_per_wi>;
+                using KernelName =
+                    disabled_sg_loadstore_wrapper_krn<InnerKernelName>;
+
+                cgh.parallel_for<KernelName>(
+                    sycl::nd_range<1>({global_size}, {local_size}),
+                    details::RngContigFunctor<EngineBuilderT,
+                                              DistributorBuilderT, items_per_wi,
+                                              disable_sg_load>(
+                        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<sycl::event, sycl::event>
+    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<sycl::event> &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<size_t>(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 <typename fnT, typename E, typename T, typename M>
+struct UniformContigFactory
+{
+    fnT get()
+    {
+        if constexpr (dispatch::UniformTypePairSupportFactory<T, M>::is_defined) {
+            return uniform_impl<E, T, M>;
+        }
+        else {
+            return nullptr;
+        }
+    }
+};
+
+void init_uniform_dispatch_3d_table(void)
+{
+    dispatch::Dispatch3DTableBuilder<uniform_impl_fn_ptr_t,
+                                     UniformContigFactory, no_of_engines,
+                                     dpctl_td_ns::num_types, no_of_methods>
+        contig;
+    contig.populate<mkl_rng_dev::uniform_method::standard, mkl_rng_dev::uniform_method::accurate>(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 000000000000..e45bc45f31da
--- /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 <dpctl4pybind11.hpp>
+
+#include "engine/base_engine.hpp"
+
+namespace dpnp::backend::ext::rng::device
+{
+extern std::pair<sycl::event, sycl::event>
+    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<sycl::event> &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 000000000000..f3db23c2a23e
--- /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 000000000000..eb7aa43450fd
--- /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 <oneapi/mkl/rng/device.hpp>
+
+#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 <typename Ty, typename ArgTy, typename Method, typename argMethod>
+struct TypePairDefinedEntry
+    : std::bool_constant<std::is_same_v<Ty, ArgTy> &&
+                         std::is_same_v<Method, argMethod>>
+{
+    static constexpr bool is_defined = true;
+};
+
+template <typename T, typename M>
+struct GaussianTypePairSupportFactory
+{
+    static constexpr bool is_defined = std::disjunction<
+        TypePairDefinedEntry<T,
+                             double,
+                             M,
+                             mkl_rng::gaussian_method::by_default>,
+        TypePairDefinedEntry<T,
+                             double,
+                             M,
+                             mkl_rng::gaussian_method::box_muller2>,
+        TypePairDefinedEntry<T, float, M, mkl_rng::gaussian_method::by_default>,
+        TypePairDefinedEntry<T,
+                             float,
+                             M,
+                             mkl_rng::gaussian_method::box_muller2>,
+        // 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 000000000000..772fa8fc9f3b
--- /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 <oneapi/mkl/rng.hpp>
+
+namespace dpnp::backend::ext::rng::host::dispatch
+{
+namespace mkl_rng = oneapi::mkl::rng;
+
+template <typename funcPtrT,
+          template <typename fnT, typename E, typename T, typename M>
+          typename factory,
+          int _no_of_engines,
+          int _no_of_types,
+          int _no_of_methods>
+class Dispatch3DTableBuilder
+{
+private:
+    template <typename E, typename T>
+    const std::vector<funcPtrT> row_per_method() const
+    {
+        std::vector<funcPtrT> per_method = {
+            factory<funcPtrT, E, T, mkl_rng::gaussian_method::by_default>{}
+                .get(),
+            factory<funcPtrT, E, T, mkl_rng::gaussian_method::box_muller2>{}
+                .get(),
+        };
+        assert(per_method.size() == _no_of_methods);
+        return per_method;
+    }
+
+    template <typename E>
+    auto table_per_type_and_method() const
+    {
+        std::vector<std::vector<funcPtrT>> table_by_type = {
+            row_per_method<E, bool>(),
+            row_per_method<E, int8_t>(),
+            row_per_method<E, uint8_t>(),
+            row_per_method<E, int16_t>(),
+            row_per_method<E, uint16_t>(),
+            row_per_method<E, int32_t>(),
+            row_per_method<E, uint32_t>(),
+            row_per_method<E, int64_t>(),
+            row_per_method<E, uint64_t>(),
+            row_per_method<E, sycl::half>(),
+            row_per_method<E, float>(),
+            row_per_method<E, double>(),
+            row_per_method<E, std::complex<float>>(),
+            row_per_method<E, std::complex<double>>()};
+        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<mkl_rng::mrg32k3a>(),
+            table_per_type_and_method<mkl_rng::philox4x32x10>(),
+            table_per_type_and_method<mkl_rng::mcg31m1>(),
+            table_per_type_and_method<mkl_rng::mcg59>()};
+        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 000000000000..d6b4a2387023
--- /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 <pybind11/pybind11.h>
+
+#include <oneapi/mkl/rng.hpp>
+
+// 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<sycl::event> &);
+
+static gaussian_impl_fn_ptr_t gaussian_dispatch_table[no_of_engines]
+                                                     [dpctl_td_ns::num_types]
+                                                     [no_of_methods];
+
+template <typename EngineT, typename DataT, typename Method>
+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<sycl::event> &depends)
+{
+    auto &exec_q = engine->get_queue();
+    dpctl_tu_ns::validate_type_for_device<DataT>(exec_q);
+
+    DataT *out = reinterpret_cast<DataT *>(out_ptr);
+    DataT mean = static_cast<DataT>(mean_val);
+    DataT stddev = static_cast<DataT>(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<DataT, Method> distribution(mean, stddev);
+    mkl_rng::mcg59 eng(exec_q, seed_values[0]);
+
+    return mkl_rng::generate(distribution, eng, n, out, depends);
+}
+
+std::pair<sycl::event, sycl::event>
+    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<sycl::event> &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<size_t>(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 <typename fnT, typename E, typename T, typename M>
+struct GaussianContigFactory
+{
+    fnT get()
+    {
+        if constexpr (dispatch::GaussianTypePairSupportFactory<T,
+                                                               M>::is_defined) {
+            return gaussian_impl<E, T, M>;
+        }
+        else {
+            return nullptr;
+        }
+    }
+};
+
+void init_gaussian_dispatch_3d_table(void)
+{
+    dispatch::Dispatch3DTableBuilder<gaussian_impl_fn_ptr_t,
+                                     GaussianContigFactory, no_of_engines,
+                                     dpctl_td_ns::num_types, no_of_methods>
+        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 000000000000..2ebf5f976e2f
--- /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 <dpctl4pybind11.hpp>
+
+#include "../device/engine/base_engine.hpp"
+
+namespace dpnp::backend::ext::rng::host
+{
+extern std::pair<sycl::event, sycl::event>
+    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<sycl::event> &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 000000000000..d76b07242e7f
--- /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 <pybind11/pybind11.h>
+#include <pybind11/stl.h>
+
+#include <dpctl4pybind11.hpp>
+#include <oneapi/mkl/rng.hpp>
+
+#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_<rng_dev_engine::EngineBase, PyEngineBase /* <---
+    //     trampoline */>(
+    //         m, "EngineBase")
+    //         .def(py::init<>())
+    //         .def("get_queue", &rng_dev_engine::EngineBase::get_queue);
+
+    //     py::class_<rng_dev_engine::MRG32k3a, rng_dev_engine::EngineBase>(m,
+    //                                                                      "MRG32k3a")
+    //         .def(py::init<sycl::queue &, std::uint32_t, std::uint64_t>(),
+    //              py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") =
+    //              0)
+    //         .def(py::init<sycl::queue &, std::vector<std::uint32_t> &,
+    //                       std::uint64_t>(),
+    //              py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") =
+    //              0)
+    //         .def(py::init<sycl::queue &, std::uint32_t,
+    //                       std::vector<std::uint64_t> &>(),
+    //              py::arg("sycl_queue"), py::arg("seed"),
+    //              py::arg("offset") = py::list())
+    //         .def(py::init<sycl::queue &, std::vector<std::uint32_t> &,
+    //                       std::vector<std::uint64_t> &>(),
+    //              py::arg("sycl_queue"), py::arg("seed"),
+    //              py::arg("offset") = py::list());
+
+    //     py::class_<rng_dev_engine::PHILOX4x32x10,
+    //     rng_dev_engine::EngineBase>(
+    //         m, "PHILOX4x32x10")
+    //         .def(py::init<sycl::queue &, std::uint64_t, std::uint64_t>(),
+    //              py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") =
+    //              0)
+    //         .def(py::init<sycl::queue &, std::vector<std::uint64_t> &,
+    //                       std::uint64_t>(),
+    //              py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") =
+    //              0)
+    //         .def(py::init<sycl::queue &, std::uint64_t,
+    //                       std::vector<std::uint64_t> &>(),
+    //              py::arg("sycl_queue"), py::arg("seed"),
+    //              py::arg("offset") = py::list())
+    //         .def(py::init<sycl::queue &, std::vector<std::uint64_t> &,
+    //                       std::vector<std::uint64_t> &>(),
+    //              py::arg("sycl_queue"), py::arg("seed"),
+    //              py::arg("offset") = py::list());
+
+    //     py::class_<rng_dev_engine::MCG31M1, rng_dev_engine::EngineBase>(m,
+    //                                                                     "MCG31M1")
+    //         .def(py::init<sycl::queue &, std::uint32_t, std::uint64_t>(),
+    //              py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") =
+    //              0)
+    //         .def(py::init<sycl::queue &, std::vector<std::uint32_t> &,
+    //                       std::uint64_t>(),
+    //              py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") =
+    //              0);
+
+    //     py::class_<rng_dev_engine::MCG59, rng_dev_engine::EngineBase>(m,
+    //     "MCG59")
+    //         .def(py::init<sycl::queue &, std::uint32_t, std::uint64_t>(),
+    //              py::arg("sycl_queue"), py::arg("seed"), py::arg("offset") =
+    //              0)
+    //         .def(py::init<sycl::queue &, std::vector<std::uint32_t> &,
+    //                       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());
+}