From 67a5235c22f1e8e24b7e28bc6fc6826183cd1975 Mon Sep 17 00:00:00 2001 From: foreman Date: Fri, 22 Feb 2019 19:43:36 -0500 Subject: [PATCH] P4 to Git Change 1747723 by cpaquot@cpaquot-ocl-lc-lnx on 2019/02/22 19:30:33 SWDEV-145570 - [HIP] Texture object implementation hipTextureObject2D test passes on PAL backend. Still needs to have sampler handled properly on ROCm backend. Affected files ... ... //depot/stg/opencl/drivers/opencl/api/hip/hip_internal.hpp#22 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_texture.cpp#10 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/device.hpp#333 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldevice.hpp#36 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palmemory.hpp#10 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocmemory.hpp#14 edit [ROCm/hip commit: c49cb353ca5471fa26391218d1689db0d1d4c353] --- projects/hip/api/hip/hip_internal.hpp | 1 + projects/hip/api/hip/hip_texture.cpp | 123 +++++++++++++++++++++++--- 2 files changed, 113 insertions(+), 11 deletions(-) diff --git a/projects/hip/api/hip/hip_internal.hpp b/projects/hip/api/hip/hip_internal.hpp index 8d918c608e..ed32d90476 100644 --- a/projects/hip/api/hip/hip_internal.hpp +++ b/projects/hip/api/hip/hip_internal.hpp @@ -88,6 +88,7 @@ namespace hip { extern std::vector g_devices; extern hipError_t ihipDeviceGetCount(int* count); extern int ihipGetDevice(); +extern hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags); extern amd::Memory* getMemoryObject(const void* ptr, size_t& offset); #define HIP_RETURN(ret) \ diff --git a/projects/hip/api/hip/hip_texture.cpp b/projects/hip/api/hip/hip_texture.cpp index b3b5c16c10..7fca3f506d 100644 --- a/projects/hip/api/hip/hip_texture.cpp +++ b/projects/hip/api/hip/hip_texture.cpp @@ -23,6 +23,16 @@ THE SOFTWARE. #include #include #include "hip_internal.hpp" +#include "platform/sampler.hpp" + +namespace hip { + struct TextureObject { + uint32_t imageSRD[HIP_IMAGE_OBJECT_SIZE_DWORD]; + uint32_t samplerSRD[HIP_SAMPLER_OBJECT_SIZE_DWORD]; + amd::Image* image; + amd::Sampler* sampler; + }; +}; void getDrvChannelOrderAndType(const enum hipArray_Format Format, unsigned int NumChannels, cl_channel_order* channelOrder, @@ -137,6 +147,68 @@ void getChannelOrderAndType(const hipChannelFormatDesc& desc, enum hipTextureRea } } +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; + } + + 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; + } + amd::Sampler* sampler = new amd::Sampler(*hip::getCurrentContext(), + normalizedCoords == CL_TRUE, + address_mode, filter_mode, CL_FILTER_NONE, 0.f, CL_MAXFLOAT); + if (sampler == nullptr) { + return nullptr; + } + if (!sampler->create()) { + delete sampler; + return nullptr; + } + return sampler; +} + +hip::TextureObject* ihipCreateTextureObject(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::getCurrentContext()->devices()[0]); + memcpy(texture->imageSRD, imageMem->cpuSrd(), sizeof(uint32_t)*HIP_IMAGE_OBJECT_SIZE_DWORD); + texture->image = ℑ + + device::Sampler* devSampler = sampler.getDeviceSampler(*hip::getCurrentContext()->devices()[0]); + memcpy(texture->samplerSRD, devSampler->hwState(), sizeof(uint32_t)*HIP_SAMPLER_OBJECT_SIZE_DWORD); + texture->sampler = &sampler; + + return texture; +} + hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResourceDesc* pResDesc, const hipTextureDesc* pTexDesc, const hipResourceViewDesc* pResViewDesc) { @@ -186,13 +258,19 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou assert(0); break; case hipResourceTypeLinear: - assert(pResViewDesc == nullptr); - memory = getMemoryObject(pResDesc->res.linear.devPtr, offset); + { + assert(pResViewDesc == nullptr); + memory = getMemoryObject(pResDesc->res.linear.devPtr, offset); - image = new (*hip::getCurrentContext()) amd::Image(*memory->asBuffer(), - CL_MEM_OBJECT_IMAGE1D, memory->getMemFlags(), imageFormat, - pResDesc->res.linear.sizeInBytes / imageFormat.getElementSize(), 1, 1, - pResDesc->res.linear.sizeInBytes, 0); + 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::getCurrentContext()) amd::Image(*memory->asBuffer(), + CL_MEM_OBJECT_IMAGE2D, memory->getMemFlags(), imageFormat, + pResDesc->res.linear.sizeInBytes / imageFormat.getElementSize(), 1, 1, + pResDesc->res.linear.sizeInBytes, 0); + } break; case hipResourceTypePitch2D: assert(pResViewDesc == nullptr); @@ -205,15 +283,32 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou break; default: HIP_RETURN(hipErrorInvalidValue); } - *pTexObject = reinterpret_cast(as_cl(image)); - HIP_RETURN(hipErrorUnknown); + if (!image->create()) { + delete image; + HIP_RETURN(hipErrorUnknown); + } + + amd::Sampler* sampler = fillSamplerDescriptor(pTexDesc->addressMode[0], pTexDesc->filterMode, pTexDesc->normalizedCoords); + + *pTexObject = reinterpret_cast(ihipCreateTextureObject(*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(textureObject); - as_amd(reinterpret_cast(textureObject))->release(); + hip::TextureObject* texture = reinterpret_cast(textureObject); + + ihipDestroyTextureObject(texture); HIP_RETURN(hipSuccess); } @@ -267,11 +362,17 @@ hipError_t ihipBindTexture(cl_mem_object_type type, amd::Image* image = new (*hip::getCurrentContext()) amd::Image(*memory->asBuffer(), type, memory->getMemFlags(), imageFormat, width, height, 1, pitch, 0); + if (!image->create()) { + delete image; + return hipErrorUnknown; + } + *offset = 0; if (tex->textureObject) { - as_amd(reinterpret_cast(tex->textureObject))->release(); + ihipDestroyTextureObject(reinterpret_cast(tex->textureObject)); } - tex->textureObject = reinterpret_cast(as_cl(image)); + amd::Sampler* sampler = fillSamplerDescriptor(tex->addressMode[0], tex->filterMode, tex->normalized); + tex->textureObject = reinterpret_cast(ihipCreateTextureObject(*image, *sampler)); return hipSuccess; } return hipErrorUnknown;