From cdc8f6b9cc7736375719612423d4cf58d537c23d Mon Sep 17 00:00:00 2001 From: foreman Date: Tue, 1 May 2018 18:46:32 -0400 Subject: [PATCH] P4 to Git Change 1548558 by skudchad@skudchad_test2_win_opencl on 2018/05/01 18:39:58 SWDEV-145570 - [HIP] - Implement some of hipMem*Async and hipMemSet2D APIs. Part 4. ReviewBoardURL = http://ocltc.amd.com/reviews/r/14743/diff/ Affected files ... ... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#16 edit [ROCm/clr commit: 4a081cbc2f09af48b9c6bf74f291c7537fc1d416] --- projects/clr/hipamd/api/hip/hip_memory.cpp | 157 +++++++++++++++++++-- 1 file changed, 147 insertions(+), 10 deletions(-) diff --git a/projects/clr/hipamd/api/hip/hip_memory.cpp b/projects/clr/hipamd/api/hip/hip_memory.cpp index 3a4af23dee..10299de971 100644 --- a/projects/clr/hipamd/api/hip/hip_memory.cpp +++ b/projects/clr/hipamd/api/hip/hip_memory.cpp @@ -26,8 +26,15 @@ THE SOFTWARE. #include "platform/command.hpp" #include "platform/memory.hpp" -extern void getChannelOrderAndType(const hipChannelFormatDesc& desc, enum hipTextureReadMode readMode, - cl_channel_order* channelOrder, cl_channel_type* channelType); +extern void getChannelOrderAndType(const hipChannelFormatDesc& desc, + enum hipTextureReadMode readMode, + cl_channel_order* channelOrder, + cl_channel_type* channelType); + +extern void getDrvChannelOrderAndType(const enum hipArray_Format Format, + unsigned int NumChannels, + cl_channel_order* channelOrder, + cl_channel_type* channelType); hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { @@ -122,9 +129,35 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) { HIP_INIT_API(dst, value, sizeBytes, stream); - assert(0 && "Unimplemented"); + amd::HostQueue* queue; - return hipErrorUnknown; + if (stream == nullptr) { + queue = hip::getNullStream(); + } else { + queue = as_amd(reinterpret_cast(stream))->asHostQueue(); + } + + 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->release(); + + return hipSuccess; } hipError_t hipMemset(void* dst, int value, size_t sizeBytes) { @@ -294,9 +327,34 @@ hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) { hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) { HIP_INIT_API(array, pAllocateArray); - assert(0 && "Unimplemented"); + if (array[0]->width == 0) { + return hipErrorInvalidValue; + } - return hipErrorUnknown; + *array = (hipArray*)malloc(sizeof(hipArray)); + array[0]->drvDesc = *pAllocateArray; + array[0]->width = pAllocateArray->width; + array[0]->height = pAllocateArray->height; + array[0]->isDrv = true; + array[0]->textureType = hipTextureType2D; + void** ptr = &array[0]->data; + + cl_channel_order channelOrder; + cl_channel_type channelType; + getDrvChannelOrderAndType(pAllocateArray->format, pAllocateArray->numChannels, + &channelOrder, &channelType); + + const cl_image_format image_format = { channelOrder, channelType }; + size_t size = pAllocateArray->width; + if (pAllocateArray->height > 0) { + size = size * pAllocateArray->height; + } + + size_t pitch = 0; + hipError_t status = ihipMallocPitch(ptr, &pitch, array[0]->width, array[0]->height, 1, CL_MEM_OBJECT_IMAGE2D, + &image_format); + + return status; } hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, @@ -477,9 +535,44 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream) { HIP_INIT_API(dst, src, sizeBytes, kind, stream); - assert(0 && "Unimplemented"); + amd::Command* command = nullptr; + amd::Command::EventWaitList waitList; + amd::Memory* memory; + amd::HostQueue* queue; - return hipErrorUnknown; + if (stream == nullptr) { + queue = hip::getNullStream(); + } else { + queue = as_amd(reinterpret_cast(stream))->asHostQueue(); + } + + if (!queue) { + return hipErrorOutOfMemory; + } + + switch (kind) { + case hipMemcpyDeviceToHost: + memory = amd::SvmManager::FindSvmBuffer(src); + command = new amd::ReadMemoryCommand(*queue, CL_COMMAND_READ_BUFFER, waitList, + *memory->asBuffer(), 0, sizeBytes, dst); + break; + case hipMemcpyHostToDevice: + memory = amd::SvmManager::FindSvmBuffer(dst); + command = new amd::WriteMemoryCommand(*queue, CL_COMMAND_WRITE_BUFFER, waitList, + *memory->asBuffer(), 0, sizeBytes, src); + break; + default: + assert(!"Shouldn't reach here"); + break; + } + if (!command) { + return hipErrorOutOfMemory; + } + + command->enqueue(); + command->release(); + + return hipSuccess; } @@ -514,9 +607,52 @@ hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t height, hipMemcpyKind kind) { HIP_INIT_API(dst, dpitch, src, spitch, width, height, kind); - assert(0 && "Unimplemented"); + amd::HostQueue* queue = hip::getNullStream(); + if (!queue) { + return hipErrorOutOfMemory; + } - return hipErrorUnknown; + // Create buffer rectangle info structure + amd::BufferRect srcRect; + amd::BufferRect dstRect; + amd::Memory* srcPtr = amd::SvmManager::FindSvmBuffer(src); + amd::Memory* dstPtr = amd::SvmManager::FindSvmBuffer(dst); + size_t region[3] = {width, height, 0}; + size_t src_slice_pitch = spitch * height; + size_t dst_slice_pitch = dpitch * height; + size_t origin[3] = { }; + + if (!srcRect.create(origin, region, spitch, src_slice_pitch) || + !dstRect.create(origin, region, dpitch, dst_slice_pitch)) { + return hipErrorInvalidValue; + } + + amd::Coord3D srcStart(srcRect.start_, 0, 0); + amd::Coord3D dstStart(dstRect.start_, 0, 0); + amd::Coord3D srcEnd(srcRect.end_, 1, 1); + amd::Coord3D dstEnd(dstRect.end_, 1, 1); + + if (!srcPtr->asBuffer()->validateRegion(srcStart, srcEnd) || + !dstPtr->asBuffer()->validateRegion(dstStart, dstEnd)) { + return hipErrorInvalidValue; + } + + amd::Command::EventWaitList waitList; + amd::Coord3D size(region[0], region[1], region[2]); + + amd::CopyMemoryCommand* command = + new amd::CopyMemoryCommand(*queue, CL_COMMAND_COPY_BUFFER_RECT, waitList, *srcPtr->asBuffer(), + *dstPtr->asBuffer(), srcStart, dstStart, size, srcRect, dstRect); + + if (!command) { + return hipErrorOutOfMemory; + } + + command->enqueue(); + command->awaitCompletion(); + command->release(); + + return hipSuccess; } hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) { @@ -679,6 +815,7 @@ hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p) { HIP_INIT_API(p); amd::HostQueue* queue = hip::getNullStream(); + if (!queue) { return hipErrorOutOfMemory; }