From 24ec56cba147860017708e329dbdff83466b0be0 Mon Sep 17 00:00:00 2001 From: foreman Date: Fri, 1 Jun 2018 15:01:45 -0400 Subject: [PATCH] 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: 2c4c7ab56460af71c2f5d9e2341d955183f6414f] --- projects/clr/hipamd/api/hip/hip_internal.hpp | 2 + projects/clr/hipamd/api/hip/hip_memory.cpp | 40 +++++++++++++++++--- projects/clr/hipamd/api/hip/hip_texture.cpp | 11 +++--- 3 files changed, 41 insertions(+), 12 deletions(-) diff --git a/projects/clr/hipamd/api/hip/hip_internal.hpp b/projects/clr/hipamd/api/hip/hip_internal.hpp index 197e607597..e82cb13a2a 100644 --- a/projects/clr/hipamd/api/hip/hip_internal.hpp +++ b/projects/clr/hipamd/api/hip/hip_internal.hpp @@ -62,5 +62,7 @@ namespace hip { }; extern std::vector g_devices; extern hipError_t ihipDeviceGetCount(int* count); +extern amd::Memory* getMemoryObject(const void* ptr, size_t& offset); + #endif // HIP_SRC_HIP_INTERNAL_H diff --git a/projects/clr/hipamd/api/hip/hip_memory.cpp b/projects/clr/hipamd/api/hip/hip_memory.cpp index 90c060d25b..922b8abe8f 100644 --- a/projects/clr/hipamd/api/hip/hip_memory.cpp +++ b/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(ptr) - reinterpret_cast(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); diff --git a/projects/clr/hipamd/api/hip/hip_texture.cpp b/projects/clr/hipamd/api/hip/hip_texture.cpp index e619065f02..f8bdbe1c89 100644 --- a/projects/clr/hipamd/api/hip/hip_texture.cpp +++ b/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);