Add SYCL CUDA backend

AuroraPerego · Sep 11, 2024 · 94e3262 · 94e3262
1 parent 53392f4
commit 94e3262
Show file tree

Hide file tree

Showing 18 changed files with 243 additions and 5 deletions.
diff --git a/cmake/alpakaCommon.cmake b/cmake/alpakaCommon.cmake
@@ -661,6 +661,7 @@ if(alpaka_ACC_SYCL_ENABLE)
         cmake_dependent_option(alpaka_SYCL_ONEAPI_CPU "Enable oneAPI CPU targets for the SYCL back-end" OFF "alpaka_ACC_SYCL_ENABLE" OFF)
         cmake_dependent_option(alpaka_SYCL_ONEAPI_FPGA "Enable oneAPI FPGA targets for the SYCL back-end" OFF "alpaka_ACC_SYCL_ENABLE" OFF)
         cmake_dependent_option(alpaka_SYCL_ONEAPI_GPU "Enable oneAPI GPU targets for the SYCL back-end" OFF "alpaka_ACC_SYCL_ENABLE" OFF)
+        cmake_dependent_option(alpaka_SYCL_ONEAPI_GPU_NVIDIA "Enable NVIDIA GPU targets for the SYCL back-end" OFF "alpaka_ACC_SYCL_ENABLE" OFF)
         # Intel FPGA emulation / simulation
         if(alpaka_SYCL_ONEAPI_FPGA)
             set(alpaka_SYCL_ONEAPI_FPGA_MODE "emulation" CACHE STRING "Synthesis type for oneAPI FPGA targets")
@@ -686,7 +687,7 @@ if(alpaka_ACC_SYCL_ENABLE)
             list(APPEND alpaka_SYCL_TARGETS ${alpaka_SYCL_ONEAPI_FPGA_TARGET})
         endif()
 
-        if(alpaka_SYCL_ONEAPI_GPU)
+        if(alpaka_SYCL_ONEAPI_GPU OR alpaka_SYCL_ONEAPI_GPU_NVIDIA)
             list(APPEND alpaka_SYCL_TARGETS ${alpaka_SYCL_ONEAPI_GPU_TARGET})
         endif()
 
@@ -745,6 +746,18 @@ if(alpaka_ACC_SYCL_ENABLE)
             target_compile_definitions(alpaka INTERFACE "ALPAKA_SYCL_ONEAPI_GPU")
         endif()
 
+        if(alpaka_SYCL_ONEAPI_GPU_NVIDIA)
+            # Create a drop-down list (in cmake-gui) of valid Intel GPU targets. On the command line the user can specifiy
+            # additional targets, such as ranges: "Gen8-Gen12LP" or lists: "icllp;skl".
+            set(alpaka_SYCL_ONEAPI_GPU_DEVICES "nvidia_gpu_sm_89" CACHE STRING "NVIDIA GPU devices / generations to compile for")
+            set_property(CACHE alpaka_SYCL_ONEAPI_GPU_DEVICES
+                        PROPERTY STRINGS "nvidia_gpu_sm_50;nvidia_gpu_sm_52;nvidia_gpu_sm_53;nvidia_gpu_sm_60;nvidia_gpu_sm_61;nvidia_gpu_sm_62;nvidia_gpu_sm_70;nvidia_gpu_sm_72;nvidia_gpu_sm_75;nvidia_gpu_sm_80;nvidia_gpu_sm_86;nvidia_gpu_sm_87;nvidia_gpu_sm_89;nvidia_gpu_sm_90")
+            # If the user has given us a list turn all ';' into ',' to pacify the Intel OpenCL compiler.
+            string(REPLACE ";" "," alpaka_SYCL_ONEAPI_GPU_DEVICES "${alpaka_SYCL_ONEAPI_GPU_DEVICES}")
+
+            target_compile_definitions(alpaka INTERFACE "ALPAKA_SYCL_ONEAPI_GPU_NVIDIA")
+        endif()
+
         #-----------------------------------------------------------------------------------------------------------------
         # Generic SYCL options
         alpaka_set_compiler_options(DEVICE target alpaka "-fsycl-unnamed-lambda") # Compiler default but made explicit here
@@ -821,6 +834,9 @@ if(alpaka_ACC_SYCL_ENABLE)
     if(alpaka_SYCL_ONEAPI_GPU)
         target_compile_definitions(alpaka INTERFACE "ALPAKA_SYCL_TARGET_GPU")
     endif()
+    if(alpaka_SYCL_ONEAPI_GPU_NVIDIA)
+        target_compile_definitions(alpaka INTERFACE "ALPAKA_SYCL_TARGET_GPU_NVIDIA")
+    endif()
 
     message(STATUS alpaka_ACC_SYCL_ENABLED)
 endif()

diff --git a/include/alpaka/acc/AccGpuSyclNvidia.hpp b/include/alpaka/acc/AccGpuSyclNvidia.hpp
@@ -0,0 +1,38 @@
+/* Copyright 2024 Jan Stephan, Aurora Perego
+ * SPDX-License-Identifier: MPL-2.0
+ */
+
+#pragma once
+
+#include "alpaka/acc/AccGenericSycl.hpp"
+#include "alpaka/acc/Tag.hpp"
+#include "alpaka/core/Sycl.hpp"
+
+#if defined(ALPAKA_ACC_SYCL_ENABLED) && defined(ALPAKA_SYCL_ONEAPI_GPU_NVIDIA)
+
+namespace alpaka
+{
+    //! The Nvidia GPU SYCL accelerator.
+    //!
+    //! This accelerator allows parallel kernel execution on a oneAPI-capable Nvidia GPU target device.
+    template<typename TDim, typename TIdx>
+    using AccGpuSyclNvidia = AccGenericSycl<TagGpuSyclNvidia, TDim, TIdx>;
+
+    namespace trait
+    {
+        template<typename TDim, typename TIdx>
+        struct AccToTag<alpaka::AccGpuSyclNvidia<TDim, TIdx>>
+        {
+            using type = alpaka::TagGpuSyclNvidia;
+        };
+
+        template<typename TDim, typename TIdx>
+        struct TagToAcc<alpaka::TagGpuSyclNvidia, TDim, TIdx>
+        {
+            using type = alpaka::AccGpuSyclNvidia<TDim, TIdx>;
+        };
+    } // namespace trait
+
+} // namespace alpaka
+
+#endif
diff --git a/include/alpaka/acc/Tag.hpp b/include/alpaka/acc/Tag.hpp
@@ -31,6 +31,7 @@ namespace alpaka
     CREATE_ACC_TAG(TagGpuCudaRt);
     CREATE_ACC_TAG(TagGpuHipRt);
     CREATE_ACC_TAG(TagGpuSyclIntel);
+    CREATE_ACC_TAG(TagGpuSyclNvidia);
 
     namespace trait
     {
@@ -67,6 +68,7 @@ namespace alpaka
         alpaka::TagGpuHipRt,
         alpaka::TagCpuSycl,
         alpaka::TagFpgaSyclIntel,
-        alpaka::TagGpuSyclIntel>;
+        alpaka::TagGpuSyclIntel,
+        alpaka::TagGpuSyclNvidia>;
 
 } // namespace alpaka
diff --git a/include/alpaka/alpaka.hpp b/include/alpaka/alpaka.hpp
@@ -22,6 +22,7 @@
 #include "alpaka/acc/AccGpuCudaRt.hpp"
 #include "alpaka/acc/AccGpuHipRt.hpp"
 #include "alpaka/acc/AccGpuSyclIntel.hpp"
+#include "alpaka/acc/AccGpuSyclNvidia.hpp"
 #include "alpaka/acc/Tag.hpp"
 #include "alpaka/acc/TagAccIsEnabled.hpp"
 #include "alpaka/acc/Traits.hpp"
@@ -81,6 +82,7 @@
 #include "alpaka/dev/DevFpgaSyclIntel.hpp"
 #include "alpaka/dev/DevGenericSycl.hpp"
 #include "alpaka/dev/DevGpuSyclIntel.hpp"
+#include "alpaka/dev/DevGpuSyclNvidia.hpp"
 #include "alpaka/dev/DevHipRt.hpp"
 #include "alpaka/dev/Traits.hpp"
 #include "alpaka/dev/cpu/Wait.hpp"
@@ -95,6 +97,7 @@
 #include "alpaka/event/EventFpgaSyclIntel.hpp"
 #include "alpaka/event/EventGenericSycl.hpp"
 #include "alpaka/event/EventGpuSyclIntel.hpp"
+#include "alpaka/event/EventGpuSyclNvidia.hpp"
 #include "alpaka/event/EventHipRt.hpp"
 #include "alpaka/event/Traits.hpp"
 // exec
@@ -128,6 +131,7 @@
 #include "alpaka/kernel/TaskKernelGpuCudaRt.hpp"
 #include "alpaka/kernel/TaskKernelGpuHipRt.hpp"
 #include "alpaka/kernel/TaskKernelGpuSyclIntel.hpp"
+#include "alpaka/kernel/TaskKernelGpuSyclNvidia.hpp"
 #include "alpaka/kernel/Traits.hpp"
 // math
 #include "alpaka/math/Complex.hpp"
@@ -144,6 +148,7 @@
 #include "alpaka/mem/buf/BufFpgaSyclIntel.hpp"
 #include "alpaka/mem/buf/BufGenericSycl.hpp"
 #include "alpaka/mem/buf/BufGpuSyclIntel.hpp"
+#include "alpaka/mem/buf/BufGpuSyclNvidia.hpp"
 #include "alpaka/mem/buf/BufHipRt.hpp"
 #include "alpaka/mem/buf/Traits.hpp"
 #include "alpaka/mem/fence/MemFenceCpu.hpp"
@@ -189,6 +194,7 @@
 #include "alpaka/platform/PlatformCudaRt.hpp"
 #include "alpaka/platform/PlatformFpgaSyclIntel.hpp"
 #include "alpaka/platform/PlatformGpuSyclIntel.hpp"
+#include "alpaka/platform/PlatformGpuSyclNvidia.hpp"
 #include "alpaka/platform/PlatformHipRt.hpp"
 #include "alpaka/platform/Traits.hpp"
 // rand
@@ -212,6 +218,8 @@
 #include "alpaka/queue/QueueFpgaSyclIntelNonBlocking.hpp"
 #include "alpaka/queue/QueueGpuSyclIntelBlocking.hpp"
 #include "alpaka/queue/QueueGpuSyclIntelNonBlocking.hpp"
+#include "alpaka/queue/QueueGpuSyclNvidiaBlocking.hpp"
+#include "alpaka/queue/QueueGpuSyclNvidiaNonBlocking.hpp"
 #include "alpaka/queue/QueueHipRtBlocking.hpp"
 #include "alpaka/queue/QueueHipRtNonBlocking.hpp"
 #include "alpaka/queue/Traits.hpp"

diff --git a/include/alpaka/dev/DevGpuSyclNvidia.hpp b/include/alpaka/dev/DevGpuSyclNvidia.hpp
@@ -0,0 +1,17 @@
+/* Copyright 2024 Jan Stephan, Aurora Perego
+ * SPDX-License-Identifier: MPL-2.0
+ */
+
+#pragma once
+
+#include "alpaka/acc/Tag.hpp"
+#include "alpaka/dev/DevGenericSycl.hpp"
+
+#if defined(ALPAKA_ACC_SYCL_ENABLED) && defined(ALPAKA_SYCL_ONEAPI_GPU_NVIDIA)
+
+namespace alpaka
+{
+    using DevGpuSyclNvidia = DevGenericSycl<TagGpuSyclNvidia>;
+} // namespace alpaka
+
+#endif
diff --git a/include/alpaka/event/EventGpuSyclNvidia.hpp b/include/alpaka/event/EventGpuSyclNvidia.hpp
@@ -0,0 +1,17 @@
+/* Copyright 2024 Jan Stephan, Aurora Perego
+ * SPDX-License-Identifier: MPL-2.0
+ */
+
+#pragma once
+
+#include "alpaka/acc/Tag.hpp"
+#include "alpaka/event/EventGenericSycl.hpp"
+
+#if defined(ALPAKA_ACC_SYCL_ENABLED) && defined(ALPAKA_SYCL_ONEAPI_GPU_NVIDIA)
+
+namespace alpaka
+{
+    using EventGpuSyclNvidia = EventGenericSycl<TagGpuSyclNvidia>;
+} // namespace alpaka
+
+#endif
diff --git a/include/alpaka/example/ExampleDefaultAcc.hpp b/include/alpaka/example/ExampleDefaultAcc.hpp
@@ -31,6 +31,8 @@ namespace alpaka
     using ExampleDefaultAcc = alpaka::AccFpgaSyclIntel<TDim, TIdx>;
 #    elif defined(ALPAKA_SYCL_ONEAPI_GPU)
     using ExampleDefaultAcc = alpaka::AccGpuSyclIntel<TDim, TIdx>;
+#    elif defined(ALPAKA_SYCL_ONEAPI_GPU_NVIDIA)
+    using ExampleDefaultAcc = alpaka::AccGpuSyclNvidia<TDim, TIdx>;
 #    endif
 #elif defined(ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED)
     using ExampleDefaultAcc = alpaka::AccCpuSerial<TDim, TIdx>;

diff --git a/include/alpaka/kernel/TaskKernelGpuSyclNvidia.hpp b/include/alpaka/kernel/TaskKernelGpuSyclNvidia.hpp
@@ -0,0 +1,20 @@
+/* Copyright 2024 Jan Stephan, Aurora Perego
+ * SPDX-License-Identifier: MPL-2.0
+ */
+
+#pragma once
+
+#include "alpaka/acc/Tag.hpp"
+#include "alpaka/kernel/TaskKernelGenericSycl.hpp"
+
+#if defined(ALPAKA_ACC_SYCL_ENABLED) && defined(ALPAKA_SYCL_ONEAPI_GPU_NVIDIA)
+
+namespace alpaka
+{
+    template<typename TDim, typename TIdx, typename TKernelFnObj, typename... TArgs>
+    using TaskKernelGpuSyclNvidia
+        = TaskKernelGenericSycl<TagGpuSyclNvidia, AccGpuSyclNvidia<TDim, TIdx>, TDim, TIdx, TKernelFnObj, TArgs...>;
+
+} // namespace alpaka
+
+#endif
diff --git a/include/alpaka/mem/buf/BufGpuSyclNvidia.hpp b/include/alpaka/mem/buf/BufGpuSyclNvidia.hpp
@@ -0,0 +1,19 @@
+/* Copyright 2024 Jan Stephan, Luca Ferragina, Aurora Perego
+ * SPDX-License-Identifier: MPL-2.0
+ */
+
+#pragma once
+
+#include "alpaka/dev/DevGpuSyclNvidia.hpp"
+#include "alpaka/mem/buf/BufGenericSycl.hpp"
+#include "alpaka/platform/PlatformGpuSyclNvidia.hpp"
+
+#if defined(ALPAKA_ACC_SYCL_ENABLED) && defined(ALPAKA_SYCL_ONEAPI_GPU_NVIDIA)
+
+namespace alpaka
+{
+    template<typename TElem, typename TDim, typename TIdx>
+    using BufGpuSyclNvidia = BufGenericSycl<TElem, TDim, TIdx, PlatformGpuSyclNvidia>;
+} // namespace alpaka
+
+#endif
diff --git a/include/alpaka/mem/global/DeviceGlobalGenericSycl.hpp b/include/alpaka/mem/global/DeviceGlobalGenericSycl.hpp
@@ -30,6 +30,13 @@ namespace alpaka
             using Type = sycl::ext::oneapi::experimental::device_global<T>;
         };
 
+        template<typename T>
+        struct DevGlobalTrait<TagGpuSyclNvidia, T>
+        {
+            // SYCL GPU implementation
+            using Type = sycl::ext::oneapi::experimental::device_global<T>;
+        };
+
         template<typename T>
         struct DevGlobalTrait<TagFpgaSyclIntel, T>
         {

diff --git a/include/alpaka/platform/PlatformGpuSyclNvidia.hpp b/include/alpaka/platform/PlatformGpuSyclNvidia.hpp
@@ -0,0 +1,36 @@
+/* Copyright 2024 Jan Stephan, Luca Ferragina, Andrea Bocci, Aurora Perego
+ * SPDX-License-Identifier: MPL-2.0
+ */
+
+#pragma once
+
+#include "alpaka/dev/DevGenericSycl.hpp"
+#include "alpaka/dev/Traits.hpp"
+#include "alpaka/platform/PlatformGenericSycl.hpp"
+
+#if defined(ALPAKA_ACC_SYCL_ENABLED) && defined(ALPAKA_SYCL_ONEAPI_GPU_NVIDIA)
+
+#    include <sycl/sycl.hpp>
+
+namespace alpaka
+{
+    namespace detail
+    {
+        template<>
+        struct SYCLDeviceSelector<TagGpuSyclNvidia>
+        {
+            auto operator()(sycl::device const& dev) const -> int
+            {
+                auto const& vendor = dev.get_info<sycl::info::device::vendor>();
+                auto const is_intel_gpu = dev.is_gpu() && (vendor.find("NVIDIA") != std::string::npos);
+
+                return is_intel_gpu ? 1 : -1;
+            }
+        };
+    } // namespace detail
+
+    //! The SYCL device manager.
+    using PlatformGpuSyclNvidia = PlatformGenericSycl<TagGpuSyclNvidia>;
+} // namespace alpaka
+
+#endif
diff --git a/include/alpaka/queue/QueueGpuSyclNvidiaBlocking.hpp b/include/alpaka/queue/QueueGpuSyclNvidiaBlocking.hpp
@@ -0,0 +1,17 @@
+/* Copyright 2024 Jan Stephan, Aurora Perego
+ * SPDX-License-Identifier: MPL-2.0
+ */
+
+#pragma once
+
+#include "alpaka/acc/Tag.hpp"
+#include "alpaka/queue/QueueGenericSyclBlocking.hpp"
+
+#if defined(ALPAKA_ACC_SYCL_ENABLED) && defined(ALPAKA_SYCL_ONEAPI_GPU_NVIDIA)
+
+namespace alpaka
+{
+    using QueueGpuSyclNvidiaBlocking = QueueGenericSyclBlocking<TagGpuSyclNvidia>;
+} // namespace alpaka
+
+#endif
diff --git a/include/alpaka/queue/QueueGpuSyclNvidiaNonBlocking.hpp b/include/alpaka/queue/QueueGpuSyclNvidiaNonBlocking.hpp
@@ -0,0 +1,17 @@
+/* Copyright 2024 Jan Stephan, Aurora Perego
+ * SPDX-License-Identifier: MPL-2.0
+ */
+
+#pragma once
+
+#include "alpaka/acc/Tag.hpp"
+#include "alpaka/queue/QueueGenericSyclNonBlocking.hpp"
+
+#if defined(ALPAKA_ACC_SYCL_ENABLED) && defined(ALPAKA_SYCL_ONEAPI_GPU_NVIDIA)
+
+namespace alpaka
+{
+    using QueueGpuSyclNvidiaNonBlocking = QueueGenericSyclNonBlocking<TagGpuSyclNvidia>;
+} // namespace alpaka
+
+#endif
diff --git a/include/alpaka/test/acc/TestAccs.hpp b/include/alpaka/test/acc/TestAccs.hpp
@@ -99,6 +99,13 @@ namespace alpaka::test
         template<typename TDim, typename TIdx>
         using AccGpuSyclIntelIfAvailableElseInt = int;
 #endif
+#if defined(ALPAKA_ACC_SYCL_ENABLED) && defined(ALPAKA_SYCL_TARGET_GPU_NVIDIA)
+        template<typename TDim, typename TIdx>
+        using AccGpuSyclNvidiaIfAvailableElseInt = AccGpuSyclNvidia<TDim, TIdx>;
+#else
+        template<typename TDim, typename TIdx>
+        using AccGpuSyclNvidiaIfAvailableElseInt = int;
+#endif
 
         //! A vector containing all available accelerators and int's.
         template<typename TDim, typename TIdx>
@@ -112,7 +119,8 @@ namespace alpaka::test
             AccGpuHipRtIfAvailableElseInt<TDim, TIdx>,
             AccCpuSyclIfAvailableElseInt<TDim, TIdx>,
             AccFpgaSyclIntelIfAvailableElseInt<TDim, TIdx>,
-            AccGpuSyclIntelIfAvailableElseInt<TDim, TIdx>>;
+            AccGpuSyclIntelIfAvailableElseInt<TDim, TIdx>,
+            AccGpuSyclNvidiaIfAvailableElseInt<TDim, TIdx>>;
     } // namespace detail
 
     //! A vector containing all available accelerators.

diff --git a/include/alpaka/test/queue/Queue.hpp b/include/alpaka/test/queue/Queue.hpp
@@ -141,6 +141,11 @@ namespace alpaka::test
         std::tuple<alpaka::DevGpuSyclIntel, alpaka::QueueGpuSyclIntelBlocking>,
         std::tuple<alpaka::DevGpuSyclIntel, alpaka::QueueGpuSyclIntelNonBlocking>
 #    endif
+#    ifdef ALPAKA_SYCL_ONEAPI_GPU_NVIDIA
+        ,
+        std::tuple<alpaka::DevGpuSyclNvidia, alpaka::QueueGpuSyclNvidiaBlocking>,
+        std::tuple<alpaka::DevGpuSyclNvidia, alpaka::QueueGpuSyclNvidiaNonBlocking>
+#    endif
 #endif
         >;
 } // namespace alpaka::test
diff --git a/test/unit/acc/src/AccTagTest.cpp b/test/unit/acc/src/AccTagTest.cpp
@@ -29,7 +29,8 @@ using TagList = std::tuple<
     alpaka::TagGpuHipRt,
     alpaka::TagCpuSycl,
     alpaka::TagFpgaSyclIntel,
-    alpaka::TagGpuSyclIntel>;
+    alpaka::TagGpuSyclIntel,
+    alpaka::TagGpuSyclNvidia>;
 
 using AccToTagMap = std::tuple<
     std::pair<alpaka::test::detail::AccCpuSerialIfAvailableElseInt<Dim, Idx>, alpaka::TagCpuSerial>,
@@ -41,7 +42,8 @@ using AccToTagMap = std::tuple<
     std::pair<alpaka::test::detail::AccGpuHipRtIfAvailableElseInt<Dim, Idx>, alpaka::TagGpuHipRt>,
     std::pair<alpaka::test::detail::AccCpuSyclIfAvailableElseInt<Dim, Idx>, alpaka::TagCpuSycl>,
     std::pair<alpaka::test::detail::AccFpgaSyclIntelIfAvailableElseInt<Dim, Idx>, alpaka::TagFpgaSyclIntel>,
-    std::pair<alpaka::test::detail::AccGpuSyclIntelIfAvailableElseInt<Dim, Idx>, alpaka::TagGpuSyclIntel>>;
+    std::pair<alpaka::test::detail::AccGpuSyclIntelIfAvailableElseInt<Dim, Idx>, alpaka::TagGpuSyclIntel>,
+    std::pair<alpaka::test::detail::AccGpuSyclNvidiaIfAvailableElseInt<Dim, Idx>, alpaka::TagGpuSyclNvidia>>;
 
 using AccTagTestMatrix = alpaka::meta::CartesianProduct<std::tuple, AccToTagMap, TagList>;