From bafa586cf5128b41433717c6abde8979707b30a9 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 19 Sep 2016 16:34:10 -0500 Subject: [PATCH] Cleanup, remove preCopyCommand. Change-Id: I3768d3789a99be8136b43179d4152fa1875665cb [ROCm/clr commit: b0509c5a58431a7a98d6e8ca70c51bc44e1bbb8e] --- .../clr/hipamd/include/hcc_detail/hip_hcc.h | 11 ---- projects/clr/hipamd/src/hip_hcc.cpp | 50 ------------------- 2 files changed, 61 deletions(-) diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index fcf5c558c6..01f1b01313 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h @@ -56,7 +56,6 @@ extern int HIP_VISIBLE_DEVICES; /* Contains a comma-separated sequence of GPU id //--- // Chicken bits for disabling functionality to work around potential issues: extern int HIP_DISABLE_HW_KERNEL_DEP; -extern int HIP_DISABLE_HW_COPY_DEP; //--- //Extern tls @@ -90,15 +89,6 @@ extern const char *API_COLOR_END; #define CTX_THREAD_SAFE 1 -// If FORCE_COPY_DEP=1 , HIP runtime will add -// synchronization for copy commands in the same stream, regardless of command type. -// If FORCE_COPY_DEP=0 data copies of the same kind (H2H, H2D, D2H, D2D) are assumed to be implicitly ordered. -// ROCR runtime implementation currently provides this guarantee when using SDMA queues but not -// when using shader queues. -// TODO - measure if this matters for performance, in particular for back-to-back small copies. -// If not, we can simplify the copy dependency tracking by collapsing to a single Copy type, and always forcing dependencies for copy commands. -#define FORCE_SAMEDIR_COPY_DEP 1 - // Compile debug trace mode - this prints debug messages to stderr when env var HIP_DB is set. // May be set to 0 to remove debug if checks - possible code size and performance difference? @@ -447,7 +437,6 @@ typedef uint64_t SeqNum_t ; void copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind); - int preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType); //--- // Member functions that begin with locked_ are thread-safe accessors - these acquire / release the critical mutex. diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index e4365dab6f..ad76dbce77 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -81,7 +81,6 @@ int HIP_VISIBLE_DEVICES = 0; /* Contains a comma-separated sequence of GPU ident //--- // Chicken bits for disabling functionality to work around potential issues: int HIP_DISABLE_HW_KERNEL_DEP = 0; -int HIP_DISABLE_HW_COPY_DEP = 0; @@ -481,54 +480,6 @@ void ihipStream_t::lockclose_postKernelCommand(hc::completion_future &kernelFutu -//--- -// Called whenever a copy command is set to the stream. -// Examines the last command sent to this stream and returns a signal to wait on, if required. -int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *copyCompletionSignal, hsa_signal_t *waitSignal, ihipCommand_t copyType) -{ - int needSync = 0; - - waitSignal->handle = 0; - - // If switching command types, we need to add a barrier packet to synchronize things. - if (FORCE_SAMEDIR_COPY_DEP || (crit->_last_command_type != copyType)) { - - - if (crit->_last_command_type == ihipCommandKernel) { - tprintf (DB_SYNC, "stream %p switch %s to %s (async copy dep on prev kernel)\n", - this, ihipCommandName[crit->_last_command_type], ihipCommandName[copyType]); - needSync = 1; - ihipSignal_t *depSignal = allocSignal(crit); - hsa_signal_store_relaxed(depSignal->_hsaSignal,1); - this->enqueueBarrier(static_cast(crit->_av.get_hsa_queue()), NULL, depSignal); - *waitSignal = depSignal->_hsaSignal; - } else if (crit->_last_copy_signal) { - needSync = 1; - tprintf (DB_SYNC, "stream %p switch %s to %s (async copy dep on other copy #%lu)\n", - this, ihipCommandName[crit->_last_command_type], ihipCommandName[copyType], crit->_last_copy_signal->_sigId); - *waitSignal = crit->_last_copy_signal->_hsaSignal; - } - - if (HIP_DISABLE_HW_COPY_DEP && needSync) { - if (HIP_DISABLE_HW_COPY_DEP == -1) { - tprintf (DB_SYNC, "IGNORE copy dependency\n") - - } else { - tprintf (DB_SYNC, "HOST-wait for copy dependency\n") - // do the wait here on the host, and disable the device-side command resolution. - hsa_signal_wait_acquire(*waitSignal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); - needSync = 0; - } - } - - crit->_last_command_type = copyType; - } - - crit->_last_copy_signal = copyCompletionSignal; - - return needSync; -} - // Precursor: the stream is already locked,specifically so this routine can enqueue work into the specified av. void ihipStream_t::launchModuleKernel( @@ -1286,7 +1237,6 @@ void ihipInit() READ_ENV_I(release, HIP_VISIBLE_DEVICES, CUDA_VISIBLE_DEVICES, "Only devices whose index is present in the secquence are visible to HIP applications and they are enumerated in the order of secquence" ); READ_ENV_I(release, HIP_DISABLE_HW_KERNEL_DEP, 0, "Disable HW dependencies before kernel commands - instead wait for dependency on host. -1 means ignore these dependencies. (debug mode)"); - READ_ENV_I(release, HIP_DISABLE_HW_COPY_DEP, 0, "Disable HW dependencies before copy commands - instead wait for dependency on host. -1 means ifnore these dependencies (debug mode)"); READ_ENV_I(release, HIP_NUM_KERNELS_INFLIGHT, 128, "Number of kernels per stream ");