added support pinned dma memcpy between host and device

[ROCm/hip commit: f74b7a3636]

projects/hip/src/hip_hcc.cpp

@@ -1259,6 +1259,7 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const
 
     if (kind == hipMemcpyHostToDevice) {
         int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, ihipCommandCopyH2D);
+        if(!srcTracked){
             if (HIP_STAGING_BUFFERS) {
                 tprintf(DB_COPY1, "D2H && !dstTracked: staged copy H2D dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
 
@@ -1279,24 +1280,65 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const
                 hc::am_copy(dst, src, sizeBytes);
 #endif
             }
+        }else{
+            hsa_agent_t dstAgent = *(static_cast<hsa_agent_t*>(dstPtrInfo._acc.get_hsa_agent()));
+            hsa_agent_t srcAgent = *(static_cast<hsa_agent_t*>(srcPtrInfo._acc.get_hsa_agent()));
+
+            ihipSignal_t *ihipSignal = allocSignal(crit);
+            hsa_signal_t copyCompleteSignal = ihipSignal->_hsa_signal;
+
+            hsa_signal_store_relaxed(copyCompleteSignal, 1);
+            void *devPtrSrc = srcPtrInfo._devicePointer;
+            tprintf(DB_COPY1, "HSA Async_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
+
+            hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, dstAgent, devPtrSrc, srcAgent, sizeBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, copyCompleteSignal);
+
+        // This is sync copy, so let's wait for copy right here:
+            if (hsa_status == HSA_STATUS_SUCCESS) {
+                waitCopy(crit, ihipSignal); // wait for copy, and return to pool.
+            } else {
+                throw ihipException(hipErrorInvalidValue);
+            }
+        }
     } else if (kind == hipMemcpyDeviceToHost) {
         int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, ihipCommandCopyD2H);
-        if (HIP_STAGING_BUFFERS) {
-            tprintf(DB_COPY1, "D2H && !dstTracked: staged copy D2H dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
-            //printf ("staged-copy- read dep signals\n");
-            device->_staging_buffer[1]->CopyDeviceToHost(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL);
+        if (!dstTracked){
+            if (HIP_STAGING_BUFFERS) {
+                tprintf(DB_COPY1, "D2H && !dstTracked: staged copy D2H dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
+                //printf ("staged-copy- read dep signals\n");
+                device->_staging_buffer[1]->CopyDeviceToHost(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL);
+    
+                // The copy completes before returning so can reset queue to empty:
+                this->wait(crit, true);
 
-            // The copy completes before returning so can reset queue to empty:
-            this->wait(crit, true);
-
-        } else {
+            } else {
             // TODO - remove, slow path.
-            tprintf(DB_COPY1, "D2H && !dstTracked: am_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
+                tprintf(DB_COPY1, "D2H && !dstTracked: am_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
 #if USE_AV_COPY
-            _av.copy(src, dst, sizeBytes);
+                _av.copy(src, dst, sizeBytes);
 #else
-            hc::am_copy(dst, src, sizeBytes);
+                hc::am_copy(dst, src, sizeBytes);
 #endif
+            }
+        }else{
+            hsa_agent_t dstAgent = *(static_cast<hsa_agent_t*>(dstPtrInfo._acc.get_hsa_agent()));
+            hsa_agent_t srcAgent = *(static_cast<hsa_agent_t*>(srcPtrInfo._acc.get_hsa_agent()));
+
+            ihipSignal_t *ihipSignal = allocSignal(crit);
+            hsa_signal_t copyCompleteSignal = ihipSignal->_hsa_signal;
+
+            hsa_signal_store_relaxed(copyCompleteSignal, 1);
+            void *devPtrDst = dstPtrInfo._devicePointer;
+            tprintf(DB_COPY1, "HSA Async_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
+
+            hsa_status_t hsa_status = hsa_amd_memory_async_copy(devPtrDst, dstAgent, src, srcAgent, sizeBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, copyCompleteSignal);
+
+        // This is sync copy, so let's wait for copy right here:
+            if (hsa_status == HSA_STATUS_SUCCESS) {
+                waitCopy(crit, ihipSignal); // wait for copy, and return to pool.
+            } else {
+                throw ihipException(hipErrorInvalidValue);
+            }
         }
     } else if (kind == hipMemcpyHostToHost)  {
         int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, ihipCommandCopyH2H);
@@ -1305,9 +1347,7 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const
             // host waits before doing host memory copy.
             hsa_signal_wait_acquire(depSignal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE);
         }
-        tprintf(DB_COPY1, "H2H memcpy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
         memcpy(dst, src, sizeBytes);
-
     } else if ((kind == hipMemcpyDeviceToDevice) && !copyEngineCanSeeSrcAndDest)  {
         int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, ihipCommandCopyP2P);
         if (HIP_STAGING_BUFFERS) {

projects/hip/tests/src/CMakeLists.txt

@@ -145,7 +145,7 @@ make_hip_executable (hip_popc hip_popc.cpp)
 make_hip_executable (hip_clz hip_clz.cpp) 
 make_hip_executable (hip_brev hip_brev.cpp) 
 make_hip_executable (hip_ffs hip_ffs.cpp) 
-make_hip_executable (hip_ldg hip_ldg.cpp)
+make_hip_executable (hip_test_ldg hip_test_ldg.cpp)
 make_hip_executable (hipGetDeviceAttribute hipGetDeviceAttribute.cpp) 
 make_hip_executable (hipEnvVar hipEnvVar.cpp) 
 make_hip_executable (hipEnvVarDriver hipEnvVarDriver.cpp) 
@@ -178,6 +178,7 @@ make_hip_executable (hipFuncDeviceSynchronize hipFuncDeviceSynchronize.cpp)
 make_hip_executable (hipPeerToPeer_simple hipPeerToPeer_simple.cpp) 
 make_hip_executable (hipMemcpyAll hipMemcpyAll.cpp)
 make_hip_executable (hipMultiThreadDevice hipMultiThreadDevice.cpp) 
+make_hip_executable (hipTestMemcpyPin hipTestMemcpyPin.cpp)
 
 make_test(hip_ballot " " )
 make_test(hip_anyall " " )
@@ -185,7 +186,7 @@ make_test(hip_popc " " )
 make_test(hip_brev " " )
 make_test(hip_clz " " )
 make_test(hip_ffs " " )
-make_test(hip_ldg " " )
+make_test(hip_test_ldg " " )
 make_test(hipEventRecord --iterations 10)
 make_test(hipMemset " " )
 make_test(hipMemset --N 10    --memsetval 0x42 )  # small copy, just 10 bytes.
@@ -222,6 +223,7 @@ make_test(hipFuncSetDeviceFlags " ")
 make_test(hipFuncGetDevice " ")
 make_test(hipFuncSetDevice " ")
 make_test(hipFuncDeviceSynchronize " ")
+make_test(hipTestMemcpyPin " ")
 make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-serial" --tests 0x1) 
 make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-pyramid" --tests 0x4) 
 make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-nearzero" --tests 0x10) 

projects/hip/tests/src/hipTestMemcpyPin.cpp

@@ -0,0 +1,33 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include<hip_runtime.h>
+#include"test_common.h"
+
+#define len 1024*1024
+#define size len * sizeof(float)
+
+int main(){
+    float *Ad, *A;
+    hipHostMalloc((void**)&A, size);
+    hipMalloc((void**)&Ad, size);
+    assert(hipSuccess == hipMemcpy(Ad, A, size, hipMemcpyHostToDevice));
+    assert(hipSuccess == hipMemcpy(A, Ad, size, hipMemcpyDeviceToHost));
+    passed();
+}

projects/hip/tests/src/hip_test_ldg.cpp

@@ -25,7 +25,7 @@ THE SOFTWARE.
 #include <stdlib.h>
 #include<iostream>
 #include "hip_runtime.h"
-
+#include "test_common.h"
 
 #define HIP_ASSERT(x) (assert((x)==hipSuccess))
 
@@ -125,7 +125,6 @@ bool dataTypesRun(){
     printf("FAILED: %d errors\n",errors);
     ret = false;
   } else {
-      printf ("PASSED!\n");
       ret = true;
   }
 
@@ -150,13 +149,14 @@ int main() {
   cout << " System major " << devProp.major << endl;
   cout << " agent prop name " << devProp.name << endl;
 
-  int errors;
-    errors = dataTypesRun<char>();
-    errors = dataTypesRun<signed char>();
-    errors = dataTypesRun<short>();
-    errors = dataTypesRun<int>();
-  cout << "__ldg " << endl ;
+   int errors = dataTypesRun<char>() &
+           dataTypesRun<signed char>() &
+           dataTypesRun<short>() &
+           dataTypesRun<int>();
   //hipResetDefaultAccelerator();
+    if(errors == 1){
+        passed();
+        return 0;
+    }
 
-  return errors;
 }