Add SYCL AMD backend

cmake/alpakaCommon.cmake

@@ -662,6 +662,7 @@ if(alpaka_ACC_SYCL_ENABLE)
         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)
+        cmake_dependent_option(alpaka_SYCL_ONEAPI_GPU_AMD "Enable AMD 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")
@@ -687,7 +688,7 @@ if(alpaka_ACC_SYCL_ENABLE)
             list(APPEND alpaka_SYCL_TARGETS ${alpaka_SYCL_ONEAPI_FPGA_TARGET})
         endif()
 
-        if(alpaka_SYCL_ONEAPI_GPU OR alpaka_SYCL_ONEAPI_GPU_NVIDIA)
+        if(alpaka_SYCL_ONEAPI_GPU OR alpaka_SYCL_ONEAPI_GPU_NVIDIA OR alpaka_SYCL_ONEAPI_GPU_AMD)
             list(APPEND alpaka_SYCL_TARGETS ${alpaka_SYCL_ONEAPI_GPU_TARGET})
         endif()
 
@@ -758,6 +759,18 @@ if(alpaka_ACC_SYCL_ENABLE)
             target_compile_definitions(alpaka INTERFACE "ALPAKA_SYCL_ONEAPI_GPU_NVIDIA")
         endif()
 
+        if(alpaka_SYCL_ONEAPI_GPU_AMD)
+            # 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 "amd_gpu_gfx1200" CACHE STRING "AMD GPU devices / generations to compile for")
+            set_property(CACHE alpaka_SYCL_ONEAPI_GPU_DEVICES
+                        PROPERTY STRINGS "amd_gpu_gfx700;amd_gpu_gfx701;amd_gpu_gfx702;amd_gpu_gfx801;amd_gpu_gfx802;amd_gpu_gfx803;amd_gpu_gfx805;amd_gpu_gfx810;amd_gpu_gfx900;amd_gpu_gfx902;amd_gpu_gfx904;amd_gpu_gfx906;amd_gpu_gfx908;amd_gpu_gfx909;amd_gpu_gfx90a;amd_gpu_gfx90c;amd_gpu_gfx940;amd_gpu_gfx941;amd_gpu_gfx942;amd_gpu_gfx1010;amd_gpu_gfx1011;amd_gpu_gfx1012;amd_gpu_gfx1013;amd_gpu_gfx1030;amd_gpu_gfx1031;amd_gpu_gfx1032;amd_gpu_gfx1033;amd_gpu_gfx1034;amd_gpu_gfx1035;amd_gpu_gfx1036;amd_gpu_gfx1100;amd_gpu_gfx1101;amd_gpu_gfx1102;amd_gpu_gfx1103;amd_gpu_gfx1150;amd_gpu_gfx1151;amd_gpu_gfx1200;amd_gpu_gfx1201")
+            # 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_AMD")
+        endif()
+
         #-----------------------------------------------------------------------------------------------------------------
         # Generic SYCL options
         alpaka_set_compiler_options(DEVICE target alpaka "-fsycl-unnamed-lambda") # Compiler default but made explicit here
@@ -837,6 +850,9 @@ if(alpaka_ACC_SYCL_ENABLE)
     if(alpaka_SYCL_ONEAPI_GPU_NVIDIA)
         target_compile_definitions(alpaka INTERFACE "ALPAKA_SYCL_TARGET_GPU_NVIDIA")
     endif()
+    if(alpaka_SYCL_ONEAPI_GPU_AMD)
+        target_compile_definitions(alpaka INTERFACE "ALPAKA_SYCL_TARGET_GPU_AMD")
+    endif()
 
     message(STATUS alpaka_ACC_SYCL_ENABLED)
 endif()

include/alpaka/acc/AccGpuSyclAmd.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_AMD)
+
+namespace alpaka
+{
+    //! The Amd GPU SYCL accelerator.
+    //!
+    //! This accelerator allows parallel kernel execution on a oneAPI-capable Amd GPU target device.
+    template<typename TDim, typename TIdx>
+    using AccGpuSyclAmd = AccGenericSycl<TagGpuSyclAmd, TDim, TIdx>;
+
+    namespace trait
+    {
+        template<typename TDim, typename TIdx>
+        struct AccToTag<alpaka::AccGpuSyclAmd<TDim, TIdx>>
+        {
+            using type = alpaka::TagGpuSyclAmd;
+        };
+
+        template<typename TDim, typename TIdx>
+        struct TagToAcc<alpaka::TagGpuSyclAmd, TDim, TIdx>
+        {
+            using type = alpaka::AccGpuSyclAmd<TDim, TIdx>;
+        };
+    } // namespace trait
+
+} // namespace alpaka
+
+#endif

include/alpaka/acc/Tag.hpp

@@ -32,6 +32,7 @@ namespace alpaka
     CREATE_ACC_TAG(TagGpuHipRt);
     CREATE_ACC_TAG(TagGpuSyclIntel);
     CREATE_ACC_TAG(TagGpuSyclNvidia);
+    CREATE_ACC_TAG(TagGpuSyclAmd);
 
     namespace trait
     {
@@ -69,6 +70,7 @@ namespace alpaka
         alpaka::TagCpuSycl,
         alpaka::TagFpgaSyclIntel,
         alpaka::TagGpuSyclIntel,
-        alpaka::TagGpuSyclNvidia>;
+        alpaka::TagGpuSyclNvidia,
+        alpaka::TagGpuSyclAmd>;
 
 } // namespace alpaka

include/alpaka/alpaka.hpp

@@ -23,6 +23,7 @@
 #include "alpaka/acc/AccGpuHipRt.hpp"
 #include "alpaka/acc/AccGpuSyclIntel.hpp"
 #include "alpaka/acc/AccGpuSyclNvidia.hpp"
+#include "alpaka/acc/AccGpuSyclAmd.hpp"
 #include "alpaka/acc/Tag.hpp"
 #include "alpaka/acc/TagAccIsEnabled.hpp"
 #include "alpaka/acc/Traits.hpp"
@@ -83,6 +84,7 @@
 #include "alpaka/dev/DevGenericSycl.hpp"
 #include "alpaka/dev/DevGpuSyclIntel.hpp"
 #include "alpaka/dev/DevGpuSyclNvidia.hpp"
+#include "alpaka/dev/DevGpuSyclAmd.hpp"
 #include "alpaka/dev/DevHipRt.hpp"
 #include "alpaka/dev/Traits.hpp"
 #include "alpaka/dev/cpu/Wait.hpp"
@@ -98,6 +100,7 @@
 #include "alpaka/event/EventGenericSycl.hpp"
 #include "alpaka/event/EventGpuSyclIntel.hpp"
 #include "alpaka/event/EventGpuSyclNvidia.hpp"
+#include "alpaka/event/EventGpuSyclAmd.hpp"
 #include "alpaka/event/EventHipRt.hpp"
 #include "alpaka/event/Traits.hpp"
 // exec
@@ -132,6 +135,7 @@
 #include "alpaka/kernel/TaskKernelGpuHipRt.hpp"
 #include "alpaka/kernel/TaskKernelGpuSyclIntel.hpp"
 #include "alpaka/kernel/TaskKernelGpuSyclNvidia.hpp"
+#include "alpaka/kernel/TaskKernelGpuSyclAmd.hpp"
 #include "alpaka/kernel/Traits.hpp"
 // math
 #include "alpaka/math/Complex.hpp"
@@ -149,6 +153,7 @@
 #include "alpaka/mem/buf/BufGenericSycl.hpp"
 #include "alpaka/mem/buf/BufGpuSyclIntel.hpp"
 #include "alpaka/mem/buf/BufGpuSyclNvidia.hpp"
+#include "alpaka/mem/buf/BufGpuSyclAmd.hpp"
 #include "alpaka/mem/buf/BufHipRt.hpp"
 #include "alpaka/mem/buf/Traits.hpp"
 #include "alpaka/mem/fence/MemFenceCpu.hpp"
@@ -195,6 +200,7 @@
 #include "alpaka/platform/PlatformFpgaSyclIntel.hpp"
 #include "alpaka/platform/PlatformGpuSyclIntel.hpp"
 #include "alpaka/platform/PlatformGpuSyclNvidia.hpp"
+#include "alpaka/platform/PlatformGpuSyclAmd.hpp"
 #include "alpaka/platform/PlatformHipRt.hpp"
 #include "alpaka/platform/Traits.hpp"
 // rand
@@ -220,6 +226,8 @@
 #include "alpaka/queue/QueueGpuSyclIntelNonBlocking.hpp"
 #include "alpaka/queue/QueueGpuSyclNvidiaBlocking.hpp"
 #include "alpaka/queue/QueueGpuSyclNvidiaNonBlocking.hpp"
+#include "alpaka/queue/QueueGpuSyclAmdBlocking.hpp"
+#include "alpaka/queue/QueueGpuSyclAmdNonBlocking.hpp"
 #include "alpaka/queue/QueueHipRtBlocking.hpp"
 #include "alpaka/queue/QueueHipRtNonBlocking.hpp"
 #include "alpaka/queue/Traits.hpp"

include/alpaka/dev/DevGpuSyclAmd.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_AMD)
+
+namespace alpaka
+{
+    using DevGpuSyclAmd = DevGenericSycl<TagGpuSyclAmd>;
+} // namespace alpaka
+
+#endif

include/alpaka/event/EventGpuSyclAmd.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_AMD)
+
+namespace alpaka
+{
+    using EventGpuSyclAmd = EventGenericSycl<TagGpuSyclAmd>;
+} // namespace alpaka
+
+#endif

include/alpaka/example/ExampleDefaultAcc.hpp

@@ -33,6 +33,8 @@ namespace alpaka
     using ExampleDefaultAcc = alpaka::AccGpuSyclIntel<TDim, TIdx>;
 #    elif defined(ALPAKA_SYCL_ONEAPI_GPU_NVIDIA)
     using ExampleDefaultAcc = alpaka::AccGpuSyclNvidia<TDim, TIdx>;
+#    elif defined(ALPAKA_SYCL_ONEAPI_GPU_AMD)
+    using ExampleDefaultAcc = alpaka::AccGpuSyclAmd<TDim, TIdx>;
 #    endif
 #elif defined(ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED)
     using ExampleDefaultAcc = alpaka::AccCpuSerial<TDim, TIdx>;

include/alpaka/kernel/TaskKernelGenericSycl.hpp

@@ -33,6 +33,7 @@
 #        pragma clang diagnostic push
 #        pragma clang diagnostic ignored "-Wunused-lambda-capture"
 #        pragma clang diagnostic ignored "-Wunused-parameter"
+#        pragma clang diagnostic ignored "-Wcuda-compat"
 #    endif
 
 #    include <sycl/sycl.hpp>

include/alpaka/kernel/TaskKernelGpuSyclAmd.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_AMD)
+
+namespace alpaka
+{
+    template<typename TDim, typename TIdx, typename TKernelFnObj, typename... TArgs>
+    using TaskKernelGpuSyclAmd
+        = TaskKernelGenericSycl<TagGpuSyclAmd, AccGpuSyclAmd<TDim, TIdx>, TDim, TIdx, TKernelFnObj, TArgs...>;
+
+} // namespace alpaka
+
+#endif

include/alpaka/mem/buf/BufGpuSyclAmd.hpp

@@ -0,0 +1,19 @@
+/* Copyright 2024 Jan Stephan, Luca Ferragina, Aurora Perego
+ * SPDX-License-Identifier: MPL-2.0
+ */
+
+#pragma once
+
+#include "alpaka/dev/DevGpuSyclAmd.hpp"
+#include "alpaka/mem/buf/BufGenericSycl.hpp"
+#include "alpaka/platform/PlatformGpuSyclAmd.hpp"
+
+#if defined(ALPAKA_ACC_SYCL_ENABLED) && defined(ALPAKA_SYCL_ONEAPI_GPU_AMD)
+
+namespace alpaka
+{
+    template<typename TElem, typename TDim, typename TIdx>
+    using BufGpuSyclAmd = BufGenericSycl<TElem, TDim, TIdx, PlatformGpuSyclAmd>;
+} // namespace alpaka
+
+#endif

include/alpaka/mem/global/DeviceGlobalGenericSycl.hpp

@@ -37,6 +37,13 @@ namespace alpaka
             using Type = sycl::ext::oneapi::experimental::device_global<T>;
         };
 
+        template<typename T>
+        struct DevGlobalTrait<TagGpuSyclAmd, T>
+        {
+            // SYCL GPU implementation
+            using Type = sycl::ext::oneapi::experimental::device_global<T>;
+        };
+
         template<typename T>
         struct DevGlobalTrait<TagFpgaSyclIntel, T>
         {

include/alpaka/platform/PlatformGpuSyclAmd.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_AMD)
+
+#    include <sycl/sycl.hpp>
+
+namespace alpaka
+{
+    namespace detail
+    {
+        template<>
+        struct SYCLDeviceSelector<TagGpuSyclAmd>
+        {
+            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("Amd(R) Corporation") != std::string::npos);
+
+                return is_intel_gpu ? 1 : -1;
+            }
+        };
+    } // namespace detail
+
+    //! The SYCL device manager.
+    using PlatformGpuSyclAmd = PlatformGenericSycl<TagGpuSyclAmd>;
+} // namespace alpaka
+
+#endif

include/alpaka/queue/QueueGpuSyclAmdBlocking.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_AMD)
+
+namespace alpaka
+{
+    using QueueGpuSyclAmdBlocking = QueueGenericSyclBlocking<TagGpuSyclAmd>;
+} // namespace alpaka
+
+#endif

include/alpaka/queue/QueueGpuSyclAmdNonBlocking.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_AMD)
+
+namespace alpaka
+{
+    using QueueGpuSyclAmdNonBlocking = QueueGenericSyclNonBlocking<TagGpuSyclAmd>;
+} // namespace alpaka
+
+#endif

include/alpaka/standalone/GpuSyclAmd.hpp

@@ -0,0 +1,13 @@
+/* Copyright 2023 Jan Stephan
+ * SPDX-License-Identifier: MPL-2.0
+ */
+
+#include "alpaka/standalone/GenericSycl.hpp"
+
+#ifndef ALPAKA_SYCL_ONEAPI_GPU_AMD
+#    define ALPAKA_SYCL_ONEAPI_GPU_AMD
+#endif
+
+#ifndef ALPAKA_SYCL_TARGET_GPU_AMD
+#    define ALPAKA_SYCL_TARGET_GPU_AMD
+#endif

include/alpaka/standalone/GpuSyclNvidia.hpp

@@ -0,0 +1,13 @@
+/* Copyright 2023 Jan Stephan
+ * SPDX-License-Identifier: MPL-2.0
+ */
+
+#include "alpaka/standalone/GenericSycl.hpp"
+
+#ifndef ALPAKA_SYCL_ONEAPI_GPU_NVIDIA
+#    define ALPAKA_SYCL_ONEAPI_GPU_NVIDIA
+#endif
+
+#ifndef ALPAKA_SYCL_TARGET_GPU_NVIDIA
+#    define ALPAKA_SYCL_TARGET_GPU_NVIDIA
+#endif

include/alpaka/test/acc/TestAccs.hpp

@@ -106,6 +106,13 @@ namespace alpaka::test
         template<typename TDim, typename TIdx>
         using AccGpuSyclNvidiaIfAvailableElseInt = int;
 #endif
+#if defined(ALPAKA_ACC_SYCL_ENABLED) && defined(ALPAKA_SYCL_TARGET_GPU_AMD)
+        template<typename TDim, typename TIdx>
+        using AccGpuSyclAmdIfAvailableElseInt = AccGpuSyclAmd<TDim, TIdx>;
+#else
+        template<typename TDim, typename TIdx>
+        using AccGpuSyclAmdIfAvailableElseInt = int;
+#endif
 
         //! A vector containing all available accelerators and int's.
         template<typename TDim, typename TIdx>
@@ -120,7 +127,8 @@ namespace alpaka::test
             AccCpuSyclIfAvailableElseInt<TDim, TIdx>,
             AccFpgaSyclIntelIfAvailableElseInt<TDim, TIdx>,
             AccGpuSyclIntelIfAvailableElseInt<TDim, TIdx>,
-            AccGpuSyclNvidiaIfAvailableElseInt<TDim, TIdx>>;
+            AccGpuSyclNvidiaIfAvailableElseInt<TDim, TIdx>,
+            AccGpuSyclAmdIfAvailableElseInt<TDim, TIdx>>;
     } // namespace detail
 
     //! A vector containing all available accelerators.