Add HIP_INIT_ALLOC to init allocated memory.

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");
 
 

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;

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 )
 {