SWDEV-408473 - Add wait time of 10 us if the waiting signal copy was < 24K.
Change-Id: I438ec9eb07e5034042a4a9a5e6e51d74daba2c83
[ROCm/clr commit: 6f5277c701]
Dieser Commit ist enthalten in:
kjayapra-amd
committet von
Karthik Jayaprakash
Karthik Jayaprakash
Ursprung
105212ef57
Commit
b432ced424
@@ -696,8 +696,12 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory,
|
||||
copyMask = kUseRegularCopyApi ? 0 : dev().fetchSDMAMask(this, true);
|
||||
}
|
||||
|
||||
// Check if host wait has to be forced
|
||||
bool forceHostWait = forceHostWaitFunc(size[0]);
|
||||
|
||||
auto wait_events = gpu().Barriers().WaitingSignal(engine);
|
||||
hsa_signal_t active = gpu().Barriers().ActiveSignal(kInitSignalValueOne, gpu().timestamp());
|
||||
hsa_signal_t active = gpu().Barriers().ActiveSignal(kInitSignalValueOne, gpu().timestamp(),
|
||||
forceHostWait);
|
||||
|
||||
if (!kUseRegularCopyApi && engine != HwQueueEngine::Unknown) {
|
||||
if (copyMask == 0) {
|
||||
@@ -2670,6 +2674,26 @@ amd::Memory* DmaBlitManager::pinHostMemory(const void* hostMem, size_t pinSize,
|
||||
return amdMemory;
|
||||
}
|
||||
|
||||
bool DmaBlitManager::forceHostWaitFunc(size_t copy_size) const {
|
||||
// 10us wait is true for all other targets.
|
||||
bool forceHostWait = true;
|
||||
// Based on the profiled results, do not wait for copy size > 24 KB.
|
||||
static constexpr size_t kGfx90aCopyThreshold = 24;
|
||||
|
||||
if ((dev().isa().versionMajor() == 9 && dev().isa().versionMinor() == 0
|
||||
&& dev().isa().versionStepping() == 10) && (copy_size >= kGfx90aCopyThreshold * Ki)) {
|
||||
// Check if this is gfx90a and restrict small copy to 24K.
|
||||
forceHostWait = false;
|
||||
} else if ((dev().isa().versionMajor() == 9) && (dev().isa().versionMinor() == 4)
|
||||
&& (dev().isa().versionStepping() == 0 || dev().isa().versionStepping() == 1
|
||||
|| dev().isa().versionStepping() == 2)) {
|
||||
// for gfx940, gfx941, gfx942, dependency signal resolution is fast, so no Host wait at all.
|
||||
forceHostWait = false;
|
||||
}
|
||||
|
||||
return forceHostWait;
|
||||
}
|
||||
|
||||
Memory* KernelBlitManager::createView(const Memory& parent, cl_image_format format,
|
||||
cl_mem_flags flags) const {
|
||||
assert((parent.owner()->asBuffer() == nullptr) && "View supports images only");
|
||||
@@ -2781,4 +2805,4 @@ bool KernelBlitManager::RunGwsInit(
|
||||
return result;
|
||||
}
|
||||
|
||||
} // namespace pal
|
||||
} // namespace roc
|
||||
|
||||
@@ -271,6 +271,8 @@ class DmaBlitManager : public device::HostBlitManager {
|
||||
address staging, //!< Staging resource
|
||||
bool hostToDev //!< True if data is copied from Host To Device
|
||||
) const;
|
||||
|
||||
bool forceHostWaitFunc(size_t copy_size) const;
|
||||
};
|
||||
|
||||
//! Kernel Blit Manager
|
||||
|
||||
@@ -81,14 +81,28 @@ public:
|
||||
hsa_signal_t signal_; //!< HSA signal to track profiling information
|
||||
Timestamp* ts_; //!< Timestamp object associated with the signal
|
||||
HwQueueEngine engine_; //!< Engine used with this signal
|
||||
bool done_; //!< True if signal is done
|
||||
amd::Monitor lock_; //!< Signal lock for update
|
||||
|
||||
typedef union {
|
||||
struct {
|
||||
uint32_t done_ : 1; //!< True if signal is done
|
||||
uint32_t forceHostWait_ : 1; //!< Force Host Wait for dependency signals
|
||||
uint32_t reserved_ : 30;
|
||||
};
|
||||
uint32_t data_;
|
||||
} Flags;
|
||||
|
||||
Flags flags_;
|
||||
|
||||
ProfilingSignal()
|
||||
: ts_(nullptr)
|
||||
, engine_(HwQueueEngine::Compute)
|
||||
, done_(true)
|
||||
, lock_("Signal Ops Lock", true)
|
||||
{ signal_.handle = 0; }
|
||||
{
|
||||
signal_.handle = 0;
|
||||
flags_.done_ = true;
|
||||
flags_.forceHostWait_ = true;
|
||||
}
|
||||
|
||||
virtual ~ProfilingSignal();
|
||||
amd::Monitor& LockSignalOps() { return lock_; }
|
||||
|
||||
@@ -151,7 +151,7 @@ void Timestamp::checkGpuTime() {
|
||||
ClPrint(amd::LOG_INFO, amd::LOG_SIG, "Signal = (0x%lx), start = %ld, "
|
||||
"end = %ld time taken= %ld ns", it->signal_.handle, start, end, end - start);
|
||||
}
|
||||
it->done_ = true;
|
||||
it->flags_.done_ = true;
|
||||
}
|
||||
signals_.clear();
|
||||
if (end != 0) {
|
||||
@@ -361,7 +361,7 @@ bool VirtualGPU::HwQueueTracker::Create() {
|
||||
|
||||
// ================================================================================================
|
||||
hsa_signal_t VirtualGPU::HwQueueTracker::ActiveSignal(
|
||||
hsa_signal_value_t init_val, Timestamp* ts) {
|
||||
hsa_signal_value_t init_val, Timestamp* ts, bool forceHostWait) {
|
||||
bool new_signal = false;
|
||||
|
||||
// Peep signal +2 ahead to see if its done
|
||||
@@ -422,8 +422,9 @@ hsa_signal_t VirtualGPU::HwQueueTracker::ActiveSignal(
|
||||
ProfilingSignal* prof_signal = signal_list_[current_id_];
|
||||
// Reset the signal and return
|
||||
hsa_signal_silent_store_relaxed(prof_signal->signal_, init_val);
|
||||
prof_signal->done_ = false;
|
||||
prof_signal->flags_.done_ = false;
|
||||
prof_signal->engine_ = engine_;
|
||||
prof_signal->flags_.forceHostWait_ = forceHostWait;
|
||||
if (ts != 0) {
|
||||
// Save HSA signal earlier to make sure the possible callback will have a valid
|
||||
// value for processing
|
||||
@@ -473,13 +474,9 @@ hsa_signal_t VirtualGPU::HwQueueTracker::ActiveSignal(
|
||||
// ================================================================================================
|
||||
std::vector<hsa_signal_t>& VirtualGPU::HwQueueTracker::WaitingSignal(HwQueueEngine engine) {
|
||||
bool explicit_wait = false;
|
||||
bool sdma_wait = false;
|
||||
// Reset all current waiting signals
|
||||
waiting_signals_.clear();
|
||||
|
||||
if(engine != HwQueueEngine::Compute)
|
||||
sdma_wait = true;
|
||||
|
||||
// Does runtime switch the active engine?
|
||||
if (engine != engine_) {
|
||||
// Yes, return the signal from the previous operation for a wait
|
||||
@@ -498,6 +495,7 @@ std::vector<hsa_signal_t>& VirtualGPU::HwQueueTracker::WaitingSignal(HwQueueEngi
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// Check if a wait is required
|
||||
if (explicit_wait) {
|
||||
bool skip_internal_signal = false;
|
||||
@@ -520,8 +518,10 @@ std::vector<hsa_signal_t>& VirtualGPU::HwQueueTracker::WaitingSignal(HwQueueEngi
|
||||
|
||||
if (hsa_signal_load_relaxed(external_signals_[i]->signal_) > 0) {
|
||||
const Settings& settings = gpu_.dev().settings();
|
||||
// Actively wait on CPU to avoid extra overheads of signal tracking on GPU
|
||||
if (!WaitForSignal<true>(external_signals_[i]->signal_, false, sdma_wait)) {
|
||||
// Actively wait on CPU to avoid extra overheads of signal tracking on GPU.
|
||||
// For small copies set forced wait
|
||||
if (!WaitForSignal<true>(external_signals_[i]->signal_, false,
|
||||
external_signals_[i]->flags_.forceHostWait_)) {
|
||||
if (settings.cpu_wait_for_signal_) {
|
||||
// Wait on CPU for completion if requested
|
||||
CpuWaitForSignal(external_signals_[i]);
|
||||
@@ -556,7 +556,7 @@ bool VirtualGPU::HwQueueTracker::CpuWaitForSignal(ProfilingSignal* signal) {
|
||||
LogPrintfError("Failed signal [0x%lx] wait", signal->signal_);
|
||||
return false;
|
||||
}
|
||||
signal->done_ = true;
|
||||
signal->flags_.done_ = true;
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
@@ -46,7 +46,7 @@ constexpr static uint64_t kUnlimitedWait = std::numeric_limits<uint64_t>::max();
|
||||
|
||||
// Active wait time out incase same sdma engine is used again,
|
||||
// then just wait instead of adding dependency wait signal.
|
||||
constexpr static uint64_t kForcedTimeout = 10;
|
||||
constexpr static uint64_t kForcedTimeout10us = 10;
|
||||
|
||||
template <bool active_wait_timeout = false>
|
||||
inline bool WaitForSignal(hsa_signal_t signal, bool active_wait = false, bool forced_wait = false) {
|
||||
@@ -56,8 +56,8 @@ inline bool WaitForSignal(hsa_signal_t signal, bool active_wait = false, bool fo
|
||||
timeout = kUnlimitedWait;
|
||||
}
|
||||
if (active_wait_timeout) {
|
||||
// If diff engine, wait to 10 ms. Otherwise no wait
|
||||
timeout = (forced_wait ? kForcedTimeout : ROC_ACTIVE_WAIT_TIMEOUT) * K;
|
||||
// If forced wait is set, then wait for 10us, else dont wait. (ns * K = us)
|
||||
timeout = (forced_wait ? kForcedTimeout10us : ROC_ACTIVE_WAIT_TIMEOUT) * K;
|
||||
if (timeout == 0) {
|
||||
return false;
|
||||
}
|
||||
@@ -232,7 +232,7 @@ class VirtualGPU : public device::VirtualDevice {
|
||||
|
||||
//! Finds a free signal for the upcomming operation
|
||||
hsa_signal_t ActiveSignal(hsa_signal_value_t init_val = kInitSignalValueOne,
|
||||
Timestamp* ts = nullptr);
|
||||
Timestamp* ts = nullptr, bool forceHostWait = true);
|
||||
|
||||
//! Wait for the curent active signal. Can idle the queue
|
||||
bool WaitCurrent() {
|
||||
|
||||
In neuem Issue referenzieren
Einen Benutzer sperren