From cea80cd8b3cfaba4c7a2d43b47264eb3e28f954b Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Wed, 13 Sep 2017 23:31:48 +0000
Subject: [PATCH] Add HIP_INIT_ALLOC to init allocated memory.

---
 src/hip_hcc.cpp        | 5 +++--
 src/hip_hcc_internal.h | 1 +
 src/hip_memory.cpp     | 6 ++++++
 3 files changed, 10 insertions(+), 2 deletions(-)

diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp
index dc72714e3e..5e297465ec 100644
--- a/src/hip_hcc.cpp
+++ b/src/hip_hcc.cpp
@@ -88,6 +88,8 @@ int HIP_HOST_COHERENT = 1;
 
 int HIP_SYNC_HOST_ALLOC = 1;
 
+int HIP_INIT_ALLOC=-1;
+
 
 #if (__hcc_workweek__ >= 17300)
 // Make sure we have required bug fix in HCC
@@ -267,8 +269,6 @@ ihipStream_t::ihipStream_t(ihipCtx_t *ctx, hc::accelerator_view av, unsigned int
         case hipDeviceScheduleBlockingSync  : _scheduleMode = Yield; break;
         default:_scheduleMode = Auto;
     };
-
-
 };
 
 
@@ -1248,6 +1248,7 @@ void HipReadEnv()
     READ_ENV_I(release, HIP_FAIL_SOC, 0, "Fault on Sub-Optimal-Copy, rather than use a slower but functional implementation.  Bit 0x1=Fail on async copy with unpinned memory.  Bit 0x2=Fail peer copy rather than use staging buffer copy");
 
     READ_ENV_I(release, HIP_SYNC_HOST_ALLOC, 0, "Sync before and after all host memory allocations.  May help stability");
+    READ_ENV_I(release, HIP_INIT_ALLOC, 0, "If not -1, initialize allocated memory to specified byte");
     READ_ENV_I(release, HIP_SYNC_NULL_STREAM, 0, "Synchronize on host for null stream submissions");
 
 
diff --git a/src/hip_hcc_internal.h b/src/hip_hcc_internal.h
index 56ca37b3e2..159a447081 100644
--- a/src/hip_hcc_internal.h
+++ b/src/hip_hcc_internal.h
@@ -67,6 +67,7 @@ extern int HIP_HIDDEN_FREE_MEM;
 extern int HIP_SYNC_HOST_ALLOC;
 
 extern int HIP_SYNC_NULL_STREAM;
+extern int HIP_INIT_ALLOC;
 
 // TODO - remove when this is standard behavior.
 extern int HCC_OPT_FLUSH;
diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp
index 4b3a4fcb12..b2f795ba75 100644
--- a/src/hip_memory.cpp
+++ b/src/hip_memory.cpp
@@ -119,6 +119,11 @@ void * allocAndSharePtr(const char *msg, size_t sizeBytes, ihipCtx_t *ctx, bool
     tprintf(DB_MEM, " alloc %s ptr:%p-%p size:%zu on dev:%d\n",
             msg, ptr, static_cast<char*>(ptr)+sizeBytes, sizeBytes, device->_deviceId);
 
+    if (HIP_INIT_ALLOC != -1) {
+        // TODO , dont' call HIP API directly here:
+        hipMemset(ptr, HIP_INIT_ALLOC, sizeBytes);
+    }	
+
     if (ptr != nullptr) {
         int r = sharePtr(ptr, ctx, shareWithAll, hipFlags);
         if (r != 0) {
@@ -1180,6 +1185,7 @@ ihipMemsetKernel(hipStream_t stream,
 
 }
 
+
 // TODO-sync: function is async unless target is pinned host memory - then these are fully sync.
 hipError_t hipMemsetAsync(void* dst, int  value, size_t sizeBytes, hipStream_t stream )
 {