Implement build for NVIDIA GPU

CMakeLists.txt

@@ -67,6 +67,62 @@ find_package(Dpctl REQUIRED)
 message(STATUS "Dpctl_INCLUDE_DIR=" ${Dpctl_INCLUDE_DIR})
 message(STATUS "Dpctl_TENSOR_INCLUDE_DIR=" ${Dpctl_TENSOR_INCLUDE_DIR})
 
+option(DPNP_TARGET_CUDA
+    "Build DPNP to target CUDA devices"
+    OFF
+)
+option(DPNP_USE_ONEMKL_INTERFACES
+    "Build DPNP with oneMKL Interfaces"
+    OFF
+)
+set(_dpnp_sycl_targets)
+set(_use_onemkl_interfaces_cuda OFF)
+if ("x${DPNP_SYCL_TARGETS}" STREQUAL "x")
+   if(DPNP_TARGET_CUDA)
+      set(_dpnp_sycl_targets "nvptx64-nvidia-cuda,spir64-unknown-unknown")
+      set(_use_onemkl_interfaces_cuda ON)
+   else()
+      if(DEFINED ENV{DPNP_TARGET_CUDA})
+          set(_dpnp_sycl_targets "nvptx64-nvidia-cuda,spir64-unknown-unknown")
+          set(_use_onemkl_interfaces_cuda ON)
+      endif()
+   endif()
+else()
+   set(_dpnp_sycl_targets ${DPNP_SYCL_TARGETS})
+endif()
+
+if(_dpnp_sycl_targets)
+   message(STATUS "Compiling for -fsycl-targets=${_dpnp_sycl_targets}")
+endif()
+
+set(_use_onemkl_interfaces OFF)
+if(DPNP_USE_ONEMKL_INTERFACES)
+    set(_use_onemkl_interfaces ON)
+else()
+    if(DEFINED ENV{DPNP_USE_ONEMKL_INTERFACES})
+        set(_use_onemkl_interfaces ON)
+    endif()
+endif()
+
+if(_use_onemkl_interfaces OR _use_onemkl_interfaces_cuda)
+    set(BUILD_FUNCTIONAL_TESTS False)
+    set(BUILD_EXAMPLES False)
+    if(_use_onemkl_interfaces_cuda)
+        # set(ENABLE_CUBLAS_BACKEND True)
+        set(ENABLE_CUSOLVER_BACKEND True)
+        set(ENABLE_CUFFT_BACKEND True)
+        # set(ENABLE_CURAND_BACKEND True)
+        set(ENABLE_MKLGPU_BACKEND False)
+        set(ENABLE_MKLCPU_BACKEND False)
+    endif()
+    FetchContent_Declare(
+            onemkl_interfaces_library
+            GIT_REPOSITORY https://github.com/oneapi-src/oneMKL.git
+            GIT_TAG develop
+    )
+    FetchContent_MakeAvailable(onemkl_interfaces_library)
+endif()
+
 if(WIN32)
     string(CONCAT WARNING_FLAGS
         "-Wall "

dpnp/backend/extensions/fft/CMakeLists.txt

@@ -34,6 +34,20 @@ set(_module_src
 pybind11_add_module(${python_module_name} MODULE ${_module_src})
 add_sycl_to_target(TARGET ${python_module_name} SOURCES ${_module_src})
 
+if(_dpnp_sycl_targets)
+    # make fat binary
+    target_compile_options(
+        ${python_module_name}
+        PRIVATE
+        -fsycl-targets=${_dpnp_sycl_targets}
+    )
+    target_link_options(
+        ${python_module_name}
+        PRIVATE
+        -fsycl-targets=${_dpnp_sycl_targets}
+    )
+endif()
+
 if (WIN32)
     if (${CMAKE_VERSION} VERSION_LESS "3.27")
         # this is a work-around for target_link_options inserting option after -link option, cause
@@ -68,7 +82,12 @@ 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_SYCL::DFT)
+if(_use_onemkl_interfaces)
+    target_link_libraries(${python_module_name} PUBLIC onemkl)
+    target_compile_options(${python_module_name} PRIVATE -DUSE_ONEMKL_INTERFACES)
+else()
+    target_link_libraries(${python_module_name} PUBLIC MKL::MKL_SYCL::DFT)
+endif()
 
 install(TARGETS ${python_module_name}
   DESTINATION "dpnp/backend/extensions/fft"

dpnp/backend/extensions/fft/common.hpp

@@ -61,10 +61,7 @@ class DescriptorWrapper
         queue_ptr_ = std::make_unique<sycl::queue>(q);
     }
 
-    descr_type &get_descriptor()
-    {
-        return descr_;
-    }
+    descr_type &get_descriptor() { return descr_; }
 
     const sycl::queue &get_queue() const
     {
@@ -187,26 +184,32 @@ class DescriptorWrapper
     // config_param::PLACEMENT
     bool get_in_place()
     {
-        // TODO: replace when MKLD-10506 is implemented
-        // mkl_dft::config_value placement;
+#if defined(USE_ONEMKL_INTERFACES)
+        mkl_dft::config_value placement;
+        descr_.get_value(mkl_dft::config_param::PLACEMENT, &placement);
+        return (placement == mkl_dft::config_value::INPLACE);
+#else
+        // TODO: remove branch when MKLD-10506 is implemented
         DFTI_CONFIG_VALUE placement;
-
         descr_.get_value(mkl_dft::config_param::PLACEMENT, &placement);
-        // TODO: replace when MKLD-10506 is implemented
-        // return (placement == mkl_dft::config_value::INPLACE);
         return (placement == DFTI_CONFIG_VALUE::DFTI_INPLACE);
+#endif // USE_ONEMKL_INTERFACES
     }
 
     void set_in_place(const bool &in_place_request)
     {
-        // TODO: replace when MKLD-10506 is implemented
-        // descr_.set_value(mkl_dft::config_param::PLACEMENT, (in_place_request)
-        // ? mkl_dft::config_value::INPLACE :
-        // mkl_dft::config_value::NOT_INPLACE);
+#if defined(USE_ONEMKL_INTERFACES)
+        descr_.set_value(mkl_dft::config_param::PLACEMENT,
+                         (in_place_request)
+                             ? mkl_dft::config_value::INPLACE
+                             : mkl_dft::config_value::NOT_INPLACE);
+#else
+        // TODO: remove branch when MKLD-10506 is implemented
         descr_.set_value(mkl_dft::config_param::PLACEMENT,
                          (in_place_request)
                              ? DFTI_CONFIG_VALUE::DFTI_INPLACE
                              : DFTI_CONFIG_VALUE::DFTI_NOT_INPLACE);
+#endif // USE_ONEMKL_INTERFACES
     }
 
     // config_param::PRECISION
@@ -221,14 +224,16 @@ class DescriptorWrapper
     // config_param::COMMIT_STATUS
     bool is_committed()
     {
-        // TODO: replace when MKLD-10506 is implemented
-        // mkl_dft::config_value committed;
+#if defined(USE_ONEMKL_INTERFACES)
+        mkl_dft::config_value committed;
+        descr_.get_value(mkl_dft::config_param::COMMIT_STATUS, &committed);
+        return (committed == mkl_dft::config_value::COMMITTED);
+#else
+        // TODO: remove branch when MKLD-10506 is implemented
         DFTI_CONFIG_VALUE committed;
-
         descr_.get_value(mkl_dft::config_param::COMMIT_STATUS, &committed);
-        // TODO: replace when MKLD-10506 is implemented
-        // return (committed == mkl_dft::config_value::COMMITTED);
         return (committed == DFTI_CONFIG_VALUE::DFTI_COMMITTED);
+#endif // USE_ONEMKL_INTERFACES
     }
 
 private:

dpnp/backend/extensions/fft/out_of_place.cpp

@@ -23,7 +23,10 @@
 // THE POSSIBILITY OF SUCH DAMAGE.
 //*****************************************************************************
 
+#if not defined(USE_ONEMKL_INTERFACES)
 #include <oneapi/mkl/dfti.hpp>
+#endif // USE_ONEMKL_INTERFACES
+
 #include <sycl/sycl.hpp>
 
 #include <dpctl4pybind11.hpp>

dpnp/backend/extensions/lapack/CMakeLists.txt

@@ -48,6 +48,20 @@ set(_module_src
 pybind11_add_module(${python_module_name} MODULE ${_module_src})
 add_sycl_to_target(TARGET ${python_module_name} SOURCES ${_module_src})
 
+if(_dpnp_sycl_targets)
+    # make fat binary
+    target_compile_options(
+        ${python_module_name}
+        PRIVATE
+        -fsycl-targets=${_dpnp_sycl_targets}
+    )
+    target_link_options(
+        ${python_module_name}
+        PRIVATE
+        -fsycl-targets=${_dpnp_sycl_targets}
+    )
+endif()
+
 if (WIN32)
     if (${CMAKE_VERSION} VERSION_LESS "3.27")
         # this is a work-around for target_link_options inserting option after -link option, cause
@@ -82,7 +96,12 @@ 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_SYCL::LAPACK)
+if(_use_onemkl_interfaces)
+    target_link_libraries(${python_module_name} PUBLIC onemkl)
+    target_compile_options(${python_module_name} PRIVATE -DUSE_ONEMKL_INTERFACES)
+else()
+    target_link_libraries(${python_module_name} PUBLIC MKL::MKL_SYCL::LAPACK)
+endif()
 
 install(TARGETS ${python_module_name}
   DESTINATION "dpnp/backend/extensions/lapack"

dpnp/backend/extensions/lapack/gesv.cpp

@@ -65,6 +65,11 @@ static sycl::event gesv_impl(sycl::queue exec_q,
                              std::vector<sycl::event> &host_task_events,
                              const std::vector<sycl::event> &depends)
 {
+#if defined(USE_ONEMKL_INTERFACES)
+    // Temporary flag for build only
+    // FIXME: Need to implement by using lapack::getrf and lapack::getrs
+    std::logic_error("Not Implemented");
+#else
     type_utils::validate_type_for_device<T>(exec_q);
 
     T *a = reinterpret_cast<T *>(in_a);
@@ -167,6 +172,7 @@ static sycl::event gesv_impl(sycl::queue exec_q,
     host_task_events.push_back(clean_up_event);
 
     return gesv_event;
+#endif // USE_ONEMKL_INTERFACES
 }
 
 std::pair<sycl::event, sycl::event>

dpnp/backend/extensions/vm/CMakeLists.txt

@@ -23,46 +23,48 @@
 # THE POSSIBILITY OF SUCH DAMAGE.
 # *****************************************************************************
 
-set(_elementwise_sources
-    ${CMAKE_CURRENT_SOURCE_DIR}/abs.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/acos.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/acosh.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/add.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/asin.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/asinh.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/atan.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/atan2.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/atanh.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/cbrt.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/ceil.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/conj.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/cos.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/cosh.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/div.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/exp.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/exp2.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/expm1.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/floor.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/fmax.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/fmin.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/fmod.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/hypot.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/ln.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/log10.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/log1p.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/log2.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/mul.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/pow.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/rint.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/sin.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/sinh.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/sqr.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/sqrt.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/sub.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/tan.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/tanh.cpp
-    ${CMAKE_CURRENT_SOURCE_DIR}/trunc.cpp
+if(NOT _use_onemkl_interfaces)
+    set(_elementwise_sources
+        ${CMAKE_CURRENT_SOURCE_DIR}/abs.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/acos.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/acosh.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/add.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/asin.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/asinh.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/atan.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/atan2.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/atanh.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/cbrt.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/ceil.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/conj.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/cos.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/cosh.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/div.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/exp.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/exp2.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/expm1.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/floor.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/fmax.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/fmin.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/fmod.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/hypot.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/ln.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/log10.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/log1p.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/log2.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/mul.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/pow.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/rint.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/sin.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/sinh.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/sqr.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/sqrt.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/sub.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/tan.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/tanh.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/trunc.cpp
 )
+endif()
 
 set(_module_src
     # TODO: remove sources from `elementwise_functions` folder
@@ -111,7 +113,11 @@ 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_SYCL::VM)
+if(_use_onemkl_interfaces)
+    target_compile_options(${python_module_name} PRIVATE -DUSE_ONEMKL_INTERFACES)
+else()
+    target_link_libraries(${python_module_name} PUBLIC MKL::MKL_SYCL::VM)
+endif()
 
 install(TARGETS ${python_module_name}
   DESTINATION "dpnp/backend/extensions/vm"

dpnp/backend/extensions/vm/vm_py.cpp

@@ -23,10 +23,11 @@
 // THE POSSIBILITY OF SUCH DAMAGE.
 //*****************************************************************************
 //
-// This file defines functions of dpnp.backend._lapack_impl extensions
+// This file defines functions of dpnp.backend._vm_impl extensions
 //
 //*****************************************************************************
 
+#if not defined(USE_ONEMKL_INTERFACES)
 #include "abs.hpp"
 #include "acos.hpp"
 #include "acosh.hpp"
@@ -67,9 +68,13 @@
 #include "trunc.hpp"
 
 namespace vm_ns = dpnp::extensions::vm;
+#endif // USE_ONEMKL_INTERFACES
+
+#include <pybind11/pybind11.h>
 
 PYBIND11_MODULE(_vm_impl, m)
 {
+#if not defined(USE_ONEMKL_INTERFACES)
     vm_ns::init_abs(m);
     vm_ns::init_acos(m);
     vm_ns::init_acosh(m);
@@ -108,4 +113,15 @@ PYBIND11_MODULE(_vm_impl, m)
     vm_ns::init_tan(m);
     vm_ns::init_tanh(m);
     vm_ns::init_trunc(m);
+#endif // USE_ONEMKL_INTERFACES
+    m.def(
+        "_is_available",
+        [](void) {
+#if defined(USE_ONEMKL_INTERFACES)
+            return false;
+#else
+            return true;
+#endif // USE_ONEMKL_INTERFACES
+        },
+        "Check if the OneMKL VM library can be used.");
 }