AI/rocm-systems
2
0
Fork 0
Code Issues Pull-Requests Actions Pakete Projekte Releases Wiki Aktivität

Set forceHostCopyEngine for other copy dirs. Support HIP_FORCE_P2P_HOST

Quellcode durchsuchen 
Also: more debug for copy and P2p.

Change-Id: I87030c525410e041b2a00baaf6c68e6c0977ff42
Dieser Commit ist enthalten in:
Ben Sander
2016-11-04 18:40:10 -05:00
Ursprung ff3298f40e
Commit 092b3dacda
3 geänderte Dateien mit 48 neuen und 24 gelöschten Zeilen
Patch-Datei herunterladen Vergleichs-Datei herunterladen Alle Dateien ausklappen Alle Dateien einklappen
src/hip_hcc.cpp
+28 -15
Datei anzeigen
@@ -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
}
}
}
src/hip_hcc.h
+1
Datei anzeigen
@@ -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; };
src/hip_memory.cpp
+19 -9
Datei anzeigen
@@ -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)
{