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

@@ -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;

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,

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;

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);
 }

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<float, hipTextureType1D, hipReadModeElementType> 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;