From 69dcd2b862fa1d7fccdfbf147896d7962aa0f119 Mon Sep 17 00:00:00 2001 From: foreman Date: Mon, 2 Dec 2019 13:37:35 -0600 Subject: [PATCH] P4 to Git Change 2039536 by jujiang@JJ-HIP on 2019/12/02 14:31:13 SWDEV-214490 - Update HIP RT for texture3D in HIP/PAL on Windows -Update ihipBindTexture http://ocltc.amd.com/reviews/r/18333/ Affected files ... ... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#89 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_texture.cpp#28 edit --- api/hip/hip_memory.cpp | 39 +++---------------------- api/hip/hip_texture.cpp | 65 +++++++++++++++++++++++++++++++++++++++-- 2 files changed, 66 insertions(+), 38 deletions(-) diff --git a/api/hip/hip_memory.cpp b/api/hip/hip_memory.cpp index 825745dca9..bd82e2634d 100644 --- a/api/hip/hip_memory.cpp +++ b/api/hip/hip_memory.cpp @@ -38,6 +38,8 @@ extern void getDrvChannelOrderAndType(const enum hipArray_Format Format, 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) { @@ -881,24 +883,7 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con amd::HostQueue* queue = hip::getNullStream(); size_t dpitch = dst->width; - - switch (dst[0].desc.f) { - case hipChannelFormatKindSigned: - dpitch *= sizeof(int); - break; - case hipChannelFormatKindUnsigned: - dpitch *= sizeof(unsigned int); - break; - case hipChannelFormatKindFloat: - dpitch *= sizeof(float); - break; - case hipChannelFormatKindNone: - dpitch *= sizeof(size_t); - break; - default: - dpitch *= 1; - break; - } + getByteSizeFromChannelFormatKind(dst[0].desc.f, &dpitch); if ((wOffset + width > (dpitch)) || width > spitch) { HIP_RETURN(hipErrorInvalidDevicePointer); @@ -1133,23 +1118,7 @@ hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p) { size_t dstOrigin[3]; size_t region[3]; if (p->dstArray != nullptr) { - switch (p->dstArray->desc.f) { - 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; - } + getByteSizeFromChannelFormatKind(p->dstArray->desc.f, &byteSize); region[2] = p->extent.depth; region[1] = p->extent.height; region[0] = p->extent.width; diff --git a/api/hip/hip_texture.cpp b/api/hip/hip_texture.cpp index d0dac5d910..1edd0b8aca 100644 --- a/api/hip/hip_texture.cpp +++ b/api/hip/hip_texture.cpp @@ -193,6 +193,27 @@ void getChannelOrderAndType(const hipChannelFormatDesc& desc, enum hipTextureRea } } +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; + } +} + amd::Sampler* fillSamplerDescriptor(enum hipTextureAddressMode addressMode, enum hipTextureFilterMode filterMode, int normalizedCoords) { #ifndef CL_FILTER_NONE @@ -401,14 +422,33 @@ hipError_t ihipBindTexture(cl_mem_object_type type, } if (hip::getCurrentContext()) { 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); - amd::Image* image = new (*hip::getCurrentContext()) amd::Image(*memory->asBuffer(), - type, memory->getMemFlags(), imageFormat, width, height, 1, pitch, 0); + 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::getCurrentContext()) amd::Image(*memory->asBuffer(), + type, memory->getMemFlags(), imageFormat, width, height, depth, rowPitch, slicePitch); if (!image->create()) { delete image; return hipErrorMemoryAllocation; @@ -437,6 +477,19 @@ hipError_t ihipBindTexture(cl_mem_object_type type, 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; @@ -444,7 +497,10 @@ hipError_t ihipBindTexture(cl_mem_object_type type, } 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; } return hipErrorInvalidValue; @@ -508,6 +564,9 @@ hipError_t ihipBindTextureToArrayImpl(TlsData* tls, int dim, enum hipTextureRead case 2: clType = CL_MEM_OBJECT_IMAGE2D; break; + case 3: + clType = CL_MEM_OBJECT_IMAGE3D; + break; default: HIP_RETURN(hipErrorInvalidValue); }