Merge pull request #40678 from fwyzard/ROCm_developments

Improve support for HIP/ROCm and Alpaka
cms-sw · Feb 5, 2023 · 06f284e · 06f284e
2 parents 3e10dc1 + a14e8fc
commit 06f284e
Show file tree

Hide file tree

Showing 28 changed files with 406 additions and 154 deletions.
diff --git a/DataFormats/Portable/interface/PortableCollection.h b/DataFormats/Portable/interface/PortableCollection.h
@@ -6,13 +6,13 @@
 namespace traits {
 
   // trait for a generic SoA-based product
-  template <typename T, typename TDev, typename = std::enable_if_t<cms::alpakatools::is_device_v<TDev>>>
+  template <typename T, typename TDev, typename = std::enable_if_t<alpaka::isDevice<TDev>>>
   class PortableCollectionTrait;
 
 }  // namespace traits
 
 // type alias for a generic SoA-based product
-template <typename T, typename TDev, typename = std::enable_if_t<cms::alpakatools::is_device_v<TDev>>>
+template <typename T, typename TDev, typename = std::enable_if_t<alpaka::isDevice<TDev>>>
 using PortableCollection = typename traits::PortableCollectionTrait<T, TDev>::CollectionType;
 
 #endif  // DataFormats_Portable_interface_PortableCollection_h
diff --git a/DataFormats/Portable/interface/PortableDeviceCollection.h b/DataFormats/Portable/interface/PortableDeviceCollection.h
@@ -5,12 +5,13 @@
 #include <optional>
 #include <type_traits>
 
+#include <alpaka/alpaka.hpp>
+
 #include "HeterogeneousCore/AlpakaInterface/interface/config.h"
 #include "HeterogeneousCore/AlpakaInterface/interface/memory.h"
-#include "HeterogeneousCore/AlpakaInterface/interface/traits.h"
 
 // generic SoA-based product in device memory
-template <typename T, typename TDev, typename = std::enable_if_t<cms::alpakatools::is_device_v<TDev>>>
+template <typename T, typename TDev, typename = std::enable_if_t<alpaka::isDevice<TDev>>>
 class PortableDeviceCollection {
   static_assert(not std::is_same_v<TDev, alpaka_common::DevHost>,
                 "Use PortableHostCollection<T> instead of PortableDeviceCollection<T, DevHost>");
@@ -32,7 +33,7 @@ class PortableDeviceCollection {
     assert(reinterpret_cast<uintptr_t>(buffer_->data()) % Layout::alignment == 0);
   }
 
-  template <typename TQueue, typename = std::enable_if_t<cms::alpakatools::is_queue_v<TQueue>>>
+  template <typename TQueue, typename = std::enable_if_t<alpaka::isQueue<TQueue>>>
   PortableDeviceCollection(int32_t elements, TQueue const& queue)
       : buffer_{cms::alpakatools::make_device_buffer<std::byte[]>(queue, Layout::computeDataSize(elements))},
         layout_{buffer_->data(), elements},

diff --git a/DataFormats/Portable/interface/PortableHostCollection.h b/DataFormats/Portable/interface/PortableHostCollection.h
@@ -4,10 +4,11 @@
 #include <cassert>
 #include <optional>
 
+#include <alpaka/alpaka.hpp>
+
 #include "HeterogeneousCore/AlpakaInterface/interface/config.h"
 #include "HeterogeneousCore/AlpakaInterface/interface/host.h"
 #include "HeterogeneousCore/AlpakaInterface/interface/memory.h"
-#include "HeterogeneousCore/AlpakaInterface/interface/traits.h"
 
 // generic SoA-based product in host memory
 template <typename T>
@@ -30,7 +31,7 @@ class PortableHostCollection {
     assert(reinterpret_cast<uintptr_t>(buffer_->data()) % Layout::alignment == 0);
   }
 
-  template <typename TQueue, typename = std::enable_if_t<cms::alpakatools::is_queue_v<TQueue>>>
+  template <typename TQueue, typename = std::enable_if_t<alpaka::isQueue<TQueue>>>
   PortableHostCollection(int32_t elements, TQueue const& queue)
       // allocate pinned host memory associated to the given work queue, accessible by the queue's device
       : buffer_{cms::alpakatools::make_host_buffer<std::byte[]>(queue, Layout::computeDataSize(elements))},

diff --git a/DataFormats/Portable/interface/Product.h b/DataFormats/Portable/interface/Product.h
@@ -4,10 +4,11 @@
 #include <memory>
 #include <utility>
 
+#include <alpaka/alpaka.hpp>
+
 #include "DataFormats/Portable/interface/ProductBase.h"
 #include "HeterogeneousCore/AlpakaInterface/interface/ScopedContextFwd.h"
 #include "HeterogeneousCore/AlpakaInterface/interface/config.h"
-#include "HeterogeneousCore/AlpakaInterface/interface/traits.h"
 
 namespace edm {
   template <typename T>
@@ -30,7 +31,7 @@ namespace cms::alpakatools {
      * it. Here is a somewhat natural place. If the overhead is too much, we
      * can use them only where synchronization between queues is needed.
      */
-  template <typename TQueue, typename T, typename = std::enable_if_t<cms::alpakatools::is_queue_v<TQueue>>>
+  template <typename TQueue, typename T, typename = std::enable_if_t<alpaka::isQueue<TQueue>>>
   class Product : public ProductBase<TQueue> {
   public:
     using Queue = TQueue;

diff --git a/DataFormats/Portable/interface/ProductBase.h b/DataFormats/Portable/interface/ProductBase.h
@@ -8,15 +8,14 @@
 #include <alpaka/alpaka.hpp>
 
 #include "HeterogeneousCore/AlpakaInterface/interface/ScopedContextFwd.h"
-#include "HeterogeneousCore/AlpakaInterface/interface/traits.h"
 
 namespace cms::alpakatools {
 
   /**
    * Base class for all instantiations of Product<TQueue, T> to hold the
    * non-T-dependent members.
    */
-  template <typename TQueue, typename = std::enable_if_t<cms::alpakatools::is_queue_v<TQueue>>>
+  template <typename TQueue, typename = std::enable_if_t<alpaka::isQueue<TQueue>>>
   class ProductBase {
   public:
     using Queue = TQueue;

diff --git a/DataFormats/SoATemplate/interface/SoACommon.h b/DataFormats/SoATemplate/interface/SoACommon.h
@@ -16,7 +16,7 @@
 #include "FWCore/Utilities/interface/typedefs.h"
 
 // CUDA attributes
-#ifdef __CUDACC__
+#if defined(__CUDACC__) || defined(__HIPCC__)
 #define SOA_HOST_ONLY __host__
 #define SOA_DEVICE_ONLY __device__
 #define SOA_HOST_DEVICE __host__ __device__
@@ -35,6 +35,12 @@
     printf("%s\n", (A));          \
     __trap();                     \
   }
+#elif defined(__HIPCC__) && defined(__HIP_DEVICE_COMPILE__)
+#define SOA_THROW_OUT_OF_RANGE(A) \
+  {                               \
+    printf("%s\n", (A));          \
+    abort();                      \
+  }
 #else
 #define SOA_THROW_OUT_OF_RANGE(A) \
   { throw std::out_of_range(A); }

diff --git a/DataFormats/SoATemplate/test/BuildFile.xml b/DataFormats/SoATemplate/test/BuildFile.xml
@@ -1,17 +1,26 @@
 <iftool name="cuda-gcc-support">
-  <bin file="SoALayoutAndView_t.cu" name="SoALayoutAndView_t">
+  <bin file="SoALayoutAndView_t.cu" name="testCudaSoALayoutAndView_t">
     <use name="boost"/>
     <use name="cuda"/>
     <use name="eigen"/>
     <use name="HeterogeneousCore/CUDAUtilities"/>
   </bin>
+</iftool>
 
-  <!-- dictionaries for FakeSoA -->
-  <library file="" name="FakeSoADict"/>
-
-  <!-- This test triggers a bug in ROOT, so it's kept disabled
-  <bin file="SoAStreamer_t.cpp" name="SoAStreamer_t">
-    <use name="root"/>
+<iftool name="rocm">
+  <bin file="SoALayoutAndView_t.hip.cc" name="testRocmSoALayoutAndView_t">
+    <use name="boost"/>
+    <use name="eigen"/>
+    <use name="rocm"/>
+    <use name="HeterogeneousCore/ROCmUtilities"/>
   </bin>
-  -->
 </iftool>
+
+<!-- This test triggers a bug in ROOT, so it's kept disabled
+<bin file="SoAStreamer_t.cpp" name="SoAStreamer_t">
+  <use name="root"/>
+</bin>
+-->
+
+<!-- dictionaries for FakeSoA -->
+<library file="" name="FakeSoADict"/>