From aa6cf2e8ac74a849d98fdf667a9a87f637beefae Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Wed, 11 Oct 2023 16:22:31 +0200 Subject: [PATCH 01/19] implement memory visibility --- include/alpaka/acc/AccCpuOmp2Blocks.hpp | 6 +++ include/alpaka/acc/AccCpuOmp2Threads.hpp | 6 +++ include/alpaka/acc/AccCpuSerial.hpp | 7 +++ include/alpaka/acc/AccCpuSycl.hpp | 6 +++ include/alpaka/acc/AccCpuTbbBlocks.hpp | 6 +++ include/alpaka/acc/AccCpuThreads.hpp | 6 +++ include/alpaka/acc/AccFpgaSyclIntel.hpp | 6 +++ include/alpaka/acc/AccGenericSycl.hpp | 5 +++ include/alpaka/acc/AccGpuCudaRt.hpp | 6 +++ include/alpaka/acc/AccGpuHipRt.hpp | 6 +++ include/alpaka/acc/AccGpuSyclIntel.hpp | 6 +++ include/alpaka/mem/Visibility.hpp | 47 +++++++++++++++++++++ include/alpaka/mem/buf/BufCpu.hpp | 8 ++++ include/alpaka/mem/buf/BufCpuSycl.hpp | 6 +++ include/alpaka/mem/buf/BufCudaRt.hpp | 6 +++ include/alpaka/mem/buf/BufFpgaSyclIntel.hpp | 6 +++ include/alpaka/mem/buf/BufGenericSycl.hpp | 6 +++ include/alpaka/mem/buf/BufGpuSyclIntel.hpp | 6 +++ include/alpaka/mem/buf/BufHipRt.hpp | 6 +++ 19 files changed, 157 insertions(+) create mode 100644 include/alpaka/mem/Visibility.hpp diff --git a/include/alpaka/acc/AccCpuOmp2Blocks.hpp b/include/alpaka/acc/AccCpuOmp2Blocks.hpp index a5c59e6446ae..3a532ad8582b 100644 --- a/include/alpaka/acc/AccCpuOmp2Blocks.hpp +++ b/include/alpaka/acc/AccCpuOmp2Blocks.hpp @@ -220,6 +220,12 @@ namespace alpaka { using type = alpaka::AccCpuOmp2Blocks; }; + + template + struct MemVisibility> + { + using type = alpaka::MemVisibleCPU; + }; } // namespace trait } // namespace alpaka diff --git a/include/alpaka/acc/AccCpuOmp2Threads.hpp b/include/alpaka/acc/AccCpuOmp2Threads.hpp index bc326bc05c4f..e8a9acf7bcc3 100644 --- a/include/alpaka/acc/AccCpuOmp2Threads.hpp +++ b/include/alpaka/acc/AccCpuOmp2Threads.hpp @@ -231,6 +231,12 @@ namespace alpaka { using type = alpaka::AccCpuOmp2Threads; }; + + template + struct MemVisibility> + { + using type = alpaka::MemVisibleCPU; + }; } // namespace trait } // namespace alpaka diff --git a/include/alpaka/acc/AccCpuSerial.hpp b/include/alpaka/acc/AccCpuSerial.hpp index 4a4e8f0621a6..bd493c517344 100644 --- a/include/alpaka/acc/AccCpuSerial.hpp +++ b/include/alpaka/acc/AccCpuSerial.hpp @@ -33,6 +33,7 @@ #include "alpaka/acc/Tag.hpp" #include "alpaka/core/Concepts.hpp" #include "alpaka/dev/DevCpu.hpp" +#include "alpaka/mem/Visibility.hpp" #include #include @@ -214,6 +215,12 @@ namespace alpaka { using type = alpaka::AccCpuSerial; }; + + template + struct MemVisibility> + { + using type = alpaka::MemVisibleCPU; + }; } // namespace trait } // namespace alpaka diff --git a/include/alpaka/acc/AccCpuSycl.hpp b/include/alpaka/acc/AccCpuSycl.hpp index 7a2615fc9de2..079b75786453 100644 --- a/include/alpaka/acc/AccCpuSycl.hpp +++ b/include/alpaka/acc/AccCpuSycl.hpp @@ -87,6 +87,12 @@ namespace alpaka::trait { using type = alpaka::AccCpuSycl; }; + + template + struct MemVisibility> + { + using type = alpaka::MemVisibleCPU; + }; } // namespace alpaka::trait #endif diff --git a/include/alpaka/acc/AccCpuTbbBlocks.hpp b/include/alpaka/acc/AccCpuTbbBlocks.hpp index 3ef4283d7b63..32867eb4e5fd 100644 --- a/include/alpaka/acc/AccCpuTbbBlocks.hpp +++ b/include/alpaka/acc/AccCpuTbbBlocks.hpp @@ -212,6 +212,12 @@ namespace alpaka { using type = alpaka::AccCpuTbbBlocks; }; + + template + struct MemVisibility> + { + using type = alpaka::MemVisibleCPU; + }; } // namespace trait } // namespace alpaka diff --git a/include/alpaka/acc/AccCpuThreads.hpp b/include/alpaka/acc/AccCpuThreads.hpp index f4984b63d734..8dc2f58a5dfa 100644 --- a/include/alpaka/acc/AccCpuThreads.hpp +++ b/include/alpaka/acc/AccCpuThreads.hpp @@ -239,6 +239,12 @@ namespace alpaka { using type = alpaka::AccCpuThreads; }; + + template + struct MemVisibility> + { + using type = alpaka::MemVisibleCPU; + }; } // namespace trait } // namespace alpaka diff --git a/include/alpaka/acc/AccFpgaSyclIntel.hpp b/include/alpaka/acc/AccFpgaSyclIntel.hpp index db4c0b94c09d..26b15a83f129 100644 --- a/include/alpaka/acc/AccFpgaSyclIntel.hpp +++ b/include/alpaka/acc/AccFpgaSyclIntel.hpp @@ -87,6 +87,12 @@ namespace alpaka::trait { using type = alpaka::AccFpgaSyclIntel; }; + + template + struct MemVisibility> + { + using type = alpaka::MemVisibleFpgaSyclIntel; + }; } // namespace alpaka::trait #endif diff --git a/include/alpaka/acc/AccGenericSycl.hpp b/include/alpaka/acc/AccGenericSycl.hpp index fc8ba1da760f..9e533f892701 100644 --- a/include/alpaka/acc/AccGenericSycl.hpp +++ b/include/alpaka/acc/AccGenericSycl.hpp @@ -118,6 +118,11 @@ namespace alpaka::trait { }; + struct MemVisibility> + { + using type = alpaka::MemVisibleGenericSycl; + }; + //! The SYCL accelerator device properties get trait specialization. template typename TAcc, typename TDim, typename TIdx> struct GetAccDevProps< diff --git a/include/alpaka/acc/AccGpuCudaRt.hpp b/include/alpaka/acc/AccGpuCudaRt.hpp index 5f27e519722d..30cf9ffdd456 100644 --- a/include/alpaka/acc/AccGpuCudaRt.hpp +++ b/include/alpaka/acc/AccGpuCudaRt.hpp @@ -28,6 +28,12 @@ namespace alpaka { using type = alpaka::AccGpuCudaRt; }; + + template + struct MemVisibility> + { + using type = alpaka::MemVisibleGpuCudaRt; + }; } // namespace trait } // namespace alpaka diff --git a/include/alpaka/acc/AccGpuHipRt.hpp b/include/alpaka/acc/AccGpuHipRt.hpp index 43c94ab831bc..215a4848e318 100644 --- a/include/alpaka/acc/AccGpuHipRt.hpp +++ b/include/alpaka/acc/AccGpuHipRt.hpp @@ -28,6 +28,12 @@ namespace alpaka { using type = alpaka::AccGpuHipRt; }; + + template + struct MemVisibility> + { + using type = alpaka::MemVisibleGpuHipRt; + }; } // namespace trait } // namespace alpaka diff --git a/include/alpaka/acc/AccGpuSyclIntel.hpp b/include/alpaka/acc/AccGpuSyclIntel.hpp index bc60307ee17e..622b8ea0a02e 100644 --- a/include/alpaka/acc/AccGpuSyclIntel.hpp +++ b/include/alpaka/acc/AccGpuSyclIntel.hpp @@ -87,6 +87,12 @@ namespace alpaka::trait { using type = alpaka::AccGpuSyclIntel; }; + + template + struct MemVisibility> + { + using type = alpaka::MemVisibleGpuSyclIntel; + }; } // namespace alpaka::trait #endif diff --git a/include/alpaka/mem/Visibility.hpp b/include/alpaka/mem/Visibility.hpp new file mode 100644 index 000000000000..1188c459a7ad --- /dev/null +++ b/include/alpaka/mem/Visibility.hpp @@ -0,0 +1,47 @@ +/* Copyright 2023 Simeon Ehrig + * SPDX-License-Identifier: MPL-2.0 + */ + +#pragma once + +#include "alpaka/meta/TypeListOps.hpp" + +#include + +#define CREATE_MEM_VISIBILITY(mem_name) \ + struct mem_name \ + { \ + static std::string get_name() \ + { \ + return #mem_name; \ + } \ + } + +namespace alpaka +{ + CREATE_MEM_VISIBILITY(MemVisibleCPU); + CREATE_MEM_VISIBILITY(MemVisibleFpgaSyclIntel); + CREATE_MEM_VISIBILITY(MemVisibleGenericSycl); + CREATE_MEM_VISIBILITY(MemVisibleGpuCudaRt); + CREATE_MEM_VISIBILITY(MemVisibleGpuHipRt); + CREATE_MEM_VISIBILITY(MemVisibleGpuSyclIntel); + + namespace trait + { + //! Get memory visibility from a type. + //! Normally it is acc or buffer type. + //! + //! \tparam Type which implements the trait + //! \return Memory visibility type + template + struct MemVisibility; + } // namespace trait + + template + inline constexpr bool hasSameMemView() + { + return alpaka::meta::Contains< + typename alpaka::trait::MemVisibility::type, + typename alpaka::trait::MemVisibility::type>::value; + } +} // namespace alpaka diff --git a/include/alpaka/mem/buf/BufCpu.hpp b/include/alpaka/mem/buf/BufCpu.hpp index 4bfc91c73332..282ef672cfb9 100644 --- a/include/alpaka/mem/buf/BufCpu.hpp +++ b/include/alpaka/mem/buf/BufCpu.hpp @@ -11,6 +11,7 @@ #include "alpaka/core/Vectorize.hpp" #include "alpaka/dev/DevCpu.hpp" #include "alpaka/dev/Traits.hpp" +#include "alpaka/mem/Visibility.hpp" #include "alpaka/mem/alloc/AllocCpuAligned.hpp" #include "alpaka/mem/buf/Traits.hpp" #include "alpaka/mem/view/ViewAccessOps.hpp" @@ -307,6 +308,13 @@ namespace alpaka { using type = TIdx; }; + + template + struct MemVisibility> + { + using type = std::tuple; + }; + } // namespace trait } // namespace alpaka diff --git a/include/alpaka/mem/buf/BufCpuSycl.hpp b/include/alpaka/mem/buf/BufCpuSycl.hpp index d63eebf540ca..a04a94e66629 100644 --- a/include/alpaka/mem/buf/BufCpuSycl.hpp +++ b/include/alpaka/mem/buf/BufCpuSycl.hpp @@ -13,6 +13,12 @@ namespace alpaka { template using BufCpuSycl = BufGenericSycl; + + template + struct MemVisibility> + { + using type = std::tuple; + }; } // namespace alpaka #endif diff --git a/include/alpaka/mem/buf/BufCudaRt.hpp b/include/alpaka/mem/buf/BufCudaRt.hpp index a5e0020bdb62..62bc8316eb9f 100644 --- a/include/alpaka/mem/buf/BufCudaRt.hpp +++ b/include/alpaka/mem/buf/BufCudaRt.hpp @@ -13,6 +13,12 @@ namespace alpaka { template using BufCudaRt = BufUniformCudaHipRt; + + template + struct MemVisibility> + { + using type = std::tuple; + }; } // namespace alpaka #endif // ALPAKA_ACC_GPU_CUDA_ENABLED diff --git a/include/alpaka/mem/buf/BufFpgaSyclIntel.hpp b/include/alpaka/mem/buf/BufFpgaSyclIntel.hpp index 2dca26f1984f..712406fa8b0e 100644 --- a/include/alpaka/mem/buf/BufFpgaSyclIntel.hpp +++ b/include/alpaka/mem/buf/BufFpgaSyclIntel.hpp @@ -13,6 +13,12 @@ namespace alpaka { template using BufFpgaSyclIntel = BufGenericSycl; + + template + struct MemVisibility> + { + using type = std::tuple; + }; } // namespace alpaka #endif diff --git a/include/alpaka/mem/buf/BufGenericSycl.hpp b/include/alpaka/mem/buf/BufGenericSycl.hpp index b4a5fd94ed54..a0180c46a7fd 100644 --- a/include/alpaka/mem/buf/BufGenericSycl.hpp +++ b/include/alpaka/mem/buf/BufGenericSycl.hpp @@ -258,6 +258,12 @@ namespace alpaka::trait return getPtrNative(buf); } }; + + template + struct MemVisibility> + { + using type = std::tuple; + }; } // namespace alpaka::trait # include "alpaka/mem/buf/sycl/Copy.hpp" diff --git a/include/alpaka/mem/buf/BufGpuSyclIntel.hpp b/include/alpaka/mem/buf/BufGpuSyclIntel.hpp index dd20f8a39648..584e799f4e28 100644 --- a/include/alpaka/mem/buf/BufGpuSyclIntel.hpp +++ b/include/alpaka/mem/buf/BufGpuSyclIntel.hpp @@ -13,6 +13,12 @@ namespace alpaka { template using BufGpuSyclIntel = BufGenericSycl; + + template + struct MemVisibility> + { + using type = std::tuple; + }; } // namespace alpaka #endif diff --git a/include/alpaka/mem/buf/BufHipRt.hpp b/include/alpaka/mem/buf/BufHipRt.hpp index 4a59bc46e5d5..75015cbc7401 100644 --- a/include/alpaka/mem/buf/BufHipRt.hpp +++ b/include/alpaka/mem/buf/BufHipRt.hpp @@ -13,6 +13,12 @@ namespace alpaka { template using BufHipRt = BufUniformCudaHipRt; + + template + struct MemVisibility> + { + using type = std::tuple; + }; } // namespace alpaka #endif // ALPAKA_ACC_GPU_HIP_ENABLED From 57d2034a1757809dd830e73caf3b43126f5dfbc7 Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Mon, 22 Apr 2024 17:38:22 +0200 Subject: [PATCH 02/19] use device type instead acc type --- include/alpaka/acc/AccCpuOmp2Blocks.hpp | 6 --- include/alpaka/acc/AccCpuOmp2Threads.hpp | 6 --- include/alpaka/acc/AccCpuSerial.hpp | 7 --- include/alpaka/acc/AccCpuSycl.hpp | 6 --- include/alpaka/acc/AccCpuTbbBlocks.hpp | 6 --- include/alpaka/acc/AccCpuThreads.hpp | 6 --- include/alpaka/acc/AccFpgaSyclIntel.hpp | 6 --- include/alpaka/acc/AccGpuCudaRt.hpp | 6 --- include/alpaka/acc/AccGpuHipRt.hpp | 6 --- include/alpaka/acc/AccGpuSyclIntel.hpp | 6 --- include/alpaka/alpaka.hpp | 2 + include/alpaka/dev/DevCpu.hpp | 7 +++ include/alpaka/dev/DevCpuSycl.hpp | 9 ++++ include/alpaka/dev/DevCudaRt.hpp | 9 ++++ include/alpaka/dev/DevFpgaSyclIntel.hpp | 9 ++++ include/alpaka/dev/DevGpuSyclIntel.hpp | 9 ++++ include/alpaka/dev/DevHipRt.hpp | 9 ++++ include/alpaka/mem/Visibility.hpp | 67 ++++++++++++++++++++++-- include/alpaka/mem/buf/BufCudaRt.hpp | 13 +++-- include/alpaka/meta/IsTuple.hpp | 29 ++++++++++ 20 files changed, 154 insertions(+), 70 deletions(-) create mode 100644 include/alpaka/meta/IsTuple.hpp diff --git a/include/alpaka/acc/AccCpuOmp2Blocks.hpp b/include/alpaka/acc/AccCpuOmp2Blocks.hpp index 3a532ad8582b..a5c59e6446ae 100644 --- a/include/alpaka/acc/AccCpuOmp2Blocks.hpp +++ b/include/alpaka/acc/AccCpuOmp2Blocks.hpp @@ -220,12 +220,6 @@ namespace alpaka { using type = alpaka::AccCpuOmp2Blocks; }; - - template - struct MemVisibility> - { - using type = alpaka::MemVisibleCPU; - }; } // namespace trait } // namespace alpaka diff --git a/include/alpaka/acc/AccCpuOmp2Threads.hpp b/include/alpaka/acc/AccCpuOmp2Threads.hpp index e8a9acf7bcc3..bc326bc05c4f 100644 --- a/include/alpaka/acc/AccCpuOmp2Threads.hpp +++ b/include/alpaka/acc/AccCpuOmp2Threads.hpp @@ -231,12 +231,6 @@ namespace alpaka { using type = alpaka::AccCpuOmp2Threads; }; - - template - struct MemVisibility> - { - using type = alpaka::MemVisibleCPU; - }; } // namespace trait } // namespace alpaka diff --git a/include/alpaka/acc/AccCpuSerial.hpp b/include/alpaka/acc/AccCpuSerial.hpp index bd493c517344..4a4e8f0621a6 100644 --- a/include/alpaka/acc/AccCpuSerial.hpp +++ b/include/alpaka/acc/AccCpuSerial.hpp @@ -33,7 +33,6 @@ #include "alpaka/acc/Tag.hpp" #include "alpaka/core/Concepts.hpp" #include "alpaka/dev/DevCpu.hpp" -#include "alpaka/mem/Visibility.hpp" #include #include @@ -215,12 +214,6 @@ namespace alpaka { using type = alpaka::AccCpuSerial; }; - - template - struct MemVisibility> - { - using type = alpaka::MemVisibleCPU; - }; } // namespace trait } // namespace alpaka diff --git a/include/alpaka/acc/AccCpuSycl.hpp b/include/alpaka/acc/AccCpuSycl.hpp index 079b75786453..7a2615fc9de2 100644 --- a/include/alpaka/acc/AccCpuSycl.hpp +++ b/include/alpaka/acc/AccCpuSycl.hpp @@ -87,12 +87,6 @@ namespace alpaka::trait { using type = alpaka::AccCpuSycl; }; - - template - struct MemVisibility> - { - using type = alpaka::MemVisibleCPU; - }; } // namespace alpaka::trait #endif diff --git a/include/alpaka/acc/AccCpuTbbBlocks.hpp b/include/alpaka/acc/AccCpuTbbBlocks.hpp index 32867eb4e5fd..3ef4283d7b63 100644 --- a/include/alpaka/acc/AccCpuTbbBlocks.hpp +++ b/include/alpaka/acc/AccCpuTbbBlocks.hpp @@ -212,12 +212,6 @@ namespace alpaka { using type = alpaka::AccCpuTbbBlocks; }; - - template - struct MemVisibility> - { - using type = alpaka::MemVisibleCPU; - }; } // namespace trait } // namespace alpaka diff --git a/include/alpaka/acc/AccCpuThreads.hpp b/include/alpaka/acc/AccCpuThreads.hpp index 8dc2f58a5dfa..f4984b63d734 100644 --- a/include/alpaka/acc/AccCpuThreads.hpp +++ b/include/alpaka/acc/AccCpuThreads.hpp @@ -239,12 +239,6 @@ namespace alpaka { using type = alpaka::AccCpuThreads; }; - - template - struct MemVisibility> - { - using type = alpaka::MemVisibleCPU; - }; } // namespace trait } // namespace alpaka diff --git a/include/alpaka/acc/AccFpgaSyclIntel.hpp b/include/alpaka/acc/AccFpgaSyclIntel.hpp index 26b15a83f129..db4c0b94c09d 100644 --- a/include/alpaka/acc/AccFpgaSyclIntel.hpp +++ b/include/alpaka/acc/AccFpgaSyclIntel.hpp @@ -87,12 +87,6 @@ namespace alpaka::trait { using type = alpaka::AccFpgaSyclIntel; }; - - template - struct MemVisibility> - { - using type = alpaka::MemVisibleFpgaSyclIntel; - }; } // namespace alpaka::trait #endif diff --git a/include/alpaka/acc/AccGpuCudaRt.hpp b/include/alpaka/acc/AccGpuCudaRt.hpp index 30cf9ffdd456..5f27e519722d 100644 --- a/include/alpaka/acc/AccGpuCudaRt.hpp +++ b/include/alpaka/acc/AccGpuCudaRt.hpp @@ -28,12 +28,6 @@ namespace alpaka { using type = alpaka::AccGpuCudaRt; }; - - template - struct MemVisibility> - { - using type = alpaka::MemVisibleGpuCudaRt; - }; } // namespace trait } // namespace alpaka diff --git a/include/alpaka/acc/AccGpuHipRt.hpp b/include/alpaka/acc/AccGpuHipRt.hpp index 215a4848e318..43c94ab831bc 100644 --- a/include/alpaka/acc/AccGpuHipRt.hpp +++ b/include/alpaka/acc/AccGpuHipRt.hpp @@ -28,12 +28,6 @@ namespace alpaka { using type = alpaka::AccGpuHipRt; }; - - template - struct MemVisibility> - { - using type = alpaka::MemVisibleGpuHipRt; - }; } // namespace trait } // namespace alpaka diff --git a/include/alpaka/acc/AccGpuSyclIntel.hpp b/include/alpaka/acc/AccGpuSyclIntel.hpp index 622b8ea0a02e..bc60307ee17e 100644 --- a/include/alpaka/acc/AccGpuSyclIntel.hpp +++ b/include/alpaka/acc/AccGpuSyclIntel.hpp @@ -87,12 +87,6 @@ namespace alpaka::trait { using type = alpaka::AccGpuSyclIntel; }; - - template - struct MemVisibility> - { - using type = alpaka::MemVisibleGpuSyclIntel; - }; } // namespace alpaka::trait #endif diff --git a/include/alpaka/alpaka.hpp b/include/alpaka/alpaka.hpp index e06dede53d48..b9396429c4f9 100644 --- a/include/alpaka/alpaka.hpp +++ b/include/alpaka/alpaka.hpp @@ -130,6 +130,7 @@ #include "alpaka/math/MathStdLib.hpp" #include "alpaka/math/MathUniformCudaHipBuiltIn.hpp" // mem +#include "alpaka/mem/Visibility.hpp" #include "alpaka/mem/alloc/AllocCpuAligned.hpp" #include "alpaka/mem/alloc/AllocCpuNew.hpp" #include "alpaka/mem/alloc/Traits.hpp" @@ -171,6 +172,7 @@ #include "alpaka/meta/Integral.hpp" #include "alpaka/meta/IsArrayOrVector.hpp" #include "alpaka/meta/IsStrictBase.hpp" +#include "alpaka/meta/IsTuple.hpp" #include "alpaka/meta/NdLoop.hpp" #include "alpaka/meta/NonZero.hpp" #include "alpaka/meta/Set.hpp" diff --git a/include/alpaka/dev/DevCpu.hpp b/include/alpaka/dev/DevCpu.hpp index e36c263072fa..ac294cab229a 100644 --- a/include/alpaka/dev/DevCpu.hpp +++ b/include/alpaka/dev/DevCpu.hpp @@ -8,6 +8,7 @@ #include "alpaka/dev/Traits.hpp" #include "alpaka/dev/common/QueueRegistry.hpp" #include "alpaka/dev/cpu/SysInfo.hpp" +#include "alpaka/mem/Visibility.hpp" #include "alpaka/mem/buf/Traits.hpp" #include "alpaka/platform/Traits.hpp" #include "alpaka/queue/Properties.hpp" @@ -185,6 +186,12 @@ namespace alpaka { using type = PlatformCpu; }; + + template<> + struct MemVisibility + { + using type = alpaka::MemVisibleCPU; + }; } // namespace trait using QueueCpuNonBlocking = QueueGenericThreadsNonBlocking; diff --git a/include/alpaka/dev/DevCpuSycl.hpp b/include/alpaka/dev/DevCpuSycl.hpp index 04b15a867558..fd58370e8931 100644 --- a/include/alpaka/dev/DevCpuSycl.hpp +++ b/include/alpaka/dev/DevCpuSycl.hpp @@ -12,6 +12,15 @@ namespace alpaka { using DevCpuSycl = DevGenericSycl; + + namespace trait + { + template<> + struct MemVisibility + { + using type = alpaka::MemVisibleCPU; + }; + } // namespace trait } // namespace alpaka #endif diff --git a/include/alpaka/dev/DevCudaRt.hpp b/include/alpaka/dev/DevCudaRt.hpp index 92dcba3a89b4..a487fa26aee1 100644 --- a/include/alpaka/dev/DevCudaRt.hpp +++ b/include/alpaka/dev/DevCudaRt.hpp @@ -13,6 +13,15 @@ namespace alpaka { //! The CUDA RT device handle. using DevCudaRt = DevUniformCudaHipRt; + + namespace trait + { + template<> + struct MemVisibility + { + using type = alpaka::MemVisibleGpuCudaRt; + }; + } // namespace trait } // namespace alpaka #endif // ALPAKA_ACC_GPU_CUDA_ENABLED diff --git a/include/alpaka/dev/DevFpgaSyclIntel.hpp b/include/alpaka/dev/DevFpgaSyclIntel.hpp index 516027db6b2a..5e4c17ad2502 100644 --- a/include/alpaka/dev/DevFpgaSyclIntel.hpp +++ b/include/alpaka/dev/DevFpgaSyclIntel.hpp @@ -12,6 +12,15 @@ namespace alpaka { using DevFpgaSyclIntel = DevGenericSycl; + + namespace trait + { + template<> + struct MemVisibility + { + using type = alpaka::MemVisibleFpgaSyclIntel; + }; + } // namespace trait } // namespace alpaka #endif diff --git a/include/alpaka/dev/DevGpuSyclIntel.hpp b/include/alpaka/dev/DevGpuSyclIntel.hpp index 9897d40ebbc5..0b8786a35eb7 100644 --- a/include/alpaka/dev/DevGpuSyclIntel.hpp +++ b/include/alpaka/dev/DevGpuSyclIntel.hpp @@ -12,6 +12,15 @@ namespace alpaka { using DevGpuSyclIntel = DevGenericSycl; + + namespace trait + { + template<> + struct MemVisibility + { + using type = alpaka::MemVisibleGpuSyclIntel; + }; + } // namespace trait } // namespace alpaka #endif diff --git a/include/alpaka/dev/DevHipRt.hpp b/include/alpaka/dev/DevHipRt.hpp index 819c2f5c4543..075fb3cf5efd 100644 --- a/include/alpaka/dev/DevHipRt.hpp +++ b/include/alpaka/dev/DevHipRt.hpp @@ -13,6 +13,15 @@ namespace alpaka { //! The HIP RT device handle. using DevHipRt = DevUniformCudaHipRt; + + namespace trait + { + template<> + struct MemVisibility + { + using type = alpaka::MemVisibleGpuHipRt; + }; + } // namespace trait } // namespace alpaka #endif // ALPAKA_ACC_GPU_HIP_ENABLED diff --git a/include/alpaka/mem/Visibility.hpp b/include/alpaka/mem/Visibility.hpp index 1188c459a7ad..3e4fea635523 100644 --- a/include/alpaka/mem/Visibility.hpp +++ b/include/alpaka/mem/Visibility.hpp @@ -4,14 +4,18 @@ #pragma once +#include "alpaka/meta/ForEachType.hpp" +#include "alpaka/meta/IsTuple.hpp" #include "alpaka/meta/TypeListOps.hpp" #include +#include +#include #define CREATE_MEM_VISIBILITY(mem_name) \ struct mem_name \ { \ - static std::string get_name() \ + static std::string name() \ { \ return #mem_name; \ } \ @@ -31,17 +35,70 @@ namespace alpaka //! Get memory visibility from a type. //! Normally it is acc or buffer type. //! - //! \tparam Type which implements the trait - //! \return Memory visibility type + //! \tparam TType which implements the trait template struct MemVisibility; } // namespace trait - template + namespace detail + { + struct AppendMemTypeName + { + template + void operator()(std::vector& vs) + { + vs.push_back(TTYPE::name()); + } + }; + } // namespace detail + + template + static std::string getMemVisiblityName() + { + using MemVisibilityType = typename alpaka::trait::MemVisibility>::type; + if constexpr(alpaka::meta::isTuple()) + { + std::vector vs; + alpaka::meta::forEachType(detail::AppendMemTypeName{}, vs); + + std::stringstream ss; + ss << "<"; + for(auto i = 0; i < vs.size(); ++i) + { + if(i == (vs.size() - 1)) + { + ss << vs[i] << ">"; + } + else + { + ss << vs[i] << ", "; + } + } + return ss.str(); + } + else + { + return MemVisibilityType::name(); + } + } + + template + static std::string getMemVisiblityName(TType) + { + return getMemVisiblityName(); + } + + template inline constexpr bool hasSameMemView() { return alpaka::meta::Contains< typename alpaka::trait::MemVisibility::type, - typename alpaka::trait::MemVisibility::type>::value; + typename alpaka::trait::MemVisibility::type>::value; + } + + template + inline constexpr bool hasSameMemView(TDev, TBuf) + { + return hasSameMemView, std::decay_t>(); } } // namespace alpaka diff --git a/include/alpaka/mem/buf/BufCudaRt.hpp b/include/alpaka/mem/buf/BufCudaRt.hpp index 62bc8316eb9f..79288b689c3c 100644 --- a/include/alpaka/mem/buf/BufCudaRt.hpp +++ b/include/alpaka/mem/buf/BufCudaRt.hpp @@ -5,8 +5,10 @@ #pragma once #include "alpaka/core/ApiCudaRt.hpp" +#include "alpaka/mem/Visibility.hpp" #include "alpaka/mem/buf/BufUniformCudaHipRt.hpp" + #ifdef ALPAKA_ACC_GPU_CUDA_ENABLED namespace alpaka @@ -14,11 +16,14 @@ namespace alpaka template using BufCudaRt = BufUniformCudaHipRt; - template - struct MemVisibility> + namespace trait { - using type = std::tuple; - }; + template + struct MemVisibility> + { + using type = std::tuple; + }; + } // namespace trait } // namespace alpaka #endif // ALPAKA_ACC_GPU_CUDA_ENABLED diff --git a/include/alpaka/meta/IsTuple.hpp b/include/alpaka/meta/IsTuple.hpp new file mode 100644 index 000000000000..ee4a920923ba --- /dev/null +++ b/include/alpaka/meta/IsTuple.hpp @@ -0,0 +1,29 @@ +/* Copyright 2024 Simeon Ehrig + * SPDX-License-Identifier: MPL-2.0 + */ + +#pragma once + +#include +#include + +// copied from https://stackoverflow.com/a/51073558/22035743 +namespace alpaka::meta +{ + template + struct IsTuple : std::false_type + { + }; + + template + struct IsTuple> : std::true_type + { + }; + + template + constexpr bool isTuple() + { + return IsTuple>::value; + } + +} // namespace alpaka::meta From 4e0d945efece7f9e3da6387b6d62e1303598d644 Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Wed, 24 Apr 2024 16:28:37 +0200 Subject: [PATCH 03/19] fixes warnings --- include/alpaka/mem/Visibility.hpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/include/alpaka/mem/Visibility.hpp b/include/alpaka/mem/Visibility.hpp index 3e4fea635523..39239ada5d88 100644 --- a/include/alpaka/mem/Visibility.hpp +++ b/include/alpaka/mem/Visibility.hpp @@ -53,7 +53,7 @@ namespace alpaka } // namespace detail template - static std::string getMemVisiblityName() + [[maybe_unused]] static std::string getMemVisiblityName() { using MemVisibilityType = typename alpaka::trait::MemVisibility>::type; if constexpr(alpaka::meta::isTuple()) @@ -63,7 +63,7 @@ namespace alpaka std::stringstream ss; ss << "<"; - for(auto i = 0; i < vs.size(); ++i) + for(std::size_t i = 0; i < vs.size(); ++i) { if(i == (vs.size() - 1)) { @@ -83,7 +83,7 @@ namespace alpaka } template - static std::string getMemVisiblityName(TType) + [[maybe_unused]] static std::string getMemVisiblityName(TType) { return getMemVisiblityName(); } From d7e7b8eebf3ad2fc6af3688d4121f987b0d5c8f1 Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Tue, 30 Apr 2024 10:43:43 +0200 Subject: [PATCH 04/19] add small fixes --- include/alpaka/mem/Visibility.hpp | 3 ++- include/alpaka/mem/buf/BufGpuSyclIntel.hpp | 11 +++++++---- include/alpaka/mem/buf/BufHipRt.hpp | 11 +++++++---- 3 files changed, 16 insertions(+), 9 deletions(-) diff --git a/include/alpaka/mem/Visibility.hpp b/include/alpaka/mem/Visibility.hpp index 39239ada5d88..4775215677f8 100644 --- a/include/alpaka/mem/Visibility.hpp +++ b/include/alpaka/mem/Visibility.hpp @@ -11,6 +11,7 @@ #include #include #include +#include #define CREATE_MEM_VISIBILITY(mem_name) \ struct mem_name \ @@ -97,7 +98,7 @@ namespace alpaka } template - inline constexpr bool hasSameMemView(TDev, TBuf) + inline constexpr bool hasSameMemView(TDev&, TBuf&) { return hasSameMemView, std::decay_t>(); } diff --git a/include/alpaka/mem/buf/BufGpuSyclIntel.hpp b/include/alpaka/mem/buf/BufGpuSyclIntel.hpp index 584e799f4e28..91d6bc9193b9 100644 --- a/include/alpaka/mem/buf/BufGpuSyclIntel.hpp +++ b/include/alpaka/mem/buf/BufGpuSyclIntel.hpp @@ -14,11 +14,14 @@ namespace alpaka template using BufGpuSyclIntel = BufGenericSycl; - template - struct MemVisibility> + namespace trait { - using type = std::tuple; - }; + template + struct MemVisibility> + { + using type = std::tuple; + }; + } // namespace trait } // namespace alpaka #endif diff --git a/include/alpaka/mem/buf/BufHipRt.hpp b/include/alpaka/mem/buf/BufHipRt.hpp index 75015cbc7401..99c57424fe3a 100644 --- a/include/alpaka/mem/buf/BufHipRt.hpp +++ b/include/alpaka/mem/buf/BufHipRt.hpp @@ -14,11 +14,14 @@ namespace alpaka template using BufHipRt = BufUniformCudaHipRt; - template - struct MemVisibility> + namespace trait { - using type = std::tuple; - }; + template + struct MemVisibility> + { + using type = std::tuple; + }; + } // namespace trait } // namespace alpaka #endif // ALPAKA_ACC_GPU_HIP_ENABLED From d1dce1fe2412305f8f1927916fb2140d9f4ab496 Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Mon, 6 May 2024 15:58:26 +0200 Subject: [PATCH 05/19] implement first memory visibility test --- test/unit/CMakeLists.txt | 6 +- test/unit/mem/CMakeLists.txt | 33 +++++++++ test/unit/mem/src/Visibility.cpp | 119 +++++++++++++++++++++++++++++++ 3 files changed, 153 insertions(+), 5 deletions(-) create mode 100644 test/unit/mem/CMakeLists.txt create mode 100644 test/unit/mem/src/Visibility.cpp diff --git a/test/unit/CMakeLists.txt b/test/unit/CMakeLists.txt index bcd108ff8d4c..209a41d9ccbf 100644 --- a/test/unit/CMakeLists.txt +++ b/test/unit/CMakeLists.txt @@ -25,11 +25,7 @@ add_subdirectory("idx/") add_subdirectory("intrinsic/") add_subdirectory("kernel/") add_subdirectory("math/") -add_subdirectory("mem/buf/") -add_subdirectory("mem/copy/") -add_subdirectory("mem/fence/") -add_subdirectory("mem/p2p/") -add_subdirectory("mem/view/") +add_subdirectory("mem/") add_subdirectory("meta/") add_subdirectory("queue/") add_subdirectory("rand/") diff --git a/test/unit/mem/CMakeLists.txt b/test/unit/mem/CMakeLists.txt new file mode 100644 index 000000000000..16399eca3414 --- /dev/null +++ b/test/unit/mem/CMakeLists.txt @@ -0,0 +1,33 @@ +# +# Copyright 2024 Simeon Ehrig +# SPDX-License-Identifier: MPL-2.0 +# + +################################################################################ +# Required CMake version. +################################################################################ + +cmake_minimum_required(VERSION 3.22) + +set(_TARGET_NAME "memTest") + +append_recursive_files_add_to_src_group("src/" "src/" "cpp" _FILES_SOURCE) + +alpaka_add_executable( + ${_TARGET_NAME} + ${_FILES_SOURCE}) +target_link_libraries( + ${_TARGET_NAME} + PRIVATE common) + +set_target_properties(${_TARGET_NAME} PROPERTIES FOLDER "test/unit") +target_compile_definitions(${_TARGET_NAME} PRIVATE "-DTEST_UNIT_MEM") + +add_test(NAME ${_TARGET_NAME} COMMAND ${_TARGET_NAME} ${_alpaka_TEST_OPTIONS}) + + +add_subdirectory("buf/") +add_subdirectory("copy/") +add_subdirectory("fence/") +add_subdirectory("p2p/") +add_subdirectory("view/") diff --git a/test/unit/mem/src/Visibility.cpp b/test/unit/mem/src/Visibility.cpp new file mode 100644 index 000000000000..e26400b586e6 --- /dev/null +++ b/test/unit/mem/src/Visibility.cpp @@ -0,0 +1,119 @@ +/* Copyright 2024 Simeon Ehrig + * SPDX-License-Identifier: MPL-2.0 + */ + +#include + +#include + +#include +#include + +using Dim = alpaka::DimInt<1>; +using Idx = std::size_t; + +template +constexpr bool isCPUTag() +{ + if constexpr( + std::is_same_v || std::is_same_v + || std::is_same_v || std::is_same_v + || std::is_same_v) + { + return true; + } + else + { + return false; + } +} + +using TagList = std::tuple< + alpaka::TagCpuSerial, + alpaka::TagCpuThreads, + alpaka::TagCpuTbbBlocks, + alpaka::TagCpuOmp2Blocks, + alpaka::TagCpuOmp2Threads, + alpaka::TagGpuCudaRt, + alpaka::TagGpuHipRt, + alpaka::TagCpuSycl, + alpaka::TagFpgaSyclIntel, + alpaka::TagGpuSyclIntel>; + +TEMPLATE_LIST_TEST_CASE("memoryVisibilityType", "[mem][visibility]", TagList) +{ + using Tag = TestType; + if constexpr(alpaka::AccIsEnabled::value) + { + using DevType = decltype(alpaka::getDevByIdx(alpaka::Platform>{}, 0)); + if constexpr(isCPUTag()) + { + STATIC_REQUIRE( + std::is_same_v::type, alpaka::MemVisibleCPU>); + } + else if(std::is_same_v) + { + STATIC_REQUIRE( + std::is_same_v::type, alpaka::MemVisibleGpuCudaRt>); + } + } +} + +using TagTagList = alpaka::meta::CartesianProduct; + +template +void do_job(TDev dev, TBuf buffer) +{ + STATIC_REQUIRE(alpaka::hasSameMemView(dev, buffer)); +} + +TEMPLATE_LIST_TEST_CASE("printDefines", "[mem][visibility]", TagTagList) +{ + using Tag1 = std::tuple_element_t<0, TestType>; + using Tag2 = std::tuple_element_t<1, TestType>; + + if constexpr(alpaka::AccIsEnabled::value && alpaka::AccIsEnabled::value) + { + SUCCEED(Tag1::get_name() << " + " << Tag2::get_name()); + + using Acc1 = alpaka::TagToAcc; + using Acc2 = alpaka::TagToAcc; + + auto const plt1 = alpaka::Platform{}; + auto const plt2 = alpaka::Platform{}; + + auto const dev1 = alpaka::getDevByIdx(plt1, 0); + auto const dev2 = alpaka::getDevByIdx(plt2, 0); + + using BufAcc1 = alpaka::Buf; + using BufAcc2 = alpaka::Buf; + + BufAcc1 bufDev1(alpaka::allocBuf(dev1, Idx(1))); + BufAcc2 bufDev2(alpaka::allocBuf(dev2, Idx(1))); + + STATIC_REQUIRE(alpaka::hasSameMemView(dev1, bufDev1)); + STATIC_REQUIRE(alpaka::hasSameMemView(dev2, bufDev2)); + + // at the moment, only the cpu platform has different accelerator types + // therefore all cpu accelerators can access the memory of other cpu accelerators + // if the accelerator is not a cpu accelerator, both accelerators needs to be the + // same to support access to the memory of each other + if constexpr((isCPUTag() && isCPUTag()) || std::is_same_v) + { + STATIC_REQUIRE(alpaka::hasSameMemView(dev1, bufDev2)); + STATIC_REQUIRE(alpaka::hasSameMemView(dev2, bufDev1)); + } + else + { + STATIC_REQUIRE_FALSE(alpaka::hasSameMemView(dev1, bufDev2)); + STATIC_REQUIRE_FALSE(alpaka::hasSameMemView(dev2, bufDev1)); + } + + do_job(dev1, bufDev1); + do_job(dev2, bufDev2); + // do_job(dev1, bufDev2); + + // std::cout << std::boolalpha << "tag 1 is cpu: " << isCPUTag() << "\n"; + // std::cout << std::boolalpha << "tag 2 is cpu: " << isCPUTag() << "\n"; + } +} From 9bf634fc0a90b7d03592518d4dd1a1a514f45cb1 Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Tue, 7 May 2024 08:55:20 +0200 Subject: [PATCH 06/19] move memory visibility from device to platform --- include/alpaka/dev/DevCpu.hpp | 7 ------ include/alpaka/dev/DevCpuSycl.hpp | 9 ------- include/alpaka/dev/DevCudaRt.hpp | 9 ------- include/alpaka/dev/DevFpgaSyclIntel.hpp | 9 ------- include/alpaka/dev/DevGpuSyclIntel.hpp | 9 ------- include/alpaka/dev/DevHipRt.hpp | 9 ------- include/alpaka/platform/PlatformCpu.hpp | 7 ++++++ include/alpaka/platform/PlatformCpuSycl.hpp | 7 ++++++ include/alpaka/platform/PlatformCudaRt.hpp | 10 ++++++++ .../alpaka/platform/PlatformFpgaSyclIntel.hpp | 7 ++++++ .../alpaka/platform/PlatformGpuSyclIntel.hpp | 7 ++++++ include/alpaka/platform/PlatformHipRt.hpp | 10 ++++++++ test/unit/mem/src/Visibility.cpp | 24 +++++++++---------- 13 files changed, 60 insertions(+), 64 deletions(-) diff --git a/include/alpaka/dev/DevCpu.hpp b/include/alpaka/dev/DevCpu.hpp index ac294cab229a..e36c263072fa 100644 --- a/include/alpaka/dev/DevCpu.hpp +++ b/include/alpaka/dev/DevCpu.hpp @@ -8,7 +8,6 @@ #include "alpaka/dev/Traits.hpp" #include "alpaka/dev/common/QueueRegistry.hpp" #include "alpaka/dev/cpu/SysInfo.hpp" -#include "alpaka/mem/Visibility.hpp" #include "alpaka/mem/buf/Traits.hpp" #include "alpaka/platform/Traits.hpp" #include "alpaka/queue/Properties.hpp" @@ -186,12 +185,6 @@ namespace alpaka { using type = PlatformCpu; }; - - template<> - struct MemVisibility - { - using type = alpaka::MemVisibleCPU; - }; } // namespace trait using QueueCpuNonBlocking = QueueGenericThreadsNonBlocking; diff --git a/include/alpaka/dev/DevCpuSycl.hpp b/include/alpaka/dev/DevCpuSycl.hpp index fd58370e8931..04b15a867558 100644 --- a/include/alpaka/dev/DevCpuSycl.hpp +++ b/include/alpaka/dev/DevCpuSycl.hpp @@ -12,15 +12,6 @@ namespace alpaka { using DevCpuSycl = DevGenericSycl; - - namespace trait - { - template<> - struct MemVisibility - { - using type = alpaka::MemVisibleCPU; - }; - } // namespace trait } // namespace alpaka #endif diff --git a/include/alpaka/dev/DevCudaRt.hpp b/include/alpaka/dev/DevCudaRt.hpp index a487fa26aee1..92dcba3a89b4 100644 --- a/include/alpaka/dev/DevCudaRt.hpp +++ b/include/alpaka/dev/DevCudaRt.hpp @@ -13,15 +13,6 @@ namespace alpaka { //! The CUDA RT device handle. using DevCudaRt = DevUniformCudaHipRt; - - namespace trait - { - template<> - struct MemVisibility - { - using type = alpaka::MemVisibleGpuCudaRt; - }; - } // namespace trait } // namespace alpaka #endif // ALPAKA_ACC_GPU_CUDA_ENABLED diff --git a/include/alpaka/dev/DevFpgaSyclIntel.hpp b/include/alpaka/dev/DevFpgaSyclIntel.hpp index 5e4c17ad2502..516027db6b2a 100644 --- a/include/alpaka/dev/DevFpgaSyclIntel.hpp +++ b/include/alpaka/dev/DevFpgaSyclIntel.hpp @@ -12,15 +12,6 @@ namespace alpaka { using DevFpgaSyclIntel = DevGenericSycl; - - namespace trait - { - template<> - struct MemVisibility - { - using type = alpaka::MemVisibleFpgaSyclIntel; - }; - } // namespace trait } // namespace alpaka #endif diff --git a/include/alpaka/dev/DevGpuSyclIntel.hpp b/include/alpaka/dev/DevGpuSyclIntel.hpp index 0b8786a35eb7..9897d40ebbc5 100644 --- a/include/alpaka/dev/DevGpuSyclIntel.hpp +++ b/include/alpaka/dev/DevGpuSyclIntel.hpp @@ -12,15 +12,6 @@ namespace alpaka { using DevGpuSyclIntel = DevGenericSycl; - - namespace trait - { - template<> - struct MemVisibility - { - using type = alpaka::MemVisibleGpuSyclIntel; - }; - } // namespace trait } // namespace alpaka #endif diff --git a/include/alpaka/dev/DevHipRt.hpp b/include/alpaka/dev/DevHipRt.hpp index 075fb3cf5efd..819c2f5c4543 100644 --- a/include/alpaka/dev/DevHipRt.hpp +++ b/include/alpaka/dev/DevHipRt.hpp @@ -13,15 +13,6 @@ namespace alpaka { //! The HIP RT device handle. using DevHipRt = DevUniformCudaHipRt; - - namespace trait - { - template<> - struct MemVisibility - { - using type = alpaka::MemVisibleGpuHipRt; - }; - } // namespace trait } // namespace alpaka #endif // ALPAKA_ACC_GPU_HIP_ENABLED diff --git a/include/alpaka/platform/PlatformCpu.hpp b/include/alpaka/platform/PlatformCpu.hpp index c431fd418785..aab9f3a3790f 100644 --- a/include/alpaka/platform/PlatformCpu.hpp +++ b/include/alpaka/platform/PlatformCpu.hpp @@ -6,6 +6,7 @@ #include "alpaka/core/Concepts.hpp" #include "alpaka/dev/DevCpu.hpp" +#include "alpaka/mem/Visibility.hpp" #include "alpaka/platform/Traits.hpp" #include @@ -65,5 +66,11 @@ namespace alpaka return {}; } }; + + template<> + struct MemVisibility + { + using type = alpaka::MemVisibleCPU; + }; } // namespace trait } // namespace alpaka diff --git a/include/alpaka/platform/PlatformCpuSycl.hpp b/include/alpaka/platform/PlatformCpuSycl.hpp index db055f9689b2..486076851c73 100644 --- a/include/alpaka/platform/PlatformCpuSycl.hpp +++ b/include/alpaka/platform/PlatformCpuSycl.hpp @@ -6,6 +6,7 @@ #include "alpaka/dev/DevGenericSycl.hpp" #include "alpaka/dev/Traits.hpp" +#include "alpaka/mem/Visibility.hpp" #include "alpaka/platform/PlatformGenericSycl.hpp" #include @@ -39,6 +40,12 @@ namespace alpaka::trait { using type = DevGenericSycl; // = DevCpuSycl }; + + template<> + struct MemVisibility + { + using type = alpaka::MemVisibleCPU; + }; } // namespace alpaka::trait #endif diff --git a/include/alpaka/platform/PlatformCudaRt.hpp b/include/alpaka/platform/PlatformCudaRt.hpp index 9bf76fa66682..011f782c2d2d 100644 --- a/include/alpaka/platform/PlatformCudaRt.hpp +++ b/include/alpaka/platform/PlatformCudaRt.hpp @@ -5,6 +5,7 @@ #pragma once #include "alpaka/core/ApiCudaRt.hpp" +#include "alpaka/mem/Visibility.hpp" #include "alpaka/platform/PlatformUniformCudaHipRt.hpp" #ifdef ALPAKA_ACC_GPU_CUDA_ENABLED @@ -13,6 +14,15 @@ namespace alpaka { //! The CUDA RT platform. using PlatformCudaRt = PlatformUniformCudaHipRt; + + namespace trait + { + template<> + struct MemVisibility + { + using type = alpaka::MemVisibleGpuCudaRt; + }; + } // namespace trait } // namespace alpaka #endif // ALPAKA_ACC_GPU_CUDA_ENABLED diff --git a/include/alpaka/platform/PlatformFpgaSyclIntel.hpp b/include/alpaka/platform/PlatformFpgaSyclIntel.hpp index 21ee2c257c5f..b6190b75312c 100644 --- a/include/alpaka/platform/PlatformFpgaSyclIntel.hpp +++ b/include/alpaka/platform/PlatformFpgaSyclIntel.hpp @@ -6,6 +6,7 @@ #include "alpaka/dev/DevGenericSycl.hpp" #include "alpaka/dev/Traits.hpp" +#include "alpaka/mem/Visibility.hpp" #include "alpaka/platform/PlatformGenericSycl.hpp" #if defined(ALPAKA_ACC_SYCL_ENABLED) && defined(ALPAKA_SYCL_ONEAPI_FPGA) @@ -57,6 +58,12 @@ namespace alpaka::trait { using type = DevGenericSycl; // = DevFpgaSyclIntel }; + + template<> + struct MemVisibility + { + using type = alpaka::MemVisibleFpgaSyclIntel; + }; } // namespace alpaka::trait #endif diff --git a/include/alpaka/platform/PlatformGpuSyclIntel.hpp b/include/alpaka/platform/PlatformGpuSyclIntel.hpp index 216bb5ae26a1..7cfabd09e303 100644 --- a/include/alpaka/platform/PlatformGpuSyclIntel.hpp +++ b/include/alpaka/platform/PlatformGpuSyclIntel.hpp @@ -6,6 +6,7 @@ #include "alpaka/dev/DevGenericSycl.hpp" #include "alpaka/dev/Traits.hpp" +#include "alpaka/mem/Visibility.hpp" #include "alpaka/platform/PlatformGenericSycl.hpp" #include @@ -42,6 +43,12 @@ namespace alpaka::trait { using type = DevGenericSycl; // = DevGpuSyclIntel }; + + template<> + struct MemVisibility + { + using type = alpaka::MemVisibleGpuSyclIntel; + }; } // namespace alpaka::trait #endif diff --git a/include/alpaka/platform/PlatformHipRt.hpp b/include/alpaka/platform/PlatformHipRt.hpp index 25303aeaf10e..4179bd8f0967 100644 --- a/include/alpaka/platform/PlatformHipRt.hpp +++ b/include/alpaka/platform/PlatformHipRt.hpp @@ -5,6 +5,7 @@ #pragma once #include "alpaka/core/ApiHipRt.hpp" +#include "alpaka/mem/Visibility.hpp" #include "alpaka/platform/PlatformUniformCudaHipRt.hpp" #ifdef ALPAKA_ACC_GPU_HIP_ENABLED @@ -13,6 +14,15 @@ namespace alpaka { //! The HIP RT platform. using PlatformHipRt = PlatformUniformCudaHipRt; + + namespace trait + { + template<> + struct MemVisibility + { + using type = alpaka::MemVisibleGpuHipRt; + }; + } // namespace trait } // namespace alpaka #endif // ALPAKA_ACC_GPU_HIP_ENABLED diff --git a/test/unit/mem/src/Visibility.cpp b/test/unit/mem/src/Visibility.cpp index e26400b586e6..8e4ecd11539a 100644 --- a/test/unit/mem/src/Visibility.cpp +++ b/test/unit/mem/src/Visibility.cpp @@ -45,16 +45,16 @@ TEMPLATE_LIST_TEST_CASE("memoryVisibilityType", "[mem][visibility]", TagList) using Tag = TestType; if constexpr(alpaka::AccIsEnabled::value) { - using DevType = decltype(alpaka::getDevByIdx(alpaka::Platform>{}, 0)); + using PltfType = alpaka::Platform>; if constexpr(isCPUTag()) { STATIC_REQUIRE( - std::is_same_v::type, alpaka::MemVisibleCPU>); + std::is_same_v::type, alpaka::MemVisibleCPU>); } else if(std::is_same_v) { STATIC_REQUIRE( - std::is_same_v::type, alpaka::MemVisibleGpuCudaRt>); + std::is_same_v::type, alpaka::MemVisibleGpuCudaRt>); } } } @@ -91,8 +91,8 @@ TEMPLATE_LIST_TEST_CASE("printDefines", "[mem][visibility]", TagTagList) BufAcc1 bufDev1(alpaka::allocBuf(dev1, Idx(1))); BufAcc2 bufDev2(alpaka::allocBuf(dev2, Idx(1))); - STATIC_REQUIRE(alpaka::hasSameMemView(dev1, bufDev1)); - STATIC_REQUIRE(alpaka::hasSameMemView(dev2, bufDev2)); + STATIC_REQUIRE(alpaka::hasSameMemView(plt1, bufDev1)); + STATIC_REQUIRE(alpaka::hasSameMemView(plt2, bufDev2)); // at the moment, only the cpu platform has different accelerator types // therefore all cpu accelerators can access the memory of other cpu accelerators @@ -100,18 +100,18 @@ TEMPLATE_LIST_TEST_CASE("printDefines", "[mem][visibility]", TagTagList) // same to support access to the memory of each other if constexpr((isCPUTag() && isCPUTag()) || std::is_same_v) { - STATIC_REQUIRE(alpaka::hasSameMemView(dev1, bufDev2)); - STATIC_REQUIRE(alpaka::hasSameMemView(dev2, bufDev1)); + STATIC_REQUIRE(alpaka::hasSameMemView(plt1, bufDev2)); + STATIC_REQUIRE(alpaka::hasSameMemView(plt2, bufDev1)); } else { - STATIC_REQUIRE_FALSE(alpaka::hasSameMemView(dev1, bufDev2)); - STATIC_REQUIRE_FALSE(alpaka::hasSameMemView(dev2, bufDev1)); + STATIC_REQUIRE_FALSE(alpaka::hasSameMemView(plt1, bufDev2)); + STATIC_REQUIRE_FALSE(alpaka::hasSameMemView(plt2, bufDev1)); } - do_job(dev1, bufDev1); - do_job(dev2, bufDev2); - // do_job(dev1, bufDev2); + // do_job(dev1, bufDev1); + // do_job(dev2, bufDev2); + // do_job(dev1, bufDev2); // std::cout << std::boolalpha << "tag 1 is cpu: " << isCPUTag() << "\n"; // std::cout << std::boolalpha << "tag 2 is cpu: " << isCPUTag() << "\n"; From d5dbce5d23ec9e86968cf5cacaf1687f22d1b506 Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Tue, 7 May 2024 17:24:34 +0200 Subject: [PATCH 07/19] extend hasSameMemView to supported devices and accelerators --- include/alpaka/mem/Visibility.hpp | 32 +++++-- include/alpaka/mem/view/ViewAccessOps.hpp | 1 + test/unit/mem/CMakeLists.txt | 2 +- test/unit/mem/src/Visibility.cpp | 111 +++++++++++----------- 4 files changed, 81 insertions(+), 65 deletions(-) diff --git a/include/alpaka/mem/Visibility.hpp b/include/alpaka/mem/Visibility.hpp index 4775215677f8..568c959da010 100644 --- a/include/alpaka/mem/Visibility.hpp +++ b/include/alpaka/mem/Visibility.hpp @@ -4,6 +4,7 @@ #pragma once +#include "alpaka/mem/view/ViewAccessOps.hpp" #include "alpaka/meta/ForEachType.hpp" #include "alpaka/meta/IsTuple.hpp" #include "alpaka/meta/TypeListOps.hpp" @@ -53,10 +54,14 @@ namespace alpaka }; } // namespace detail - template + template< + typename T, + typename = std::enable_if_t< + alpaka::isPlatform> || alpaka::isDevice> + || alpaka::isAccelerator> || alpaka::internal::isView>>> [[maybe_unused]] static std::string getMemVisiblityName() { - using MemVisibilityType = typename alpaka::trait::MemVisibility>::type; + using MemVisibilityType = typename alpaka::trait::MemVisibility>::type; if constexpr(alpaka::meta::isTuple()) { std::vector vs; @@ -89,12 +94,27 @@ namespace alpaka return getMemVisiblityName(); } - template + template< + typename T, + typename TBuf, + typename = std::enable_if_t< + (alpaka::isPlatform> || alpaka::isDevice> + || alpaka::isAccelerator>) &&alpaka::internal::isView>>> inline constexpr bool hasSameMemView() { - return alpaka::meta::Contains< - typename alpaka::trait::MemVisibility::type, - typename alpaka::trait::MemVisibility::type>::value; + if constexpr(alpaka::isDevice> || alpaka::isAccelerator>) + { + using Platform = alpaka::Platform; + return alpaka::meta::Contains< + typename alpaka::trait::MemVisibility::type, + typename alpaka::trait::MemVisibility::type>::value; + } + else + { + return alpaka::meta::Contains< + typename alpaka::trait::MemVisibility::type, + typename alpaka::trait::MemVisibility::type>::value; + } } template diff --git a/include/alpaka/mem/view/ViewAccessOps.hpp b/include/alpaka/mem/view/ViewAccessOps.hpp index 27056678dd17..36723d8902d6 100644 --- a/include/alpaka/mem/view/ViewAccessOps.hpp +++ b/include/alpaka/mem/view/ViewAccessOps.hpp @@ -20,6 +20,7 @@ namespace alpaka::internal inline constexpr bool isView = false; // TODO(bgruber): replace this by a concept in C++20 + // TODO(simeonehrig): extend the trait by memory Visiblity type template inline constexpr bool isView< TView, diff --git a/test/unit/mem/CMakeLists.txt b/test/unit/mem/CMakeLists.txt index 16399eca3414..f4e76bcb7262 100644 --- a/test/unit/mem/CMakeLists.txt +++ b/test/unit/mem/CMakeLists.txt @@ -9,7 +9,7 @@ cmake_minimum_required(VERSION 3.22) -set(_TARGET_NAME "memTest") +set(_TARGET_NAME "memVisibilityTest") append_recursive_files_add_to_src_group("src/" "src/" "cpp" _FILES_SOURCE) diff --git a/test/unit/mem/src/Visibility.cpp b/test/unit/mem/src/Visibility.cpp index 8e4ecd11539a..5d58bc996480 100644 --- a/test/unit/mem/src/Visibility.cpp +++ b/test/unit/mem/src/Visibility.cpp @@ -28,7 +28,7 @@ constexpr bool isCPUTag() } } -using TagList = std::tuple< +using AccTags = std::tuple< alpaka::TagCpuSerial, alpaka::TagCpuThreads, alpaka::TagCpuTbbBlocks, @@ -40,26 +40,25 @@ using TagList = std::tuple< alpaka::TagFpgaSyclIntel, alpaka::TagGpuSyclIntel>; -TEMPLATE_LIST_TEST_CASE("memoryVisibilityType", "[mem][visibility]", TagList) +using EnabledAccTags = alpaka::meta::Filter; + +TEMPLATE_LIST_TEST_CASE("memoryVisibilityType", "[mem][visibility]", EnabledAccTags) { using Tag = TestType; - if constexpr(alpaka::AccIsEnabled::value) + + using PltfType = alpaka::Platform>; + if constexpr(isCPUTag()) + { + STATIC_REQUIRE(std::is_same_v::type, alpaka::MemVisibleCPU>); + } + else if(std::is_same_v) { - using PltfType = alpaka::Platform>; - if constexpr(isCPUTag()) - { - STATIC_REQUIRE( - std::is_same_v::type, alpaka::MemVisibleCPU>); - } - else if(std::is_same_v) - { - STATIC_REQUIRE( - std::is_same_v::type, alpaka::MemVisibleGpuCudaRt>); - } + STATIC_REQUIRE( + std::is_same_v::type, alpaka::MemVisibleGpuCudaRt>); } } -using TagTagList = alpaka::meta::CartesianProduct; +using EnabledTagTagList = alpaka::meta::CartesianProduct; template void do_job(TDev dev, TBuf buffer) @@ -67,53 +66,49 @@ void do_job(TDev dev, TBuf buffer) STATIC_REQUIRE(alpaka::hasSameMemView(dev, buffer)); } -TEMPLATE_LIST_TEST_CASE("printDefines", "[mem][visibility]", TagTagList) +TEMPLATE_LIST_TEST_CASE("testHasSameMemView", "[mem][visibility]", EnabledTagTagList) { using Tag1 = std::tuple_element_t<0, TestType>; using Tag2 = std::tuple_element_t<1, TestType>; - if constexpr(alpaka::AccIsEnabled::value && alpaka::AccIsEnabled::value) + SUCCEED(Tag1::get_name() << " + " << Tag2::get_name()); + + using Acc1 = alpaka::TagToAcc; + using Acc2 = alpaka::TagToAcc; + + auto const plt1 = alpaka::Platform{}; + auto const plt2 = alpaka::Platform{}; + + auto const dev1 = alpaka::getDevByIdx(plt1, 0); + auto const dev2 = alpaka::getDevByIdx(plt2, 0); + + using BufAcc1 = alpaka::Buf; + using BufAcc2 = alpaka::Buf; + + BufAcc1 bufDev1(alpaka::allocBuf(dev1, Idx(1))); + BufAcc2 bufDev2(alpaka::allocBuf(dev2, Idx(1))); + + STATIC_REQUIRE(alpaka::hasSameMemView(plt1, bufDev1)); + STATIC_REQUIRE(alpaka::hasSameMemView(dev1, bufDev1)); + STATIC_REQUIRE(alpaka::hasSameMemView()); + STATIC_REQUIRE(alpaka::hasSameMemView(plt2, bufDev2)); + STATIC_REQUIRE(alpaka::hasSameMemView(dev2, bufDev2)); + STATIC_REQUIRE(alpaka::hasSameMemView()); + + // at the moment, only the cpu platform has different accelerator types + // therefore all cpu accelerators can access the memory of other cpu accelerators + // if the accelerator is not a cpu accelerator, both accelerators needs to be the + // same to support access to the memory of each other + if constexpr((isCPUTag() && isCPUTag()) || std::is_same_v) + { + STATIC_REQUIRE(alpaka::hasSameMemView(plt1, bufDev2)); + STATIC_REQUIRE(alpaka::hasSameMemView(plt2, bufDev1)); + STATIC_REQUIRE(alpaka::hasSameMemView()); + } + else { - SUCCEED(Tag1::get_name() << " + " << Tag2::get_name()); - - using Acc1 = alpaka::TagToAcc; - using Acc2 = alpaka::TagToAcc; - - auto const plt1 = alpaka::Platform{}; - auto const plt2 = alpaka::Platform{}; - - auto const dev1 = alpaka::getDevByIdx(plt1, 0); - auto const dev2 = alpaka::getDevByIdx(plt2, 0); - - using BufAcc1 = alpaka::Buf; - using BufAcc2 = alpaka::Buf; - - BufAcc1 bufDev1(alpaka::allocBuf(dev1, Idx(1))); - BufAcc2 bufDev2(alpaka::allocBuf(dev2, Idx(1))); - - STATIC_REQUIRE(alpaka::hasSameMemView(plt1, bufDev1)); - STATIC_REQUIRE(alpaka::hasSameMemView(plt2, bufDev2)); - - // at the moment, only the cpu platform has different accelerator types - // therefore all cpu accelerators can access the memory of other cpu accelerators - // if the accelerator is not a cpu accelerator, both accelerators needs to be the - // same to support access to the memory of each other - if constexpr((isCPUTag() && isCPUTag()) || std::is_same_v) - { - STATIC_REQUIRE(alpaka::hasSameMemView(plt1, bufDev2)); - STATIC_REQUIRE(alpaka::hasSameMemView(plt2, bufDev1)); - } - else - { - STATIC_REQUIRE_FALSE(alpaka::hasSameMemView(plt1, bufDev2)); - STATIC_REQUIRE_FALSE(alpaka::hasSameMemView(plt2, bufDev1)); - } - - // do_job(dev1, bufDev1); - // do_job(dev2, bufDev2); - // do_job(dev1, bufDev2); - - // std::cout << std::boolalpha << "tag 1 is cpu: " << isCPUTag() << "\n"; - // std::cout << std::boolalpha << "tag 2 is cpu: " << isCPUTag() << "\n"; + STATIC_REQUIRE_FALSE(alpaka::hasSameMemView(plt1, bufDev2)); + STATIC_REQUIRE_FALSE(alpaka::hasSameMemView(plt2, bufDev1)); + STATIC_REQUIRE_FALSE(alpaka::hasSameMemView()); } } From 8e9ce27505e575264582ab8ccab8fed26e10e491 Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Wed, 8 May 2024 15:28:16 +0200 Subject: [PATCH 08/19] improve test --- test/unit/mem/src/Visibility.cpp | 69 ++++++++++++-------------------- 1 file changed, 26 insertions(+), 43 deletions(-) diff --git a/test/unit/mem/src/Visibility.cpp b/test/unit/mem/src/Visibility.cpp index 5d58bc996480..6615893f2a37 100644 --- a/test/unit/mem/src/Visibility.cpp +++ b/test/unit/mem/src/Visibility.cpp @@ -12,59 +12,42 @@ using Dim = alpaka::DimInt<1>; using Idx = std::size_t; -template -constexpr bool isCPUTag() -{ - if constexpr( - std::is_same_v || std::is_same_v - || std::is_same_v || std::is_same_v - || std::is_same_v) - { - return true; - } - else - { - return false; - } -} - -using AccTags = std::tuple< - alpaka::TagCpuSerial, - alpaka::TagCpuThreads, - alpaka::TagCpuTbbBlocks, - alpaka::TagCpuOmp2Blocks, - alpaka::TagCpuOmp2Threads, - alpaka::TagGpuCudaRt, - alpaka::TagGpuHipRt, - alpaka::TagCpuSycl, - alpaka::TagFpgaSyclIntel, - alpaka::TagGpuSyclIntel>; - -using EnabledAccTags = alpaka::meta::Filter; - -TEMPLATE_LIST_TEST_CASE("memoryVisibilityType", "[mem][visibility]", EnabledAccTags) +TEMPLATE_LIST_TEST_CASE("memoryVisibilityType", "[mem][visibility]", alpaka::EnabledAccTags) { using Tag = TestType; + + REQUIRE(true); using PltfType = alpaka::Platform>; - if constexpr(isCPUTag()) + + if constexpr(alpaka::isCpuTag::value) { - STATIC_REQUIRE(std::is_same_v::type, alpaka::MemVisibleCPU>); + REQUIRE(std::is_same_v::type, alpaka::MemVisibleCPU>); } else if(std::is_same_v) { - STATIC_REQUIRE( - std::is_same_v::type, alpaka::MemVisibleGpuCudaRt>); + REQUIRE(std::is_same_v::type, alpaka::MemVisibleGpuCudaRt>); + } + else if(std::is_same_v) + { + REQUIRE(std::is_same_v::type, alpaka::MemVisibleGpuSyclIntel>); + } + else if(std::is_same_v) + { + REQUIRE(std::is_same_v::type, alpaka::MemVisibleGpuHipRt>); + } + else if(std::is_same_v) + { + REQUIRE(std::is_same_v::type, alpaka::MemVisibleGenericSycl>); + } + else if(std::is_same_v) + { + REQUIRE( + std::is_same_v::type, alpaka::MemVisibleFpgaSyclIntel>); } } -using EnabledTagTagList = alpaka::meta::CartesianProduct; - -template -void do_job(TDev dev, TBuf buffer) -{ - STATIC_REQUIRE(alpaka::hasSameMemView(dev, buffer)); -} +using EnabledTagTagList = alpaka::meta::CartesianProduct; TEMPLATE_LIST_TEST_CASE("testHasSameMemView", "[mem][visibility]", EnabledTagTagList) { @@ -99,7 +82,7 @@ TEMPLATE_LIST_TEST_CASE("testHasSameMemView", "[mem][visibility]", EnabledTagTag // therefore all cpu accelerators can access the memory of other cpu accelerators // if the accelerator is not a cpu accelerator, both accelerators needs to be the // same to support access to the memory of each other - if constexpr((isCPUTag() && isCPUTag()) || std::is_same_v) + if constexpr((alpaka::isCpuTag::value && alpaka::isCpuTag::value) || std::is_same_v) { STATIC_REQUIRE(alpaka::hasSameMemView(plt1, bufDev2)); STATIC_REQUIRE(alpaka::hasSameMemView(plt2, bufDev1)); From e2b827d4c939ffcdd5a59bd393a386775c7a539f Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Wed, 8 May 2024 15:46:46 +0200 Subject: [PATCH 09/19] improve memoryVisibilityType again --- test/unit/mem/src/Visibility.cpp | 47 ++++++++++++-------------------- 1 file changed, 18 insertions(+), 29 deletions(-) diff --git a/test/unit/mem/src/Visibility.cpp b/test/unit/mem/src/Visibility.cpp index 6615893f2a37..65657de70ddd 100644 --- a/test/unit/mem/src/Visibility.cpp +++ b/test/unit/mem/src/Visibility.cpp @@ -12,38 +12,27 @@ using Dim = alpaka::DimInt<1>; using Idx = std::size_t; -TEMPLATE_LIST_TEST_CASE("memoryVisibilityType", "[mem][visibility]", alpaka::EnabledAccTags) +using ExpectedTagsMemVisibilities = std::tuple< + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple>; + +TEMPLATE_LIST_TEST_CASE("memoryVisibilityType", "[mem][visibility]", ExpectedTagsMemVisibilities) { - using Tag = TestType; + using Tag = std::tuple_element_t<0, TestType>; + using ExpectedMemVisibility = std::tuple_element_t<1, TestType>; - - REQUIRE(true); - using PltfType = alpaka::Platform>; - - if constexpr(alpaka::isCpuTag::value) - { - REQUIRE(std::is_same_v::type, alpaka::MemVisibleCPU>); - } - else if(std::is_same_v) - { - REQUIRE(std::is_same_v::type, alpaka::MemVisibleGpuCudaRt>); - } - else if(std::is_same_v) - { - REQUIRE(std::is_same_v::type, alpaka::MemVisibleGpuSyclIntel>); - } - else if(std::is_same_v) - { - REQUIRE(std::is_same_v::type, alpaka::MemVisibleGpuHipRt>); - } - else if(std::is_same_v) - { - REQUIRE(std::is_same_v::type, alpaka::MemVisibleGenericSycl>); - } - else if(std::is_same_v) + if constexpr(alpaka::AccIsEnabled::value) { - REQUIRE( - std::is_same_v::type, alpaka::MemVisibleFpgaSyclIntel>); + using PltfType = alpaka::Platform>; + STATIC_REQUIRE(std::is_same_v::type, ExpectedMemVisibility>); } } From f20c2a108bfcd8554379caa0f034613c1ddab739 Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Tue, 14 May 2024 11:20:52 +0200 Subject: [PATCH 10/19] add memory visibility to raw pointer view --- include/alpaka/alpaka.hpp | 1 - include/alpaka/mem/Visibility.hpp | 11 +- include/alpaka/mem/view/ViewPlainPtr.hpp | 70 ++++++++----- include/alpaka/meta/IsTuple.hpp | 29 ------ include/alpaka/meta/TypeListOps.hpp | 33 +++++- test/unit/mem/src/Visibility.cpp | 125 ++++++++++++++++++++--- 6 files changed, 193 insertions(+), 76 deletions(-) delete mode 100644 include/alpaka/meta/IsTuple.hpp diff --git a/include/alpaka/alpaka.hpp b/include/alpaka/alpaka.hpp index b9396429c4f9..b0458c00249e 100644 --- a/include/alpaka/alpaka.hpp +++ b/include/alpaka/alpaka.hpp @@ -172,7 +172,6 @@ #include "alpaka/meta/Integral.hpp" #include "alpaka/meta/IsArrayOrVector.hpp" #include "alpaka/meta/IsStrictBase.hpp" -#include "alpaka/meta/IsTuple.hpp" #include "alpaka/meta/NdLoop.hpp" #include "alpaka/meta/NonZero.hpp" #include "alpaka/meta/Set.hpp" diff --git a/include/alpaka/mem/Visibility.hpp b/include/alpaka/mem/Visibility.hpp index 568c959da010..90f941457125 100644 --- a/include/alpaka/mem/Visibility.hpp +++ b/include/alpaka/mem/Visibility.hpp @@ -4,10 +4,12 @@ #pragma once +#include "alpaka/acc/Traits.hpp" +#include "alpaka/dev/Traits.hpp" #include "alpaka/mem/view/ViewAccessOps.hpp" #include "alpaka/meta/ForEachType.hpp" -#include "alpaka/meta/IsTuple.hpp" #include "alpaka/meta/TypeListOps.hpp" +#include "alpaka/platform/Traits.hpp" #include #include @@ -17,7 +19,7 @@ #define CREATE_MEM_VISIBILITY(mem_name) \ struct mem_name \ { \ - static std::string name() \ + static std::string get_name() \ { \ return #mem_name; \ } \ @@ -49,7 +51,7 @@ namespace alpaka template void operator()(std::vector& vs) { - vs.push_back(TTYPE::name()); + vs.push_back(TTYPE::get_name()); } }; } // namespace detail @@ -84,7 +86,7 @@ namespace alpaka } else { - return MemVisibilityType::name(); + return MemVisibilityType::get_name(); } } @@ -115,6 +117,7 @@ namespace alpaka typename alpaka::trait::MemVisibility::type, typename alpaka::trait::MemVisibility::type>::value; } + ALPAKA_UNREACHABLE({}); } template diff --git a/include/alpaka/mem/view/ViewPlainPtr.hpp b/include/alpaka/mem/view/ViewPlainPtr.hpp index ceb4d95aed32..97817822f460 100644 --- a/include/alpaka/mem/view/ViewPlainPtr.hpp +++ b/include/alpaka/mem/view/ViewPlainPtr.hpp @@ -8,9 +8,18 @@ #include "alpaka/dev/DevCpu.hpp" #include "alpaka/dev/DevGenericSycl.hpp" #include "alpaka/dev/DevUniformCudaHipRt.hpp" +#include "alpaka/mem/Visibility.hpp" #include "alpaka/mem/view/Traits.hpp" #include "alpaka/mem/view/ViewAccessOps.hpp" #include "alpaka/meta/DependentFalseType.hpp" +#include "alpaka/meta/TypeListOps.hpp" +#include "alpaka/platform/PlatformCpu.hpp" +#include "alpaka/platform/PlatformCpuSycl.hpp" +#include "alpaka/platform/PlatformCudaRt.hpp" +#include "alpaka/platform/PlatformFpgaSyclIntel.hpp" +#include "alpaka/platform/PlatformGenericSycl.hpp" +#include "alpaka/platform/PlatformGpuSyclIntel.hpp" +#include "alpaka/platform/PlatformHipRt.hpp" #include "alpaka/vec/Vec.hpp" #include @@ -19,8 +28,14 @@ namespace alpaka { //! The memory view to wrap plain pointers. - template - struct ViewPlainPtr final : internal::ViewAccessOps> + template< + typename TDev, + typename TElem, + typename TDim, + typename TIdx, + typename TMemVisibility = + typename alpaka::meta::toTuple>::type>::type> + struct ViewPlainPtr final : internal::ViewAccessOps> { static_assert(!std::is_const_v, "The idx type of the view can not be const!"); @@ -49,15 +64,21 @@ namespace alpaka namespace trait { //! The ViewPlainPtr device type trait specialization. - template - struct DevType> + template + struct DevType> { using type = alpaka::Dev; }; + template + struct MemVisibility> + { + using type = TMemVisibility; + }; + //! The ViewPlainPtr device get trait specialization. - template - struct GetDev> + template + struct GetDev> { static auto getDev(ViewPlainPtr const& view) -> alpaka::Dev { @@ -66,15 +87,15 @@ namespace alpaka }; //! The ViewPlainPtr dimension getter trait. - template - struct DimType> + template + struct DimType> { using type = TDim; }; //! The ViewPlainPtr memory element type get trait specialization. - template - struct ElemType> + template + struct ElemType> { using type = TElem; }; @@ -82,32 +103,32 @@ namespace alpaka namespace trait { - template - struct GetExtents> + template + struct GetExtents> { - ALPAKA_FN_HOST auto operator()(ViewPlainPtr const& view) const + ALPAKA_FN_HOST auto operator()(ViewPlainPtr const& view) const { return view.m_extentElements; } }; //! The ViewPlainPtr native pointer get trait specialization. - template - struct GetPtrNative> + template + struct GetPtrNative> { - static auto getPtrNative(ViewPlainPtr const& view) -> TElem const* + static auto getPtrNative(ViewPlainPtr const& view) -> TElem const* { return view.m_pMem; } - static auto getPtrNative(ViewPlainPtr& view) -> TElem* + static auto getPtrNative(ViewPlainPtr& view) -> TElem* { return view.m_pMem; } }; - template - struct GetPitchesInBytes> + template + struct GetPitchesInBytes> { ALPAKA_FN_HOST auto operator()(ViewPlainPtr const& view) const { @@ -174,18 +195,19 @@ namespace alpaka }; #endif //! The ViewPlainPtr offset get trait specialization. - template - struct GetOffsets> + template + struct GetOffsets> { - ALPAKA_FN_HOST auto operator()(ViewPlainPtr const&) const -> Vec + ALPAKA_FN_HOST auto operator()(ViewPlainPtr const&) const + -> Vec { return Vec::zeros(); } }; //! The ViewPlainPtr idx type trait specialization. - template - struct IdxType> + template + struct IdxType> { using type = TIdx; }; diff --git a/include/alpaka/meta/IsTuple.hpp b/include/alpaka/meta/IsTuple.hpp deleted file mode 100644 index ee4a920923ba..000000000000 --- a/include/alpaka/meta/IsTuple.hpp +++ /dev/null @@ -1,29 +0,0 @@ -/* Copyright 2024 Simeon Ehrig - * SPDX-License-Identifier: MPL-2.0 - */ - -#pragma once - -#include -#include - -// copied from https://stackoverflow.com/a/51073558/22035743 -namespace alpaka::meta -{ - template - struct IsTuple : std::false_type - { - }; - - template - struct IsTuple> : std::true_type - { - }; - - template - constexpr bool isTuple() - { - return IsTuple>::value; - } - -} // namespace alpaka::meta diff --git a/include/alpaka/meta/TypeListOps.hpp b/include/alpaka/meta/TypeListOps.hpp index 2d6bcfe7f45f..de9a59bb2453 100644 --- a/include/alpaka/meta/TypeListOps.hpp +++ b/include/alpaka/meta/TypeListOps.hpp @@ -1,9 +1,10 @@ -/* Copyright 2022 Bernhard Manfred Gruber +/* Copyright 2024 Bernhard Manfred Gruber, Simeon Ehrig * SPDX-License-Identifier: MPL-2.0 */ #pragma once +#include #include namespace alpaka::meta @@ -35,4 +36,34 @@ namespace alpaka::meta { static constexpr bool value = std::is_same_v || Contains, Value>::value; }; + + // copied from https://stackoverflow.com/a/51073558/22035743 + template + struct IsTuple : std::false_type + { + }; + + template + struct IsTuple> : std::true_type + { + }; + + template + constexpr bool isTuple() + { + return IsTuple>::value; + } + + template + struct toTuple + { + using type = std::tuple; + }; + + template + struct toTuple> + { + using type = std::tuple; + }; + } // namespace alpaka::meta diff --git a/test/unit/mem/src/Visibility.cpp b/test/unit/mem/src/Visibility.cpp index 65657de70ddd..3098c76aa249 100644 --- a/test/unit/mem/src/Visibility.cpp +++ b/test/unit/mem/src/Visibility.cpp @@ -12,28 +12,60 @@ using Dim = alpaka::DimInt<1>; using Idx = std::size_t; -using ExpectedTagsMemVisibilities = std::tuple< - std::tuple, - std::tuple, - std::tuple, - std::tuple, - std::tuple, - std::tuple, - std::tuple, - std::tuple, - std::tuple, - std::tuple>; +// TODO(SimeonEhrig): Replace implementation. Instead using a list, specialize `alpaka::Platform` for +// tags to get the Memory Visiblity + +//! \brief check if the accelerator related to the tag is bounded to the cpu platform +//! \tparam TTag alpaka tag type +template +struct isCpuTag : std::false_type +{ +}; + +template +struct isCpuTag< + TTag, + std::enable_if_t< + // TAGCpuSycl is not included because it has it's own platform + std::is_same_v || std::is_same_v + || std::is_same_v || std::is_same_v + || std::is_same_v>> : std::true_type +{ +}; + +template +struct AccIsEnabledMemVisibilities : std::false_type +{ +}; + +template +struct AccIsEnabledMemVisibilities< + TTagMemView, + std::void_t, alpaka::DimInt<1>, int>>> : std::true_type +{ +}; + +using ExpectedTagsMemVisibilities = alpaka::meta::Filter< + std::tuple< + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple>, + AccIsEnabledMemVisibilities>; TEMPLATE_LIST_TEST_CASE("memoryVisibilityType", "[mem][visibility]", ExpectedTagsMemVisibilities) { using Tag = std::tuple_element_t<0, TestType>; using ExpectedMemVisibility = std::tuple_element_t<1, TestType>; - if constexpr(alpaka::AccIsEnabled::value) - { - using PltfType = alpaka::Platform>; - STATIC_REQUIRE(std::is_same_v::type, ExpectedMemVisibility>); - } + using PltfType = alpaka::Platform>; + STATIC_REQUIRE(std::is_same_v::type, ExpectedMemVisibility>); } using EnabledTagTagList = alpaka::meta::CartesianProduct; @@ -71,7 +103,7 @@ TEMPLATE_LIST_TEST_CASE("testHasSameMemView", "[mem][visibility]", EnabledTagTag // therefore all cpu accelerators can access the memory of other cpu accelerators // if the accelerator is not a cpu accelerator, both accelerators needs to be the // same to support access to the memory of each other - if constexpr((alpaka::isCpuTag::value && alpaka::isCpuTag::value) || std::is_same_v) + if constexpr((isCpuTag::value && isCpuTag::value) || std::is_same_v) { STATIC_REQUIRE(alpaka::hasSameMemView(plt1, bufDev2)); STATIC_REQUIRE(alpaka::hasSameMemView(plt2, bufDev1)); @@ -84,3 +116,62 @@ TEMPLATE_LIST_TEST_CASE("testHasSameMemView", "[mem][visibility]", EnabledTagTag STATIC_REQUIRE_FALSE(alpaka::hasSameMemView()); } } + +using EnabledTagTagMemVisibilityList + = alpaka::meta::CartesianProduct; + +TEMPLATE_LIST_TEST_CASE("testMemView", "[mem][visibility]", EnabledTagTagMemVisibilityList) +{ + using Tag1 = std::tuple_element_t<0, std::tuple_element_t<0, TestType>>; + using ExpectedMemVisibilityForTag1 = std::tuple_element_t<1, std::tuple_element_t<0, TestType>>; + using Tag2 = std::tuple_element_t<0, std::tuple_element_t<1, TestType>>; + using ExpectedMemVisibilityForTag2 = std::tuple_element_t<1, std::tuple_element_t<1, TestType>>; + + SUCCEED( + "Tag1: " << Tag1::get_name() << " + " << ExpectedMemVisibilityForTag1::get_name() + << "\nTag2: " << Tag2::get_name() << " + " << ExpectedMemVisibilityForTag1::get_name()); + + + constexpr Idx data_size = 10; + + using Acc1 = alpaka::TagToAcc; + using Acc2 = alpaka::TagToAcc; + + auto const plt1 = alpaka::Platform{}; + auto const plt2 = alpaka::Platform{}; + + auto const dev1 = alpaka::getDevByIdx(plt1, 0); + auto const dev2 = alpaka::getDevByIdx(plt2, 0); + + using Vec1D = alpaka::Vec, Idx>; + Vec1D const extents(Vec1D::all(data_size)); + + std::array data; + + // STATIC_REQUIRE(std::is_same_v>::type, std::tuple>); STATIC_REQUIRE(std::is_same_v::type, std::tuple>); + // STATIC_REQUIRE(std::is_same_v::type, std::tuple>); + + auto data_view1 = alpaka::createView(dev1, data.data(), extents); + STATIC_REQUIRE(std::is_same_v< + typename alpaka::trait::MemVisibility::type, + std::tuple>); + STATIC_REQUIRE(alpaka::hasSameMemView(plt1, data_view1)); + + auto data_view2 = alpaka::createView(dev2, data.data(), extents); + STATIC_REQUIRE(std::is_same_v< + typename alpaka::trait::MemVisibility::type, + std::tuple>); + STATIC_REQUIRE(alpaka::hasSameMemView(plt2, data_view2)); + + if constexpr(std::is_same_v) + { + STATIC_REQUIRE(alpaka::hasSameMemView(plt1, data_view2)); + STATIC_REQUIRE(alpaka::hasSameMemView(plt2, data_view1)); + } + else + { + STATIC_REQUIRE_FALSE(alpaka::hasSameMemView(plt1, data_view2)); + STATIC_REQUIRE_FALSE(alpaka::hasSameMemView(plt2, data_view1)); + } +} From fa496fd9337d34b3c40fa43e5d9ece1922d2476c Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Tue, 14 May 2024 13:23:35 +0200 Subject: [PATCH 11/19] sycl fixes CI_FILTER: ^linux_icpx --- include/alpaka/mem/buf/BufCpuSycl.hpp | 11 +++++++---- include/alpaka/mem/buf/BufGenericSycl.hpp | 4 ++-- test/unit/mem/src/Visibility.cpp | 4 ---- 3 files changed, 9 insertions(+), 10 deletions(-) diff --git a/include/alpaka/mem/buf/BufCpuSycl.hpp b/include/alpaka/mem/buf/BufCpuSycl.hpp index a04a94e66629..b3560e571b7f 100644 --- a/include/alpaka/mem/buf/BufCpuSycl.hpp +++ b/include/alpaka/mem/buf/BufCpuSycl.hpp @@ -14,11 +14,14 @@ namespace alpaka template using BufCpuSycl = BufGenericSycl; - template - struct MemVisibility> + namespace trait { - using type = std::tuple; - }; + template + struct MemVisibility> + { + using type = std::tuple; + }; + } // namespace trait } // namespace alpaka #endif diff --git a/include/alpaka/mem/buf/BufGenericSycl.hpp b/include/alpaka/mem/buf/BufGenericSycl.hpp index a0180c46a7fd..f26c45ba9bf6 100644 --- a/include/alpaka/mem/buf/BufGenericSycl.hpp +++ b/include/alpaka/mem/buf/BufGenericSycl.hpp @@ -259,8 +259,8 @@ namespace alpaka::trait } }; - template - struct MemVisibility> + template + struct MemVisibility> { using type = std::tuple; }; diff --git a/test/unit/mem/src/Visibility.cpp b/test/unit/mem/src/Visibility.cpp index 3098c76aa249..fbb5066101af 100644 --- a/test/unit/mem/src/Visibility.cpp +++ b/test/unit/mem/src/Visibility.cpp @@ -148,10 +148,6 @@ TEMPLATE_LIST_TEST_CASE("testMemView", "[mem][visibility]", EnabledTagTagMemVisi std::array data; - // STATIC_REQUIRE(std::is_same_v>::type, std::tuple>); STATIC_REQUIRE(std::is_same_v::type, std::tuple>); - // STATIC_REQUIRE(std::is_same_v::type, std::tuple>); - auto data_view1 = alpaka::createView(dev1, data.data(), extents); STATIC_REQUIRE(std::is_same_v< typename alpaka::trait::MemVisibility::type, From 341d76573d179b5652cd895ed91981f17b0c7092 Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Tue, 14 May 2024 13:56:10 +0200 Subject: [PATCH 12/19] fix CI_FILTER: ^linux_icpx --- include/alpaka/mem/Visibility.hpp | 1 + include/alpaka/platform/PlatformCpuSycl.hpp | 2 +- test/unit/mem/src/Visibility.cpp | 2 +- 3 files changed, 3 insertions(+), 2 deletions(-) diff --git a/include/alpaka/mem/Visibility.hpp b/include/alpaka/mem/Visibility.hpp index 90f941457125..381687f13a59 100644 --- a/include/alpaka/mem/Visibility.hpp +++ b/include/alpaka/mem/Visibility.hpp @@ -30,6 +30,7 @@ namespace alpaka CREATE_MEM_VISIBILITY(MemVisibleCPU); CREATE_MEM_VISIBILITY(MemVisibleFpgaSyclIntel); CREATE_MEM_VISIBILITY(MemVisibleGenericSycl); + CREATE_MEM_VISIBILITY(MemVisibleCpuSycl); CREATE_MEM_VISIBILITY(MemVisibleGpuCudaRt); CREATE_MEM_VISIBILITY(MemVisibleGpuHipRt); CREATE_MEM_VISIBILITY(MemVisibleGpuSyclIntel); diff --git a/include/alpaka/platform/PlatformCpuSycl.hpp b/include/alpaka/platform/PlatformCpuSycl.hpp index 486076851c73..893ce8e39dc1 100644 --- a/include/alpaka/platform/PlatformCpuSycl.hpp +++ b/include/alpaka/platform/PlatformCpuSycl.hpp @@ -44,7 +44,7 @@ namespace alpaka::trait template<> struct MemVisibility { - using type = alpaka::MemVisibleCPU; + using type = alpaka::MemVisibleCpuSycl; }; } // namespace alpaka::trait diff --git a/test/unit/mem/src/Visibility.cpp b/test/unit/mem/src/Visibility.cpp index fbb5066101af..2849486d2ea6 100644 --- a/test/unit/mem/src/Visibility.cpp +++ b/test/unit/mem/src/Visibility.cpp @@ -54,7 +54,7 @@ using ExpectedTagsMemVisibilities = alpaka::meta::Filter< std::tuple, std::tuple, std::tuple, - std::tuple, + std::tuple, std::tuple, std::tuple>, AccIsEnabledMemVisibilities>; From 3a3fbf519d29edff0f214d22129f8ed760d98ebe Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Tue, 14 May 2024 15:00:17 +0200 Subject: [PATCH 13/19] fix 2 CI_FILTER: ^linux_icpx --- include/alpaka/mem/buf/BufCpuSycl.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/alpaka/mem/buf/BufCpuSycl.hpp b/include/alpaka/mem/buf/BufCpuSycl.hpp index b3560e571b7f..dc1c767d4952 100644 --- a/include/alpaka/mem/buf/BufCpuSycl.hpp +++ b/include/alpaka/mem/buf/BufCpuSycl.hpp @@ -19,7 +19,7 @@ namespace alpaka template struct MemVisibility> { - using type = std::tuple; + using type = std::tuple; }; } // namespace trait } // namespace alpaka From db2b4c6f57fddb6aeaea17c4b84949fb5dad3a38 Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Tue, 14 May 2024 16:26:05 +0200 Subject: [PATCH 14/19] implement memory visibility as template type for bufCPU CI_FILTER: ^nope --- include/alpaka/dev/DevCpu.hpp | 8 +-- include/alpaka/mem/buf/BufCpu.hpp | 94 +++++++++++++++++++------------ include/alpaka/mem/buf/Traits.hpp | 12 +++- 3 files changed, 71 insertions(+), 43 deletions(-) diff --git a/include/alpaka/dev/DevCpu.hpp b/include/alpaka/dev/DevCpu.hpp index e36c263072fa..b1cda711541b 100644 --- a/include/alpaka/dev/DevCpu.hpp +++ b/include/alpaka/dev/DevCpu.hpp @@ -167,16 +167,16 @@ namespace alpaka }; } // namespace trait - template + template class BufCpu; namespace trait { //! The CPU device memory buffer type trait specialization. - template - struct BufType + template + struct BufType { - using type = BufCpu; + using type = BufCpu; }; //! The CPU device platform type trait specialization. diff --git a/include/alpaka/mem/buf/BufCpu.hpp b/include/alpaka/mem/buf/BufCpu.hpp index 282ef672cfb9..178d90716892 100644 --- a/include/alpaka/mem/buf/BufCpu.hpp +++ b/include/alpaka/mem/buf/BufCpu.hpp @@ -86,8 +86,8 @@ namespace alpaka } // namespace detail //! The CPU memory buffer. - template - class BufCpu : public internal::ViewAccessOps> + template + class BufCpu : public internal::ViewAccessOps> { public: template @@ -104,67 +104,69 @@ namespace alpaka namespace trait { //! The BufCpu device type trait specialization. - template - struct DevType> + template + struct DevType> { using type = DevCpu; }; //! The BufCpu device get trait specialization. - template - struct GetDev> + template + struct GetDev> { - ALPAKA_FN_HOST static auto getDev(BufCpu const& buf) -> DevCpu + ALPAKA_FN_HOST static auto getDev(BufCpu const& buf) -> DevCpu { return buf.m_spBufCpuImpl->m_dev; } }; //! The BufCpu dimension getter trait. - template - struct DimType> + template + struct DimType> { using type = TDim; }; //! The BufCpu memory element type get trait specialization. - template - struct ElemType> + template + struct ElemType> { using type = TElem; }; //! The BufCpu width get trait specialization. - template - struct GetExtents> + template + struct GetExtents> { - ALPAKA_FN_HOST auto operator()(BufCpu const& buf) + ALPAKA_FN_HOST auto operator()(BufCpu const& buf) { return buf.m_spBufCpuImpl->m_extentElements; } }; //! The BufCpu native pointer get trait specialization. - template - struct GetPtrNative> + template + struct GetPtrNative> { - ALPAKA_FN_HOST static auto getPtrNative(BufCpu const& buf) -> TElem const* + ALPAKA_FN_HOST static auto getPtrNative(BufCpu const& buf) + -> TElem const* { return buf.m_spBufCpuImpl->m_pMem; } - ALPAKA_FN_HOST static auto getPtrNative(BufCpu& buf) -> TElem* + ALPAKA_FN_HOST static auto getPtrNative(BufCpu& buf) -> TElem* { return buf.m_spBufCpuImpl->m_pMem; } }; //! The BufCpu pointer on device get trait specialization. - template - struct GetPtrDev, DevCpu> + template + struct GetPtrDev, DevCpu, TMemVisibility> { - ALPAKA_FN_HOST static auto getPtrDev(BufCpu const& buf, DevCpu const& dev) - -> TElem const* + ALPAKA_FN_HOST static auto getPtrDev( + BufCpu const& buf, + DevCpu const& dev) -> TElem const* { if(dev == getDev(buf)) { @@ -176,7 +178,8 @@ namespace alpaka } } - ALPAKA_FN_HOST static auto getPtrDev(BufCpu& buf, DevCpu const& dev) -> TElem* + ALPAKA_FN_HOST static auto getPtrDev(BufCpu& buf, DevCpu const& dev) + -> TElem* { if(dev == getDev(buf)) { @@ -194,7 +197,8 @@ namespace alpaka struct BufAlloc { template - ALPAKA_FN_HOST static auto allocBuf(DevCpu const& dev, TExtent const& extent) -> BufCpu + ALPAKA_FN_HOST static auto allocBuf(DevCpu const& dev, TExtent const& extent) + -> BufCpu>::type> { ALPAKA_DEBUG_MINIMAL_LOG_SCOPE; @@ -216,7 +220,15 @@ namespace alpaka auto* memPtr = alpaka::malloc(Allocator{}, static_cast(getExtentProduct(extent))); auto deleter = [](TElem* ptr) { alpaka::free(Allocator{}, ptr); }; - return BufCpu(dev, memPtr, std::move(deleter), extent); + return BufCpu< + TElem, + TDim, + TIdx, + typename alpaka::trait::MemVisibility>::type>( + dev, + memPtr, + std::move(deleter), + extent); } }; @@ -225,7 +237,8 @@ namespace alpaka struct AsyncBufAlloc { template - ALPAKA_FN_HOST static auto allocAsyncBuf(TQueue queue, TExtent const& extent) -> BufCpu + ALPAKA_FN_HOST static auto allocAsyncBuf(TQueue queue, TExtent const& extent) + -> BufCpu>::type> { ALPAKA_DEBUG_MINIMAL_LOG_SCOPE; @@ -261,7 +274,15 @@ namespace alpaka }); }; - return BufCpu(dev, memPtr, std::move(deleter), extent); + return BufCpu< + TElem, + TDim, + TIdx, + typename alpaka::trait::MemVisibility>::type>( + dev, + memPtr, + std::move(deleter), + extent); } }; @@ -279,7 +300,8 @@ namespace alpaka ALPAKA_FN_HOST static auto allocMappedBuf( DevCpu const& host, PlatformCpu const& /*platform*/, - TExtent const& extent) -> BufCpu + // TODO: needs to Visibility of DevCpu and PlatformCpu + TExtent const& extent) -> BufCpu::type> { // Allocate standard host memory. return allocBuf(host, extent); @@ -293,26 +315,26 @@ namespace alpaka }; //! The BufCpu offset get trait specialization. - template - struct GetOffsets> + template + struct GetOffsets> { - ALPAKA_FN_HOST auto operator()(BufCpu const&) const -> Vec + ALPAKA_FN_HOST auto operator()(BufCpu const&) const -> Vec { return Vec::zeros(); } }; //! The BufCpu idx type trait specialization. - template - struct IdxType> + template + struct IdxType> { using type = TIdx; }; - template - struct MemVisibility> + template + struct MemVisibility> { - using type = std::tuple; + using type = TMemVisibility; }; } // namespace trait diff --git a/include/alpaka/mem/buf/Traits.hpp b/include/alpaka/mem/buf/Traits.hpp index 33e7c9bda7f1..fca1a07417ad 100644 --- a/include/alpaka/mem/buf/Traits.hpp +++ b/include/alpaka/mem/buf/Traits.hpp @@ -19,7 +19,13 @@ namespace alpaka namespace trait { //! The memory buffer type trait. - template + template< + typename TDev, + typename TElem, + typename TDim, + typename TIdx, + typename TMemVisibility, + typename TSfinae = void> struct BufType; //! The memory allocator trait. @@ -48,8 +54,8 @@ namespace alpaka } // namespace trait //! The memory buffer type trait alias template to remove the ::type. - template - using Buf = typename trait::BufType, TElem, TDim, TIdx>::type; + template + using Buf = typename trait::BufType, TElem, TDim, TIdx, TMemVisibility>::type; //! Allocates memory on the given device. //! From 527f0d2c85bc60a4a1294eb35652d694b4172fa8 Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Thu, 16 May 2024 17:23:10 +0200 Subject: [PATCH 15/19] finish implementing visibility as template type for buffer types - tests if the correct visibility type is set for allocated memory is missing, except for `allocBuf` - the state of the commit is, it should compile with all backends and does not break existing tests --- example/bufferCopy/src/bufferCopy.cpp | 5 +- example/convolution1D/src/convolution1D.cpp | 2 +- .../counterBasedRng/src/counterBasedRng.cpp | 3 +- example/heatEquation/src/heatEquation.cpp | 5 +- .../src/monteCarloIntegration.cpp | 6 +- example/randomCells2D/src/randomCells2D.cpp | 12 +- .../randomStrategies/src/randomStrategies.cpp | 8 +- example/reduce/src/reduce.cpp | 9 +- example/vectorAdd/src/vectorAdd.cpp | 8 +- include/alpaka/dev/DevGenericSycl.hpp | 8 +- include/alpaka/dev/DevUniformCudaHipRt.hpp | 8 +- include/alpaka/mem/Visibility.hpp | 25 +++- include/alpaka/mem/buf/BufCpu.hpp | 26 ++-- include/alpaka/mem/buf/BufCpuSycl.hpp | 13 +- include/alpaka/mem/buf/BufCudaRt.hpp | 13 +- include/alpaka/mem/buf/BufFpgaSyclIntel.hpp | 10 +- include/alpaka/mem/buf/BufGenericSycl.hpp | 99 +++++++++------ include/alpaka/mem/buf/BufGpuSyclIntel.hpp | 13 +- include/alpaka/mem/buf/BufHipRt.hpp | 13 +- .../alpaka/mem/buf/BufUniformCudaHipRt.hpp | 120 ++++++++++++------ include/alpaka/mem/view/ViewPlainPtr.hpp | 3 +- include/alpaka/meta/TypeListOps.hpp | 49 +++++-- test/unit/math/src/Buffer.hpp | 29 ++++- test/unit/mem/copy/src/BufSlicing.cpp | 4 +- test/unit/mem/src/Visibility.cpp | 7 +- 25 files changed, 290 insertions(+), 208 deletions(-) diff --git a/example/bufferCopy/src/bufferCopy.cpp b/example/bufferCopy/src/bufferCopy.cpp index 12bd4a1a2acb..ad9d7cc08143 100644 --- a/example/bufferCopy/src/bufferCopy.cpp +++ b/example/bufferCopy/src/bufferCopy.cpp @@ -148,8 +148,7 @@ auto main() -> int // // The `alloc` method returns a reference counted buffer handle. // When the last such handle is destroyed, the memory is freed automatically. - using BufHost = alpaka::Buf; - BufHost hostBuffer(alpaka::allocBuf(devHost, extents)); + auto hostBuffer(alpaka::allocBuf(devHost, extents)); // You can also use already allocated memory and wrap it within a view (irrespective of the device type). // The view does not own the underlying memory. So you have to make sure that // the view does not outlive its underlying memory. @@ -159,7 +158,7 @@ auto main() -> int // Allocate accelerator memory buffers // // The interface to allocate a buffer is the same on the host and on the device. - using BufAcc = alpaka::Buf; + using BufAcc = alpaka::Buf>; BufAcc deviceBuffer1(alpaka::allocBuf(devAcc, extents)); BufAcc deviceBuffer2(alpaka::allocBuf(devAcc, extents)); diff --git a/example/convolution1D/src/convolution1D.cpp b/example/convolution1D/src/convolution1D.cpp index 047f462ef0a7..ed3885bfff27 100644 --- a/example/convolution1D/src/convolution1D.cpp +++ b/example/convolution1D/src/convolution1D.cpp @@ -83,7 +83,7 @@ auto main() -> int using DevAcc = alpaka::ExampleDefaultAcc; using QueueProperty = alpaka::Blocking; using QueueAcc = alpaka::Queue; - using BufAcc = alpaka::Buf; + using BufAcc = alpaka::Buf>; std::cout << "Using alpaka accelerator: " << alpaka::getAccName() << '\n'; diff --git a/example/counterBasedRng/src/counterBasedRng.cpp b/example/counterBasedRng/src/counterBasedRng.cpp index 86da223f1d0a..0f2489639323 100644 --- a/example/counterBasedRng/src/counterBasedRng.cpp +++ b/example/counterBasedRng/src/counterBasedRng.cpp @@ -162,8 +162,7 @@ auto main() -> int CounterBasedRngKernel::Key key = {rd(), rd()}; // Allocate buffer on the accelerator - using BufAcc = alpaka::Buf; - BufAcc bufAcc(alpaka::allocBuf(devAcc, extent)); + auto bufAcc(alpaka::allocBuf(devAcc, extent)); // Create the kernel execution task. auto const taskKernelAcc = alpaka::createTaskKernel( diff --git a/example/heatEquation/src/heatEquation.cpp b/example/heatEquation/src/heatEquation.cpp index ff7ee6c7dafe..4e97e15f6974 100644 --- a/example/heatEquation/src/heatEquation.cpp +++ b/example/heatEquation/src/heatEquation.cpp @@ -121,9 +121,8 @@ auto main() -> int double* const pNextHost = std::data(uNextBufHost); // Accelerator buffer - using BufAcc = alpaka::Buf; - auto uNextBufAcc = BufAcc{alpaka::allocBuf(devAcc, extent)}; - auto uCurrBufAcc = BufAcc{alpaka::allocBuf(devAcc, extent)}; + auto uNextBufAcc{alpaka::allocBuf(devAcc, extent)}; + auto uCurrBufAcc{alpaka::allocBuf(devAcc, extent)}; double* pCurrAcc = std::data(uCurrBufAcc); double* pNextAcc = std::data(uNextBufAcc); diff --git a/example/monteCarloIntegration/src/monteCarloIntegration.cpp b/example/monteCarloIntegration/src/monteCarloIntegration.cpp index 52e050785c88..1059c75b8a7b 100644 --- a/example/monteCarloIntegration/src/monteCarloIntegration.cpp +++ b/example/monteCarloIntegration/src/monteCarloIntegration.cpp @@ -88,8 +88,6 @@ auto main() -> int using QueueAcc = alpaka::Queue; QueueAcc queue{devAcc}; - using BufHost = alpaka::Buf; - using BufAcc = alpaka::Buf; using WorkDiv = alpaka::WorkDivMembers; // Problem parameter. constexpr size_t numPoints = 1'000'000u; @@ -104,8 +102,8 @@ auto main() -> int alpaka::GridBlockExtentSubDivRestrictions::Unrestricted)}; // Setup buffer. - BufHost bufHost{alpaka::allocBuf(devHost, extent)}; - BufAcc bufAcc{alpaka::allocBuf(devAcc, extent)}; + auto bufHost{alpaka::allocBuf(devHost, extent)}; + auto bufAcc{alpaka::allocBuf(devAcc, extent)}; uint32_t* const ptrBufAcc{std::data(bufAcc)}; // Initialize the global count to 0. diff --git a/example/randomCells2D/src/randomCells2D.cpp b/example/randomCells2D/src/randomCells2D.cpp index a0a21370cecc..1f614f7ce516 100644 --- a/example/randomCells2D/src/randomCells2D.cpp +++ b/example/randomCells2D/src/randomCells2D.cpp @@ -156,12 +156,12 @@ auto main() -> int using QueueAcc = alpaka::Queue; QueueAcc queue{devAcc}; - using BufHost = alpaka::Buf; - using BufAcc = alpaka::Buf; - using BufHostRand = alpaka::Buf; - using BufAccRand = alpaka::Buf; - using BufHostRandVec = alpaka::Buf; - using BufAccRandVec = alpaka::Buf; + using BufHost = alpaka::Buf>; + using BufAcc = alpaka::Buf>; + using BufHostRand = alpaka::Buf>; + using BufAccRand = alpaka::Buf>; + using BufHostRandVec = alpaka::Buf>; + using BufAccRandVec = alpaka::Buf>; using WorkDiv = alpaka::WorkDivMembers; constexpr Idx numX = NUM_X; diff --git a/example/randomStrategies/src/randomStrategies.cpp b/example/randomStrategies/src/randomStrategies.cpp index 3dd0c9efde5e..43f01d0fe00f 100644 --- a/example/randomStrategies/src/randomStrategies.cpp +++ b/example/randomStrategies/src/randomStrategies.cpp @@ -44,8 +44,8 @@ struct Box QueueAcc queue; ///< default accelerator queue // buffers holding the PRNG states - using BufHostRand = alpaka::Buf; - using BufAccRand = alpaka::Buf; + using BufHostRand = alpaka::Buf>; + using BufAccRand = alpaka::Buf>; Vec const extentRand; ///< size of the buffer of PRNG states WorkDiv workdivRand; ///< work division for PRNG buffer initialization @@ -53,8 +53,8 @@ struct Box BufAccRand bufAccRand; ///< device side PRNG states buffer // buffers holding the "simulation" results - using BufHost = alpaka::Buf; - using BufAcc = alpaka::Buf; + using BufHost = alpaka::Buf>; + using BufAcc = alpaka::Buf>; Vec const extentResult; ///< size of the results buffer WorkDiv workdivResult; ///< work division of the result calculation diff --git a/example/reduce/src/reduce.cpp b/example/reduce/src/reduce.cpp index 2d5fe2c1dc21..d132995ef9d7 100644 --- a/example/reduce/src/reduce.cpp +++ b/example/reduce/src/reduce.cpp @@ -49,7 +49,7 @@ auto reduce( DevAcc devAcc, QueueAcc queue, uint64_t n, - alpaka::Buf hostMemory, + alpaka::Buf> hostMemory, TFunc func) -> T { static constexpr uint64_t blockSize = getMaxBlockSize(); @@ -62,10 +62,11 @@ auto reduce( if(blockCount > maxBlockCount) blockCount = maxBlockCount; - alpaka::Buf sourceDeviceMemory = alpaka::allocBuf(devAcc, n); + using DevBuf = alpaka::Buf>; - alpaka::Buf destinationDeviceMemory - = alpaka::allocBuf(devAcc, static_cast(blockCount)); + DevBuf sourceDeviceMemory = alpaka::allocBuf(devAcc, n); + + DevBuf destinationDeviceMemory = alpaka::allocBuf(devAcc, static_cast(blockCount)); // copy the data to the GPU alpaka::memcpy(queue, sourceDeviceMemory, hostMemory, n); diff --git a/example/vectorAdd/src/vectorAdd.cpp b/example/vectorAdd/src/vectorAdd.cpp index 5eca205279c8..1639a7602616 100644 --- a/example/vectorAdd/src/vectorAdd.cpp +++ b/example/vectorAdd/src/vectorAdd.cpp @@ -111,8 +111,8 @@ auto main() -> int auto const devHost = alpaka::getDevByIdx(platformHost, 0); // Allocate 3 host memory buffers - using BufHost = alpaka::Buf; - BufHost bufHostA(alpaka::allocBuf(devHost, extent)); + auto bufHostA(alpaka::allocBuf(devHost, extent)); + using BufHost = decltype(bufHostA); BufHost bufHostB(alpaka::allocBuf(devHost, extent)); BufHost bufHostC(alpaka::allocBuf(devHost, extent)); @@ -129,8 +129,8 @@ auto main() -> int } // Allocate 3 buffers on the accelerator - using BufAcc = alpaka::Buf; - BufAcc bufAccA(alpaka::allocBuf(devAcc, extent)); + auto bufAccA(alpaka::allocBuf(devAcc, extent)); + using BufAcc = decltype(bufAccA); BufAcc bufAccB(alpaka::allocBuf(devAcc, extent)); BufAcc bufAccC(alpaka::allocBuf(devAcc, extent)); diff --git a/include/alpaka/dev/DevGenericSycl.hpp b/include/alpaka/dev/DevGenericSycl.hpp index 729090f8f2d3..55902d840f15 100644 --- a/include/alpaka/dev/DevGenericSycl.hpp +++ b/include/alpaka/dev/DevGenericSycl.hpp @@ -32,7 +32,7 @@ namespace alpaka { - template + template class BufGenericSycl; namespace detail @@ -219,10 +219,10 @@ namespace alpaka::trait }; //! The SYCL device memory buffer type trait specialization. - template - struct BufType, TElem, TDim, TIdx> + template + struct BufType, TElem, TDim, TIdx, TMemVisibility> { - using type = BufGenericSycl; + using type = BufGenericSycl; }; //! The SYCL device platform type trait specialization. diff --git a/include/alpaka/dev/DevUniformCudaHipRt.hpp b/include/alpaka/dev/DevUniformCudaHipRt.hpp index 876d8ca5a434..dfbc1dc62d77 100644 --- a/include/alpaka/dev/DevUniformCudaHipRt.hpp +++ b/include/alpaka/dev/DevUniformCudaHipRt.hpp @@ -48,7 +48,7 @@ namespace alpaka template struct PlatformUniformCudaHipRt; - template + template struct BufUniformCudaHipRt; //! The CUDA/HIP RT device handle. @@ -222,10 +222,10 @@ namespace alpaka }; //! The CUDA/HIP RT device memory buffer type trait specialization. - template - struct BufType, TElem, TDim, TIdx> + template + struct BufType, TElem, TDim, TIdx, TMemVisibility> { - using type = BufUniformCudaHipRt; + using type = BufUniformCudaHipRt; }; //! The CUDA/HIP RT device platform type trait specialization. diff --git a/include/alpaka/mem/Visibility.hpp b/include/alpaka/mem/Visibility.hpp index 381687f13a59..f5cf7c0d024a 100644 --- a/include/alpaka/mem/Visibility.hpp +++ b/include/alpaka/mem/Visibility.hpp @@ -65,7 +65,7 @@ namespace alpaka [[maybe_unused]] static std::string getMemVisiblityName() { using MemVisibilityType = typename alpaka::trait::MemVisibility>::type; - if constexpr(alpaka::meta::isTuple()) + if constexpr(alpaka::meta::isList) { std::vector vs; alpaka::meta::forEachType(detail::AppendMemTypeName{}, vs); @@ -126,4 +126,27 @@ namespace alpaka { return hasSameMemView, std::decay_t>(); } + + namespace detail + { + template + struct MemVisibilityHelper + { + using type = typename alpaka::trait::MemVisibility::type; + }; + + template + struct MemVisibilityHelper< + T, + std::enable_if_t> || alpaka::isAccelerator>>> + { + using type = typename alpaka::trait::MemVisibility>>::type; + }; + } // namespace detail + + template + using MemVisibility = typename alpaka::detail::MemVisibilityHelper>::type; + + template + using MemVisibilityTypeList = alpaka::meta::toTuple>>; } // namespace alpaka diff --git a/include/alpaka/mem/buf/BufCpu.hpp b/include/alpaka/mem/buf/BufCpu.hpp index 178d90716892..1045933965b4 100644 --- a/include/alpaka/mem/buf/BufCpu.hpp +++ b/include/alpaka/mem/buf/BufCpu.hpp @@ -16,6 +16,7 @@ #include "alpaka/mem/buf/Traits.hpp" #include "alpaka/mem/view/ViewAccessOps.hpp" #include "alpaka/meta/DependentFalseType.hpp" +#include "alpaka/meta/Unique.hpp" #include "alpaka/platform/PlatformCpu.hpp" #include "alpaka/vec/Vec.hpp" @@ -198,7 +199,7 @@ namespace alpaka { template ALPAKA_FN_HOST static auto allocBuf(DevCpu const& dev, TExtent const& extent) - -> BufCpu>::type> + -> BufCpu>> { ALPAKA_DEBUG_MINIMAL_LOG_SCOPE; @@ -220,11 +221,7 @@ namespace alpaka auto* memPtr = alpaka::malloc(Allocator{}, static_cast(getExtentProduct(extent))); auto deleter = [](TElem* ptr) { alpaka::free(Allocator{}, ptr); }; - return BufCpu< - TElem, - TDim, - TIdx, - typename alpaka::trait::MemVisibility>::type>( + return BufCpu>>( dev, memPtr, std::move(deleter), @@ -238,7 +235,7 @@ namespace alpaka { template ALPAKA_FN_HOST static auto allocAsyncBuf(TQueue queue, TExtent const& extent) - -> BufCpu>::type> + -> BufCpu>> { ALPAKA_DEBUG_MINIMAL_LOG_SCOPE; @@ -274,11 +271,7 @@ namespace alpaka }); }; - return BufCpu< - TElem, - TDim, - TIdx, - typename alpaka::trait::MemVisibility>::type>( + return BufCpu>>( dev, memPtr, std::move(deleter), @@ -300,8 +293,13 @@ namespace alpaka ALPAKA_FN_HOST static auto allocMappedBuf( DevCpu const& host, PlatformCpu const& /*platform*/, - // TODO: needs to Visibility of DevCpu and PlatformCpu - TExtent const& extent) -> BufCpu::type> + TExtent const& extent) + -> BufCpu< + TElem, + TDim, + TIdx, + alpaka::meta::Unique< + std::tuple, alpaka::MemVisibility>>> { // Allocate standard host memory. return allocBuf(host, extent); diff --git a/include/alpaka/mem/buf/BufCpuSycl.hpp b/include/alpaka/mem/buf/BufCpuSycl.hpp index dc1c767d4952..d6d05f487cb4 100644 --- a/include/alpaka/mem/buf/BufCpuSycl.hpp +++ b/include/alpaka/mem/buf/BufCpuSycl.hpp @@ -11,17 +11,8 @@ namespace alpaka { - template - using BufCpuSycl = BufGenericSycl; - - namespace trait - { - template - struct MemVisibility> - { - using type = std::tuple; - }; - } // namespace trait + template + using BufCpuSycl = BufGenericSycl; } // namespace alpaka #endif diff --git a/include/alpaka/mem/buf/BufCudaRt.hpp b/include/alpaka/mem/buf/BufCudaRt.hpp index 79288b689c3c..f974b2490f3f 100644 --- a/include/alpaka/mem/buf/BufCudaRt.hpp +++ b/include/alpaka/mem/buf/BufCudaRt.hpp @@ -13,17 +13,8 @@ namespace alpaka { - template - using BufCudaRt = BufUniformCudaHipRt; - - namespace trait - { - template - struct MemVisibility> - { - using type = std::tuple; - }; - } // namespace trait + template + using BufCudaRt = BufUniformCudaHipRt; } // namespace alpaka #endif // ALPAKA_ACC_GPU_CUDA_ENABLED diff --git a/include/alpaka/mem/buf/BufFpgaSyclIntel.hpp b/include/alpaka/mem/buf/BufFpgaSyclIntel.hpp index 712406fa8b0e..30238b69cc62 100644 --- a/include/alpaka/mem/buf/BufFpgaSyclIntel.hpp +++ b/include/alpaka/mem/buf/BufFpgaSyclIntel.hpp @@ -11,14 +11,8 @@ namespace alpaka { - template - using BufFpgaSyclIntel = BufGenericSycl; - - template - struct MemVisibility> - { - using type = std::tuple; - }; + template + using BufFpgaSyclIntel = BufGenericSycl; } // namespace alpaka #endif diff --git a/include/alpaka/mem/buf/BufGenericSycl.hpp b/include/alpaka/mem/buf/BufGenericSycl.hpp index f26c45ba9bf6..ae64487dac35 100644 --- a/include/alpaka/mem/buf/BufGenericSycl.hpp +++ b/include/alpaka/mem/buf/BufGenericSycl.hpp @@ -24,8 +24,8 @@ namespace alpaka { //! The SYCL memory buffer. - template - class BufGenericSycl : public internal::ViewAccessOps> + template + class BufGenericSycl : public internal::ViewAccessOps> { public: static_assert( @@ -62,67 +62,68 @@ namespace alpaka namespace alpaka::trait { //! The BufGenericSycl device type trait specialization. - template - struct DevType> + template + struct DevType> { using type = DevGenericSycl; }; //! The BufGenericSycl device get trait specialization. - template - struct GetDev> + template + struct GetDev> { - static auto getDev(BufGenericSycl const& buf) + static auto getDev(BufGenericSycl const& buf) { return buf.m_dev; } }; //! The BufGenericSycl dimension getter trait specialization. - template - struct DimType> + template + struct DimType> { using type = TDim; }; //! The BufGenericSycl memory element type get trait specialization. - template - struct ElemType> + template + struct ElemType> { using type = TElem; }; //! The BufGenericSycl extent get trait specialization. - template - struct GetExtents> + template + struct GetExtents> { - auto operator()(BufGenericSycl const& buf) const + auto operator()(BufGenericSycl const& buf) const { return buf.m_extentElements; } }; //! The BufGenericSycl native pointer get trait specialization. - template - struct GetPtrNative> + template + struct GetPtrNative> { - static auto getPtrNative(BufGenericSycl const& buf) -> TElem const* + static auto getPtrNative(BufGenericSycl const& buf) + -> TElem const* { return buf.m_spMem.get(); } - static auto getPtrNative(BufGenericSycl& buf) -> TElem* + static auto getPtrNative(BufGenericSycl& buf) -> TElem* { return buf.m_spMem.get(); } }; //! The BufGenericSycl pointer on device get trait specialization. - template - struct GetPtrDev, DevGenericSycl> + template + struct GetPtrDev, DevGenericSycl> { static auto getPtrDev( - BufGenericSycl const& buf, + BufGenericSycl const& buf, DevGenericSycl const& dev) -> TElem const* { if(dev == getDev(buf)) @@ -135,8 +136,9 @@ namespace alpaka::trait } } - static auto getPtrDev(BufGenericSycl& buf, DevGenericSycl const& dev) - -> TElem* + static auto getPtrDev( + BufGenericSycl& buf, + DevGenericSycl const& dev) -> TElem* { if(dev == getDev(buf)) { @@ -155,7 +157,7 @@ namespace alpaka::trait { template static auto allocBuf(DevGenericSycl const& dev, TExtent const& extent) - -> BufGenericSycl + -> BufGenericSycl> { ALPAKA_DEBUG_MINIMAL_LOG_SCOPE; @@ -197,7 +199,11 @@ namespace alpaka::trait nativeContext); auto deleter = [ctx = nativeContext](TElem* ptr) { sycl::free(ptr, ctx); }; - return BufGenericSycl(dev, memPtr, std::move(deleter), extent); + return BufGenericSycl>( + dev, + memPtr, + std::move(deleter), + extent); } }; @@ -208,10 +214,10 @@ namespace alpaka::trait }; //! The BufGenericSycl offset get trait specialization. - template - struct GetOffsets> + template + struct GetOffsets> { - auto operator()(BufGenericSycl const&) const -> Vec + auto operator()(BufGenericSycl const&) const -> Vec { return Vec::zeros(); } @@ -222,8 +228,11 @@ namespace alpaka::trait struct BufAllocMapped { template - static auto allocMappedBuf(DevCpu const& host, TPlatform const& platform, TExtent const& extent) - -> BufCpu + static auto allocMappedBuf(DevCpu const& host, TPlatform const& platform, TExtent const& extent) -> BufCpu< + TElem, + TDim, + TIdx, + alpaka::meta::Unique, alpaka::MemVisibility>>> { ALPAKA_DEBUG_MINIMAL_LOG_SCOPE; @@ -233,36 +242,46 @@ namespace alpaka::trait TElem* memPtr = sycl::malloc_host(static_cast(getExtentProduct(extent)), ctx); auto deleter = [ctx](TElem* ptr) { sycl::free(ptr, ctx); }; - return BufCpu(host, memPtr, std::move(deleter), extent); + return BufCpu< + TElem, + TDim, + TIdx, + alpaka::meta::Unique, alpaka::MemVisibility>>>( + host, + memPtr, + std::move(deleter), + extent); } }; //! The BufGenericSycl idx type trait specialization. - template - struct IdxType> + template + struct IdxType> { using type = TIdx; }; //! The BufCpu pointer on SYCL device get trait specialization. - template - struct GetPtrDev, DevGenericSycl> + template + struct GetPtrDev, DevGenericSycl> { - static auto getPtrDev(BufCpu const& buf, DevGenericSycl const&) -> TElem const* + static auto getPtrDev(BufCpu const& buf, DevGenericSycl const&) + -> TElem const* { return getPtrNative(buf); } - static auto getPtrDev(BufCpu& buf, DevGenericSycl const&) -> TElem* + static auto getPtrDev(BufCpu& buf, DevGenericSycl const&) + -> TElem* { return getPtrNative(buf); } }; - template - struct MemVisibility> + template + struct MemVisibility> { - using type = std::tuple; + using type = TMemVisibility; }; } // namespace alpaka::trait diff --git a/include/alpaka/mem/buf/BufGpuSyclIntel.hpp b/include/alpaka/mem/buf/BufGpuSyclIntel.hpp index 91d6bc9193b9..f46d90971387 100644 --- a/include/alpaka/mem/buf/BufGpuSyclIntel.hpp +++ b/include/alpaka/mem/buf/BufGpuSyclIntel.hpp @@ -11,17 +11,8 @@ namespace alpaka { - template - using BufGpuSyclIntel = BufGenericSycl; - - namespace trait - { - template - struct MemVisibility> - { - using type = std::tuple; - }; - } // namespace trait + template + using BufGpuSyclIntel = BufGenericSycl; } // namespace alpaka #endif diff --git a/include/alpaka/mem/buf/BufHipRt.hpp b/include/alpaka/mem/buf/BufHipRt.hpp index 99c57424fe3a..cf392d6c94a9 100644 --- a/include/alpaka/mem/buf/BufHipRt.hpp +++ b/include/alpaka/mem/buf/BufHipRt.hpp @@ -11,17 +11,8 @@ namespace alpaka { - template - using BufHipRt = BufUniformCudaHipRt; - - namespace trait - { - template - struct MemVisibility> - { - using type = std::tuple; - }; - } // namespace trait + template + using BufHipRt = BufUniformCudaHipRt; } // namespace alpaka #endif // ALPAKA_ACC_GPU_HIP_ENABLED diff --git a/include/alpaka/mem/buf/BufUniformCudaHipRt.hpp b/include/alpaka/mem/buf/BufUniformCudaHipRt.hpp index 826edaba7b1e..f52c258405fa 100644 --- a/include/alpaka/mem/buf/BufUniformCudaHipRt.hpp +++ b/include/alpaka/mem/buf/BufUniformCudaHipRt.hpp @@ -29,7 +29,7 @@ namespace alpaka struct ApiCudaRt; struct ApiHipRt; - template + template class BufCpu; namespace detail @@ -50,10 +50,10 @@ namespace alpaka } // namespace detail //! The CUDA/HIP memory buffer. - template + template struct BufUniformCudaHipRt : detail::PitchHolder - , internal::ViewAccessOps> + , internal::ViewAccessOps> { static_assert(!std::is_const_v, "The elem type of the buffer must not be const"); static_assert(!std::is_const_v, "The idx type of the buffer must not be const!"); @@ -90,17 +90,17 @@ namespace alpaka namespace trait { //! The BufUniformCudaHipRt device type trait specialization. - template - struct DevType> + template + struct DevType> { using type = DevUniformCudaHipRt; }; //! The BufUniformCudaHipRt device get trait specialization. - template - struct GetDev> + template + struct GetDev> { - ALPAKA_FN_HOST static auto getDev(BufUniformCudaHipRt const& buf) + ALPAKA_FN_HOST static auto getDev(BufUniformCudaHipRt const& buf) -> DevUniformCudaHipRt { return buf.m_dev; @@ -108,51 +108,53 @@ namespace alpaka }; //! The BufUniformCudaHipRt dimension getter trait specialization. - template - struct DimType> + template + struct DimType> { using type = TDim; }; //! The BufUniformCudaHipRt memory element type get trait specialization. - template - struct ElemType> + template + struct ElemType> { using type = TElem; }; //! The BufUniformCudaHipRt extent get trait specialization. - template - struct GetExtents> + template + struct GetExtents> { - ALPAKA_FN_HOST auto operator()(BufUniformCudaHipRt const& buffer) const + ALPAKA_FN_HOST auto operator()( + BufUniformCudaHipRt const& buffer) const { return buffer.m_extentElements; } }; //! The BufUniformCudaHipRt native pointer get trait specialization. - template - struct GetPtrNative> + template + struct GetPtrNative> { - ALPAKA_FN_HOST static auto getPtrNative(BufUniformCudaHipRt const& buf) - -> TElem const* + ALPAKA_FN_HOST static auto getPtrNative( + BufUniformCudaHipRt const& buf) -> TElem const* { return buf.m_spMem.get(); } - ALPAKA_FN_HOST static auto getPtrNative(BufUniformCudaHipRt& buf) -> TElem* + ALPAKA_FN_HOST static auto getPtrNative(BufUniformCudaHipRt& buf) + -> TElem* { return buf.m_spMem.get(); } }; //! The BufUniformCudaHipRt pointer on device get trait specialization. - template - struct GetPtrDev, DevUniformCudaHipRt> + template + struct GetPtrDev, DevUniformCudaHipRt> { ALPAKA_FN_HOST static auto getPtrDev( - BufUniformCudaHipRt const& buf, + BufUniformCudaHipRt const& buf, DevUniformCudaHipRt const& dev) -> TElem const* { if(dev == getDev(buf)) @@ -166,7 +168,7 @@ namespace alpaka } ALPAKA_FN_HOST static auto getPtrDev( - BufUniformCudaHipRt& buf, + BufUniformCudaHipRt& buf, DevUniformCudaHipRt const& dev) -> TElem* { if(dev == getDev(buf)) @@ -180,11 +182,11 @@ namespace alpaka } }; - template - struct GetPitchesInBytes> + template + struct GetPitchesInBytes> { - ALPAKA_FN_HOST auto operator()(BufUniformCudaHipRt const& buf) const - -> Vec + ALPAKA_FN_HOST auto operator()( + BufUniformCudaHipRt const& buf) const -> Vec { Vec v{}; if constexpr(TDim::value > 0) @@ -207,7 +209,12 @@ namespace alpaka { template ALPAKA_FN_HOST static auto allocBuf(DevUniformCudaHipRt const& dev, TExtent const& extent) - -> BufUniformCudaHipRt + -> BufUniformCudaHipRt< + TApi, + TElem, + Dim, + TIdx, + alpaka::MemVisibilityTypeList>> { ALPAKA_DEBUG_MINIMAL_LOG_SCOPE; @@ -289,7 +296,12 @@ namespace alpaka template ALPAKA_FN_HOST static auto allocAsyncBuf(TQueue queue, [[maybe_unused]] TExtent const& extent) - -> BufUniformCudaHipRt + -> BufUniformCudaHipRt< + TApi, + TElem, + TDim, + TIdx, + alpaka::MemVisibilityTypeList>> { ALPAKA_DEBUG_MINIMAL_LOG_SCOPE; @@ -342,7 +354,14 @@ namespace alpaka ALPAKA_FN_HOST static auto allocMappedBuf( DevCpu const& host, PlatformUniformCudaHipRt const& /*platform*/, - TExtent const& extent) -> BufCpu + TExtent const& extent) + -> BufCpu< + TElem, + TDim, + TIdx, + alpaka::meta::Unique, + alpaka::MemVisibility>>>> { ALPAKA_DEBUG_MINIMAL_LOG_SCOPE; @@ -355,7 +374,17 @@ namespace alpaka TApi::hostMallocMapped | TApi::hostMallocPortable)); auto deleter = [](TElem* ptr) { ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK_NOEXCEPT(TApi::hostFree(ptr)); }; - return BufCpu(host, memPtr, std::move(deleter), extent); + return BufCpu< + TElem, + TDim, + TIdx, + alpaka::meta::Unique, + alpaka::MemVisibility>>>>( + host, + memPtr, + std::move(deleter), + extent); } }; @@ -366,10 +395,10 @@ namespace alpaka }; //! The BufUniformCudaHipRt offset get trait specialization. - template - struct GetOffsets> + template + struct GetOffsets> { - ALPAKA_FN_HOST auto operator()(BufUniformCudaHipRt const&) const + ALPAKA_FN_HOST auto operator()(BufUniformCudaHipRt const&) const -> Vec { return Vec::zeros(); @@ -377,18 +406,24 @@ namespace alpaka }; //! The BufUniformCudaHipRt idx type trait specialization. - template - struct IdxType> + template + struct IdxType> { using type = TIdx; }; + template + struct MemVisibility> + { + using type = alpaka::meta::toTuple; + }; + //! The BufCpu pointer on CUDA/HIP device get trait specialization. - template - struct GetPtrDev, DevUniformCudaHipRt> + template + struct GetPtrDev, DevUniformCudaHipRt> { ALPAKA_FN_HOST static auto getPtrDev( - BufCpu const& buf, + BufCpu const& buf, DevUniformCudaHipRt const&) -> TElem const* { // TODO: Check if the memory is mapped at all! @@ -402,8 +437,9 @@ namespace alpaka return pDev; } - ALPAKA_FN_HOST static auto getPtrDev(BufCpu& buf, DevUniformCudaHipRt const&) - -> TElem* + ALPAKA_FN_HOST static auto getPtrDev( + BufCpu& buf, + DevUniformCudaHipRt const&) -> TElem* { // TODO: Check if the memory is mapped at all! TElem* pDev(nullptr); diff --git a/include/alpaka/mem/view/ViewPlainPtr.hpp b/include/alpaka/mem/view/ViewPlainPtr.hpp index 97817822f460..adeff7558bfc 100644 --- a/include/alpaka/mem/view/ViewPlainPtr.hpp +++ b/include/alpaka/mem/view/ViewPlainPtr.hpp @@ -33,8 +33,7 @@ namespace alpaka typename TElem, typename TDim, typename TIdx, - typename TMemVisibility = - typename alpaka::meta::toTuple>::type>::type> + typename TMemVisibility = alpaka::meta::toTuple>> struct ViewPlainPtr final : internal::ViewAccessOps> { static_assert(!std::is_const_v, "The idx type of the view can not be const!"); diff --git a/include/alpaka/meta/TypeListOps.hpp b/include/alpaka/meta/TypeListOps.hpp index de9a59bb2453..e77ccb5cade8 100644 --- a/include/alpaka/meta/TypeListOps.hpp +++ b/include/alpaka/meta/TypeListOps.hpp @@ -39,31 +39,56 @@ namespace alpaka::meta // copied from https://stackoverflow.com/a/51073558/22035743 template - struct IsTuple : std::false_type + struct IsList : std::false_type { }; - template - struct IsTuple> : std::true_type + template class TList, typename... TTypes> + struct IsList> : std::true_type { }; + //! \brief Checks whether the specified type is a list. List is a type with a variadic number of template types. template - constexpr bool isTuple() + constexpr bool isList = IsList>::value; + + namespace detail { - return IsTuple>::value; - } + template class TListType, typename TType, typename = void> + struct ToListImpl + { + using type = TListType; + }; - template - struct toTuple + template class TListType, typename TList> + struct ToListImpl>> + { + using type = TList; + }; + } // namespace detail + + //! \brief Takes an arbitrary number of types (T) and creates a type list of type TListType with the types (T). If + //! T is a single template parameter and it satisfies alpaka::meta::isList, the type of the structure is T (no type + //! change). + //! \tparam TListType type of the created list + //! \tparam T possible list types or type list + template class TListType, typename... T> + struct ToList; + + template class TListType, typename T> + struct ToList : detail::ToListImpl { - using type = std::tuple; }; - template - struct toTuple> + template class TListType, typename T, typename... Ts> + struct ToList { - using type = std::tuple; + using type = TListType; }; + //! \brief If T is a single argument and a type list (fullfil alpaka::meta::isList), the return type is T. + //! Otherwise, std::tuple is returned with T types as template parameters. + template + using toTuple = typename ToList::type; + } // namespace alpaka::meta diff --git a/test/unit/math/src/Buffer.hpp b/test/unit/math/src/Buffer.hpp index 7687fcfd3ca1..8adcb189fcaf 100644 --- a/test/unit/math/src/Buffer.hpp +++ b/test/unit/math/src/Buffer.hpp @@ -6,12 +6,36 @@ #include "Defines.hpp" +#include +#include #include #include namespace mathtest { + namespace detail + { + template + struct MemVisibilityMappedBufferImpl + { + using type = typename alpaka::MemVisibilityTypeList; + }; + + template + struct MemVisibilityMappedBufferImpl< + THost, + TPlatform, + std::enable_if_t>> + { + using type = typename alpaka::meta::Unique< + std::tuple, alpaka::MemVisibility>>; + }; + } // namespace detail + + template + using MemVisibilityMappedBuffer = typename detail::MemVisibilityMappedBufferImpl::type; + //! Provides alpaka-style buffer with arguments' data. //! TData can be a plain value or a complex data-structure. //! The operator() is overloaded and returns the value from the correct Buffer, @@ -32,11 +56,12 @@ namespace mathtest // Defines using's for alpaka-buffer. using DevHost = alpaka::DevCpu; using PlatformHost = alpaka::Platform; - using BufHost = alpaka::Buf; using DevAcc = alpaka::Dev; using PlatformAcc = alpaka::Platform; - using BufAcc = alpaka::Buf; + + using BufHost = alpaka::Buf>; + using BufAcc = alpaka::Buf>; PlatformHost platformHost; DevHost devHost; diff --git a/test/unit/mem/copy/src/BufSlicing.cpp b/test/unit/mem/copy/src/BufSlicing.cpp index 6169fdaf5ff2..ed2a4dc6f1f9 100644 --- a/test/unit/mem/copy/src/BufSlicing.cpp +++ b/test/unit/mem/copy/src/BufSlicing.cpp @@ -28,8 +28,8 @@ struct TestContainer using DevHost = alpaka::DevCpu; using PlatformHost = alpaka::Platform; - using BufHost = alpaka::Buf; - using BufDevice = alpaka::Buf; + using BufHost = alpaka::Buf>; + using BufDevice = alpaka::Buf>; using SubView = alpaka::ViewSubView; diff --git a/test/unit/mem/src/Visibility.cpp b/test/unit/mem/src/Visibility.cpp index 2849486d2ea6..2a46ab35577a 100644 --- a/test/unit/mem/src/Visibility.cpp +++ b/test/unit/mem/src/Visibility.cpp @@ -83,11 +83,14 @@ TEMPLATE_LIST_TEST_CASE("testHasSameMemView", "[mem][visibility]", EnabledTagTag auto const plt1 = alpaka::Platform{}; auto const plt2 = alpaka::Platform{}; + using Plt1 = decltype(plt1); + using Plt2 = decltype(plt2); + auto const dev1 = alpaka::getDevByIdx(plt1, 0); auto const dev2 = alpaka::getDevByIdx(plt2, 0); - using BufAcc1 = alpaka::Buf; - using BufAcc2 = alpaka::Buf; + using BufAcc1 = alpaka::Buf>; + using BufAcc2 = alpaka::Buf>; BufAcc1 bufDev1(alpaka::allocBuf(dev1, Idx(1))); BufAcc2 bufDev2(alpaka::allocBuf(dev2, Idx(1))); From 4a09464998e2ae0e8b0d74e60042888b3d07af75 Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Tue, 21 May 2024 09:30:23 +0200 Subject: [PATCH 16/19] fix --- benchmarks/babelstream/src/AlpakaStream.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/benchmarks/babelstream/src/AlpakaStream.h b/benchmarks/babelstream/src/AlpakaStream.h index ba556b028dba..0524fb37a0ce 100644 --- a/benchmarks/babelstream/src/AlpakaStream.h +++ b/benchmarks/babelstream/src/AlpakaStream.h @@ -42,8 +42,8 @@ struct AlpakaStream : Stream using DevHost = alpaka::Dev; using PlatformAcc = alpaka::Platform; using DevAcc = alpaka::Dev; - using BufHost = alpaka::Buf; - using BufAcc = alpaka::Buf; + using BufHost = alpaka::Buf>; + using BufAcc = alpaka::Buf>; using Queue = alpaka::Queue; using WorkDiv = alpaka::WorkDivMembers; From 6ae951d2dc0967ec8276a7e713356b6cdb1f35d6 Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Tue, 21 May 2024 11:12:38 +0200 Subject: [PATCH 17/19] fix 2 --- include/alpaka/mem/buf/BufUniformCudaHipRt.hpp | 1 + 1 file changed, 1 insertion(+) diff --git a/include/alpaka/mem/buf/BufUniformCudaHipRt.hpp b/include/alpaka/mem/buf/BufUniformCudaHipRt.hpp index f52c258405fa..919fe18847da 100644 --- a/include/alpaka/mem/buf/BufUniformCudaHipRt.hpp +++ b/include/alpaka/mem/buf/BufUniformCudaHipRt.hpp @@ -14,6 +14,7 @@ #include "alpaka/mem/buf/Traits.hpp" #include "alpaka/mem/view/ViewAccessOps.hpp" #include "alpaka/meta/DependentFalseType.hpp" +#include "alpaka/meta/Unique.hpp" #include "alpaka/vec/Vec.hpp" #include From a5f73362f533b2250a92606728baf83a6e823d9f Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Tue, 21 May 2024 13:26:45 +0200 Subject: [PATCH 18/19] fix 3 --- include/alpaka/mem/buf/BufGenericSycl.hpp | 6 ++++-- include/alpaka/mem/buf/BufUniformCudaHipRt.hpp | 11 +++++------ 2 files changed, 9 insertions(+), 8 deletions(-) diff --git a/include/alpaka/mem/buf/BufGenericSycl.hpp b/include/alpaka/mem/buf/BufGenericSycl.hpp index ae64487dac35..7f9d63c05cb7 100644 --- a/include/alpaka/mem/buf/BufGenericSycl.hpp +++ b/include/alpaka/mem/buf/BufGenericSycl.hpp @@ -9,9 +9,11 @@ #include "alpaka/dev/Traits.hpp" #include "alpaka/dim/DimIntegralConst.hpp" #include "alpaka/dim/Traits.hpp" +#include "alpaka/mem/Visibility.hpp" #include "alpaka/mem/buf/BufCpu.hpp" #include "alpaka/mem/buf/Traits.hpp" #include "alpaka/mem/view/ViewAccessOps.hpp" +#include "alpaka/meta/Unique.hpp" #include "alpaka/vec/Vec.hpp" #include @@ -232,7 +234,7 @@ namespace alpaka::trait TElem, TDim, TIdx, - alpaka::meta::Unique, alpaka::MemVisibility>>> + alpaka::meta::Unique>>> { ALPAKA_DEBUG_MINIMAL_LOG_SCOPE; @@ -246,7 +248,7 @@ namespace alpaka::trait TElem, TDim, TIdx, - alpaka::meta::Unique, alpaka::MemVisibility>>>( + alpaka::meta::Unique>>>( host, memPtr, std::move(deleter), diff --git a/include/alpaka/mem/buf/BufUniformCudaHipRt.hpp b/include/alpaka/mem/buf/BufUniformCudaHipRt.hpp index 919fe18847da..d15e80763f47 100644 --- a/include/alpaka/mem/buf/BufUniformCudaHipRt.hpp +++ b/include/alpaka/mem/buf/BufUniformCudaHipRt.hpp @@ -11,6 +11,7 @@ #include "alpaka/dev/DevUniformCudaHipRt.hpp" #include "alpaka/dev/Traits.hpp" #include "alpaka/dim/DimIntegralConst.hpp" +#include "alpaka/mem/Visibility.hpp" #include "alpaka/mem/buf/Traits.hpp" #include "alpaka/mem/view/ViewAccessOps.hpp" #include "alpaka/meta/DependentFalseType.hpp" @@ -360,9 +361,8 @@ namespace alpaka TElem, TDim, TIdx, - alpaka::meta::Unique, - alpaka::MemVisibility>>>> + alpaka::meta::Unique< + std::tuple>>>> { ALPAKA_DEBUG_MINIMAL_LOG_SCOPE; @@ -379,9 +379,8 @@ namespace alpaka TElem, TDim, TIdx, - alpaka::meta::Unique, - alpaka::MemVisibility>>>>( + alpaka::meta::Unique< + std::tuple>>>>( host, memPtr, std::move(deleter), From c3eebea23127f4f96992537329cb983e0daceda1 Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Mon, 27 May 2024 16:57:40 +0200 Subject: [PATCH 19/19] add buffer visibility tests --- include/alpaka/acc/Tag.hpp | 1 + include/alpaka/mem/Visibility.hpp | 3 +- test/unit/mem/src/Visibility.cpp | 67 +++++++++++++++++++++++++++++++ 3 files changed, 70 insertions(+), 1 deletion(-) diff --git a/include/alpaka/acc/Tag.hpp b/include/alpaka/acc/Tag.hpp index f7880afd6f15..73575daa0a4d 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); +#undef CREATE_ACC_TAG namespace trait { diff --git a/include/alpaka/mem/Visibility.hpp b/include/alpaka/mem/Visibility.hpp index f5cf7c0d024a..349beb719e20 100644 --- a/include/alpaka/mem/Visibility.hpp +++ b/include/alpaka/mem/Visibility.hpp @@ -34,6 +34,7 @@ namespace alpaka CREATE_MEM_VISIBILITY(MemVisibleGpuCudaRt); CREATE_MEM_VISIBILITY(MemVisibleGpuHipRt); CREATE_MEM_VISIBILITY(MemVisibleGpuSyclIntel); +#undef CREATE_MEM_VISIBILITY namespace trait { @@ -62,7 +63,7 @@ namespace alpaka typename = std::enable_if_t< alpaka::isPlatform> || alpaka::isDevice> || alpaka::isAccelerator> || alpaka::internal::isView>>> - [[maybe_unused]] static std::string getMemVisiblityName() + inline std::string getMemVisiblityName() { using MemVisibilityType = typename alpaka::trait::MemVisibility>::type; if constexpr(alpaka::meta::isList) diff --git a/test/unit/mem/src/Visibility.cpp b/test/unit/mem/src/Visibility.cpp index 2a46ab35577a..097f7ec60b88 100644 --- a/test/unit/mem/src/Visibility.cpp +++ b/test/unit/mem/src/Visibility.cpp @@ -174,3 +174,70 @@ TEMPLATE_LIST_TEST_CASE("testMemView", "[mem][visibility]", EnabledTagTagMemVisi STATIC_REQUIRE_FALSE(alpaka::hasSameMemView(plt2, data_view1)); } } + +TEMPLATE_LIST_TEST_CASE("testMemBuf", "[mem][visibility]", EnabledTagTagMemVisibilityList) +{ + using Tag1 = std::tuple_element_t<0, std::tuple_element_t<0, TestType>>; + using ExpectedMemVisibilityForTag1 = std::tuple_element_t<1, std::tuple_element_t<0, TestType>>; + using Tag2 = std::tuple_element_t<0, std::tuple_element_t<1, TestType>>; + using ExpectedMemVisibilityForTag2 = std::tuple_element_t<1, std::tuple_element_t<1, TestType>>; + + SUCCEED( + "Tag1: " << Tag1::get_name() << " + " << ExpectedMemVisibilityForTag1::get_name() + << "\nTag2: " << Tag2::get_name() << " + " << ExpectedMemVisibilityForTag1::get_name()); + + constexpr Idx data_size = 10; + + using Acc1 = alpaka::TagToAcc; + using Acc2 = alpaka::TagToAcc; + + auto const plt1 = alpaka::Platform{}; + auto const plt2 = alpaka::Platform{}; + + auto const dev1 = alpaka::getDevByIdx(plt1, 0); + auto const dev2 = alpaka::getDevByIdx(plt2, 0); + + using Vec1D = alpaka::Vec, Idx>; + Vec1D const extents(Vec1D::all(data_size)); + + // we need only to test the first tag, because the second tag contains the same acc's + auto buf1 = alpaka::allocBuf(dev1, extents); + + STATIC_REQUIRE(std::is_same_v< + typename alpaka::trait::MemVisibility::type, + std::tuple>); + + STATIC_REQUIRE(alpaka::hasSameMemView(plt1, buf1)); + + if constexpr(!std::is_same_v) + { + alpaka::Queue queue1(dev1); + + auto buf1Async = alpaka::allocAsyncBuf(queue1, extents); + + alpaka::wait(queue1); + + STATIC_REQUIRE(std::is_same_v< + typename alpaka::trait::MemVisibility::type, + std::tuple>); + + STATIC_REQUIRE(alpaka::hasSameMemView(plt1, buf1Async)); + } + + if constexpr(isCpuTag::value && alpaka::hasMappedBufSupport>) + { + auto mappedBuffer = alpaka::allocMappedBuf, float, Idx>(dev1, plt2, extents); + + using expectedMappedBufferMemView + = alpaka::meta::Unique>; + + STATIC_REQUIRE(std::is_same_v< + typename alpaka::trait::MemVisibility::type, + expectedMappedBufferMemView>); + + STATIC_REQUIRE(alpaka::hasSameMemView(plt1, mappedBuffer)); + STATIC_REQUIRE(alpaka::hasSameMemView(plt2, mappedBuffer)); + STATIC_REQUIRE(alpaka::hasSameMemView(dev1, mappedBuffer)); + STATIC_REQUIRE(alpaka::hasSameMemView(dev2, mappedBuffer)); + } +}