Add initial HIP_SYNC_NULL_STREAM=0 mode.

This eliminates host-synchronization for null stream.  Instead, the
null-stream uses GPU-side events to wait for other streams.
Default is OFF pending additional testing.

Add enhanced null-stream test.

Also refine HIP_TRACE_API.

hipamd/src/hip_stream.cpp

@@ -150,7 +150,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream)
 
     if (stream == NULL) {
         ihipCtx_t *ctx = ihipGetTlsDefaultCtx();
-        ctx->locked_syncDefaultStream(true/*waitOnSelf*/);
+        ctx->locked_syncDefaultStream(true/*waitOnSelf*/, true/*syncToHost*/);
     } else {
         stream->locked_wait();
         e = hipSuccess;
@@ -174,7 +174,7 @@ hipError_t hipStreamDestroy(hipStream_t stream)
     //--- Drain the stream:
     if (stream == NULL) {
         ihipCtx_t *ctx = ihipGetTlsDefaultCtx();
-        ctx->locked_syncDefaultStream(true/*waitOnSelf*/);
+        ctx->locked_syncDefaultStream(true/*waitOnSelf*/, true /*syncToHost*/);
     } else {
         stream->locked_wait();
         e = hipSuccess;