From 092b3dacda39e041a822a5273bf68fb843a7c08d Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 4 Nov 2016 18:40:10 -0500 Subject: [PATCH] Set forceHostCopyEngine for other copy dirs. Support HIP_FORCE_P2P_HOST Also: more debug for copy and P2p. Change-Id: I87030c525410e041b2a00baaf6c68e6c0977ff42 --- src/hip_hcc.cpp | 43 ++++++++++++++++++++++++++++--------------- src/hip_hcc.h | 1 + src/hip_memory.cpp | 28 +++++++++++++++++++--------- 3 files changed, 48 insertions(+), 24 deletions(-) diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index d74307918e..d3f87a15c9 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -1141,7 +1141,7 @@ void ihipReadEnv_Callback(void *var_ptr, const char *var_name1, const char *var_ env = getenv(var_name2); } - std::string var_string = "TBD"; + std::string var_string = "0"; if (env) { var_string = setterCallback(var_ptr, env); } @@ -1828,9 +1828,14 @@ void ihipStream_t::resolveHcMemcpyDirection(unsigned hipMemKind, const hc::AmPoi if (*hcCopyDir == hc::hcMemcpyDeviceToDevice) { if (!canSeePeerMemory(ctx, ihipGetPrimaryCtx(dstPtrInfo->_appId), ihipGetPrimaryCtx(srcPtrInfo->_appId))) { *forceHostCopyEngine = true; - tprintf (DB_COPY, "Forcing use of host copy engine.\n"); + tprintf (DB_COPY, "P2P D2D : copy engine cannot see both host and device pointers - forcing copy through staging buffers.\n"); } else { - tprintf (DB_COPY, "Will use SDMA engine on streamDevice=%s.\n", ctx->toString().c_str()); + if (HIP_FORCE_P2P_HOST ) { + *forceHostCopyEngine = true; + tprintf (DB_COPY, "P2P D2D. Copy engine can see src and dst but HIP_FORCE_P2P_HOST=0, forcing copy through staging buffers.\n"); + } else { + tprintf (DB_COPY, "P2P D2D. Copy engine can see src and dst, Will use SDMA engine on streamDevice=%s.\n", ctx->toString().c_str()); + } } }; } @@ -1858,15 +1863,19 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes, resolveHcMemcpyDirection(kind, &dstPtrInfo, &srcPtrInfo, &hcCopyDir, &forceHostCopyEngine); - tprintf (DB_COPY, "locked_copy dir=%s dst=%p src=%p sz=%zu\n", hcMemcpyStr(hcCopyDir), src, dst, sizeBytes); { LockedAccessor_StreamCrit_t crit (_criticalData); #if DISABLE_COPY_EXT #warning ("Disabled copy_ext path, P2P host staging copies will not work") + tprintf (DB_COPY, "copySync copyEngine_dev:%d dst=%p(home_dev:%d) src=%p(home_dev:%d) sz=%zu dstTracked=%d srcTracked=%d dir=%s forceHostCopyEngine=%d. Call HCC copy\n", + ctx->getDeviceNum(), dst, dstPtrInfo._appId, src, srcPtrInfo._appId, sizeBytes, dstTracked, srcTracked, hcMemcpyStr(hcCopyDir), forceHostCopyEngine); // Note - peer-to-peer copies which require host staging will not work in this path. crit->_av.copy(src, dst, sizeBytes); #else + // If srcTracked == dstTracked =1 and forceHostCopyEngine=0 then we wil use async SDMA. (assuming HCC implementation doesn't override somehow) + tprintf (DB_COPY, "copySync copyEngine_dev:%d dst=%p(home_dev:%d) src=%p(home_dev:%d) sz=%zu dstTracked=%d srcTracked=%d dir=%s forceHostCopyEngine=%d. Call HCC copy_ext.\n", + ctx->getDeviceNum(), dst, dstPtrInfo._appId, src, srcPtrInfo._appId, sizeBytes, dstTracked, srcTracked, hcMemcpyStr(hcCopyDir), forceHostCopyEngine); crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, forceHostCopyEngine); #endif } @@ -1904,19 +1913,18 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS); - bool copyEngineCanSeeSrcAndDest = true; - if ((kind == hipMemcpyDeviceToDevice) || - ((kind == hipMemcpyDefault) && srcTracked && dstTracked)) { - copyEngineCanSeeSrcAndDest = canSeePeerMemory(ctx, ihipGetPrimaryCtx(dstPtrInfo._appId), ihipGetPrimaryCtx(srcPtrInfo._appId)); - } + hc::hcCommandKind hcCopyDir; + bool forceHostCopyEngine; + resolveHcMemcpyDirection(kind, &dstPtrInfo, &srcPtrInfo, &hcCopyDir, &forceHostCopyEngine); - tprintf (DB_COPY, "locked_copyAsync: async memcpy dstTracked=%d srcTracked=%d copyEngineCanSeeSrcAndDest=%d\n", - dstTracked, srcTracked, copyEngineCanSeeSrcAndDest); + + tprintf (DB_COPY, "copyAsync dst=%p(home_dev:%d) src=%p(home_dev:%d) sz=%zu dstTracked=%d srcTracked=%d dir=%s forceHostCopyEngine=%d\n", + dst, dstPtrInfo._appId, src, srcPtrInfo._appId, sizeBytes, dstTracked, srcTracked, hcMemcpyStr(hcCopyDir), forceHostCopyEngine); // "tracked" really indicates if the pointer's virtual address is available in the GPU address space. // If both pointers are not tracked, we need to fall back to a sync copy. - if (dstTracked && srcTracked && copyEngineCanSeeSrcAndDest) { + if (dstTracked && srcTracked && !forceHostCopyEngine) { LockedAccessor_StreamCrit_t crit(_criticalData); // Perform asynchronous copy: @@ -1933,9 +1941,14 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes } } else { - // TODO - call copy_ext directly here? - locked_copySync(dst, src, sizeBytes, kind); - //crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, forceHostCopyEngine); + LockedAccessor_StreamCrit_t crit(_criticalData); +#if DISABLE_COPY_EXT +#warning ("Disabled copy_ext path, P2P host staging copies will not work") + // Note - peer-to-peer copies which require host staging will not work in this path. + crit->_av.copy(src, dst, sizeBytes); +#else + crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, forceHostCopyEngine); +#endif } } } diff --git a/src/hip_hcc.h b/src/hip_hcc.h index ecec00485d..30512a9bd2 100644 --- a/src/hip_hcc.h +++ b/src/hip_hcc.h @@ -657,6 +657,7 @@ public: // Functions: ihipCtxCritical_t &criticalData() { return _criticalData; }; // TODO, move private. Fix P2P. const ihipDevice_t *getDevice() const { return _device; }; + int getDeviceNum() const { return _device->_deviceId; }; // TODO - review uses of getWriteableDevice(), can these be converted to getDevice() ihipDevice_t *getWriteableDevice() const { return _device; }; diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index 4f45370c5a..85b22d86d0 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -123,23 +123,26 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) hip_status = hipErrorMemoryAllocation; } else { hc::am_memtracker_update(*ptr, device->_deviceId, 0); + int peerCnt=0; { LockedAccessor_CtxCrit_t crit(ctx->criticalData()); // the peerCnt always stores self so make sure the trace actually - if (crit->peerCnt() > 1) { + peerCnt = crit->peerCnt(); + if (peerCnt > 1) { hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr); } } + tprintf(DB_MEM, " allocated %p (size=%zu) on dev:%d and allowed %d other peer(s) access\n", *ptr, sizeBytes, device->_deviceId, peerCnt-1); } } else { hip_status = hipErrorMemoryAllocation; } - //printf (" hipMalloc allocated %p\n", *ptr); return ihipLogStatus(hip_status); } + hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { HIP_INIT_API(ptr, sizeBytes, flags); @@ -153,26 +156,28 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) auto device = ctx->getWriteableDevice(); if((flags == hipHostMallocDefault)|| (flags == hipHostMallocPortable)){ *ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned); - if(sizeBytes < 1 && (*ptr == NULL)){ + if (sizeBytes < 1 && (*ptr == NULL)) { hip_status = hipErrorMemoryAllocation; } else { hc::am_memtracker_update(*ptr, device->_deviceId, amHostPinned); } - tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr); - } else if(flags & hipHostMallocMapped){ + tprintf(DB_MEM, "allocated pinned host ptr=%p on dev=%d\n", *ptr, device->_deviceId); + } else if(flags & hipHostMallocMapped) { *ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned); - if(sizeBytes && (*ptr == NULL)){ + if (sizeBytes && (*ptr == NULL)) { hip_status = hipErrorMemoryAllocation; - }else{ + } else { hc::am_memtracker_update(*ptr, device->_deviceId, flags); + int peerCnt=0; { LockedAccessor_CtxCrit_t crit(ctx->criticalData()); - if (crit->peerCnt()) { + peerCnt = crit->peerCnt(); + if (peerCnt) { hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr); } } + tprintf(DB_MEM, "allocated pinned host ptr=%p on dev=%d, allow access to %d peer(s)\n", *ptr, device->_deviceId, peerCnt); } - tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr); } } return ihipLogStatus(hip_status); @@ -355,6 +360,8 @@ hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags) vecAcc.push_back(ihipGetDevice(i)->_acc); } am_status = hc::am_memory_host_lock(device->_acc, hostPtr, sizeBytes, &vecAcc[0], vecAcc.size()); + + tprintf(DB_MEM, " %s registered ptr=%p\n", __func__, hostPtr); if(am_status == AM_SUCCESS){ hip_status = hipSuccess; } else { @@ -378,6 +385,7 @@ hipError_t hipHostUnregister(void *hostPtr) }else{ auto device = ctx->getWriteableDevice(); am_status_t am_status = hc::am_memory_host_unlock(device->_acc, hostPtr); + tprintf(DB_MEM, " %s unregistered ptr=%p\n", __func__, hostPtr); if(am_status != AM_SUCCESS){ hip_status = hipErrorHostMemoryNotRegistered; } @@ -399,6 +407,7 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou hc::accelerator acc = ctx->getDevice()->_acc; void *ptr = acc.get_symbol_address(symbolName); + tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, ptr); if(ptr == nullptr) { @@ -428,6 +437,7 @@ hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_ hc::accelerator acc = ctx->getDevice()->_acc; void *ptr = acc.get_symbol_address(symbolName); + tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, ptr); if(ptr == nullptr) {