Add cuRAND support (uxlfoundation#75)

This commit adds a cuRAND backend to oneMKL.

The implementation follows closely to that of the existing x86 and
Intel GPU, with the following limitations:

(1) the cuRAND Host API does not support copy-construction for
    generators;
(2) cuRAND does not support initializer lists for seeding
    generators;
(3) cuRAND uses inverse cumulative distribution function (ICDF)
    methods only for quasi-random number generators; and
(4) cuRAND does not support custom ranges on the output sequence
    of random numbers.

Functionality requiring any of 1-3 above are not implemented here
but could be added at a later time. Additional kernels were added to
support (4).

Author: vrpascuzzi

CMakeLists.txt

@@ -47,6 +47,7 @@ if(ENABLE_MKLCPU_BACKEND)
   option(ENABLE_MKLCPU_THREAD_TBB "" ON)
 endif()
 option(ENABLE_CUBLAS_BACKEND "" OFF)
+option(ENABLE_CURAND_BACKEND "" OFF)
 option(ENABLE_NETLIB_BACKEND "" OFF)
 
 ## Domains
@@ -58,7 +59,8 @@ if(ENABLE_MKLCPU_BACKEND
   list(APPEND DOMAINS_LIST "blas")
 endif()
 if(ENABLE_MKLCPU_BACKEND
-        OR ENABLE_MKLGPU_BACKEND)
+        OR ENABLE_MKLGPU_BACKEND
+        OR ENABLE_CURAND_BACKEND)
   list(APPEND DOMAINS_LIST "rng")
 endif()
 

cmake/FindCompiler.cmake

@@ -31,10 +31,25 @@ if(is_dpcpp)
 
   add_library(ONEMKL::SYCL::SYCL INTERFACE IMPORTED)
   if(UNIX)
-    set_target_properties(ONEMKL::SYCL::SYCL PROPERTIES
-      INTERFACE_COMPILE_OPTIONS "-fsycl"
-      INTERFACE_LINK_OPTIONS "-fsycl"
-      INTERFACE_LINK_LIBRARIES ${SYCL_LIBRARY})
+    set(UNIX_INTERFACE_COMPILE_OPTIONS -fsycl)
+    set(UNIX_INTERFACE_LINK_OPTIONS -fsycl)
+    if(ENABLE_CURAND_BACKEND)
+      list(APPEND UNIX_INTERFACE_COMPILE_OPTIONS
+        -fsycl-targets=nvptx64-nvidia-cuda-sycldevice -fsycl-unnamed-lambda)
+      list(APPEND UNIX_INTERFACE_LINK_OPTIONS
+        -fsycl-targets=nvptx64-nvidia-cuda-sycldevice)
+    endif()
+    if(ENABLE_CURAND_BACKEND)
+      set_target_properties(ONEMKL::SYCL::SYCL PROPERTIES
+        INTERFACE_COMPILE_OPTIONS "${UNIX_INTERFACE_COMPILE_OPTIONS}"
+        INTERFACE_LINK_OPTIONS "${UNIX_INTERFACE_LINK_OPTIONS}"
+        INTERFACE_LINK_LIBRARIES ${SYCL_LIBRARY})
+    else()
+      set_target_properties(ONEMKL::SYCL::SYCL PROPERTIES
+        INTERFACE_COMPILE_OPTIONS "-fsycl"
+        INTERFACE_LINK_OPTIONS "-fsycl"
+        INTERFACE_LINK_LIBRARIES ${SYCL_LIBRARY})
+    endif()
   else()
     set_target_properties(ONEMKL::SYCL::SYCL PROPERTIES
       INTERFACE_COMPILE_OPTIONS "-fsycl"

cmake/FindcuRAND.cmake

@@ -0,0 +1,51 @@
+#===============================================================================
+# Copyright 2021 Intel Corporation
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions
+# and limitations under the License.
+#
+#
+# SPDX-License-Identifier: Apache-2.0
+#===============================================================================
+
+find_package(CUDA 10.0 REQUIRED)
+get_filename_component(SYCL_BINARY_DIR ${CMAKE_CXX_COMPILER} DIRECTORY)
+# the OpenCL include file from cuda is opencl 1.1 and it is not compatible with DPC++
+# the OpenCL include headers 1.2 onward is required. This is used to bypass NVIDIA OpenCL headers
+find_path(OPENCL_INCLUDE_DIR CL/cl.h OpenCL/cl.h 
+HINTS 
+${OPENCL_INCLUDE_DIR}
+${SYCL_BINARY_DIR}/../include/sycl/
+)
+# this is work around to avoid duplication half creation in both cuda and SYCL
+add_compile_definitions(CUDA_NO_HALF)
+
+find_package(Threads REQUIRED)
+
+include(FindPackageHandleStandardArgs)
+find_package_handle_standard_args(cuRAND
+    REQUIRED_VARS
+	CUDA_TOOLKIT_INCLUDE
+	CUDA_curand_LIBRARY
+        CUDA_LIBRARIES
+        CUDA_CUDA_LIBRARY
+        OPENCL_INCLUDE_DIR
+)
+if(NOT TARGET ONEMKL::cuRAND::cuRAND)
+  add_library(ONEMKL::cuRAND::cuRAND SHARED IMPORTED)
+  set_target_properties(ONEMKL::cuRAND::cuRAND PROPERTIES
+      IMPORTED_LOCATION ${CUDA_curand_LIBRARY}
+      INTERFACE_INCLUDE_DIRECTORIES "${OPENCL_INCLUDE_DIR};${CUDA_TOOLKIT_INCLUDE}"
+      INTERFACE_LINK_LIBRARIES "Threads::Threads;${CUDA_CUDA_LIBRARY};${CUDA_LIBRARIES}"
+  )
+
+endif()

include/oneapi/mkl/detail/backends.hpp

@@ -26,15 +26,15 @@
 namespace oneapi {
 namespace mkl {
 
-enum class backend { mklcpu, mklgpu, cublas, netlib, unsupported };
+enum class backend { mklcpu, mklgpu, cublas, curand, netlib, unsupported };
 
 typedef std::map<backend, std::string> backendmap;
 
-static backendmap backend_map = { { backend::mklcpu, "mklcpu" },
-                                  { backend::mklgpu, "mklgpu" },
-                                  { backend::cublas, "cublas" },
-                                  { backend::netlib, "netlib" },
-                                  { backend::unsupported, "unsupported" } };
+static backendmap backend_map = {
+    { backend::mklcpu, "mklcpu" }, { backend::mklgpu, "mklgpu" },
+    { backend::cublas, "cublas" }, { backend::curand, "curand" },
+    { backend::netlib, "netlib" }, { backend::unsupported, "unsupported" }
+};
 
 } //namespace mkl
 } //namespace oneapi

include/oneapi/mkl/detail/backends_table.hpp

@@ -74,6 +74,12 @@ static std::map<domain, std::map<device, std::vector<const char*>>> libraries =
           {
 #ifdef ENABLE_MKLGPU_BACKEND
               LIB_NAME("rng_mklgpu")
+#endif
+          } },
+        { device::nvidiagpu,
+          {
+#ifdef ENABLE_CURAND_BACKEND
+              LIB_NAME("rng_curand")
 #endif
           } } } }
 };

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

@@ -0,0 +1,51 @@
+/*******************************************************************************
+* Copyright 2020 Intel Corporation
+*
+* Licensed under the Apache License, Version 2.0 (the "License");
+* you may not use this file except in compliance with the License.
+* You may obtain a copy of the License at
+*
+* http://www.apache.org/licenses/LICENSE-2.0
+*
+* Unless required by applicable law or agreed to in writing,
+* software distributed under the License is distributed on an "AS IS" BASIS,
+* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+* See the License for the specific language governing permissions
+* and limitations under the License.
+*
+*
+* SPDX-License-Identifier: Apache-2.0
+*******************************************************************************/
+
+#ifndef _ONEMKL_RNG_CURAND_HPP_
+#define _ONEMKL_RNG_CURAND_HPP_
+
+#include <cstdint>
+#include <CL/sycl.hpp>
+
+#include "oneapi/mkl/detail/export.hpp"
+#include "oneapi/mkl/rng/detail/engine_impl.hpp"
+
+namespace oneapi {
+namespace mkl {
+namespace rng {
+namespace curand {
+
+ONEMKL_EXPORT oneapi::mkl::rng::detail::engine_impl* create_philox4x32x10(cl::sycl::queue queue,
+                                                                          std::uint64_t seed);
+
+ONEMKL_EXPORT oneapi::mkl::rng::detail::engine_impl* create_philox4x32x10(
+    cl::sycl::queue queue, std::initializer_list<std::uint64_t> seed);
+
+ONEMKL_EXPORT oneapi::mkl::rng::detail::engine_impl* create_mrg32k3a(cl::sycl::queue queue,
+                                                                     std::uint32_t seed);
+
+ONEMKL_EXPORT oneapi::mkl::rng::detail::engine_impl* create_mrg32k3a(
+    cl::sycl::queue queue, std::initializer_list<std::uint32_t> seed);
+
+} // namespace curand
+} // namespace rng
+} // namespace mkl
+} // namespace oneapi
+
+#endif //_ONEMKL_RNG_CURAND_HPP_

include/oneapi/mkl/rng/engines.hpp

@@ -37,6 +37,9 @@
 #ifdef ENABLE_MKLGPU_BACKEND
 #include "oneapi/mkl/rng/detail/mklgpu/onemkl_rng_mklgpu.hpp"
 #endif
+#ifdef ENABLE_CURAND_BACKEND
+#include "oneapi/mkl/rng/detail/curand/onemkl_rng_curand.hpp"
+#endif
 
 namespace oneapi {
 namespace mkl {
@@ -76,6 +79,15 @@ class philox4x32x10 {
             : pimpl_(mklgpu::create_philox4x32x10(selector.get_queue(), seed)) {}
 #endif
 
+#ifdef ENABLE_CURAND_BACKEND
+    philox4x32x10(backend_selector<backend::curand> selector, std::uint64_t seed = default_seed)
+            : pimpl_(curand::create_philox4x32x10(selector.get_queue(), seed)) {}
+
+    philox4x32x10(backend_selector<backend::curand> selector,
+                  std::initializer_list<std::uint64_t> seed)
+            : pimpl_(curand::create_philox4x32x10(selector.get_queue(), seed)) {}
+#endif
+
     philox4x32x10(const philox4x32x10& other) {
         pimpl_.reset(other.pimpl_.get()->copy_state());
     }
@@ -149,6 +161,14 @@ class mrg32k3a {
             : pimpl_(mklgpu::create_mrg32k3a(selector.get_queue(), seed)) {}
 #endif
 
+#ifdef ENABLE_CURAND_BACKEND
+    mrg32k3a(backend_selector<backend::curand> selector, std::uint32_t seed = default_seed)
+            : pimpl_(curand::create_mrg32k3a(selector.get_queue(), seed)) {}
+
+    mrg32k3a(backend_selector<backend::curand> selector, std::initializer_list<std::uint32_t> seed)
+            : pimpl_(curand::create_mrg32k3a(selector.get_queue(), seed)) {}
+#endif
+
     mrg32k3a(const mrg32k3a& other) {
         pimpl_.reset(other.pimpl_.get()->copy_state());
     }

src/config.hpp.in

@@ -21,6 +21,7 @@
 #define ONEMKL_CONFIG_H
 
 #cmakedefine ENABLE_CUBLAS_BACKEND
+#cmakedefine ENABLE_CURAND_BACKEND
 #cmakedefine ENABLE_MKLCPU_BACKEND
 #cmakedefine ENABLE_MKLGPU_BACKEND
 #cmakedefine ENABLE_NETLIB_BACKEND

src/rng/backends/CMakeLists.txt

@@ -24,3 +24,7 @@ endif()
 if(ENABLE_MKLGPU_BACKEND)
   add_subdirectory(mklgpu)
 endif()
+
+if(ENABLE_CURAND_BACKEND AND UNIX)
+  add_subdirectory(curand)
+endif()

src/rng/backends/curand/CMakeLists.txt

@@ -0,0 +1,60 @@
+#===============================================================================
+# Copyright 2021 Intel Corporation
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions
+# and limitations under the License.
+#
+#
+# SPDX-License-Identifier: Apache-2.0
+#===============================================================================
+
+set(LIB_NAME onemkl_rng_curand)
+set(LIB_OBJ ${LIB_NAME}_obj)
+find_package(cuRAND REQUIRED)
+
+add_library(${LIB_NAME})
+add_library(${LIB_OBJ} OBJECT
+  philox4x32x10.cpp
+  mrg32k3a.cpp
+  $<$<BOOL:${BUILD_SHARED_LIBS}>: mkl_rng_curand_wrappers.cpp>
+)
+
+target_include_directories(${LIB_OBJ}
+  PRIVATE ${PROJECT_SOURCE_DIR}/include
+          ${PROJECT_SOURCE_DIR}/src
+          ${CMAKE_BINARY_DIR}/bin
+          ${MKL_INCLUDE}
+)
+
+target_link_libraries(${LIB_OBJ} PUBLIC ONEMKL::SYCL::SYCL ONEMKL::cuRAND::cuRAND)
+target_compile_features(${LIB_OBJ} PUBLIC cxx_std_11)
+set_target_properties(${LIB_OBJ} PROPERTIES
+  POSITION_INDEPENDENT_CODE ON
+)
+
+target_link_libraries(${LIB_NAME} PUBLIC ${LIB_OBJ})
+
+# Add major version to the library
+set_target_properties(${LIB_NAME} PROPERTIES
+  SOVERSION ${PROJECT_VERSION_MAJOR}
+)
+
+# Add dependencies rpath to the library
+list(APPEND CMAKE_BUILD_RPATH $<TARGET_FILE_DIR:${LIB_NAME}>)
+
+# Add the library to install package
+install(TARGETS ${LIB_OBJ} EXPORT oneMKLTargets)
+install(TARGETS ${LIB_NAME} EXPORT oneMKLTargets
+  RUNTIME DESTINATION bin
+  ARCHIVE DESTINATION lib
+  LIBRARY DESTINATION lib
+)