diff --git a/cmake/alpakaCommon.cmake b/cmake/alpakaCommon.cmake index 28f58b7933af..e8ded6d46df8 100644 --- a/cmake/alpakaCommon.cmake +++ b/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() diff --git a/include/alpaka/acc/AccGpuSyclAmd.hpp b/include/alpaka/acc/AccGpuSyclAmd.hpp new file mode 100644 index 000000000000..18fdb9205903 --- /dev/null +++ b/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 + using AccGpuSyclAmd = AccGenericSycl; + + namespace trait + { + template + struct AccToTag> + { + using type = alpaka::TagGpuSyclAmd; + }; + + template + struct TagToAcc + { + using type = alpaka::AccGpuSyclAmd; + }; + } // namespace trait + +} // namespace alpaka + +#endif diff --git a/include/alpaka/acc/Tag.hpp b/include/alpaka/acc/Tag.hpp index b6e85cdce077..e9eb61f0e93c 100644 --- a/include/alpaka/acc/Tag.hpp +++ b/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 diff --git a/include/alpaka/alpaka.hpp b/include/alpaka/alpaka.hpp index ca5bd54b31ea..dcd281ec3b29 100644 --- a/include/alpaka/alpaka.hpp +++ b/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" diff --git a/include/alpaka/dev/DevGpuSyclAmd.hpp b/include/alpaka/dev/DevGpuSyclAmd.hpp new file mode 100644 index 000000000000..fe43d905fa48 --- /dev/null +++ b/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; +} // namespace alpaka + +#endif diff --git a/include/alpaka/event/EventGpuSyclAmd.hpp b/include/alpaka/event/EventGpuSyclAmd.hpp new file mode 100644 index 000000000000..5c3bd58e1e06 --- /dev/null +++ b/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; +} // namespace alpaka + +#endif diff --git a/include/alpaka/example/ExampleDefaultAcc.hpp b/include/alpaka/example/ExampleDefaultAcc.hpp index f90c0b7fc713..035fde130736 100644 --- a/include/alpaka/example/ExampleDefaultAcc.hpp +++ b/include/alpaka/example/ExampleDefaultAcc.hpp @@ -33,6 +33,8 @@ namespace alpaka using ExampleDefaultAcc = alpaka::AccGpuSyclIntel; # elif defined(ALPAKA_SYCL_ONEAPI_GPU_NVIDIA) using ExampleDefaultAcc = alpaka::AccGpuSyclNvidia; +# elif defined(ALPAKA_SYCL_ONEAPI_GPU_AMD) + using ExampleDefaultAcc = alpaka::AccGpuSyclAmd; # endif #elif defined(ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED) using ExampleDefaultAcc = alpaka::AccCpuSerial; diff --git a/include/alpaka/kernel/TaskKernelGenericSycl.hpp b/include/alpaka/kernel/TaskKernelGenericSycl.hpp index 11cc2cae4590..069373b66d6d 100644 --- a/include/alpaka/kernel/TaskKernelGenericSycl.hpp +++ b/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 diff --git a/include/alpaka/kernel/TaskKernelGpuSyclAmd.hpp b/include/alpaka/kernel/TaskKernelGpuSyclAmd.hpp new file mode 100644 index 000000000000..92bab083ff5f --- /dev/null +++ b/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 + using TaskKernelGpuSyclAmd + = TaskKernelGenericSycl, TDim, TIdx, TKernelFnObj, TArgs...>; + +} // namespace alpaka + +#endif diff --git a/include/alpaka/mem/buf/BufGpuSyclAmd.hpp b/include/alpaka/mem/buf/BufGpuSyclAmd.hpp new file mode 100644 index 000000000000..96762ffe0ed2 --- /dev/null +++ b/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 + using BufGpuSyclAmd = BufGenericSycl; +} // namespace alpaka + +#endif diff --git a/include/alpaka/mem/global/DeviceGlobalGenericSycl.hpp b/include/alpaka/mem/global/DeviceGlobalGenericSycl.hpp index 7112620aa15c..6310588c022f 100644 --- a/include/alpaka/mem/global/DeviceGlobalGenericSycl.hpp +++ b/include/alpaka/mem/global/DeviceGlobalGenericSycl.hpp @@ -37,6 +37,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/PlatformGpuSyclAmd.hpp b/include/alpaka/platform/PlatformGpuSyclAmd.hpp new file mode 100644 index 000000000000..103cb50f8fa2 --- /dev/null +++ b/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 + +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("Amd(R) Corporation") != std::string::npos); + + return is_intel_gpu ? 1 : -1; + } + }; + } // namespace detail + + //! The SYCL device manager. + using PlatformGpuSyclAmd = PlatformGenericSycl; +} // namespace alpaka + +#endif diff --git a/include/alpaka/queue/QueueGpuSyclAmdBlocking.hpp b/include/alpaka/queue/QueueGpuSyclAmdBlocking.hpp new file mode 100644 index 000000000000..699e5eb3863d --- /dev/null +++ b/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; +} // namespace alpaka + +#endif diff --git a/include/alpaka/queue/QueueGpuSyclAmdNonBlocking.hpp b/include/alpaka/queue/QueueGpuSyclAmdNonBlocking.hpp new file mode 100644 index 000000000000..feab369a9fe4 --- /dev/null +++ b/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; +} // namespace alpaka + +#endif diff --git a/include/alpaka/standalone/GpuSyclAmd.hpp b/include/alpaka/standalone/GpuSyclAmd.hpp new file mode 100644 index 000000000000..2dd67068642d --- /dev/null +++ b/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 diff --git a/include/alpaka/standalone/GpuSyclNvidia.hpp b/include/alpaka/standalone/GpuSyclNvidia.hpp new file mode 100644 index 000000000000..202e52867b6b --- /dev/null +++ b/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 diff --git a/include/alpaka/test/acc/TestAccs.hpp b/include/alpaka/test/acc/TestAccs.hpp index 4b1e4f761021..4af9c76b4952 100644 --- a/include/alpaka/test/acc/TestAccs.hpp +++ b/include/alpaka/test/acc/TestAccs.hpp @@ -106,6 +106,13 @@ namespace alpaka::test template using AccGpuSyclNvidiaIfAvailableElseInt = int; #endif +#if defined(ALPAKA_ACC_SYCL_ENABLED) && defined(ALPAKA_SYCL_TARGET_GPU_AMD) + template + using AccGpuSyclAmdIfAvailableElseInt = AccGpuSyclAmd; +#else + template + using AccGpuSyclAmdIfAvailableElseInt = int; +#endif //! A vector containing all available accelerators and int's. template @@ -120,7 +127,8 @@ namespace alpaka::test AccCpuSyclIfAvailableElseInt, AccFpgaSyclIntelIfAvailableElseInt, AccGpuSyclIntelIfAvailableElseInt, - AccGpuSyclNvidiaIfAvailableElseInt>; + AccGpuSyclNvidiaIfAvailableElseInt, + AccGpuSyclAmdIfAvailableElseInt>; } // 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 a2d8c90687c8..402993a38705 100644 --- a/include/alpaka/test/queue/Queue.hpp +++ b/include/alpaka/test/queue/Queue.hpp @@ -146,6 +146,11 @@ namespace alpaka::test std::tuple, std::tuple # endif +# ifdef ALPAKA_SYCL_ONEAPI_GPU_AMD + , + 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 6a91e5f8d0a6..50de87ff1183 100644 --- a/test/unit/acc/src/AccTagTest.cpp +++ b/test/unit/acc/src/AccTagTest.cpp @@ -30,7 +30,8 @@ using TagList = std::tuple< alpaka::TagCpuSycl, alpaka::TagFpgaSyclIntel, alpaka::TagGpuSyclIntel, - alpaka::TagGpuSyclNvidia>; + alpaka::TagGpuSyclNvidia, + alpaka::TagGpuSyclAmd>; using AccToTagMap = std::tuple< std::pair, alpaka::TagCpuSerial>, @@ -43,7 +44,8 @@ using AccToTagMap = std::tuple< std::pair, alpaka::TagCpuSycl>, std::pair, alpaka::TagFpgaSyclIntel>, std::pair, alpaka::TagGpuSyclIntel>, - std::pair, alpaka::TagGpuSyclNvidia>>; + std::pair, alpaka::TagGpuSyclNvidia>, + std::pair, alpaka::TagGpuSyclAmd>>; using AccTagTestMatrix = alpaka::meta::CartesianProduct; diff --git a/test/unit/mem/view/src/ViewStaticAccMem.cpp b/test/unit/mem/view/src/ViewStaticAccMem.cpp index d00e643a1fb5..aba25d7ff105 100644 --- a/test/unit/mem/view/src/ViewStaticAccMem.cpp +++ b/test/unit/mem/view/src/ViewStaticAccMem.cpp @@ -72,6 +72,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_AMD) +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 baae5dc4afd9..3c95d49b1726 100644 --- a/test/unit/warp/src/Activemask.cpp +++ b/test/unit/warp/src/Activemask.cpp @@ -70,6 +70,7 @@ TEMPLATE_LIST_TEST_CASE("activemask", "[warp]", alpaka::test::TestAccs) alpaka::TagCpuSycl, alpaka::TagGpuSyclIntel, alpaka::TagGpuSyclNvidia, + alpaka::TagGpuSyclAmd, alpaka::TagFpgaSyclIntel, alpaka::TagGenericSycl>) {