From 397ad08faecb2de3fc705c56c2e47f1a14da195b Mon Sep 17 00:00:00 2001
From: foreman <dl.swbuild@amd.com>
Date: Tue, 28 May 2019 19:06:58 -0400
Subject: [PATCH] P4 to Git Change 1788615 by cpaquot@cpaquot-ocl-lc-lnx on
 2019/05/28 18:54:55

	SWDEV-190565 - [HIP] Allow null stream as argument to hipStreamWaitEvent.

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_context.cpp#19 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_stream.cpp#19 edit
---
 hipamd/api/hip/hip_context.cpp | 24 ++++++++++++------------
 hipamd/api/hip/hip_stream.cpp  |  6 +++++-
 2 files changed, 17 insertions(+), 13 deletions(-)

diff --git a/hipamd/api/hip/hip_context.cpp b/hipamd/api/hip/hip_context.cpp
index 47bf651b8c..d519370971 100644
--- a/hipamd/api/hip/hip_context.cpp
+++ b/hipamd/api/hip/hip_context.cpp
@@ -255,7 +255,7 @@ hipError_t hipCtxGetDevice(hipDevice_t* device) {
     HIP_RETURN(hipErrorInvalidValue);
   }
 
-  return hipErrorUnknown;
+  HIP_RETURN(hipErrorUnknown);
 }
 
 hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) {
@@ -263,7 +263,7 @@ hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) {
 
   assert(0 && "Unimplemented");
 
-  return hipErrorUnknown;
+  HIP_RETURN(hipErrorUnknown);
 }
 
 hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) {
@@ -271,7 +271,7 @@ hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) {
 
   assert(0 && "Unimplemented");
 
-  return hipErrorUnknown;
+  HIP_RETURN(hipErrorUnknown);
 }
 
 hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) {
@@ -279,7 +279,7 @@ hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) {
 
   assert(0 && "Unimplemented");
 
-  return hipErrorUnknown;
+  HIP_RETURN(hipErrorUnknown);
 }
 
 hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) {
@@ -287,7 +287,7 @@ hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) {
 
   assert(0 && "Unimplemented");
 
-  return hipErrorUnknown;
+  HIP_RETURN(hipErrorUnknown);
 }
 
 hipError_t hipCtxSynchronize(void) {
@@ -295,7 +295,7 @@ hipError_t hipCtxSynchronize(void) {
 
   assert(0 && "Unimplemented");
 
-  return hipErrorUnknown;
+  HIP_RETURN(hipErrorUnknown);
 }
 
 hipError_t hipCtxGetFlags(unsigned int* flags) {
@@ -303,7 +303,7 @@ hipError_t hipCtxGetFlags(unsigned int* flags) {
 
   assert(0 && "Unimplemented");
 
-  return hipErrorUnknown;
+  HIP_RETURN(hipErrorUnknown);
 }
 
 hipError_t hipDevicePrimaryCtxGetState(hipDevice_t dev, unsigned int* flags, int* active) {
@@ -311,7 +311,7 @@ hipError_t hipDevicePrimaryCtxGetState(hipDevice_t dev, unsigned int* flags, int
 
   assert(0 && "Unimplemented");
 
-  return hipErrorUnknown;
+  HIP_RETURN(hipErrorUnknown);
 }
 
 hipError_t hipDevicePrimaryCtxRelease(hipDevice_t dev) {
@@ -319,7 +319,7 @@ hipError_t hipDevicePrimaryCtxRelease(hipDevice_t dev) {
 
   assert(0 && "Unimplemented");
 
-  return hipErrorUnknown;
+  HIP_RETURN(hipErrorUnknown);
 }
 
 hipError_t hipDevicePrimaryCtxRetain(hipCtx_t* pctx, hipDevice_t dev) {
@@ -327,7 +327,7 @@ hipError_t hipDevicePrimaryCtxRetain(hipCtx_t* pctx, hipDevice_t dev) {
 
   assert(0 && "Unimplemented");
 
-  return hipErrorUnknown;
+  HIP_RETURN(hipErrorUnknown);
 }
 
 hipError_t hipDevicePrimaryCtxReset(hipDevice_t dev) {
@@ -335,7 +335,7 @@ hipError_t hipDevicePrimaryCtxReset(hipDevice_t dev) {
 
   assert(0 && "Unimplemented");
 
-  return hipErrorUnknown;
+  HIP_RETURN(hipErrorUnknown);
 }
 
 hipError_t hipDevicePrimaryCtxSetFlags(hipDevice_t dev, unsigned int flags) {
@@ -343,5 +343,5 @@ hipError_t hipDevicePrimaryCtxSetFlags(hipDevice_t dev, unsigned int flags) {
 
   assert(0 && "Unimplemented");
 
-  return hipErrorUnknown;
+  HIP_RETURN(hipErrorUnknown);
 }
diff --git a/hipamd/api/hip/hip_stream.cpp b/hipamd/api/hip/hip_stream.cpp
index 72ede63ecb..59456303fe 100644
--- a/hipamd/api/hip/hip_stream.cpp
+++ b/hipamd/api/hip/hip_stream.cpp
@@ -188,7 +188,11 @@ hipError_t hipStreamDestroy(hipStream_t stream) {
 hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) {
   HIP_INIT_API(stream, event, flags);
 
-  if (stream == nullptr || event == nullptr) {
+  if (stream == nullptr) {
+    stream = reinterpret_cast<hipStream_t>(as_cl(hip::getNullStream()));
+  }
+
+  if (event == nullptr) {
     HIP_RETURN(hipErrorInvalidResourceHandle);
   }