diff --git a/bin/hipcc b/bin/hipcc index 3e88b77693..ccea650776 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -244,6 +244,8 @@ my $toolArgs = ""; # arguments to pass to the hcc or nvcc tool foreach $arg (@ARGV) { + $trimarg = $arg; + $trimarg =~ s/^\s+|\s+$//g; # Remive whitespace my $swallowArg = 0; if ($arg eq '-c') { $compileOnly = 1; @@ -254,38 +256,37 @@ foreach $arg (@ARGV) $needLDFLAGS = 1; } - if(($arg eq '-stdlib=libc++') and ($setStdLib eq 0)) + if(($trimarg eq '-stdlib=libc++') and ($setStdLib eq 0)) { $HIPCXXFLAGS .= " -stdlib=libc++"; $setStdLib = 1; } - if(($arg eq '-stdlib=libstdc++') and ($setStdLib eq 0)) + if(($trimarg eq '-stdlib=libstdc++') and ($setStdLib eq 0)) { $HIPCXXFLAGS .= " -stdlib=libstdc++"; $HIPCXXFLAGS .= $HCC_WA_FLAGS; $setStdLib = 1; } - if($arg eq '--version') { + if($trimarg eq '--version') { $printHipVersion = 1; } - if($arg eq '--short-version') { + if($trimarg eq '--short-version') { $printHipVersion = 1; $runCmd = 0; } - if($arg eq '-M') { + if($trimarg eq '-M') { $compileOnly = 1; $buildDeps = 1; } - if($arg eq '-use_fast_math') { - print "In fast Math"; + if($trimarg eq '-use_fast_math') { $HIPCXXFLAGS .= " -DHIP_FAST_MATH "; } - if(($arg eq '-use-staticlib') and ($setLinkType eq 0)) + if(($trimarg eq '-use-staticlib') and ($setLinkType eq 0)) { $linkType = 0; $setLinkType = 1; } - if(($arg eq '-use-sharedlib') and ($setLinkType eq 0)) + if(($trimarg eq '-use-sharedlib') and ($setLinkType eq 0)) { $linkType = 1; $setLinkType = 1; diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index d760ade15d..c87e201c0c 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -69,6 +69,8 @@ std::string HIP_LAUNCH_BLOCKING_KERNELS; std::vector g_hipLaunchBlockingKernels; int HIP_API_BLOCKING = 0; +int HIP_MAX_QUEUES = 0; + int HIP_PRINT_ENV = 0; int HIP_TRACE_API= 0; std::string HIP_TRACE_API_COLOR("green"); @@ -254,7 +256,7 @@ ihipStream_t::ihipStream_t(ihipCtx_t *ctx, hc::accelerator_view av, unsigned int }; - tprintf(DB_SYNC, " streamCreate: stream=%p\n", this); + tprintf(DB_SYNC, " streamCreate: stream=%s\n", ToString(this).c_str()); }; @@ -269,7 +271,7 @@ ihipStream_t::~ihipStream_t() void ihipStream_t::wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty) { if (! assertQueueEmpty) { - tprintf (DB_SYNC, "stream %p wait for queue-empty..\n", this); + tprintf (DB_SYNC, "stream %s wait for queue-empty..\n", ToString(this).c_str()); hc::hcWaitMode waitMode = hc::hcWaitModeActive; if (_scheduleMode == Auto) { @@ -353,6 +355,7 @@ ihipCtx_t * ihipStream_t::getCtx() const // Lock the stream to prevent other threads from intervening. LockedAccessor_StreamCrit_t ihipStream_t::lockopen_preKernelCommand() { + LockedAccessor_StreamCrit_t crit(_criticalData, false/*no unlock at destruction*/); if(crit->_kernelCnt > HIP_NUM_KERNELS_INFLIGHT){ @@ -361,6 +364,17 @@ LockedAccessor_StreamCrit_t ihipStream_t::lockopen_preKernelCommand() } crit->_kernelCnt++; + if (HIP_MAX_QUEUES && !crit->_hasQueue) { + // Obtain mutex access to the device critical data, release by destructor + LockedAccessor_CtxCrit_t ctxCrit(this->_ctx->criticalData()); + crit->_av = this->_ctx->stealActiveQueue(ctxCrit, this); + crit->_hasQueue = true; + } + + + + assert(crit->_hasQueue); + return crit; } @@ -391,21 +405,22 @@ void ihipStream_t::lockclose_postKernelCommand(const char * kernelName, hc::acce }; -//============================================================================= -// Recompute the peercnt and the packed _peerAgents whenever a peer is added or deleted. -// The packed _peerAgents can efficiently be used on each memory allocation. -template<> -void ihipCtxCriticalBase_t::recomputePeerAgents() -{ - _peerCnt = 0; - std::for_each (_peers.begin(), _peers.end(), [this](ihipCtx_t* ctx) { - _peerAgents[_peerCnt++] = ctx->getDevice()->_hsaAgent; - }); -} + + //============================================================================= + // Recompute the peercnt and the packed _peerAgents whenever a peer is added or deleted. + // The packed _peerAgents can efficiently be used on each memory allocation. + template<> + void ihipCtxCriticalBase_t::recomputePeerAgents() + { + _peerCnt = 0; + std::for_each (_peers.begin(), _peers.end(), [this](ihipCtx_t* ctx) { + _peerAgents[_peerCnt++] = ctx->getDevice()->_hsaAgent; + }); + } -template<> -bool ihipCtxCriticalBase_t::isPeerWatcher(const ihipCtx_t *peer) + template<> + bool ihipCtxCriticalBase_t::isPeerWatcher(const ihipCtx_t *peer) { auto match = std::find(_peers.begin(), _peers.end(), peer); return (match != std::end(_peers)); @@ -880,6 +895,44 @@ std::string ihipCtx_t::toString() const return ss.str(); }; + +hc::accelerator_view +ihipCtx_t::stealActiveQueue(LockedAccessor_CtxCrit_t &ctxCrit, ihipStream_t *needyStream ) +{ + + // TODO - review handling if queue can't be found. + while (1) { + for (auto iter=ctxCrit->streams().begin(); iter != ctxCrit->streams().end(); iter++) { + if (*iter != needyStream) { + auto victimCritPtr = (*iter)->_criticalData.mtry_lock(); + if (victimCritPtr && victimCritPtr->_hasQueue && (victimCritPtr->_kernelCnt == 0)) { + + + victimCritPtr->_hasQueue = false; + + tprintf(DB_SYNC, " stealActiveQueue move queue from victim:%s to needy:%s\n", + ToString(*iter).c_str(), ToString(needyStream).c_str()); + + return victimCritPtr->_av; + } + } + } + } +} + + +hc::accelerator_view +ihipCtx_t::createOrStealQueue(LockedAccessor_CtxCrit_t &ctxCrit) +{ + if (HIP_MAX_QUEUES && (ctxCrit->streams().size() >= HIP_MAX_QUEUES)) { + // Steal a queue from an existing stream: + return this->stealActiveQueue (ctxCrit, nullptr); + } else { + // Create a new view + return getWriteableDevice()->_acc.create_view(); + } +} + //---- @@ -921,13 +974,6 @@ void ihipCtx_t::locked_syncDefaultStream(bool waitOnSelf) } } -//--- -void ihipCtx_t::locked_addStream(ihipStream_t *s) -{ - LockedAccessor_CtxCrit_t crit(_criticalData); - - crit->addStream(s); -} //--- void ihipCtx_t::locked_removeStream(ihipStream_t *s) @@ -1217,6 +1263,7 @@ void ihipInit() READ_ENV_I(release, HIP_API_BLOCKING, 0, "Make HIP APIs 'host-synchronous', so they block until completed. Impacts hipMemcpyAsync, hipMemsetAsync." ); + READ_ENV_I(release, HIP_MAX_QUEUES, 0, "Maximum number of queues that this app will use per-device. Additional streams will share the specified number of queues. 0=no limit."); READ_ENV_C(release, HIP_DB, 0, "Print debug info. Bitmask (HIP_DB=0xff) or flags separated by '+' (HIP_DB=api+sync+mem+copy)", HIP_DB_callback); if ((HIP_DB & (1< public: ihipStreamCriticalBase_t(hc::accelerator_view av) : _kernelCnt(0), - _av(av) + _av(av), + _hasQueue(true) { }; @@ -410,11 +410,20 @@ public: } ihipStreamCriticalBase_t * mlock() { LockedBase::lock(); return this;}; + ihipStreamCriticalBase_t * mtry_lock() { + return LockedBase::try_lock() ? this: nullptr; + }; public: - // TODO - remove _kernelCnt mechanism: uint32_t _kernelCnt; // Count of inflight kernels in this stream. Reset at ::wait(). + hc::accelerator_view _av; + + // True if the stream has an allocated queue (accelerato_view) for its use: + // Always true at ihipStream creation but queue may later be stolen. + // This acts as a valid bit for the _av. + bool _hasQueue; +private: }; @@ -422,6 +431,7 @@ public: // for the ihipCtx_t and then for the individual streams. The locks should not be acquired in reverse order // or deadlock may occur. In some cases, it may be possible to reduce the range where the locks must be held. // HIP routines should avoid acquiring and releasing the same lock during the execution of a single HIP API. +// Another option is to use try_lock in the innermost lock query. typedef ihipStreamCriticalBase_t ihipStreamCritical_t; @@ -436,6 +446,7 @@ public: enum ScheduleMode {Auto, Spin, Yield}; typedef uint64_t SeqNum_t ; + // TODOD -make av a reference to avoid shared_ptr overhead? ihipStream_t(ihipCtx_t *ctx, hc::accelerator_view av, unsigned int flags); ~ihipStream_t(); @@ -499,11 +510,14 @@ private: bool canSeeMemory(const ihipCtx_t *thisCtx, const hc::AmPointerInfo *dstInfo, const hc::AmPointerInfo *srcInfo); - -private: // Data +public: // TODO - move private // Critical Data - MUST be accessed through LockedAccessor_StreamCrit_t ihipStreamCritical_t _criticalData; +private: // Data + + std::mutex _hasQueueLock; + ihipCtx_t *_ctx; // parent context that owns this stream. // Friends: @@ -602,6 +616,7 @@ public: const std::list &const_streams() const { return _streams; }; + // Peer Accessor classes: bool isPeerWatcher(const ihipCtx_t *peer); // returns True if peer has access to memory physically located on this device. bool addPeerWatcher(const ihipCtx_t *thisCtx, ihipCtx_t *peer); @@ -651,17 +666,22 @@ public: // Functions: ihipCtx_t(ihipDevice_t *device, unsigned deviceCnt, unsigned flags); // note: calls constructor for _criticalData ~ihipCtx_t(); - // Functions which read or write the critical data are named locked_. + // Functions which read or write the critical data are named locked_. + // (might be better called "locking_" // ihipCtx_t does not use recursive locks so the ihip implementation must avoid calling a locked_ function from within a locked_ function. // External functions which call several locked_ functions will acquire and release the lock for each function. if this occurs in // performance-sensitive code we may want to refactor by adding non-locked functions and creating a new locked_ member function to call them all. - void locked_addStream(ihipStream_t *s); void locked_removeStream(ihipStream_t *s); void locked_reset(); void locked_waitAllStreams(); void locked_syncDefaultStream(bool waitOnSelf); - ihipCtxCritical_t &criticalData() { return _criticalData; }; // TODO, move private. Fix P2P. + // Will allocate a queue and assign it to the needyStream: + hc::accelerator_view stealActiveQueue(LockedAccessor_CtxCrit_t &ctxCrit, + ihipStream_t *needyStream); + hc::accelerator_view createOrStealQueue(LockedAccessor_CtxCrit_t &ctxCrit); + + ihipCtxCritical_t &criticalData() { return _criticalData; }; const ihipDevice_t *getDevice() const { return _device; }; int getDeviceNum() const { return _device->_deviceId; }; diff --git a/src/hip_stream.cpp b/src/hip_stream.cpp index 8350035357..d754ffe5f6 100644 --- a/src/hip_stream.cpp +++ b/src/hip_stream.cpp @@ -45,11 +45,15 @@ hipError_t ihipStreamCreate(hipStream_t *stream, unsigned int flags) //Note this is an execute_in_order queue, so all kernels submitted will atuomatically wait for prev to complete: //This matches CUDA stream behavior: - auto istream = new ihipStream_t(ctx, acc.create_view(), flags); + { + // Obtain mutex access to the device critical data, release by destructor + LockedAccessor_CtxCrit_t ctxCrit(ctx->criticalData()); + auto istream = new ihipStream_t(ctx, ctx->createOrStealQueue(ctxCrit), flags); - ctx->locked_addStream(istream); + ctxCrit->addStream(istream); + *stream = istream; + } - *stream = istream; tprintf(DB_SYNC, "hipStreamCreate, stream=%p\n", *stream); } else { e = hipErrorInvalidDevice;