diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 91a26bf971..44cccf14fa 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/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(dstPtrInfo._acc.get_hsa_agent())); + hsa_agent_t srcAgent = *(static_cast(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(dstPtrInfo._acc.get_hsa_agent())); + hsa_agent_t srcAgent = *(static_cast(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) { diff --git a/projects/hip/tests/src/CMakeLists.txt b/projects/hip/tests/src/CMakeLists.txt index e58298248b..563cd2c738 100644 --- a/projects/hip/tests/src/CMakeLists.txt +++ b/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) diff --git a/projects/hip/tests/src/hipTestMemcpyPin.cpp b/projects/hip/tests/src/hipTestMemcpyPin.cpp new file mode 100644 index 0000000000..c36170003c --- /dev/null +++ b/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 +#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(); +} diff --git a/projects/hip/tests/src/hip_ldg.cpp b/projects/hip/tests/src/hip_test_ldg.cpp similarity index 94% rename from projects/hip/tests/src/hip_ldg.cpp rename to projects/hip/tests/src/hip_test_ldg.cpp index 612eae32e8..a58652d240 100644 --- a/projects/hip/tests/src/hip_ldg.cpp +++ b/projects/hip/tests/src/hip_test_ldg.cpp @@ -25,7 +25,7 @@ THE SOFTWARE. #include #include #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(); - errors = dataTypesRun(); - errors = dataTypesRun(); - errors = dataTypesRun(); - cout << "__ldg " << endl ; + int errors = dataTypesRun() & + dataTypesRun() & + dataTypesRun() & + dataTypesRun(); //hipResetDefaultAccelerator(); + if(errors == 1){ + passed(); + return 0; + } - return errors; }