From a5ca15a59929d0335fe01efd54deca3f9ae992ee Mon Sep 17 00:00:00 2001
From: German Andryeyev <German.Andryeyev@amd.com>
Date: Mon, 10 Aug 2020 15:22:14 -0400
Subject: [PATCH] Enable prefetch async functionality

Fix a typo with the name define, when compilation wasn't enabled.
Force CPU prefetch if system was forced in runtime

Change-Id: Id4b578f9fa44a45426fdb5d8ecb1da803aa42313


[ROCm/clr commit: 6e69258b6930b40d7959882d64856ec3226e7ae1]
---
 projects/clr/rocclr/device/rocm/rocdevice.cpp |  2 +-
 .../clr/rocclr/device/rocm/rocsettings.cpp    |  3 +--
 .../clr/rocclr/device/rocm/rocvirtual.cpp     | 20 +++++++++++--------
 3 files changed, 14 insertions(+), 11 deletions(-)

diff --git a/projects/clr/rocclr/device/rocm/rocdevice.cpp b/projects/clr/rocclr/device/rocm/rocdevice.cpp
index 9f8a476fcf..7fc422d8fc 100755
--- a/projects/clr/rocclr/device/rocm/rocdevice.cpp
+++ b/projects/clr/rocclr/device/rocm/rocdevice.cpp
@@ -2007,7 +2007,7 @@ bool Device::GetSvmAttributes(void** data, size_t* data_sizes, int* attributes,
 #if AMD_HMM_SUPPORT
   std::vector<hsa_amd_svm_attribute_pair_t> attr;
 
-  for (int i = 0; i < num_attributes; ++i) {
+  for (size_t i = 0; i < num_attributes; ++i) {
     switch (attributes[i]) {
       case amd::MemRangeAttribute::ReadMostly:
         attr.push_back({HSA_AMD_SVM_ATTRIB_READ_ONLY, 0});
diff --git a/projects/clr/rocclr/device/rocm/rocsettings.cpp b/projects/clr/rocclr/device/rocm/rocsettings.cpp
index 05c527a205..d123dec363 100644
--- a/projects/clr/rocclr/device/rocm/rocsettings.cpp
+++ b/projects/clr/rocclr/device/rocm/rocsettings.cpp
@@ -87,8 +87,7 @@ Settings::Settings() {
   lcWavefrontSize64_ = true;
   imageBufferWar_ = false;
 
-  hmmFlags_ = (!flagIsDefault(ROC_HMM_FLAGS)) ? ROC_HMM_FLAGS :
-      Hmm::EnableSystemMemory | Hmm::EnableMallocPrefetch;
+  hmmFlags_ = (!flagIsDefault(ROC_HMM_FLAGS)) ? ROC_HMM_FLAGS : Hmm::EnableSystemMemory;
 
   rocr_backend_ = true;
 }
diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp
index 55103e546b..b3a02aedda 100644
--- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp
+++ b/projects/clr/rocclr/device/rocm/rocvirtual.cpp
@@ -1206,26 +1206,30 @@ void VirtualGPU::submitSvmFreeMemory(amd::SvmFreeMemoryCommand& cmd) {
 
 // ================================================================================================
 void VirtualGPU::submitSvmPrefetchAsync(amd::SvmPrefetchAsyncCommand& cmd) {
-#if ROCR_HMM_SUPPORT
+#if AMD_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();
+  hsa_agent_t agent = (cmd.cpu_access() ||
+      (dev().settings().hmmFlags_ & Settings::Hmm::EnableSystemMemory)) ?
+      dev().getCpuAgent() : gpu_device();
 
   // Initiate a prefetch command
-  hsa_amd_svm_prefetch_async(cmd.dev_prt(), cmd.count(), agent, 0, nullptr, barrier_signal_);
+  hsa_status_t status = hsa_amd_svm_prefetch_async(
+      const_cast<void*>(cmd.dev_ptr()), cmd.count(), agent, 0, nullptr, barrier_signal_);
 
   // Wait for the prefetch
-  if (hsa_signal_wait_scacquire(barrier_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1),
-                              HSA_WAIT_STATE_BLOCKED) != 0) {
-    LogError("Barrier packet submission failed");
-    return false;
+  if ((status != HSA_STATUS_SUCCESS) ||
+      hsa_signal_wait_scacquire(barrier_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1),
+                                HSA_WAIT_STATE_BLOCKED) != 0) {
+    LogError("hsa_amd_svm_prefetch_async failed");
+    cmd.setStatus(CL_INVALID_OPERATION);
   }
 
   // Add system scope, since the prefetch scope is unclear
   addSystemScope();
-#endif // ROCR_HMM_SUPPORT
+#endif // AMD_HMM_SUPPORT
 }
 
 // ================================================================================================