Implement build for NVIDIA GPU

CMakeLists.txt

@@ -67,6 +67,51 @@ 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
+)
+set(CMAKE_INTERPROCEDURAL_OPTIMIZATION OFF)
+set(_dpnp_sycl_targets)
+set(_dpnp_use_onemkl OFF)
+if ("x${DPNP_SYCL_TARGETS}" STREQUAL "x")
+   if(DPNP_TARGET_CUDA)
+      set(_dpnp_use_onemkl ON)
+      set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DDPNP_TARGET_CUDA=1")
+      set(_dpnp_sycl_targets "nvptx64-nvidia-cuda,spir64-unknown-unknown")
+   else()
+      if(DEFINED ENV{DPNP_TARGET_CUDA})
+          set(_dpnp_use_onemkl ON)
+          set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DDPNP_TARGET_CUDA=1")
+          set(_dpnp_sycl_targets "nvptx64-nvidia-cuda,spir64-unknown-unknown")
+      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()
+
+if(_dpnp_use_onemkl)
+    set(BUILD_FUNCTIONAL_TESTS False)
+    set(BUILD_EXAMPLES False)
+    set(ENABLE_CUBLAS_BACKEND False)
+    set(ENABLE_CUSOLVER_BACKEND True)
+    set(ENABLE_CUFFT_BACKEND True)
+    set(ENABLE_CURAND_BACKEND False)
+    set(ENABLE_MKLGPU_BACKEND False)
+    set(ENABLE_MKLCPU_BACKEND False)
+    FetchContent_Declare(
+            onemkl_interface_library
+            GIT_REPOSITORY https://github.com/oneapi-src/oneMKL.git
+            GIT_TAG develop
+    )
+    FetchContent_MakeAvailable(onemkl_interface_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,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::DFT)
+if(_dpnp_use_onemkl)
+    target_link_libraries(${python_module_name} PUBLIC onemkl)
+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

@@ -187,26 +187,32 @@ class DescriptorWrapper
     // config_param::PLACEMENT
     bool get_in_place()
     {
+#if DPNP_TARGET_CUDA
+        mkl_dft::config_value placement;
+        descr_.get_value(mkl_dft::config_param::PLACEMENT, &placement);
+        return (placement == mkl_dft::config_value::INPLACE);
+#else
         // TODO: replace when MKLD-10506 is implemented
-        // mkl_dft::config_value placement;
         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
     }
 
     void set_in_place(const bool &in_place_request)
     {
+#if DPNP_TARGET_CUDA
+        descr_.set_value(mkl_dft::config_param::PLACEMENT,
+                         (in_place_request)
+                             ? mkl_dft::config_value::INPLACE
+                             : mkl_dft::config_value::NOT_INPLACE);
+#else
         // 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);
         descr_.set_value(mkl_dft::config_param::PLACEMENT,
                          (in_place_request)
                              ? DFTI_CONFIG_VALUE::DFTI_INPLACE
                              : DFTI_CONFIG_VALUE::DFTI_NOT_INPLACE);
+#endif
     }
 
     // config_param::PRECISION
@@ -221,14 +227,16 @@ class DescriptorWrapper
     // config_param::COMMIT_STATUS
     bool is_committed()
     {
+#if DPNP_TARGET_CUDA
+        mkl_dft::config_value committed;
+        descr_.get_value(mkl_dft::config_param::COMMIT_STATUS, &committed);
+        return (committed == mkl_dft::config_value::COMMITTED);
+#else
         // TODO: replace when MKLD-10506 is implemented
-        // mkl_dft::config_value committed;
         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
     }
 
 private:

dpnp/backend/extensions/fft/out_of_place.cpp

@@ -23,7 +23,10 @@
 // THE POSSIBILITY OF SUCH DAMAGE.
 //*****************************************************************************
 
+#if !DPNP_TARGET_CUDA
 #include <oneapi/mkl/dfti.hpp>
+#endif
+
 #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,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::LAPACK)
+if(_dpnp_use_onemkl)
+    target_link_libraries(${python_module_name} PUBLIC onemkl)
+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 DPNP_TARGET_CUDA
+    // 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
 }
 
 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 _dpnp_use_onemkl)
+    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,9 @@ 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(NOT _dpnp_use_onemkl)
+    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 !DPNP_TARGET_CUDA
 #include "abs.hpp"
 #include "acos.hpp"
 #include "acosh.hpp"
@@ -67,9 +68,22 @@
 #include "trunc.hpp"
 
 namespace vm_ns = dpnp::extensions::vm;
+#endif
+
+#include <pybind11/pybind11.h>
+
+bool mkl_vm_is_defined()
+{
+#if DPNP_TARGET_CUDA
+    return false;
+#else
+    return true;
+#endif
+}
 
 PYBIND11_MODULE(_vm_impl, m)
 {
+#if !DPNP_TARGET_CUDA
     vm_ns::init_abs(m);
     vm_ns::init_acos(m);
     vm_ns::init_acosh(m);
@@ -108,4 +122,7 @@ PYBIND11_MODULE(_vm_impl, m)
     vm_ns::init_tan(m);
     vm_ns::init_tanh(m);
     vm_ns::init_trunc(m);
+#endif
+    m.def("mkl_vm_is_defined", mkl_vm_is_defined,
+          "Check if the OneMKL VM library can be used.");
 }

dpnp/dpnp_algo/dpnp_elementwise_common.py

@@ -32,6 +32,7 @@
 )
 
 import dpnp
+import dpnp.backend.extensions.vm._vm_impl as vmi
 from dpnp.dpnp_array import dpnp_array
 
 __all__ = [
@@ -111,11 +112,12 @@ def _call_func(src, dst, sycl_queue, depends=None):
             if depends is None:
                 depends = []
 
-            if mkl_fn_to_call is not None and mkl_fn_to_call(
-                sycl_queue, src, dst
-            ):
-                # call pybind11 extension for unary function from OneMKL VM
-                return mkl_impl_fn(sycl_queue, src, dst, depends)
+            if vmi.mkl_vm_is_defined() and mkl_fn_to_call is not None:
+                if getattr(vmi, mkl_fn_to_call)(sycl_queue, src, dst):
+                    # call pybind11 extension for unary function from OneMKL VM
+                    return getattr(vmi, mkl_impl_fn)(
+                        sycl_queue, src, dst, depends
+                    )
             return unary_dp_impl_fn(src, dst, sycl_queue, depends)
 
         super().__init__(
@@ -264,11 +266,12 @@ def _call_func(src1, src2, dst, sycl_queue, depends=None):
             if depends is None:
                 depends = []
 
-            if mkl_fn_to_call is not None and mkl_fn_to_call(
-                sycl_queue, src1, src2, dst
-            ):
-                # call pybind11 extension for binary function from OneMKL VM
-                return mkl_impl_fn(sycl_queue, src1, src2, dst, depends)
+            if vmi.mkl_vm_is_defined() and mkl_fn_to_call is not None:
+                if getattr(vmi, mkl_fn_to_call)(sycl_queue, src1, src2, dst):
+                    # call pybind11 extension for binary function from OneMKL VM
+                    return getattr(vmi, mkl_impl_fn)(
+                        sycl_queue, src1, src2, dst, depends
+                    )
             return binary_dp_impl_fn(src1, src2, dst, sycl_queue, depends)
 
         super().__init__(