From dd3fbeba4d01c9373d878c53edd53cb85cc377d2 Mon Sep 17 00:00:00 2001 From: foreman Date: Fri, 4 Oct 2019 19:02:35 -0400 Subject: [PATCH] P4 to Git Change 2008906 by axie@axie-hip-vdi-pal2 on 2019/10/04 18:55:34 SWDEV-189650 - [HIP-CLANG][HIP/VDI/PAL] Hangs on test hip_threadfence_system 1. In HIP + VDI + ROCm, allow SVM atomic in VEGA10 and later ASIC. GFX8 (Tonga) was enabled before. 2. In HIP + VDI + PAL Linux driver, allow SVM atomic in VEGA10 and later ASIC. Tests: 1. In HIP + VDI + ROCm, hip_threadfence_system test passed. 2. In HIP + VDI + PAL + Linux , hip_threadfence_system test passed. 3. OpenCL + PAL, clinfo and ocltest runtime test pass. 4. OpenCL + ROCM, clinfo and ocltest runtime test pass. 5. Windows 10, VEGA 10, clinfo and and ocltest runtime test pass. hip_threadfence_system test passed by skipping the test. Teamcity presubmission test: http://ocltc.amd.com:8111/viewModification.html?modId=127083&personal=true&tab=vcsModificationBuilds http://ocltc.amd.com:8111/viewModification.html?modId=127076&personal=true&tab=vcsModificationBuilds ReviewBoard: http://ocltc.amd.com/reviews/r/18077/ Affected files ... ... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#73 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldevice.cpp#171 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palresource.cpp#80 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palresource.hpp#31 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.cpp#134 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocmemory.cpp#44 edit ... //depot/stg/opencl/drivers/opencl/runtime/utils/flags.hpp#320 edit --- hipamd/api/hip/hip_memory.cpp | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) diff --git a/hipamd/api/hip/hip_memory.cpp b/hipamd/api/hip/hip_memory.cpp index 19262ecc54..fad9d2aefe 100644 --- a/hipamd/api/hip/hip_memory.cpp +++ b/hipamd/api/hip/hip_memory.cpp @@ -231,7 +231,13 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(ihipMalloc(ptr, sizeBytes, CL_MEM_SVM_FINE_GRAIN_BUFFER | (flags << 16))); + unsigned int ihipFlags = CL_MEM_SVM_FINE_GRAIN_BUFFER | (flags << 16); + if (flags & hipHostMallocCoherent || + (!(flags & hipHostMallocNonCoherent) && HIP_HOST_COHERENT)) { + ihipFlags |= CL_MEM_SVM_ATOMICS; + } + + HIP_RETURN(ihipMalloc(ptr, sizeBytes, ihipFlags)); } hipError_t hipFree(void* ptr) {