tolerate spaces in hip args

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;

src/hip_hcc.cpp

@@ -69,6 +69,8 @@ std::string HIP_LAUNCH_BLOCKING_KERNELS;
 std::vector<std::string> 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<CtxMutex>::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<CtxMutex>::recomputePeerAgents()
+    {
+        _peerCnt = 0;
+        std::for_each (_peers.begin(), _peers.end(), [this](ihipCtx_t* ctx) {
+            _peerAgents[_peerCnt++] = ctx->getDevice()->_hsaAgent;
+        });
+    }
 
 
-template<>
-bool ihipCtxCriticalBase_t<CtxMutex>::isPeerWatcher(const ihipCtx_t *peer)
+    template<>
+    bool ihipCtxCriticalBase_t<CtxMutex>::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<<DB_API))  && (HIP_TRACE_API == 0)) {

src/hip_hcc.h

@@ -235,8 +235,7 @@ extern void recordApiTrace(std::string *fullStr, const std::string &apiStr);
 #define DB_SYNC   1 /* 0x02 - trace synchronization pieces */
 #define DB_MEM    2 /* 0x04 - trace memory allocation / deallocation */
 #define DB_COPY   3 /* 0x08 - trace memory copy and peer commands. . */
-#define DB_SIGNAL 4 /* 0x10 - trace signal pool commands */
-#define DB_MAX_FLAG 5
+#define DB_MAX_FLAG 4
 // When adding a new debug flag, also add to the char name table below.
 //
 
@@ -251,7 +250,6 @@ static const DbName dbName [] =
     {KYEL, "sync"},
     {KCYN, "mem"},
     {KMAG, "copy"},
-    {KRED, "signal"},
 };
 
 
@@ -366,6 +364,7 @@ struct LockedBase {
     // Most uses should use the lock-accessor.
     void lock() { _mutex.lock(); }
     void unlock() { _mutex.unlock(); }
+    bool try_lock() { return _mutex.try_lock(); }
 
     MUTEX_TYPE  _mutex;
 };
@@ -402,7 +401,8 @@ class ihipStreamCriticalBase_t : public LockedBase<MUTEX_TYPE>
 public:
     ihipStreamCriticalBase_t(hc::accelerator_view av) :
         _kernelCnt(0),
-        _av(av)
+        _av(av),
+        _hasQueue(true)
     {
     };
 
@@ -410,11 +410,20 @@ public:
     }
 
     ihipStreamCriticalBase_t<StreamMutex>  * mlock() { LockedBase<MUTEX_TYPE>::lock(); return this;};
+    ihipStreamCriticalBase_t<StreamMutex>  * mtry_lock() { 
+        return LockedBase<MUTEX_TYPE>::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<StreamMutex> 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<ihipStream_t*> &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; };

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;