P4 to Git Change 1562965 by skudchad@skudchad_test2_win_opencl on 2018/06/01 14:48:31

SWDEV-145570 - [HIP] - Implement hipHostRegister/Unregister, hipHostAlloc.

	ReviewBoardURL = http://ocltc.amd.com/reviews/r/15041/diff/

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_internal.hpp#12 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#30 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_texture.cpp#8 edit


[ROCm/clr commit: 2c4c7ab564]

projects/clr/hipamd/api/hip/hip_internal.hpp

@@ -62,5 +62,7 @@ namespace hip {
 };
 extern std::vector<amd::Context*> g_devices;
 extern hipError_t ihipDeviceGetCount(int* count);
+extern amd::Memory* getMemoryObject(const void* ptr, size_t& offset);
+
 
 #endif // HIP_SRC_HIP_INTERNAL_H

projects/clr/hipamd/api/hip/hip_memory.cpp

@@ -37,7 +37,7 @@ extern void getDrvChannelOrderAndType(const enum hipArray_Format Format,
                                       cl_channel_type* channelType);
 
 inline amd::Memory* getMemoryObject(const void* ptr, size_t& offset) {
-  amd::Memory *memObj = amd::SvmManager::FindSvmBuffer(ptr);
+  amd::Memory *memObj = amd::MemObjMap::FindMemObj(ptr);
   if (memObj != nullptr) {
     offset = reinterpret_cast<size_t>(ptr) - reinterpret_cast<size_t>(memObj->getSvmPtr());
   }
@@ -500,20 +500,48 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) {
 
 hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) {
   HIP_INIT_API(hostPtr, sizeBytes, flags);
+  if(hostPtr != nullptr) {
+    amd::Context *amdContext = hip::getCurrentContext();
+    amd::Memory* mem = new (*amdContext) amd::Buffer(*amdContext, CL_MEM_USE_HOST_PTR, sizeBytes);
 
-  assert(0 && "Unimplemented");
-
-  return hipErrorUnknown;
+    if (!mem->create(hostPtr)) {
+      mem->release();
+      return hipErrorMemoryAllocation;
+    }
+    amd::MemObjMap::AddMemObj(hostPtr, mem);
+    return hipSuccess;
+  } else {
+    return ihipMalloc(&hostPtr, sizeBytes, flags);
+  }
 }
 
 hipError_t hipHostUnregister(void* hostPtr) {
   HIP_INIT_API(hostPtr);
 
-  assert(0 && "Unimplemented");
+  if (amd::SvmBuffer::malloced(hostPtr)) {
+    hip::syncStreams();
+    hip::getNullStream()->finish();
+    amd::SvmBuffer::free(*hip::getCurrentContext(), hostPtr);
+    return hipSuccess;
+  } else {
+    size_t offset = 0;
+    amd::Memory* mem = getMemoryObject(hostPtr, offset);
 
-  return hipErrorUnknown;
+    if(mem) {
+      mem->release();
+      return hipSuccess;
+    }
+  }
+
+  return hipErrorInvalidValue;
 }
 
+// Deprecated function:
+hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) {
+  return ihipMalloc(ptr, sizeBytes, flags);
+};
+
+
 hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t count,
                              size_t offset, hipMemcpyKind kind) {
   HIP_INIT_API(symbolName, src, count, offset, kind);

projects/clr/hipamd/api/hip/hip_texture.cpp

@@ -157,11 +157,11 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou
   const amd::Image::Format imageFormat(image_format);
 
   amd::Memory* memory = nullptr;
-
+  size_t offset = 0;
   switch (pResDesc->resType) {
     case hipResourceTypeArray:
       {
-        memory = amd::SvmManager::FindSvmBuffer(pResDesc->res.array.array->data);
+        memory = getMemoryObject(pResDesc->res.array.array->data, offset);
 
         getChannelOrderAndType(pResDesc->res.array.array->desc, pTexDesc->readMode,
                              &image_format.image_channel_order, &image_format.image_channel_data_type);
@@ -187,8 +187,8 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou
       break;
     case hipResourceTypeLinear:
       assert(pResViewDesc == nullptr);
+      memory = getMemoryObject(pResDesc->res.linear.devPtr, offset);
 
-      memory = amd::SvmManager::FindSvmBuffer(pResDesc->res.linear.devPtr);
       image = new (*hip::getCurrentContext()) amd::Image(*memory->asBuffer(),
         CL_MEM_OBJECT_IMAGE1D, memory->getMemFlags(), imageFormat,
         pResDesc->res.linear.sizeInBytes / imageFormat.getElementSize(), 1, 1,
@@ -196,8 +196,8 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou
       break;
     case hipResourceTypePitch2D:
       assert(pResViewDesc == nullptr);
+      memory = getMemoryObject(pResDesc->res.pitch2D.devPtr, offset);
 
-      memory = amd::SvmManager::FindSvmBuffer(pResDesc->res.pitch2D.devPtr);
       image = new (*hip::getCurrentContext()) amd::Image(*memory->asBuffer(),
         CL_MEM_OBJECT_IMAGE2D, memory->getMemFlags(), imageFormat,
         pResDesc->res.pitch2D.width, pResDesc->res.pitch2D.height, 1,
@@ -263,8 +263,7 @@ hipError_t ihipBindTexture(cl_mem_object_type type,
         &image_format.image_channel_order, &image_format.image_channel_data_type);
     }
     const amd::Image::Format imageFormat(image_format);
-
-    amd::Memory* memory = amd::SvmManager::FindSvmBuffer(devPtr);
+    amd::Memory* memory = getMemoryObject(devPtr, *offset);
     amd::Image* image = new (*hip::getCurrentContext()) amd::Image(*memory->asBuffer(),
       type, memory->getMemFlags(), imageFormat, width, height, 1, pitch, 0);