From 3ce752065a6f4e354ee4add8b8ece9592dbbd2d1 Mon Sep 17 00:00:00 2001 From: foreman Date: Thu, 19 Apr 2018 18:35:00 -0400 Subject: [PATCH] P4 to Git Change 1544061 by skudchad@skudchad_test2_win_opencl on 2018/04/19 18:24:45 SWDEV-145570 - [HIP] - Add some hip_mem* APIs ReviewBoardURL = http://ocltc.amd.com/reviews/r/14647/diff/ Affected files ... ... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#10 edit --- api/hip/hip_memory.cpp | 125 +++++++++++++++++++++++++++++++++++------ 1 file changed, 108 insertions(+), 17 deletions(-) diff --git a/api/hip/hip_memory.cpp b/api/hip/hip_memory.cpp index 8cb51be08d..c0cc21367a 100644 --- a/api/hip/hip_memory.cpp +++ b/api/hip/hip_memory.cpp @@ -21,8 +21,10 @@ THE SOFTWARE. */ #include - #include "hip_internal.hpp" +#include "platform/context.hpp" +#include "platform/command.hpp" +#include "platform/memory.hpp" hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { @@ -119,7 +121,7 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind return hipSuccess; } -hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) { +hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) { HIP_INIT_API(dst, value, sizeBytes, stream); assert(0 && "Unimplemented"); @@ -130,9 +132,36 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t s hipError_t hipMemset(void* dst, int value, size_t sizeBytes) { HIP_INIT_API(dst, value, sizeBytes); - assert(0 && "Unimplemented"); + amd::Device* device = g_context->devices()[0]; - return hipErrorUnknown; + amd::HostQueue* queue = new amd::HostQueue(*g_context, *device, 0, + amd::CommandQueue::RealTimeDisabled, + amd::CommandQueue::Priority::Normal); + if (!queue) { + return hipErrorOutOfMemory; + } + + amd::Command::EventWaitList waitList; + amd::Memory* memory = amd::SvmManager::FindSvmBuffer(dst); + + + amd::Coord3D fillOffset(0, 0, 0); + amd::Coord3D fillSize(sizeBytes, 1, 1); + amd::FillMemoryCommand* command = + new amd::FillMemoryCommand(*queue, CL_COMMAND_FILL_BUFFER, waitList, *memory->asBuffer(), + &value, sizeof(int), fillOffset, fillSize); + + if (!command) { + return hipErrorOutOfMemory; + } + + command->enqueue(); + command->awaitCompletion(); + command->release(); + + queue->release(); + + return hipSuccess; } hipError_t hipMemPtrGetInfo(void *ptr, size_t *size) { @@ -146,17 +175,21 @@ hipError_t hipMemPtrGetInfo(void *ptr, size_t *size) { hipError_t hipHostFree(void* ptr) { HIP_INIT_API(ptr); - assert(0 && "Unimplemented"); - - return hipErrorUnknown; + if (amd::SvmBuffer::malloced(ptr)) { + amd::SvmBuffer::free(*g_context, ptr); + return hipSuccess; + } + return hipErrorInvalidValue; } hipError_t hipFreeArray(hipArray* array) { HIP_INIT_API(array); - assert(0 && "Unimplemented"); - - return hipErrorUnknown; + if (amd::SvmBuffer::malloced(array->data)) { + amd::SvmBuffer::free(*g_context, array->data); + return hipSuccess; + } + return hipErrorInvalidValue; } hipError_t hipMemGetAddressRange(hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr) { @@ -170,25 +203,83 @@ hipError_t hipMemGetAddressRange(hipDeviceptr_t* pbase, size_t* psize, hipDevice hipError_t hipMemGetInfo(size_t* free, size_t* total) { HIP_INIT_API(free, total); - assert(0 && "Unimplemented"); + size_t freeMemory[2]; + amd::Device* device = g_context->devices()[0]; + if(!device) { + return hipErrorInvalidDevice; + } - return hipErrorUnknown; + if(!device->globalFreeMemory(freeMemory)) { + return hipErrorInvalidValue; + } + + *free = freeMemory[0]; + *total = device->info().globalMemSize_; + +return hipSuccess; } +hipError_t ihipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height, size_t depth, + cl_mem_object_type imageType) { + + amd::Device* device = g_context->devices()[0]; + + if ((width == 0) || (height == 0)) { + *ptr = nullptr; + return hipSuccess; + } + else if (!(device->info().image2DMaxWidth_ >= width && + device->info().image2DMaxHeight_ >= height ) || (ptr == nullptr)) { + return hipErrorInvalidValue; + } + + if (g_context->devices()[0]->info().maxMemAllocSize_ < (width * height)) { + return hipErrorOutOfMemory; + } + + const cl_image_format image_format = { CL_R, CL_UNSIGNED_INT8 }; + const amd::Image::Format imageFormat(image_format); + + *pitch = width * imageFormat.getElementSize(); + + size_t sizeBytes = *pitch * height; + *ptr = amd::SvmBuffer::malloc(*g_context, CL_MEM_SVM_FINE_GRAIN_BUFFER, sizeBytes, + g_context->devices()[0]->info().memBaseAddrAlign_); + + if (!*ptr) { + return hipErrorOutOfMemory; + } + + return hipSuccess; +} + + hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height) { HIP_INIT_API(ptr, pitch, width, height); - assert(0 && "Unimplemented"); - - return hipErrorUnknown; + return ihipMallocPitch(ptr, pitch, width, height, 1, CL_MEM_OBJECT_IMAGE2D); } hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) { HIP_INIT_API(pitchedDevPtr, &extent); - assert(0 && "Unimplemented"); + size_t pitch = 0; - return hipErrorUnknown; + if (pitchedDevPtr == nullptr) { + return hipErrorInvalidValue; + } + + hipError_t status = hipSuccess; + status = ihipMallocPitch(&pitchedDevPtr->ptr, &pitch, extent.width, extent.height, extent.depth, + CL_MEM_OBJECT_IMAGE3D); + + if (status == hipSuccess) { + pitchedDevPtr->pitch = pitch; + pitchedDevPtr->xsize = extent.width; + pitchedDevPtr->ysize = extent.height; + } + + return status; } hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) {