From 4bc4f3e0ab36f8fa3580826af4c83093f7bee6e8 Mon Sep 17 00:00:00 2001
From: foreman <dl.constructicon@amd.com>
Date: Wed, 16 May 2018 16:35:53 -0400
Subject: [PATCH] P4 to Git Change 1555866 by cpaquot@cpaquot-ocl-lc-lnx on
 2018/05/16 16:27:00

	SWDEV-145570 - [HIP] Store HIP mem flags inside amd::Buffer's flags
	Use the 16 upper bits of amd::Buffer's flags field instead of adding a new field.

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#26 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.cpp#86 edit


[ROCm/clr commit: 48a85c01e52055a0e333e35a9ed7cbe966c0b23c]
---
 projects/clr/hipamd/api/hip/hip_memory.cpp | 30 +++++++++++++++++++---
 1 file changed, 27 insertions(+), 3 deletions(-)

diff --git a/projects/clr/hipamd/api/hip/hip_memory.cpp b/projects/clr/hipamd/api/hip/hip_memory.cpp
index a754745ebb..05fcfbe21a 100644
--- a/projects/clr/hipamd/api/hip/hip_memory.cpp
+++ b/projects/clr/hipamd/api/hip/hip_memory.cpp
@@ -169,7 +169,19 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) {
 hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
   HIP_INIT_API(ptr, sizeBytes, flags);
 
-  return ihipMalloc(ptr, sizeBytes, CL_MEM_SVM_FINE_GRAIN_BUFFER);
+  if (ptr == nullptr) {
+    return hipErrorInvalidValue;
+  }
+  *ptr = nullptr;
+
+  const unsigned int coherentFlags = hipHostMallocCoherent | hipHostMallocNonCoherent;
+
+  // can't have both Coherent and NonCoherent flags set at the same time
+  if ((flags & coherentFlags) == coherentFlags) {
+    return hipErrorInvalidValue;
+  }
+
+  return ihipMalloc(ptr, sizeBytes, CL_MEM_SVM_FINE_GRAIN_BUFFER | (flags << 16));
 }
 
 hipError_t hipFree(void* ptr) {
@@ -467,9 +479,21 @@ hipError_t hipMalloc3DArray(hipArray_t* array, const struct hipChannelFormatDesc
 hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) {
   HIP_INIT_API(flagsPtr, hostPtr);
 
-  assert(0 && "Unimplemented");
+  if (flagsPtr == nullptr ||
+      hostPtr == nullptr) {
+    return hipErrorInvalidValue;
+  }
 
-  return hipErrorUnknown;
+  size_t offset = 0;
+  amd::Memory* svmMem = getMemoryObject(hostPtr, offset);
+
+  if (svmMem == nullptr) {
+    return hipErrorInvalidValue;
+  }
+
+  *flagsPtr = svmMem->getMemFlags() >> 16;
+
+  return hipSuccess;
 }
 
 hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) {