[cudauvm] Update code base to CMSSW_11_2_0_pre8_Patatrack

Makefile

@@ -195,8 +195,6 @@ endif
 
 # Targets and their dependencies on externals
 TARGETS_ALL := $(notdir $(wildcard $(SRC_DIR)/*))
-# Temporarily filter out programs that do not run (yet) with CUDA 11
-TARGETS_ALL := $(filter-out cudauvm,$(TARGETS_ALL))
 
 # Split targets by required toolchain
 TARGETS_GCC := fwtest

README.md

@@ -78,7 +78,7 @@ downloaded automatically during the build process.
 | `cudatest`   | CUDA FW test                     | :heavy_check_mark: | :heavy_check_mark: | :heavy_check_mark: |                    |                    |                    |                    |                    |                    |                    |
 | `cuda`       | CUDA version (frozen)            | :heavy_check_mark: | :heavy_check_mark: |                    | :heavy_check_mark: | :heavy_check_mark: | :heavy_check_mark: | :heavy_check_mark: | :heavy_check_mark: | :heavy_check_mark: | :heavy_check_mark: |
 | `cudadev`    | CUDA version (development)       | :heavy_check_mark: | :heavy_check_mark: |                    | :heavy_check_mark: | :heavy_check_mark: | :heavy_check_mark: | :heavy_check_mark: | :heavy_check_mark: | :heavy_check_mark: | :heavy_check_mark: |
-| `cudauvm`    | CUDA version with managed memory | :heavy_check_mark: | :heavy_check_mark: |                    | :white_check_mark: | :white_check_mark: | :white_check_mark: | :white_check_mark: | :white_check_mark: |                    |                    |
+| `cudauvm`    | CUDA version with managed memory | :heavy_check_mark: | :heavy_check_mark: |                    | :heavy_check_mark: | :white_check_mark: | :white_check_mark: | :white_check_mark: | :white_check_mark: | :heavy_check_mark: | :heavy_check_mark: |
 | `kokkostest` | Kokkos FW test                   | :heavy_check_mark: | :white_check_mark: | :heavy_check_mark: |                    |                    |                    |                    |                    |                    |                    |
 | `kokkos`     | Kokkos version                   | :heavy_check_mark: |                    | :heavy_check_mark: | :heavy_check_mark: | :heavy_check_mark: | :heavy_check_mark: | :heavy_check_mark: | :heavy_check_mark: | :heavy_check_mark: |                    |
 | `alpakatest` | Alpaka FW test                   | :heavy_check_mark: |                    | :white_check_mark: |                    |                    |                    |                    |                    |                    |                    |
@@ -155,7 +155,7 @@ This program is frozen to correspond to CMSSW_11_2_0_pre8_Patatrack.
 
 #### `cudadev`
 
-This program currently contains code from CMSSW_11_2_0_pre8_Patatrack (currently equivalent to `cuda`).
+This program is currently equivalent to `cuda`.
 
 #### `cudauvm`
 

src/cudauvm/CUDACore/AtomicPairCounter.h

@@ -5,48 +5,54 @@
 
 #include "CUDACore/cudaCompat.h"
 
-class AtomicPairCounter {
-public:
-  using c_type = unsigned long long int;
+namespace cms {
+  namespace cuda {
 
-  AtomicPairCounter() {}
-  AtomicPairCounter(c_type i) { counter.ac = i; }
+    class AtomicPairCounter {
+    public:
+      using c_type = unsigned long long int;
 
-  __device__ __host__ AtomicPairCounter& operator=(c_type i) {
-    counter.ac = i;
-    return *this;
-  }
+      AtomicPairCounter() {}
+      AtomicPairCounter(c_type i) { counter.ac = i; }
 
-  struct Counters {
-    uint32_t n;  // in a "One to Many" association is the number of "One"
-    uint32_t m;  // in a "One to Many" association is the total number of associations
-  };
+      __device__ __host__ AtomicPairCounter& operator=(c_type i) {
+        counter.ac = i;
+        return *this;
+      }
 
-  union Atomic2 {
-    Counters counters;
-    c_type ac;
-  };
+      struct Counters {
+        uint32_t n;  // in a "One to Many" association is the number of "One"
+        uint32_t m;  // in a "One to Many" association is the total number of associations
+      };
 
-  static constexpr c_type incr = 1UL << 32;
+      union Atomic2 {
+        Counters counters;
+        c_type ac;
+      };
 
-  __device__ __host__ Counters get() const { return counter.counters; }
+      static constexpr c_type incr = 1UL << 32;
 
-  // increment n by 1 and m by i.  return previous value
-  __host__ __device__ __forceinline__ Counters add(uint32_t i) {
-    c_type c = i;
-    c += incr;
-    Atomic2 ret;
+      __device__ __host__ Counters get() const { return counter.counters; }
+
+      // increment n by 1 and m by i.  return previous value
+      __host__ __device__ __forceinline__ Counters add(uint32_t i) {
+        c_type c = i;
+        c += incr;
+        Atomic2 ret;
 #ifdef __CUDA_ARCH__
-    ret.ac = atomicAdd(&counter.ac, c);
+        ret.ac = atomicAdd(&counter.ac, c);
 #else
-    ret.ac = counter.ac;
-    counter.ac += c;
+        ret.ac = counter.ac;
+        counter.ac += c;
 #endif
-    return ret.counters;
-  }
+        return ret.counters;
+      }
+
+    private:
+      Atomic2 counter;
+    };
 
-private:
-  Atomic2 counter;
-};
+  }  // namespace cuda
+}  // namespace cms
 
 #endif  // HeterogeneousCore_CUDAUtilities_interface_AtomicPairCounter_h