P4 to Git Change 2012424 by axie@axie_win_opencl_ssd on 2019/10/11 14:32:12

SWDEV-206239 - [hipclang-vdi-rocm][perf]RCCL: finegrain VRAM does not work
	xGMI may not work. Otherwise, several GPUs and CPU can access the same host memory atomically by API hipExtMallocWithFlags

	Tests:
	http://ocltc:8111/viewModification.html?modId=127344&personal=true&init=1&tab=vcsModificationBuilds

	ReviewBoard: http://ocltc.amd.com/reviews/r/18109/

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#76 edit

hipamd/api/hip/hip_memory.cpp

@@ -207,7 +207,7 @@ hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flag
     HIP_RETURN(hipErrorInvalidValue);
   }
 
-  HIP_RETURN(ihipMalloc(ptr, sizeBytes, (flags & hipDeviceMallocFinegrained)? CL_MEM_SVM_FINE_GRAIN_BUFFER : 0));
+  HIP_RETURN(ihipMalloc(ptr, sizeBytes, (flags & hipDeviceMallocFinegrained)? CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_SVM_ATOMICS: 0));
 }
 
 hipError_t hipMalloc(void** ptr, size_t sizeBytes) {