From a3f4d9410cea30c299e1a8d638e04c3d9842262d Mon Sep 17 00:00:00 2001
From: Saleel Kudchadker <Saleel.Kudchadker@amd.com>
Date: Wed, 23 Feb 2022 18:21:04 -0800
Subject: [PATCH] SWDEV-321723 - Pass correct flags

Pass CL_MEM_SVM_ATOMICS with hipHostMallocNumaUser.

Change-Id: I7b8010c5584ffd2fbebb964832dcfa270d430830
---
 hipamd/src/hip_memory.cpp | 5 +++--
 1 file changed, 3 insertions(+), 2 deletions(-)

diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp
index fbcbaafc25..6c4fee0826 100644
--- a/hipamd/src/hip_memory.cpp
+++ b/hipamd/src/hip_memory.cpp
@@ -518,7 +518,7 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
   }
 
   if (flags & hipHostMallocNumaUser) {
-    ihipFlags |= CL_MEM_FOLLOW_USER_NUMA_POLICY;
+    ihipFlags |= CL_MEM_FOLLOW_USER_NUMA_POLICY | CL_MEM_SVM_ATOMICS;
   }
 
   HIP_RETURN_DURATION(ihipMalloc(ptr, sizeBytes, ihipFlags), *ptr);
@@ -978,7 +978,8 @@ hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags)
   HIP_INIT_API(hipHostRegister, hostPtr, sizeBytes, flags);
   CHECK_STREAM_CAPTURE_SUPPORTED();
   if(hostPtr != nullptr) {
-    amd::Memory* mem = new (*hip::host_device->asContext()) amd::Buffer(*hip::host_device->asContext(), CL_MEM_USE_HOST_PTR | CL_MEM_SVM_ATOMICS, sizeBytes);
+    amd::Memory* mem = new (*hip::host_device->asContext()) amd::Buffer(*hip::host_device->asContext(),
+                            CL_MEM_USE_HOST_PTR | CL_MEM_SVM_ATOMICS, sizeBytes);
 
     constexpr bool sysMemAlloc = false;
     constexpr bool skipAlloc = false;