Initial HMM support

- Expose ROCclr interfaces for HIP usage
- ROCr interfaces aren't available in staging, thus control the
build with AMD_HMM_SUPPORT define

Change-Id: Iadc2bcc230e78d3b0dc22b235189c8cc80843446


[ROCm/clr commit: c5afd5d412]

projects/clr/rocclr/device/device.hpp

@@ -76,6 +76,7 @@ class SvmCopyMemoryCommand;
 class SvmFillMemoryCommand;
 class SvmMapMemoryCommand;
 class SvmUnmapMemoryCommand;
+class SvmPrefetchAsyncCommand;
 class TransferBufferFileCommand;
 class HwDebugManager;
 class Device;
@@ -86,6 +87,30 @@ namespace option {
 class Options;
 }  // namespace option
 
+//! @note: the defines match hip values
+enum MemoryAdvice : uint32_t {
+  SetReadMostly = 1,          ///< Data will mostly be read and only occassionally be written to
+  UnsetReadMostly = 2,        ///< Undo the effect of hipMemAdviseSetReadMostly
+  SetPreferredLocation = 3,   ///< Set the preferred location for the data as the specified device
+  UnsetPreferredLocation = 4, ///< Clear the preferred location for the data
+  SetAccessedBy = 5,          ///< Data will be accessed by the specified device,
+                              ///< so prevent page faults as much as possible
+  UnsetAccessedBy = 6         ///< Let the Unified Memory subsystem decide on
+                              ///< the page faulting policy for the specified device
+};
+
+enum MemRangeAttribute : uint32_t {
+    ReadMostly = 1,           ///< Whether the range will mostly be read and only
+                              ///< occassionally be written to
+    PreferredLocation = 2,    ///< The preferred location of the range
+    AccessedBy = 3,           ///< Memory range has hipMemAdviseSetAccessedBy
+                              ///< set for specified device
+    LastPrefetchLocation = 4, ///< The last location to which the range was prefetched
+};
+
+constexpr int CpuDeviceId = static_cast<int>(-1);
+constexpr int InvalidDeviceId = static_cast<int>(-2);
+
 }  // namespace amd
 
 enum OclExtensions {
@@ -1111,7 +1136,9 @@ class VirtualDevice : public amd::HeapObject {
   virtual void submitTransferBufferFromFile(amd::TransferBufferFileCommand& cmd) {
     ShouldNotReachHere();
   }
-
+  virtual void submitSvmPrefetchAsync(amd::SvmPrefetchAsyncCommand& cmd) {
+    ShouldNotReachHere();
+  }
   //! Get the blit manager object
   device::BlitManager& blitMgr() const { return *blitMgr_; }
 
@@ -1341,6 +1368,24 @@ class Device : public RuntimeObject {
    */
   virtual void svmFree(void* ptr) const = 0;
 
+  /**
+   * @return True if the device successfully applied the SVM attributes in HMM for device memory
+   */
+  virtual bool SetSvmAttributes(const void* dev_ptr, size_t count,
+                                amd::MemoryAdvice advice, bool first_alloc = false) const {
+    ShouldNotCallThis();
+    return false;
+  }
+
+  /**
+   * @return True if the device successfully retrieved the SVM attributes from HMM for device memory
+   */
+  virtual bool GetSvmAttributes(void** data, size_t* data_sizes, int* attributes,
+                                size_t num_attributes, const void* dev_ptr, size_t count) const {
+    ShouldNotCallThis();
+    return false;
+  }
+
   //! Validate kernel
   virtual bool validateKernel(const amd::Kernel& kernel,
                               const device::VirtualDevice* vdev,

projects/clr/rocclr/device/rocm/CMakeLists.txt

@@ -36,6 +36,11 @@ target_include_directories(oclrocm
     ${ROCM_OCL_INCLUDES}
     ${ROCR_INCLUDES})
 
+option(BUILD_HMM "Build HMM support" OFF)
+if (BUILD_HMM)
+  target_compile_definitions(oclrocm
+    PRIVATE AMD_HMM_SUPPORT)
+endif()
 
 if(USE_COMGR_LIBRARY)
   target_compile_definitions(oclrocm

projects/clr/rocclr/device/rocm/rocdevice.cpp

@@ -178,6 +178,7 @@ Device::Device(hsa_agent_t bkendDevice)
   system_coarse_segment_.handle = 0;
   gpuvm_segment_.handle = 0;
   gpu_fine_grained_segment_.handle = 0;
+  prefetch_signal_.handle = 0;
 }
 
 void Device::setupCpuAgent() {
@@ -786,6 +787,11 @@ bool Device::create(bool sramEccEnabled) {
     }
   }
 
+  // Create signal for HMM prefetch operation on device
+  if (HSA_STATUS_SUCCESS != hsa_signal_create(InitSignalValue, 0, nullptr, &prefetch_signal_)) {
+    return false;
+  }
+
   return true;
 }
 
@@ -1904,6 +1910,187 @@ void* Device::svmAlloc(amd::Context& context, size_t size, size_t alignment, cl_
   return svmPtr;
 }
 
+// ================================================================================================
+bool Device::SetSvmAttributes(const void* dev_ptr, size_t count,
+                              amd::MemoryAdvice advice, bool first_alloc) const {
+  if ((settings().hmmFlags_ & Settings::Hmm::EnableSvmTracking) && !first_alloc) {
+    amd::Memory* svmMem = svmMem = amd::MemObjMap::FindMemObj(dev_ptr);
+    if (nullptr == svmMem) {
+      LogPrintfError("SetSvmAttributes received unknown memory for update: %p!", dev_ptr);
+      return false;
+    }
+  }
+#if AMD_HMM_SUPPORT
+  std::vector<hsa_amd_svm_attribute_pair_t> attr;
+  if (first_alloc) {
+    attr.push_back({HSA_AMD_SVM_ATTRIB_GLOBAL_FLAG, HSA_AMD_SVM_GLOBAL_FLAG_COARSE_GRAINED});
+  }
+
+  switch (advice) {
+    case amd::MemoryAdvice::SetReadMostly:
+      attr.push_back({HSA_AMD_SVM_ATTRIB_READ_ONLY, true});
+      break;
+    case amd::MemoryAdvice::UnsetReadMostly:
+      attr.push_back({HSA_AMD_SVM_ATTRIB_READ_ONLY, false});
+      break;
+    case amd::MemoryAdvice::SetPreferredLocation:
+      attr.push_back({HSA_AMD_SVM_ATTRIB_PREFERRED_LOCATION, getBackendDevice().handle});
+      break;
+    case amd::MemoryAdvice::UnsetPreferredLocation:
+      // Note: The current behavior doesn't match hip spec precisely
+      attr.push_back({HSA_AMD_SVM_ATTRIB_PREFERRED_LOCATION, getCpuAgent().handle});
+      break;
+    case amd::MemoryAdvice::SetAccessedBy:
+      attr.push_back({HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE, getBackendDevice().handle});
+      break;
+    case amd::MemoryAdvice::UnsetAccessedBy:
+      // @note: The current behavior doesn't match hip spec precisely
+      attr.push_back({HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE, getCpuAgent().handle});
+      break;
+    default:
+      return false;
+    break;
+  }
+
+  hsa_status_t status = hsa_amd_svm_attributes_set(const_cast<void*>(dev_ptr), count,
+                                                   attr.data(), attr.size());
+  if (status != HSA_STATUS_SUCCESS) {
+    LogError("hsa_amd_svm_attributes_set() failed");
+    return false;
+  }
+#endif // AMD_HMM_SUPPORT
+  return true;
+}
+
+// ================================================================================================
+bool Device::GetSvmAttributes(void** data, size_t* data_sizes, int* attributes,
+                              size_t num_attributes, const void* dev_ptr, size_t count) const {
+  if (settings().hmmFlags_ & Settings::Hmm::EnableSvmTracking) {
+    amd::Memory* svmMem = svmMem = amd::MemObjMap::FindMemObj(dev_ptr);
+    if (nullptr == svmMem) {
+      LogPrintfError("GetSvmAttributes received unknown memory %p for state!", dev_ptr);
+      return false;
+    }
+  }
+#if AMD_HMM_SUPPORT
+  std::vector<hsa_amd_svm_attribute_pair_t> attr;
+
+  for (int i = 0; i < num_attributes; ++i) {
+    switch (attributes[i]) {
+      case amd::MemRangeAttribute::ReadMostly:
+        attr.push_back({HSA_AMD_SVM_ATTRIB_READ_ONLY, 0});
+        break;
+      case amd::MemRangeAttribute::PreferredLocation:
+        attr.push_back({HSA_AMD_SVM_ATTRIB_PREFERRED_LOCATION, 0});
+        break;
+      case amd::MemRangeAttribute::AccessedBy:
+        attr.push_back({HSA_AMD_SVM_ATTRIB_ACCESS_QUERY, 0});
+        break;
+      case amd::MemRangeAttribute::LastPrefetchLocation:
+        attr.push_back({HSA_AMD_SVM_ATTRIB_PREFETCH_LOCATION, 0});
+        break;
+      default:
+        return false;
+      break;
+    }
+  }
+
+  hsa_status_t status = hsa_amd_svm_attributes_get(const_cast<void*>(dev_ptr), count,
+                                                   attr.data(), attr.size());
+  if (status != HSA_STATUS_SUCCESS) {
+    LogError("hsa_amd_svm_attributes_get() failed");
+    return false;
+  }
+
+  uint32_t idx = 0;
+  for (auto& it : attr) {
+    switch (it.attribute) {
+      case HSA_AMD_SVM_ATTRIB_READ_ONLY:
+        if (data_sizes[idx] != sizeof(uint32_t)) {
+          return false;
+        }
+        // Cast ROCr value into the hip format
+        *reinterpret_cast<uint32_t*>(data[idx]) = static_cast<uint32_t>(it.value);
+        break;
+      // The logic should be identical for the both queries
+      case HSA_AMD_SVM_ATTRIB_PREFERRED_LOCATION:
+      case HSA_AMD_SVM_ATTRIB_PREFETCH_LOCATION:
+        if (data_sizes[idx] != sizeof(uint32_t)) {
+          return false;
+        }
+        *reinterpret_cast<int32_t*>(data[idx]) = static_cast<int32_t>(amd::InvalidDeviceId);
+        // Find device agent returned by ROCr
+        for (auto& device : devices()) {
+          if (static_cast<Device*>(device)->getBackendDevice().handle == it.value) {
+            *reinterpret_cast<uint32_t*>(data[idx]) = static_cast<uint32_t>(device->index());
+          }
+        }
+        // Find CPU agent returned by ROCr
+        for (auto& agent_info : getCpuAgents()) {
+          if (agent_info.agent.handle == it.value) {
+            *reinterpret_cast<int32_t*>(data[idx]) = static_cast<int32_t>(amd::CpuDeviceId);
+          }
+        }
+        break;
+      case HSA_AMD_SVM_ATTRIB_ACCESS_QUERY:
+        // Make sure it's multiple of 4
+        if (data_sizes[idx] % 4 != 0) {
+          return false;
+        }
+        // @note currently it's a nop
+        break;
+      default:
+        return false;
+      break;
+    }
+    // Find the next location in the query
+    ++idx;
+  }
+#endif // AMD_HMM_SUPPORT
+  return true;
+}
+
+// ================================================================================================
+bool Device::SvmAllocInit(void* memory, size_t size) const {
+  amd::MemoryAdvice advice = amd::MemoryAdvice::SetAccessedBy;
+  constexpr bool kFirstAlloc = true;
+  SetSvmAttributes(memory, size, advice, kFirstAlloc);
+
+  if (settings().hmmFlags_ & Settings::Hmm::EnableSystemMemory) {
+    advice = amd::MemoryAdvice::UnsetPreferredLocation;
+    SetSvmAttributes(memory, size, advice);
+  } else {
+    advice = amd::MemoryAdvice::SetPreferredLocation;
+    SetSvmAttributes(memory, size, advice);
+  }
+
+  if ((settings().hmmFlags_ & Settings::Hmm::EnableMallocPrefetch) == 0) {
+    return true;
+  }
+
+#if AMD_HMM_SUPPORT
+  // Initialize signal for the barrier
+  hsa_signal_store_relaxed(prefetch_signal_, InitSignalValue);
+
+  // Initiate a prefetch command which should force memory update in HMM
+  hsa_status_t status = hsa_amd_svm_prefetch_async(memory, size, getBackendDevice(),
+                                                   0, nullptr, prefetch_signal_);
+  if (status != HSA_STATUS_SUCCESS) {
+    LogError("hsa_amd_svm_attributes_get() failed");
+    return false;
+  }
+
+  // Wait for the prefetch
+  if (hsa_signal_wait_acquire(prefetch_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1),
+                              HSA_WAIT_STATE_BLOCKED) != 0) {
+    LogError("Barrier packet submission failed");
+    return false;
+  }
+#endif // AMD_HMM_SUPPORT
+  return true;
+}
+
+// ================================================================================================
 void Device::svmFree(void* ptr) const {
   amd::Memory* svmMem = nullptr;
   svmMem = amd::MemObjMap::FindMemObj(ptr);

projects/clr/rocclr/device/rocm/rocdevice.hpp

@@ -79,7 +79,7 @@ class IProDevice;
 class Sampler : public device::Sampler {
  public:
   //! Constructor
-    Sampler(const Device& dev) : dev_(dev) {}
+  Sampler(const Device& dev) : dev_(dev) {}
 
   //! Default destructor for the device memory object
   virtual ~Sampler();
@@ -381,7 +381,13 @@ class Device : public NullDevice {
 
   virtual void svmFree(void* ptr) const;
 
-  virtual bool SetClockMode(const cl_set_device_clock_mode_input_amd setClockModeInput, cl_set_device_clock_mode_output_amd* pSetClockModeOutput);
+  virtual bool SetSvmAttributes(const void* dev_ptr, size_t count,
+                                amd::MemoryAdvice advice, bool first_alloc = false) const;
+  virtual bool GetSvmAttributes(void** data, size_t* data_sizes, int* attributes,
+                                size_t num_attributes, const void* dev_ptr, size_t count) const;
+
+  virtual bool SetClockMode(const cl_set_device_clock_mode_input_amd setClockModeInput,
+                            cl_set_device_clock_mode_output_amd* pSetClockModeOutput);
 
   //! Returns transfer engine object
   const device::BlitManager& xferMgr() const { return xferQueue()->blitMgr(); }
@@ -464,7 +470,12 @@ class Device : public NullDevice {
   roc::Memory* getGpuMemory(amd::Memory* mem  //!< Pointer to AMD memory object
                             ) const;
 
+  //! Initialize memory in AMD HMM on the current device or keeps it in the host memory
+  bool SvmAllocInit(void* memory, size_t size) const;
+
  private:
+  static const hsa_signal_value_t InitSignalValue = 1;
+
   static hsa_ven_amd_loader_1_00_pfn_t amd_loader_ext_table;
 
   amd::Monitor* mapCacheOps_;            //!< Lock to serialise cache for the map resources
@@ -485,6 +496,8 @@ class Device : public NullDevice {
   hsa_amd_memory_pool_t system_coarse_segment_;
   hsa_amd_memory_pool_t gpuvm_segment_;
   hsa_amd_memory_pool_t gpu_fine_grained_segment_;
+  hsa_signal_t prefetch_signal_;    //!< Prefetch signal, used to explicitly prefetch SVM on device
+
   size_t gpuvm_segment_max_alloc_;
   size_t alloc_granularity_;
   static const bool offlineDevice_;

projects/clr/rocclr/device/rocm/rocmemory.cpp

@@ -42,7 +42,7 @@
 
 namespace roc {
 
-/////////////////////////////////roc::Memory//////////////////////////////
+// ======================================= roc::Memory ============================================
 Memory::Memory(const roc::Device& dev, amd::Memory& owner)
     : device::Memory(owner),
       dev_(dev),
@@ -620,8 +620,7 @@ void Memory::mgpuCacheWriteBack() {
   }
 }
 
-/////////////////////////////////roc::Buffer//////////////////////////////
-
+// ==================================== roc::Buffer ===============================================
 Buffer::Buffer(const roc::Device& dev, amd::Memory& owner) : roc::Memory(dev, owner) {}
 
 Buffer::Buffer(const roc::Device& dev, size_t size) : roc::Memory(dev, size) {}
@@ -634,6 +633,7 @@ Buffer::~Buffer() {
   }
 }
 
+// ================================================================================================
 void Buffer::destroy() {
   if (owner()->parent() != nullptr) {
     return;
@@ -647,15 +647,24 @@ void Buffer::destroy() {
   cl_mem_flags memFlags = owner()->getMemFlags();
 
   if (owner()->getSvmPtr() != nullptr) {
-    if (dev().forceFineGrain(owner()) ||
-        dev().isFineGrainedSystem(true)) {
+    if (dev().forceFineGrain(owner()) || dev().isFineGrainedSystem(true)) {
       memFlags |= CL_MEM_SVM_FINE_GRAIN_BUFFER;
     }
     const bool isFineGrain = memFlags & CL_MEM_SVM_FINE_GRAIN_BUFFER;
 
     if (kind_ != MEMORY_KIND_PTRGIVEN) {
       if (isFineGrain) {
-        dev().hostFree(deviceMemory_, size());
+        if (memFlags & CL_MEM_ALLOC_HOST_PTR) {
+  #if AMD_HMM_SUPPORT
+          // AMD HMM path. Destroy system memory
+          amd::Os::uncommitMemory(deviceMemory_, size());
+          amd::Os::releaseMemory(deviceMemory_, size());
+  #else
+          dev().hostFree(deviceMemory_, size());;
+  #endif // AMD_HMM_SUPPORT
+        } else {
+          dev().hostFree(deviceMemory_, size());
+        }
       } else {
         dev().memFree(deviceMemory_, size());
       }
@@ -705,6 +714,7 @@ void Buffer::destroy() {
   }
 }
 
+// ================================================================================================
 bool Buffer::create() {
   if (owner() == nullptr) {
     deviceMemory_ = dev().hostAlloc(size(), 1, false);
@@ -731,7 +741,19 @@ bool Buffer::create() {
 
     if (owner()->getSvmPtr() == reinterpret_cast<void*>(1)) {
       if (isFineGrain) {
-        if (memFlags & CL_MEM_SVM_ATOMICS) {
+        if (memFlags & CL_MEM_ALLOC_HOST_PTR) {
+#if AMD_HMM_SUPPORT
+          // AMD HMM path. Just allocate system memory and KFD will manage it
+          deviceMemory_ =  amd::Os::reserveMemory(
+              0, size(), amd::Os::pageSize(), amd::Os::MEM_PROT_RW);
+          amd::Os::commitMemory(deviceMemory_, size(), amd::Os::MEM_PROT_RW);
+          // Currently HMM requires cirtain initial calls to mark sysmem allocation as
+          // GPU accessible or prefetch memory into GPU
+          dev().SvmAllocInit(deviceMemory_, size());
+#else
+          deviceMemory_ = dev().hostAlloc(size(), 1, false);
+#endif // AMD_HMM_SUPPORT
+        } else if (memFlags & CL_MEM_SVM_ATOMICS) {
           deviceMemory_ = dev().hostAlloc(size(), 1, true);
         }
         else {
@@ -745,6 +767,14 @@ bool Buffer::create() {
     } else {
       deviceMemory_ = owner()->getSvmPtr();
       kind_ = MEMORY_KIND_PTRGIVEN;
+#if AMD_HMM_SUPPORT
+      if (memFlags & CL_MEM_ALLOC_HOST_PTR) {
+        // Currently HMM requires cirtain initial calls to mark sysmem allocation as
+        // GPU accessible or prefetch memory into the current device
+        // @note: Skip any allocaiton here, since sysmem was allocated on another device.
+        dev().SvmAllocInit(deviceMemory_, size());
+      }
+#endif // AMD_HMM_SUPPORT
     }
 
     if (!isFineGrain && (owner()->parent() != nullptr) &&
@@ -870,9 +900,10 @@ bool Buffer::create() {
 
   if (owner()->getSvmPtr() != owner()->getHostMem()) {
     if (memFlags & (CL_MEM_USE_HOST_PTR | CL_MEM_ALLOC_HOST_PTR)) {
-      hsa_amd_memory_pool_t pool = (memFlags & CL_MEM_SVM_ATOMICS)? dev().SystemSegment() : dev().SystemCoarseSegment();
-      hsa_status_t status = hsa_amd_memory_lock_to_pool(owner()->getHostMem(), owner()->getSize(), nullptr,
-                                                0, pool, 0, &deviceMemory_);
+      hsa_amd_memory_pool_t pool = (memFlags & CL_MEM_SVM_ATOMICS) ?
+                                    dev().SystemSegment() : dev().SystemCoarseSegment();
+      hsa_status_t status = hsa_amd_memory_lock_to_pool(owner()->getHostMem(),
+          owner()->getSize(), nullptr, 0, pool, 0, &deviceMemory_);
       if (status != HSA_STATUS_SUCCESS) {
         DevLogPrintfError("Failed to lock memory to pool, failed with hsa_status: %d \n", status);
         deviceMemory_ = nullptr;
@@ -887,7 +918,7 @@ bool Buffer::create() {
   return deviceMemory_ != nullptr;
 }
 
-/////////////////////////////////roc::Image//////////////////////////////
+// ======================================= roc::Image =============================================
 typedef struct ChannelOrderMap {
   uint32_t cl_channel_order;
   hsa_ext_image_channel_order_t hsa_channel_order;

projects/clr/rocclr/device/rocm/rocsettings.cpp

@@ -86,6 +86,9 @@ Settings::Settings() {
 
   lcWavefrontSize64_ = true;
   imageBufferWar_ = false;
+
+  hmmFlags_ = (!flagIsDefault(ROC_HMM_FLAGS)) ? ROC_HMM_FLAGS :
+      Hmm::EnableSystemMemory | Hmm::EnableMallocPrefetch;
 }
 
 bool Settings::create(bool fullProfile, int gfxipMajor, int gfxipMinor, bool coop_groups) {

projects/clr/rocclr/device/rocm/rocsettings.hpp

@@ -34,6 +34,13 @@ namespace roc {
 //! Device settings
 class Settings : public device::Settings {
  public:
+  enum Hmm : uint32_t {
+    EnableSystemMemory    = 0x01, //!< Forces system memory preference by default
+    EnableMallocPrefetch  = 0x02, //!< Skips default prefetch after allocation
+    EnableSvmTracking     = 0x04, //!< Enables SW SVM tracking
+    EnableDebugSvm        = 0x08  //!< Extra debug flag (reserved for runtime developers)
+  };
+
   union {
     struct {
       uint doublePrecision_ : 1;        //!< Enables double precision support
@@ -75,6 +82,8 @@ class Settings : public device::Settings {
 
   size_t sdmaCopyThreshold_;  //!< Use SDMA to copy above this size
 
+  uint32_t      hmmFlags_;    //!< HMM functionality control flags
+
   //! Default constructor
   Settings();
 

projects/clr/rocclr/device/rocm/rocvirtual.cpp

@@ -1170,6 +1170,7 @@ void VirtualGPU::submitWriteMemory(amd::WriteMemoryCommand& cmd) {
   profilingEnd(cmd);
 }
 
+// ================================================================================================
 void VirtualGPU::submitSvmFreeMemory(amd::SvmFreeMemoryCommand& cmd) {
   // Make sure VirtualGPU has an exclusive access to the resources
   amd::ScopedLock lock(execution());
@@ -1191,6 +1192,31 @@ void VirtualGPU::submitSvmFreeMemory(amd::SvmFreeMemoryCommand& cmd) {
   profilingEnd(cmd);
 }
 
+// ================================================================================================
+void VirtualGPU::submitSvmPrefetchAsync(amd::SvmPrefetchAsyncCommand& cmd) {
+#if ROCR_HMM_SUPPORT
+  // Initialize signal for the barrier
+  hsa_signal_store_relaxed(barrier_signal_, InitSignalValue);
+
+  // Find the requested agent for the transfer
+  hsa_agent_t agent = cmd.cpu_access() ? dev().cpu_agent() ? gpu_device();
+
+  // Initiate a prefetch command
+  hsa_amd_svm_prefetch_async(cmd.dev_prt(), cmd.count(), agent, 0, nullptr, barrier_signal_);
+
+  // Wait for the prefetch
+  if (hsa_signal_wait_acquire(barrier_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1),
+                              HSA_WAIT_STATE_BLOCKED) != 0) {
+    LogError("Barrier packet submission failed");
+    return false;
+  }
+
+  // Add system scope, since the prefetch scope is unclear
+  addSystemScope();
+#endif // ROCR_HMM_SUPPORT
+}
+
+// ================================================================================================
 bool VirtualGPU::copyMemory(cl_command_type type, amd::Memory& srcMem, amd::Memory& dstMem,
                             bool entire, const amd::Coord3D& srcOrigin,
                             const amd::Coord3D& dstOrigin, const amd::Coord3D& size,

projects/clr/rocclr/device/rocm/rocvirtual.hpp

@@ -208,6 +208,7 @@ class VirtualGPU : public device::VirtualDevice {
   void submitSvmFillMemory(amd::SvmFillMemoryCommand& cmd);
   void submitSvmMapMemory(amd::SvmMapMemoryCommand& cmd);
   void submitSvmUnmapMemory(amd::SvmUnmapMemoryCommand& cmd);
+  void submitSvmPrefetchAsync(amd::SvmPrefetchAsyncCommand& cmd);
 
   // { roc OpenCL integration
   // Added these stub (no-ops) implementation of pure virtual methods,

projects/clr/rocclr/platform/command.cpp

@@ -203,7 +203,6 @@ bool Event::awaitCompletion() {
         lock_.wait();
       }
     }
-
     ClPrint(LOG_DEBUG, LOG_WAIT, "event %p wait completed", this);
   }
 
@@ -633,4 +632,13 @@ bool CopyMemoryP2PCommand::validateMemory() {
   return true;
 }
 
+// ================================================================================================
+bool SvmPrefetchAsyncCommand::validateMemory() {
+  amd::Memory* svmMem = svmMem = amd::MemObjMap::FindMemObj(dev_ptr());
+  if (nullptr == svmMem) {
+    LogPrintfError("SvmPrefetchAsync received unknown memory for prefetch: %p!", dev_ptr());
+    return false;
+  }
+}
+
 }  // namespace amd

projects/clr/rocclr/platform/command.hpp

@@ -1364,7 +1364,6 @@ class TransferBufferFileCommand : public OneMemoryArgCommand {
  *              as 1D structures so origin_[0] and size_[0] are
  *              equivalent to offset_ and count_ respectively.
  */
-
 class CopyMemoryP2PCommand : public CopyMemoryCommand {
  public:
   CopyMemoryP2PCommand(HostQueue& queue, cl_command_type cmdType,
@@ -1378,6 +1377,29 @@ class CopyMemoryP2PCommand : public CopyMemoryCommand {
   bool validateMemory();
 };
 
+/*! \brief      Prefetch command for SVM memory
+ *
+ *  \details    Prefetches SVM memory into the current device or CPU
+ */
+class SvmPrefetchAsyncCommand : public Command {
+  const void* dev_ptr_;   //!< Device pointer to memory for prefetch
+  size_t count_;          //!< the size for prefetch
+  bool   cpu_access_;     //!< Prefetch data into CPU location
+
+ public:
+  SvmPrefetchAsyncCommand(HostQueue& queue, const EventWaitList& eventWaitList,
+                          const void* dev_ptr, size_t count, bool cpu_access)
+      : Command(queue, 1, eventWaitList), dev_ptr_(dev_ptr), count_(count), cpu_access_(cpu_access) {}
+
+  virtual void submit(device::VirtualDevice& device) { device.submitSvmPrefetchAsync(*this); }
+
+  bool validateMemory();
+
+  const void* dev_ptr() const { return dev_ptr_; }
+  size_t count() const { return count_; }
+  size_t cpu_access() const { return cpu_access_; }
+};
+
 /*! @}
  *  @}
  */

projects/clr/rocclr/utils/flags.hpp

@@ -56,6 +56,8 @@ debug(size_t, PARAMETERS_MIN_ALIGNMENT, 16,                                   \
         "Minimum alignment required for the abstract parameters stack")       \
 debug(size_t, MEMOBJ_BASE_ADDR_ALIGN, 4*Ki,                                   \
         "Alignment of the base address of any allocate memory object")        \
+release(uint, ROC_HMM_FLAGS, 0,                                               \
+        "ROCm HMM configuration flags")                                       \
 release(cstring, GPU_DEVICE_ORDINAL, "",                                      \
         "Select the device ordinal (comma seperated list of available devices)") \
 release(bool, REMOTE_ALLOC, false,                                            \