From ef514eef71af16a05be833caa3090f41e58a084e Mon Sep 17 00:00:00 2001 From: vsytch Date: Wed, 5 Feb 2020 10:26:18 -0500 Subject: [PATCH] Device texture functions should not normalize the sampled pixel (#1826) * Device texture functions should not normalize the sampled pixel. This is already done by HW. * Add support to use h/w capability for normalized float data convertion for driver API's Co-authored-by: ansurya <50609411+ansurya@users.noreply.github.com> --- include/hip/hcc_detail/hip_texture_types.h | 2 ++ include/hip/hcc_detail/texture_functions.h | 24 +++---------- include/hip/hcc_detail/texture_types.h | 1 + src/hip_texture.cpp | 35 ++++++++++++------- .../texture/hipNormalizedFloatValueTex.cpp | 22 ++++++++++-- 5 files changed, 49 insertions(+), 35 deletions(-) diff --git a/include/hip/hcc_detail/hip_texture_types.h b/include/hip/hcc_detail/hip_texture_types.h index b229f4e696..fcd6d69dbe 100644 --- a/include/hip/hcc_detail/hip_texture_types.h +++ b/include/hip/hcc_detail/hip_texture_types.h @@ -57,6 +57,7 @@ struct __HIP_TEXTURE_ATTRIB texture : public textureReference { texture(int norm = 0, enum hipTextureFilterMode fMode = hipFilterModePoint, enum hipTextureAddressMode aMode = hipAddressModeClamp) { normalized = norm; + readMode = hipReadModeNormalizedFloat; filterMode = fMode; addressMode[0] = aMode; addressMode[1] = aMode; @@ -68,6 +69,7 @@ struct __HIP_TEXTURE_ATTRIB texture : public textureReference { texture(int norm, enum hipTextureFilterMode fMode, enum hipTextureAddressMode aMode, struct hipChannelFormatDesc desc) { normalized = norm; + readMode = hipReadModeNormalizedFloat; filterMode = fMode; addressMode[0] = aMode; addressMode[1] = aMode; diff --git a/include/hip/hcc_detail/texture_functions.h b/include/hip/hcc_detail/texture_functions.h index 8cb1088e0e..4a845079fb 100644 --- a/include/hip/hcc_detail/texture_functions.h +++ b/include/hip/hcc_detail/texture_functions.h @@ -157,32 +157,16 @@ union TData { #define TEXTURE_RETURN_UINT_XYZW return make_uint4(texel.u.x, texel.u.y, texel.u.z, texel.u.w); -#define TEXTURE_RETURN_FLOAT return (texFormatToSize[texRef.format] == 1)? texel.f.x : (float)texel.u.x/texFormatToSize[texRef.format]; +#define TEXTURE_RETURN_FLOAT return texel.f.x; -#define TEXTURE_RETURN_FLOAT_X return (texFormatToSize[texRef.format] == 1)? make_float1(texel.f.x) : make_float1((float)texel.u.x/texFormatToSize[texRef.format]); +#define TEXTURE_RETURN_FLOAT_X return make_float1(texel.f.x); -#define TEXTURE_RETURN_FLOAT_XY return (texFormatToSize[texRef.format] == 1)? make_float2(texel.f.x, texel.f.y) : make_float2((float)texel.u.x/texFormatToSize[texRef.format], (float)texel.u.y/texFormatToSize[texRef.format]); +#define TEXTURE_RETURN_FLOAT_XY return make_float2(texel.f.x, texel.f.y); -#define TEXTURE_RETURN_FLOAT_XYZW return (texFormatToSize[texRef.format] == 1)? make_float4(texel.f.x, texel.f.y, texel.f.z, texel.f.w) : make_float4((float)texel.u.x/texFormatToSize[texRef.format], (float)texel.u.y/texFormatToSize[texRef.format], (float)texel.u.z/texFormatToSize[texRef.format], (float)texel.u.w/texFormatToSize[texRef.format]) ; +#define TEXTURE_RETURN_FLOAT_XYZW return make_float4(texel.f.x, texel.f.y, texel.f.z, texel.f.w); extern "C" { -// this is really a sparse array with only valid values being the ones indexed by the enum hipArray_Format(e.g. texFormatToSize[HIP_AD_FORMAT_UNSIGNED_INT8] = UCHAR_MAX) -__device__ __constant__ static int texFormatToSize[] = { - 1, /* HIP_AD_FORMAT_NOT_INITIALIZED */ - UCHAR_MAX, /* HIP_AD_FORMAT_UNSIGNED_INT8 */ - USHRT_MAX, /* HIP_AD_FORMAT_UNSIGNED_INT16 */ - 1, /* HIP_AD_FORMAT_UNSIGNED_INT32 */ - 1,1,1,1, /* Invalid values */ - SCHAR_MAX, /* HIP_AD_FORMAT_SIGNED_INT8 */ - SHRT_MAX, /* HIP_AD_FORMAT_SIGNED_INT16 */ - 1, /* HIP_AD_FORMAT_SIGNED_INT32 */ - 1,1,1,1,1, /* Invalid values */ - 1, /* HIP_AD_FORMAT_HALF */ - 1,1,1,1,1,1,1,1,1,1,1,1,1,1,1, /* Invalid values */ - 1 /* HIP_AD_FORMAT_FLOAT */ -}; - __device__ __hip_float4_vector_value_type __ockl_image_sample_1D( unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, diff --git a/include/hip/hcc_detail/texture_types.h b/include/hip/hcc_detail/texture_types.h index ec3c78532c..832b9095bd 100644 --- a/include/hip/hcc_detail/texture_types.h +++ b/include/hip/hcc_detail/texture_types.h @@ -73,6 +73,7 @@ enum hipTextureReadMode { hipReadModeElementType = 0, hipReadModeNormalizedFloat */ typedef struct textureReference { int normalized; + enum hipTextureReadMode readMode;// used only for driver API's enum hipTextureFilterMode filterMode; enum hipTextureAddressMode addressMode[3]; // Texture address mode for up to 3 dimensions struct hipChannelFormatDesc channelDesc; diff --git a/src/hip_texture.cpp b/src/hip_texture.cpp index 81eb5ad272..27cf321fbc 100644 --- a/src/hip_texture.cpp +++ b/src/hip_texture.cpp @@ -29,24 +29,32 @@ void saveTextureInfo(const hipTexture* pTexture, const hipResourceDesc* pResDesc } } -void getDrvChannelOrderAndType(const enum hipArray_Format Format, unsigned int NumChannels, +void getDrvChannelOrderAndType(const enum hipArray_Format Format, enum hipTextureReadMode readMode, unsigned int NumChannels, hsa_ext_image_channel_order_t* channelOrder, hsa_ext_image_channel_type_t* channelType) { switch (Format) { case HIP_AD_FORMAT_UNSIGNED_INT8: - *channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8; + *channelType = readMode == hipReadModeNormalizedFloat + ? HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT8 + : HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8; break; case HIP_AD_FORMAT_UNSIGNED_INT16: - *channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16; + *channelType = readMode == hipReadModeNormalizedFloat + ? HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT16 + : HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16; break; case HIP_AD_FORMAT_UNSIGNED_INT32: *channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32; break; case HIP_AD_FORMAT_SIGNED_INT8: - *channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8; + *channelType = readMode == hipReadModeNormalizedFloat + ? HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT8 + : HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8; break; case HIP_AD_FORMAT_SIGNED_INT16: - *channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16; + *channelType = readMode == hipReadModeNormalizedFloat + ? HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT16 + : HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16; break; case HIP_AD_FORMAT_SIGNED_INT32: *channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32; @@ -422,7 +430,7 @@ hipError_t ihipBindTextureImpl(TlsData *tls_, int dim, enum hipTextureReadMode r hsa_ext_image_channel_order_t channelOrder; hsa_ext_image_channel_type_t channelType; if (NULL == desc) { - getDrvChannelOrderAndType(tex->format, tex->numChannels, &channelOrder, &channelType); + getDrvChannelOrderAndType(tex->format, readMode, tex->numChannels, &channelOrder, &channelType); } else { getChannelOrderAndType(*desc, readMode, &channelOrder, &channelType); } @@ -497,7 +505,7 @@ hipError_t ihipBindTexture2DImpl(TlsData *tls, int dim, enum hipTextureReadMode hsa_ext_image_channel_type_t channelType; if (NULL == desc) { - getDrvChannelOrderAndType(tex->format, tex->numChannels, &channelOrder, &channelType); + getDrvChannelOrderAndType(tex->format, readMode, tex->numChannels, &channelOrder, &channelType); } else { getChannelOrderAndType(*desc, readMode, &channelOrder, &channelType); } @@ -602,7 +610,7 @@ hipError_t ihipBindTextureToArrayImpl(TlsData *tls_, int dim, enum hipTextureRea hsa_ext_image_channel_order_t channelOrder; hsa_ext_image_channel_type_t channelType; if (array->isDrv) { - getDrvChannelOrderAndType(array->Format, array->NumChannels, + getDrvChannelOrderAndType(array->Format, readMode, array->NumChannels, &channelOrder, &channelType); } else { getChannelOrderAndType(desc, readMode, &channelOrder, &channelType); @@ -724,7 +732,10 @@ hipError_t hipTexRefSetFormat(textureReference* tex, hipArray_Format fmt, int Nu hipError_t hipTexRefSetFlags(textureReference* tex, unsigned int flags) { HIP_INIT_API(hipTexRefSetFlags, tex, flags); hipError_t hip_status = hipSuccess; - tex->normalized = flags; + if(flags == HIP_TRSF_READ_AS_INTEGER) + tex->readMode = hipReadModeElementType; + else if(flags == HIP_TRSF_NORMALIZED_COORDINATES) + tex->normalized = flags; return ihipLogStatus(hip_status); } @@ -757,7 +768,7 @@ hipError_t hipTexRefSetArray(textureReference* tex, hipArray_const_t array, unsi HIP_INIT_API(hipTexRefSetArray, tex, array, flags); hipError_t hip_status = hipSuccess; - hip_status = ihipBindTextureToArrayImpl(tls, array->textureType, hipReadModeElementType, array, + hip_status = ihipBindTextureToArrayImpl(tls, array->textureType, tex->readMode, array, array->desc, tex); return ihipLogStatus(hip_status); } @@ -785,7 +796,7 @@ hipError_t hipTexRefSetAddress(size_t* offset, textureReference* tex, hipDevicep HIP_INIT_API(hipTexRefSetAddress, offset, tex, devPtr, size); hipError_t hip_status = hipSuccess; // TODO: hipReadModeElementType is default. - hip_status = ihipBindTextureImpl(tls, hipTextureType1D, hipReadModeElementType, offset, devPtr, NULL, + hip_status = ihipBindTextureImpl(tls, hipTextureType1D, tex->readMode, offset, devPtr, NULL, size, tex); return ihipLogStatus(hip_status); } @@ -816,7 +827,7 @@ hipError_t hipTexRefSetAddress2D(textureReference* tex, const HIP_ARRAY_DESCRIPT //TODO: Fix when HSA accepts user defined pitch if(pitch % 64) pitch =0; - hip_status = ihipBindTexture2DImpl(tls, hipTextureType2D, hipReadModeElementType, &offset, devPtr, + hip_status = ihipBindTexture2DImpl(tls, hipTextureType2D, tex->readMode, &offset, devPtr, NULL, desc->Width, desc->Height, tex, pitch); return ihipLogStatus(hip_status); } diff --git a/tests/src/texture/hipNormalizedFloatValueTex.cpp b/tests/src/texture/hipNormalizedFloatValueTex.cpp index 0a9c879376..3179f7412e 100644 --- a/tests/src/texture/hipNormalizedFloatValueTex.cpp +++ b/tests/src/texture/hipNormalizedFloatValueTex.cpp @@ -27,8 +27,24 @@ THE SOFTWARE. */ #include "test_common.h" - #define SIZE 10 + +static float getNormalizedValue(const float value, + const enum hipArray_Format texFormat) { + switch (texFormat) { + case HIP_AD_FORMAT_SIGNED_INT8: + return (value / SCHAR_MAX); + case HIP_AD_FORMAT_UNSIGNED_INT8: + return (value / UCHAR_MAX); + case HIP_AD_FORMAT_SIGNED_INT16: + return (value / SHRT_MAX); + case HIP_AD_FORMAT_UNSIGNED_INT16: + return (value / USHRT_MAX); + default: + return value; + } +} + texture textureNormalizedVal_1D; __global__ void normalizedValTextureTest(unsigned int numElements, float* pDst) @@ -47,7 +63,6 @@ bool textureTest(enum hipArray_Format texFormat) T *dData = NULL; HIPCHECK(hipMalloc((void **) &dData, sizeof(T)*SIZE)); HIPCHECK(hipMemcpyHtoD((hipDeviceptr_t)dData, hData, sizeof(T)*SIZE)); - textureReference* texRef = &textureNormalizedVal_1D; HIPCHECK(hipTexRefSetAddressMode(texRef, 0, hipAddressModeClamp)); HIPCHECK(hipTexRefSetAddressMode(texRef, 1, hipAddressModeClamp)); @@ -73,7 +88,8 @@ bool textureTest(enum hipArray_Format texFormat) for(int i = 0; i < SIZE; i++) { - if((float)hData[i]/texFormatToSize[texFormat] != hOutputData[i]) + float expected = getNormalizedValue(float(hData[i]), texFormat); + if(expected != hOutputData[i]) { printf("mismatch at index:%d for texType:%d output:%f\n",i,texFormat,hOutputData[i]); testResult = false;