more temporary memory in tmop

fem/bilininteg_convection_pa.cpp

@@ -900,11 +900,18 @@ void ConvectionIntegrator::AssemblePA(const FiniteElementSpace &fes)
    dim = mesh->Dimension();
    ne = fes.GetNE();
    const DofToQuad::Mode mode = DofToQuad::TENSOR;
-   geom = mesh->GetGeometricFactors(*ir, GeometricFactors::JACOBIANS, mode);
+#ifdef MFEM_USE_UMPIRE
+   const MemoryType temp_type = Device::GetDeviceMemoryType() == MemoryType::DEVICE_UMPIRE
+      ? MemoryType::DEVICE_UMPIRE_2 : Device::GetDeviceMemoryType();
+#else
+   const MemoryType temp_type = Device::GetDeviceMemoryType();
+#endif
+
+   geom = mesh->GetGeometricFactors(*ir, GeometricFactors::JACOBIANS, mode, temp_type);
    maps = &el.GetDofToQuad(*ir, mode);
    dofs1D = maps->ndof;
    quad1D = maps->nqpt;
-   pa_data.SetSize(symmDims * nq * ne, Device::GetMemoryType());
+   pa_data.SetSize(symmDims * nq * ne, temp_type);
    Vector vel;
    if (VectorConstantCoefficient *cQ =
           dynamic_cast<VectorConstantCoefficient*>(Q))
@@ -915,7 +922,7 @@ void ConvectionIntegrator::AssemblePA(const FiniteElementSpace &fes)
                dynamic_cast<VectorGridFunctionCoefficient*>(Q))
    {
       Vector xe;
-      vel.SetSize(dim * nq * ne);
+      vel.SetSize(dim * nq * ne, temp_type);
 
       const GridFunction *gf = vgfQ->GetGridFunction();
       const ElementDofOrdering ordering = ElementDofOrdering::LEXICOGRAPHIC;

fem/bilininteg_mass_pa.cpp

@@ -49,7 +49,13 @@ void MassIntegrator::SetupPA(const FiniteElementSpace &fes)
    const DofToQuad::Mode mode = DofToQuad::TENSOR;
    const int flags = GeometricFactors::JACOBIANS |
                      GeometricFactors::COORDINATES;
-   geom = mesh->GetGeometricFactors(*ir, flags, mode);
+#ifdef MFEM_USE_UMPIRE
+   const MemoryType temp_type = Device::GetDeviceMemoryType() == MemoryType::DEVICE_UMPIRE
+      ? MemoryType::DEVICE_UMPIRE_2 : Device::GetDeviceMemoryType();
+#else
+   const MemoryType temp_type = Device::GetDeviceMemoryType();
+#endif
+   geom = mesh->GetGeometricFactors(*ir, flags, mode, temp_type);
    maps = &el.GetDofToQuad(*ir, mode);
    dofs1D = maps->ndof;
    quad1D = maps->nqpt;

fem/pfespace.cpp

@@ -3147,7 +3147,7 @@ static void SetSubVector(const int N,
                          const Array<int> &indices,
                          const Vector &in, Vector &out)
 {
-   auto y = out.Write();
+   auto y = out.ReadWrite();
    const auto x = in.Read();
    const auto I = indices.Read();
    MFEM_FORALL(i, N, y[I[i]] = x[i];);
@@ -3234,7 +3234,7 @@ static void AddSubVector(const int num_unique_dst_indices,
                          const Vector &src,
                          Vector &dst)
 {
-   auto y = dst.Write();
+   auto y = dst.ReadWrite();
    const auto x = src.Read();
    const auto DST_I = unique_dst_indices.Read();
    const auto SRC_O = unique_to_src_offsets.Read();

fem/tmop.cpp

@@ -1872,6 +1872,17 @@ void AdaptivityEvaluator::SetParMetaInfo(const ParMesh &m,
 }
 #endif
 
+void AdaptivityEvaluator::ClearGeometricFactors()
+{
+#ifdef MFEM_USE_MPI
+   if (pmesh) pmesh->DeleteGeometricFactors();
+   if (pfes) pfes->GetParMesh()->DeleteGeometricFactors();
+#else
+   if (mesh) mesh->DeleteGeometricFactors();
+   if (fes) fes->GetMesh()->DeleteGeometricFactors();
+#endif
+}
+
 AdaptivityEvaluator::~AdaptivityEvaluator()
 {
    delete fes;

fem/tmop.hpp

@@ -606,6 +606,8 @@ public:
 
    virtual void ComputeAtNewPosition(const Vector &new_nodes,
                                      Vector &new_field) = 0;
+
+   void ClearGeometricFactors();
 };
 
 /** @brief Base class representing target-matrix construction algorithms for
@@ -1216,6 +1218,8 @@ public:
       {
          PA.H.GetMemory().DeleteDevice();
          PA.H0.GetMemory().DeleteDevice();
+         //PA.Jtr.GetMemory().DeleteDevice();
+         //PA.setup_Jtr = false;
       }
    }
 };

fem/tmop_pa.cpp

@@ -207,12 +207,18 @@ void TMOP_Integrator::AssemblePA(const FiniteElementSpace &fes)
    PA.setup_Jtr = false;
    PA.setup_Grad = false;
 
+
+#ifdef MFEM_USE_UMPIRE
+   const MemoryType temp_type = Device::GetDeviceMemoryType() == MemoryType::DEVICE_UMPIRE
+      ? MemoryType::DEVICE_UMPIRE_2 : Device::GetDeviceMemoryType();
+#else
+   const MemoryType temp_type = Device::GetDeviceMemoryType();
+#endif
+
    // H for Grad
-   PA.H.SetSize(dim*dim * dim*dim * nq*ne);
-   PA.H.GetMemory().UseTemporary(true);
+   PA.H.SetSize(dim*dim * dim*dim * nq*ne, temp_type);
    // H0 for coeff0
-   PA.H0.SetSize(dim * dim * nq*ne);
-   PA.H0.GetMemory().UseTemporary(true);
+   PA.H0.SetSize(dim * dim * nq*ne, temp_type);
 
    // Restriction setup
    const ElementDofOrdering ordering = ElementDofOrdering::LEXICOGRAPHIC;
@@ -232,7 +238,7 @@ void TMOP_Integrator::AssemblePA(const FiniteElementSpace &fes)
    PA.O = 1.0;
 
    // TargetConstructor TargetType setup
-   PA.Jtr.SetSize(dim, dim, PA.ne*PA.nq);
+   PA.Jtr.SetSize(dim, dim, PA.ne*PA.nq);//, temp_type);
    ComputeElementTargetsPA();
 
    // Coeff0 PA.C0

general/cuda.cpp

@@ -1,4 +1,4 @@
-// Copyright (c) 2010-2020, Lawrence Livermore National Security, LLC. Produced
+// Copyright (c) 2010-2021, Lawrence Livermore National Security, LLC. Produced
 // at the Lawrence Livermore National Laboratory. All Rights reserved. See files
 // LICENSE and NOTICE for details. LLNL-CODE-806117.
 //
@@ -61,11 +61,11 @@ void* CuMallocManaged(void** dptr, size_t bytes)
    return *dptr;
 }
 
-void* CuMemAllocHost(void** ptr, size_t bytes)
+void* CuMemAllocHostPinned(void** ptr, size_t bytes)
 {
 #ifdef MFEM_USE_CUDA
 #ifdef MFEM_TRACK_CUDA_MEM
-   mfem::out << "CuMemAllocHost(): allocating " << bytes << " bytes ... "
+   mfem::out << "CuMemAllocHostPinned(): allocating " << bytes << " bytes ... "
              << std::flush;
 #endif
    MFEM_GPU_CHECK(cudaMallocHost(ptr, bytes));
@@ -91,11 +91,11 @@ void* CuMemFree(void *dptr)
    return dptr;
 }
 
-void* CuMemFreeHost(void *ptr)
+void* CuMemFreeHostPinned(void *ptr)
 {
 #ifdef MFEM_USE_CUDA
 #ifdef MFEM_TRACK_CUDA_MEM
-   mfem::out << "CuMemFreeHost(): deallocating memory @ " << ptr << " ... "
+   mfem::out << "CuMemFreeHostPinned(): deallocating memory @ " << ptr << " ... "
              << std::flush;
 #endif
    MFEM_GPU_CHECK(cudaFreeHost(ptr));

general/cuda.hpp

@@ -1,4 +1,4 @@
-// Copyright (c) 2010-2020, Lawrence Livermore National Security, LLC. Produced
+// Copyright (c) 2010-2021, Lawrence Livermore National Security, LLC. Produced
 // at the Lawrence Livermore National Laboratory. All Rights reserved. See files
 // LICENSE and NOTICE for details. LLNL-CODE-806117.
 //
@@ -43,7 +43,6 @@
 #if defined(MFEM_USE_CUDA) && defined(__CUDA_ARCH__)
 #define MFEM_SHARED __shared__
 #define MFEM_SYNC_THREAD __syncthreads()
-#define MFEM_BLOCK_ID(k) blockIdx.k
 #define MFEM_THREAD_ID(k) threadIdx.k
 #define MFEM_THREAD_SIZE(k) blockDim.k
 #define MFEM_FOREACH_THREAD(i,k,N) for(int i=threadIdx.k; i<N; i+=blockDim.k)
@@ -64,14 +63,14 @@ void* CuMemAlloc(void **d_ptr, size_t bytes);
 /// Allocates managed device memory
 void* CuMallocManaged(void **d_ptr, size_t bytes);
 
-/// Allocate page-locked (pinned) host memory
-void* CuMemAllocHost(void **ptr, size_t bytes);
+/// Allocates page-locked (pinned) host memory
+void* CuMemAllocHostPinned(void **ptr, size_t bytes);
 
 /// Frees device memory and returns destination ptr.
 void* CuMemFree(void *d_ptr);
 
 /// Frees page-locked (pinned) host memory and returns destination ptr.
-void* CuMemFreeHost(void *ptr);
+void* CuMemFreeHostPinned(void *ptr);
 
 /// Copies memory from Host to Device and returns destination ptr.
 void* CuMemcpyHtoD(void *d_dst, const void *h_src, size_t bytes);

general/device.cpp

@@ -1,4 +1,4 @@
-// Copyright (c) 2010-2020, Lawrence Livermore National Security, LLC. Produced
+// Copyright (c) 2010-2021, Lawrence Livermore National Security, LLC. Produced
 // at the Lawrence Livermore National Laboratory. All Rights reserved. See files
 // LICENSE and NOTICE for details. LLNL-CODE-806117.
 //
@@ -12,9 +12,10 @@
 #include "forall.hpp"
 #include "occa.hpp"
 #ifdef MFEM_USE_CEED
-#include <ceed.h>
+#include "../fem/ceed/util.hpp"
 #endif
 
+#include <unordered_map>
 #include <string>
 #include <map>
 
@@ -33,13 +34,16 @@ occa::device occaDevice;
 
 #ifdef MFEM_USE_CEED
 Ceed ceed = NULL;
+
+ceed::BasisMap ceed_basis_map;
+ceed::RestrMap ceed_restr_map;
 #endif
 
 // Backends listed by priority, high to low:
 static const Backend::Id backend_list[Backend::NUM_BACKENDS] =
 {
    Backend::CEED_CUDA, Backend::OCCA_CUDA, Backend::RAJA_CUDA, Backend::CUDA,
-   Backend::HIP, Backend::DEBUG,
+   Backend::CEED_HIP, Backend::RAJA_HIP, Backend::HIP, Backend::DEBUG,
    Backend::OCCA_OMP, Backend::RAJA_OMP, Backend::OMP,
    Backend::CEED_CPU, Backend::OCCA_CPU, Backend::RAJA_CPU, Backend::CPU
 };
@@ -48,7 +52,7 @@ static const Backend::Id backend_list[Backend::NUM_BACKENDS] =
 static const char *backend_name[Backend::NUM_BACKENDS] =
 {
    "ceed-cuda", "occa-cuda", "raja-cuda", "cuda",
-   "hip", "debug",
+   "ceed-hip", "raja-hip", "hip", "debug",
    "occa-omp", "raja-omp", "omp",
    "ceed-cpu", "occa-cpu", "raja-cpu", "cpu"
 };
@@ -61,10 +65,7 @@ Device Device::device_singleton;
 bool Device::device_env = false;
 bool Device::mem_host_env = false;
 bool Device::mem_device_env = false;
-#ifdef MFEM_USE_UMPIRE
-bool Device::use_host_umpire = true;
-bool Device::use_device_umpire = true;
-#endif
+bool Device::mem_types_set = false;
 
 Device::Device()
 {
@@ -132,7 +133,7 @@ Device::Device()
       {
          MFEM_ABORT("Unknown memory backend!");
       }
-      mm.Configure(host_mem_type, device_mem_type, device_mem_type);
+      mm.Configure(host_mem_type, device_mem_type);
    }
 
    if (getenv("MFEM_DEVICE"))
@@ -151,6 +152,18 @@ Device::~Device()
    {
       free(device_option);
 #ifdef MFEM_USE_CEED
+      // Destroy FES -> CeedBasis, CeedElemRestriction hash table contents
+      for (auto entry : internal::ceed_basis_map)
+      {
+         CeedBasisDestroy(&entry.second);
+      }
+      internal::ceed_basis_map.clear();
+      for (auto entry : internal::ceed_restr_map)
+      {
+         CeedElemRestrictionDestroy(&entry.second);
+      }
+      internal::ceed_restr_map.clear();
+      // Destroy Ceed context
       CeedDestroy(&internal::ceed);
 #endif
       mm.Destroy();
@@ -162,8 +175,6 @@ Device::~Device()
    Get().host_mem_class = MemoryClass::HOST;
    Get().device_mem_type = MemoryType::HOST;
    Get().device_mem_class = MemoryClass::HOST;
-   Get().device_temp_mem_type = MemoryType::HOST;
-   Get().device_temp_mem_class = MemoryClass::HOST;
 }
 
 void Device::Configure(const std::string &device, const int dev)
@@ -209,15 +220,24 @@ void Device::Configure(const std::string &device, const int dev)
       beg = end + 1;
    }
 
-   // OCCA_CUDA needs CUDA or RAJA_CUDA:
-   if (Allows(Backend::OCCA_CUDA) && !Allows(Backend::RAJA_CUDA))
+   // OCCA_CUDA and CEED_CUDA need CUDA or RAJA_CUDA:
+   if (Allows(Backend::OCCA_CUDA|Backend::CEED_CUDA) &&
+       !Allows(Backend::RAJA_CUDA))
    {
       Get().MarkBackend(Backend::CUDA);
    }
-   if (Allows(Backend::CEED_CUDA))
+   // CEED_HIP needs HIP:
+   if (Allows(Backend::CEED_HIP))
    {
-      Get().MarkBackend(Backend::CUDA);
+      Get().MarkBackend(Backend::HIP);
    }
+   // OCCA_OMP will use OMP or RAJA_OMP unless MFEM_USE_OPENMP=NO:
+#ifdef MFEM_USE_OPENMP
+   if (Allows(Backend::OCCA_OMP) && !Allows(Backend::RAJA_OMP))
+   {
+      Get().MarkBackend(Backend::OMP);
+   }
+#endif
 
    // Perform setup.
    Get().Setup(dev);
@@ -232,6 +252,30 @@ void Device::Configure(const std::string &device, const int dev)
    destroy_mm = true;
 }
 
+// static method
+void Device::SetMemoryTypes(MemoryType h_mt, MemoryType d_mt)
+{
+   // If the device and/or the MemoryTypes are configured through the
+   // environment (variables 'MFEM_DEVICE', 'MFEM_MEMORY'), ignore calls to this
+   // method.
+   if (mem_host_env || mem_device_env || device_env) { return; }
+
+   MFEM_VERIFY(!IsConfigured(), "the default MemoryTypes can only be set before"
+               " Device construction and configuration");
+   MFEM_VERIFY(IsHostMemory(h_mt),
+               "invalid host MemoryType, h_mt = " << (int)h_mt);
+   MFEM_VERIFY(IsDeviceMemory(d_mt) || d_mt == h_mt,
+               "invalid device MemoryType, d_mt = " << (int)d_mt
+               << " (h_mt = " << (int)h_mt << ')');
+
+   Get().host_mem_type = h_mt;
+   Get().device_mem_type = d_mt;
+   mem_types_set = true;
+
+   // h_mt and d_mt will be set as dual to each other during configuration by
+   // the call mm.Configure(...) in UpdateMemoryTypeAndClass()
+}
+
 void Device::Print(std::ostream &out)
 {
    out << "Device configuration: ";
@@ -259,10 +303,6 @@ void Device::Print(std::ostream &out)
    if (Device::Allows(Backend::DEVICE_MASK))
    {
       out << ',' << MemoryTypeName[static_cast<int>(device_mem_type)];
-      if (device_temp_mem_type != device_mem_type)
-      {
-         out << ',' << MemoryTypeName[static_cast<int>(device_temp_mem_type)];
-      }
    }
    out << std::endl;
 }
@@ -275,7 +315,14 @@ void Device::UpdateMemoryTypeAndClass()
 
 #ifdef MFEM_USE_UMPIRE
    // If MFEM has been compiled with Umpire support, use it as the default
-   if (!mem_host_env && use_host_umpire) { host_mem_type = MemoryType::HOST_UMPIRE; }
+   if (!mem_host_env && !mem_types_set)
+   {
+      host_mem_type = MemoryType::HOST_UMPIRE;
+      if (!mem_device_env)
+      {
+         device_mem_type = MemoryType::HOST_UMPIRE;
+      }
+   }
 #endif
 
    // Enable the device memory type
@@ -297,18 +344,13 @@ void Device::UpdateMemoryTypeAndClass()
                   device_mem_type = MemoryType::DEVICE;
             }
          }
-         else
+         else if (!mem_types_set)
          {
-#ifdef MFEM_USE_UMPIRE
-            if (use_device_umpire)
-            {
-               device_mem_type = MemoryType::DEVICE_UMPIRE;
-            }
-            else
+#ifndef MFEM_USE_UMPIRE
+            device_mem_type = MemoryType::DEVICE;
+#else
+            device_mem_type = MemoryType::DEVICE_UMPIRE;
 #endif
-            {
-               device_mem_type = MemoryType::DEVICE;
-            }
          }
       }
       device_mem_class = MemoryClass::DEVICE;
@@ -328,21 +370,11 @@ void Device::UpdateMemoryTypeAndClass()
       device_mem_type = MemoryType::DEVICE_DEBUG;
    }
 
-   // Setup device_temp_mem_{type,class}
-   switch (device_mem_type)
-   {
-      case MemoryType::DEVICE_UMPIRE:
-         device_temp_mem_type = MemoryType::DEVICE_TEMP_UMPIRE;
-         device_temp_mem_class = MemoryClass::DEVICE_TEMP;
-         break;
-      default:
-         device_temp_mem_type = device_mem_type;
-         device_temp_mem_class = device_mem_class;
-         break;
-   }
+   MFEM_VERIFY(!device || IsDeviceMemory(device_mem_type),
+               "invalid device memory configuration!");
 
    // Update the memory manager with the new settings
-   mm.Configure(host_mem_type, device_mem_type, device_temp_mem_type);
+   mm.Configure(host_mem_type, device_mem_type);
 }
 
 void Device::Enable()
@@ -390,6 +422,8 @@ static void RajaDeviceSetup(const int dev, int &ngpu)
 {
 #ifdef MFEM_USE_CUDA
    if (ngpu <= 0) { DeviceSetup(dev, ngpu); }
+#elif defined(MFEM_USE_HIP)
+   HipDeviceSetup(dev, ngpu);
 #else
    MFEM_CONTRACT_VAR(dev);
    MFEM_CONTRACT_VAR(ngpu);
@@ -456,7 +490,8 @@ static void CeedDeviceSetup(const char* ceed_spec)
    CeedInit(ceed_spec, &internal::ceed);
    const char *ceed_backend;
    CeedGetResource(internal::ceed, &ceed_backend);
-   if (strcmp(ceed_spec, ceed_backend) && strcmp(ceed_spec, "/cpu/self"))
+   if (strcmp(ceed_spec, ceed_backend) && strcmp(ceed_spec, "/cpu/self") &&
+       strcmp(ceed_spec, "/gpu/hip"))
    {
       mfem::out << std::endl << "WARNING!!!\n"
                 "libCEED is not using the requested backend!!!\n"
@@ -494,12 +529,16 @@ void Device::Setup(const int device)
    MFEM_VERIFY(!Allows(Backend::CEED_MASK),
                "the CEED backends require MFEM built with MFEM_USE_CEED=YES");
 #else
-   MFEM_VERIFY(!Allows(Backend::CEED_CPU) || !Allows(Backend::CEED_CUDA),
+   int ceed_cpu  = Allows(Backend::CEED_CPU);
+   int ceed_cuda = Allows(Backend::CEED_CUDA);
+   int ceed_hip  = Allows(Backend::CEED_HIP);
+   MFEM_VERIFY(ceed_cpu + ceed_cuda + ceed_hip <= 1,
                "Only one CEED backend can be enabled at a time!");
 #endif
    if (Allows(Backend::CUDA)) { CudaDeviceSetup(dev, ngpu); }
    if (Allows(Backend::HIP)) { HipDeviceSetup(dev, ngpu); }
-   if (Allows(Backend::RAJA_CUDA)) { RajaDeviceSetup(dev, ngpu); }
+   if (Allows(Backend::RAJA_CUDA) || Allows(Backend::RAJA_HIP))
+   { RajaDeviceSetup(dev, ngpu); }
    // The check for MFEM_USE_OCCA is in the function OccaDeviceSetup().
    if (Allows(Backend::OCCA_MASK)) { OccaDeviceSetup(dev); }
    if (Allows(Backend::CEED_CPU))
@@ -525,6 +564,17 @@ void Device::Setup(const int device)
          CeedDeviceSetup(device_option);
       }
    }
+   if (Allows(Backend::CEED_HIP))
+   {
+      if (!device_option)
+      {
+         CeedDeviceSetup("/gpu/hip");
+      }
+      else
+      {
+         CeedDeviceSetup(device_option);
+      }
+   }
    if (Allows(Backend::DEBUG)) { ngpu = 1; }
 }
 

general/device.hpp

@@ -1,4 +1,4 @@
-// Copyright (c) 2010-2020, Lawrence Livermore National Security, LLC. Produced
+// Copyright (c) 2010-2021, Lawrence Livermore National Security, LLC. Produced
 // at the Lawrence Livermore National Laboratory. All Rights reserved. See files
 // LICENSE and NOTICE for details. LLNL-CODE-806117.
 //
@@ -46,26 +46,33 @@ struct Backend
       /** @brief [device] RAJA CUDA backend. Enabled when MFEM_USE_RAJA = YES
           and MFEM_USE_CUDA = YES. */
       RAJA_CUDA = 1 << 6,
+      /** @brief [device] RAJA HIP backend. Enabled when MFEM_USE_RAJA = YES
+          and MFEM_USE_HIP = YES. */
+      RAJA_HIP = 1 << 7,
       /** @brief [host] OCCA CPU backend: sequential execution on each MPI rank.
           Enabled when MFEM_USE_OCCA = YES. */
-      OCCA_CPU = 1 << 7,
+      OCCA_CPU = 1 << 8,
       /// [host] OCCA OpenMP backend. Enabled when MFEM_USE_OCCA = YES.
-      OCCA_OMP = 1 << 8,
+      OCCA_OMP = 1 << 9,
       /** @brief [device] OCCA CUDA backend. Enabled when MFEM_USE_OCCA = YES
           and MFEM_USE_CUDA = YES. */
-      OCCA_CUDA = 1 << 9,
+      OCCA_CUDA = 1 << 10,
       /** @brief [host] CEED CPU backend. GPU backends can still be used, but
           with expensive memory transfers. Enabled when MFEM_USE_CEED = YES. */
-      CEED_CPU  = 1 << 10,
+      CEED_CPU  = 1 << 11,
       /** @brief [device] CEED CUDA backend working together with the CUDA
           backend. Enabled when MFEM_USE_CEED = YES and MFEM_USE_CUDA = YES.
-          NOTE: The current default libCEED GPU backend is non-deterministic! */
-      CEED_CUDA = 1 << 11,
+          NOTE: The current default libCEED CUDA backend is non-deterministic! */
+      CEED_CUDA = 1 << 12,
+      /** @brief [device] CEED HIP backend working together with the HIP
+          backend. Enabled when MFEM_USE_CEED = YES and MFEM_USE_HIP = YES. */
+      CEED_HIP = 1 << 13,
       /** @brief [device] Debug backend: host memory is READ/WRITE protected
           while a device is in use. It allows to test the "device" code-path
           (using separate host/device memory pools and host <-> device
-          transfers) without any GPU hardware. */
-      DEBUG = 1 << 12
+          transfers) without any GPU hardware. As 'DEBUG' is sometimes used
+          as a macro, `_DEVICE` has been added to avoid conflicts. */
+      DEBUG = 1 << 14
    };
 
    /** @brief Additional useful constants. For example, the *_MASK constants can
@@ -73,23 +80,23 @@ struct Backend
    enum
    {
       /// Number of backends: from (1 << 0) to (1 << (NUM_BACKENDS-1)).
-      NUM_BACKENDS = 13,
+      NUM_BACKENDS = 15,
 
       /// Biwise-OR of all CPU backends
       CPU_MASK = CPU | RAJA_CPU | OCCA_CPU | CEED_CPU,
       /// Biwise-OR of all CUDA backends
       CUDA_MASK = CUDA | RAJA_CUDA | OCCA_CUDA | CEED_CUDA,
       /// Biwise-OR of all HIP backends
-      HIP_MASK = HIP,
+      HIP_MASK = HIP | RAJA_HIP | CEED_HIP,
       /// Biwise-OR of all OpenMP backends
       OMP_MASK = OMP | RAJA_OMP | OCCA_OMP,
       /// Bitwise-OR of all CEED backends
-      CEED_MASK = CEED_CPU | CEED_CUDA,
+      CEED_MASK = CEED_CPU | CEED_CUDA | CEED_HIP,
       /// Biwise-OR of all device backends
       DEVICE_MASK = CUDA_MASK | HIP_MASK | DEBUG,
 
       /// Biwise-OR of all RAJA backends
-      RAJA_MASK = RAJA_CPU | RAJA_OMP | RAJA_CUDA,
+      RAJA_MASK = RAJA_CPU | RAJA_OMP | RAJA_CUDA | RAJA_HIP,
       /// Biwise-OR of all OCCA backends
       OCCA_MASK = OCCA_CPU | OCCA_OMP | OCCA_CUDA
    };
@@ -117,29 +124,25 @@ private:
    friend class MemoryManager;
    enum MODES {SEQUENTIAL, ACCELERATED};
 
-   static bool device_env, mem_host_env, mem_device_env;
+   static bool device_env, mem_host_env, mem_device_env, mem_types_set;
    static Device device_singleton;
-#ifdef MFEM_USE_UMPIRE
-   static bool use_host_umpire;
-   static bool use_device_umpire;
-#endif
 
-   MODES mode{Device::SEQUENTIAL};
-   int dev = 0; ///< Device ID of the configured device.
+   MODES mode = Device::SEQUENTIAL;
+   int dev = 0;   ///< Device ID of the configured device.
    int ngpu = -1; ///< Number of detected devices; -1: not initialized.
-   unsigned long backends{Backend::CPU}; ///< Bitwise-OR of all configured backends.
+   /// Bitwise-OR of all configured backends.
+   unsigned long backends = Backend::CPU;
    /// Set to true during configuration, except in 'device_singleton'.
-   bool destroy_mm{false};
-   bool mpi_gpu_aware{false};
+   bool destroy_mm = false;
+   bool mpi_gpu_aware = false;
 
-   MemoryType host_mem_type{MemoryType::HOST};      ///< Current Host MemoryType
-   MemoryClass host_mem_class{MemoryClass::HOST};    ///< Current Host MemoryClass
+   MemoryType host_mem_type = MemoryType::HOST;    ///< Current Host MemoryType
+   MemoryClass host_mem_class = MemoryClass::HOST; ///< Current Host MemoryClass
 
-   MemoryType device_mem_type{MemoryType::HOST};    ///< Current Device MemoryType
-   MemoryClass device_mem_class{MemoryClass::HOST};  ///< Current Device MemoryClass
-
-   MemoryType device_temp_mem_type{MemoryType::HOST};    ///< Current Device MemoryType
-   MemoryClass device_temp_mem_class{MemoryClass::HOST};  ///< Current Device MemoryClass
+   /// Current Device MemoryType
+   MemoryType device_mem_type = MemoryType::HOST;
+   /// Current Device MemoryClass
+   MemoryClass device_mem_class = MemoryClass::HOST;
 
    char *device_option = NULL;
    Device(Device const&);
@@ -192,25 +195,44 @@ public:
        * The available backends are described by the Backend class.
        * The string name of a backend is the lowercase version of the
          Backend::Id enumeration constant with '_' replaced by '-', e.g. the
-         string name of 'RAJA_CPU' is 'raja-cpu'.
+         string name of 'RAJA_CPU' is 'raja-cpu'. The string name of the debug
+         backend (Backend::Id 'DEBUG_DEVICE') is exceptionally set to 'debug'.
        * The 'cpu' backend is always enabled with lowest priority.
        * The current backend priority from highest to lowest is:
-         'ceed-cuda', 'occa-cuda', 'raja-cuda', 'cuda', 'hip', 'debug',
+         'ceed-cuda', 'occa-cuda', 'raja-cuda', 'cuda',
+         'ceed-hip', 'hip', 'debug',
          'occa-omp', 'raja-omp', 'omp',
          'ceed-cpu', 'occa-cpu', 'raja-cpu', 'cpu'.
        * Multiple backends can be configured at the same time.
        * Only one 'occa-*' backend can be configured at a time.
        * The backend 'occa-cuda' enables the 'cuda' backend unless 'raja-cuda'
          is already enabled.
+       * The backend 'occa-omp' enables the 'omp' backend (if MFEM was built
+         with MFEM_USE_OPENMP=YES) unless 'raja-omp' is already enabled.
+       * Only one 'ceed-*' backend can be configured at a time.
        * The backend 'ceed-cpu' delegates to a libCEED CPU backend the setup and
          evaluation of the operator.
        * The backend 'ceed-cuda' delegates to a libCEED CUDA backend the setup
-         and evaluation of the operator and enables the 'cuda' backend to avoid
-         transfer between host and device.
+         and evaluation of operators and enables the 'cuda' backend to avoid
+         transfers between host and device.
+       * The backend 'ceed-hip' delegates to a libCEED HIP backend the setup
+         and evaluation of operators and enables the 'hip' backend to avoid
+         transfers between host and device.
        * The 'debug' backend should not be combined with other device backends.
    */
    void Configure(const std::string &device, const int dev = 0);
 
+   /// Set the default host and device MemoryTypes, @a h_mt and @a d_mt.
+   /** The host and device MemoryTypes are also set to be dual to each other.
+
+       These two MemoryType%s are used by most MFEM classes when allocating
+       memory used on host and device, respectively.
+
+       This method can only be called before Device construction and
+       configuration, and the specified memory types must be compatible with
+       the subsequent Device configuration. */
+   static void SetMemoryTypes(MemoryType h_mt, MemoryType d_mt);
+
    /// Print the configuration of the MFEM virtual device object.
    void Print(std::ostream &out = mfem::out);
 
@@ -259,26 +281,10 @@ public:
    /** @deprecated Use GetDeviceMemoryClass() instead. */
    static inline MemoryClass GetMemoryClass() { return Get().device_mem_class; }
 
-   /** @brief Get the current Device Temporary MemoryType. This is the MemoryType used by
-       MFEM classes when allocating temporary memory to be used with device kernels.
-   */
-   static inline MemoryType GetDeviceTempMemoryType() { return Get().device_temp_mem_type; }
-
-   /** @brief Get the current Device Temporary MemoryClass. This is the MemoryClass used
-       by MFEM device kernels when they need to access temporary Memory objects. */
-   static inline MemoryClass GetDeviceTempMemoryClass() { return Get().device_temp_mem_class; }
-
    static void SetGPUAwareMPI(const bool force = true)
    { Get().mpi_gpu_aware = force; }
 
    static bool GetGPUAwareMPI() { return Get().mpi_gpu_aware; }
-
-#ifdef MFEM_USE_UMPIRE
-   static bool UseHostUmpire() { return Get().use_host_umpire; }
-   static void UseHostUmpire(bool use) { Get().use_host_umpire = use; }
-   static bool UseDeviceUmpire() { return Get().use_device_umpire; }
-   static void UseDeviceUmpire(bool use) { Get().use_device_umpire = use; }
-#endif
 };
 
 
@@ -298,8 +304,7 @@ MemoryClass GetMemoryClass(const Memory<T> &mem, bool on_dev)
    else
    {
       mem.UseDevice(true);
-      if (mem.UseTemporary()) { return Device::GetDeviceTempMemoryClass(); }
-      else { return Device::GetDeviceMemoryClass(); }
+      return Device::GetDeviceMemoryClass();
    }
 }
 

general/hip.cpp

@@ -1,4 +1,4 @@
-// Copyright (c) 2010-2020, Lawrence Livermore National Security, LLC. Produced
+// Copyright (c) 2010-2021, Lawrence Livermore National Security, LLC. Produced
 // at the Lawrence Livermore National Laboratory. All Rights reserved. See files
 // LICENSE and NOTICE for details. LLNL-CODE-806117.
 //
@@ -61,6 +61,21 @@ void* HipMallocManaged(void** dptr, size_t bytes)
    return *dptr;
 }
 
+void* HipMemAllocHostPinned(void** ptr, size_t bytes)
+{
+#ifdef MFEM_USE_HIP
+#ifdef MFEM_TRACK_HIP_MEM
+   mfem::out << "HipMemAllocHostPinned(): allocating " << bytes << " bytes ... "
+             << std::flush;
+#endif
+   MFEM_GPU_CHECK(hipHostMalloc(ptr, bytes, hipHostMallocDefault));
+#ifdef MFEM_TRACK_HIP_MEM
+   mfem::out << "done: " << *ptr << std::endl;
+#endif
+#endif
+   return *ptr;
+}
+
 void* HipMemFree(void *dptr)
 {
 #ifdef MFEM_USE_HIP
@@ -76,6 +91,21 @@ void* HipMemFree(void *dptr)
    return dptr;
 }
 
+void* HipMemFreeHostPinned(void *ptr)
+{
+#ifdef MFEM_USE_HIP
+#ifdef MFEM_TRACK_HIP_MEM
+   mfem::out << "HipMemFreeHostPinned(): deallocating memory @ " << ptr << " ... "
+             << std::flush;
+#endif
+   MFEM_GPU_CHECK(hipHostFree(ptr));
+#ifdef MFEM_TRACK_HIP_MEM
+   mfem::out << "done." << std::endl;
+#endif
+#endif
+   return ptr;
+}
+
 void* HipMemcpyHtoD(void* dst, const void* src, size_t bytes)
 {
 #ifdef MFEM_USE_HIP
@@ -125,7 +155,7 @@ void* HipMemcpyDtoDAsync(void* dst, const void *src, size_t bytes)
 void* HipMemcpyDtoH(void *dst, const void *src, size_t bytes)
 {
 #ifdef MFEM_USE_HIP
-#ifdef MFEM_TRACK_HPI_MEM
+#ifdef MFEM_TRACK_HIP_MEM
    mfem::out << "HipMemcpyDtoH(): copying " << bytes << " bytes from "
              << src << " to " << dst << " ... " << std::flush;
 #endif

general/hip.hpp

@@ -1,4 +1,4 @@
-// Copyright (c) 2010-2020, Lawrence Livermore National Security, LLC. Produced
+// Copyright (c) 2010-2021, Lawrence Livermore National Security, LLC. Produced
 // at the Lawrence Livermore National Laboratory. All Rights reserved. See files
 // LICENSE and NOTICE for details. LLNL-CODE-806117.
 //
@@ -43,7 +43,6 @@
 #if defined(MFEM_USE_HIP) && defined(__HIP_DEVICE_COMPILE__)
 #define MFEM_SHARED __shared__
 #define MFEM_SYNC_THREAD __syncthreads()
-#define MFEM_BLOCK_ID(k) hipBlockIdx_ ##k
 #define MFEM_THREAD_ID(k) hipThreadIdx_ ##k
 #define MFEM_THREAD_SIZE(k) hipBlockDim_ ##k
 #define MFEM_FOREACH_THREAD(i,k,N) \
@@ -65,9 +64,15 @@ void* HipMemAlloc(void **d_ptr, size_t bytes);
 /// Allocates managed device memory
 void* HipMallocManaged(void **d_ptr, size_t bytes);
 
+/// Allocates page-locked (pinned) host memory
+void* HipMemAllocHostPinned(void **ptr, size_t bytes);
+
 /// Frees device memory
 void* HipMemFree(void *d_ptr);
 
+/// Frees page-locked (pinned) host memory and returns destination ptr.
+void* HipMemFreeHostPinned(void *ptr);
+
 /// Copies memory from Host to Device
 void* HipMemcpyHtoD(void *d_dst, const void *h_src, size_t bytes);
 

general/mem_manager.cpp

@@ -1,4 +1,4 @@
-// Copyright (c) 2010-2020, Lawrence Livermore National Security, LLC. Produced
+// Copyright (c) 2010-2021, Lawrence Livermore National Security, LLC. Produced
 // at the Lawrence Livermore National Laboratory. All Rights reserved. See files
 // LICENSE and NOTICE for details. LLNL-CODE-806117.
 //
@@ -56,58 +56,45 @@ MemoryType GetMemoryType(MemoryClass mc)
       case MemoryClass::HOST_32: return MemoryType::HOST_32;
       case MemoryClass::HOST_64: return MemoryType::HOST_64;
       case MemoryClass::DEVICE:  return mm.GetDeviceMemoryType();
-      case MemoryClass::DEVICE_TEMP:  return mm.GetDeviceTempMemoryType();
       case MemoryClass::MANAGED: return MemoryType::MANAGED;
    }
    MFEM_VERIFY(false,"");
    return MemoryType::HOST;
 }
 
-// We want to keep this pairs, as it is checked in MFEM_VERIFY_TYPES
-MemoryType MemoryManager::GetDualMemoryType_(MemoryType mt)
-{
-   switch (mt)
-   {
-      // TODO TMS: temporary
-      case MemoryType::HOST:           return MemoryType::DEVICE_UMPIRE;
-      case MemoryType::HOST_32:        return MemoryType::DEVICE;
-      case MemoryType::HOST_64:        return MemoryType::DEVICE;
-      case MemoryType::HOST_DEBUG:     return MemoryType::DEVICE_DEBUG;
-      case MemoryType::HOST_UMPIRE:    return MemoryType::DEVICE_UMPIRE;
-      case MemoryType::HOST_PINNED:    return MemoryType::DEVICE_TEMP_UMPIRE;
-      case MemoryType::MANAGED:        return MemoryType::MANAGED;
-      case MemoryType::DEVICE:         return MemoryType::HOST;
-      case MemoryType::DEVICE_DEBUG:   return MemoryType::HOST_DEBUG;
-      //case MemoryType::DEVICE_UMPIRE:  return MemoryType::HOST_UMPIRE;
-      case MemoryType::DEVICE_UMPIRE:  return MemoryType::HOST;
-      //case MemoryType::DEVICE_TEMP_UMPIRE:  return MemoryType::HOST_UMPIRE;
-      case MemoryType::DEVICE_TEMP_UMPIRE:  return MemoryType::HOST;
-      default: mfem_error("Unknown memory type!");
-   }
-   MFEM_VERIFY(false,"");
-   return MemoryType::HOST;
-}
 
 static void MFEM_VERIFY_TYPES(const MemoryType h_mt, const MemoryType d_mt)
 {
-   MFEM_ASSERT(IsHostMemory(h_mt),"");
-   MFEM_ASSERT(IsDeviceMemory(d_mt),"");
+   MFEM_VERIFY(IsHostMemory(h_mt), "h_mt = " << (int)h_mt);
+   MFEM_VERIFY(IsDeviceMemory(d_mt) || d_mt == MemoryType::DEFAULT,
+               "d_mt = " << (int)d_mt);
+   // If h_mt == MemoryType::HOST_DEBUG, then d_mt == MemoryType::DEVICE_DEBUG
+   //                                      or d_mt == MemoryType::DEFAULT
+   MFEM_VERIFY(h_mt != MemoryType::HOST_DEBUG ||
+               d_mt == MemoryType::DEVICE_DEBUG ||
+               d_mt == MemoryType::DEFAULT,
+               "d_mt = " << MemoryTypeName[(int)d_mt]);
+   // If d_mt == MemoryType::DEVICE_DEBUG, then h_mt == MemoryType::HOST_DEBUG
+   MFEM_VERIFY(d_mt != MemoryType::DEVICE_DEBUG ||
+               h_mt == MemoryType::HOST_DEBUG,
+               "h_mt = " << MemoryTypeName[(int)h_mt]);
+#if 0
    const bool sync =
-      // TODO TMS: temporary
+      (h_mt == MemoryType::HOST_PINNED && d_mt == MemoryType::DEVICE) ||
+      (h_mt == MemoryType::HOST_PINNED && d_mt == MemoryType::DEVICE_UMPIRE) ||
+      (h_mt == MemoryType::HOST_PINNED && d_mt == MemoryType::DEVICE_UMPIRE_2) ||
+      (h_mt == MemoryType::HOST_UMPIRE && d_mt == MemoryType::DEVICE) ||
       (h_mt == MemoryType::HOST_UMPIRE && d_mt == MemoryType::DEVICE_UMPIRE) ||
-      (h_mt == MemoryType::HOST_UMPIRE && d_mt == MemoryType::DEVICE_TEMP_UMPIRE) ||
-
-      (h_mt == MemoryType::HOST && d_mt == MemoryType::DEVICE_UMPIRE) ||
-      (h_mt == MemoryType::HOST && d_mt == MemoryType::DEVICE_TEMP_UMPIRE) ||
-
-      (h_mt == MemoryType::HOST_PINNED && d_mt == MemoryType::DEVICE_TEMP_UMPIRE) ||
-
+      (h_mt == MemoryType::HOST_UMPIRE && d_mt == MemoryType::DEVICE_UMPIRE_2) ||
       (h_mt == MemoryType::HOST_DEBUG && d_mt == MemoryType::DEVICE_DEBUG) ||
       (h_mt == MemoryType::MANAGED && d_mt == MemoryType::MANAGED) ||
       (h_mt == MemoryType::HOST_64 && d_mt == MemoryType::DEVICE) ||
       (h_mt == MemoryType::HOST_32 && d_mt == MemoryType::DEVICE) ||
-      (h_mt == MemoryType::HOST && d_mt == MemoryType::DEVICE);
+      (h_mt == MemoryType::HOST && d_mt == MemoryType::DEVICE) ||
+      (h_mt == MemoryType::HOST && d_mt == MemoryType::DEVICE_UMPIRE) ||
+      (h_mt == MemoryType::HOST && d_mt == MemoryType::DEVICE_UMPIRE_2);
    MFEM_VERIFY(sync, "");
+#endif
 }
 
 MemoryClass operator*(MemoryClass mc1, MemoryClass mc2)
@@ -149,9 +136,12 @@ struct Memory
    void *const h_ptr;
    void *d_ptr;
    const size_t bytes;
-   const MemoryType h_mt, d_mt;
+   const MemoryType h_mt;
+   MemoryType d_mt;
+   mutable bool h_rw, d_rw;
    Memory(void *p, size_t b, MemoryType h, MemoryType d):
-      h_ptr(p), d_ptr(nullptr), bytes(b), h_mt(h), d_mt(d) { }
+      h_ptr(p), d_ptr(nullptr), bytes(b), h_mt(h), d_mt(d),
+      h_rw(true), d_rw(true) { }
 };
 
 /// Alias class that holds the base memory region and the offset
@@ -187,8 +177,8 @@ public:
    virtual ~HostMemorySpace() { }
    virtual void Alloc(void **ptr, size_t bytes) { *ptr = std::malloc(bytes); }
    virtual void Dealloc(void *ptr) { std::free(ptr); }
-   virtual void Protect(const void*, size_t) { }
-   virtual void Unprotect(const void*, size_t) { }
+   virtual void Protect(const Memory&, size_t) { }
+   virtual void Unprotect(const Memory&, size_t) { }
    virtual void AliasProtect(const void*, size_t) { }
    virtual void AliasUnprotect(const void*, size_t) { }
 };
@@ -336,18 +326,16 @@ inline void MmuDealloc(void *ptr, const size_t bytes)
 /// MMU protection, through ::mprotect with no read/write accesses
 inline void MmuProtect(const void *ptr, const size_t bytes)
 {
-   static const bool mmu_protect_error = getenv("MFEM_MMU_PROTECT_ERROR");
    if (!::mprotect(const_cast<void*>(ptr), bytes, PROT_NONE)) { return; }
-   if (mmu_protect_error) { mfem_error("MMU protection (NONE) error"); }
+   mfem_error("MMU protection (NONE) error");
 }
 
 /// MMU un-protection, through ::mprotect with read/write accesses
 inline void MmuAllow(const void *ptr, const size_t bytes)
 {
    const int RW = PROT_READ | PROT_WRITE;
-   static const bool mmu_protect_error = getenv("MFEM_MMU_PROTECT_ERROR");
    if (!::mprotect(const_cast<void*>(ptr), bytes, RW)) { return; }
-   if (mmu_protect_error) { mfem_error("MMU protection (R/W) error"); }
+   mfem_error("MMU protection (R/W) error");
 }
 #else
 inline void MmuInit() { }
@@ -368,8 +356,10 @@ public:
    MmuHostMemorySpace(): HostMemorySpace() { MmuInit(); }
    void Alloc(void **ptr, size_t bytes) { MmuAlloc(ptr, bytes); }
    void Dealloc(void *ptr) { MmuDealloc(ptr, maps->memories.at(ptr).bytes); }
-   void Protect(const void *ptr, size_t bytes) { MmuProtect(ptr, bytes); }
-   void Unprotect(const void *ptr, size_t bytes) { MmuAllow(ptr, bytes); }
+   void Protect(const Memory& mem, size_t bytes)
+   { if (mem.h_rw) { mem.h_rw = false; MmuProtect(mem.h_ptr, bytes); } }
+   void Unprotect(const Memory &mem, size_t bytes)
+   { if (!mem.h_rw) { mem.h_rw = true; MmuAllow(mem.h_ptr, bytes); } }
    /// Aliases need to be restricted during protection
    void AliasProtect(const void *ptr, size_t bytes)
    { MmuProtect(MmuAddrR(ptr), MmuLengthR(ptr, bytes)); }
@@ -416,13 +406,29 @@ public:
    { return CuMemcpyDtoH(dst, src, bytes); }
 };
 
-/// The CUDA page-locked host memory space
-class CudaHostMemorySpace: public HostMemorySpace
+/// The CUDA/HIP page-locked host memory space
+class HostPinnedMemorySpace: public HostMemorySpace
 {
 public:
-   CudaHostMemorySpace(): HostMemorySpace() { }
-   void Alloc(void ** ptr, size_t bytes) override { CuMemAllocHost(ptr, bytes); }
-   void Dealloc(void *ptr) override { CuMemFreeHost(ptr); }
+   HostPinnedMemorySpace(): HostMemorySpace() { }
+   void Alloc(void ** ptr, size_t bytes) override
+   {
+#ifdef MFEM_USE_CUDA
+      CuMemAllocHostPinned(ptr, bytes);
+#endif
+#ifdef MFEM_USE_HIP
+      HipMemAllocHostPinned(ptr, bytes);
+#endif
+   }
+   void Dealloc(void *ptr) override
+   {
+#ifdef MFEM_USE_CUDA
+      CuMemFreeHostPinned(ptr);
+#endif
+#ifdef MFEM_USE_HIP
+      HipMemFreeHostPinned(ptr);
+#endif
+   }
 };
 
 /// The HIP device memory space
@@ -467,8 +473,10 @@ public:
    MmuDeviceMemorySpace(): DeviceMemorySpace() { }
    void Alloc(Memory &m) { MmuAlloc(&m.d_ptr, m.bytes); }
    void Dealloc(Memory &m) { MmuDealloc(m.d_ptr, m.bytes); }
-   void Protect(const Memory &m) { MmuProtect(m.d_ptr, m.bytes); }
-   void Unprotect(const Memory &m) { MmuAllow(m.d_ptr, m.bytes); }
+   void Protect(const Memory &m)
+   { if (m.d_rw) { m.d_rw = false; MmuProtect(m.d_ptr, m.bytes); } }
+   void Unprotect(const Memory &m)
+   { if (!m.d_rw) { m.d_rw = true; MmuAllow(m.d_ptr, m.bytes); } }
    /// Aliases need to be restricted during protection
    void AliasProtect(const void *ptr, size_t bytes)
    { MmuProtect(MmuAddrR(ptr), MmuLengthR(ptr, bytes)); }
@@ -483,101 +491,63 @@ public:
    { return std::memcpy(dst, src, bytes); }
 };
 
-#ifndef MFEM_USE_UMPIRE
-class UmpireHostMemorySpace : public NoHostMemorySpace { };
-class UmpireDeviceMemorySpace : public NoDeviceMemorySpace { };
-class UmpireDeviceTempMemorySpace : public NoDeviceMemorySpace { };
-#else
-
-// TODO TMS: replace with um.hasAllocatorId(int) when it exists
-bool UmpireHasId(const umpire::ResourceManager & rm, int id)
+#ifdef MFEM_USE_UMPIRE
+class UmpireMemorySpace
 {
-   const auto & ids = rm.getAllocatorIds();
-   return std::find(ids.begin(), ids.end(), id) != ids.end();
-}
-
-/// The Umpire host memory space
-class UmpireHostMemorySpace : public HostMemorySpace
-{
-private:
+protected:
    umpire::ResourceManager &rm;
-   umpire::Allocator h_allocator;
+   umpire::Allocator allocator;
    bool owns_allocator{false};
+
 public:
    // TODO: this only releases unused memory
-   ~UmpireHostMemorySpace() { if (owns_allocator) { h_allocator.release(); } }
-   UmpireHostMemorySpace(): HostMemorySpace(),
-      rm(umpire::ResourceManager::getInstance())
+   virtual ~UmpireMemorySpace() { if (owns_allocator) { allocator.release(); } }
+   UmpireMemorySpace(const char * name, const char * space)
+      : rm(umpire::ResourceManager::getInstance())
    {
-      const int id = MemoryManager::GetUmpireHostAllocatorId();
-      if (!UmpireHasId(rm, id))
+      if (!rm.isAllocator(name))
       {
-         h_allocator = rm.makeAllocator<umpire::strategy::DynamicPool>("MFEM_HOST",
-                                                                       rm.getAllocator("HOST"));
+         allocator = rm.makeAllocator<umpire::strategy::DynamicPool>(
+                        name, rm.getAllocator(space));
          owns_allocator = true;
       }
       else
       {
-         h_allocator = rm.getAllocator(id);
+         allocator = rm.getAllocator(name);
+         owns_allocator = false;
       }
-      MemoryManager::SetUmpireHostAllocatorId(id);
    }
-   void Alloc(void **ptr, size_t bytes) override { *ptr = h_allocator.allocate(bytes); }
-   void Dealloc(void *ptr) override { h_allocator.deallocate(ptr); }
+};
+
+/// The Umpire host memory space
+class UmpireHostMemorySpace : public HostMemorySpace, public UmpireMemorySpace
+{
+private:
+   umpire::strategy::AllocationStrategy *strat;
+public:
+   UmpireHostMemorySpace(const char * name)
+      : HostMemorySpace(),
+        UmpireMemorySpace(name, "HOST"),
+        strat(allocator.getAllocationStrategy()) {}
+   void Alloc(void **ptr, size_t bytes) override
+   { *ptr = allocator.allocate(bytes); }
+   void Dealloc(void *ptr) override { allocator.deallocate(ptr); }
    void Insert(void *ptr, size_t bytes)
-   { mfem_error("UmpireHostMemorySpace::Insert is unsupported"); }
+   { rm.registerAllocation(ptr, {ptr, bytes, strat}); }
 };
 
 /// The Umpire device memory space
-#ifdef MFEM_USE_CUDA
-class UmpireDeviceMemorySpaceImpl : public DeviceMemorySpace
+#if defined(MFEM_USE_CUDA) || defined(MFEM_USE_HIP)
+class UmpireDeviceMemorySpace : public DeviceMemorySpace,
+   public UmpireMemorySpace
 {
 public:
-   enum class AllocatorType { TEMPORARY, PERMANENT };
-private:
-   umpire::ResourceManager &rm;
-   umpire::Allocator d_allocator;
-   bool owns_allocator{false};
-
-   int SetupAllocator(int possible_id, const char * allocator_name)
-   {
-      if (!UmpireHasId(rm, possible_id))
-      {
-         d_allocator = rm.makeAllocator<umpire::strategy::DynamicPool>(allocator_name,
-                                                                       rm.getAllocator("DEVICE"));
-         owns_allocator = true;
-      }
-      else
-      {
-         d_allocator = rm.getAllocator(possible_id);
-      }
-
-      return d_allocator.getId();
-   }
-public:
-   // TODO: this only releases unused memory
-   ~UmpireDeviceMemorySpaceImpl() { if (owns_allocator) { d_allocator.release(); } }
-   UmpireDeviceMemorySpaceImpl(AllocatorType t): DeviceMemorySpace(),
-      rm(umpire::ResourceManager::getInstance())
-   {
-      switch (t)
-      {
-         case AllocatorType::PERMANENT:
-            MemoryManager::SetUmpireDeviceAllocatorId(SetupAllocator(
-                                                         MemoryManager::GetUmpireDeviceAllocatorId(),
-                                                         "MFEM_DEVICE"));
-            break;
-         case AllocatorType::TEMPORARY:
-            MemoryManager::SetUmpireDeviceTempAllocatorId(SetupAllocator(
-                                                             MemoryManager::GetUmpireDeviceTempAllocatorId(),
-                                                             "MFEM_DEVICE_TEMPORARY"));
-            break;
-         default:
-            mfem_error("Unknown Umpire AllocatorType");
-      }
-   }
-   void Alloc(Memory &base) override { base.d_ptr = d_allocator.allocate(base.bytes); }
-   void Dealloc(Memory &base) override { d_allocator.deallocate(base.d_ptr); }
+   UmpireDeviceMemorySpace(const char * name)
+      : DeviceMemorySpace(),
+        UmpireMemorySpace(name, "DEVICE") {}
+   void Alloc(Memory &base) override
+   { base.d_ptr = allocator.allocate(base.bytes); }
+   void Dealloc(Memory &base) override { rm.deallocate(base.d_ptr); }
    void *HtoD(void *dst, const void *src, size_t bytes) override
    {
 #ifdef MFEM_USE_CUDA
@@ -609,24 +579,13 @@ public:
       //rm.copy(dst, const_cast<void*>(src), bytes); return dst;
    }
 };
-
-class UmpireDeviceMemorySpace : public UmpireDeviceMemorySpaceImpl
-{
-public:
-   UmpireDeviceMemorySpace() : UmpireDeviceMemorySpaceImpl(
-         AllocatorType::PERMANENT) {}
-};
-
-class UmpireDeviceTempMemorySpace : public UmpireDeviceMemorySpaceImpl
-{
-public:
-   UmpireDeviceTempMemorySpace() : UmpireDeviceMemorySpaceImpl(
-         AllocatorType::TEMPORARY) {}
-};
 #else
-class UmpireDeviceMemorySpace : public NoDeviceMemorySpace { };
-class UmpireDeviceTempMemorySpace : public NoDeviceMemorySpace { };
-#endif // MFEM_USE_CUDA
+class UmpireDeviceMemorySpace : public NoDeviceMemorySpace
+{
+public:
+   UmpireDeviceMemorySpace(const char * /*unused*/) {}
+};
+#endif // MFEM_USE_CUDA || MFEM_USE_HIP
 #endif // MFEM_USE_UMPIRE
 
 /// Memory space controller class
@@ -666,6 +625,7 @@ public:
       device[static_cast<int>(MemoryType::DEVICE)-shift] = nullptr;
       device[static_cast<int>(MT::DEVICE_DEBUG)-shift] = nullptr;
       device[static_cast<int>(MT::DEVICE_UMPIRE)-shift] = nullptr;
+      device[static_cast<int>(MT::DEVICE_UMPIRE_2)-shift] = nullptr;
    }
 
    HostMemorySpace* Host(const MemoryType mt)
@@ -701,9 +661,14 @@ private:
       switch (mt)
       {
          case MT::HOST_DEBUG: return new MmuHostMemorySpace();
-         case MT::HOST_UMPIRE: return new UmpireHostMemorySpace();
-         case MT::HOST_PINNED: return new CudaHostMemorySpace();
-
+#ifdef MFEM_USE_UMPIRE
+         case MT::HOST_UMPIRE:
+            return new UmpireHostMemorySpace(
+                      MemoryManager::GetUmpireHostAllocatorName());
+#else
+         case MT::HOST_UMPIRE: return new NoHostMemorySpace();
+#endif
+         case MT::HOST_PINNED: return new HostPinnedMemorySpace();
          default: MFEM_ABORT("Unknown host memory controller!");
       }
       return nullptr;
@@ -713,8 +678,17 @@ private:
    {
       switch (mt)
       {
-         case MT::DEVICE_UMPIRE: return new UmpireDeviceMemorySpace();
-         case MT::DEVICE_TEMP_UMPIRE: return new UmpireDeviceTempMemorySpace();
+#ifdef MFEM_USE_UMPIRE
+         case MT::DEVICE_UMPIRE:
+            return new UmpireDeviceMemorySpace(
+                      MemoryManager::GetUmpireDeviceAllocatorName());
+         case MT::DEVICE_UMPIRE_2:
+            return new UmpireDeviceMemorySpace(
+                      MemoryManager::GetUmpireDevice2AllocatorName());
+#else
+         case MT::DEVICE_UMPIRE: return new NoDeviceMemorySpace();
+         case MT::DEVICE_UMPIRE_2: return new NoDeviceMemorySpace();
+#endif
          case MT::DEVICE_DEBUG: return new MmuDeviceMemorySpace();
          case MT::DEVICE:
          {
@@ -741,20 +715,50 @@ void *MemoryManager::New_(void *h_tmp, size_t bytes, MemoryType mt,
                           unsigned &flags)
 {
    MFEM_ASSERT(exists, "Internal error!");
-   MFEM_ASSERT(mt != MemoryType::HOST, "Internal error!");
-   const bool is_host_mem = IsHostMemory(mt);
-   const MemType dual_mt = GetDualMemoryType_(mt);
-   const MemType h_mt = is_host_mem ? mt : dual_mt;
-   const MemType d_mt = is_host_mem ? dual_mt : mt;
-   MFEM_VERIFY_TYPES(h_mt, d_mt);
-   void *h_ptr = h_tmp;
+   if (IsHostMemory(mt))
+   {
+      MFEM_ASSERT(mt != MemoryType::HOST && h_tmp == nullptr,
+                  "Internal error!");
+      // d_mt = MemoryType::DEFAULT means d_mt = GetDualMemoryType(h_mt),
+      // evaluated at the time when the device pointer is allocated, see
+      // GetDevicePtr() and GetAliasDevicePtr()
+      const MemoryType d_mt = MemoryType::DEFAULT;
+      // We rely on the next call using lazy dev alloc
+      return New_(h_tmp, bytes, mt, d_mt, Mem::VALID_HOST, flags);
+   }
+   else
+   {
+      const MemoryType h_mt = GetDualMemoryType(mt);
+      return New_(h_tmp, bytes, h_mt, mt, Mem::VALID_DEVICE, flags);
+   }
+}
+
+void *MemoryManager::New_(void *h_tmp, size_t bytes, MemoryType h_mt,
+                          MemoryType d_mt, unsigned valid_flags,
+                          unsigned &flags)
+{
+   MFEM_ASSERT(exists, "Internal error!");
+   MFEM_ASSERT(IsHostMemory(h_mt), "h_mt must be host type");
+   MFEM_ASSERT(IsDeviceMemory(d_mt) || d_mt == h_mt ||
+               d_mt == MemoryType::DEFAULT,
+               "d_mt must be device type, the same is h_mt, or DEFAULT");
+   MFEM_ASSERT((h_mt != MemoryType::HOST || h_tmp != nullptr) &&
+               (h_mt == MemoryType::HOST || h_tmp == nullptr),
+               "Internal error");
+   MFEM_ASSERT((valid_flags & ~(Mem::VALID_HOST | Mem::VALID_DEVICE)) == 0,
+               "Internal error");
+   void *h_ptr;
    if (h_tmp == nullptr) { ctrl->Host(h_mt)->Alloc(&h_ptr, bytes); }
-   flags = Mem::REGISTERED;
-   flags |= Mem::OWNS_INTERNAL | Mem::OWNS_HOST | Mem::OWNS_DEVICE;
-   flags |= is_host_mem ? Mem::VALID_HOST : Mem::VALID_DEVICE;
-   if (is_host_mem) { mm.Insert(h_ptr, bytes, h_mt, d_mt); }
-   else { mm.InsertDevice(nullptr, h_ptr, bytes, h_mt, d_mt); }
+   else { h_ptr = h_tmp; }
+   flags = Mem::REGISTERED | Mem::OWNS_INTERNAL | Mem::OWNS_HOST |
+           Mem::OWNS_DEVICE | valid_flags;
+   // The other New_() method relies on this lazy allocation behavior.
+   mm.Insert(h_ptr, bytes, h_mt, d_mt); // lazy dev alloc
+   // mm.InsertDevice(nullptr, h_ptr, bytes, h_mt, d_mt); // non-lazy dev alloc
+
+   // MFEM_VERIFY_TYPES(h_mt, mt); // done by mm.Insert() above
    CheckHostMemoryType_(h_mt, h_ptr);
+
    return h_ptr;
 }
 
@@ -766,9 +770,12 @@ void *MemoryManager::Register_(void *ptr, void *h_tmp, size_t bytes,
    MFEM_ASSERT(exists, "Internal error!");
    MFEM_ASSERT(!alias, "Cannot register an alias!");
    const bool is_host_mem = IsHostMemory(mt);
-   const MemType dual_mt = GetDualMemoryType_(mt);
-   const MemType h_mt = is_host_mem ? mt : dual_mt;
-   const MemType d_mt = is_host_mem ? dual_mt : mt;
+   const MemType h_mt = is_host_mem ? mt : GetDualMemoryType(mt);
+   const MemType d_mt = is_host_mem ? MemoryType::DEFAULT : mt;
+   // d_mt = MemoryType::DEFAULT means d_mt = GetDualMemoryType(h_mt),
+   // evaluated at the time when the device pointer is allocated, see
+   // GetDevicePtr() and GetAliasDevicePtr()
+
    MFEM_VERIFY_TYPES(h_mt, d_mt);
 
    if (ptr == nullptr && h_tmp == nullptr)
@@ -789,17 +796,42 @@ void *MemoryManager::Register_(void *ptr, void *h_tmp, size_t bytes,
    }
    else // DEVICE TYPES
    {
-      h_ptr = h_tmp;
-      if (own && h_tmp == nullptr) { ctrl->Host(h_mt)->Alloc(&h_ptr, bytes); }
+      MFEM_VERIFY(ptr, "cannot register NULL device pointer");
+      if (h_tmp == nullptr) { ctrl->Host(h_mt)->Alloc(&h_ptr, bytes); }
+      else { h_ptr = h_tmp; }
       mm.InsertDevice(ptr, h_ptr, bytes, h_mt, d_mt);
       flags = own ? flags | Mem::OWNS_DEVICE : flags & ~Mem::OWNS_DEVICE;
-      flags = own ? flags | Mem::OWNS_HOST   : flags & ~Mem::OWNS_HOST;
-      flags |= Mem::VALID_DEVICE;
+      flags |= (Mem::OWNS_HOST | Mem::VALID_DEVICE);
    }
    CheckHostMemoryType_(h_mt, h_ptr);
    return h_ptr;
 }
 
+void MemoryManager::Register_(void *h_ptr, void *d_ptr, size_t bytes,
+                              MemoryType h_mt, MemoryType d_mt,
+                              bool own, bool alias, unsigned &flags)
+{
+   MFEM_CONTRACT_VAR(alias);
+   MFEM_ASSERT(exists, "Internal error!");
+   MFEM_ASSERT(!alias, "Cannot register an alias!");
+   MFEM_VERIFY_TYPES(h_mt, d_mt);
+
+   if (h_ptr == nullptr && d_ptr == nullptr)
+   {
+      MFEM_VERIFY(bytes == 0, "internal error");
+      return;
+   }
+
+   flags |= Mem::REGISTERED | Mem::OWNS_INTERNAL;
+
+   mm.InsertDevice(d_ptr, h_ptr, bytes, h_mt, d_mt);
+   flags = (own ? flags | (Mem::OWNS_HOST | Mem::OWNS_DEVICE) :
+            flags & ~(Mem::OWNS_HOST | Mem::OWNS_DEVICE)) |
+           Mem::VALID_HOST;
+
+   CheckHostMemoryType_(h_mt, h_ptr);
+}
+
 void MemoryManager::Alias_(void *base_h_ptr, size_t offset, size_t bytes,
                            unsigned base_flags, unsigned &flags)
 {
@@ -809,6 +841,34 @@ void MemoryManager::Alias_(void *base_h_ptr, size_t offset, size_t bytes,
            ~(Mem::OWNS_HOST | Mem::OWNS_DEVICE);
 }
 
+void MemoryManager::SetDeviceMemoryType_(void *h_ptr, unsigned flags,
+                                         MemoryType d_mt)
+{
+   MFEM_VERIFY(h_ptr, "cannot set the device memory type: Memory is empty!");
+   if (!(flags & Mem::ALIAS))
+   {
+      auto mem_iter = maps->memories.find(h_ptr);
+      MFEM_VERIFY(mem_iter != maps->memories.end(), "internal error");
+      internal::Memory &mem = mem_iter->second;
+      if (mem.d_mt == d_mt) { return; }
+      MFEM_VERIFY(mem.d_ptr == nullptr, "cannot set the device memory type:"
+                  " device memory is allocated!");
+      mem.d_mt = d_mt;
+   }
+   else
+   {
+      auto alias_iter = maps->aliases.find(h_ptr);
+      MFEM_VERIFY(alias_iter != maps->aliases.end(), "internal error");
+      internal::Alias &alias = alias_iter->second;
+      internal::Memory &base_mem = *alias.mem;
+      if (base_mem.d_mt == d_mt) { return; }
+      MFEM_VERIFY(base_mem.d_ptr == nullptr,
+                  "cannot set the device memory type:"
+                  " alias' base device memory is allocated!");
+      base_mem.d_mt = d_mt;
+   }
+}
+
 MemoryType MemoryManager::Delete_(void *h_ptr, MemoryType mt, unsigned flags)
 {
    const bool alias = flags & Mem::ALIAS;
@@ -865,11 +925,12 @@ bool MemoryManager::MemoryClassCheck_(MemoryClass mc, void *h_ptr,
    const bool known = mm.IsKnown(h_ptr);
    const bool alias = mm.IsAlias(h_ptr);
    const bool check = known || ((flags & Mem::ALIAS) && alias);
-   MFEM_VERIFY(check,"Unknown host pointer: " << h_ptr);
+   MFEM_VERIFY(check, "Unknown host pointer: " << h_ptr);
    const internal::Memory &mem =
       (flags & Mem::ALIAS) ?
       *maps->aliases.at(h_ptr).mem : maps->memories.at(h_ptr);
-   const MemoryType &d_mt = mem.d_mt;
+   MemoryType d_mt = mem.d_mt;
+   if (d_mt == MemoryType::DEFAULT) { d_mt = GetDualMemoryType(h_mt); }
    switch (mc)
    {
       case MemoryClass::HOST_32:
@@ -888,7 +949,7 @@ bool MemoryManager::MemoryClassCheck_(MemoryClass mc, void *h_ptr,
          MFEM_VERIFY(d_mt == MemoryType::DEVICE ||
                      d_mt == MemoryType::DEVICE_DEBUG ||
                      d_mt == MemoryType::DEVICE_UMPIRE ||
-                     d_mt == MemoryType::DEVICE_TEMP_UMPIRE ||
+                     d_mt == MemoryType::DEVICE_UMPIRE_2 ||
                      d_mt == MemoryType::MANAGED,"");
          return true;
       }
@@ -923,12 +984,7 @@ void *MemoryManager::ReadWrite_(void *h_ptr, MemoryType h_mt, MemoryClass mc,
       flags = (flags | Mem::VALID_DEVICE) & ~Mem::VALID_HOST;
       if (flags & Mem::ALIAS)
       { return mm.GetAliasDevicePtr(h_ptr, bytes, copy); }
-      else
-      {
-         // TODO TMS
-         if (flags & Mem::USE_TEMPORARY) { mm.UpdateDeviceMemoryType(h_ptr, GetDeviceTempMemoryType()); }
-         return mm.GetDevicePtr(h_ptr, bytes, copy);
-      }
+      else { return mm.GetDevicePtr(h_ptr, bytes, copy); }
    }
 }
 
@@ -952,12 +1008,7 @@ const void *MemoryManager::Read_(void *h_ptr, MemoryType h_mt, MemoryClass mc,
       flags |= Mem::VALID_DEVICE;
       if (flags & Mem::ALIAS)
       { return mm.GetAliasDevicePtr(h_ptr, bytes, copy); }
-      else
-      {
-         // TODO TMS
-         if (flags & Mem::USE_TEMPORARY) { mm.UpdateDeviceMemoryType(h_ptr, GetDeviceTempMemoryType()); }
-         return mm.GetDevicePtr(h_ptr, bytes, copy);
-      }
+      else { return mm.GetDevicePtr(h_ptr, bytes, copy); }
    }
 }
 
@@ -979,13 +1030,7 @@ void *MemoryManager::Write_(void *h_ptr, MemoryType h_mt, MemoryClass mc,
       flags = (flags | Mem::VALID_DEVICE) & ~Mem::VALID_HOST;
       if (flags & Mem::ALIAS)
       { return mm.GetAliasDevicePtr(h_ptr, bytes, false); }
-      else
-      {
-         // TODO TMS
-         if (flags & Mem::USE_TEMPORARY) { mm.UpdateDeviceMemoryType(h_ptr, GetDeviceTempMemoryType()); }
-         return mm.GetDevicePtr(h_ptr, bytes, false);
-      }
-
+      else { return mm.GetDevicePtr(h_ptr, bytes, false); }
    }
 }
 
@@ -1090,11 +1135,8 @@ void MemoryManager::Copy_(void *dst_h_ptr, const void *src_h_ptr,
       {
          if (dst_h_ptr != src_d_ptr && bytes != 0)
          {
-            internal::Memory &dst_h_base = maps->memories.at(dst_h_ptr);
             internal::Memory &src_d_base = maps->memories.at(src_d_ptr);
-            MemoryType dst_h_mt = dst_h_base.h_mt;
             MemoryType src_d_mt = src_d_base.d_mt;
-            ctrl->Host(dst_h_mt)->Unprotect(dst_h_ptr, bytes);
             ctrl->Device(src_d_mt)->DtoH(dst_h_ptr, src_d_ptr, bytes);
          }
       }
@@ -1219,7 +1261,7 @@ void MemoryManager::Insert(void *h_ptr, size_t bytes,
 void MemoryManager::InsertDevice(void *d_ptr, void *h_ptr, size_t bytes,
                                  MemoryType h_mt, MemoryType d_mt)
 {
-   MFEM_VERIFY_TYPES(h_mt, d_mt);
+   // MFEM_VERIFY_TYPES(h_mt, d_mt); // done by Insert() below
    MFEM_ASSERT(h_ptr != NULL, "internal error");
    Insert(h_ptr, bytes, h_mt, d_mt);
    internal::Memory &mem = maps->memories.at(h_ptr);
@@ -1246,7 +1288,6 @@ void MemoryManager::InsertAlias(const void *base_ptr, void *alias_ptr,
       offset += alias.offset;
    }
    internal::Memory &mem = maps->memories.at(base_ptr);
-   MFEM_VERIFY(mem.d_mt != MemoryType::DEVICE_TEMP_UMPIRE, "aliasing temp mem");
    auto res =
       maps->aliases.emplace(alias_ptr,
                             internal::Alias{&mem, offset, bytes, 1, mem.h_mt});
@@ -1278,8 +1319,10 @@ void MemoryManager::EraseDevice(void *h_ptr)
    if (!h_ptr) { return; }
    auto mem_map_iter = maps->memories.find(h_ptr);
    if (mem_map_iter == maps->memories.end()) { mfem_error("Unknown pointer!"); }
-   if (maps->aliases.find(h_ptr) != maps->aliases.end())
+   auto it = maps->aliases.find(h_ptr);
+   if (it != maps->aliases.end())
    {
+      fprintf(stderr, "count = %lu\n", it->second.counter);
       mfem_error("cannot delete aliased obj!");
    }
    internal::Memory &mem = mem_map_iter->second;
@@ -1287,20 +1330,6 @@ void MemoryManager::EraseDevice(void *h_ptr)
    mem.d_ptr = nullptr;
 }
 
-void MemoryManager::UpdateDeviceMemoryType(void *h_ptr, const MemoryType mt)
-{
-   if (!h_ptr) { return; }
-   MFEM_VERIFY(IsDeviceMemory(mt), "Invalid MemoryType");
-   auto mem_map_iter = maps->memories.find(h_ptr);
-   if (mem_map_iter == maps->memories.end()) { mfem_error("Unknown pointer!"); }
-   internal::Memory &mem = mem_map_iter->second;
-   if (mem.d_mt == mt) { return; }
-   MFEM_VERIFY(mem.d_ptr == nullptr,
-               "Cannot change the memory type if d_ptr != nullptr");
-   auto old = maps->memories.erase(mem_map_iter);
-   maps->memories.emplace(h_ptr, internal::Memory(h_ptr, mem.bytes, mem.h_mt, mt));
-}
-
 void MemoryManager::EraseAlias(void *alias_ptr)
 {
    if (!alias_ptr) { return; }
@@ -1321,16 +1350,21 @@ void *MemoryManager::GetDevicePtr(const void *h_ptr, size_t bytes,
    }
    internal::Memory &mem = maps->memories.at(h_ptr);
    const MemoryType &h_mt = mem.h_mt;
-   const MemoryType &d_mt = mem.d_mt;
+   MemoryType &d_mt = mem.d_mt;
    MFEM_VERIFY_TYPES(h_mt, d_mt);
-   if (!mem.d_ptr) { ctrl->Device(d_mt)->Alloc(mem); }
+   if (!mem.d_ptr)
+   {
+      if (d_mt == MemoryType::DEFAULT) { d_mt = GetDualMemoryType(h_mt); }
+      ctrl->Device(d_mt)->Alloc(mem);
+   }
+   // Aliases might have done some protections
    ctrl->Device(d_mt)->Unprotect(mem);
    if (copy_data)
    {
       MFEM_ASSERT(bytes <= mem.bytes, "invalid copy size");
       ctrl->Device(d_mt)->HtoD(mem.d_ptr, h_ptr, bytes);
    }
-   ctrl->Host(h_mt)->Protect(h_ptr, bytes);
+   ctrl->Host(h_mt)->Protect(mem, bytes);
    return mem.d_ptr;
 }
 
@@ -1349,13 +1383,18 @@ void *MemoryManager::GetAliasDevicePtr(const void *alias_ptr, size_t bytes,
    const size_t offset = alias.offset;
    internal::Memory &mem = *alias.mem;
    const MemoryType &h_mt = mem.h_mt;
-   const MemoryType &d_mt = mem.d_mt;
+   MemoryType &d_mt = mem.d_mt;
    MFEM_VERIFY_TYPES(h_mt, d_mt);
-   if (!mem.d_ptr) { ctrl->Device(d_mt)->Alloc(mem); }
+   if (!mem.d_ptr)
+   {
+      if (d_mt == MemoryType::DEFAULT) { d_mt = GetDualMemoryType(h_mt); }
+      ctrl->Device(d_mt)->Alloc(mem);
+   }
    void *alias_h_ptr = static_cast<char*>(mem.h_ptr) + offset;
    void *alias_d_ptr = static_cast<char*>(mem.d_ptr) + offset;
    MFEM_ASSERT(alias_h_ptr == alias_ptr, "internal error");
    MFEM_ASSERT(bytes <= alias.bytes, "internal error");
+   mem.d_rw = false;
    ctrl->Device(d_mt)->AliasUnprotect(alias_d_ptr, bytes);
    ctrl->Host(h_mt)->AliasUnprotect(alias_ptr, bytes);
    if (copy) { ctrl->Device(d_mt)->HtoD(alias_d_ptr, alias_h_ptr, bytes); }
@@ -1371,8 +1410,8 @@ void *MemoryManager::GetHostPtr(const void *ptr, size_t bytes, bool copy)
    const MemoryType &h_mt = mem.h_mt;
    const MemoryType &d_mt = mem.d_mt;
    MFEM_VERIFY_TYPES(h_mt, d_mt);
-   ctrl->Host(h_mt)->Unprotect(mem.h_ptr, bytes);
    // Aliases might have done some protections
+   ctrl->Host(h_mt)->Unprotect(mem, bytes);
    if (mem.d_ptr) { ctrl->Device(d_mt)->Unprotect(mem); }
    if (copy && mem.d_ptr) { ctrl->Device(d_mt)->DtoH(mem.h_ptr, mem.d_ptr, bytes); }
    if (mem.d_ptr) { ctrl->Device(d_mt)->Protect(mem); }
@@ -1390,6 +1429,7 @@ void *MemoryManager::GetAliasHostPtr(const void *ptr, size_t bytes,
    void *alias_h_ptr = static_cast<char*>(mem->h_ptr) + alias.offset;
    void *alias_d_ptr = static_cast<char*>(mem->d_ptr) + alias.offset;
    MFEM_ASSERT(alias_h_ptr == ptr,  "internal error");
+   mem->h_rw = false;
    ctrl->Host(h_mt)->AliasUnprotect(alias_h_ptr, bytes);
    if (mem->d_ptr) { ctrl->Device(d_mt)->AliasUnprotect(alias_d_ptr, bytes); }
    if (copy_data && mem->d_ptr)
@@ -1411,14 +1451,46 @@ MemoryManager::MemoryManager() { Init(); }
 
 MemoryManager::~MemoryManager() { if (exists) { Destroy(); } }
 
-void MemoryManager::Configure(const MemoryType host_mt,
-                              const MemoryType device_mt,
-                              const MemoryType device_tmt)
+void MemoryManager::SetDualMemoryType(MemoryType mt, MemoryType dual_mt)
 {
+   MFEM_VERIFY(!configured, "changing the dual MemoryTypes is not allowed after"
+               " MemoryManager configuration!");
+   UpdateDualMemoryType(mt, dual_mt);
+}
+
+void MemoryManager::UpdateDualMemoryType(MemoryType mt, MemoryType dual_mt)
+{
+   MFEM_VERIFY((int)mt < MemoryTypeSize,
+               "invalid MemoryType, mt = " << (int)mt);
+   MFEM_VERIFY((int)dual_mt < MemoryTypeSize,
+               "invalid dual MemoryType, dual_mt = " << (int)dual_mt);
+
+   if ((IsHostMemory(mt) && IsDeviceMemory(dual_mt)) ||
+       (IsDeviceMemory(mt) && IsHostMemory(dual_mt)))
+   {
+      dual_map[(int)mt] = dual_mt;
+   }
+   else
+   {
+      // mt + dual_mt is not a pair of host + device types: this is only allowed
+      // when mt == dual_mt and mt is a host type; in this case we do not
+      // actually update the dual
+      MFEM_VERIFY(mt == dual_mt && IsHostMemory(mt),
+                  "invalid (mt, dual_mt) pair: ("
+                  << MemoryTypeName[(int)mt] << ", "
+                  << MemoryTypeName[(int)dual_mt] << ')');
+   }
+}
+
+void MemoryManager::Configure(const MemoryType host_mt,
+                              const MemoryType device_mt)
+{
+   MemoryManager::UpdateDualMemoryType(host_mt, device_mt);
+   MemoryManager::UpdateDualMemoryType(device_mt, host_mt);
    Init();
    host_mem_type = host_mt;
    device_mem_type = device_mt;
-   device_temp_mem_type = device_tmt;
+   configured = true;
 }
 
 void MemoryManager::Destroy()
@@ -1436,6 +1508,7 @@ void MemoryManager::Destroy()
    host_mem_type = MemoryType::HOST;
    device_mem_type = MemoryType::HOST;
    exists = false;
+   configured = false;
 }
 
 void MemoryManager::RegisterCheck(void *ptr)
@@ -1522,16 +1595,32 @@ void MemoryManager::CheckHostMemoryType_(MemoryType h_mt, void *h_ptr)
 MemoryManager mm;
 
 bool MemoryManager::exists = false;
-
-#ifdef MFEM_USE_UMPIRE
-int MemoryManager::h_umpire_id = -1;
-int MemoryManager::d_umpire_id = -1;
-int MemoryManager::d_umpire_temp_id = -1;
-#endif
+bool MemoryManager::configured = false;
 
 MemoryType MemoryManager::host_mem_type = MemoryType::HOST;
 MemoryType MemoryManager::device_mem_type = MemoryType::HOST;
-MemoryType MemoryManager::device_temp_mem_type = MemoryType::HOST;
+
+MemoryType MemoryManager::dual_map[MemoryTypeSize] =
+{
+   /* HOST            */  MemoryType::DEVICE,
+   /* HOST_32         */  MemoryType::DEVICE,
+   /* HOST_64         */  MemoryType::DEVICE,
+   /* HOST_DEBUG      */  MemoryType::DEVICE_DEBUG,
+   /* HOST_UMPIRE     */  MemoryType::DEVICE_UMPIRE,
+   /* HOST_PINNED     */  MemoryType::DEVICE,
+   /* MANAGED         */  MemoryType::MANAGED,
+   /* DEVICE          */  MemoryType::HOST,
+   /* DEVICE_DEBUG    */  MemoryType::HOST_DEBUG,
+   /* DEVICE_UMPIRE   */  MemoryType::HOST_UMPIRE,
+   /* DEVICE_UMPIRE_2 */  MemoryType::HOST_UMPIRE
+};
+
+#ifdef MFEM_USE_UMPIRE
+const char * MemoryManager::h_umpire_name = "MFEM_HOST";
+const char * MemoryManager::d_umpire_name = "MFEM_DEVICE";
+const char * MemoryManager::d_umpire_2_name = "MFEM_DEVICE_2";
+#endif
+
 
 const char *MemoryTypeName[MemoryTypeSize] =
 {
@@ -1549,13 +1638,13 @@ const char *MemoryTypeName[MemoryTypeSize] =
    "device-debug",
 #if defined(MFEM_USE_CUDA)
    "cuda-umpire",
-   "cuda-umpire-temp"
+   "cuda-umpire-2",
 #elif defined(MFEM_USE_HIP)
    "hip-umpire",
-   "hip-umpire-temp"
+   "hip-umpire-2",
 #else
    "device-umpire",
-   "device-umpire-temp"
+   "device-umpire-2",
 #endif
 };
 

general/mem_manager.hpp

@@ -1,4 +1,4 @@
-// Copyright (c) 2010-2020, Lawrence Livermore National Security, LLC. Produced
+// Copyright (c) 2010-2021, Lawrence Livermore National Security, LLC. Produced
 // at the Lawrence Livermore National Laboratory. All Rights reserved. See files
 // LICENSE and NOTICE for details. LLNL-CODE-806117.
 //
@@ -27,20 +27,30 @@ namespace mfem
 /// Memory types supported by MFEM.
 enum class MemoryType
 {
-   HOST,                ///< Host memory; using new[] and delete[]
-   HOST_32,             ///< Host memory; aligned at 32 bytes
-   HOST_64,             ///< Host memory; aligned at 64 bytes
-   HOST_DEBUG,          ///< Host memory; allocated from a "host-debug" pool
-   HOST_UMPIRE,         ///< Host memory; using Umpire
-   HOST_PINNED,         ///< Host memory: pinned (page-locked)
-   MANAGED,             /**< Managed memory; using CUDA or HIP *MallocManaged
-                             and *Free */
-   DEVICE,              ///< Device memory; using CUDA or HIP *Malloc and *Free
-   DEVICE_DEBUG,        /**< Pseudo-device memory; allocated on host from a
-                             "device-debug" pool */
-   DEVICE_UMPIRE,       ///< Device memory; using Umpire
-   DEVICE_TEMP_UMPIRE,  ///< Temporary Device memory; using Umpire
-   SIZE                 ///< Number of host and device memory types
+   HOST,           ///< Host memory; using new[] and delete[]
+   HOST_32,        ///< Host memory; aligned at 32 bytes
+   HOST_64,        ///< Host memory; aligned at 64 bytes
+   HOST_DEBUG,     ///< Host memory; allocated from a "host-debug" pool
+   HOST_UMPIRE,    /**< Host memory; using an Umpire allocator which can be set
+                        with MemoryManager::SetUmpireHostAllocatorName */
+   HOST_PINNED,    ///< Host memory: pinned (page-locked)
+   MANAGED,        /**< Managed memory; using CUDA or HIP *MallocManaged
+                        and *Free */
+   DEVICE,         ///< Device memory; using CUDA or HIP *Malloc and *Free
+   DEVICE_DEBUG,   /**< Pseudo-device memory; allocated on host from a
+                        "device-debug" pool */
+   DEVICE_UMPIRE,  /**< Device memory; using an Umpire allocator which can be
+                        set with MemoryManager::SetUmpireDeviceAllocatorName */
+   DEVICE_UMPIRE_2, /**< Device memory; using a second Umpire allocator settable
+                         with MemoryManager::SetUmpireDevice2AllocatorName */
+   SIZE,           ///< Number of host and device memory types
+
+   PRESERVE,       /**< Pseudo-MemoryType used as default value for MemoryType
+                        parameters to request preservation of existing
+                        MemoryType, e.g. in copy constructors. */
+   DEFAULT         /**< Pseudo-MemoryType used as default value for MemoryType
+                        parameters to request the use of the default host or
+                        device MemoryType. */
 };
 
 /// Static casts to 'int' and sizes of some useful memory types.
@@ -60,17 +70,22 @@ extern const char *MemoryTypeName[MemoryTypeSize];
 enum class MemoryClass
 {
    HOST,    /**< Memory types: { HOST, HOST_32, HOST_64, HOST_DEBUG,
-                                 HOST_UMPIRE, MANAGED } */
+                                 HOST_UMPIRE, HOST_PINNED, MANAGED } */
    HOST_32, ///< Memory types: { HOST_32, HOST_64, HOST_DEBUG }
    HOST_64, ///< Memory types: { HOST_64, HOST_DEBUG }
-   DEVICE,  ///< Memory types: { DEVICE, DEVICE_DEBUG, DEVICE_UMPIRE, MANAGED }
-   DEVICE_TEMP, ///< Memory types: { DEVICE_TEMP_UMPIRE }
+   DEVICE,  /**< Memory types: { DEVICE, DEVICE_DEBUG, DEVICE_UMPIRE,
+                                 DEVICE_UMPIRE_2, MANAGED } */
    MANAGED  ///< Memory types: { MANAGED }
 };
 
 /// Return true if the given memory type is in MemoryClass::HOST.
 inline bool IsHostMemory(MemoryType mt) { return mt <= MemoryType::MANAGED; }
-inline bool IsDeviceMemory(MemoryType mt) { return mt >= MemoryType::MANAGED; }
+
+/// Return true if the given memory type is in MemoryClass::DEVICE
+inline bool IsDeviceMemory(MemoryType mt)
+{
+   return mt >= MemoryType::MANAGED && mt < MemoryType::SIZE;
+}
 
 /// Return a suitable MemoryType for a given MemoryClass.
 MemoryType GetMemoryType(MemoryClass mc);
@@ -140,9 +155,7 @@ protected:
       VALID_DEVICE  = 1 << 5, ///< %Device pointer is valid
       USE_DEVICE    = 1 << 6, /**< Internal device flag, see e.g.
                                    Vector::UseDevice() */
-      ALIAS         = 1 << 7, ///< Pointer is an alias
-      USE_TEMPORARY = 1 << 8  ///< Temporary Device memory flag
-
+      ALIAS         = 1 << 7  ///< Pointer is an alias
    };
 
    /// Pointer to host memory. Not owned.
@@ -182,6 +195,12 @@ public:
        MemoryType is still set as valid. */
    Memory(int size, MemoryType mt) { New(size, mt); }
 
+   /** @brief Allocate memory for @a size entries with the given host MemoryType
+       @a h_mt and device MemoryType @a d_mt. */
+   /** The newly allocated memory is not initialized. The host pointer is set as
+       valid. */
+   Memory(int size, MemoryType h_mt, MemoryType d_mt) { New(size, h_mt, d_mt); }
+
    /** @brief Wrap an externally allocated host pointer, @a ptr with the current
        host memory type returned by MemoryManager::GetHostMemoryType(). */
    /** The parameter @a own determines whether @a ptr will be deleted when the
@@ -239,18 +258,6 @@ public:
    void UseDevice(bool use_dev) const
    { flags = use_dev ? (flags | USE_DEVICE) : (flags & ~USE_DEVICE); }
 
-   bool UseTemporary() const { return flags & USE_TEMPORARY; }
-
-   void UseTemporary(bool use_temp)
-   {
-      if (use_temp != UseTemporary())
-      {
-         MFEM_VERIFY(!(flags & VALID_DEVICE),
-                     "Cannot change temporary status when the device pointer already exists");
-      }
-      flags = use_temp ? (flags | USE_TEMPORARY) : (flags & ~USE_TEMPORARY);
-   }
-
    /// Return the size of the allocated memory.
    int Capacity() const { return capacity; }
 
@@ -278,9 +285,24 @@ public:
    /** The newly allocated memory is not initialized, however the given
        MemoryType is still set as valid.
 
+       When @a mt is a host type, the device MemoryType will be set later, if
+       requested, using the dual type of @a mt, see
+       MemoryManager::GetDualMemoryType().
+
+       When @a mt is a device type, the host MemoryType will be set immediately
+       to be the dual of @a mt, see MemoryManager::GetDualMemoryType().
+
        @note The current memory is NOT deleted by this method. */
    inline void New(int size, MemoryType mt);
 
+   /** @brief Allocate memory for @a size entries with the given host MemoryType
+       @a h_mt and device MemoryType @a d_mt. */
+   /** The newly allocated memory is not initialized. The host pointer is set as
+       valid.
+
+       @note The current memory is NOT deleted by this method. */
+   inline void New(int size, MemoryType h_mt, MemoryType d_mt);
+
    /** @brief Wrap an externally allocated host pointer, @a ptr with the current
        host memory type returned by MemoryManager::GetHostMemoryType(). */
    /** The parameter @a own determines whether @a ptr will be deleted when the
@@ -301,21 +323,18 @@ public:
        @note The current memory is NOT deleted by this method. */
    inline void Wrap(T *ptr, int size, MemoryType mt, bool own);
 
-   /** Wrap an externally pair of allocated pointers, @a h_ptr and @ d_ptr,
+   /** Wrap an externally pair of allocated pointers, @a h_ptr and @a d_ptr,
        of the given host MemoryType @a h_mt. */
    /** The new memory object will have the device MemoryType set as valid.
 
        The given @a h_ptr and @a d_ptr must be allocated appropriately for the
-       given host MemoryType and its associated device MemoryType:
-          - MANAGED => MANAGED,
-          - HOST_DEBUG => DEVICE_DEBUG,
-          - HOST_UMPIRE => DEVICE_UMPIRE,
-          - HOST, HOST_32, HOST_64 => DEVICE.
+       given host MemoryType and its dual device MemoryType as defined by
+       MemoryManager::GetDualMemoryType().
 
        The parameter @a own determines whether both @a h_ptr and @a d_ptr will
        be deleted when the method Delete() is called.
 
-       @note Ownership can also be controled by using the folowing methods:
+       @note Ownership can also be controlled by using the following methods:
          - ClearOwnerFlags,
          - SetHostPtrOwner,
          - SetDevicePtrOwner.
@@ -329,12 +348,28 @@ public:
        @note The current memory is NOT deleted by this method. */
    inline void MakeAlias(const Memory &base, int offset, int size);
 
+   /// Set the device MemoryType to be used by the Memory object.
+   /** If the specified @a d_mt is not a device MemoryType, i.e. not one of the
+       types in MemoryClass::DEVICE, then this method will return immediately.
+
+       If the device MemoryType has been previously set to a different type and
+       the actual device memory has been allocated, this method will trigger an
+       error. This method will not perform the actual device memory allocation,
+       however, the allocation may already exist if the MemoryType is the same
+       as the current one.
+
+       If the Memory is an alias Memory, the device MemoryType of its base will
+       be updated as described above. */
+   inline void SetDeviceMemoryType(MemoryType d_mt);
+
    /** @brief Delete the owned pointers. The Memory is not reset by this method,
        i.e. it will, generally, not be Empty() after this call. */
    inline void Delete();
 
-   /** @brief Delete the owned device pointer. */
-   inline void DeleteDevice();
+   /** @brief Delete the device pointer, if owned. If @a copy_to_host is true
+       and the data is valid only on device, move it to host before deleting.
+       Invalidates the device memory. */
+   inline void DeleteDevice(bool copy_to_host = true);
 
    /// Array subscript operator for host memory.
    inline T &operator[](int idx);
@@ -488,6 +523,12 @@ private:
       static inline T *New(std::size_t size) { return new T[size]; }
    };
 #endif
+
+   // Shortcut for Alloc<new_align_bytes>::New(size)
+   static inline T *NewHOST(std::size_t size)
+   {
+      return Alloc<new_align_bytes>::New(size);
+   }
 };
 
 
@@ -509,20 +550,28 @@ private:
    /// Device memory type set during the Setup.
    static MemoryType device_mem_type;
 
-   /// Device temporary memory type set during the Setup.
-   static MemoryType device_temp_mem_type;
-
    /// Allow to detect if a global memory manager instance exists.
    static bool exists;
 
    /// Return true if the global memory manager instance exists.
    static bool Exists() { return exists; }
 
+   /// Array defining the dual MemoryType for each MemoryType
+   /** The dual of a host MemoryType is a device MemoryType and vice versa: the
+       dual of a device MemoryType is a host MemoryType. */
+   static MemoryType dual_map[MemoryTypeSize];
+
+   /// Update the dual memory type of @a mt to be @a dual_mt.
+   static void UpdateDualMemoryType(MemoryType mt, MemoryType dual_mt);
+
+   /// True if Configure() was called.
+   static bool configured;
+
    /// Host and device allocator names for Umpire.
 #ifdef MFEM_USE_UMPIRE
-   static int h_umpire_id;
-   static int d_umpire_id;
-   static int d_umpire_temp_id;
+   static const char * h_umpire_name;
+   static const char * d_umpire_name;
+   static const char * d_umpire_2_name;
 #endif
 
 private: // Static methods used by the Memory<T> class
@@ -532,15 +581,26 @@ private: // Static methods used by the Memory<T> class
    /// memory type, e.g. CUDA (mt will not be HOST).
    static void *New_(void *h_tmp, size_t bytes, MemoryType mt, unsigned &flags);
 
+   static void *New_(void *h_tmp, size_t bytes, MemoryType h_mt,
+                     MemoryType d_mt, unsigned valid_flags, unsigned &flags);
+
    /// Register an external pointer of the given MemoryType.
    /// Return the host pointer.
    static void *Register_(void *ptr, void *h_ptr, size_t bytes, MemoryType mt,
                           bool own, bool alias, unsigned &flags);
 
+   /// Register a pair of external host and device pointers
+   static void Register_(void *h_ptr, void *d_ptr, size_t bytes,
+                         MemoryType h_mt, MemoryType d_mt,
+                         bool own, bool alias, unsigned &flags);
+
    /// Register an alias. Note: base_h_ptr may be an alias.
    static void Alias_(void *base_h_ptr, size_t offset, size_t bytes,
                       unsigned base_flags, unsigned &flags);
 
+   static void SetDeviceMemoryType_(void *h_ptr, unsigned flags,
+                                    MemoryType d_mt);
+
    /// Un-register and free memory identified by its host pointer. Returns the
    /// memory type of the host pointer.
    static MemoryType Delete_(void *h_ptr, MemoryType mt, unsigned flags);
@@ -552,9 +612,6 @@ private: // Static methods used by the Memory<T> class
    static bool MemoryClassCheck_(MemoryClass mc, void *h_ptr,
                                  MemoryType h_mt, size_t bytes, unsigned flags);
 
-   /// Return the dual memory type of the given one.
-   static MemoryType GetDualMemoryType_(MemoryType mt);
-
    /// Return a pointer to the memory identified by the host pointer h_ptr for
    /// access with the given MemoryClass.
    static void *ReadWrite_(void *h_ptr, MemoryType h_mt, MemoryClass mc,
@@ -628,8 +685,6 @@ private:
    /// Erase an alias from the aliases map
    void EraseAlias(void *alias_ptr);
 
-   void UpdateDeviceMemoryType(void *h_ptr, const MemoryType mt);
-
    /// Return the corresponding device pointer of h_ptr,
    /// allocating and moving the data if needed
    void *GetDevicePtr(const void *h_ptr, size_t bytes, bool copy_data);
@@ -653,21 +708,59 @@ public:
    /// Initialize the memory manager.
    void Init();
 
-   /// Configure the Memory manager with given default host, device, and device temporary types
-   /// This method will be called when configuring a device.
-   void Configure(const MemoryType h_mt, const MemoryType d_mt,
-                  const MemoryType d_tmt);
+   /// Return the dual MemoryType of the given one, @a mt.
+   /** The default dual memory types are:
+
+       memory type     | dual type
+       --------------- | ---------
+       HOST            | DEVICE
+       HOST_32         | DEVICE
+       HOST_64         | DEVICE
+       HOST_DEBUG      | DEVICE_DEBUG
+       HOST_UMPIRE     | DEVICE_UMPIRE
+       HOST_PINNED     | DEVICE
+       MANAGED         | MANAGED
+       DEVICE          | HOST
+       DEVICE_DEBUG    | HOST_DEBUG
+       DEVICE_UMPIRE   | HOST_UMPIRE
+       DEVICE_UMPIRE_2 | HOST_UMPIRE
+
+       The dual types can be modified before device configuration using the
+       method SetDualMemoryType() or by calling Device::SetMemoryTypes(). */
+   static inline MemoryType GetDualMemoryType(MemoryType mt)
+   { return dual_map[(int)mt]; }
+
+   /// Set the dual memory type of @a mt to be @a dual_mt.
+   /** This method can only be called before configuration, i.e. before calling
+       Configure(), which is typically done during Device construction.
+
+       One of the types must be a host MemoryType and the other must be a device
+       MemoryType or both types must be the same host memory type. The latter
+       case is only allowed for convenience in setting up pure host execution,
+       so the actual dual is not updated. */
+   static void SetDualMemoryType(MemoryType mt, MemoryType dual_mt);
+
+   /** @brief Configure the Memory manager with given default host and device
+       types. This method will be called when configuring a device.
+
+       The host and device MemoryType%s, @a h_mt and @a d_mt, are set to be dual
+       to each other. */
+   void Configure(const MemoryType h_mt, const MemoryType d_mt);
 
 #ifdef MFEM_USE_UMPIRE
-   /// Set the host and device Umpire allocator ids
-   static void SetUmpireHostAllocatorId(int h_id) { h_umpire_id = h_id; }
-   static void SetUmpireDeviceAllocatorId(int d_id) { d_umpire_id = d_id; }
-   static void SetUmpireDeviceTempAllocatorId(int d_id) { d_umpire_temp_id = d_id; }
+   /// Set the host Umpire allocator name used with MemoryType::HOST_UMPIRE
+   static void SetUmpireHostAllocatorName(const char * h_name) { h_umpire_name = h_name; }
+   /// Set the device Umpire allocator name used with MemoryType::DEVICE_UMPIRE
+   static void SetUmpireDeviceAllocatorName(const char * d_name) { d_umpire_name = d_name; }
+   /// Set the device Umpire allocator name used with MemoryType::DEVICE_UMPIRE_2
+   static void SetUmpireDevice2AllocatorName(const char * d_name) { d_umpire_2_name = d_name; }
 
-   /// Get the host and device Umpire allocator ids
-   static int GetUmpireHostAllocatorId() { return h_umpire_id; }
-   static int GetUmpireDeviceAllocatorId() { return d_umpire_id; }
-   static int GetUmpireDeviceTempAllocatorId() { return d_umpire_temp_id; }
+   /// Get the host Umpire allocator name used with MemoryType::HOST_UMPIRE
+   static const char * GetUmpireHostAllocatorName() { return h_umpire_name; }
+   /// Get the device Umpire allocator name used with MemoryType::DEVICE_UMPIRE
+   static const char * GetUmpireDeviceAllocatorName() { return d_umpire_name; }
+   /// Get the device Umpire allocator name used with MemoryType::DEVICE_UMPIRE_2
+   static const char * GetUmpireDevice2AllocatorName() { return d_umpire_2_name; }
 #endif
 
    /// Free all the device memories
@@ -692,7 +785,6 @@ public:
 
    static MemoryType GetHostMemoryType() { return host_mem_type; }
    static MemoryType GetDeviceMemoryType() { return device_mem_type; }
-   static MemoryType GetDeviceTempMemoryType() { return device_temp_mem_type; }
 };
 
 
@@ -702,7 +794,7 @@ template <typename T>
 inline void Memory<T>::Reset()
 {
    h_ptr = NULL;
-   h_mt = MemoryManager::host_mem_type;
+   h_mt = MemoryManager::GetHostMemoryType();
    capacity = 0;
    flags = 0;
 }
@@ -721,8 +813,8 @@ inline void Memory<T>::New(int size)
 {
    capacity = size;
    flags = OWNS_HOST | VALID_HOST;
-   h_mt = MemoryManager::host_mem_type;
-   h_ptr = (h_mt == MemoryType::HOST) ? Alloc<new_align_bytes>::New(size) :
+   h_mt = MemoryManager::GetHostMemoryType();
+   h_ptr = (h_mt == MemoryType::HOST) ? NewHOST(size) :
            (T*)MemoryManager::New_(nullptr, size*sizeof(T), h_mt, flags);
 }
 
@@ -733,12 +825,21 @@ inline void Memory<T>::New(int size, MemoryType mt)
    const size_t bytes = size*sizeof(T);
    const bool mt_host = mt == MemoryType::HOST;
    if (mt_host) { flags = OWNS_HOST | VALID_HOST; }
-   h_mt = IsHostMemory(mt) ? mt : MemoryManager::GetDualMemoryType_(mt);
-   T *h_tmp = (h_mt == MemoryType::HOST) ?
-              Alloc<new_align_bytes>::New(size) : nullptr;
+   h_mt = IsHostMemory(mt) ? mt : MemoryManager::GetDualMemoryType(mt);
+   T *h_tmp = (h_mt == MemoryType::HOST) ? NewHOST(size) : nullptr;
    h_ptr = (mt_host) ? h_tmp : (T*)MemoryManager::New_(h_tmp, bytes, mt, flags);
 }
 
+template <typename T>
+inline void Memory<T>::New(int size, MemoryType h_mt, MemoryType d_mt)
+{
+   capacity = size;
+   const size_t bytes = size*sizeof(T);
+   this->h_mt = h_mt;
+   T *h_tmp = (h_mt == MemoryType::HOST) ? NewHOST(size) : nullptr;
+   h_ptr = (T*)MemoryManager::New_(h_tmp, bytes, h_mt, d_mt, VALID_HOST, flags);
+}
+
 template <typename T>
 inline void Memory<T>::Wrap(T *ptr, int size, bool own)
 {
@@ -746,7 +847,7 @@ inline void Memory<T>::Wrap(T *ptr, int size, bool own)
    capacity = size;
    const size_t bytes = size*sizeof(T);
    flags = (own ? OWNS_HOST : 0) | VALID_HOST;
-   h_mt = MemoryManager::host_mem_type;
+   h_mt = MemoryManager::GetHostMemoryType();
 #ifdef MFEM_DEBUG
    if (own && MemoryManager::Exists())
    { MFEM_VERIFY(h_mt == MemoryManager::GetHostMemoryType_(h_ptr),""); }
@@ -765,14 +866,14 @@ inline void Memory<T>::Wrap(T *ptr, int size, MemoryType mt, bool own)
       h_ptr = ptr;
       if (mt == MemoryType::HOST || !own)
       {
-         // Skip restration
+         // Skip registration
          flags = (own ? OWNS_HOST : 0) | VALID_HOST;
          return;
       }
    }
    else
    {
-      h_mt = MemoryManager::GetDualMemoryType_(mt);
+      h_mt = MemoryManager::GetDualMemoryType(mt);
       h_ptr = (h_mt == MemoryType::HOST) ? new T[size] : nullptr;
    }
    flags = 0;
@@ -789,8 +890,8 @@ inline void Memory<T>::Wrap(T *ptr, T *d_ptr, int size, MemoryType mt, bool own)
    capacity = size;
    MFEM_ASSERT(IsHostMemory(h_mt),"");
    const size_t bytes = size*sizeof(T);
-   const MemoryType d_mt = MemoryManager::GetDualMemoryType_(h_mt);
-   MemoryManager::Register_(d_ptr, h_ptr, bytes, d_mt, own, false, flags);
+   const MemoryType d_mt = MemoryManager::GetDualMemoryType(h_mt);
+   MemoryManager::Register_(h_ptr, d_ptr, bytes, h_mt, d_mt, own, false, flags);
 }
 
 template <typename T>
@@ -809,6 +910,18 @@ inline void Memory<T>::MakeAlias(const Memory &base, int offset, int size)
    }
 }
 
+template <typename T>
+inline void Memory<T>::SetDeviceMemoryType(MemoryType d_mt)
+{
+   if (!IsDeviceMemory(d_mt)) { return; }
+   if (!(flags & REGISTERED))
+   {
+      MemoryManager::Register_(h_ptr, nullptr, capacity*sizeof(T), h_mt,
+                               flags & OWNS_HOST, flags & ALIAS, flags);
+   }
+   MemoryManager::SetDeviceMemoryType_(h_ptr, flags, d_mt);
+}
+
 template <typename T>
 inline void Memory<T>::Delete()
 {
@@ -824,12 +937,11 @@ inline void Memory<T>::Delete()
 }
 
 template <typename T>
-inline void Memory<T>::DeleteDevice()
+inline void Memory<T>::DeleteDevice(bool copy_to_host)
 {
-   const bool registered = flags & REGISTERED;
-
-   if (registered)
+   if (flags & REGISTERED)
    {
+      if (copy_to_host) { Read(MemoryClass::HOST, capacity); }
       MemoryManager::DeleteDevice_((void*)h_ptr, flags);
    }
 }
@@ -961,7 +1073,7 @@ inline void Memory<T>::CopyFrom(const Memory &src, int size)
    {
       if (h_ptr != src.h_ptr && size != 0)
       {
-         MFEM_ASSERT(h_ptr + size <= src || src + size <= h_ptr,
+         MFEM_ASSERT(h_ptr + size <= src.h_ptr || src.h_ptr + size <= h_ptr,
                      "data overlaps!");
          std::memcpy(h_ptr, src, size*sizeof(T));
       }

mesh/mesh.cpp

@@ -755,7 +755,9 @@ void Mesh::GetLocalQuadToWdgTransformation(
 
 const GeometricFactors* Mesh::GetGeometricFactors(const IntegrationRule& ir,
                                                   const int flags,
-                                                  mfem::DofToQuad::Mode mode)
+                                                  mfem::DofToQuad::Mode mode,
+                                                  MemoryType d_mt
+                                               )
 {
    for (int i = 0; i < geom_factors.Size(); i++)
    {
@@ -767,7 +769,7 @@ const GeometricFactors* Mesh::GetGeometricFactors(const IntegrationRule& ir,
    }
 
    this->EnsureNodes();
-   GeometricFactors *gf = new GeometricFactors(this, ir, flags, mode);
+   GeometricFactors *gf = new GeometricFactors(this, ir, flags, mode, d_mt);
    geom_factors.Append(gf);
    return gf;
 }
@@ -10512,7 +10514,7 @@ int Mesh::FindPoints(DenseMatrix &point_mat, Array<int>& elem_ids,
 
 
 GeometricFactors::GeometricFactors(const Mesh *mesh, const IntegrationRule &ir,
-                                   int flags, DofToQuad::Mode mode)
+                                   int flags, DofToQuad::Mode mode, MemoryType d_mt)
 {
    this->mesh = mesh;
    IntRule = &ir;
@@ -10528,19 +10530,21 @@ GeometricFactors::GeometricFactors(const Mesh *mesh, const IntegrationRule &ir,
    const int NQ   = ir.GetNPoints();
 
    unsigned eval_flags = 0;
+   MemoryType my_d_mt = (d_mt != MemoryType::DEFAULT) ? d_mt :
+                         Device::GetDeviceMemoryType();
    if (flags & GeometricFactors::COORDINATES)
    {
-      X.SetSize(vdim*NQ*NE, Device::GetDeviceTempMemoryType());
+      X.SetSize(vdim*NQ*NE, my_d_mt);
       eval_flags |= QuadratureInterpolator::VALUES;
    }
    if (flags & GeometricFactors::JACOBIANS)
    {
-      J.SetSize(dim*vdim*NQ*NE, Device::GetDeviceTempMemoryType());
+      J.SetSize(dim*vdim*NQ*NE, my_d_mt);
       eval_flags |= QuadratureInterpolator::DERIVATIVES;
    }
    if (flags & GeometricFactors::DETERMINANTS)
    {
-      detJ.SetSize(NQ*NE, Device::GetDeviceTempMemoryType());
+      detJ.SetSize(NQ*NE, my_d_mt);
       eval_flags |= QuadratureInterpolator::DETERMINANTS;
    }
 
@@ -10559,7 +10563,7 @@ GeometricFactors::GeometricFactors(const Mesh *mesh, const IntegrationRule &ir,
 
    if (elem_restr)
    {
-      Vector Enodes(vdim*ND*NE, Device::GetDeviceTempMemoryType());
+      Vector Enodes(vdim*ND*NE, my_d_mt);
       elem_restr->Mult(*nodes, Enodes);
       qi->Mult(Enodes, eval_flags, X, J, detJ);
    }
@@ -10587,19 +10591,20 @@ GeometricFactors::GeometricFactors(const GridFunction *nodes_,
    const int NQ   = ir.GetNPoints();
 
    unsigned eval_flags = 0;
+   MemoryType d_mt = Device::GetDeviceMemoryType();
    if (flags & GeometricFactors::COORDINATES)
    {
-      X.SetSize(vdim*NQ*NE, Device::GetDeviceTempMemoryType());
+      X.SetSize(vdim*NQ*NE, d_mt);
       eval_flags |= QuadratureInterpolator::VALUES;
    }
    if (flags & GeometricFactors::JACOBIANS)
    {
-      J.SetSize(dim*vdim*NQ*NE, Device::GetDeviceTempMemoryType());
+      J.SetSize(dim*vdim*NQ*NE, d_mt);
       eval_flags |= QuadratureInterpolator::DERIVATIVES;
    }
    if (flags & GeometricFactors::DETERMINANTS)
    {
-      detJ.SetSize(NQ*NE, Device::GetDeviceTempMemoryType());
+      detJ.SetSize(NQ*NE, d_mt);
       eval_flags |= QuadratureInterpolator::DETERMINANTS;
    }
 
@@ -10618,7 +10623,7 @@ GeometricFactors::GeometricFactors(const GridFunction *nodes_,
 
    if (elem_restr)
    {
-      Vector Enodes(vdim*ND*NE, Device::GetDeviceTempMemoryType());
+      Vector Enodes(vdim*ND*NE, d_mt);
       elem_restr->Mult(*nodes, Enodes);
       qi->Mult(Enodes, eval_flags, X, J, detJ);
    }

mesh/mesh.hpp

@@ -743,7 +743,8 @@ public:
        integration rule. */
    const GeometricFactors* GetGeometricFactors(const IntegrationRule& ir,
                                                const int flags,
-                                               DofToQuad::Mode = DofToQuad::FULL);
+                                               DofToQuad::Mode = DofToQuad::FULL,
+                                               MemoryType d_mt = MemoryType::DEFAULT);
 
    /** @brief Return the mesh geometric factors for the faces corresponding
         to the given integration rule. */
@@ -1374,7 +1375,8 @@ public:
    };
 
    GeometricFactors(const Mesh *mesh, const IntegrationRule &ir, int flags,
-                    DofToQuad::Mode = DofToQuad::FULL);
+                    DofToQuad::Mode = DofToQuad::FULL,
+                    MemoryType d_mt = MemoryType::DEFAULT);
 
    GeometricFactors(const GridFunction *nodes_, const IntegrationRule &ir,
                     int flags,