diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp
index 860c6f87a8..445e74cbeb 100644
--- a/projects/clr/hipamd/src/hip_hcc.cpp
+++ b/projects/clr/hipamd/src/hip_hcc.cpp
@@ -88,11 +88,13 @@ int HIP_HOST_COHERENT = 1;
 
 int HIP_SYNC_HOST_ALLOC = 1;
 
-int HIP_SYNC_STREAM_WAIT = 0;
 
+int HIP_INIT_ALLOC=-1;
+int HIP_SYNC_STREAM_WAIT = 0;
 int HIP_FORCE_NULL_STREAM=0;
 
 
+
 #if (__hcc_workweek__ >= 17300)
 // Make sure we have required bug fix in HCC
 // Perform resolution on the GPU:
@@ -271,8 +273,6 @@ ihipStream_t::ihipStream_t(ihipCtx_t *ctx, hc::accelerator_view av, unsigned int
         case hipDeviceScheduleBlockingSync  : _scheduleMode = Yield; break;
         default:_scheduleMode = Auto;
     };
-
-
 };
 
 
@@ -1252,6 +1252,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");
     READ_ENV_I(release, HIP_FORCE_NULL_STREAM,  0, "Force all stream allocations to secretly return the null stream");
 
diff --git a/projects/clr/hipamd/src/hip_hcc_internal.h b/projects/clr/hipamd/src/hip_hcc_internal.h
index 88717e48bf..4b7e533a4c 100644
--- a/projects/clr/hipamd/src/hip_hcc_internal.h
+++ b/projects/clr/hipamd/src/hip_hcc_internal.h
@@ -68,8 +68,10 @@ extern int HIP_SYNC_HOST_ALLOC;
 extern int HIP_SYNC_STREAM_WAIT;
 
 extern int HIP_SYNC_NULL_STREAM;
+extern int HIP_INIT_ALLOC;
 extern int HIP_FORCE_NULL_STREAM;
 
+
 // TODO - remove when this is standard behavior.
 extern int HCC_OPT_FLUSH;
 
diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp
index 4b3a4fcb12..b2f795ba75 100644
--- a/projects/clr/hipamd/src/hip_memory.cpp
+++ b/projects/clr/hipamd/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 )
 {