From b48d6279efe4c4e89efde49c45d05f5352b71397 Mon Sep 17 00:00:00 2001 From: Vladislav Sytchenko Date: Sun, 23 Feb 2020 15:34:31 -0500 Subject: [PATCH] HIP-VDI texture rework The current texture implementation is based off the one for HIP-HCC. There's a lot of problems with it - only creating images from buffers, hard coding logic and ignoring user parameters. This leads to a whole lot of UB even with simple examples (as seen with RedShift's code). This CL is aimed to bring the HIP-VDI texture implementation closer to what is described by Cuda. hipMemcpyAtoA() - image to image copy. hipMemcpyHtoA()/hipMemcpyDtoA() - buffer to image copy. hipMemcpyAtoH()/hipMemcpyAtoD() - image to buffer copy. hipArrayCreate()/hipArray3DCreate()/hipMallocArray()/hipMalloc3DArray() - creates 1D/2D/3D/1D Array/2D Array images. hipCreateTextureObject() - creates sampler, (optional) creates 1D/2D image from buffer, (optional) creates image views. hipBindTexture() - creates 1D image from buffer (should create a typed buffer, however this is not compatible with HIP-HCC). hipBindTexture2D() - creates 2D image form buffer. hipBindTextureToArray() - creates image view. hipTexRefSetAddress() - creates 1D image from buffer (should create a typed buffer, however this is not compatible with HIP-HCC). hipTexRefSetAddress2D() - creates 2D image from buffer. hipTexRefSetArray() - creates image view. There are still a lot of TODOs in the code, here's a few important ones: 1. VDI doesn't support a lot of sampler flags. 2. VDI doesn't support device to image 2D/3D copy. 3. Mipmaps implementation is incomplete. 4. Image view implementation is incomplete. Change-Id: Ia374ee27aa14f76451fee7667495036f4419a487 [ROCm/clr commit: f71817a34278e508e68138a13747a28114c63195] --- projects/clr/hipamd/vdi/hip_conversions.hpp | 583 ++++++ projects/clr/hipamd/vdi/hip_memory.cpp | 1833 +++++++++++-------- projects/clr/hipamd/vdi/hip_texture.cpp | 1745 +++++++++++------- 3 files changed, 2718 insertions(+), 1443 deletions(-) create mode 100644 projects/clr/hipamd/vdi/hip_conversions.hpp diff --git a/projects/clr/hipamd/vdi/hip_conversions.hpp b/projects/clr/hipamd/vdi/hip_conversions.hpp new file mode 100644 index 0000000000..321ef7e024 --- /dev/null +++ b/projects/clr/hipamd/vdi/hip_conversions.hpp @@ -0,0 +1,583 @@ +/* +Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +#include +#include + +namespace hip +{ +inline +cl_channel_type getCLChannelType(const hipArray_Format hipFormat, + const hipTextureReadMode hipReadMode) { + if (hipReadMode == hipReadModeElementType) { + switch (hipFormat) { + case HIP_AD_FORMAT_UNSIGNED_INT8: + return CL_UNSIGNED_INT8; + case HIP_AD_FORMAT_SIGNED_INT8: + return CL_SIGNED_INT8; + case HIP_AD_FORMAT_UNSIGNED_INT16: + return CL_UNSIGNED_INT16; + case HIP_AD_FORMAT_SIGNED_INT16: + return CL_SIGNED_INT16; + case HIP_AD_FORMAT_UNSIGNED_INT32: + return CL_UNSIGNED_INT32; + case HIP_AD_FORMAT_SIGNED_INT32: + return CL_SIGNED_INT32; + case HIP_AD_FORMAT_HALF: + return CL_HALF_FLOAT; + case HIP_AD_FORMAT_FLOAT: + return CL_FLOAT; + } + } else if (hipReadMode == hipReadModeNormalizedFloat) { + switch (hipFormat) { + case HIP_AD_FORMAT_UNSIGNED_INT8: + return CL_UNORM_INT8; + case HIP_AD_FORMAT_SIGNED_INT8: + return CL_SNORM_INT8; + case HIP_AD_FORMAT_UNSIGNED_INT16: + return CL_UNORM_INT16; + case HIP_AD_FORMAT_SIGNED_INT16: + return CL_SNORM_INT16; + case HIP_AD_FORMAT_UNSIGNED_INT32: + return CL_UNSIGNED_INT32; + case HIP_AD_FORMAT_SIGNED_INT32: + return CL_SIGNED_INT32; + case HIP_AD_FORMAT_HALF: + return CL_HALF_FLOAT; + case HIP_AD_FORMAT_FLOAT: + return CL_FLOAT; + } + } + + ShouldNotReachHere(); +} + +inline +cl_channel_order getCLChannelOrder(const unsigned int hipNumChannels) { + switch (hipNumChannels) { + case 1: + return CL_R; + case 2: + return CL_RG; + case 4: + return CL_RGBA; + } + + ShouldNotReachHere(); +} + +inline +cl_mem_object_type getCLMemObjectType(const unsigned int hipWidth, + const unsigned int hipHeight, + const unsigned int hipDepth, + const unsigned int flags) { + if (flags == hipArrayDefault) { + if ((hipWidth != 0) && (hipHeight == 0) && (hipDepth == 0)) { + return CL_MEM_OBJECT_IMAGE1D; + } else if ((hipWidth != 0) && (hipHeight != 0) && (hipDepth == 0)) { + return CL_MEM_OBJECT_IMAGE2D; + } else if ((hipWidth != 0) && (hipHeight != 0) && (hipDepth != 0)) { + return CL_MEM_OBJECT_IMAGE3D; + } + } else if (flags == hipArrayLayered) { + if ((hipWidth != 0) && (hipHeight == 0) && (hipDepth != 0)) { + return CL_MEM_OBJECT_IMAGE1D_ARRAY; + } else if ((hipWidth != 0) && (hipHeight != 0) && (hipDepth != 0)) { + return CL_MEM_OBJECT_IMAGE2D_ARRAY; + } + } + + ShouldNotReachHere(); +} + +inline +cl_addressing_mode getCLAddressingMode(const hipTextureAddressMode hipAddressMode) { + switch (hipAddressMode) { + case hipAddressModeWrap: + return CL_ADDRESS_REPEAT; + case hipAddressModeClamp: + return CL_ADDRESS_CLAMP; + case hipAddressModeMirror: + return CL_ADDRESS_MIRRORED_REPEAT; + case hipAddressModeBorder: + return CL_ADDRESS_CLAMP_TO_EDGE; + } + + ShouldNotReachHere(); +} + +inline +cl_filter_mode getCLFilterMode(const hipTextureFilterMode hipFilterMode) { + switch (hipFilterMode) { + case hipFilterModePoint: + return CL_FILTER_NEAREST; + case hipFilterModeLinear: + return CL_FILTER_LINEAR; + } + + ShouldNotReachHere(); +} + +inline +cl_mem_object_type getCLMemObjectType(const hipResourceType hipResType) { + switch (hipResType) { + case hipResourceTypeLinear: + return CL_MEM_OBJECT_IMAGE1D; + case hipResourceTypePitch2D: + return CL_MEM_OBJECT_IMAGE2D; + } + + ShouldNotReachHere(); +} + +inline +size_t getElementSize(const hipArray_Format arrayFormat) { + switch (arrayFormat) { + case HIP_AD_FORMAT_UNSIGNED_INT8: + case HIP_AD_FORMAT_SIGNED_INT8: + return 1; + case HIP_AD_FORMAT_UNSIGNED_INT16: + case HIP_AD_FORMAT_SIGNED_INT16: + case HIP_AD_FORMAT_HALF: + return 2; + case HIP_AD_FORMAT_UNSIGNED_INT32: + case HIP_AD_FORMAT_SIGNED_INT32: + case HIP_AD_FORMAT_FLOAT: + return 4; + } + + ShouldNotReachHere(); +} + +inline +hipChannelFormatDesc getChannelFormatDesc(int numChannels, + hipArray_Format arrayFormat) { + switch (arrayFormat) { + case HIP_AD_FORMAT_UNSIGNED_INT8: + switch (numChannels) { + case 1: + return {8, 0, 0, 0, hipChannelFormatKindUnsigned}; + case 2: + return {8, 8, 0, 0, hipChannelFormatKindUnsigned}; + case 4: + return {8, 8, 8, 8, hipChannelFormatKindUnsigned}; + } + case HIP_AD_FORMAT_SIGNED_INT8: + switch (numChannels) { + case 1: + return {8, 0, 0, 0, hipChannelFormatKindSigned}; + case 2: + return {8, 8, 0, 0, hipChannelFormatKindSigned}; + case 4: + return {8, 8, 8, 8, hipChannelFormatKindSigned}; + } + case HIP_AD_FORMAT_UNSIGNED_INT16: + switch (numChannels) { + case 1: + return {16, 0, 0, 0, hipChannelFormatKindUnsigned}; + case 2: + return {16, 16, 0, 0, hipChannelFormatKindUnsigned}; + case 4: + return {16, 16, 16, 16, hipChannelFormatKindUnsigned}; + } + case HIP_AD_FORMAT_SIGNED_INT16: + switch (numChannels) { + case 1: + return {16, 0, 0, 0, hipChannelFormatKindSigned}; + case 2: + return {16, 16, 0, 0, hipChannelFormatKindSigned}; + case 4: + return {16, 16, 16, 16, hipChannelFormatKindSigned}; + } + case HIP_AD_FORMAT_UNSIGNED_INT32: + switch (numChannels) { + case 1: + return {32, 0, 0, 0, hipChannelFormatKindUnsigned}; + case 2: + return {32, 32, 0, 0, hipChannelFormatKindUnsigned}; + case 4: + return {32, 32, 32, 32, hipChannelFormatKindUnsigned}; + } + case HIP_AD_FORMAT_SIGNED_INT32: + switch (numChannels) { + case 1: + return {32, 0, 0, 0, hipChannelFormatKindSigned}; + case 2: + return {32, 32, 0, 0, hipChannelFormatKindSigned}; + case 4: + return {32, 32, 32, 32, hipChannelFormatKindSigned}; + } + case HIP_AD_FORMAT_HALF: + switch (numChannels) { + case 1: + return {16, 0, 0, 0, hipChannelFormatKindFloat}; + case 2: + return {16, 16, 0, 0, hipChannelFormatKindFloat}; + case 4: + return {16, 16, 16, 16, hipChannelFormatKindFloat}; + } + case HIP_AD_FORMAT_FLOAT: + switch (numChannels) { + case 1: + return {32, 0, 0, 0, hipChannelFormatKindFloat}; + case 2: + return {32, 32, 0, 0, hipChannelFormatKindFloat}; + case 4: + return {32, 32, 32, 32, hipChannelFormatKindFloat}; + } + } + + ShouldNotReachHere(); +} + +inline +unsigned int getNumChannels(const hipChannelFormatDesc& desc) { + return ((desc.x != 0) + (desc.y != 0) + (desc.z != 0) + (desc.w != 0)); +} + +inline +hipArray_Format getArrayFormat(const hipChannelFormatDesc& desc) { + switch (desc.f) { + case hipChannelFormatKindUnsigned: + switch (desc.x) { + case 8: + return HIP_AD_FORMAT_UNSIGNED_INT8; + case 16: + return HIP_AD_FORMAT_UNSIGNED_INT16; + case 32: + return HIP_AD_FORMAT_UNSIGNED_INT32; + } + case hipChannelFormatKindSigned: + switch (desc.x) { + case 8: + return HIP_AD_FORMAT_SIGNED_INT8; + case 16: + return HIP_AD_FORMAT_SIGNED_INT16; + case 32: + return HIP_AD_FORMAT_SIGNED_INT32; + } + case hipChannelFormatKindFloat: + switch (desc.x) { + case 16: + return HIP_AD_FORMAT_HALF; + case 32: + return HIP_AD_FORMAT_FLOAT; + } + } + + ShouldNotReachHere(); +} + +inline +int getNumChannels(const hipResourceViewFormat hipFormat) { + switch (hipFormat) { + case hipResViewFormatUnsignedChar1: + case hipResViewFormatSignedChar1: + case hipResViewFormatUnsignedShort1: + case hipResViewFormatSignedShort1: + case hipResViewFormatUnsignedInt1: + case hipResViewFormatHalf1: + case hipResViewFormatFloat1: + return 1; + case hipResViewFormatUnsignedChar2: + case hipResViewFormatSignedChar2: + case hipResViewFormatUnsignedShort2: + case hipResViewFormatSignedShort2: + case hipResViewFormatUnsignedInt2: + case hipResViewFormatHalf2: + case hipResViewFormatFloat2: + return 2; + case hipResViewFormatUnsignedChar4: + case hipResViewFormatSignedChar4: + case hipResViewFormatUnsignedShort4: + case hipResViewFormatSignedShort4: + case hipResViewFormatUnsignedInt4: + case hipResViewFormatHalf4: + case hipResViewFormatFloat4: + return 4; + } + + ShouldNotReachHere(); +} + +inline +hipArray_Format getArrayFormat(const hipResourceViewFormat hipFormat) { + switch (hipFormat) { + case hipResViewFormatUnsignedChar1: + case hipResViewFormatUnsignedChar2: + case hipResViewFormatUnsignedChar4: + return HIP_AD_FORMAT_UNSIGNED_INT8; + case hipResViewFormatSignedChar1: + case hipResViewFormatSignedChar2: + case hipResViewFormatSignedChar4: + return HIP_AD_FORMAT_SIGNED_INT8; + case hipResViewFormatUnsignedShort1: + case hipResViewFormatUnsignedShort2: + case hipResViewFormatUnsignedShort4: + return HIP_AD_FORMAT_UNSIGNED_INT16; + case hipResViewFormatSignedShort1: + case hipResViewFormatSignedShort2: + case hipResViewFormatSignedShort4: + return HIP_AD_FORMAT_SIGNED_INT16; + case hipResViewFormatUnsignedInt1: + case hipResViewFormatUnsignedInt2: + case hipResViewFormatUnsignedInt4: + return HIP_AD_FORMAT_UNSIGNED_INT32; + case hipResViewFormatSignedInt1: + case hipResViewFormatSignedInt2: + case hipResViewFormatSignedInt4: + return HIP_AD_FORMAT_SIGNED_INT32; + case hipResViewFormatHalf1: + case hipResViewFormatHalf2: + case hipResViewFormatHalf4: + return HIP_AD_FORMAT_HALF; + case hipResViewFormatFloat1: + case hipResViewFormatFloat2: + case hipResViewFormatFloat4: + return HIP_AD_FORMAT_FLOAT; + } + + ShouldNotReachHere(); +} + +inline +hipResourceViewFormat getResourceViewFormat(const hipChannelFormatDesc& desc) { + switch (desc.f) { + case hipChannelFormatKindUnsigned: + switch (getNumChannels(desc)) { + case 1: + switch (desc.x) { + case 8: + return hipResViewFormatUnsignedChar1; + case 16: + return hipResViewFormatUnsignedShort1; + case 32: + return hipResViewFormatUnsignedInt1; + } + case 2: + switch (desc.x) { + case 8: + return hipResViewFormatUnsignedChar2; + case 16: + return hipResViewFormatUnsignedShort2; + case 32: + return hipResViewFormatUnsignedInt2; + } + case 4: + switch (desc.x) { + case 8: + return hipResViewFormatUnsignedChar4; + case 16: + return hipResViewFormatUnsignedShort4; + case 32: + return hipResViewFormatUnsignedInt4; + } + } + case hipChannelFormatKindSigned: + switch (getNumChannels(desc)) { + case 1: + switch (desc.x) { + case 8: + return hipResViewFormatSignedChar1; + case 16: + return hipResViewFormatSignedShort1; + case 32: + return hipResViewFormatSignedInt1; + } + case 2: + switch (desc.x) { + case 8: + return hipResViewFormatSignedChar2; + case 16: + return hipResViewFormatSignedShort2; + case 32: + return hipResViewFormatSignedInt2; + } + case 4: + switch (desc.x) { + case 8: + return hipResViewFormatSignedChar4; + case 16: + return hipResViewFormatSignedShort4; + case 32: + return hipResViewFormatSignedInt4; + } + } + case hipChannelFormatKindFloat: + switch (getNumChannels(desc)) { + case 1: + switch (desc.x) { + case 16: + return hipResViewFormatHalf1; + case 32: + return hipResViewFormatFloat1; + } + case 2: + switch (desc.x) { + case 16: + return hipResViewFormatHalf2; + case 32: + return hipResViewFormatFloat2; + } + case 4: + switch (desc.x) { + case 16: + return hipResViewFormatHalf4; + case 32: + return hipResViewFormatFloat4; + } + } + } + + ShouldNotReachHere(); +} + +inline +hipTextureReadMode getReadMode(unsigned int flags) { + if (flags & HIP_TRSF_READ_AS_INTEGER) { + return hipReadModeElementType; + } else { + return hipReadModeNormalizedFloat; + } +} + +inline +int getNormalizedCoords(unsigned int flags) { + if (flags & HIP_TRSF_NORMALIZED_COORDINATES) { + return 1; + } else { + return 0; + } +} + +inline +int getSRGB(unsigned int flags) { + if (flags & HIP_TRSF_SRGB) { + return 1; + } else { + return 0; + } +} + +inline +hipTextureDesc getTextureDesc(const textureReference* texRef, + const hipTextureReadMode readMode) { + hipTextureDesc texDesc = {}; + std::memcpy(texDesc.addressMode, texRef->addressMode, sizeof(texDesc.addressMode)); + texDesc.filterMode = texRef->filterMode; + texDesc.readMode = readMode; + texDesc.sRGB = texRef->sRGB; + texDesc.normalizedCoords = texRef->normalized; + texDesc.maxAnisotropy = texRef->maxAnisotropy; + texDesc.mipmapFilterMode = texRef->mipmapFilterMode; + texDesc.mipmapLevelBias = texRef->mipmapLevelBias; + texDesc.minMipmapLevelClamp = texRef->minMipmapLevelClamp; + texDesc.maxMipmapLevelClamp = texRef->maxMipmapLevelClamp; + + return texDesc; +} + +inline +hipResourceViewDesc getResourceViewDesc(hipArray_const_t array, + const hipResourceViewFormat format) { + hipResourceViewDesc resViewDesc = {}; + resViewDesc.format = format; + resViewDesc.width = array->width; + resViewDesc.height = array->height; + resViewDesc.depth = array->depth; + resViewDesc.firstMipmapLevel = 0; + resViewDesc.lastMipmapLevel = 0; + resViewDesc.firstLayer = 0; + resViewDesc.lastLayer = 0; /* TODO add hipArray::numLayers */ + + return resViewDesc; +} + +inline +hipResourceViewDesc getResourceViewDesc(hipMipmappedArray_const_t array, + const hipResourceViewFormat format) { + hipResourceViewDesc resViewDesc = {}; + resViewDesc.format = format; + resViewDesc.width = array->width; + resViewDesc.height = array->height; + resViewDesc.depth = array->depth; + resViewDesc.firstMipmapLevel = 0; + resViewDesc.lastMipmapLevel = 0; /* TODO add hipMipmappedArray::numMipLevels */ + resViewDesc.firstLayer = 0; + resViewDesc.lastLayer = 0; /* TODO add hipArray::numLayers */ + + return resViewDesc; +} + +inline +std::pair getMemoryType(const hipMemcpyKind kind) { + switch (kind) { + case hipMemcpyHostToHost: + return {hipMemoryTypeHost, hipMemoryTypeHost}; + case hipMemcpyHostToDevice: + return {hipMemoryTypeHost, hipMemoryTypeDevice}; + case hipMemcpyDeviceToHost: + return {hipMemoryTypeDevice, hipMemoryTypeHost}; + case hipMemcpyDeviceToDevice: + return {hipMemoryTypeDevice, hipMemoryTypeDevice}; + case hipMemcpyDefault: + return {hipMemoryTypeUnified, hipMemoryTypeUnified}; + } + + ShouldNotReachHere(); +} + +inline +hipMemcpy3DParms getMemcpy3DParms(const hip_Memcpy2D& desc2D) { + hipMemcpy3DParms desc3D = {}; + + desc3D.srcXInBytes = desc2D.srcXInBytes; + desc3D.srcY = desc2D.srcY; + desc3D.srcZ = 0; + desc3D.srcLOD = 0; + desc3D.srcMemoryType = desc2D.srcMemoryType; + desc3D.srcHost = desc2D.srcHost; + desc3D.srcDevice = desc2D.srcDevice; + desc3D.srcArray = desc2D.srcArray; + desc3D.srcPitch = desc2D.srcPitch; + desc3D.srcHeight = 0; + + desc3D.dstXInBytes = desc2D.dstXInBytes; + desc3D.dstY = desc2D.dstY; + desc3D.dstZ = 0; + desc3D.dstLOD = 0; + desc3D.dstMemoryType = desc2D.dstMemoryType; + desc3D.dstHost = desc2D.dstHost; + desc3D.dstDevice = desc2D.dstDevice; + desc3D.dstArray = desc2D.dstArray; + desc3D.dstPitch = desc2D.dstPitch; + desc3D.dstHeight = 0; + + desc3D.WidthInBytes = desc2D.WidthInBytes; + desc3D.Height = desc2D.Height; + desc3D.Depth = 0; + + return desc3D; +} +}; \ No newline at end of file diff --git a/projects/clr/hipamd/vdi/hip_memory.cpp b/projects/clr/hipamd/vdi/hip_memory.cpp index 4b36a70b6e..63dd2c712e 100644 --- a/projects/clr/hipamd/vdi/hip_memory.cpp +++ b/projects/clr/hipamd/vdi/hip_memory.cpp @@ -20,24 +20,11 @@ #include #include "hip_internal.hpp" +#include "hip_conversions.hpp" #include "platform/context.hpp" #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 getDrvChannelOrderAndType(const enum hipArray_Format Format, - unsigned int NumChannels, - cl_channel_order* channelOrder, - cl_channel_type* channelType); - -extern void setDescFromChannelType(cl_channel_type channelType, hipChannelFormatDesc* desc); - -extern void getByteSizeFromChannelFormatKind(enum hipChannelFormatKind channelFormatKind, size_t* byteSize); - amd::Memory* getMemoryObject(const void* ptr, size_t& offset) { amd::Memory *memObj = amd::MemObjMap::FindMemObj(ptr); if (memObj != nullptr) { @@ -282,14 +269,27 @@ hipError_t hipHostFree(void* ptr) { HIP_RETURN(hipErrorInvalidValue); } +hipError_t ihipArrayDestroy(hipArray* array) { + if (array == nullptr) { + return hipErrorInvalidValue; + } + + cl_mem memObj = reinterpret_cast(array->data); + if (is_valid(memObj) == false) { + return hipErrorInvalidValue; + } + + as_amd(memObj)->release(); + + delete array; + + return hipSuccess; +} + hipError_t hipFreeArray(hipArray* array) { HIP_INIT_API(hipFreeArray, array); - if (amd::SvmBuffer::malloced(array->data)) { - amd::SvmBuffer::free(*hip::getCurrentDevice()->asContext(), array->data); - HIP_RETURN(hipSuccess); - } - HIP_RETURN(hipErrorInvalidValue); + HIP_RETURN(ihipArrayDestroy(array)); } hipError_t hipMemGetAddressRange(hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr) { @@ -393,142 +393,228 @@ hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) { HIP_RETURN(status); } -hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) { +amd::Image* ihipImageCreate(const cl_channel_order channelOrder, + const cl_channel_type channelType, + const cl_mem_object_type imageType, + const size_t imageWidth, + const size_t imageHeight, + const size_t imageDepth, + const size_t imageArraySize, + const size_t imageRowPitch, + const size_t imageSlicePitch, + const uint32_t numMipLevels, + amd::Memory* buffer) { + const amd::Image::Format imageFormat({channelOrder, channelType}); + if (!imageFormat.isValid()) { + return nullptr; + } + + amd::Context& context = *hip::getCurrentDevice()->asContext(); + if (!imageFormat.isSupported(context, imageType)) { + return nullptr; + } + + const std::vector& devices = context.devices(); + if (!devices[0]->info().imageSupport_) { + return nullptr; + } + + if (!amd::Image::validateDimensions(devices, + imageType, + imageWidth, + imageHeight, + imageDepth, + imageArraySize)) { + return nullptr; + } + + // TODO validate the image descriptor. + + amd::Image* image = nullptr; + if (buffer != nullptr) { + switch (imageType) { + case CL_MEM_OBJECT_IMAGE1D: + case CL_MEM_OBJECT_IMAGE2D: + image = new (context) amd::Image(*buffer->asBuffer(), + imageType, + CL_MEM_READ_WRITE, + imageFormat, + imageWidth, + (imageHeight == 0) ? 1 : imageHeight, + (imageDepth == 0) ? 1 : imageDepth, + imageRowPitch, + imageSlicePitch); + break; + default: + ShouldNotReachHere(); + } + } else { + switch (imageType) { + case CL_MEM_OBJECT_IMAGE1D: + case CL_MEM_OBJECT_IMAGE2D: + case CL_MEM_OBJECT_IMAGE3D: + image = new (context) amd::Image(context, + imageType, + CL_MEM_READ_WRITE, + imageFormat, + imageWidth, + (imageHeight == 0) ? 1 : imageHeight, + (imageDepth == 0) ? 1 : imageDepth, + imageWidth * imageFormat.getElementSize(), /* row pitch */ + imageWidth * imageHeight * imageFormat.getElementSize(), /* slice pitch */ + numMipLevels); + break; + case CL_MEM_OBJECT_IMAGE1D_ARRAY: + image = new (context) amd::Image(context, + imageType, + CL_MEM_READ_WRITE, + imageFormat, + imageWidth, + imageArraySize, + 1, /* image depth */ + imageWidth * imageFormat.getElementSize(), + imageWidth * imageHeight * imageFormat.getElementSize(), + numMipLevels); + break; + case CL_MEM_OBJECT_IMAGE2D_ARRAY: + image = new (context) amd::Image(context, + imageType, + CL_MEM_READ_WRITE, + imageFormat, + imageWidth, + imageHeight, + imageArraySize, + imageWidth * imageFormat.getElementSize(), + imageWidth * imageHeight * imageFormat.getElementSize(), + numMipLevels); + break; + default: + ShouldNotReachHere(); + } + } + + if (image == nullptr) { + return nullptr; + } + + if (!image->create(nullptr)) { + delete image; + return nullptr; + } + + return image; +} + +hipError_t ihipArrayCreate(hipArray** array, + const HIP_ARRAY3D_DESCRIPTOR* pAllocateArray, + unsigned int numMipmapLevels) { + // NumChannels specifies the number of packed components per HIP array element; it may be 1, 2, or 4; + if ((pAllocateArray->NumChannels != 1) && + (pAllocateArray->NumChannels != 2) && + (pAllocateArray->NumChannels != 4)) { + return hipErrorInvalidValue; + } + + if ((pAllocateArray->Flags & hipArraySurfaceLoadStore) || + (pAllocateArray->Flags & hipArrayCubemap) || + (pAllocateArray->Flags & hipArrayTextureGather)) { + return hipErrorNotSupported; + } + + const cl_channel_order channelOrder = hip::getCLChannelOrder(pAllocateArray->NumChannels); + const cl_channel_type channelType = hip::getCLChannelType(pAllocateArray->Format, hipReadModeElementType); + const cl_mem_object_type imageType = hip::getCLMemObjectType(pAllocateArray->Width, + pAllocateArray->Height, + pAllocateArray->Depth, + pAllocateArray->Flags); + + amd::Image* image = ihipImageCreate(channelOrder, + channelType, + imageType, + pAllocateArray->Width, + pAllocateArray->Height, + pAllocateArray->Depth, + // The number of layers is determined by the depth extent. + pAllocateArray->Depth, /* array size */ + 0, /* row pitch */ + 0, /* slice pitch */ + numMipmapLevels, + nullptr /* buffer */); + + if (image == nullptr) { + return hipErrorInvalidValue; + } + + cl_mem memObj = as_cl(image); + *array = new hipArray{reinterpret_cast(memObj)}; + + // It is UB to call hipGet*() on an array created via hipArrayCreate()/hipArray3DCreate(). + // This is due to hip not differentiating between runtime and driver types. + // TODO change the hipArray struct in driver_types.h. + (*array)->desc = hip::getChannelFormatDesc(pAllocateArray->NumChannels, pAllocateArray->Format); + (*array)->width = pAllocateArray->Width; + (*array)->height = pAllocateArray->Height; + (*array)->depth = pAllocateArray->Depth; + (*array)->Format = pAllocateArray->Format; + (*array)->NumChannels = pAllocateArray->NumChannels; + + return hipSuccess; +} + +hipError_t hipArrayCreate(hipArray** array, + const HIP_ARRAY_DESCRIPTOR* pAllocateArray) { HIP_INIT_API(hipArrayCreate, array, pAllocateArray); - if (pAllocateArray->Width == 0) { - HIP_RETURN(hipErrorInvalidValue); - } + HIP_ARRAY3D_DESCRIPTOR desc = {pAllocateArray->Width, + pAllocateArray->Height, + 0, /* Depth */ + pAllocateArray->Format, + pAllocateArray->NumChannels, + hipArrayDefault /* Flags */}; - *array = (hipArray*)malloc(sizeof(hipArray)); - 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 }; - setDescFromChannelType(channelType, &(array[0]->desc)); - - size_t pitch = 0; - hipError_t status = ihipMallocPitch(ptr, &pitch, array[0]->width, array[0]->height, 1, CL_MEM_OBJECT_IMAGE2D, - &image_format); - - HIP_RETURN(status); + HIP_RETURN(ihipArrayCreate(array, &desc, 0)); } -hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, - size_t width, size_t height, unsigned int flags) { + +hipError_t hipMallocArray(hipArray** array, + const hipChannelFormatDesc* desc, + size_t width, + size_t height, + unsigned int flags) { HIP_INIT_API(hipMallocArray, array, desc, width, height, flags); - if (width == 0) { - HIP_RETURN(hipErrorInvalidValue); - } + HIP_ARRAY3D_DESCRIPTOR allocateArray = {width, + height, + 0, /* Depth */ + hip::getArrayFormat(*desc), + hip::getNumChannels(*desc), + flags}; - *array = (hipArray*)malloc(sizeof(hipArray)); - array[0]->type = flags; - array[0]->width = width; - array[0]->height = height; - array[0]->depth = 1; - array[0]->desc = *desc; - array[0]->isDrv = false; - array[0]->textureType = hipTextureType2D; - void** ptr = &array[0]->data; - - cl_channel_order channelOrder; - cl_channel_type channelType; - getChannelOrderAndType(*desc, hipReadModeElementType, &channelOrder, &channelType); - - const cl_image_format image_format = { channelOrder, channelType }; - - // Dummy flags check - switch (flags) { - case hipArrayLayered: - case hipArrayCubemap: - case hipArraySurfaceLoadStore: - case hipArrayTextureGather: - assert(0 && "Unspported"); - break; - case hipArrayDefault: - default: - break; - } - size_t pitch = 0; - hipError_t status = ihipMallocPitch(ptr, &pitch, width, height, 1, CL_MEM_OBJECT_IMAGE2D, - &image_format); - - HIP_RETURN(status); + HIP_RETURN(ihipArrayCreate(array, &allocateArray, 0 /* numMipLevels */)); } -hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAllocateArray) { +hipError_t hipArray3DCreate(hipArray** array, + const HIP_ARRAY3D_DESCRIPTOR* pAllocateArray) { HIP_INIT_API(hipArray3DCreate, array, pAllocateArray); - *array = (hipArray*)malloc(sizeof(hipArray)); - array[0]->type = pAllocateArray->Flags; - array[0]->width = pAllocateArray->Width; - array[0]->height = pAllocateArray->Height; - array[0]->depth = pAllocateArray->Depth; - array[0]->Format = pAllocateArray->Format; - array[0]->NumChannels = pAllocateArray->NumChannels; - array[0]->isDrv = true; - array[0]->textureType = hipTextureType3D; - 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 pitch = 0; - hipError_t status = ihipMallocPitch(ptr, &pitch, array[0]->width, array[0]->height, array[0]->depth, CL_MEM_OBJECT_IMAGE3D, - &image_format); - - HIP_RETURN(status); + HIP_RETURN(ihipArrayCreate(array, pAllocateArray, 0 /* numMipLevels */)); } -hipError_t hipMalloc3DArray(hipArray_t* array, const struct hipChannelFormatDesc* desc, - struct hipExtent extent, unsigned int flags) { +hipError_t hipMalloc3DArray(hipArray_t* array, + const hipChannelFormatDesc* desc, + hipExtent extent, + unsigned int flags) { + // TODO overload operator<<(ostream&, hipExtent&). HIP_INIT_API(hipMalloc3DArray, array, desc, &extent, flags); - *array = (hipArray*)malloc(sizeof(hipArray)); - array[0]->type = flags; - array[0]->width = extent.width; - array[0]->height = extent.height; - array[0]->depth = extent.depth; - array[0]->desc = *desc; - array[0]->isDrv = false; - array[0]->textureType = hipTextureType3D; - void** ptr = &array[0]->data; + HIP_ARRAY3D_DESCRIPTOR allocateArray = {extent.width, + extent.height, + extent.depth, + hip::getArrayFormat(*desc), + hip::getNumChannels(*desc), + flags}; - cl_channel_order channelOrder; - cl_channel_type channelType; - getChannelOrderAndType(*desc, hipReadModeElementType, &channelOrder, &channelType); - - const cl_image_format image_format = { channelOrder, channelType }; - - // Dummy flags check - switch (flags) { - case hipArrayCubemap: - case hipArraySurfaceLoadStore: - case hipArrayTextureGather: - assert(0 && "Unspported"); - break; - case hipArrayLayered: - case hipArrayDefault: - default: - break; - } - size_t pitch = 0; - hipError_t status = ihipMallocPitch(ptr, &pitch, extent.width, extent.height, extent.depth, - CL_MEM_OBJECT_IMAGE3D, &image_format); - - HIP_RETURN(status); + HIP_RETURN(ihipArrayCreate(array, &allocateArray, 0)); } hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) { @@ -708,40 +794,28 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t co HIP_RETURN(hipMemcpyAsync(dst, device_ptr, count, kind, stream)); } -hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) { - HIP_INIT_API(hipMemcpyHtoD, dst, src, sizeBytes); +hipError_t hipMemcpyHtoD(hipDeviceptr_t dstDevice, + void* srcHost, + size_t ByteCount) { + HIP_INIT_API(hipMemcpyHtoD, dstDevice, srcHost, ByteCount); - hip::syncStreams(); - amd::HostQueue* queue = hip::getNullStream(); - - HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyHostToDevice, *queue)); + HIP_RETURN(ihipMemcpy(dstDevice, srcHost, ByteCount, hipMemcpyHostToDevice, *hip::getQueue(nullptr))); } -hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) { - HIP_INIT_API(hipMemcpyDtoH, dst, src, sizeBytes); +hipError_t hipMemcpyDtoH(void* dstHost, + hipDeviceptr_t srcDevice, + size_t ByteCount) { + HIP_INIT_API(hipMemcpyDtoH, dstHost, srcDevice, ByteCount); - hip::syncStreams(); - amd::HostQueue* queue = hip::getNullStream(); - - HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToHost, *queue)); + HIP_RETURN(ihipMemcpy(dstHost, srcDevice, ByteCount, hipMemcpyDeviceToHost, *hip::getQueue(nullptr))); } -hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes) { - HIP_INIT_API(hipMemcpyDtoD, dst, src, sizeBytes); +hipError_t hipMemcpyDtoD(hipDeviceptr_t dstDevice, + hipDeviceptr_t srcDevice, + size_t ByteCount) { + HIP_INIT_API(hipMemcpyDtoD, dstDevice, srcDevice, ByteCount); - hip::syncStreams(); - amd::HostQueue* queue = hip::getNullStream(); - - HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToDevice, *queue)); -} - -hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) { - HIP_INIT_API(NONE, dst, src, sizeBytes); - - hip::syncStreams(); - amd::HostQueue* queue = hip::getNullStream(); - - HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyHostToHost, *queue)); + HIP_RETURN(ihipMemcpy(dstDevice, srcDevice, ByteCount, hipMemcpyDeviceToDevice, *hip::getQueue(nullptr))); } hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, @@ -753,90 +827,69 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, HIP_RETURN(ihipMemcpy(dst, src, sizeBytes, kind, *queue, true)); } - -hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, +hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dstDevice, + void* srcHost, + size_t ByteCount, hipStream_t stream) { - HIP_INIT_API(hipMemcpyHtoDAsync, dst, src, sizeBytes, stream); + HIP_INIT_API(hipMemcpyHtoDAsync, dstDevice, srcHost, ByteCount, stream); - amd::HostQueue* queue = hip::getQueue(stream); - - HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyHostToDevice, - *queue, true)); + HIP_RETURN(ihipMemcpy(dstDevice, srcHost, ByteCount, hipMemcpyHostToDevice, *hip::getQueue(stream), true)); } -hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes, +hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dstDevice, + hipDeviceptr_t srcDevice, + size_t ByteCount, hipStream_t stream) { - HIP_INIT_API(hipMemcpyDtoDAsync, dst, src, sizeBytes, stream); + HIP_INIT_API(hipMemcpyDtoDAsync, dstDevice, srcDevice, ByteCount, stream); - amd::HostQueue* queue = hip::getQueue(stream); - - HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToDevice, - *queue, true)); + HIP_RETURN(ihipMemcpy(dstDevice, srcDevice, ByteCount, hipMemcpyDeviceToDevice, *hip::getQueue(stream), true)); } -hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, +hipError_t hipMemcpyDtoHAsync(void* dstHost, + hipDeviceptr_t srcDevice, + size_t ByteCount, hipStream_t stream) { - HIP_INIT_API(hipMemcpyDtoHAsync, dst, src, sizeBytes, stream); + HIP_INIT_API(hipMemcpyDtoHAsync, dstHost, srcDevice, ByteCount, stream); - amd::HostQueue* queue = hip::getQueue(stream); - - HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToHost, - *queue, true)); + HIP_RETURN(ihipMemcpy(dstHost, srcDevice, ByteCount, hipMemcpyDeviceToHost, *hip::getQueue(stream), true)); } -hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, - size_t height, hipMemcpyKind kind, amd::HostQueue& queue, - bool isAsync = false) { - // Create buffer rectangle info structure - amd::BufferRect srcRect; - amd::BufferRect dstRect; +hipError_t ihipMemcpyAtoD(hipArray* srcArray, + void* dstDevice, + amd::Coord3D srcOrigin, + amd::Coord3D dstOrigin, + amd::Coord3D region, + size_t dstRowPitch, + size_t dstSlicePitch, + hipStream_t stream, + bool isAsync = false) { + // TODO VDI doesn't support 2D/3D image to buffer copy. + (void)dstRowPitch; + (void)dstSlicePitch; - size_t region[3] = {width, height, 1}; - size_t src_slice_pitch = spitch * height; - size_t dst_slice_pitch = dpitch * height; - size_t sOrigin[3] = { }; - size_t dOrigin[3] = { }; - amd::Memory* srcMemory = getMemoryObject(src, sOrigin[0]); - amd::Memory* dstMemory = getMemoryObject(dst, dOrigin[0]); - - if (src_slice_pitch == 0 || - dst_slice_pitch == 0 || - dst == nullptr || - src == nullptr) { - return hipSuccess; - } - - if (!srcRect.create(sOrigin, region, spitch, src_slice_pitch) || - !dstRect.create(dOrigin, region, dpitch, dst_slice_pitch)) { + cl_mem srcMemObj = reinterpret_cast(srcArray->data); + if (is_valid(srcMemObj) == false) { return hipErrorInvalidValue; } - amd::Command* command = nullptr; - amd::Command::EventWaitList waitList; + amd::Image* srcImage = as_amd(srcMemObj)->asImage(); + size_t offset = 0; + amd::Memory* dstMemory = getMemoryObject(dstDevice, offset); + assert(offset != 0); - amd::Coord3D srcStart(srcRect.start_, 0, 0); - amd::Coord3D dstStart(dstRect.start_, 0, 0); - amd::Coord3D size(region[0], region[1], region[2]); - - if (((srcMemory == nullptr) && (dstMemory == nullptr)) || - (kind == hipMemcpyHostToHost)) { - for(unsigned int y = 0; y < height; y++) { - void* pDst = reinterpret_cast(reinterpret_cast(dst) + y * dpitch); - void* pSrc = reinterpret_cast(reinterpret_cast(src) + y * spitch); - memcpy(pDst, pSrc, width); - } - return hipSuccess; - } else if ((srcMemory == nullptr) && (dstMemory != nullptr)) { - command = new amd::WriteMemoryCommand(queue, CL_COMMAND_WRITE_BUFFER_RECT, waitList, - *dstMemory->asBuffer(), dstStart, size, src, dstRect, srcRect); - } else if ((srcMemory != nullptr) && (dstMemory == nullptr)) { - command = new amd::ReadMemoryCommand(queue, CL_COMMAND_READ_BUFFER_RECT, waitList, - *srcMemory->asBuffer(), srcStart, size, dst, srcRect, dstRect); - } else if ((srcMemory != nullptr) && (dstMemory != nullptr)) { - command = new amd::CopyMemoryCommand(queue, CL_COMMAND_COPY_BUFFER_RECT, waitList, *srcMemory->asBuffer(), - *dstMemory->asBuffer(), srcStart, dstStart, size, srcRect, dstRect); + if (!srcImage->validateRegion(srcOrigin, region)) { + return hipErrorInvalidValue; } + amd::CopyMemoryCommand* command = new amd::CopyMemoryCommand(*hip::getQueue(stream), + CL_COMMAND_COPY_IMAGE_TO_BUFFER, + amd::Command::EventWaitList{}, + *srcImage, + *dstMemory, + srcOrigin, + dstOrigin, + region); + if (command == nullptr) { return hipErrorOutOfMemory; } @@ -850,578 +903,686 @@ hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch return hipSuccess; } +hipError_t ihipMemcpyDtoA(void* srcDevice, + hipArray* dstArray, + amd::Coord3D srcOrigin, + amd::Coord3D dstOrigin, + amd::Coord3D region, + size_t srcRowPitch, + size_t srcSlicePitch, + hipStream_t stream, + bool isAsync = false) { + // TODO VDI doesn't support 2D/3D buffer to image copy. + (void)srcRowPitch; + (void)srcSlicePitch; + + cl_mem dstMemObj = reinterpret_cast(dstArray->data); + if (is_valid(dstMemObj) == false) { + return hipErrorInvalidValue; + } + + size_t offset = 0; + amd::Memory* srcMemory = getMemoryObject(srcDevice, offset); + assert(offset != 0); + amd::Image* dstImage = as_amd(dstMemObj)->asImage(); + + const size_t copySizeInBytes = region[0] * region[1] * region[2] * dstImage->getImageFormat().getElementSize(); + if (!srcMemory->validateRegion(srcOrigin, {copySizeInBytes, 0, 0}) || + !dstImage->validateRegion(dstOrigin, region)) { + return hipErrorInvalidValue; + } + + amd::CopyMemoryCommand* command = new amd::CopyMemoryCommand(*hip::getQueue(stream), + CL_COMMAND_COPY_BUFFER_TO_IMAGE, + amd::Command::EventWaitList{}, + *srcMemory, + *dstImage, + srcOrigin, + dstOrigin, + region); + + if (command == nullptr) { + return hipErrorOutOfMemory; + } + + command->enqueue(); + if (!isAsync) { + command->awaitCompletion(); + } + command->release(); + + return hipSuccess; +} + +hipError_t ihipMemcpyDtoD(void* srcDevice, + void* dstDevice, + amd::Coord3D srcOrigin, + amd::Coord3D dstOrigin, + amd::Coord3D copyRegion, + size_t srcRowPitch, + size_t srcSlicePitch, + size_t dstRowPitch, + size_t dstSlicePitch, + hipStream_t stream, + bool isAsync = false) { + size_t srcOffset = 0; + amd::Memory *srcMemory = getMemoryObject(srcDevice, srcOffset); + size_t dstOffset = 0; + amd::Memory *dstMemory = getMemoryObject(dstDevice, dstOffset); + + amd::BufferRect srcRect; + if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), srcRowPitch, srcSlicePitch)) { + return hipErrorInvalidValue; + } + + amd::Coord3D srcStart(srcRect.start_ + srcOffset, 0, 0); + amd::Coord3D srcEnd(srcRect.end_ + srcOffset, 1, 1); + if (!srcMemory->validateRegion(srcStart, srcEnd)) { + return hipErrorInvalidValue; + } + + amd::BufferRect dstRect; + if (!dstRect.create(static_cast(dstOrigin), static_cast(copyRegion), dstRowPitch, dstSlicePitch)) { + return hipErrorInvalidValue; + } + + amd::Coord3D dstStart(dstRect.start_ + dstOffset, 0, 0); + amd::Coord3D dstEnd(dstRect.end_ + dstOffset, 1, 1); + if (!dstMemory->validateRegion(dstStart, dstEnd)) { + return hipErrorInvalidValue; + } + + amd::CopyMemoryCommand* command = new amd::CopyMemoryCommand(*hip::getQueue(stream), + CL_COMMAND_COPY_BUFFER_RECT, + amd::Command::EventWaitList{}, + *srcMemory, + *dstMemory, + srcStart, + dstStart, + copyRegion, + srcRect, + dstRect); + + if (command == nullptr) { + return hipErrorOutOfMemory; + } + + command->enqueue(); + if (!isAsync) { + command->awaitCompletion(); + } + command->release(); + + return hipSuccess; +} + +hipError_t ihipMemcpyDtoH(void* srcDevice, + void* dstHost, + amd::Coord3D srcOrigin, + amd::Coord3D dstOrigin, + amd::Coord3D copyRegion, + size_t srcRowPitch, + size_t srcSlicePitch, + size_t dstRowPitch, + size_t dstSlicePitch, + hipStream_t stream, + bool isAsync = false) { + size_t srcOffset = 0; + amd::Memory *srcMemory = getMemoryObject(srcDevice, srcOffset); + + amd::BufferRect srcRect; + if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), srcRowPitch, srcSlicePitch)) { + return hipErrorInvalidValue; + } + + amd::Coord3D srcStart(srcRect.start_ + srcOffset, 0, 0); + amd::Coord3D srcEnd(srcRect.end_ + srcOffset, 1, 1); + if (!srcMemory->validateRegion(srcStart, srcEnd)) { + return hipErrorInvalidValue; + } + + amd::BufferRect dstRect; + if (!dstRect.create(static_cast(dstOrigin), static_cast(copyRegion), dstRowPitch, dstSlicePitch)) { + return hipErrorInvalidValue; + } + + amd::ReadMemoryCommand* command = new amd::ReadMemoryCommand(*hip::getQueue(stream), + CL_COMMAND_READ_BUFFER_RECT, + amd::Command::EventWaitList{}, + *srcMemory, + srcStart, + copyRegion, + dstHost, + srcRect, + dstRect); + + if (command == nullptr) { + return hipErrorOutOfMemory; + } + + command->enqueue(); + if (!isAsync) { + command->awaitCompletion(); + } + command->release(); + + return hipSuccess; +} + +hipError_t ihipMemcpyHtoD(const void* srcHost, + void* dstDevice, + amd::Coord3D srcOrigin, + amd::Coord3D dstOrigin, + amd::Coord3D copyRegion, + size_t srcRowPitch, + size_t srcSlicePitch, + size_t dstRowPitch, + size_t dstSlicePitch, + hipStream_t stream, + bool isAsync = false) { + size_t dstOffset = 0; + amd::Memory *dstMemory = getMemoryObject(dstDevice, dstOffset); + + amd::BufferRect srcRect; + if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), srcRowPitch, srcSlicePitch)) { + return hipErrorInvalidValue; + } + + amd::BufferRect dstRect; + if (!dstRect.create(static_cast(dstOrigin), static_cast(copyRegion), dstRowPitch, dstSlicePitch)) { + return hipErrorInvalidValue; + } + + amd::Coord3D dstStart(dstRect.start_ + dstOffset, 0, 0); + amd::Coord3D dstEnd(dstRect.end_ + dstOffset, 1, 1); + if (!dstMemory->validateRegion(dstStart, dstEnd)) { + return hipErrorInvalidValue; + } + + amd::WriteMemoryCommand* command = new amd::WriteMemoryCommand(*hip::getQueue(stream), + CL_COMMAND_WRITE_BUFFER_RECT, + amd::Command::EventWaitList{}, + *dstMemory, + dstStart, + copyRegion, + srcHost, + dstRect, + srcRect); + + if (command == nullptr) { + return hipErrorOutOfMemory; + } + + command->enqueue(); + if (!isAsync) { + command->awaitCompletion(); + } + command->release(); + + return hipSuccess; +} + +hipError_t ihipMemcpyHtoH(const void* srcHost, + void* dstHost, + amd::Coord3D srcOrigin, + amd::Coord3D dstOrigin, + amd::Coord3D copyRegion, + size_t srcRowPitch, + size_t srcSlicePitch, + size_t dstRowPitch, + size_t dstSlicePitch) { + amd::BufferRect srcRect; + if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), srcRowPitch, srcSlicePitch)) { + return hipErrorInvalidValue; + } + + amd::BufferRect dstRect; + if (!dstRect.create(static_cast(dstOrigin), static_cast(copyRegion), dstRowPitch, dstSlicePitch)) { + return hipErrorInvalidValue; + } + + for (size_t slice = 0; slice < copyRegion[2]; slice++) { + for (size_t row = 0; row < copyRegion[1]; row++) { + const void* srcRow = static_cast(srcHost) + srcRect.start_ + row * srcRect.rowPitch_ + slice * srcRect.slicePitch_; + void* dstRow = static_cast(dstHost) + dstRect.start_ + row * dstRect.rowPitch_ + slice * dstRect.slicePitch_; + std::memcpy(dstRow, srcRow, copyRegion[0]); + } + } + + return hipSuccess; +} + +hipError_t ihipMemcpyAtoA(hipArray* srcArray, + hipArray* dstArray, + amd::Coord3D srcOrigin, + amd::Coord3D dstOrigin, + amd::Coord3D region, + hipStream_t stream, + bool isAsync = false) { + cl_mem srcMemObj = reinterpret_cast(srcArray->data); + cl_mem dstMemObj = reinterpret_cast(dstArray->data); + if (!is_valid(srcMemObj) || !is_valid(dstMemObj)) { + return hipErrorInvalidValue; + } + + amd::Image* srcImage = as_amd(srcMemObj)->asImage(); + amd::Image* dstImage = as_amd(dstMemObj)->asImage(); + + if (!srcImage->validateRegion(srcOrigin, region) || + !dstImage->validateRegion(dstOrigin, region)) { + return hipErrorInvalidValue; + } + + amd::CopyMemoryCommand* command = new amd::CopyMemoryCommand(*hip::getQueue(stream), + CL_COMMAND_COPY_IMAGE, + amd::Command::EventWaitList{}, + *srcImage, + *dstImage, + srcOrigin, + dstOrigin, + region); + + if (command == nullptr) { + return hipErrorOutOfMemory; + } + + command->enqueue(); + if (!isAsync) { + command->awaitCompletion(); + } + command->release(); + + return hipSuccess; +} + +hipError_t ihipMemcpyHtoA(const void* srcHost, + hipArray* dstArray, + amd::Coord3D srcOrigin, + amd::Coord3D dstOrigin, + amd::Coord3D copyRegion, + size_t srcRowPitch, + size_t srcSlicePitch, + hipStream_t stream, + bool isAsync = false) { + if (srcHost == nullptr) { + return hipErrorInvalidValue; + } + + cl_mem dstMemObj = reinterpret_cast(dstArray->data); + if (is_valid(dstMemObj) == false) { + return hipErrorInvalidValue; + } + + amd::BufferRect srcRect; + if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), srcRowPitch, srcSlicePitch)) { + return hipErrorInvalidValue; + } + + amd::Image* dstImage = as_amd(dstMemObj)->asImage(); + // HIP assumes the width is in bytes, but OCL assumes it's in pixels. + const size_t elementSize = dstImage->getImageFormat().getElementSize(); + static_cast(dstOrigin)[0] /= elementSize; + static_cast(copyRegion)[0] /= elementSize; + + if (!dstImage->validateRegion(dstOrigin, copyRegion)) { + return hipErrorInvalidValue; + } + + amd::WriteMemoryCommand* command = new amd::WriteMemoryCommand(*hip::getQueue(stream), + CL_COMMAND_WRITE_IMAGE, + amd::Command::EventWaitList{}, + *dstImage, + dstOrigin, + copyRegion, + static_cast(srcHost) + srcRect.start_, + srcRowPitch, + srcSlicePitch); + + if (command == nullptr) { + return hipErrorOutOfMemory; + } + + command->enqueue(); + if (!isAsync) { + command->awaitCompletion(); + } + command->release(); + + return hipSuccess; +} + +hipError_t ihipMemcpyAtoH(hipArray* srcArray, + void* dstHost, + amd::Coord3D srcOrigin, + amd::Coord3D dstOrigin, + amd::Coord3D copyRegion, + size_t dstRowPitch, + size_t dstSlicePitch, + hipStream_t stream, + bool isAsync = false) { + cl_mem srcMemObj = reinterpret_cast(srcArray->data); + if (!is_valid(srcMemObj)) { + return hipErrorInvalidValue; + } + + if (dstHost == nullptr) { + return hipErrorInvalidValue; + } + + amd::BufferRect dstRect; + if (!dstRect.create(static_cast(dstOrigin), static_cast(copyRegion), dstRowPitch, dstSlicePitch)) { + return hipErrorInvalidValue; + } + + + amd::Image* srcImage = as_amd(srcMemObj)->asImage(); + // HIP assumes the width is in bytes, but OCL assumes it's in pixels. + const size_t elementSize = srcImage->getImageFormat().getElementSize(); + static_cast(srcOrigin)[0] /= elementSize; + static_cast(copyRegion)[0] /= elementSize; + + if (!srcImage->validateRegion(srcOrigin, copyRegion) || + !srcImage->isRowSliceValid(dstRowPitch, dstSlicePitch, copyRegion[0], copyRegion[1])) { + return hipErrorInvalidValue; + } + + amd::ReadMemoryCommand* command = new amd::ReadMemoryCommand(*hip::getQueue(stream), + CL_COMMAND_READ_IMAGE, + amd::Command::EventWaitList{}, + *srcImage, + srcOrigin, + copyRegion, + static_cast(dstHost) + dstRect.start_, + dstRowPitch, + dstSlicePitch); + + if (command == nullptr) { + return hipErrorOutOfMemory; + } + + command->enqueue(); + if (!isAsync) { + command->awaitCompletion(); + } + command->release(); + + return hipSuccess; +} + +hipError_t ihipMemcpyParam3D(const hipMemcpy3DParms* pCopy, + hipStream_t stream, + bool isAsync = false) { + // If {src/dst}MemoryType is hipMemoryTypeUnified, {src/dst}Device and {src/dst}Pitch specify the (unified virtual address space) + // base address of the source data and the bytes per row to apply. {src/dst}Array is ignored. + hipMemoryType srcMemoryType = pCopy->srcMemoryType; + if (srcMemoryType == hipMemoryTypeUnified) { + srcMemoryType = amd::MemObjMap::FindMemObj(pCopy->srcDevice) ? hipMemoryTypeDevice : hipMemoryTypeHost; + } + hipMemoryType dstMemoryType = pCopy->dstMemoryType; + if (dstMemoryType == hipMemoryTypeUnified) { + dstMemoryType = amd::MemObjMap::FindMemObj(pCopy->dstDevice) ? hipMemoryTypeDevice : hipMemoryTypeHost; + } + + amd::Coord3D srcOrigin = {pCopy->srcXInBytes, pCopy->srcY, pCopy->srcZ}; + amd::Coord3D dstOrigin = {pCopy->dstXInBytes, pCopy->dstY, pCopy->dstZ}; + amd::Coord3D copyRegion = {pCopy->WidthInBytes, (pCopy->Height != 0) ? pCopy->Height : 1, (pCopy->Depth != 0) ? pCopy->Depth : 1}; + + if ((srcMemoryType == hipMemoryTypeHost) && (dstMemoryType == hipMemoryTypeHost)) { + // Host to Host. + return ihipMemcpyHtoH(pCopy->srcHost, pCopy->dstHost, srcOrigin, dstOrigin, copyRegion, pCopy->srcPitch, pCopy->srcPitch * pCopy->srcHeight, pCopy->dstPitch, pCopy->dstPitch * pCopy->dstHeight); + } else if ((srcMemoryType == hipMemoryTypeHost) && (dstMemoryType == hipMemoryTypeDevice)) { + // Host to Device. + return ihipMemcpyHtoD(pCopy->srcHost, pCopy->dstDevice, srcOrigin, dstOrigin, copyRegion, pCopy->srcPitch, pCopy->srcPitch * pCopy->srcHeight, pCopy->dstPitch, pCopy->dstPitch * pCopy->dstHeight, stream, isAsync); + } else if ((srcMemoryType == hipMemoryTypeDevice) && (dstMemoryType == hipMemoryTypeHost)) { + // Device to Host. + return ihipMemcpyDtoH(pCopy->srcDevice, pCopy->dstHost, srcOrigin, dstOrigin, copyRegion, pCopy->srcPitch, pCopy->srcPitch * pCopy->srcHeight, pCopy->dstPitch, pCopy->dstPitch * pCopy->dstHeight, stream, isAsync); + } else if ((srcMemoryType == hipMemoryTypeDevice) && (dstMemoryType == hipMemoryTypeDevice)) { + // Device to Device. + return ihipMemcpyDtoD(pCopy->srcDevice, pCopy->dstDevice, srcOrigin, dstOrigin, copyRegion, pCopy->srcPitch, pCopy->srcPitch * pCopy->srcHeight, pCopy->dstPitch, pCopy->dstPitch * pCopy->dstHeight, stream, isAsync); + } else if ((srcMemoryType == hipMemoryTypeHost) && (dstMemoryType == hipMemoryTypeArray)) { + // Host to Image. + return ihipMemcpyHtoA(pCopy->srcHost, pCopy->dstArray, srcOrigin, dstOrigin, copyRegion, pCopy->srcPitch, pCopy->srcPitch * pCopy->srcHeight, stream, isAsync); + } else if ((srcMemoryType == hipMemoryTypeArray) && (dstMemoryType == hipMemoryTypeHost)) { + // Image to Host. + return ihipMemcpyAtoH(pCopy->srcArray, pCopy->dstHost, srcOrigin, dstOrigin, copyRegion, pCopy->dstPitch, pCopy->dstPitch * pCopy->dstHeight, stream, isAsync); + } else if ((srcMemoryType == hipMemoryTypeDevice) && (dstMemoryType == hipMemoryTypeArray)) { + // Device to Image. + return ihipMemcpyDtoA(pCopy->srcDevice, pCopy->dstArray, srcOrigin, dstOrigin, copyRegion, pCopy->srcPitch, pCopy->srcPitch * pCopy->srcHeight, stream, isAsync); + } else if ((srcMemoryType == hipMemoryTypeArray) && (dstMemoryType == hipMemoryTypeDevice)) { + // Image to Device. + return ihipMemcpyAtoD(pCopy->srcArray, pCopy->dstDevice, srcOrigin, dstOrigin, copyRegion, pCopy->dstPitch, pCopy->dstPitch * pCopy->dstHeight, stream, isAsync); + } else if ((srcMemoryType == hipMemoryTypeArray) && (dstMemoryType == hipMemoryTypeArray)) { + // Image to Image. + return ihipMemcpyAtoA(pCopy->srcArray, pCopy->dstArray, srcOrigin, dstOrigin, copyRegion, stream, isAsync); + } else { + ShouldNotReachHere(); + } + + return hipSuccess; +} + +hipError_t ihipMemcpyParam2D(const hip_Memcpy2D* pCopy, + hipStream_t stream, + bool isAsync = false) { + hipMemcpy3DParms desc = hip::getMemcpy3DParms(*pCopy); + + return ihipMemcpyParam3D(&desc, stream, isAsync); +} + +hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, + size_t height, hipMemcpyKind kind, hipStream_t stream, bool isAsync = false) { + hip_Memcpy2D desc = {}; + + desc.srcXInBytes = 0; + desc.srcY = 0; + desc.srcMemoryType = std::get<0>(hip::getMemoryType(kind)); + desc.srcHost = src; + desc.srcDevice = const_cast(src); + desc.srcArray = nullptr; // Ignored. + desc.srcPitch = spitch; + + desc.dstXInBytes = 0; + desc.dstY = 0; + desc.dstMemoryType = std::get<1>(hip::getMemoryType(kind)); + desc.dstHost = dst; + desc.dstDevice = dst; + desc.dstArray = nullptr; // Ignored. + desc.dstPitch = dpitch; + + desc.WidthInBytes = width; + desc.Height = height; + + return ihipMemcpyParam2D(&desc, stream, isAsync); +} + hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) { HIP_INIT_API(hipMemcpyParam2D, pCopy); - hipError_t e = hipSuccess; - if (pCopy == nullptr) { - e = hipErrorInvalidValue; - } else { - hip::syncStreams(); - amd::HostQueue* queue = hip::getNullStream(); - e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->WidthInBytes, pCopy->srcHost, pCopy->srcPitch, - pCopy->WidthInBytes, pCopy->Height, hipMemcpyDefault, *queue); - } - HIP_RETURN(e); + + HIP_RETURN(ihipMemcpyParam2D(pCopy, nullptr)); } hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { HIP_INIT_API(hipMemcpy2D, dst, dpitch, src, spitch, width, height, kind); - hip::syncStreams(); - amd::HostQueue* queue = hip::getNullStream(); - - HIP_RETURN(ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind, *queue)); + HIP_RETURN(ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind, nullptr)); } - hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) { HIP_INIT_API(hipMemcpy2DAsync, dst, dpitch, src, spitch, width, height, kind, stream); - amd::HostQueue* queue = hip::getQueue(stream); - - HIP_RETURN(ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind, *queue, true)); + HIP_RETURN(ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind, stream, true)); } -hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, - size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { +hipError_t ihipMemcpy2DToArray(hipArray_t dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream, bool isAsync = false) { + hip_Memcpy2D desc = {}; + + desc.srcXInBytes = 0; + desc.srcY = 0; + desc.srcMemoryType = std::get<0>(hip::getMemoryType(kind)); + desc.srcHost = const_cast(src); + desc.srcDevice = const_cast(src); + desc.srcArray = nullptr; + desc.srcPitch = spitch; + + desc.dstXInBytes = wOffset; + desc.dstY = hOffset; + desc.dstMemoryType = hipMemoryTypeArray; + desc.dstHost = nullptr; + desc.dstDevice = nullptr; + desc.dstArray = dst; + desc.dstPitch = 0; // Ignored. + + desc.WidthInBytes = width; + desc.Height = height; + + return ihipMemcpyParam2D(&desc, stream, isAsync); +} + +hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { HIP_INIT_API(hipMemcpy2DToArray, dst, wOffset, hOffset, src, spitch, width, height, kind); - if (dst->data == nullptr) { + HIP_RETURN(ihipMemcpy2DToArray(dst, wOffset, hOffset, src, spitch, width, height, kind, nullptr)); +} + +hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t count, hipMemcpyKind kind) { + HIP_INIT_API(hipMemcpyToArray, dst, wOffset, hOffset, src, count, kind); + + if (dst == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - hip::syncStreams(); - amd::HostQueue* queue = hip::getNullStream(); + const size_t arrayHeight = (dst->height != 0) ? dst->height : 1; + const size_t witdthInBytes = count / arrayHeight; - size_t dpitch = dst->width; - getByteSizeFromChannelFormatKind(dst[0].desc.f, &dpitch); + const size_t height = (count / dst->width) / (hip::getElementSize(dst->Format) * dst->NumChannels); - if ((wOffset + width > (dpitch)) || width > spitch) { - HIP_RETURN(hipErrorInvalidDevicePointer); - } + HIP_RETURN(ihipMemcpy2DToArray(dst, wOffset, hOffset, src, 0 /* spitch */, witdthInBytes, height, kind, nullptr)); +} - // Create buffer rectangle info structure - amd::BufferRect srcRect; - amd::BufferRect dstRect; +hipError_t ihipMemcpy2DFromArray(void* dst, size_t dpitch, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream, bool isAsync = false) { + hip_Memcpy2D desc = {}; - size_t region[3] = {width, height, 1}; - size_t src_slice_pitch = spitch * height; - size_t dst_slice_pitch = dpitch * height; - size_t sOrigin[3] = { }; - size_t dOrigin[3] = {wOffset, hOffset, 0}; - size_t offset = 0; - amd::Memory* srcMemory = getMemoryObject(src, offset); - amd::Memory* dstMemory = getMemoryObject(dst->data, offset); + desc.srcXInBytes = wOffsetSrc; + desc.srcY = hOffsetSrc; + desc.srcMemoryType = hipMemoryTypeArray; + desc.srcHost = nullptr; + desc.srcDevice = nullptr; + desc.srcArray = const_cast(src); + desc.srcPitch = 0; // Ignored. - assert(offset == 0); + desc.dstXInBytes = 0; + desc.dstY = 0; + desc.dstMemoryType = std::get<1>(hip::getMemoryType(kind)); + desc.dstHost = dst; + desc.dstDevice = dst; + desc.dstArray = nullptr; + desc.dstPitch = dpitch; - if (!srcRect.create(sOrigin, region, spitch, src_slice_pitch) || - !dstRect.create(dOrigin, region, dpitch, dst_slice_pitch)) { + desc.WidthInBytes = width; + desc.Height = height; + + return ihipMemcpyParam2D(&desc, stream, isAsync); +} + +hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t src, size_t wOffsetSrc, size_t hOffset, size_t count, hipMemcpyKind kind) { + HIP_INIT_API(hipMemcpyFromArray, dst, src, wOffsetSrc, hOffset, count, kind); + + if (src == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - amd::Command* command = nullptr; - amd::Command::EventWaitList waitList; + const size_t arrayHeight = (src->height != 0) ? src->height : 1; + const size_t witdthInBytes = count / arrayHeight; - amd::Coord3D srcStart(srcRect.start_, 0, 0); - amd::Coord3D dstStart(dstRect.start_, 0, 0); - amd::Coord3D size(region[0], region[1], region[2]); - - if (((srcMemory == nullptr) && (dstMemory == nullptr)) || - (kind == hipMemcpyHostToHost)) { - void* newDst = reinterpret_cast(reinterpret_cast(dst->data) - + dpitch * hOffset + wOffset); - for(unsigned int y = 0; y < height; y++) { - void* pDst = reinterpret_cast(reinterpret_cast(newDst) + y * dpitch); - void* pSrc = reinterpret_cast(reinterpret_cast(src) + y * spitch); - memcpy(pDst, pSrc, width); - } - HIP_RETURN(hipSuccess); - } else if ((srcMemory == nullptr) && (dstMemory != nullptr)) { - command = new amd::WriteMemoryCommand(*queue, CL_COMMAND_WRITE_BUFFER_RECT, waitList, - *dstMemory->asBuffer(), dstStart, size, src, dstRect, srcRect); - } else if ((srcMemory != nullptr) && (dstMemory == nullptr)) { - command = new amd::ReadMemoryCommand(*queue, CL_COMMAND_READ_BUFFER_RECT, waitList, - *srcMemory->asBuffer(), srcStart, size, dst, srcRect, dstRect); - } else if ((srcMemory != nullptr) && (dstMemory != nullptr)) { - command = new amd::CopyMemoryCommand(*queue, CL_COMMAND_COPY_BUFFER_RECT, waitList, *srcMemory->asBuffer(), - *dstMemory->asBuffer(), srcStart, dstStart, size, srcRect, dstRect); - } - - if (command == nullptr) { - HIP_RETURN(hipErrorOutOfMemory); - } - - command->enqueue(); - command->awaitCompletion(); - command->release(); - - HIP_RETURN(hipSuccess); + const size_t height = (count / src->width) / (hip::getElementSize(src->Format) * src->NumChannels); + HIP_RETURN(ihipMemcpy2DFromArray(dst, 0 /* dpitch */, src, wOffsetSrc, hOffset, witdthInBytes, height, kind, nullptr)); } -hipError_t hipMemcpyToArray(hipArray* dstArray, size_t wOffset, size_t hOffset, const void* src, - size_t count, hipMemcpyKind kind) { - HIP_INIT_API(hipMemcpyToArray, dstArray, wOffset, hOffset, src, count, kind); +hipError_t hipMemcpyHtoA(hipArray* dstArray, + size_t dstOffset, + const void* srcHost, + size_t ByteCount) { + HIP_INIT_API(hipMemcpyHtoA, dstArray, dstOffset, srcHost, ByteCount); - hip::syncStreams(); - amd::HostQueue* queue = hip::getNullStream(); - - amd::Command* command = nullptr; - amd::Command::EventWaitList waitList; - - size_t sOffset = 0; - amd::Memory* srcMemory = getMemoryObject(src, sOffset); - size_t dOffset = 0; - amd::Memory* dstMemory = getMemoryObject(dstArray->data, dOffset); - - assert(dOffset == 0); - - assert((kind == hipMemcpyHostToDevice) && "Invalid case"); - - if ((srcMemory == nullptr) && (dstMemory != nullptr)) { - command = new amd::WriteMemoryCommand(*queue, CL_COMMAND_WRITE_BUFFER, waitList, - *dstMemory->asBuffer(), {wOffset, hOffset}, count, src); - } else if ((srcMemory != nullptr) && (dstMemory != nullptr)) { - command = new amd::CopyMemoryCommand(*queue, CL_COMMAND_COPY_BUFFER, waitList, - *srcMemory->asBuffer(),*dstMemory->asBuffer(), sOffset, {wOffset, hOffset}, count); - } else { - ShouldNotReachHere(); - } - - if (command == nullptr) { - HIP_RETURN(hipErrorOutOfMemory); - } - - command->enqueue(); - command->awaitCompletion(); - command->release(); - - HIP_RETURN(hipSuccess); + HIP_RETURN(ihipMemcpyHtoA(srcHost, dstArray, {0, 0, 0}, {dstOffset, 0, 0}, {ByteCount, 1, 1}, 0, 0, nullptr)); } -hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffset, size_t hOffset, - size_t count, hipMemcpyKind kind) { - HIP_INIT_API(hipMemcpyFromArray, dst, srcArray, wOffset, hOffset, count, kind); +hipError_t hipMemcpyAtoH(void* dstHost, + hipArray* srcArray, + size_t srcOffset, + size_t ByteCount) { + HIP_INIT_API(hipMemcpyAtoH, dstHost, srcArray, srcOffset, ByteCount); - hip::syncStreams(); - amd::HostQueue* queue = hip::getNullStream(); - - amd::Command* command = nullptr; - amd::Command::EventWaitList waitList; - - size_t sOffset = 0; - amd::Memory* srcMemory = getMemoryObject(srcArray->data, sOffset); - size_t dOffset = 0; - amd::Memory* dstMemory = getMemoryObject(dst, dOffset); - - assert(sOffset == 0); - - assert((kind == hipMemcpyDeviceToHost) && "Invalid case"); - - if ((srcMemory != nullptr) && (dstMemory == nullptr)) { - command = new amd::ReadMemoryCommand(*queue, CL_COMMAND_READ_BUFFER, waitList, - *srcMemory->asBuffer(), {wOffset, hOffset}, count, dst); - } else if ((srcMemory != nullptr) && (dstMemory != nullptr)) { - command = new amd::CopyMemoryCommand(*queue, CL_COMMAND_COPY_BUFFER, waitList, - *srcMemory->asBuffer(), *dstMemory->asBuffer(), {wOffset, hOffset}, dOffset, count); - } else { - ShouldNotReachHere(); - } - - if (command == nullptr) { - HIP_RETURN(hipErrorOutOfMemory); - } - - command->enqueue(); - command->awaitCompletion(); - command->release(); - - HIP_RETURN(hipSuccess); + HIP_RETURN(ihipMemcpyAtoH(srcArray, dstHost, {srcOffset, 0, 0}, {0, 0, 0}, {ByteCount, 1, 1}, 0, 0, nullptr)); } -hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHost, size_t count) { - HIP_INIT_API(hipMemcpyHtoA, dstArray, dstOffset, srcHost, count); - - hip::syncStreams(); - amd::HostQueue* queue = hip::getNullStream(); - - amd::Command* command = nullptr; - amd::Command::EventWaitList waitList; - - size_t sOffset = 0; - amd::Memory* srcMemory = getMemoryObject(srcHost, sOffset); - size_t dOffset = 0; - amd::Memory* dstMemory = getMemoryObject(dstArray->data, dOffset); - - assert(dOffset == 0); - - if ((srcMemory == nullptr) && (dstMemory != nullptr)) { - command = new amd::WriteMemoryCommand(*queue, CL_COMMAND_WRITE_BUFFER, waitList, - *dstMemory->asBuffer(), dstOffset, count, srcHost); - } else if ((srcMemory != nullptr) && (dstMemory != nullptr)) { - command = new amd::CopyMemoryCommand(*queue, CL_COMMAND_COPY_BUFFER, waitList, - *srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dstOffset, count); - } else { - ShouldNotReachHere(); - } - - if (command == nullptr) { - HIP_RETURN(hipErrorOutOfMemory); - } - - command->enqueue(); - command->awaitCompletion(); - command->release(); - - HIP_RETURN(hipSuccess); -} - -hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t count) { - HIP_INIT_API(hipMemcpyAtoH, dst, srcArray, srcOffset, count); - - hip::syncStreams(); - amd::HostQueue* queue = hip::getNullStream(); - - amd::Command* command = nullptr; - amd::Command::EventWaitList waitList; - - size_t sOffset = 0; - amd::Memory* srcMemory = getMemoryObject(srcArray->data, sOffset); - size_t dOffset = 0; - amd::Memory* dstMemory = getMemoryObject(dst, dOffset); - - assert(sOffset == 0); - - if ((srcMemory != nullptr) && (dstMemory == nullptr)) { - command = new amd::ReadMemoryCommand(*queue, CL_COMMAND_READ_BUFFER, waitList, - *srcMemory->asBuffer(), srcOffset, count, dst); - } else if ((srcMemory != nullptr) && (dstMemory != nullptr)) { - command = new amd::CopyMemoryCommand(*queue, CL_COMMAND_COPY_BUFFER, waitList, - *srcMemory->asBuffer(), *dstMemory->asBuffer(), srcOffset, dOffset, count); - } else { - ShouldNotReachHere(); - } - - if (command == nullptr) { - HIP_RETURN(hipErrorOutOfMemory); - } - - command->enqueue(); - command->awaitCompletion(); - command->release(); - - HIP_RETURN(hipSuccess); -} - -hipError_t ihipMemcpy3D_V1(const struct hipMemcpy3DParms* p, hipStream_t stream, bool isAsync = false) { - const void* srcPtr = nullptr; - size_t srcElementSizeInBytes = sizeof(unsigned char); - size_t srcRowPitchInBytes = 0; - size_t srcSlicePitchInBytes = 0; - if (p->srcMemoryType == hipMemoryTypeHost) { - srcPtr = p->srcHost; - srcRowPitchInBytes = p->srcPitch; - srcSlicePitchInBytes = srcRowPitchInBytes * p->srcHeight; - } else if ((p->srcMemoryType == hipMemoryTypeDevice) || - (p->srcMemoryType == hipMemoryTypeUnified)) { - srcPtr = p->srcDevice; - srcRowPitchInBytes = p->srcPitch; - srcSlicePitchInBytes = srcRowPitchInBytes * p->srcHeight; - } else if (p->srcMemoryType == hipMemoryTypeArray) { - srcPtr = p->srcArray->data; - getByteSizeFromChannelFormatKind(p->srcArray->desc.f, &srcElementSizeInBytes); - srcElementSizeInBytes *= p->srcArray->NumChannels; - srcRowPitchInBytes = srcElementSizeInBytes * p->srcArray->width; - srcSlicePitchInBytes = srcRowPitchInBytes * p->srcArray->height; - } else { - ShouldNotReachHere(); - } - - void* dstPtr = nullptr; - size_t dstElementSizeInBytes = sizeof(unsigned char); - size_t dstRowPitchInBytes = 0; - size_t dstSlicePitchInBytes = 0; - if (p->dstMemoryType == hipMemoryTypeHost) { - dstPtr = p->dstHost; - dstRowPitchInBytes = p->dstPitch; - dstSlicePitchInBytes = dstRowPitchInBytes * p->dstHeight; - } else if ((p->dstMemoryType == hipMemoryTypeDevice) || - (p->dstMemoryType == hipMemoryTypeUnified)) { - dstPtr = p->dstDevice; - dstRowPitchInBytes = p->dstPitch; - dstSlicePitchInBytes = dstRowPitchInBytes * p->dstHeight; - } else if (p->dstMemoryType == hipMemoryTypeArray) { - dstPtr = p->dstArray->data; - getByteSizeFromChannelFormatKind(p->dstArray->desc.f, &dstElementSizeInBytes); - dstElementSizeInBytes *= p->dstArray->NumChannels; - dstRowPitchInBytes = dstElementSizeInBytes * p->dstArray->width; - dstSlicePitchInBytes = dstRowPitchInBytes * p->dstArray->height; - } else { - ShouldNotReachHere(); - } - - // For HIP arrays, srcXInBytes must be evenly divisible by the array element size. - if ((p->srcMemoryType == hipMemoryTypeArray) && - ((p->srcXInBytes % srcElementSizeInBytes) != 0)) { - return hipErrorInvalidValue; - } - - // If specified, srcPitch must be greater than or equal to WidthInBytes + srcXInBytes - if ((p->srcMemoryType != hipMemoryTypeArray) && - (p->srcPitch < (p->WidthInBytes + p->srcXInBytes))) { - return hipErrorInvalidValue; - } - - // If specified, srcHeight must be greater than or equal to Height + srcY - if ((p->srcMemoryType != hipMemoryTypeArray) && - (p->srcHeight < (p->Height + p->srcY))) { - return hipErrorInvalidValue; - } - - // For HIP arrays, dstXInBytes must be evenly divisible by the array element size. - if ((p->dstMemoryType == hipMemoryTypeArray) && - ((p->dstXInBytes % dstElementSizeInBytes) != 0)) { - return hipErrorInvalidValue; - } - - // If specified, srcPitch must be greater than or equal to WidthInBytes + srcXInBytes - if ((p->dstMemoryType != hipMemoryTypeArray) && - (p->dstPitch < (p->WidthInBytes + p->dstXInBytes))) { - return hipErrorInvalidValue; - } - - // If specified, srcHeight must be greater than or equal to Height + srcY - if ((p->dstMemoryType != hipMemoryTypeArray) && - (p->dstHeight < (p->Height + p->dstY))) { - return hipErrorInvalidValue; - } - - // The srcLOD and dstLOD members of the CUDA_MEMCPY3D structure must be set to 0. - if ((p->srcLOD != 0) || (p->dstLOD != 0)) { - return hipErrorInvalidValue; - } - - size_t region[3]; - region[0] = p->WidthInBytes; - region[1] = p->Height; - region[2] = p->Depth; - - size_t srcOrigin[3]; - srcOrigin[0] = p->srcXInBytes; - srcOrigin[1] = p->srcY; - srcOrigin[2] = p->srcZ; - - size_t dstOrigin[3]; - dstOrigin[0] = p->dstXInBytes; - dstOrigin[1] = p->dstY; - dstOrigin[2] = p->dstZ; - - amd::BufferRect srcRect; - if (!srcRect.create(srcOrigin, region, srcRowPitchInBytes, srcSlicePitchInBytes)) { - return hipErrorInvalidValue; - } - - size_t srcMemoryOffset = 0; - amd::Memory* srcMemory = getMemoryObject(srcPtr, srcMemoryOffset); - amd::Coord3D srcStart(srcRect.start_ + srcMemoryOffset, 0, 0); - - amd::BufferRect dstRect; - if (!dstRect.create(dstOrigin, region, dstRowPitchInBytes, dstSlicePitchInBytes)) { - return hipErrorInvalidValue; - } - - size_t dstMemoryOffset = 0; - amd::Memory* dstMemory = getMemoryObject(dstPtr, dstMemoryOffset); - amd::Coord3D dstStart(dstRect.start_ + dstMemoryOffset, 0, 0); - - amd::Command* command = nullptr; - amd::Command::EventWaitList waitList; - amd::HostQueue* queue = hip::getQueue(stream); - amd::Coord3D regionSize(region[0], region[1], region[2]); - - if (((srcMemory == nullptr) && (dstMemory == nullptr)) || - (p->kind == hipMemcpyHostToHost)) { - memcpy(dstPtr, srcPtr, region[0] * region[1] * region[2]); - return hipSuccess; - } else if ((srcMemory == nullptr) && (dstMemory != nullptr)) { - command = new amd::WriteMemoryCommand(*queue, CL_COMMAND_WRITE_BUFFER_RECT, waitList, - *dstMemory->asBuffer(), srcStart, regionSize, srcPtr, srcRect, dstRect); - } else if ((srcMemory != nullptr) && (dstMemory == nullptr)) { - command = new amd::ReadMemoryCommand(*queue, CL_COMMAND_READ_BUFFER_RECT, waitList, - *srcMemory->asBuffer(), srcStart, regionSize, dstPtr, srcRect, dstRect); - } else if ((srcMemory != nullptr) && (dstMemory != nullptr)) { - command = new amd::CopyMemoryCommand(*queue, CL_COMMAND_COPY_BUFFER_RECT, waitList, - *srcMemory->asBuffer(),*dstMemory->asBuffer(), srcStart, dstStart, regionSize, - srcRect, dstRect); - } - - if (command == nullptr) { - return hipErrorOutOfMemory; - } - - command->enqueue(); - if (!isAsync) { - command->awaitCompletion(); - } - command->release(); - - return hipSuccess; -} - -hipError_t ihipMemcpy3D_V2(const struct hipMemcpy3DParms* p, hipStream_t stream, bool isAsync = false) { - void* srcPtr = nullptr; - size_t srcElementSizeInBytes = sizeof(unsigned char); - size_t srcRowPitchInBytes = 0; - size_t srcSlicePitchInBytes = 0; - if ((p->srcArray != nullptr) && (p->srcPtr.ptr == nullptr)) { - srcPtr = p->srcArray->data; - getByteSizeFromChannelFormatKind(p->srcArray->desc.f, &srcElementSizeInBytes); - srcElementSizeInBytes *= p->srcArray->NumChannels; - srcRowPitchInBytes = srcElementSizeInBytes * p->srcArray->width; - srcSlicePitchInBytes = srcRowPitchInBytes * p->srcArray->height; - } else if ((p->srcArray == nullptr) && (p->srcPtr.ptr != nullptr)) { - srcPtr = p->srcPtr.ptr; - srcRowPitchInBytes = p->srcPtr.pitch; - srcSlicePitchInBytes = srcRowPitchInBytes * p->srcPtr.ysize; - } else { - ShouldNotReachHere(); - } - - void* dstPtr = nullptr; - size_t dstElementSizeInBytes = sizeof(unsigned char); - size_t dstRowPitchInBytes = 0; - size_t dstSlicePitchInBytes = 0; - if ((p->dstArray != nullptr) && (p->dstPtr.ptr == nullptr)) { - dstPtr = p->dstArray->data; - getByteSizeFromChannelFormatKind(p->dstArray->desc.f, &dstElementSizeInBytes); - dstElementSizeInBytes *= p->dstArray->NumChannels; - dstRowPitchInBytes = dstElementSizeInBytes * p->dstArray->width; - dstSlicePitchInBytes = dstRowPitchInBytes * p->dstArray->height; - } else if ((p->dstArray == nullptr) && (p->dstPtr.ptr != nullptr)) { - dstPtr = p->dstPtr.ptr; - dstRowPitchInBytes = p->srcPtr.pitch; - dstSlicePitchInBytes = dstRowPitchInBytes * p->dstPtr.ysize; - } else { - ShouldNotReachHere(); - } - - // If the source and destination are both arrays, they must have the same element size. - if (((p->srcArray != nullptr) && (p->dstArray != nullptr)) && - (srcElementSizeInBytes != dstElementSizeInBytes)) { - return hipErrorInvalidValue; - } - - // If a HIP array is participating in the copy, the extent is defined in terms of that array's elements. - // If no HIP array is participating in the copy, the extent is defined in elements of unsigned char. - size_t region[3]; - if (p->srcArray != nullptr) { - region[0] = srcRowPitchInBytes; - } else if (p->dstArray != nullptr) { - region[0] = dstRowPitchInBytes; - } else { - region[0] = sizeof(unsigned char) * p->extent.width; - } - region[1] = p->extent.height; - region[2] = p->extent.depth; - - // The offset into the object is defined in units of the object's elements. - size_t srcOrigin[3]; - srcOrigin[0] = srcElementSizeInBytes * p->srcPos.x; - srcOrigin[1] = p->srcPos.y; - srcOrigin[2] = p->srcPos.z; - - amd::BufferRect srcRect; - if (!srcRect.create(srcOrigin, region, srcRowPitchInBytes, srcSlicePitchInBytes)) { - return hipErrorInvalidValue; - } - - size_t srcMemoryOffset = 0; - amd::Memory* srcMemory = getMemoryObject(srcPtr, srcMemoryOffset); - amd::Coord3D srcStart(srcRect.start_ + srcMemoryOffset, 0, 0); - - size_t dstOrigin[3]; - dstOrigin[0] = dstElementSizeInBytes * p->dstPos.x; - dstOrigin[1] = p->dstPos.y; - dstOrigin[2] = p->dstPos.z; - - amd::BufferRect dstRect; - if (!dstRect.create(dstOrigin, region, dstRowPitchInBytes, dstSlicePitchInBytes)) { - return hipErrorInvalidValue; - } - - size_t dstMemoryOffset = 0; - amd::Memory* dstMemory = getMemoryObject(dstPtr, dstMemoryOffset); - amd::Coord3D dstStart(dstRect.start_ + dstMemoryOffset, 0, 0); - - amd::Command* command = nullptr; - amd::Command::EventWaitList waitList; - amd::HostQueue* queue = hip::getQueue(stream); - amd::Coord3D regionSize(region[0], region[1], region[2]); - - if (((srcMemory == nullptr) && (dstMemory == nullptr)) || - (p->kind == hipMemcpyHostToHost)) { - memcpy(dstPtr, srcPtr, region[0] * region[1] * region[2]); - return hipSuccess; - } else if ((srcMemory == nullptr) && (dstMemory != nullptr)) { - command = new amd::WriteMemoryCommand(*queue, CL_COMMAND_WRITE_BUFFER_RECT, waitList, - *dstMemory->asBuffer(), srcStart, regionSize, srcPtr, srcRect, dstRect); - } else if ((srcMemory != nullptr) && (dstMemory == nullptr)) { - command = new amd::ReadMemoryCommand(*queue, CL_COMMAND_READ_BUFFER_RECT, waitList, - *srcMemory->asBuffer(), srcStart, regionSize, dstPtr, srcRect, dstRect); - } else if ((srcMemory != nullptr) && (dstMemory != nullptr)) { - command = new amd::CopyMemoryCommand(*queue, CL_COMMAND_COPY_BUFFER_RECT, waitList, - *srcMemory->asBuffer(),*dstMemory->asBuffer(), srcStart, dstStart, regionSize, - srcRect, dstRect); - } - - if (command == nullptr) { - return hipErrorOutOfMemory; - } - - command->enqueue(); - if (!isAsync) { - command->awaitCompletion(); - } - command->release(); - - return hipSuccess; -} - -hipError_t ihipMemcpy3D(const struct hipMemcpy3DParms* p, hipStream_t stream, bool isAsync = false) { - // Having src and dst be an array is ambigous, since we can't tell if the user intended to call hipMemcpy3D_V1() or hipMemcpy3D_V2(). +hipError_t ihipMemcpy3D(const hipMemcpy3DParms* p, + hipStream_t stream, + bool isAsync = false) { + // Having src and dst be an array is ambigous, since we can't tell + // if the user intended to call runtime or driver version of hipMemcpy3D(). // For now hope that we never encounter this case. assert((p->srcArray == nullptr) || (p->dstArray == nullptr)); - // When calling hipMemcpy3D_V1(), the user must specify - // one of srcHost, srcDevice or srcArray and - // one of dstHost, dstDevice or dstArray. - if (((p->srcHost != nullptr) || (p->srcDevice != nullptr) || (p->srcArray != nullptr)) && - ((p->dstHost != nullptr) || (p->dstDevice != nullptr) || (p->dstArray != nullptr))) { - return ihipMemcpy3D_V1(p, stream, isAsync); + // Now we need to patch the user provided struct if they intended on calling the runtime version. + hipMemcpy3DParms pCopy = {}; + std::memcpy(&pCopy, p, sizeof(hipMemcpy3DParms)); + + if ((p->srcPtr.ptr != nullptr) || (p->dstPtr.ptr != nullptr)) { + pCopy.WidthInBytes = p->extent.width; + pCopy.Height = p->extent.height; + pCopy.Depth = p->extent.depth; } - // When calling hipMemcpy3D_V2(), the user must specify - // one of srcArray or srcPtr and - // one of dstArray or dstPtr. - if (((p->srcArray != nullptr) || (p->srcPtr.ptr != nullptr)) && - ((p->dstArray != nullptr) || (p->dstPtr.ptr != nullptr))) { - return ihipMemcpy3D_V2(p, stream, isAsync); + if (p->srcPtr.ptr != nullptr) { + pCopy.srcXInBytes = p->srcPos.x; + pCopy.srcY = p->srcPos.y; + pCopy.srcZ = p->srcPos.z; + pCopy.srcLOD = 0; + pCopy.srcMemoryType = std::get<0>(hip::getMemoryType(p->kind)); + pCopy.srcHost = p->srcPtr.ptr; + pCopy.srcDevice = p->srcPtr.ptr; + pCopy.srcPitch = p->srcPtr.pitch; + pCopy.srcHeight = p->srcPtr.ysize; + + if (p->dstArray != nullptr) { + pCopy.dstMemoryType = hipMemoryTypeArray; + // When reffering to array memory, hipExtent::width is in elements. + pCopy.WidthInBytes *= hip::getElementSize(p->dstArray->Format); + } } - // If we got here, then the user specified an invalid combination of src/dst parameters. - return hipErrorInvalidValue; + if (p->dstPtr.ptr != nullptr) { + pCopy.dstXInBytes = p->dstPos.x; + pCopy.dstY = p->dstPos.y; + pCopy.dstZ = p->dstPos.z; + pCopy.dstLOD = 0; + pCopy.dstMemoryType = std::get<1>(hip::getMemoryType(p->kind)); + pCopy.dstHost = p->dstPtr.ptr; + pCopy.dstDevice = p->dstPtr.ptr; + pCopy.dstPitch = p->dstPtr.pitch; + pCopy.dstHeight = p->dstPtr.ysize; + + if (p->srcArray != nullptr) { + pCopy.srcMemoryType = hipMemoryTypeArray; + // When reffering to array memory, hipExtent::width is in elements. + pCopy.WidthInBytes *= hip::getElementSize(p->srcArray->Format); + } + } + + return ihipMemcpyParam3D(&pCopy, stream, isAsync); } -hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p) { +hipError_t hipMemcpy3D(const hipMemcpy3DParms* p) { HIP_INIT_API(hipMemcpy3D, p); HIP_RETURN(ihipMemcpy3D(p, nullptr)); } -hipError_t hipMemcpy3DAsync(const struct hipMemcpy3DParms* p, hipStream_t stream) { +hipError_t hipMemcpy3DAsync(const hipMemcpy3DParms* p, hipStream_t stream) { HIP_INIT_API(hipMemcpy3DAsync, p, stream); HIP_RETURN(ihipMemcpy3D(p, stream, true)); @@ -1655,16 +1816,6 @@ hipError_t hipIpcCloseMemHandle(void* dev_ptr) { HIP_RETURN(hipSuccess); } -hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) { - hipChannelFormatDesc cd; - cd.x = x; - cd.y = y; - cd.z = z; - cd.w = w; - cd.f = f; - return cd; -} - hipError_t hipHostGetDevicePointer(void** devicePointer, void* hostPointer, unsigned flags) { HIP_INIT_API(hipHostGetDevicePointer, devicePointer, hostPointer, flags); @@ -1674,7 +1825,7 @@ hipError_t hipHostGetDevicePointer(void** devicePointer, void* hostPointer, unsi if (!memObj) { HIP_RETURN(hipErrorInvalidValue); } - *devicePointer = reinterpret_cast(memObj->getDeviceMemory(*hip::getCurrentDevice()->devices()[0])->virtualAddress() + offset); +*devicePointer = reinterpret_cast(memObj->getDeviceMemory(*hip::getCurrentDevice()->devices()[0])->virtualAddress() + offset); HIP_RETURN(hipSuccess); } @@ -1710,3 +1861,167 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void HIP_RETURN(hipErrorInvalidValue); } + +hipError_t hipArrayDestroy(hipArray* array) { + HIP_INIT_API(hipArrayDestroy, array); + + HIP_RETURN(ihipArrayDestroy(array)); +} + +hipError_t hipArray3DGetDescriptor(HIP_ARRAY3D_DESCRIPTOR* pArrayDescriptor, + hipArray* array) { + HIP_INIT_API(hipArray3DGetDescriptor, pArrayDescriptor, array); + + assert(false && "Unimplemented"); + + HIP_RETURN(hipSuccess); +} + +hipError_t hipArrayGetDescriptor(HIP_ARRAY_DESCRIPTOR* pArrayDescriptor, + hipArray* array) { + HIP_INIT_API(hipArrayGetDescriptor, pArrayDescriptor, array); + + assert(false && "Unimplemented"); + + HIP_RETURN(hipSuccess); +} + +hipError_t hipMemcpyParam2DAsync(const hip_Memcpy2D* pCopy, + hipStream_t stream) { + HIP_INIT_API(hipMemcpyParam2D, pCopy); + + HIP_RETURN(ihipMemcpyParam2D(pCopy, stream, true)); +} + +hipError_t ihipMemcpy2DArrayToArray(hipArray_t dst, size_t wOffsetDst, size_t hOffsetDst, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream, bool isAsync = false) { + hip_Memcpy2D desc = {}; + + desc.srcXInBytes = wOffsetSrc; + desc.srcY = hOffsetSrc; + desc.srcMemoryType = hipMemoryTypeArray; + desc.srcHost = nullptr; + desc.srcDevice = nullptr; + desc.srcArray = const_cast(src); + desc.srcPitch = 0; // Ignored. + + desc.dstXInBytes = wOffsetDst; + desc.dstY = hOffsetDst; + desc.dstMemoryType = hipMemoryTypeArray; + desc.dstHost = nullptr; + desc.dstDevice = nullptr; + desc.dstArray = dst; + desc.dstPitch = 0; // Ignored. + + desc.WidthInBytes = width; + desc.Height = height; + + return ihipMemcpyParam2D(&desc, stream, isAsync); +} + +hipError_t hipMemcpy2DArrayToArray(hipArray_t dst, size_t wOffsetDst, size_t hOffsetDst, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, hipMemcpyKind kind) { + HIP_INIT_API(hipMemcpy2DArrayToArray, dst, wOffsetDst, hOffsetDst, src, wOffsetSrc, hOffsetSrc, width, height, kind); + + HIP_RETURN(ihipMemcpy2DArrayToArray(dst, wOffsetDst, hOffsetDst, src, wOffsetSrc, hOffsetSrc, width, height, kind, nullptr)); +} + +hipError_t hipMemcpyArrayToArray(hipArray_t dst, size_t wOffsetDst, size_t hOffsetDst, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, hipMemcpyKind kind) { + HIP_INIT_API(hipMemcpyArrayToArray, dst, wOffsetDst, hOffsetDst, src, wOffsetSrc, hOffsetSrc, width, height, kind); + + HIP_RETURN(ihipMemcpy2DArrayToArray(dst, wOffsetDst, hOffsetDst, src, wOffsetSrc, hOffsetSrc, width, height, kind, nullptr)); +} + +hipError_t hipMemcpy2DFromArray(void* dst, size_t dpitch, hipArray_const_t src, size_t wOffsetSrc, size_t hOffset, size_t width, size_t height, hipMemcpyKind kind) { + HIP_INIT_API(hipMemcpy2DFromArray, dst, dpitch, src, wOffsetSrc, hOffset, width, height, kind); + + HIP_RETURN(ihipMemcpy2DFromArray(dst, dpitch, src, wOffsetSrc, hOffset, width, height, kind, nullptr)); +} + +hipError_t hipMemcpy2DFromArrayAsync(void* dst, size_t dpitch, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) { + HIP_INIT_API(hipMemcpy2DFromArrayAsync, dst, dpitch, src, wOffsetSrc, hOffsetSrc, width, height, kind, stream); + + HIP_RETURN(ihipMemcpy2DFromArray(dst, dpitch, src, wOffsetSrc, hOffsetSrc, width, height, kind, stream, true)); +} + +hipError_t hipMemcpyFromArrayAsync(void* dst, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t count, hipMemcpyKind kind, hipStream_t stream) { + HIP_INIT_API(hipMemcpyFromArrayAsync, dst, src, wOffsetSrc, hOffsetSrc, count, kind, stream); + + if (src == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + const size_t arrayHeight = (src->height != 0) ? src->height : 1; + const size_t widthInBytes = count / arrayHeight; + + const size_t height = (count / src->width) / (hip::getElementSize(src->Format) * src->NumChannels); + + HIP_RETURN(ihipMemcpy2DFromArray(dst, 0 /* dpitch */, src, wOffsetSrc, hOffsetSrc, widthInBytes, height, kind, stream, true)); +} + +hipError_t hipMemcpy2DToArrayAsync(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) { + HIP_INIT_API(hipMemcpy2DToArrayAsync, dst, wOffset, hOffset, src, spitch, width, height, kind); + + HIP_RETURN(ihipMemcpy2DToArray(dst, wOffset, hOffset, src, spitch, width, height, kind, stream, true)); +} + +hipError_t hipMemcpyToArrayAsync(hipArray_t dst, size_t wOffset, size_t hOffset, const void* src, size_t count, hipMemcpyKind kind, hipStream_t stream) { + HIP_INIT_API(hipMemcpyToArrayAsync, dst, wOffset, hOffset, src, count, kind); + + if (dst == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + const size_t arrayHeight = (dst->height != 0) ? dst->height : 1; + const size_t widthInBytes = count / arrayHeight; + + const size_t height = (count / dst->width) / (hip::getElementSize(dst->Format) * dst->NumChannels); + + HIP_RETURN(ihipMemcpy2DToArray(dst, wOffset, hOffset, src, 0 /* spitch */, widthInBytes, height, kind, stream, true)); +} + +hipError_t hipMemcpyAtoA(hipArray* dstArray, + size_t dstOffset, + hipArray* srcArray, + size_t srcOffset, + size_t ByteCount) { + HIP_INIT_API(hipMemcpyAtoA, dstArray, dstOffset, srcArray, srcOffset, ByteCount); + + HIP_RETURN(ihipMemcpyAtoA(srcArray, dstArray, {srcOffset, 0, 0}, {dstOffset, 0, 0}, {ByteCount, 1, 1}, nullptr)); +} + +hipError_t hipMemcpyAtoD(hipDeviceptr_t dstDevice, + hipArray* srcArray, + size_t srcOffset, + size_t ByteCount) { + HIP_INIT_API(hipMemcpyAtoD, dstDevice, srcArray, srcOffset, ByteCount); + + HIP_RETURN(ihipMemcpyAtoD(srcArray, dstDevice, {srcOffset, 0, 0}, {0, 0, 0}, {ByteCount, 1, 1}, 0, 0, nullptr)); +} + +hipError_t hipMemcpyAtoHAsync(void* dstHost, + hipArray* srcArray, + size_t srcOffset, + size_t ByteCount, + hipStream_t stream) { + HIP_INIT_API(hipMemcpyAtoHAsync, dstHost, srcArray, srcOffset, ByteCount, stream); + + HIP_RETURN(ihipMemcpyAtoH(srcArray, dstHost, {srcOffset, 0, 0}, {0, 0, 0}, {ByteCount, 1, 1}, 0, 0, stream, true)); +} + +hipError_t hipMemcpyDtoA(hipArray* dstArray, + size_t dstOffset, + hipDeviceptr_t srcDevice, + size_t ByteCount) { + HIP_INIT_API(hipMemcpyDtoA, dstArray, dstOffset, srcDevice, ByteCount); + + HIP_RETURN(ihipMemcpyDtoA(srcDevice, dstArray, {0, 0, 0}, {dstOffset, 0, 0}, {ByteCount, 1, 1}, 0, 0, nullptr)); +} + +hipError_t hipMemcpyHtoAAsync(hipArray* dstArray, + size_t dstOffset, + const void* srcHost, + size_t ByteCount, + hipStream_t stream) { + HIP_INIT_API(hipMemcpyHtoAAsync, dstArray, dstOffset, srcHost, ByteCount, stream); + + HIP_RETURN(ihipMemcpyHtoA(srcHost, dstArray, {0, 0, 0}, {dstOffset, 0, 0}, {ByteCount, 1, 1}, 0, 0, stream, true)); +} \ No newline at end of file diff --git a/projects/clr/hipamd/vdi/hip_texture.cpp b/projects/clr/hipamd/vdi/hip_texture.cpp index 98d4a773f4..634c83557c 100644 --- a/projects/clr/hipamd/vdi/hip_texture.cpp +++ b/projects/clr/hipamd/vdi/hip_texture.cpp @@ -21,6 +21,7 @@ #include #include #include "hip_internal.hpp" +#include "hip_conversions.hpp" #include "platform/sampler.hpp" namespace hip { @@ -30,777 +31,1153 @@ namespace hip { amd::Image* image; amd::Sampler* sampler; hipResourceDesc resDesc; + hipTextureDesc texDesc; + hipResourceViewDesc resViewDesc; + + TextureObject(amd::Image* image_, + amd::Sampler* sampler_, + const hipResourceDesc& resDesc_, + const hipTextureDesc& texDesc_, + const hipResourceViewDesc& resViewDesc_) : + image(image_), + sampler(sampler_), + resDesc(resDesc_), + texDesc(texDesc_), + resViewDesc(resViewDesc_) { + amd::Context& context = *hip::getCurrentDevice()->asContext(); + amd::Device& device = *context.devices()[0]; + + device::Memory* imageMem = image->getDeviceMemory(device); + std::memcpy(imageSRD, imageMem->cpuSrd(), sizeof(imageSRD)); + + device::Sampler* samplerMem = sampler->getDeviceSampler(device); + std::memcpy(samplerSRD, samplerMem->hwState(), sizeof(samplerSRD)); + } }; }; -void getDrvChannelOrderAndType(const enum hipArray_Format Format, unsigned int NumChannels, - cl_channel_order* channelOrder, - cl_channel_type* channelType) { - switch (Format) { - case HIP_AD_FORMAT_UNSIGNED_INT8: - *channelType = CL_UNSIGNED_INT8; - break; - case HIP_AD_FORMAT_UNSIGNED_INT16: - *channelType = CL_UNSIGNED_INT16; - break; - case HIP_AD_FORMAT_UNSIGNED_INT32: - *channelType = CL_UNSIGNED_INT32; - break; - case HIP_AD_FORMAT_SIGNED_INT8: - *channelType = CL_SIGNED_INT8; - break; - case HIP_AD_FORMAT_SIGNED_INT16: - *channelType = CL_SIGNED_INT16; - break; - case HIP_AD_FORMAT_SIGNED_INT32: - *channelType = CL_SIGNED_INT32; - break; - case HIP_AD_FORMAT_HALF: - *channelType = CL_HALF_FLOAT; - break; - case HIP_AD_FORMAT_FLOAT: - *channelType = CL_FLOAT; - break; - default: - break; +amd::Image* ihipImageCreate(const cl_channel_order channelOrder, + const cl_channel_type channelType, + const cl_mem_object_type imageType, + const size_t imageWidth, + const size_t imageHeight, + const size_t imageDepth, + const size_t imageArraySize, + const size_t imageRowPitch, + const size_t imageSlicePitch, + const uint32_t numMipLevels, + amd::Memory* buffer); + +hipError_t ihipCreateTextureObject(hipTextureObject_t* pTexObject, + const hipResourceDesc* pResDesc, + const hipTextureDesc* pTexDesc, + const hipResourceViewDesc* pResViewDesc) { + amd::Device* device = hip::getCurrentDevice()->devices()[0]; + const device::Info& info = device->info(); + + // pResViewDesc can only be specified if the type of resource is a HIP array or a HIP mipmapped array. + if ((pResViewDesc != nullptr) && + ((pResDesc->resType != hipResourceTypeArray) && (pResDesc->resType != hipResourceTypeMipmappedArray))) { + return hipErrorInvalidValue; } - if (NumChannels == 4) { - *channelOrder = CL_RGBA; - } else if (NumChannels == 2) { - *channelOrder = CL_RG; - } else if (NumChannels == 1) { - *channelOrder = CL_R; - } -} - -void setDescFromChannelType(cl_channel_type channelType, hipChannelFormatDesc* desc) { - - memset(desc, 0x00, sizeof(hipChannelFormatDesc)); - - switch (channelType) { - case CL_SIGNED_INT8: - case CL_SIGNED_INT16: - case CL_SIGNED_INT32: - desc->f = hipChannelFormatKindSigned; - break; - case CL_UNSIGNED_INT8: - case CL_UNSIGNED_INT16: - case CL_UNSIGNED_INT32: - desc->f = hipChannelFormatKindUnsigned; - break; - case CL_HALF_FLOAT: - case CL_FLOAT: - desc->f = hipChannelFormatKindFloat; - break; - default: - desc->f = hipChannelFormatKindNone; - break; + // If hipResourceDesc::resType is set to hipResourceTypeArray, + // hipResourceDesc::res::array::array must be set to a valid HIP array handle. + if ((pResDesc->resType == hipResourceTypeArray) && + (pResDesc->res.array.array == nullptr)) { + return hipErrorInvalidValue; } - switch (channelType) { - case CL_SIGNED_INT8: - case CL_UNSIGNED_INT8: - desc->x = 8; - break; - case CL_SIGNED_INT16: - case CL_UNSIGNED_INT16: - case CL_HALF_FLOAT: - desc->x = 16; - break; - case CL_SIGNED_INT32: - case CL_UNSIGNED_INT32: - case CL_FLOAT: - desc->x = 32; - break; - default: - desc->x = 0; - break; + // If hipResourceDesc::resType is set to hipResourceTypeMipmappedArray, + // hipResourceDesc::res::mipmap::mipmap must be set to a valid HIP mipmapped array handle + // and hipTextureDesc::normalizedCoords must be set to true. + if ((pResDesc->resType == hipResourceTypeMipmappedArray) && + ((pResDesc->res.mipmap.mipmap == nullptr) || (pTexDesc->normalizedCoords == 0))) { + return hipErrorInvalidValue; } -} -void getChannelOrderAndType(const hipChannelFormatDesc& desc, enum hipTextureReadMode readMode, - cl_channel_order* channelOrder, cl_channel_type* channelType) { - if (desc.x != 0 && desc.y != 0 && desc.z != 0 && desc.w != 0) { - *channelOrder = CL_RGBA; - } else if (desc.x != 0 && desc.y != 0 && desc.z != 0 && desc.w == 0) { - *channelOrder = CL_RGB; - } else if (desc.x != 0 && desc.y != 0 && desc.z == 0 && desc.w == 0) { - *channelOrder = CL_RG; - } else if (desc.x != 0 && desc.y == 0 && desc.z == 0 && desc.w == 0) { - *channelOrder = CL_R; - } else { - } + // If hipResourceDesc::resType is set to hipResourceTypeLinear, + // hipResourceDesc::res::linear::devPtr must be set to a valid device pointer, that is aligned to hipDeviceProp::textureAlignment. + // The total number of elements in the linear address range cannot exceed hipDeviceProp::maxTexture1DLinear. + if ((pResDesc->resType == hipResourceTypeLinear) && + ((pResDesc->res.linear.devPtr == nullptr) || + (!amd::isMultipleOf(pResDesc->res.linear.devPtr, info.imageBaseAddressAlignment_)) || + (pResDesc->res.linear.sizeInBytes >= info.imageMaxBufferSize_))) { + return hipErrorInvalidValue; + } - switch (desc.f) { - case hipChannelFormatKindUnsigned: - switch (desc.x) { - case 32: - *channelType = CL_UNSIGNED_INT32; - break; - case 16: - *channelType = readMode == hipReadModeNormalizedFloat - ? CL_UNORM_INT16 - : CL_UNSIGNED_INT16; - break; - case 8: - *channelType = readMode == hipReadModeNormalizedFloat - ? CL_UNORM_INT8 - : CL_UNSIGNED_INT8; - break; - default: - *channelType = CL_UNSIGNED_INT32; - } - break; - case hipChannelFormatKindSigned: - switch (desc.x) { - case 32: - *channelType = CL_SIGNED_INT32; - break; - case 16: - *channelType = readMode == hipReadModeNormalizedFloat - ? CL_SNORM_INT16 - : CL_SIGNED_INT16; - break; - case 8: - *channelType = readMode == hipReadModeNormalizedFloat - ? CL_SNORM_INT8 - : CL_SIGNED_INT8; - break; - default: - *channelType = CL_SIGNED_INT32; - } - break; - case hipChannelFormatKindFloat: - switch (desc.x) { - case 32: - *channelType = CL_FLOAT; - break; - case 16: - *channelType = CL_HALF_FLOAT; - break; - case 8: - break; - default: - *channelType = CL_FLOAT; - } - break; - case hipChannelFormatKindNone: - default: - break; - } -} + // If hipResourceDesc::resType is set to hipResourceTypePitch2D, + // hipResourceDesc::res::pitch2D::devPtr must be set to a valid device pointer, that is aligned to hipDeviceProp::textureAlignment. + // hipResourceDesc::res::pitch2D::width and hipResourceDesc::res::pitch2D::height specify the width and height of the array in elements, + // and cannot exceed hipDeviceProp::maxTexture2DLinear[0] and hipDeviceProp::maxTexture2DLinear[1] respectively. + // hipResourceDesc::res::pitch2D::pitchInBytes specifies the pitch between two rows in bytes and has to be aligned to hipDeviceProp::texturePitchAlignment. + // Pitch cannot exceed hipDeviceProp::maxTexture2DLinear[2]. + if ((pResDesc->resType == hipResourceTypePitch2D) && + ((pResDesc->res.pitch2D.devPtr == nullptr) || + (!amd::isMultipleOf(pResDesc->res.pitch2D.devPtr, info.imageBaseAddressAlignment_)) || + (pResDesc->res.pitch2D.width >= info.image2DMaxWidth_) || + (pResDesc->res.pitch2D.height >= info.image2DMaxHeight_) || + (!amd::isMultipleOf(pResDesc->res.pitch2D.pitchInBytes, info.imagePitchAlignment_)))) { + // TODO check pitch limits. + return hipErrorInvalidValue; + } -void getByteSizeFromChannelFormatKind(enum hipChannelFormatKind channelFormatKind, size_t* byteSize) { - switch (channelFormatKind) - { - case hipChannelFormatKindSigned: - *byteSize = sizeof(int); - break; - case hipChannelFormatKindUnsigned: - *byteSize = sizeof(unsigned int); - break; - case hipChannelFormatKindFloat: - *byteSize = sizeof(float); - break; - case hipChannelFormatKindNone: - *byteSize = sizeof(size_t); - break; - default: - *byteSize = 1; - break; - } -} + // Mipmaps are currently not supported. + if (pResDesc->resType == hipResourceTypeMipmappedArray) { + return hipErrorNotSupported; + } + // We don't program the border_color_ptr field in the HW sampler SRD. + if (pTexDesc->addressMode[0] == hipAddressModeBorder) { + return hipErrorNotSupported; + } + // We don't program the force_degamma/skip_degamma fields in the HW sampler SRD. + if (pTexDesc->sRGB == 1) { + return hipErrorNotSupported; + } + // We don't program the max_ansio_ratio field in the the HW sampler SRD. + if (pTexDesc->maxAnisotropy != 0) { + return hipErrorNotSupported; + } + // We don't program the lod_bias field in the HW sampler SRD. + if (pTexDesc->mipmapLevelBias != 0) { + return hipErrorNotSupported; + } + // We don't program the min_lod field in the HW sampler SRD. + if (pTexDesc->minMipmapLevelClamp != 0) { + return hipErrorNotSupported; + } + // We don't program the max_lod field in the HW sampler SRD. + if (pTexDesc->maxMipmapLevelClamp != 0) { + return hipErrorNotSupported; + } + + // TODO VDI assumes all dimensions have the same addressing mode. + cl_addressing_mode addressMode = CL_ADDRESS_NONE; + // If hipTextureDesc::normalizedCoords is set to zero, + // hipAddressModeWrap and hipAddressModeMirror won't be supported + // and will be switched to hipAddressModeClamp. + if ((pTexDesc->normalizedCoords == 0) && + ((pTexDesc->addressMode[0] == hipAddressModeWrap) || (pTexDesc->addressMode[0] == hipAddressModeMirror))) { + addressMode = hip::getCLAddressingMode(hipAddressModeClamp); + } + // hipTextureDesc::addressMode is ignored if hipResourceDesc::resType is hipResourceTypeLinear + else if (pResDesc->resType != hipResourceTypeLinear) { + addressMode = hip::getCLAddressingMode(pTexDesc->addressMode[0]); + } -amd::Sampler* fillSamplerDescriptor(enum hipTextureAddressMode addressMode, - enum hipTextureFilterMode filterMode, int normalizedCoords) { #ifndef CL_FILTER_NONE #define CL_FILTER_NONE 0x1142 #endif - uint32_t filter_mode = CL_FILTER_NONE; - switch (filterMode) { - case hipFilterModePoint: - filter_mode = CL_FILTER_NEAREST; - break; - case hipFilterModeLinear: - filter_mode = CL_FILTER_LINEAR; - break; + cl_filter_mode filterMode = CL_FILTER_NONE; +#undef CL_FILTER_NONE + // hipTextureDesc::filterMode is ignored if hipResourceDesc::resType is hipResourceTypeLinear. + if (pResDesc->resType != hipResourceTypeLinear) { + filterMode = hip::getCLFilterMode(pTexDesc->filterMode); } - uint32_t address_mode = CL_ADDRESS_NONE; - switch (addressMode) { - case hipAddressModeWrap: - address_mode = CL_ADDRESS_REPEAT; - break; - case hipAddressModeClamp: - address_mode = CL_ADDRESS_CLAMP; - break; - case hipAddressModeMirror: - address_mode = CL_ADDRESS_MIRRORED_REPEAT; - break; - case hipAddressModeBorder: - address_mode = CL_ADDRESS_CLAMP_TO_EDGE; - break; +#ifndef CL_FILTER_NONE +#define CL_FILTER_NONE 0x1142 +#endif + cl_filter_mode mipFilterMode = CL_FILTER_NONE; +#undef CL_FILTER_NONE + if (pResDesc->resType == hipResourceTypeMipmappedArray) { + mipFilterMode = hip::getCLFilterMode(pTexDesc->mipmapFilterMode); } - amd::Sampler* sampler = new amd::Sampler(*hip::getCurrentDevice()->asContext(), - normalizedCoords == CL_TRUE, - address_mode, filter_mode, CL_FILTER_NONE, 0.f, CL_MAXFLOAT); + + amd::Sampler* sampler = new amd::Sampler(*hip::getCurrentDevice()->asContext(), + pTexDesc->normalizedCoords, + addressMode, + filterMode, + mipFilterMode, + pTexDesc->minMipmapLevelClamp, + pTexDesc->maxMipmapLevelClamp); + if (sampler == nullptr) { - return nullptr; + return hipErrorOutOfMemory; } + if (!sampler->create()) { delete sampler; - return nullptr; - } - return sampler; -} - -hip::TextureObject* ihipCreateTextureObject(const hipResourceDesc& resDesc, amd::Image& image, amd::Sampler& sampler) { - hip::TextureObject* texture; - ihipMalloc(reinterpret_cast(&texture), sizeof(hip::TextureObject), CL_MEM_SVM_FINE_GRAIN_BUFFER); - - if (texture == nullptr) { - return nullptr; - } - - device::Memory* imageMem = image.getDeviceMemory(*hip::getCurrentDevice()->devices()[0]); - memcpy(texture->imageSRD, imageMem->cpuSrd(), sizeof(uint32_t)*HIP_IMAGE_OBJECT_SIZE_DWORD); - texture->image = ℑ - - device::Sampler* devSampler = sampler.getDeviceSampler(*hip::getCurrentDevice()->devices()[0]); - memcpy(texture->samplerSRD, devSampler->hwState(), sizeof(uint32_t)*HIP_SAMPLER_OBJECT_SIZE_DWORD); - texture->sampler = &sampler; - - memcpy(&texture->resDesc, &resDesc, sizeof(hipResourceDesc)); - - return texture; -} - -hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResourceDesc* pResDesc, - const hipTextureDesc* pTexDesc, - const hipResourceViewDesc* pResViewDesc) { - HIP_INIT_API(NONE, pTexObject, pResDesc, pTexDesc, pResViewDesc); - - amd::Device* device = hip::getCurrentDevice()->devices()[0]; - - if (!device->info().imageSupport_) { - HIP_RETURN(hipErrorInvalidValue); + return hipErrorOutOfMemory; } amd::Image* image = nullptr; - - cl_image_format image_format; - getChannelOrderAndType(pResDesc->res.pitch2D.desc, pTexDesc->readMode, - &image_format.image_channel_order, &image_format.image_channel_data_type); - - const amd::Image::Format imageFormat(image_format); - - amd::Memory* memory = nullptr; - size_t offset = 0; - cl_mem_object_type clType; - switch (pResDesc->resType) { - case hipResourceTypeArray: - { - memory = getMemoryObject(pResDesc->res.array.array->data, offset); + case hipResourceTypeArray: { + cl_mem memObj = reinterpret_cast(pResDesc->res.array.array->data); + if (!is_valid(memObj)) { + return hipErrorInvalidValue; + } + image = as_amd(memObj)->asImage(); - getChannelOrderAndType(pResDesc->res.array.array->desc, pTexDesc->readMode, - &image_format.image_channel_order, &image_format.image_channel_data_type); - const amd::Image::Format imageFormat(image_format); - switch (pResDesc->res.array.array->type) { - case hipArrayLayered: - case hipArrayCubemap: - assert(0); - break; - case hipArraySurfaceLoadStore: - case hipArrayTextureGather: - case hipArrayDefault: - default: - switch(pResDesc->res.array.array->textureType) { - case hipTextureType3D: - clType = CL_MEM_OBJECT_IMAGE3D; - image = new (*hip::getCurrentDevice()->asContext()) amd::Image(*memory->asBuffer(), - clType, memory->getMemFlags(), imageFormat, - pResDesc->res.array.array->width, pResDesc->res.array.array->height, - pResDesc->res.array.array->depth, 0, 0); - break; - case hipTextureType2D: - clType = CL_MEM_OBJECT_IMAGE2D; - image = new (*hip::getCurrentDevice()->asContext()) amd::Image(*memory->asBuffer(), - clType, memory->getMemFlags(), imageFormat, - pResDesc->res.array.array->width, pResDesc->res.array.array->height, 1, 0, 0); - break; - default: - break; - } - break; - } + hipTextureReadMode readMode = pTexDesc->readMode; + // 32-bit integer format will not be promoted, regardless of whether or not + // this hipTextureDesc::readMode is set hipReadModeNormalizedFloat is specified. + if ((hip::getElementSize(pResDesc->res.array.array->Format) == 4) && + (pResDesc->res.array.array->Format != HIP_AD_FORMAT_FLOAT)) { + readMode = hipReadModeElementType; + } + + // We need to create an image view if the user requested to use normalized pixel values, + // due to already having the image created with a different format. + if ((pResViewDesc != nullptr) || + (readMode == hipReadModeNormalizedFloat)) { + // TODO VDI currently right now can only change the format of the image. + const cl_channel_order channelOrder = (pResViewDesc != nullptr) ? hip::getCLChannelOrder(hip::getNumChannels(pResViewDesc->format)) : + hip::getCLChannelOrder(pResDesc->res.array.array->NumChannels); + const cl_channel_type channelType = (pResViewDesc != nullptr) ? hip::getCLChannelType(hip::getArrayFormat(pResViewDesc->format), readMode) : + hip::getCLChannelType(pResDesc->res.array.array->Format, readMode); + const amd::Image::Format imageFormat(cl_image_format{channelOrder, channelType}); + if (!imageFormat.isValid()) { + return hipErrorInvalidValue; } - break; - case hipResourceTypeMipmappedArray: - assert(0); - break; - case hipResourceTypeLinear: - { - assert(pResViewDesc == nullptr); - memory = getMemoryObject(pResDesc->res.linear.devPtr, offset); - getChannelOrderAndType(pResDesc->res.linear.desc, pTexDesc->readMode, - &image_format.image_channel_order, &image_format.image_channel_data_type); - const amd::Image::Format imageFormat(image_format); - - image = new (*hip::getCurrentDevice()->asContext()) amd::Image(*memory->asBuffer(), - CL_MEM_OBJECT_IMAGE2D, memory->getMemFlags(), imageFormat, - pResDesc->res.linear.sizeInBytes / imageFormat.getElementSize(), 1, 1, - pResDesc->res.linear.sizeInBytes, 0); + image = image->createView(*hip::getCurrentDevice()->asContext(), imageFormat, nullptr); + if (image == nullptr) { + return hipErrorInvalidValue; } - break; - case hipResourceTypePitch2D: - assert(pResViewDesc == nullptr); - memory = getMemoryObject(pResDesc->res.pitch2D.devPtr, offset); - - image = new (*hip::getCurrentDevice()->asContext()) amd::Image(*memory->asBuffer(), - CL_MEM_OBJECT_IMAGE2D, memory->getMemFlags(), imageFormat, - pResDesc->res.pitch2D.width, pResDesc->res.pitch2D.height, 1, - pResDesc->res.pitch2D.pitchInBytes, 0); - break; - default: HIP_RETURN(hipErrorInvalidValue); - } - - if (!image->create()) { - delete image; - HIP_RETURN(hipErrorOutOfMemory); - } - - amd::Sampler* sampler = fillSamplerDescriptor(pTexDesc->addressMode[0], pTexDesc->filterMode, pTexDesc->normalizedCoords); - - *pTexObject = reinterpret_cast(ihipCreateTextureObject(*pResDesc, *image, *sampler)); - - HIP_RETURN(hipSuccess); -} - -void ihipDestroyTextureObject(hip::TextureObject* texture) { - texture->image->release(); - texture->sampler->release(); - - hipFree(texture); -} - -hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject) { - HIP_INIT_API(NONE, textureObject); - - hip::TextureObject* texture = reinterpret_cast(textureObject); - - ihipDestroyTextureObject(texture); - - HIP_RETURN(hipSuccess); -} - -hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc, - hipTextureObject_t textureObject) { - HIP_INIT_API(NONE, pResDesc, textureObject); - - hip::TextureObject* texture = reinterpret_cast(textureObject); - - if (pResDesc != nullptr && texture != nullptr) { - memcpy(pResDesc, &(texture->resDesc), sizeof(hipResourceDesc)); - } - - HIP_RETURN(hipErrorInvalidValue); -} - -hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc, - hipTextureObject_t textureObject) { - HIP_INIT_API(NONE, pResViewDesc, textureObject); - - assert(0 && "Unimplemented"); - - HIP_RETURN(hipErrorNotSupported); -} - -hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc, - hipTextureObject_t textureObject) { - HIP_INIT_API(NONE, pTexDesc, textureObject); - - assert(0 && "Unimplemented"); - - HIP_RETURN(hipErrorNotSupported); -} - -hipError_t ihipBindTexture(cl_mem_object_type type, - size_t* offset, textureReference* tex, const void* devPtr, - const hipChannelFormatDesc& desc, size_t width, size_t height, - size_t pitch) { - if (tex == nullptr) { - return hipErrorInvalidImage; - } - if (hip::getCurrentDevice()) { - cl_image_format image_format; - size_t byteSize; - size_t rowPitch = 0; - size_t depth = 0; - size_t slicePitch = 0; - - getChannelOrderAndType(desc, hipReadModeElementType, - &image_format.image_channel_order, &image_format.image_channel_data_type); - getByteSizeFromChannelFormatKind(desc.f, &byteSize); - const amd::Image::Format imageFormat(image_format); - amd::Memory* memory = getMemoryObject(devPtr, *offset); - - switch (type) { - case CL_MEM_OBJECT_IMAGE3D: - rowPitch = width * byteSize; - depth = pitch; - slicePitch = rowPitch * height; - break; - case CL_MEM_OBJECT_IMAGE2D: - default: - rowPitch = pitch; - depth = 1; - slicePitch = 0; - break; } - - amd::Image* image = new (*hip::getCurrentDevice()->asContext()) amd::Image(*memory->asBuffer(), - type, memory->getMemFlags(), imageFormat, width, height, depth, rowPitch, slicePitch); - if (!image->create()) { - delete image; - return hipErrorOutOfMemory; - } - - *offset = 0; - if (tex->textureObject) { - ihipDestroyTextureObject(reinterpret_cast(tex->textureObject)); - } - amd::Sampler* sampler = fillSamplerDescriptor(tex->addressMode[0], tex->filterMode, tex->normalized); - - hipResourceDesc resDesc; - memset(&resDesc, 0, sizeof(hipResourceDesc)); - switch (type) { - case CL_MEM_OBJECT_IMAGE1D: - resDesc.resType = hipResourceTypeLinear; - resDesc.res.linear.devPtr = const_cast(devPtr); - resDesc.res.linear.desc = desc; - resDesc.res.linear.sizeInBytes = image->getSize(); - break; - case CL_MEM_OBJECT_IMAGE2D: - resDesc.resType = hipResourceTypePitch2D; - resDesc.res.pitch2D.devPtr = const_cast(devPtr); - resDesc.res.pitch2D.desc = desc; - resDesc.res.pitch2D.width = width; - resDesc.res.pitch2D.height = height; - resDesc.res.pitch2D.pitchInBytes = pitch; - break; - case CL_MEM_OBJECT_IMAGE3D: - resDesc.resType = hipResourceTypeArray; - resDesc.res.array.array = (hipArray*)malloc(sizeof(hipArray)); - resDesc.res.array.array->desc = desc; - resDesc.res.array.array->width = width; - resDesc.res.array.array->height = height; - resDesc.res.array.array->depth = depth; - resDesc.res.array.array->Format = tex->format; - resDesc.res.array.array->NumChannels = tex->numChannels; - resDesc.res.array.array->isDrv = false; - resDesc.res.array.array->textureType = hipTextureType3D; - resDesc.res.array.array->data = const_cast(devPtr); - break; - default: - resDesc.resType = hipResourceTypeArray; - resDesc.res.array.array = nullptr; - break; - } - - tex->textureObject = reinterpret_cast(ihipCreateTextureObject(resDesc, *image, *sampler)); - if(type == CL_MEM_OBJECT_IMAGE3D) { - free(resDesc.res.array.array); - } - memset(&resDesc, 0, sizeof(hipResourceDesc)); - return hipSuccess; + break; } - return hipErrorInvalidValue; -} - -hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* devPtr, - const hipChannelFormatDesc* desc, size_t size) { - HIP_INIT_API(NONE, offset, tex, devPtr, desc, size); - - if (desc == nullptr) { - HIP_RETURN(hipErrorInvalidValue); + case hipResourceTypeMipmappedArray: { + ShouldNotReachHere(); + break; + } + case hipResourceTypeLinear: { + const cl_channel_order channelOrder = hip::getCLChannelOrder(hip::getNumChannels(pResDesc->res.linear.desc)); + const cl_channel_type channelType = hip::getCLChannelType(hip::getArrayFormat(pResDesc->res.linear.desc), pTexDesc->readMode); + const amd::Image::Format imageFormat({channelOrder, channelType}); + const cl_mem_object_type imageType = hip::getCLMemObjectType(pResDesc->resType); + size_t offset = 0; + image = ihipImageCreate(channelOrder, + channelType, + imageType, + (pResDesc->res.linear.sizeInBytes / imageFormat.getElementSize()), /* imageWidth */ + 0, /* imageHeight */ + 0, /* imageDepth */ + 0, /* imageArraySize */ + 0, /* imageRowPitch */ + 0, /* imageSlicePitch */ + 0, /* numMipLevels */ + getMemoryObject(pResDesc->res.linear.devPtr, offset)); + // TODO take care of non-zero offset. + assert(offset == 0); + if (image == nullptr) { + return hipErrorInvalidValue; + } + break; + } + case hipResourceTypePitch2D: { + const cl_channel_order channelOrder = hip::getCLChannelOrder(hip::getNumChannels(pResDesc->res.pitch2D.desc)); + const cl_channel_type channelType = hip::getCLChannelType(hip::getArrayFormat(pResDesc->res.pitch2D.desc), pTexDesc->readMode); + const cl_mem_object_type imageType = hip::getCLMemObjectType(pResDesc->resType); + size_t offset = 0; + image = ihipImageCreate(channelOrder, + channelType, + imageType, + pResDesc->res.pitch2D.width, /* imageWidth */ + pResDesc->res.pitch2D.height, /* imageHeight */ + 0, /* imageDepth */ + 0, /* imageArraySize */ + pResDesc->res.pitch2D.pitchInBytes, /* imageRowPitch */ + 0, /* imageSlicePitch */ + 0, /* numMipLevels */ + getMemoryObject(pResDesc->res.pitch2D.devPtr, offset)); + // TODO take care of non-zero offset. + assert(offset == 0); + if (image == nullptr) { + return hipErrorInvalidValue; + } + break; } - cl_image_format image_format; - getChannelOrderAndType(*desc, hipReadModeElementType, - &image_format.image_channel_order, &image_format.image_channel_data_type); - const amd::Image::Format imageFormat(image_format); - - HIP_RETURN(ihipBindTexture(CL_MEM_OBJECT_IMAGE1D, offset, tex, devPtr, *desc, size / imageFormat.getElementSize(), 1, size)); -} - -hipError_t hipBindTexture2D(size_t* offset, textureReference* tex, const void* devPtr, - const hipChannelFormatDesc* desc, size_t width, size_t height, - size_t pitch) { - HIP_INIT_API(NONE, offset, tex, devPtr, desc, width, height, pitch); - - HIP_RETURN(ihipBindTexture(CL_MEM_OBJECT_IMAGE2D, offset, tex, devPtr, *desc, width, height, pitch)); -} - -hipError_t hipBindTextureToArray(textureReference* tex, hipArray_const_t array, - const hipChannelFormatDesc* desc) { - HIP_INIT_API(NONE, tex, array, desc); - - assert(0 && "Unimplemented"); - - HIP_RETURN(hipErrorNotSupported); -} - -hipError_t ihipBindTextureImpl(TlsData* tls, int dim, enum hipTextureReadMode readMode, size_t* offset, - const void* devPtr, const struct hipChannelFormatDesc* desc, - size_t size, textureReference* tex) { - HIP_INIT_API(NONE, dim, readMode, offset, devPtr, size, tex); - - assert(1 == dim); - - HIP_RETURN(ihipBindTexture(CL_MEM_OBJECT_IMAGE1D, offset, tex, devPtr, *desc, size, 1, 0)); -} - -hipError_t ihipBindTextureToArrayImpl(TlsData* tls, int dim, enum hipTextureReadMode readMode, - hipArray_const_t array, - const struct hipChannelFormatDesc& desc, - textureReference* tex) { - HIP_INIT_API(NONE, dim, readMode, &desc, array, tex); - - cl_mem_object_type clType; - size_t offset = 0; - - switch (dim) { - case 1: - clType = CL_MEM_OBJECT_IMAGE1D; - break; - case 2: - clType = CL_MEM_OBJECT_IMAGE2D; - break; - case 3: - case hipTextureType2DLayered: - clType = CL_MEM_OBJECT_IMAGE3D; - break; - default: - HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(ihipBindTexture(clType, &offset, tex, array->data, desc, array->width, - array->height, array->depth)); -} - -hipError_t hipBindTextureToMipmappedArray(textureReference* tex, - hipMipmappedArray_const_t mipmappedArray, - const hipChannelFormatDesc* desc) { - HIP_INIT_API(NONE, tex, mipmappedArray, desc); - - assert(0 && "Unimplemented"); - - HIP_RETURN(hipErrorNotSupported); -} - -hipError_t ihipUnbindTextureImpl(const hipTextureObject_t& textureObject) { - - ihipDestroyTextureObject(reinterpret_cast(textureObject)); + void *texObjectBuffer = nullptr; + ihipMalloc(&texObjectBuffer, sizeof(hip::TextureObject), CL_MEM_SVM_FINE_GRAIN_BUFFER); + if (texObjectBuffer == nullptr) { + return hipErrorOutOfMemory; + } + *pTexObject = reinterpret_cast(new (texObjectBuffer) hip::TextureObject{image, sampler, *pResDesc, *pTexDesc, (pResViewDesc != nullptr) ? *pResViewDesc : hipResourceViewDesc{}}); return hipSuccess; } -hipError_t hipUnbindTexture(const textureReference* tex) { - HIP_INIT_API(NONE, tex); +hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, + const hipResourceDesc* pResDesc, + const hipTextureDesc* pTexDesc, + const hipResourceViewDesc* pResViewDesc) { + HIP_INIT_API(hipCreateTextureObject, pTexObject, pResDesc, pTexDesc, pResViewDesc); - ihipDestroyTextureObject(reinterpret_cast(tex->textureObject)); - - HIP_RETURN(hipSuccess); + HIP_RETURN(ihipCreateTextureObject(pTexObject, pResDesc, pTexDesc, pResViewDesc)); } -hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_const_t array) { - HIP_INIT_API(NONE, desc, array); - if (desc != nullptr) { - *desc = array->desc; +hipError_t ihipDestroyTextureObject(hipTextureObject_t texObject) { + if (texObject == nullptr) { + return hipErrorInvalidValue; } - HIP_RETURN(hipSuccess); + hip::TextureObject* hipTexObject = reinterpret_cast(texObject); + const hipResourceType type = hipTexObject->resDesc.resType; + const bool isImageFromBuffer = (type == hipResourceTypeLinear) || (type == hipResourceTypePitch2D); + const bool isImageView = ((type == hipResourceTypeArray) || (type == hipResourceTypeMipmappedArray)) && + !hipTexObject->image->isParent(); + if (isImageFromBuffer || isImageView) { + hipTexObject->image->release(); + } + + // TODO Should call ihipFree() to not polute the api trace. + return hipFree(hipTexObject); } -hipError_t hipGetTextureAlignmentOffset(size_t* offset, const textureReference* tex) { - HIP_INIT_API(NONE, offset, tex); +hipError_t hipDestroyTextureObject(hipTextureObject_t texObject) { + HIP_INIT_API(hipDestroyTextureObject, texObject); - if ((offset == nullptr) || (tex == nullptr)) { + HIP_RETURN(ihipDestroyTextureObject(texObject)); +} + + +hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc, + hipTextureObject_t texObject) { + HIP_INIT_API(hipGetTextureObjectResourceDesc, pResDesc, texObject); + + if ((pResDesc == nullptr) || + (texObject == nullptr)) { HIP_RETURN(hipErrorInvalidValue); } + hip::TextureObject* hipTexObject = reinterpret_cast(texObject); + *pResDesc = hipTexObject->resDesc; + + HIP_RETURN(hipSuccess); +} + +hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc, + hipTextureObject_t texObject) { + HIP_INIT_API(hipGetTextureObjectResourceViewDesc, pResViewDesc, texObject); + + if ((pResViewDesc == nullptr) || + (texObject == nullptr)) { + HIP_RETURN(hipErrorInvalidValue); + } + + hip::TextureObject* hipTexObject = reinterpret_cast(texObject); + *pResViewDesc = hipTexObject->resViewDesc; + + HIP_RETURN(hipSuccess); +} + +hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc, + hipTextureObject_t texObject) { + HIP_INIT_API(hipGetTextureObjectTextureDesc, pTexDesc, texObject); + + if ((pTexDesc == nullptr) || + (texObject == nullptr)) { + HIP_RETURN(hipErrorInvalidValue); + } + + hip::TextureObject* hipTexObject = reinterpret_cast(texObject); + *pTexDesc = hipTexObject->texDesc; + + HIP_RETURN(hipSuccess); +} + +hipError_t ihipBindTexture(size_t* offset, + const textureReference* texref, + const void* devPtr, + const hipChannelFormatDesc* desc, + size_t size, + hipTextureReadMode readMode) { + if ((offset == nullptr) || + (texref == nullptr) || + (devPtr == nullptr) || + (desc == nullptr)) { + return hipErrorInvalidValue; + } + + // Any previous address or HIP array state associated with the texture reference is superseded by this function. + // Any memory previously bound to hTexRef is unbound. + // No need to check for errors. + ihipDestroyTextureObject(texref->textureObject); + + // If the device memory pointer was returned from hipMalloc(), + // the offset is guaranteed to be 0 and NULL may be passed as the offset parameter. + // TODO enforce alignment on devPtr. + if (offset != nullptr) { + *offset = 0; + } + + hipResourceDesc resDesc = {}; + resDesc.resType = hipResourceTypeLinear; + resDesc.res.linear.devPtr = const_cast(devPtr); + resDesc.res.linear.desc = *desc; + resDesc.res.linear.sizeInBytes = size; + + hipTextureDesc texDesc = hip::getTextureDesc(texref, readMode); + + return ihipCreateTextureObject(const_cast(&texref->textureObject), &resDesc, &texDesc, nullptr); +} + +hipError_t ihipBindTexture2D(size_t* offset, + const textureReference* texref, + const void* devPtr, + const hipChannelFormatDesc* desc, + size_t width, + size_t height, + size_t pitch, + hipTextureReadMode readMode) { + if ((offset == nullptr) || + (texref == nullptr) || + (devPtr == nullptr) || + (desc == nullptr)) { + return hipErrorInvalidValue; + } + + // Any previous address or HIP array state associated with the texture reference is superseded by this function. + // Any memory previously bound to hTexRef is unbound. + // No need to check for errors. + ihipDestroyTextureObject(texref->textureObject); + + // If the device memory pointer was returned from hipMalloc(), + // the offset is guaranteed to be 0 and NULL may be passed as the offset parameter. + // TODO enforce alignment on devPtr. + if (offset != nullptr) { + *offset = 0; + } + + hipResourceDesc resDesc = {}; + resDesc.resType = hipResourceTypePitch2D; + resDesc.res.pitch2D.devPtr = const_cast(devPtr); + resDesc.res.pitch2D.desc = *desc; + resDesc.res.pitch2D.width = width; + resDesc.res.pitch2D.height = height; + resDesc.res.pitch2D.pitchInBytes = pitch; + + hipTextureDesc texDesc = hip::getTextureDesc(texref, readMode); + + return ihipCreateTextureObject(const_cast(&texref->textureObject), &resDesc, &texDesc, nullptr); +} + +hipError_t hipBindTexture2D(size_t* offset, + textureReference* texref, + const void* devPtr, + const hipChannelFormatDesc* desc, + size_t width, + size_t height, + size_t pitch) { + HIP_INIT_API(hipBindTexture2D, offset, texref, devPtr, desc, width, height, pitch); + + // TODO need compiler support to extract the read mode from textureReference. + HIP_RETURN(ihipBindTexture2D(offset, texref, devPtr, desc, width, height, pitch, hipReadModeElementType)); +} + +hipError_t ihipBindTextureToArray(const textureReference* texref, + hipArray_const_t array, + const hipChannelFormatDesc* desc, + hipTextureReadMode readMode) { + if ((texref == nullptr) || + (array == nullptr) || + (desc == nullptr)) { + return hipErrorInvalidValue; + } + + // Any previous address or HIP array state associated with the texture reference is superseded by this function. + // Any memory previously bound to hTexRef is unbound. + // No need to check for errors. + ihipDestroyTextureObject(texref->textureObject); + + hipResourceDesc resDesc = {}; + resDesc.resType = hipResourceTypeArray; + resDesc.res.array.array = const_cast(array); + + hipTextureDesc texDesc = hip::getTextureDesc(texref, readMode); + + hipResourceViewFormat format = hip::getResourceViewFormat(*desc); + hipResourceViewDesc resViewDesc = hip::getResourceViewDesc(array, format); + + return ihipCreateTextureObject(const_cast(&texref->textureObject), &resDesc, &texDesc, &resViewDesc); +} + +hipError_t hipBindTextureToArray(textureReference* texref, + hipArray_const_t array, + const hipChannelFormatDesc* desc) { + HIP_INIT_API(hipBindTextureToArray, texref, array, desc); + + // TODO need compiler support to extract the read mode from textureReference. + HIP_RETURN(ihipBindTextureToArray(texref, array, desc, hipReadModeElementType)); +} + +hipError_t ihipBindTextureImpl(TlsData *tls, + int dim, + hipTextureReadMode readMode, + size_t* offset, + const void* devPtr, + const hipChannelFormatDesc* desc, + size_t size, + textureReference* texref) { + HIP_INIT_API(ihipBindTextureImpl, tls, dim, readMode, offset, devPtr, desc, size, texref); + + (void)dim; // Silence compiler warnings. + + HIP_RETURN(ihipBindTexture(offset, texref, devPtr, desc, size, readMode)); +} + +hipError_t ihipBindTextureToArrayImpl(TlsData *tls, + int dim, + hipTextureReadMode readMode, + hipArray_const_t array, + const hipChannelFormatDesc& desc, + textureReference* texref) { + // TODO overload operator<<(ostream&, hipChannelFormatDesc&). + HIP_INIT_API(ihipBindTextureToArrayImpl, tls, dim, readMode, array, &desc, texref); + + (void)dim; // Silence compiler warnings. + + HIP_RETURN(ihipBindTextureToArray(texref, array, &desc, readMode)); +} + +hipError_t ihipBindTextureToMipmappedArray(textureReference* texref, + hipMipmappedArray_const_t mipmappedArray, + const hipChannelFormatDesc* desc, + hipTextureReadMode readMode) { + if ((texref == nullptr) || + (mipmappedArray == nullptr) || + (desc == nullptr)) { + return hipErrorInvalidValue; + } + + // Any previous address or HIP array state associated with the texture reference is superseded by this function. + // Any memory previously bound to hTexRef is unbound. + // No need to check for errors. + ihipDestroyTextureObject(texref->textureObject); + + hipResourceDesc resDesc = {}; + resDesc.resType = hipResourceTypeMipmappedArray; + resDesc.res.mipmap.mipmap = const_cast(mipmappedArray); + + hipTextureDesc texDesc = hip::getTextureDesc(texref, readMode); + + hipResourceViewFormat format = hip::getResourceViewFormat(*desc); + hipResourceViewDesc resViewDesc = hip::getResourceViewDesc(mipmappedArray, format); + + return ihipCreateTextureObject(const_cast(&texref->textureObject), &resDesc, &texDesc, &resViewDesc); +} + +hipError_t hipBindTextureToMipmappedArray(textureReference* texref, + hipMipmappedArray_const_t mipmappedArray, + const hipChannelFormatDesc* desc) { + HIP_INIT_API(hipBindTextureToMipmappedArray, texref, mipmappedArray, desc); + + // TODO need compiler support to extract the read mode from textureReference. + HIP_RETURN(ihipBindTextureToMipmappedArray(texref, mipmappedArray, desc, hipReadModeElementType)); +} + +hipError_t ihipBindTexture2DImpl(int dim, + hipTextureReadMode readMode, + size_t* offset, + const void* devPtr, + const hipChannelFormatDesc* desc, + size_t width, + size_t height, + textureReference* texref, + size_t pitch) { + HIP_INIT_API(ihipBindTexture2DImpl, dim, readMode, offset, devPtr, desc, width, height, texref, pitch); + + (void)dim; // Silence compiler warnings. + + HIP_RETURN(ihipBindTexture2D(offset, texref, devPtr, desc, width, height, pitch, readMode)); +} + +hipError_t ihipUnbindTextureImpl(const hipTextureObject_t& textureObject) { + // TODO overload operator<<(ostream&, hipTextureObject_t&). + HIP_INIT_API(ihipUnbindTextureImpl, &textureObject); + + HIP_RETURN(ihipDestroyTextureObject(textureObject)); +} + +hipError_t hipUnbindTexture(const textureReference* texref) { + HIP_INIT_API(hipUnbindTexture, texref); + + if (texref == nullptr) { + return hipErrorInvalidValue; + } + + HIP_RETURN(ihipDestroyTextureObject(texref->textureObject)); +} + +hipError_t hipBindTexture(size_t* offset, + textureReference* texref, + const void* devPtr, + const hipChannelFormatDesc* desc, + size_t size) { + HIP_INIT_API(hipBindTexture, offset, texref, devPtr, desc, size); + + // TODO need compiler support to extract the read mode from textureReference. + HIP_RETURN(ihipBindTexture(offset, texref, devPtr, desc, size, hipReadModeElementType)); +} + +hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, + hipArray_const_t array) { + HIP_INIT_API(hipGetChannelDesc, desc, array); + + if ((desc == nullptr) || + (array == nullptr)) { + HIP_RETURN(hipErrorInvalidValue); + } + + // It is UB to call hipGetChannelDesc() on an array created via hipArrayCreate()/hipArray3DCreate(). + // This is due to hip not differentiating between runtime and driver types. + *desc = array->desc; + + HIP_RETURN(hipSuccess); +} + +hipError_t hipGetTextureAlignmentOffset(size_t* offset, + const textureReference* texref) { + HIP_INIT_API(hipGetTextureAlignmentOffset, offset, texref); + + if ((offset == nullptr) || + (texref == nullptr)) { + HIP_RETURN(hipErrorInvalidValue); + } + + // TODO enforce alignment on devPtr. *offset = 0; HIP_RETURN(hipSuccess); } -hipError_t hipGetTextureReference(const textureReference** tex, const void* symbol) { - HIP_INIT_API(NONE, tex, symbol); +hipError_t hipGetTextureReference(const textureReference** texref, const void* symbol) { + HIP_INIT_API(hipGetTextureReference, texref, symbol); assert(0 && "Unimplemented"); HIP_RETURN(hipErrorNotSupported); } -hipError_t hipTexRefSetFormat(textureReference* tex, hipArray_Format fmt, int NumPackedComponents) { - HIP_INIT_API(NONE, tex, fmt, NumPackedComponents); +hipError_t hipTexRefSetFormat(textureReference* texRef, + hipArray_Format fmt, + int NumPackedComponents) { + HIP_INIT_API(hipTexRefSetFormat, texRef, fmt, NumPackedComponents); - if (tex == nullptr) { - HIP_RETURN(hipErrorInvalidImage); - } - - tex->format = fmt; - tex->numChannels = NumPackedComponents; - - HIP_RETURN(hipSuccess); -} - -hipError_t hipTexRefSetFlags(textureReference* tex, unsigned int flags) { - HIP_INIT_API(NONE, tex, flags); - - if (tex == nullptr) { - HIP_RETURN(hipErrorInvalidImage); - } - - tex->normalized = flags; - - HIP_RETURN(hipSuccess); -} - -hipError_t hipTexRefSetFilterMode(textureReference* tex, hipTextureFilterMode fm) { - HIP_INIT_API(NONE, tex, fm); - - if (tex == nullptr) { - HIP_RETURN(hipErrorInvalidImage); - } - - tex->filterMode = fm; - - HIP_RETURN(hipSuccess); -} - -hipError_t hipTexRefGetAddressMode(hipTextureAddressMode* am, textureReference tex, int dim) { - HIP_INIT_API(NONE, am, &tex, dim); - - if ((am == nullptr) || (dim >= 3)) { + if (texRef == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - *am = tex.addressMode[dim]; + texRef->format = fmt; + texRef->numChannels = NumPackedComponents; HIP_RETURN(hipSuccess); } -hipError_t hipTexRefSetAddressMode(textureReference* tex, int dim, hipTextureAddressMode am) { - HIP_INIT_API(NONE, tex, dim, am); +hipError_t hipTexRefSetFlags(textureReference* texRef, + unsigned int Flags) { + HIP_INIT_API(hipTexRefSetFlags, texRef, Flags); - if (tex == nullptr) { - HIP_RETURN(hipErrorInvalidImage); - } - - tex->addressMode[dim] = am; - - HIP_RETURN(hipSuccess); -} - -hipError_t hipTexRefGetArray(hipArray_t* array, textureReference tex) { - HIP_INIT_API(NONE, array, &tex); - - hip::TextureObject* texture = nullptr; - - if ((array == nullptr) || (*array == nullptr)) { - HIP_RETURN(hipErrorInvalidImage); - } - - texture = reinterpret_cast(tex.textureObject); - if(hipResourceTypeArray != texture->resDesc.resType){ + if (texRef == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - if (texture->resDesc.res.array.array == nullptr) { - HIP_RETURN(hipErrorInvalidValue); - } - - **array = *(texture->resDesc.res.array.array); + // TODO add textureReference::flags. + // Using textureReference::normalized for this purpose is OK for now, + // because calling hipTexRefGetFlags() on a textureRefence after hipBindTexture() is UB + // due to HIP not differentiating between runtime and driver api. + texRef->normalized = Flags; HIP_RETURN(hipSuccess); } -hipError_t hipTexRefSetArray(textureReference* tex, hipArray_const_t array, unsigned int flags) { - HIP_INIT_API(NONE, tex, array, flags); +hipError_t hipTexRefSetFilterMode(textureReference* texRef, + hipTextureFilterMode fm) { + HIP_INIT_API(hipTexRefSetFilterMode, texRef, fm); - size_t offset = 0; - cl_mem_object_type clType; - - if ((tex == nullptr) || (array == nullptr)) { - HIP_RETURN(hipErrorInvalidImage); + if (texRef == nullptr) { + HIP_RETURN(hipErrorInvalidValue); } - switch(array->textureType) { - case hipTextureType3D: - clType = CL_MEM_OBJECT_IMAGE3D; - break; - case hipTextureType2D: - clType = CL_MEM_OBJECT_IMAGE2D; - break; - case hipTextureType1D: - clType = CL_MEM_OBJECT_IMAGE1D; - break; - default: - HIP_RETURN(hipErrorInvalidValue); - } - HIP_RETURN(ihipBindTexture(clType, &offset, tex, array->data, array->desc, array->width, - array->height, array->depth)); -} - -hipError_t hipTexRefGetAddress(hipDeviceptr_t* dev_ptr, textureReference tex) { - HIP_INIT_API(NONE, dev_ptr, &tex); - - hip::TextureObject* texture = nullptr; - device::Memory* dev_mem = nullptr; - - texture = reinterpret_cast(tex.textureObject); - if ((texture == nullptr) || (texture->image == nullptr)) { - HIP_RETURN(hipErrorInvalidImage); - } - - dev_mem = texture->image->getDeviceMemory(*hip::getCurrentDevice()->devices()[0]); - if (dev_mem == nullptr) { - HIP_RETURN(hipErrorInvalidImage); - } - - *dev_ptr = reinterpret_cast(dev_mem->virtualAddress()); + texRef->filterMode = fm; HIP_RETURN(hipSuccess); } -hipError_t hipTexRefSetAddress(size_t* offset, textureReference* tex, hipDeviceptr_t devPtr, - size_t size) { - HIP_INIT_API(NONE, offset, tex, devPtr, size); +hipError_t hipTexRefGetAddressMode(hipTextureAddressMode* pam, + textureReference texRef, + int dim) { + // TODO overload operator<<(ostream&, textureReference&). + HIP_INIT_API(hipTexRefGetAddressMode, pam, &texRef, dim); - if (tex == nullptr) { - HIP_RETURN(hipErrorInvalidImage); - } - - cl_image_format image_format; - getDrvChannelOrderAndType(tex->format, tex->numChannels, - &image_format.image_channel_order, &image_format.image_channel_data_type); - const amd::Image::Format imageFormat(image_format); - - HIP_RETURN(ihipBindTexture(CL_MEM_OBJECT_IMAGE1D, offset, tex, devPtr, tex->channelDesc, size / imageFormat.getElementSize(), 1, size)); -} - -hipError_t hipTexRefSetAddress2D(textureReference* tex, const HIP_ARRAY_DESCRIPTOR* desc, - hipDeviceptr_t devPtr, size_t pitch) { - HIP_INIT_API(NONE, tex, desc, devPtr, pitch); - - if (desc == nullptr) { + if (pam == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - size_t offset; - HIP_RETURN(ihipBindTexture(CL_MEM_OBJECT_IMAGE2D, &offset, tex, devPtr, tex->channelDesc, desc->Width, desc->Height, pitch)); + // Currently, the only valid value for dim are 0 and 1. + if ((dim != 0) || (dim != 1)) { + HIP_RETURN(hipErrorInvalidValue); + } + + *pam = texRef.addressMode[dim]; + + HIP_RETURN(hipSuccess); } + +hipError_t hipTexRefSetAddressMode(textureReference* texRef, + int dim, + hipTextureAddressMode am) { + HIP_INIT_API(hipTexRefSetAddressMode, texRef, dim, am); + + if (texRef == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + if ((dim < 0) || (dim > 2)) { + HIP_RETURN(hipErrorInvalidValue); + } + + texRef->addressMode[dim] = am; + + HIP_RETURN(hipSuccess); +} + +hipError_t hipTexRefGetArray(hipArray_t* pArray, + textureReference texRef) { + // TODO overload operator<<(ostream&, textureReference&). + HIP_INIT_API(hipTexRefGetArray, pArray, &texRef); + + if (pArray == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + hipResourceDesc resDesc = {}; + // TODO use ihipGetTextureObjectResourceDesc() to not pollute the API trace. + hipError_t error = hipGetTextureObjectResourceDesc(&resDesc, texRef.textureObject); + if (error != hipSuccess) { + return HIP_RETURN(error); + } + + switch (resDesc.resType) { + case hipResourceTypeLinear: + case hipResourceTypePitch2D: + case hipResourceTypeMipmappedArray: + HIP_RETURN(hipErrorInvalidValue); + case hipResourceTypeArray: + *pArray = resDesc.res.array.array; + break; + } + + HIP_RETURN(hipSuccess); +} + +hipError_t hipTexRefSetArray(textureReference* texRef, + hipArray_const_t array, + unsigned int flags) { + HIP_INIT_API(hipTexRefSetArray, texRef, array, flags); + + if ((texRef == nullptr) || + (array == nullptr)) { + HIP_RETURN(hipErrorInvalidValue); + } + + if (flags != HIP_TRSA_OVERRIDE_FORMAT) { + HIP_RETURN(hipErrorInvalidValue); + } + + // Any previous address or HIP array state associated with the texture reference is superseded by this function. + // Any memory previously bound to hTexRef is unbound. + // No need to check for errors. + ihipDestroyTextureObject(texRef->textureObject); + + hipResourceDesc resDesc = {}; + resDesc.resType = hipResourceTypeArray; + resDesc.res.array.array = const_cast(array); + + // TODO need compiler support to extract the read mode from textureReference. + hipTextureDesc texDesc = hip::getTextureDesc(texRef, hipReadModeElementType); + + hipResourceViewFormat format = hip::getResourceViewFormat(hip::getChannelFormatDesc(texRef->numChannels, texRef->format)); + hipResourceViewDesc resViewDesc = hip::getResourceViewDesc(array, format); + + HIP_RETURN(ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, &resViewDesc)); +} + +hipError_t hipTexRefGetAddress(hipDeviceptr_t* dptr, + textureReference texRef) { + // TODO overload operator<<(ostream&, textureReference&). + HIP_INIT_API(hipTexRefGetAddress, dptr, &texRef); + + if (dptr == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + hipResourceDesc resDesc = {}; + // TODO use ihipGetTextureObjectResourceDesc() to not pollute the API trace. + hipError_t error = hipGetTextureObjectResourceDesc(&resDesc, texRef.textureObject); + if (error != hipSuccess) { + return HIP_RETURN(error); + } + + switch (resDesc.resType) { + // Need to verify. + // If the texture reference is not bound to any device memory range, + // return hipErroInvalidValue. + case hipResourceTypeArray: + case hipResourceTypeMipmappedArray: + HIP_RETURN(hipErrorInvalidValue); + case hipResourceTypeLinear: + *dptr = resDesc.res.linear.devPtr; + break; + case hipResourceTypePitch2D: + *dptr = resDesc.res.pitch2D.devPtr; + break; + } + + HIP_RETURN(hipSuccess); +} + +hipError_t hipTexRefSetAddress(size_t* ByteOffset, + textureReference* texRef, + hipDeviceptr_t dptr, + size_t bytes) { + HIP_INIT_API(hipTexRefSetAddress, ByteOffset, texRef, dptr, bytes); + + if (texRef == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + // Any previous address or HIP array state associated with the texture reference is superseded by this function. + // Any memory previously bound to hTexRef is unbound. + // No need to check for errors. + ihipDestroyTextureObject(texRef->textureObject); + + // If the device memory pointer was returned from hipMemAlloc(), + // the offset is guaranteed to be 0 and NULL may be passed as the ByteOffset parameter. + // TODO enforce alignment on devPtr. + if (ByteOffset != nullptr) { + *ByteOffset = 0; + } + + hipResourceDesc resDesc = {}; + resDesc.resType = hipResourceTypeLinear; + resDesc.res.linear.devPtr = dptr; + resDesc.res.linear.desc = hip::getChannelFormatDesc(texRef->numChannels, texRef->format); + resDesc.res.linear.sizeInBytes = bytes; + + // TODO add textureReference::flags. + // Using textureReference::normalized for this purpose is OK for now, + // because calling hipTexRefGetFlags() on a textureRefence after hipBindTexture() + // due to HIP not differentiating between runtime and driver api. + hipTextureReadMode readMode = hip::getReadMode(texRef->normalized); + texRef->sRGB = hip::getSRGB(texRef->normalized); + texRef->normalized = hip::getNormalizedCoords(texRef->normalized); + hipTextureDesc texDesc = hip::getTextureDesc(texRef, readMode); + + HIP_RETURN(ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, nullptr)); +} + +hipError_t hipTexRefSetAddress2D(textureReference* texRef, + const HIP_ARRAY_DESCRIPTOR* desc, + hipDeviceptr_t dptr, + size_t Pitch) { + HIP_INIT_API(hipTexRefSetAddress2D, texRef, desc, dptr, Pitch); + + if ((texRef == nullptr) || + (desc == nullptr)) { + HIP_RETURN(hipErrorInvalidValue); + } + + // Any previous address or HIP array state associated with the texture reference is superseded by this function. + // Any memory previously bound to hTexRef is unbound. + // No need to check for errors. + ihipDestroyTextureObject(texRef->textureObject); + + hipResourceDesc resDesc = {}; + resDesc.resType = hipResourceTypePitch2D; + resDesc.res.linear.devPtr = dptr; + resDesc.res.linear.desc = hip::getChannelFormatDesc(desc->NumChannels, desc->Format); // Need to verify. + resDesc.res.pitch2D.width = desc->Width; + resDesc.res.pitch2D.height = desc->Height; + resDesc.res.pitch2D.pitchInBytes = Pitch; + + // TODO add textureReference::flags. + // Using textureReference::normalized for this purpose is OK for now, + // because calling hipTexRefGetFlags() on a textureRefence after hipBindTexture() + // due to HIP not differentiating between runtime and driver api. + hipTextureReadMode readMode = hip::getReadMode(texRef->normalized); + texRef->sRGB = hip::getSRGB(texRef->normalized); + texRef->normalized = hip::getNormalizedCoords(texRef->normalized); + hipTextureDesc texDesc = hip::getTextureDesc(texRef, readMode); + + HIP_RETURN(ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, nullptr)); +} + +hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) { + return {x, y, z, w, f}; +} + +hipError_t ihipBindTextureToMipmappedArrayImpl(TlsData *tls, + int dim, + hipTextureReadMode readMode, + hipMipmappedArray_const_t mipmappedArray, + const struct hipChannelFormatDesc& desc, + textureReference* texref) { + // TODO overload operator<<(ostream&, hipChannelFormatDesc&). + HIP_INIT_API(ihipBindTextureToMipmappedArrayImpl, tls, dim, readMode, mipmappedArray, &desc, texref); + + (void)dim; // Silence compiler warnings. + + HIP_RETURN(ihipBindTextureToMipmappedArray(texref, mipmappedArray, &desc, readMode)); +} + +hipError_t hipTexRefGetBorderColor(float* pBorderColor, + textureReference texRef) { + // TODO overload operator<<(ostream&, textureReference&). + HIP_INIT_API(hipTexRefGetBorderColor, pBorderColor, &texRef); + + if (pBorderColor == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + // TODO add textureReference::borderColor. + assert(false && "textureReference::borderColor is missing in header"); + // std::memcpy(pBorderColor, texRef.borderColor, sizeof(texRef.borderColor)); + + HIP_RETURN(hipSuccess); +} + +hipError_t hipTexRefGetFilterMode(hipTextureFilterMode* pfm, + textureReference texRef) { + // TODO overload operator<<(ostream&, textureReference&). + HIP_INIT_API(hipTexRefGetFilterMode, pfm, &texRef); + + if (pfm == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + *pfm = texRef.filterMode; + + HIP_RETURN(hipSuccess); +} + +hipError_t hipTexRefGetFlags(unsigned int* pFlags, + textureReference texRef) { + // TODO overload operator<<(ostream&, textureReference&). + HIP_INIT_API(hipTexRefGetFlags, pFlags, &texRef); + + if (pFlags == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + // TODO add textureReference::flags. + // Using textureReference::normalized for this purpose is OK for now, + // because calling hipTexRefGetFlags() on a textureRefence after hipBindTexture() + // due to HIP not differentiating between runtime and driver api. + *pFlags = texRef.normalized; + + HIP_RETURN(hipSuccess); +} + +hipError_t hipTexRefGetFormat(hipArray_Format* pFormat, + int* pNumChannels, + textureReference texRef) { + // TODO overload operator<<(ostream&, textureReference&). + HIP_INIT_API(hipTexRefGetFormat, pFormat, pNumChannels, &texRef); + + if ((pFormat == nullptr) || + (pNumChannels == nullptr)) { + HIP_RETURN(hipErrorInvalidValue); + } + + *pFormat = texRef.format; + *pNumChannels = texRef.numChannels; + + HIP_RETURN(hipSuccess); +} + +hipError_t hipTexRefGetMaxAnisotropy(int* pmaxAnsio, + textureReference texRef) { + // TODO overload operator<<(ostream&, textureReference&). + HIP_INIT_API(hipTexRefGetMaxAnisotropy, pmaxAnsio, &texRef); + + if (pmaxAnsio == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + *pmaxAnsio = texRef.maxAnisotropy; + + HIP_RETURN(hipErrorInvalidValue); +} + +hipError_t hipTexRefGetMipmapFilterMode(hipTextureFilterMode* pfm, + textureReference texRef) { + // TODO overload operator<<(ostream&, textureReference&). + HIP_INIT_API(hipTexRefGetMipmapFilterMode, pfm, &texRef); + + if (pfm == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + *pfm = texRef.mipmapFilterMode; + + HIP_RETURN(hipErrorInvalidValue); +} + +hipError_t hipTexRefGetMipmapLevelBias(float* pbias, + textureReference texRef) { + // TODO overload operator<<(ostream&, textureReference&). + HIP_INIT_API(hipTexRefGetMipmapLevelBias, pbias, &texRef); + + if (pbias == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + *pbias = texRef.mipmapLevelBias; + + HIP_RETURN(hipErrorInvalidValue); +} + +hipError_t hipTexRefGetMipmapLevelClamp(float* pminMipmapLevelClamp, + float* pmaxMipmapLevelClamp, + textureReference texRef) { + // TODO overload operator<<(ostream&, textureReference&). + HIP_INIT_API(hipTexRefGetMipmapLevelClamp, pminMipmapLevelClamp, pmaxMipmapLevelClamp, &texRef); + + if ((pminMipmapLevelClamp == nullptr) || + (pmaxMipmapLevelClamp == nullptr)){ + HIP_RETURN(hipErrorInvalidValue); + } + + *pminMipmapLevelClamp = texRef.minMipmapLevelClamp; + *pmaxMipmapLevelClamp = texRef.maxMipmapLevelClamp; + + HIP_RETURN(hipErrorInvalidValue); +} + +hipError_t hipTexRefGetMipMappedArray(hipMipmappedArray_t* pArray, + textureReference texRef) { + // TODO overload operator<<(ostream&, textureReference&). + HIP_INIT_API(hipTexRefGetMipMappedArray, pArray, &texRef); + + if (pArray == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + hipResourceDesc resDesc = {}; + // TODO use ihipGetTextureObjectResourceDesc() to not pollute the API trace. + hipError_t error = hipGetTextureObjectResourceDesc(&resDesc, texRef.textureObject); + if (error != hipSuccess) { + return HIP_RETURN(error); + } + + switch (resDesc.resType) { + case hipResourceTypeLinear: + case hipResourceTypePitch2D: + case hipResourceTypeArray: + HIP_RETURN(hipErrorInvalidValue); + case hipResourceTypeMipmappedArray: + *pArray = resDesc.res.mipmap.mipmap; + break; + } + + HIP_RETURN(hipSuccess); +} + +hipError_t hipTexRefSetBorderColor(textureReference* texRef, + float* pBorderColor) { + HIP_INIT_API(hipTexRefSetBorderColor, texRef, pBorderColor); + + if ((texRef == nullptr) || + (pBorderColor == nullptr)) { + HIP_RETURN(hipErrorInvalidValue); + } + + // TODO add textureReference::borderColor. + assert(false && "textureReference::borderColor is missing in header"); + // std::memcpy(texRef.borderColor, pBorderColor, sizeof(texRef.borderColor)); + + HIP_RETURN(hipSuccess); +} + +hipError_t hipTexRefSetMaxAnisotropy(textureReference* texRef, + unsigned int maxAniso) { + HIP_INIT_API(hipTexRefSetMaxAnisotropy, texRef, maxAniso); + + if (texRef == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + texRef->maxAnisotropy = maxAniso; + + HIP_RETURN(hipSuccess); +} + +hipError_t hipTexRefSetMipmapFilterMode(textureReference* texRef, + hipTextureFilterMode fm) { + HIP_INIT_API(hipTexRefSetMipmapFilterMode, texRef, fm); + + if (texRef == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + texRef->mipmapFilterMode = fm; + + HIP_RETURN(hipSuccess); +} + +hipError_t hipTexRefSetMipmapLevelBias(textureReference* texRef, + float bias) { + HIP_INIT_API(hipTexRefSetMipmapLevelBias, texRef, bias); + + if (texRef == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + texRef->mipmapLevelBias = bias; + + HIP_RETURN(hipSuccess); +} + +hipError_t hipTexRefSetMipmapLevelClamp(textureReference* texRef, + float minMipMapLevelClamp, + float maxMipMapLevelClamp) { + HIP_INIT_API(hipTexRefSetMipmapLevelClamp, minMipMapLevelClamp, maxMipMapLevelClamp); + + if (texRef == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + texRef->minMipmapLevelClamp = minMipMapLevelClamp; + texRef->maxMipmapLevelClamp = maxMipMapLevelClamp; + + HIP_RETURN(hipSuccess); +} + +hipError_t hipTexRefSetMipmappedArray(textureReference* texRef, + hipMipmappedArray* mipmappedArray, + unsigned int Flags) { + HIP_INIT_API(hipTexRefSetMipmappedArray, texRef, mipmappedArray, Flags); + + if ((texRef == nullptr) || + (mipmappedArray == nullptr)) { + HIP_RETURN(hipErrorInvalidValue); + } + + if (Flags != HIP_TRSA_OVERRIDE_FORMAT) { + HIP_RETURN(hipErrorInvalidValue); + } + + // Any previous address or HIP array state associated with the texture reference is superseded by this function. + // Any memory previously bound to hTexRef is unbound. + // No need to check for errors. + ihipDestroyTextureObject(texRef->textureObject); + + hipResourceDesc resDesc = {}; + resDesc.resType = hipResourceTypeMipmappedArray; + resDesc.res.mipmap.mipmap = mipmappedArray; + + // TODO need compiler support to extract the read mode from textureReference. + hipTextureDesc texDesc = hip::getTextureDesc(texRef, hipReadModeElementType); + + hipResourceViewFormat format = hip::getResourceViewFormat(hip::getChannelFormatDesc(texRef->numChannels, texRef->format)); + hipResourceViewDesc resViewDesc = hip::getResourceViewDesc(mipmappedArray, format); + + HIP_RETURN(ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, &resViewDesc)); +} \ No newline at end of file