Various fixes for heterogeneous utilities

DataFormats/SoATemplate/interface/SoACommon.h

@@ -13,9 +13,17 @@
 
 #include <boost/preprocessor.hpp>
 
+#ifdef __CUDACC__
+#include <cuda_runtime.h>
+#endif
+
+#ifdef __HIPCC__
+#include <hip/hip_runtime_api.h>
+#endif
+
 #include "FWCore/Utilities/interface/typedefs.h"
 
-// CUDA attributes
+// CUDA/ROCm attributes
 #if defined(__CUDACC__) || defined(__HIPCC__)
 #define SOA_HOST_ONLY __host__
 #define SOA_DEVICE_ONLY __device__

DataFormats/SoATemplate/interface/SoALayout.h

@@ -7,6 +7,7 @@
  */
 
 #include <cassert>
+#include <cstring>
 
 #include "FWCore/Reflection/interface/reflex.h"
 

HeterogeneousCore/AlpakaCore/README.md

@@ -23,7 +23,7 @@ The `BuildFile.xml` must contain `<flags ALPAKA_BACKENDS="1"/>` to enable the be
   * If you need to transfer some data back to host, use `stream::SynchronizingEDProducer`
 * All code using `ALPAKA_ACCELERATOR_NAMESPACE` should be placed in `Package/SubPackage/{interface,src,plugins,test}/alpaka` directory
   * Alpaka-dependent code that uses templates instead of the namespace macro can be placed in `Package/SubPackage/interface` directory
-* All source files (not headers) using Alpaka device code (such as kernel call, functions called by kernels) must have a suffic `.dev.cc`, and be placed in the aforementioned `alpaka` subdirectory
+* All source files (not headers) using Alpaka device code (such as kernel call, functions called by kernels) must have a suffix `.dev.cc`, and be placed in the aforementioned `alpaka` subdirectory
 * Any code that `#include`s a header from the framework or from the `HeterogeneousCore/AlpakaCore` must be separated from the Alpaka device code, and have the usual `.cc` suffix.
   * Some framework headers are allowed to be used in `.dev.cc` files:
     * Any header containing only macros, e.g. `FWCore/Utilities/interface/CMSUnrollLoop.h`, `FWCore/Utilities/interface/stringize.h`

HeterogeneousCore/AlpakaInterface/interface/atomicInc.h

@@ -0,0 +1,15 @@
+#ifndef HeterogeneousCore_AlpakaInterface_interface_atomicInc_h
+#define HeterogeneousCore_AlpakaInterface_interface_atomicInc_h
+
+#include <alpaka/alpaka.hpp>
+
+// This function is similar to atomicInc, but deduces the limiting value from the type itself.
+
+ALPAKA_NO_HOST_ACC_WARNING
+template <typename TAcc, typename T, typename THierarchy = alpaka::hierarchy::Grids>
+ALPAKA_FN_HOST_ACC auto atomicInc(TAcc const& acc, T* address, THierarchy const& hierarchy = THierarchy()) -> T {
+  T limit = std::numeric_limits<T>::max();
+  return alpaka::atomicInc(acc, address, limit, hierarchy);
+}
+
+#endif  // HeterogeneousCore_AlpakaInterface_interface_atomicInc_h

HeterogeneousCore/AlpakaInterface/interface/atomicIncSaturate.h

@@ -0,0 +1,29 @@
+#ifndef HeterogeneousCore_AlpakaInterface_interface_atomicIncSaturate_h
+#define HeterogeneousCore_AlpakaInterface_interface_atomicIncSaturate_h
+
+#include <alpaka/alpaka.hpp>
+
+// This function is similar to atomicInc, but instead of wrapping around it saturates at the given value.
+
+ALPAKA_NO_HOST_ACC_WARNING
+template <typename TAcc, typename T, typename THierarchy = alpaka::hierarchy::Grids>
+ALPAKA_FN_HOST_ACC auto atomicIncSaturate(TAcc const& acc,
+                                          T* address,
+                                          T const& limit,
+                                          THierarchy const& hierarchy = THierarchy()) -> T {
+  T assumed;
+  T old = *address;
+
+  do {
+    assumed = old;
+    if (assumed >= limit) {
+      // Saturate at limit.
+      break;
+    }
+    old = alpaka::atomicCas(acc, address, assumed, assumed + 1, hierarchy);
+  } while (old != assumed);
+
+  return old;
+}
+
+#endif  // HeterogeneousCore_AlpakaInterface_interface_atomicIncSaturate_h