From b65e44205e291c808419f6ade510e06bccbd3a29 Mon Sep 17 00:00:00 2001
From: foreman <dl.constructicon@amd.com>
Date: Thu, 23 Feb 2017 17:06:16 -0500
Subject: [PATCH] P4 to Git Change 1377072 by gandryey@gera-w8 on 2017/02/23
 16:54:28

	SWDEV-79445 - OCL generic changes and code clean-up
	- Switch wait state in HSA copy functions from HSA_WAIT_STATE_ACTIVE to  HSA_WAIT_STATE_BLOCKED

Affected files ...

... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocblit.cpp#13 edit


[ROCm/clr commit: 37a5e9a169749c3f711b746b0da6d3fa9ff3ea60]
---
 .../rocclr/runtime/device/rocm/rocblit.cpp    | 21 ++++++++++---------
 1 file changed, 11 insertions(+), 10 deletions(-)

diff --git a/projects/clr/rocclr/runtime/device/rocm/rocblit.cpp b/projects/clr/rocclr/runtime/device/rocm/rocblit.cpp
index f52a45bcda..bcbb98ae87 100644
--- a/projects/clr/rocclr/runtime/device/rocm/rocblit.cpp
+++ b/projects/clr/rocclr/runtime/device/rocm/rocblit.cpp
@@ -79,7 +79,7 @@ DmaBlitManager::readBuffer(
 
         // Check if a pinned transfer can be executed
         if (pinSize && (srcSize > MinSizeForPinnedTransfer)) {
-            // Allign offset to 4K boundary (Vista/Win7 limitation)
+            // Align offset to 4K boundary
             char* tmpHost = const_cast<char*>(
                 amd::alignDown(reinterpret_cast<const char*>(dstHost),
                 PinnedMemoryAlignment));
@@ -262,7 +262,7 @@ DmaBlitManager::writeBuffer(
 
         // Check if a pinned transfer can be executed
         if (pinSize && (dstSize > MinSizeForPinnedTransfer)) {
-            // Allign offset to 4K boundary (Vista/Win7 limitation)
+            // Align offset to 4K boundary
             char* tmpHost = const_cast<char*>(
                 amd::alignDown(reinterpret_cast<const char*>(srcHost),
                 PinnedMemoryAlignment));
@@ -474,7 +474,7 @@ DmaBlitManager::copyBufferRect(
 
         hsa_signal_value_t val =
             hsa_signal_wait_acquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ,
-            0, uint64_t(-1), HSA_WAIT_STATE_ACTIVE);
+            0, uint64_t(-1), HSA_WAIT_STATE_BLOCKED);
 
         if (val != 0) {
             LogError("Async copy failed");
@@ -644,18 +644,19 @@ bool DmaBlitManager::hsaCopy(
     // Use SDMA to transfer the data
     status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent,
         size[0], 0, nullptr, completion_signal_);
+
     if (status == HSA_STATUS_SUCCESS) {
         hsa_signal_value_t val = hsa_signal_wait_acquire(
             completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0,
-            uint64_t(-1), HSA_WAIT_STATE_ACTIVE);
+            uint64_t(-1), HSA_WAIT_STATE_BLOCKED);
         if (val != (kInitVal - 1)) {
             LogError("Async copy failed");
             status = HSA_STATUS_ERROR;
         }
-   }
-   else {
+    }
+    else {
         LogPrintfError("Hsa copy from host to device failed with code %d", status);
-   }
+    }
 
     return (status == HSA_STATUS_SUCCESS);
 }
@@ -695,7 +696,7 @@ bool DmaBlitManager::hsaCopyStaged(
                 hsa_signal_value_t val =
                 hsa_signal_wait_acquire(completion_signal_,
                 HSA_SIGNAL_CONDITION_EQ, 0,
-                uint64_t(-1), HSA_WAIT_STATE_ACTIVE);
+                uint64_t(-1), HSA_WAIT_STATE_BLOCKED);
 
                 if (val != (kInitVal - 1)) {
                     LogError("Async copy failed");
@@ -718,7 +719,7 @@ bool DmaBlitManager::hsaCopyStaged(
         if (status == HSA_STATUS_SUCCESS) {
             hsa_signal_value_t val = hsa_signal_wait_acquire(
             completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1),
-            HSA_WAIT_STATE_ACTIVE);
+            HSA_WAIT_STATE_BLOCKED);
 
             if (val != (kInitVal - 1)) {
                 LogError("Async copy failed");
@@ -2371,7 +2372,7 @@ DmaBlitManager::pinHostMemory(
     const static bool SysMem = true;
     amd::Memory* amdMemory;
 
-    // Allign offset to 4K boundary (Vista/Win7 limitation)
+    // Align offset to 4K boundary
     char* tmpHost = const_cast<char*>(
         amd::alignDown(reinterpret_cast<const char*>(hostMem),
         PinnedMemoryAlignment));