diff --git a/cmake/alpakaCommon.cmake b/cmake/alpakaCommon.cmake index a9a13f883b1e..28f58b7933af 100644 --- 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 new file mode 100644 index 000000000000..3473559ee9c5 --- /dev/null +++ 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 + using AccGpuSyclNvidia = AccGenericSycl; + + namespace trait + { + template + struct AccToTag> + { + using type = alpaka::TagGpuSyclNvidia; + }; + + template + struct TagToAcc + { + using type = alpaka::AccGpuSyclNvidia; + }; + } // namespace trait + +} // namespace alpaka + +#endif diff --git a/include/alpaka/acc/Tag.hpp b/include/alpaka/acc/Tag.hpp index f7880afd6f15..b6e85cdce077 100644 --- 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 index dd99f7d741b7..ca5bd54b31ea 100644 --- 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 new file mode 100644 index 000000000000..264ce96595e9 --- /dev/null +++ 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; +} // namespace alpaka + +#endif diff --git a/include/alpaka/event/EventGpuSyclNvidia.hpp b/include/alpaka/event/EventGpuSyclNvidia.hpp new file mode 100644 index 000000000000..5f550a271721 --- /dev/null +++ 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; +} // namespace alpaka + +#endif diff --git a/include/alpaka/example/ExampleDefaultAcc.hpp b/include/alpaka/example/ExampleDefaultAcc.hpp index 22f77f953973..f90c0b7fc713 100644 --- a/include/alpaka/example/ExampleDefaultAcc.hpp +++ b/include/alpaka/example/ExampleDefaultAcc.hpp @@ -31,6 +31,8 @@ namespace alpaka using ExampleDefaultAcc = alpaka::AccFpgaSyclIntel; # elif defined(ALPAKA_SYCL_ONEAPI_GPU) using ExampleDefaultAcc = alpaka::AccGpuSyclIntel; +# elif defined(ALPAKA_SYCL_ONEAPI_GPU_NVIDIA) + using ExampleDefaultAcc = alpaka::AccGpuSyclNvidia; # endif #elif defined(ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED) using ExampleDefaultAcc = alpaka::AccCpuSerial; diff --git a/include/alpaka/kernel/TaskKernelGpuSyclNvidia.hpp b/include/alpaka/kernel/TaskKernelGpuSyclNvidia.hpp new file mode 100644 index 000000000000..9df32cb2c612 --- /dev/null +++ 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 + using TaskKernelGpuSyclNvidia + = TaskKernelGenericSycl, TDim, TIdx, TKernelFnObj, TArgs...>; + +} // namespace alpaka + +#endif diff --git a/include/alpaka/mem/buf/BufGpuSyclNvidia.hpp b/include/alpaka/mem/buf/BufGpuSyclNvidia.hpp new file mode 100644 index 000000000000..3d52ecb75cc2 --- /dev/null +++ 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 + using BufGpuSyclNvidia = BufGenericSycl; +} // namespace alpaka + +#endif diff --git a/include/alpaka/mem/global/DeviceGlobalGenericSycl.hpp b/include/alpaka/mem/global/DeviceGlobalGenericSycl.hpp index 56ee98c885e8..7112620aa15c 100644 --- 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; }; + template + struct DevGlobalTrait + { + // SYCL GPU implementation + using Type = sycl::ext::oneapi::experimental::device_global; + }; + template struct DevGlobalTrait { diff --git a/include/alpaka/platform/PlatformGpuSyclNvidia.hpp b/include/alpaka/platform/PlatformGpuSyclNvidia.hpp new file mode 100644 index 000000000000..cf1e5c174dc5 --- /dev/null +++ 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 + +namespace alpaka +{ + namespace detail + { + template<> + struct SYCLDeviceSelector + { + auto operator()(sycl::device const& dev) const -> int + { + auto const& vendor = dev.get_info(); + 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; +} // namespace alpaka + +#endif diff --git a/include/alpaka/queue/QueueGpuSyclNvidiaBlocking.hpp b/include/alpaka/queue/QueueGpuSyclNvidiaBlocking.hpp new file mode 100644 index 000000000000..e877fdb3fe08 --- /dev/null +++ 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; +} // namespace alpaka + +#endif diff --git a/include/alpaka/queue/QueueGpuSyclNvidiaNonBlocking.hpp b/include/alpaka/queue/QueueGpuSyclNvidiaNonBlocking.hpp new file mode 100644 index 000000000000..a735d58fb0ba --- /dev/null +++ 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; +} // namespace alpaka + +#endif diff --git a/include/alpaka/test/acc/TestAccs.hpp b/include/alpaka/test/acc/TestAccs.hpp index 2370fa42686b..4b1e4f761021 100644 --- a/include/alpaka/test/acc/TestAccs.hpp +++ b/include/alpaka/test/acc/TestAccs.hpp @@ -99,6 +99,13 @@ namespace alpaka::test template using AccGpuSyclIntelIfAvailableElseInt = int; #endif +#if defined(ALPAKA_ACC_SYCL_ENABLED) && defined(ALPAKA_SYCL_TARGET_GPU_NVIDIA) + template + using AccGpuSyclNvidiaIfAvailableElseInt = AccGpuSyclNvidia; +#else + template + using AccGpuSyclNvidiaIfAvailableElseInt = int; +#endif //! A vector containing all available accelerators and int's. template @@ -112,7 +119,8 @@ namespace alpaka::test AccGpuHipRtIfAvailableElseInt, AccCpuSyclIfAvailableElseInt, AccFpgaSyclIntelIfAvailableElseInt, - AccGpuSyclIntelIfAvailableElseInt>; + AccGpuSyclIntelIfAvailableElseInt, + AccGpuSyclNvidiaIfAvailableElseInt>; } // namespace detail //! A vector containing all available accelerators. diff --git a/include/alpaka/test/queue/Queue.hpp b/include/alpaka/test/queue/Queue.hpp index 0518e6d41e41..a2d8c90687c8 100644 --- a/include/alpaka/test/queue/Queue.hpp +++ b/include/alpaka/test/queue/Queue.hpp @@ -141,6 +141,11 @@ namespace alpaka::test std::tuple, std::tuple # endif +# ifdef ALPAKA_SYCL_ONEAPI_GPU_NVIDIA + , + std::tuple, + std::tuple +# endif #endif >; } // namespace alpaka::test diff --git a/test/unit/acc/src/AccTagTest.cpp b/test/unit/acc/src/AccTagTest.cpp index ac798bc48f61..6a91e5f8d0a6 100644 --- 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::TagCpuSerial>, @@ -41,7 +42,8 @@ using AccToTagMap = std::tuple< std::pair, alpaka::TagGpuHipRt>, std::pair, alpaka::TagCpuSycl>, std::pair, alpaka::TagFpgaSyclIntel>, - std::pair, alpaka::TagGpuSyclIntel>>; + std::pair, alpaka::TagGpuSyclIntel>, + std::pair, alpaka::TagGpuSyclNvidia>>; using AccTagTestMatrix = alpaka::meta::CartesianProduct; diff --git a/test/unit/mem/view/src/ViewStaticAccMem.cpp b/test/unit/mem/view/src/ViewStaticAccMem.cpp index a42a37c140d7..d00e643a1fb5 100644 --- a/test/unit/mem/view/src/ViewStaticAccMem.cpp +++ b/test/unit/mem/view/src/ViewStaticAccMem.cpp @@ -66,6 +66,12 @@ using EnabledAccsElseInt = std::tuple>; template using EnabledAccs = typename alpaka::meta::Filter, std::is_class>; using TestAccs = EnabledAccs; +#elif defined(ALPAKA_ACC_SYCL_ENABLED) && defined(ALPAKA_SYCL_TARGET_GPU_NVIDIA) +template +using EnabledAccsElseInt = std::tuple>; +template +using EnabledAccs = typename alpaka::meta::Filter, std::is_class>; +using TestAccs = EnabledAccs; #else using TestAccs = alpaka::test::EnabledAccs; #endif diff --git a/test/unit/warp/src/Activemask.cpp b/test/unit/warp/src/Activemask.cpp index 223f3535f670..baae5dc4afd9 100644 --- a/test/unit/warp/src/Activemask.cpp +++ b/test/unit/warp/src/Activemask.cpp @@ -69,6 +69,7 @@ TEMPLATE_LIST_TEST_CASE("activemask", "[warp]", alpaka::test::TestAccs) Acc, alpaka::TagCpuSycl, alpaka::TagGpuSyclIntel, + alpaka::TagGpuSyclNvidia, alpaka::TagFpgaSyclIntel, alpaka::TagGenericSycl>) {