diff --git a/hipamd/include/hcc_detail/hip_hcc.h b/hipamd/include/hcc_detail/hip_hcc.h index aacc63e1f2..7c47d1b2a9 100644 --- a/hipamd/include/hcc_detail/hip_hcc.h +++ b/hipamd/include/hcc_detail/hip_hcc.h @@ -25,7 +25,7 @@ THE SOFTWARE. #include "hip/hcc_detail/staging_buffer.h" -#if defined(__HCC__) && (__hcc_workweek__ < 16155) +#if defined(__HCC__) && (__hcc_workweek__ < 16186) #error("This version of HIP requires a newer version of HCC."); #endif @@ -37,11 +37,7 @@ THE SOFTWARE. // 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 2 - - -// Use new am_memory_host_lock APIs: -#define USE_HCC_LOCK_API 1 +#define USE_PEER_TO_PEER 3 //--- diff --git a/hipamd/src/hip_event.cpp b/hipamd/src/hip_event.cpp index 31caf84780..983b190644 100644 --- a/hipamd/src/hip_event.cpp +++ b/hipamd/src/hip_event.cpp @@ -134,11 +134,7 @@ hipError_t hipEventSynchronize(hipEvent_t event) device->locked_syncDefaultStream(true); return ihipLogStatus(hipSuccess); } else { -#if __hcc_workweek__ >= 16033 eh->_marker.wait((eh->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive); -#else - eh->_marker.wait(); -#endif eh->_stream->locked_reclaimSignals(eh->_copy_seq_id); return ihipLogStatus(hipSuccess); diff --git a/hipamd/src/hip_ldg.cpp b/hipamd/src/hip_ldg.cpp index b4b4e8f9e8..620ff91076 100644 --- a/hipamd/src/hip_ldg.cpp +++ b/hipamd/src/hip_ldg.cpp @@ -20,7 +20,6 @@ THE SOFTWARE. #include #include"hcc_detail/hip_ldg.h" -#if __hcc_workweek__ >= 16164 __device__ char __ldg(const char* ptr) { return ptr[0]; @@ -167,6 +166,5 @@ __device__ double2 __ldg(const double2* ptr) return ptr[0]; } -#endif diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 9769229ab4..cabea6f759 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -254,11 +254,7 @@ hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags) for(int i=0;i_acc, hostPtr, sizeBytes, &vecAcc[0], vecAcc.size()); -#else - am_status = AM_ERROR_MISC; -#endif if(am_status == AM_SUCCESS){ hip_status = hipSuccess; }else{ @@ -281,11 +277,7 @@ hipError_t hipHostUnregister(void *hostPtr) if(hostPtr == NULL){ hip_status = hipErrorInvalidValue; }else{ -#if USE_HCC_LOCK_API am_status_t am_status = hc::am_memory_host_unlock(device->_acc, hostPtr); -#else - am_status_t am_status = AM_ERROR_MISC; -#endif if(am_status != AM_SUCCESS){ hip_status = hipErrorHostMemoryNotRegistered; }