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: 6e69258b69]

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});

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;
 }

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
 }
 
 // ================================================================================================