[RNG] Add mcg59

include/oneapi/mkl/rng/detail/curand/onemkl_rng_curand.hpp

@@ -82,6 +82,12 @@ ONEMKL_EXPORT oneapi::mkl::rng::detail::engine_impl* create_mrg32k3a(cl::sycl::q
 ONEMKL_EXPORT oneapi::mkl::rng::detail::engine_impl* create_mrg32k3a(
     cl::sycl::queue queue, std::initializer_list<std::uint32_t> seed);
 
+ONEMKL_EXPORT oneapi::mkl::rng::detail::engine_impl* create_mcg59(cl::sycl::queue queue,
+                                                                     std::uint64_t seed);
+
+ONEMKL_EXPORT oneapi::mkl::rng::detail::engine_impl* create_mcg59(
+    cl::sycl::queue queue, std::initializer_list<std::uint64_t> seed);
+
 } // namespace curand
 } // namespace rng
 } // namespace mkl

include/oneapi/mkl/rng/detail/mklcpu/onemkl_rng_mklcpu.hpp

@@ -43,6 +43,12 @@ ONEMKL_EXPORT oneapi::mkl::rng::detail::engine_impl* create_mrg32k3a(cl::sycl::q
 ONEMKL_EXPORT oneapi::mkl::rng::detail::engine_impl* create_mrg32k3a(
     cl::sycl::queue queue, std::initializer_list<std::uint32_t> seed);
 
+ONEMKL_EXPORT oneapi::mkl::rng::detail::engine_impl* create_mcg59(cl::sycl::queue queue,
+                                                                    std::uint64_t seed);
+
+ONEMKL_EXPORT oneapi::mkl::rng::detail::engine_impl* create_mcg59(
+    cl::sycl::queue queue, std::initializer_list<std::uint64_t> seed);
+
 } // namespace mklcpu
 } // namespace rng
 } // namespace mkl

include/oneapi/mkl/rng/detail/mklgpu/onemkl_rng_mklgpu.hpp

@@ -43,6 +43,12 @@ ONEMKL_EXPORT oneapi::mkl::rng::detail::engine_impl* create_mrg32k3a(cl::sycl::q
 ONEMKL_EXPORT oneapi::mkl::rng::detail::engine_impl* create_mrg32k3a(
     cl::sycl::queue queue, std::initializer_list<std::uint32_t> seed);
 
+ONEMKL_EXPORT oneapi::mkl::rng::detail::engine_impl* create_mcg59(cl::sycl::queue queue,
+                                                                    std::uint64_t seed);
+
+ONEMKL_EXPORT oneapi::mkl::rng::detail::engine_impl* create_mcg59(
+    cl::sycl::queue queue, std::initializer_list<std::uint64_t> seed);
+
 } // namespace mklgpu
 } // namespace rng
 } // namespace mkl

include/oneapi/mkl/rng/detail/rng_loader.hpp

@@ -45,6 +45,12 @@ ONEMKL_EXPORT engine_impl* create_mrg32k3a(oneapi::mkl::device libkey, cl::sycl:
 ONEMKL_EXPORT engine_impl* create_mrg32k3a(oneapi::mkl::device libkey, cl::sycl::queue queue,
                                            std::initializer_list<std::uint32_t> seed);
 
+ONEMKL_EXPORT engine_impl* create_mcg59(oneapi::mkl::device libkey, cl::sycl::queue queue,
+                                           std::uint64_t seed);
+
+ONEMKL_EXPORT engine_impl* create_mcg59(oneapi::mkl::device libkey, cl::sycl::queue queue,
+                                           std::initializer_list<std::uint64_t> seed);
+
 } // namespace detail
 } // namespace rng
 } // namespace mkl

include/oneapi/mkl/rng/distributions.hpp

@@ -358,7 +358,8 @@ class poisson {
 template <typename UIntType = std::uint32_t>
 class bits {
 public:
-    static_assert(std::is_same<UIntType, std::uint32_t>::value, "rng bits type is not supported");
+    static_assert(std::is_same<UIntType, std::uint32_t>::value || std::is_same<UIntType, std::uint64_t>::value, 
+            "rng bits type is not supported");
     using result_type = UIntType;
 };
 

include/oneapi/mkl/rng/engines.hpp

@@ -210,6 +210,77 @@ class mrg32k3a {
                                 const std::vector<sycl::event>& dependencies);
 };
 
+// Class oneapi::mkl::rng::mcg59
+//
+// Represents Mcg59 counter-based pseudorandom number generator
+//
+// Supported parallelization methods:
+//      leapfrog
+class mcg59 {
+public:
+    static constexpr std::uint64_t default_seed = 0;
+
+    mcg59(sycl::queue queue, std::uint64_t seed = default_seed)
+            : pimpl_(detail::create_mcg59(get_device_id(queue), queue, seed)) {}
+
+#ifdef ENABLE_MKLCPU_BACKEND
+    mcg59(backend_selector<backend::mklcpu> selector, std::uint64_t seed = default_seed)
+            : pimpl_(mklcpu::create_mcg59(selector.get_queue(), seed)) {}
+
+#endif
+
+#ifdef ENABLE_MKLGPU_BACKEND
+    mcg59(backend_selector<backend::mklgpu> selector, std::uint64_t seed = default_seed)
+            : pimpl_(mklgpu::create_mcg59(selector.get_queue(), seed)) {}
+
+#endif
+
+#ifdef ENABLE_CURAND_BACKEND
+    mcg59(backend_selector<backend::curand> selector, std::uint64_t seed = default_seed)
+            : pimpl_(curand::create_mcg59(selector.get_queue(), seed)) {}
+#endif
+
+    mcg59(const mcg59& other) {
+        pimpl_.reset(other.pimpl_.get()->copy_state());
+    }
+
+    mcg59(mcg59&& other) {
+        pimpl_ = std::move(other.pimpl_);
+    }
+
+    mcg59& operator=(const mcg59& other) {
+        if (this == &other)
+            return *this;
+        pimpl_.reset(other.pimpl_.get()->copy_state());
+        return *this;
+    }
+
+    mcg59& operator=(mcg59&& other) {
+        if (this == &other)
+            return *this;
+        pimpl_ = std::move(other.pimpl_);
+        return *this;
+    }
+
+private:
+    std::unique_ptr<detail::engine_impl> pimpl_;
+
+    template <typename Engine>
+    friend void skip_ahead(Engine& engine, std::uint64_t num_to_skip);
+
+    template <typename Engine>
+    friend void leapfrog(Engine& engine, std::uint64_t idx, std::uint64_t stride);
+
+    template <typename Distr, typename Engine>
+    friend void generate(const Distr& distr, Engine& engine, std::int64_t n,
+                         sycl::buffer<typename Distr::result_type, 1>& r);
+
+    template <typename Distr, typename Engine>
+    friend sycl::event generate(const Distr& distr, Engine& engine, std::int64_t n,
+                                typename Distr::result_type* r,
+                                const std::vector<sycl::event>& dependencies);
+};
+
 // Default engine to be used for common cases
 using default_engine = philox4x32x10;
 

src/rng/backends/curand/CMakeLists.txt

@@ -62,6 +62,7 @@ find_package(cuRAND REQUIRED)
 
 set(SOURCES philox4x32x10.cpp
   mrg32k3a.cpp
+  mcg59.cpp
   $<$<BOOL:${BUILD_SHARED_LIBS}>: mkl_rng_curand_wrappers.cpp>)
 
 add_library(${LIB_NAME})

src/rng/backends/curand/mcg59.cpp

@@ -0,0 +1,110 @@
+/*******************************************************************************
+ * cuRAND back-end Copyright (c) 2021, The Regents of the University of
+ * California, through Lawrence Berkeley National Laboratory (subject to receipt
+ * of any required approvals from the U.S. Dept. of Energy). All rights
+ * reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * (1) Redistributions of source code must retain the above copyright notice,
+ * this list of conditions and the following disclaimer.
+ *
+ * (2) 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.
+ *
+ * (3) Neither the name of the University of California, Lawrence Berkeley
+ * National Laboratory, U.S. Dept. of Energy nor the names of its contributors
+ * may be used to endorse or promote products derived from this software
+ * without specific prior written permission.
+ *
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+ * ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER 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.
+ *
+ * You are under no obligation whatsoever to provide any bug fixes, patches,
+ * or upgrades to the features, functionality or performance of the source
+ * code ("Enhancements") to anyone; however, if you choose to make your
+ * Enhancements available either publicly, or directly to Lawrence Berkeley
+ * National Laboratory, without imposing a separate written license agreement
+ * for such Enhancements, then you hereby grant the following license: a
+ * non-exclusive, royalty-free perpetual license to install, use, modify,
+ * prepare derivative works, incorporate into other computer software,
+ * distribute, and sublicense such enhancements or derivative works thereof,
+ * in binary and source code form.
+ *
+ * If you have questions about your rights to use or distribute this software,
+ * please contact Berkeley Lab's Intellectual Property Office at
+ * [email protected].
+ *
+ * NOTICE.  This Software was developed under funding from the U.S. Department
+ * of Energy and the U.S. Government consequently retains certain rights.  As
+ * such, the U.S. Government has been granted for itself and others acting on
+ * its behalf a paid-up, nonexclusive, irrevocable, worldwide license in the
+ * Software to reproduce, distribute copies to the public, prepare derivative
+ * works, and perform publicly and display publicly, and to permit others to do
+ * so.
+ ******************************************************************************/
+
+#include <CL/sycl.hpp>
+#include <CL/sycl/backend/cuda.hpp>
+#include <iostream>
+
+#include "oneapi/mkl/rng/detail/engine_impl.hpp"
+// #include "oneapi/mkl/rng/engines.hpp"
+#include "curand_helper.hpp"
+#include "oneapi/mkl/exceptions.hpp"
+#include "oneapi/mkl/rng/detail/curand/onemkl_rng_curand.hpp"
+
+namespace oneapi {
+namespace mkl {
+namespace rng {
+namespace curand {
+/*
+ * Note that cuRAND consists of two pieces: a host (CPU) API and a device (GPU)
+ * API. The host API acts like any standard library; the `curand.h' header is
+ * included and the functions can be called as usual. The generator is
+ * instantiated on the host and random numbers can be generated on either the
+ * host CPU or device. For device-side generation, calls to the library happen
+ * on the host, but the actual work of RNG is done on the device. In this case,
+ * the resulting random numbers are stored in global memory on the device. These
+ * random numbers can then be used in other kernels or be copied back to the
+ * host for further processing. For host-side generation, everything is done on
+ * the host, and the random numbers are stored in host memory.
+ *
+ * The second piece is the device header, `curand_kernel.h'. Using this file
+ * permits setting up random number generator states and generating sequences of
+ * random numbers. This allows random numbers to be generated and immediately
+ * consumed in other kernels without requiring the random numbers to be written
+ * to, and read from, global memory.
+ *
+ * Here we utilize the host API since this is most aligned with how oneMKL
+ * generates random numbers.
+ *
+ */
+
+oneapi::mkl::rng::detail::engine_impl* create_mcg59(sycl::queue queue, std::uint64_t seed) {
+    throw oneapi::mkl::unimplemented("rng", "mcg59 engine");
+    return nullptr;
+}
+
+oneapi::mkl::rng::detail::engine_impl* create_mcg59(
+    cl::sycl::queue queue, std::initializer_list<std::uint64_t> seed) {
+    throw oneapi::mkl::unimplemented("rng", "mcg59 engine");
+    return nullptr;
+}
+
+} // namespace curand
+} // namespace rng
+} // namespace mkl
+} // namespace oneapi

src/rng/backends/curand/mkl_rng_curand_wrappers.cpp

@@ -64,5 +64,6 @@
 extern "C" ONEMKL_EXPORT rng_function_table_t mkl_rng_table = {
     WRAPPER_VERSION, oneapi::mkl::rng::curand::create_philox4x32x10,
     oneapi::mkl::rng::curand::create_philox4x32x10, oneapi::mkl::rng::curand::create_mrg32k3a,
-    oneapi::mkl::rng::curand::create_mrg32k3a
+    oneapi::mkl::rng::curand::create_mrg32k3a, oneapi::mkl::rng::curand::create_mcg59,
+    oneapi::mkl::rng::curand::create_mcg59
 };

src/rng/backends/mklcpu/CMakeLists.txt

@@ -25,6 +25,7 @@ find_package(MKL REQUIRED)
 set(SOURCES cpu_common.hpp
   philox4x32x10.cpp
   mrg32k3a.cpp
+  mcg59.cpp
   $<$<BOOL:${BUILD_SHARED_LIBS}>: mkl_rng_cpu_wrappers.cpp>
 )