AI/rocm-systems
2
0
Derivar 0
Código Questões Pedidos de integração Operações Pacotes Planeamentos Lançamentos Wiki Trabalho

Use HCC's synchronous accelerator_view::copy

Ver a proveniência 
Replace large block of HIP code with a call to HCC av::copy().

Change-Id: Ic32e1801cf8d4cd116ac02b72c41b1a1e4b6065c
Este cometimento está contido em:
Ben Sander
2016-09-19 11:20:51 -05:00
ascendente e843d8cb51
cometimento da44f3f907
1 ficheiros modificados com 2 adições e 146 eliminações
Descarregar ficheiro patch Descarregar ficheiro diff Expandir todos os ficheiros Colapsar todos os ficheiros
src/hip_hcc.cpp
+2 -146
Ver ficheiro
@@ -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<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->_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<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->_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<hsa_agent_t*> (dstPtrInfo._acc.get_hsa_agent()));
hsa_agent_t srcAgent = * (static_cast<hsa_agent_t*> (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);
}