Remove USE_AV_COPY, USE_PEER_TO_PEER fallback paths.

Change-Id: I9c20173e62029c4caebabc98784c6d7697758e4f


[ROCm/clr commit: 60f797cc5a]

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
 
 
 //---

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
     }
 
 

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;
             }

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.