SWDEV-389611 - Changed hipMallocAsync to return success for zero size allocation to match hipMalloc

Change-Id: Ia2cbfaf2df9be37b62f23b2bb5b619205752ba84


[ROCm/clr commit: db9d9faaeb]

projects/clr/hipamd/src/hip_mempool.cpp

@@ -67,9 +67,13 @@ hipError_t hipDeviceGetMemPool(hipMemPool_t* mem_pool, int device) {
 // ================================================================================================
 hipError_t hipMallocAsync(void** dev_ptr, size_t size, hipStream_t stream) {
   HIP_INIT_API(hipMallocAsync, dev_ptr, size, stream);
-  if ((dev_ptr == nullptr) || (size == 0) || (!hip::isValid(stream))) {
+  if ((dev_ptr == nullptr) || (!hip::isValid(stream))) {
     HIP_RETURN(hipErrorInvalidValue);
   }
+  if (size == 0) {
+    *dev_ptr = nullptr;
+    HIP_RETURN(hipSuccess);
+  }
   auto hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->NullStream() :
     reinterpret_cast<hip::Stream*>(stream);
   auto device = hip_stream->GetDevice();
@@ -235,9 +239,13 @@ hipError_t hipMallocFromPoolAsync(
     hipMemPool_t mem_pool,
     hipStream_t stream) {
   HIP_INIT_API(hipMallocFromPoolAsync, dev_ptr, size, mem_pool, stream);
-  if ((dev_ptr == nullptr) || (size == 0) || (mem_pool == nullptr) || (!hip::isValid(stream))) {
+  if ((dev_ptr == nullptr) || (mem_pool == nullptr) || (!hip::isValid(stream))) {
     HIP_RETURN(hipErrorInvalidValue);
   }
+  if (size == 0) {
+    *dev_ptr = nullptr;
+    HIP_RETURN(hipSuccess);
+  }
   STREAM_CAPTURE(hipMallocAsync, stream, mem_pool, size, dev_ptr);
 
   auto mpool = reinterpret_cast<hip::MemoryPool*>(mem_pool);