From 2f4c30ff797495c04412e7ad15479c420ffd7c77 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Sun, 11 Sep 2016 11:11:17 -0500 Subject: [PATCH] Remove USE_AV_COPY, USE_PEER_TO_PEER fallback paths. Change-Id: I9c20173e62029c4caebabc98784c6d7697758e4f [ROCm/clr commit: 60f797cc5a003311f9149bd68da91130e03e4c16] --- projects/clr/hipamd/include/hcc_detail/hip_hcc.h | 7 ------- projects/clr/hipamd/src/hip_hcc.cpp | 4 ---- projects/clr/hipamd/src/hip_peer.cpp | 12 ------------ projects/clr/hipamd/src/hip_stream.cpp | 9 --------- 4 files changed, 32 deletions(-) diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index 3ce2e7ed2d..fcf5c558c6 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h @@ -31,13 +31,6 @@ THE SOFTWARE. // #define USE_MEMCPYTOSYMBOL // -//Use the new HCC accelerator_view::copy instead of am_copy -#define USE_AV_COPY (__hcc_workweek__ >= 16351) - -// Compile peer-to-peer support. -// >= 2 : use HCC hc:accelerator::get_is_peer -// >= 3 : use hc::am_memtracker_update_peers(...) -#define USE_PEER_TO_PEER 3 //--- diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index 09fbe545cc..f5f028603e 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -293,7 +293,6 @@ void ihipStream_t::locked_wait(bool assertQueueEmpty) }; -#if USE_AV_COPY // Causes current stream to wait for specified event to complete: void ihipStream_t::locked_waitEvent(hipEvent_t event) { @@ -302,7 +301,6 @@ void ihipStream_t::locked_waitEvent(hipEvent_t event) // TODO - check state of event here: crit->_av.create_blocking_marker(event->_marker); } -#endif // Create a marker in this stream. // Save state in the event so it can track the status of the event. @@ -1728,7 +1726,6 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const bool copyEngineCanSeeSrcAndDest = false; if (kind == hipMemcpyDeviceToDevice) { -#if USE_PEER_TO_PEER>=2 // 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. @@ -1736,7 +1733,6 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const if (dcrit->isPeer(ihipGetPrimaryCtx(dstPtrInfo._appId)) && (dcrit->isPeer(ihipGetPrimaryCtx(srcPtrInfo._appId)))) { copyEngineCanSeeSrcAndDest = true; } -#endif } diff --git a/projects/clr/hipamd/src/hip_peer.cpp b/projects/clr/hipamd/src/hip_peer.cpp index 8f75d99918..731b872c74 100644 --- a/projects/clr/hipamd/src/hip_peer.cpp +++ b/projects/clr/hipamd/src/hip_peer.cpp @@ -47,11 +47,7 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, hipCtx_t thisCtx, hipCtx_ if (thisCtx == peerCtx) { *canAccessPeer = 0; } else { -#if USE_PEER_TO_PEER>=2 *canAccessPeer = peerCtx->getDevice()->_acc.get_is_peer(thisCtx->getDevice()->_acc); -#else - *canAccessPeer = 0; -#endif } } else { @@ -75,12 +71,8 @@ hipError_t ihipDisablePeerAccess (hipCtx_t peerCtx) auto thisCtx = ihipGetTlsDefaultCtx(); if ((thisCtx != NULL) && (peerCtx != NULL)) { -#if USE_PEER_TO_PEER>=2 // Return true if thisCtx can access peerCtx's memory: bool canAccessPeer = peerCtx->getDevice()->_acc.get_is_peer(thisCtx->getDevice()->_acc); -#else - bool canAccessPeer = 0; -#endif if (! canAccessPeer) { err = hipErrorInvalidDevice; // P2P not allowed between these devices. @@ -90,10 +82,8 @@ hipError_t ihipDisablePeerAccess (hipCtx_t peerCtx) LockedAccessor_CtxCrit_t peerCrit(peerCtx->criticalData()); bool changed = peerCrit->removePeer(thisCtx); if (changed) { -#if USE_PEER_TO_PEER>=3 // Update the peers for all memory already saved in the tracker: am_memtracker_update_peers(peerCtx->getDevice()->_acc, peerCrit->peerCnt(), peerCrit->peerAgents()); -#endif } else { err = hipErrorPeerAccessNotEnabled; // never enabled P2P access. } @@ -124,9 +114,7 @@ hipError_t ihipEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags) LockedAccessor_CtxCrit_t peerCrit(peerCtx->criticalData()); bool isNewPeer = peerCrit->addPeer(thisCtx); if (isNewPeer) { -#if USE_PEER_TO_PEER>=3 am_memtracker_update_peers(peerCtx->getDevice()->_acc, peerCrit->peerCnt(), peerCrit->peerAgents()); -#endif } else { err = hipErrorPeerAccessAlreadyEnabled; } diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index 7b3dc31f07..9475e61b21 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -74,12 +74,6 @@ hipError_t hipStreamCreate(hipStream_t *stream) } -#if USE_AV_COPY==0 -//--- -/** - * @bug This function conservatively waits for all work in the specified stream to complete. - */ -#endif hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) { HIP_INIT_API(stream, event, flags); @@ -93,14 +87,11 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int bool fastWait = false; -#if USE_AV_COPY if (stream != hipStreamNull) { stream->locked_waitEvent(event); fastWait = true; // don't use the slow host-side synchronization. } - // TODO - clean up if/else logic when USE_AV_COPY enabled. -#endif if (!fastWait) { // TODO-hcc Convert to use create_blocking_marker(...) functionality.