Added missing NULL checks and corrected API return values (#1188)

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

* Added missing NULL checks

src/hip_event.cpp

@@ -85,7 +85,7 @@ hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags) {
         (flags & ~supportedFlags) ||             // can't set any unsupported flags.
         (flags & releaseFlags) == releaseFlags;  // can't set both release flags
 
-    if (!illegalFlags) {
+    if (event && !illegalFlags) {
         *event = new ihipEvent_t(flags);
     } else {
         e = hipErrorInvalidValue;
@@ -109,46 +109,51 @@ hipError_t hipEventCreate(hipEvent_t* event) {
 
 hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) {
     HIP_INIT_SPECIAL_API(hipEventRecord, TRACE_SYNC, event, stream);
+ 
+    hipError_t status;
+    if (event){
+        auto ecd = event->locked_copyCrit();
+        if( ecd._state != hipEventStatusUnitialized) {
+            stream = ihipSyncAndResolveStream(stream);
 
-    auto ecd = event->locked_copyCrit();
+            if (HIP_SYNC_NULL_STREAM && stream->isDefaultStream()) {
+                // TODO-HIP_SYNC_NULL_STREAM : can remove this code when HIP_SYNC_NULL_STREAM = 0
+                //
+                // If default stream , then wait on all queues.
+                ihipCtx_t* ctx = ihipGetTlsDefaultCtx();
+                ctx->locked_syncDefaultStream(true, true);
 
-    if (event && ecd._state != hipEventStatusUnitialized) {
-        stream = ihipSyncAndResolveStream(stream);
+                { 
+                    LockedAccessor_EventCrit_t eCrit(event->criticalData());
+                    eCrit->_eventData.marker(hc::completion_future());  // reset event
+                    eCrit->_eventData._stream = stream;
+                    eCrit->_eventData._timestamp = hc::get_system_ticks();
+                    eCrit->_eventData._state = hipEventStatusComplete;
+                }
+                status = hipSuccess;
+            } else {
+                // Record the event in the stream:
+                // Keep a copy outside the critical section so we lock stream first, then event - to
+                // avoid deadlock
+                hc::completion_future cf = stream->locked_recordEvent(event);
 
-        if (HIP_SYNC_NULL_STREAM && stream->isDefaultStream()) {
-            // TODO-HIP_SYNC_NULL_STREAM : can remove this code when HIP_SYNC_NULL_STREAM = 0
-            //
-            // If default stream , then wait on all queues.
-            ihipCtx_t* ctx = ihipGetTlsDefaultCtx();
-            ctx->locked_syncDefaultStream(true, true);
+                {
+                    LockedAccessor_EventCrit_t eCrit(event->criticalData());
+                    eCrit->_eventData.marker(cf);
+                    eCrit->_eventData._stream = stream;
+                    eCrit->_eventData._timestamp = 0;
+                    eCrit->_eventData._state = hipEventStatusRecording;
+                }
 
-            {
-                LockedAccessor_EventCrit_t eCrit(event->criticalData());
-                eCrit->_eventData.marker(hc::completion_future());  // reset event
-                eCrit->_eventData._stream = stream;
-                eCrit->_eventData._timestamp = hc::get_system_ticks();
-                eCrit->_eventData._state = hipEventStatusComplete;
+                status = hipSuccess;
             }
-            return ihipLogStatus(hipSuccess);
         } else {
-            // Record the event in the stream:
-            // Keep a copy outside the critical section so we lock stream first, then event - to
-            // avoid deadlock
-            hc::completion_future cf = stream->locked_recordEvent(event);
-
-            {
-                LockedAccessor_EventCrit_t eCrit(event->criticalData());
-                eCrit->_eventData.marker(cf);
-                eCrit->_eventData._stream = stream;
-                eCrit->_eventData._timestamp = 0;
-                eCrit->_eventData._state = hipEventStatusRecording;
-            }
-
-            return ihipLogStatus(hipSuccess);
+            status = hipErrorInvalidResourceHandle;
         }
     } else {
-        return ihipLogStatus(hipErrorInvalidResourceHandle);
+            status = hipErrorInvalidResourceHandle;
     }
+    return ihipLogStatus(status);
 }
 
 
@@ -202,11 +207,13 @@ hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) {
 
     hipError_t status = hipSuccess;
 
-    *ms = 0.0f;
-
-    if ((start == nullptr) || (stop == nullptr)) {
+    if (ms == nullptr) {
+        status = hipErrorInvalidValue;
+    }
+    else if ((start == nullptr) || (stop == nullptr)) {
         status = hipErrorInvalidResourceHandle;
     } else {
+        *ms = 0.0f;
         auto startEcd = start->locked_copyCrit();
         auto stopEcd = stop->locked_copyCrit();
 
@@ -256,18 +263,26 @@ hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) {
 
 hipError_t hipEventQuery(hipEvent_t event) {
     HIP_INIT_SPECIAL_API(hipEventQuery, TRACE_QUERY, event);
+ 
+    hipError_t status = hipSuccess;
 
-    if (!(event->_flags & hipEventReleaseToSystem)) {
-        tprintf(DB_WARN,
-                "hipEventQuery on event without system-scope fence ; consider creating with "
-                "hipEventReleaseToSystem\n");
-    }
-
-    auto ecd = event->locked_copyCrit();
-
-    if ((ecd._state == hipEventStatusRecording) && !ecd._stream->locked_eventIsReady(event)) {
-        return ihipLogStatus(hipErrorNotReady);
+    if ( NULL == event)
+    {
+        status = hipErrorInvalidResourceHandle;
     } else {
-        return ihipLogStatus(hipSuccess);
+        if (!(event->_flags & hipEventReleaseToSystem)) {
+            tprintf(DB_WARN,
+                    "hipEventQuery on event without system-scope fence ; consider creating with "
+                    "hipEventReleaseToSystem\n");
+        }
+
+        auto ecd = event->locked_copyCrit();
+
+        if ((ecd._state == hipEventStatusRecording) && !ecd._stream->locked_eventIsReady(event)) {
+            status = hipErrorNotReady;
+        } else {
+            status = hipSuccess;
+        }
     }
+    return ihipLogStatus(status);
 }

src/hip_memory.cpp

@@ -243,13 +243,12 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) {
 
     auto ctx = ihipGetTlsDefaultCtx();
     // return NULL pointer when malloc size is 0
-    if (sizeBytes == 0) {
+    if ( nullptr == ctx || nullptr == ptr)  {
+        hip_status = hipErrorInvalidValue;
+    }
+    else if (sizeBytes == 0) {
         *ptr = NULL;
         hip_status = hipSuccess;
-
-    } else if ((ctx == nullptr) || (ptr == nullptr)) {
-        hip_status = hipErrorInvalidValue;
-
     } else {
         auto device = ctx->getWriteableDevice();
         *ptr = hip_internal::allocAndSharePtr("device_mem", sizeBytes, ctx, false /*shareWithAll*/,
@@ -309,12 +308,12 @@ hipError_t ihipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
     }
 
     auto ctx = ihipGetTlsDefaultCtx();
-
-    if (sizeBytes == 0) {
+    if ((ctx == nullptr) || (ptr == nullptr)) {
+        hip_status = hipErrorInvalidValue;
+    }
+    else if (sizeBytes == 0) {
         hip_status = hipSuccess;
         // TODO - should size of 0 return err or be siliently ignored?
-    } else if ((ctx == nullptr) || (ptr == nullptr)) {
-        hip_status = hipErrorInvalidValue;
     } else {
         unsigned trueFlags = flags;
         if (flags == hipHostMallocDefault) {
@@ -400,7 +399,7 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) {
 // width in bytes
 hipError_t ihipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height, size_t depth) {
     hipError_t hip_status = hipSuccess;
-    if(ptr==NULL)
+    if(ptr==NULL || pitch == NULL)
      {
 	hip_status=hipErrorInvalidValue;
        	return hip_status;
@@ -916,11 +915,8 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) {
     am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, hostPtr);
     if (status == AM_SUCCESS) {
         *flagsPtr = amPointerInfo._appAllocationFlags;
-        if (*flagsPtr == 0) {
-            hip_status = hipErrorInvalidValue;
-        } else {
-            hip_status = hipSuccess;
-        }
+        //0 is valid flag hipHostMallocDefault, and during hipHostMalloc if unsupported flags are passed as parameter it throws error
+        hip_status = hipSuccess;
         tprintf(DB_MEM, " %s: host ptr=%p\n", __func__, hostPtr);
     } else {
         hip_status = hipErrorInvalidValue;

src/hip_peer.cpp

@@ -41,8 +41,10 @@ THE SOFTWARE.
 hipError_t ihipDeviceCanAccessPeer(int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t peerCtx) {
     hipError_t err = hipSuccess;
 
-
-    if ((thisCtx != NULL) && (peerCtx != NULL)) {
+    if(canAccessPeer == NULL) {
+       err = hipErrorInvalidValue;
+    }
+    else if ((thisCtx != NULL) && (peerCtx != NULL)) {
         if (thisCtx == peerCtx) {
             *canAccessPeer = 0;
             tprintf(DB_MEM, "Can't be peer to self. (this=%s, peer=%s)\n",

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);