Added missing NULL checks and corrected API return values as per validation

[ROCm/clr commit: 4989452413]

projects/clr/hipamd/src/hip_stream.cpp

@@ -56,6 +56,8 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priorit
     if (ctx) {
         if (HIP_FORCE_NULL_STREAM) {
             *stream = 0;
+        } else if( NULL == stream ){
+            e = hipErrorInvalidValue;
         } else {
             hc::accelerator acc = ctx->getWriteableDevice()->_acc;
 
@@ -65,7 +67,7 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priorit
             // CUDA stream behavior is that all kernels submitted will automatically
             // wait for prev to complete, this behaviour will be mainatined by 
             // hipModuleLaunchKernel. execute_any_order will help 
-            // hipExtModuleLaunchKernel , which uses a special flag
+	    // hipExtModuleLaunchKernel , which uses a special flag
 
             {
                 // Obtain mutex access to the device critical data, release by destructor
@@ -80,9 +82,9 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priorit
                 ctxCrit->addStream(istream);
                 *stream = istream;
             }
+            tprintf(DB_SYNC, "hipStreamCreate, %s\n", ToString(*stream).c_str());
         }
 
-        tprintf(DB_SYNC, "hipStreamCreate, %s\n", ToString(*stream).c_str());
     } else {
         e = hipErrorInvalidDevice;
     }
@@ -94,8 +96,10 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priorit
 //---
 hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags) {
     HIP_INIT_API(hipStreamCreateWithFlags, stream, flags);
-
-    return ihipLogStatus(ihipStreamCreate(stream, flags, priority_normal));
+    if(flags == hipStreamDefault || flags == hipStreamNonBlocking)
+        return ihipLogStatus(ihipStreamCreate(stream, flags, priority_normal));
+    else
+        return ihipLogStatus(hipErrorInvalidValue);
 }
 
 //---
@@ -128,25 +132,25 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
 
     hipError_t e = hipSuccess;
 
-    auto ecd = event->locked_copyCrit();
-
     if (event == nullptr) {
         e = hipErrorInvalidResourceHandle;
 
-    } else if ((ecd._state != hipEventStatusUnitialized) && (ecd._state != hipEventStatusCreated)) {
-        if (HIP_SYNC_STREAM_WAIT || (HIP_SYNC_NULL_STREAM && (stream == 0))) {
-            // conservative wait on host for the specified event to complete:
-            // return _stream->locked_eventWaitComplete(this, waitMode);
-            //
-            ecd._stream->locked_eventWaitComplete(
+    } else {
+        auto ecd = event->locked_copyCrit(); 
+        if ((ecd._state != hipEventStatusUnitialized) && (ecd._state != hipEventStatusCreated)) {
+            if (HIP_SYNC_STREAM_WAIT || (HIP_SYNC_NULL_STREAM && (stream == 0))) {
+                // conservative wait on host for the specified event to complete:
+                // return _stream->locked_eventWaitComplete(this, waitMode);
+                //
+                ecd._stream->locked_eventWaitComplete(
                 ecd.marker(), (event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked
                                                                      : hc::hcWaitModeActive);
-        } else {
-            stream = ihipSyncAndResolveStream(stream);
-            // This will use create_blocking_marker to wait on the specified queue.
-            stream->locked_streamWaitEvent(ecd);
+            } else {
+                stream = ihipSyncAndResolveStream(stream);
+                // This will use create_blocking_marker to wait on the specified queue.
+                stream->locked_streamWaitEvent(ecd);
+            }
         }
-
     }  // else event not recorded, return immediately and don't create marker.
 
     return ihipLogStatus(e);