Add HIP_MAX_QUEUES feature.

Includes some tricky manipulation of the locks for contexts and streams.
issue is that stealing a stream requires we lock the context to
walk the streams to find a victim.  To avoid deadlock, we can't
have a stream locked when we lock the context.  This implementation
releases the stream lock, then acquires the context and selects the
victim.
A more stable implemenation might be to copy the stream list
from a context so that a lock is not required to walk all streams.
Smart shared_ptr could be used to prevent the streams from being
deallocated during the walk.

src/hip_stream.cpp

@@ -48,6 +48,7 @@ hipError_t ihipStreamCreate(hipStream_t *stream, unsigned int 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);
 
             ctxCrit->addStream(istream);
@@ -124,8 +125,14 @@ hipError_t hipStreamQuery(hipStream_t stream)
         stream =  device->_defaultStream;
     }
 
-    LockedAccessor_StreamCrit_t crit(stream->_criticalData);
-    int pendingOps = crit->_av.get_pending_async_ops();
+    int pendingOps = 0;
+
+    {
+        LockedAccessor_StreamCrit_t crit(stream->_criticalData);
+        if (crit->_hasQueue) {
+            pendingOps = crit->_av.get_pending_async_ops();
+        }
+    }
 
 
     hipError_t e = (pendingOps > 0) ? hipErrorNotReady : hipSuccess;