diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index f5f028603e..8bd1e0be3c 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -541,7 +541,7 @@ void ihipStream_t::launchModuleKernel( uint32_t gridDimY, uint32_t gridDimZ, uint32_t groupSegmentSize, - uint32_t privateSegmentSize, + uint32_t privateSegmentSize, void *kernarg, size_t kernSize, uint64_t kernel){ @@ -1708,151 +1708,7 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const throw ihipException(hipErrorInvalidDevice); } - hc::accelerator acc; - hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0); - hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0); - - bool dstTracked = (hc::am_memtracker_getinfo(&dstPtrInfo, dst) == AM_SUCCESS); - bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS); - bool srcInDeviceMem = srcPtrInfo._isInDeviceMem; - bool dstInDeviceMem = dstPtrInfo._isInDeviceMem; - - // Resolve default to a specific Kind so we know which algorithm to use: - if (kind == hipMemcpyDefault && resolveOn) { - kind = resolveMemcpyDirection(srcTracked, dstTracked, srcInDeviceMem, dstInDeviceMem); - }; - - hsa_signal_t depSignal; - - bool copyEngineCanSeeSrcAndDest = false; - if (kind == hipMemcpyDeviceToDevice) { - // Lock to prevent another thread from modifying peer list while we are trying to look at it. - LockedAccessor_CtxCrit_t dcrit(ctx->criticalData()); - // FIXME - this assumes peer access only from primary context. - // Would need to change the tracker to store a void * parameter that we could map to the ctx where the pointer is allocated. - if (dcrit->isPeer(ihipGetPrimaryCtx(dstPtrInfo._appId)) && (dcrit->isPeer(ihipGetPrimaryCtx(srcPtrInfo._appId)))) { - copyEngineCanSeeSrcAndDest = true; - } - } - - - if (kind == hipMemcpyHostToDevice) { - int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, ihipCommandCopyH2D); - if(!srcTracked){ - tprintf(DB_COPY1, "D2H && !dstTracked: staged copy H2D dst=%p src=%p sz=%zu\n", dst, src, sizeBytes); - UnpinnedCopyEngine::CopyMode copyMode = UnpinnedCopyEngine::UseStaging; - - if (HIP_OPTIMAL_MEM_TRANSFER) { - copyMode = UnpinnedCopyEngine::ChooseBest; - } else if (HIP_PININPLACE) { - copyMode = UnpinnedCopyEngine::UsePinInPlace; - } - device->_stagingBuffer[0]->CopyHostToDevice(copyMode, dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); - // The copy waits for inputs and then completes before returning so can reset queue to empty: - this->wait(crit, true); - } else { - // This is H2D copy, and source is pinned host memory : we can copy directly w/o using staging buffer. - 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->_hsaSignal; - - 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, g_cpu_agent, 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 (!dstTracked){ - tprintf(DB_COPY1, "D2H && !dstTracked: staged copy D2H dst=%p src=%p sz=%zu\n", dst, src, sizeBytes); - UnpinnedCopyEngine::CopyMode copyMode = UnpinnedCopyEngine::UseStaging; - - if (HIP_OPTIMAL_MEM_TRANSFER) { - copyMode = UnpinnedCopyEngine::ChooseBest; - } else if (HIP_PININPLACE) { - copyMode = UnpinnedCopyEngine::UsePinInPlace; - } - - device->_stagingBuffer[1]->CopyDeviceToHost(copyMode, dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); - - // The copy completes before returning so can reset queue to empty: - this->wait(crit, true); - } else { - // This is D2H copy, and destination is pinned host memory : we can copy directly w/o using staging buffer. - 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->_hsaSignal; - - 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, g_cpu_agent, 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); - - if (depSignalCnt) { - // host waits before doing host memory copy. - hsa_signal_wait_acquire(depSignal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); - } - memcpy(dst, src, sizeBytes); - } else if ((kind == hipMemcpyDeviceToDevice) && !copyEngineCanSeeSrcAndDest) { - int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, ihipCommandCopyP2P); - tprintf(DB_COPY1, "P2P but engine can't see both pointers: staged copy P2P dst=%p src=%p sz=%zu\n", dst, src, sizeBytes); - //printf ("staged-copy- read dep signals\n"); - hsa_agent_t dstAgent = * (static_cast (dstPtrInfo._acc.get_hsa_agent())); - hsa_agent_t srcAgent = * (static_cast (srcPtrInfo._acc.get_hsa_agent())); - - device->_stagingBuffer[1]->CopyPeerToPeer(dst, dstAgent, src, srcAgent, sizeBytes, depSignalCnt ? &depSignal : NULL); - - // The copy completes before returning so can reset queue to empty: - this->wait(crit, true); - } else { - // If not special case - these can all be handled by the hsa async copy: - ihipCommand_t commandType; - hsa_agent_t srcAgent, dstAgent; - setAsyncCopyAgents(kind, &commandType, &srcAgent, &dstAgent); - - int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, commandType); - - // Get a completion signal: - ihipSignal_t *ihipSignal = allocSignal(crit); - hsa_signal_t copyCompleteSignal = ihipSignal->_hsaSignal; - - hsa_signal_store_relaxed(copyCompleteSignal, 1); - - 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, 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); - } - } - + crit->_av.copy(src, dst, sizeBytes); }