diff --git a/hipamd/src/hip_event.cpp b/hipamd/src/hip_event.cpp index ea014ab292..88bee1b02e 100644 --- a/hipamd/src/hip_event.cpp +++ b/hipamd/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; @@ -202,11 +202,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 +258,22 @@ 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); - - 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) + { + return 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)) { + return ihipLogStatus(hipErrorNotReady); + } else { + return ihipLogStatus(hipSuccess); + } } } diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 36edcdb338..b85ba61584 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/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; diff --git a/hipamd/src/hip_peer.cpp b/hipamd/src/hip_peer.cpp index 32beda2caf..f2fca94404 100644 --- a/hipamd/src/hip_peer.cpp +++ b/hipamd/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", diff --git a/hipamd/src/hip_stream.cpp b/hipamd/src/hip_stream.cpp index c328e34b79..b748749538 100644 --- a/hipamd/src/hip_stream.cpp +++ b/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);