From 5c5588bf20a13f7025f196d93d020495e1e2aaa6 Mon Sep 17 00:00:00 2001
From: foreman <dl.swbuild@amd.com>
Date: Fri, 22 Nov 2019 17:56:07 -0500
Subject: [PATCH] P4 to Git Change 2035631 by cpaquot@cpaquot-ocl-lc-lnx on
 2019/11/22 17:52:20

	SWDEV-213526 - pytorch tests fail with hipErrorOutofMemory
	There's a bug in ROCr when loading a lot of kernels and not syncing
	So for now, if an allocation fails, sync devices and retry before
	returning hipErrorOutOfMemory error.

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#86 edit
---
 api/hip/hip_memory.cpp | 9 ++++++++-
 1 file changed, 8 insertions(+), 1 deletion(-)

diff --git a/api/hip/hip_memory.cpp b/api/hip/hip_memory.cpp
index a249397605..c1cc189bb7 100644
--- a/api/hip/hip_memory.cpp
+++ b/api/hip/hip_memory.cpp
@@ -73,7 +73,14 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
 
   *ptr = amd::SvmBuffer::malloc(*amdContext, flags, sizeBytes, amdContext->devices()[0]->info().memBaseAddrAlign_);
   if (*ptr == nullptr) {
-    return hipErrorOutOfMemory;
+
+    hip::syncStreams();
+    hip::getNullStream()->finish();
+
+    *ptr = amd::SvmBuffer::malloc(*amdContext, flags, sizeBytes, amdContext->devices()[0]->info().memBaseAddrAlign_);
+    if (*ptr == nullptr) {
+      return hipErrorOutOfMemory;
+    }
   }
   ClPrint(amd::LOG_INFO, amd::LOG_API, "ihipMalloc ptr=0x%zx", *ptr);
   return hipSuccess;