SWDEV-349490 - Add unit test of sRGB (#2940)

Add unit test of sRGB.
Update helper functions to accomodate uchar4 and float4.
Change-Id: Ib604cf603ab63b908130ef1c562c53d42328bacf

catch/include/hip_texture_helper.hh

@@ -1,4 +1,5 @@
 #pragma once
+#include <math.h>
 
 #define HIP_SAMPLING_VERIFY_EPSILON     0.00001
 // The internal precision varies by the GPU family and sometimes within the family.
@@ -6,15 +7,149 @@
 #define HIP_SAMPLING_VERIFY_RELATIVE_THRESHOLD  0.05  // 5% for filter mode
 #define HIP_SAMPLING_VERIFY_ABSOLUTE_THRESHOLD  0.1
 
-template<typename type, hipTextureFilterMode fMode = hipFilterModePoint>
-bool hipTextureSamplingVerify(const type outputData, const type expected) {
+#if HT_NVIDIA
+template<typename T>
+typename std::enable_if<sizeof(T) / sizeof(decltype(T::x)) == 4, T>::type
+inline __host__ __device__ operator+(const T &a, const T &b)
+{
+  return {a.x + b.x, a.y + b.y, a.z + b.z,  a.w + b.w};
+}
+
+template<typename T>
+typename std::enable_if<sizeof(T) / sizeof(decltype(T::x)) == 4, T>::type
+inline __host__ __device__ operator-(const T &a, const T &b)
+{
+  return {a.x - b.x, a.y - b.y, a.z - b.z,  a.w - b.w};
+}
+
+template<typename T>
+typename std::enable_if<sizeof(T) / sizeof(decltype(T::x)) == 4, bool>::type
+inline __host__ __device__ operator==(const T &a, const T &b)
+{
+  return a.x == b.x && a.y == b.y && a.z == b.z && a.w == b.w;
+}
+
+template<typename T>
+typename std::enable_if<sizeof(T) / sizeof(decltype(T::x)) == 4, T>::type
+inline __host__ __device__ operator*(const decltype(T::x) &a, const T &b)
+{
+  return {a * b.x, a * b.y, a * b.z, a * b.w};
+}
+
+template<typename T>
+typename std::enable_if<sizeof(T) / sizeof(decltype(T::x)) == 4, void>::type
+inline __host__ __device__ operator*=(T &a, const decltype(T::x) &b)
+{
+  a.x *= b;
+  a.y *= b;
+  a.z *= b;
+  a.w *= b;
+}
+#endif // HT_NVIDIA
+
+// See https://en.wikipedia.org/wiki/SRGB#Transformation
+// From CIE 1931 color space to sRGB
+inline float hipSRGBMap(float fc) {
+  double c = static_cast<double>(fc);
+
+#if !defined(_WIN32)
+  if (std::isnan(c))
+    c = 0.0;
+#else
+    if (_isnan(c)) c = 0.0;
+#endif
+
+  if (c > 1.0)
+    c = 1.0;
+  else if (c < 0.0)
+    c = 0.0;
+  else if (c < 0.0031308)
+    c = 12.92 * c;
+  else
+    c = 1.055 * pow(c, 5.0 / 12.0) - 0.055;
+
+  return static_cast<float>(c);
+}
+
+// From sRGB to CIE 1931 color space
+inline float hipSRGBUnmap(float fc) {
+  double c = static_cast<double>(fc);
+
+  if (c <= 0.04045)
+    c = c / 12.92;
+  else
+    c = pow((c + 0.055) / 1.055, 2.4);
+
+  return static_cast<float>(c);
+}
+
+inline float4 hipSRGBMap(float4 fc) {
+  fc.x = hipSRGBMap(fc.x);
+  fc.y = hipSRGBMap(fc.y);
+  fc.z = hipSRGBMap(fc.z);
+  // Alpha channel will keep unchanged
+  return fc;
+}
+
+inline float4 hipSRGBUnmap(float4 fc) {
+  fc.x = hipSRGBUnmap(fc.x);
+  fc.y = hipSRGBUnmap(fc.y);
+  fc.z = hipSRGBUnmap(fc.z);
+  // Alpha channel will keep unchanged
+  return fc;
+}
+
+template<typename T>
+typename std::enable_if<std::is_scalar<T>::value == true, double>::type
+hipFabs(const T &t) {
+  return fabs(t);
+}
+
+template<typename T>
+typename std::enable_if<sizeof(T) / sizeof(decltype(T::x)) == 1, double>::type
+hipFabs(const T &t) {
+  return fabs(t.x);
+}
+
+template<typename T>
+typename std::enable_if<sizeof(T) / sizeof(decltype(T::x)) == 2, double>::type
+hipFabs(const T &t) {
+  double x = static_cast<double>(t.x);
+  double y = static_cast<double>(t.y);
+  double s =  x * x +  y * y;
+  return sqrt(s);
+}
+
+template<typename T>
+typename std::enable_if<sizeof(T) / sizeof(decltype(T::x)) == 3, double>::type
+hipFabs(const T &t) {
+  double x = static_cast<double>(t.x);
+  double y = static_cast<double>(t.y);
+  double z = static_cast<double>(t.z);
+  double s =  x * x +  y * y + z * z;
+  return sqrt(s);
+}
+
+template<typename T>
+typename std::enable_if<sizeof(T) / sizeof(decltype(T::x)) == 4, double>::type
+hipFabs(const T &t) {
+  double x = static_cast<double>(t.x);
+  double y = static_cast<double>(t.y);
+  double z = static_cast<double>(t.z);
+  double w = static_cast<double>(t.w);
+  double s =  x * x +  y * y + z * z + w * w;
+  return sqrt(s);
+}
+
+template<typename T, hipTextureFilterMode fMode = hipFilterModePoint, bool sRGB = false>
+bool hipTextureSamplingVerify(T outputData, T expected) {
   bool testResult = false;
-  if (fMode == hipFilterModePoint) {
+  if (fMode == hipFilterModePoint && !sRGB) {
     testResult = outputData == expected;
-  } else if (fMode == hipFilterModeLinear) {
-    const type mean = (fabs(outputData) + fabs(expected)) / 2;
-    const type diff = fabs(outputData - expected);
-    const type ratio = diff / (mean + HIP_SAMPLING_VERIFY_EPSILON);
+  } else {
+    double mean = (hipFabs(outputData) + hipFabs(expected)) / 2;
+    double diff = hipFabs(outputData - expected);
+    double ratio = diff / (mean + HIP_SAMPLING_VERIFY_EPSILON);
     if (ratio <= HIP_SAMPLING_VERIFY_RELATIVE_THRESHOLD) {
       testResult = true;
     } else if (diff <= HIP_SAMPLING_VERIFY_ABSOLUTE_THRESHOLD) {
@@ -46,10 +181,11 @@ void hipTextureGetAddress(int &value, const int maxValue)
 
 // Simulate logics in CTS read_image_pixel_float().
 // x, y and z must be returned by hipTextureGetAddress()
-template<hipTextureAddressMode addressMode>
-float hipTextureGetValue(const float *data, const int x, const int width,
-         const int y = 0, const int height = 0,const int z = 0, const int depth = 0) {
-  float result = std::numeric_limits<float>::lowest();
+template<typename T, hipTextureAddressMode addressMode, bool sRGB = false>
+T hipTextureGetValue(const T *data, const int x, const int width,
+         const int y = 0, const int height = 0, const int z = 0, const int depth = 0) {
+  T result;
+  memset(&result, 0, sizeof(result));
   switch (addressMode) {
     case hipAddressModeClamp:
       if (width > 0) {
@@ -65,30 +201,35 @@ float hipTextureGetValue(const float *data, const int x, const int width,
     case hipAddressModeBorder:
       if (width > 0) {
         if (height == 0 && depth == 0) {
-          result = (x >= 0 && x < width) ? data[x] : 0;  // 1D
+          if (x >= 0 && x < width)
+            result = data[x];  // 1D
         } else if (depth == 0) {
-          result = (x >= 0 && x < width && y >= 0 && y < height) ?
-              data[y * width + x] : 0;  // 2D
+          if (x >= 0 && x < width && y >= 0 && y < height)
+              result = data[y * width + x];  // 2D
         } else {
-          result = (x >= 0 && x < width && y >= 0 && y < height && z >= 0 && z < depth) ?
-              data[z * width * height + y * width + x] : 0;  // 3D
+          if (x >= 0 && x < width && y >= 0 && y < height && z >= 0 && z < depth)
+              result = data[z * width * height + y * width + x];  // 3D
         }
       }
       break;
     default:
       break;
   }
+  if (sRGB && std::is_same<T, float4>::value) {
+    result = hipSRGBUnmap(result);
+  }
   return result;
 }
 
-template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode>
-float getExpectedValue(const int width, float x, const float *data) {
-  float result = std::numeric_limits<float>::lowest();
+template<typename T, hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool sRGB = false>
+T getExpectedValue(const int width, float x, const T *data) {
+  T result;
+  memset(&result, 0, sizeof(result));
   switch (filterMode) {
     case hipFilterModePoint: {
       int i1 = static_cast<int>(floor(x));
       hipTextureGetAddress < addressMode > (i1, width);
-      result = hipTextureGetValue < addressMode > (data, i1, width);
+      result = hipTextureGetValue < T, addressMode, sRGB > (data, i1, width);
     }
       break;
     case hipFilterModeLinear: {
@@ -99,8 +240,8 @@ float getExpectedValue(const int width, float x, const float *data) {
       hipTextureGetAddress < addressMode > (i1, width);
       hipTextureGetAddress < addressMode > (i2, width);
 
-      float t1 = hipTextureGetValue < addressMode > (data, i1, width);
-      float t2 = hipTextureGetValue < addressMode > (data, i2, width);
+      T t1 = hipTextureGetValue < T, addressMode, sRGB> (data, i1, width);
+      T t2 = hipTextureGetValue < T, addressMode, sRGB > (data, i2, width);
 
       return (1 - a) * t1 + a * t2;
     }
@@ -109,16 +250,17 @@ float getExpectedValue(const int width, float x, const float *data) {
   return result;
 }
 
-template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode>
-float getExpectedValue(const int width, const int height, float x, float y, const float *data) {
-  float result = std::numeric_limits<float>::lowest();
+template<typename T, hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool sRGB = false>
+T getExpectedValue(const int width, const int height, float x, float y, const T *data) {
+  T result;
+  memset(&result, 0, sizeof(result));
   switch (filterMode) {
     case hipFilterModePoint: {
       int i1 = static_cast<int>(floor(x));
       int j1 = static_cast<int>(floor(y));
       hipTextureGetAddress < addressMode > (i1, width);
       hipTextureGetAddress < addressMode > (j1, height);
-      result = hipTextureGetValue < addressMode > (data, i1, width, j1, height);
+      result = hipTextureGetValue < T, addressMode, sRGB > (data, i1, width, j1, height);
     }
       break;
     case hipFilterModeLinear: {
@@ -139,13 +281,13 @@ float getExpectedValue(const int width, const int height, float x, float y, cons
       hipTextureGetAddress < addressMode > (j1, height);
       hipTextureGetAddress < addressMode > (j2, height);
 
-      float t11 = hipTextureGetValue < addressMode
+      T t11 = hipTextureGetValue < T, addressMode, sRGB
           > (data, i1, width, j1, height);
-      float t21 = hipTextureGetValue < addressMode
+      T t21 = hipTextureGetValue < T, addressMode, sRGB
           > (data, i2, width, j1, height);
-      float t12 = hipTextureGetValue < addressMode
+      T t12 = hipTextureGetValue < T, addressMode, sRGB
           > (data, i1, width, j2, height);
-      float t22 = hipTextureGetValue < addressMode
+      T t22 = hipTextureGetValue < T, addressMode, sRGB
           > (data, i2, width, j2, height);
 
       result = (1 - a) * (1 - b) * t11 + a * (1 - b) * t21 + (1 - a) * b * t12
@@ -156,10 +298,11 @@ float getExpectedValue(const int width, const int height, float x, float y, cons
   return result;
 }
 
-template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode>
-float getExpectedValue(const int width, const int height, const int depth,
-                      float x, float y, float z, const float *data) {
-  float result = std::numeric_limits<float>::lowest();
+template<class T, hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool sRGB = false>
+T getExpectedValue(const int width, const int height, const int depth,
+                      float x, float y, float z, const T *data) {
+  T result;
+  memset(&result, 0, sizeof(result));
   switch (filterMode) {
     case hipFilterModePoint: {
       int i1 = static_cast<int>(floor(x));
@@ -170,7 +313,7 @@ float getExpectedValue(const int width, const int height, const int depth,
       hipTextureGetAddress < addressMode > (j1, height);
       hipTextureGetAddress < addressMode > (k1, depth);
 
-      result = hipTextureGetValue < addressMode > (data, i1, width, j1, height, k1, depth);
+      result = hipTextureGetValue < T, addressMode, sRGB > (data, i1, width, j1, height, k1, depth);
     }
       break;
     case hipFilterModeLinear: {
@@ -197,21 +340,21 @@ float getExpectedValue(const int width, const int height, const int depth,
       hipTextureGetAddress < addressMode > (k1, depth);
       hipTextureGetAddress < addressMode > (k2, depth);
 
-      float t111 = hipTextureGetValue < addressMode
+      T t111 = hipTextureGetValue < T, addressMode, sRGB
           > (data, i1, width, j1, height, k1, depth);
-      float t211 = hipTextureGetValue < addressMode
+      T t211 = hipTextureGetValue < T, addressMode, sRGB
           > (data, i2, width, j1, height, k1, depth);
-      float t121 = hipTextureGetValue < addressMode
+      T t121 = hipTextureGetValue < T, addressMode, sRGB
           > (data, i1, width, j2, height, k1, depth);
-      float t112 = hipTextureGetValue < addressMode
+      T t112 = hipTextureGetValue < T, addressMode, sRGB
           > (data, i1, width, j1, height, k2, depth);
-      float t122 = hipTextureGetValue < addressMode
+      T t122 = hipTextureGetValue < T, addressMode, sRGB
           > (data, i1, width, j2, height, k2, depth);
-      float t212 = hipTextureGetValue < addressMode
+      T t212 = hipTextureGetValue < T, addressMode, sRGB
           > (data, i2, width, j1, height, k2, depth);
-      float t221 = hipTextureGetValue < addressMode
+      T t221 = hipTextureGetValue < T, addressMode, sRGB
           > (data, i2, width, j2, height, k1, depth);
-      float t222 = hipTextureGetValue < addressMode
+      T t222 = hipTextureGetValue < T, addressMode, sRGB
           > (data, i2, width, j2, height, k2, depth);
 
       result =
@@ -224,4 +367,4 @@ float getExpectedValue(const int width, const int height, const int depth,
       break;
   }
   return result;
-}
+}

catch/unit/texture/CMakeLists.txt

@@ -43,6 +43,7 @@ set(TEST_SRC
     hipTextureObj1DCheckModes.cc
     hipTextureObj2DCheckModes.cc
     hipTextureObj3DCheckModes.cc
+    hipTextureObj2DCheckSRGBModes.cc
 )
 
 hip_add_exe_to_target(NAME TextureTest

catch/unit/texture/hipTextureObj1DCheckModes.cc

@@ -31,7 +31,7 @@ __global__ void tex1DKernel(float *outputData, hipTextureObject_t textureObject,
 }
 
 template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool normalizedCoords>
-void runTest(const int width, const float offsetX) {
+static void runTest(const int width, const float offsetX) {
   //printf("%s(addressMode=%d, filterMode=%d, normalizedCoords=%d, width=%d, offsetX=%f)\n", __FUNCTION__,
   //       addressMode, filterMode, normalizedCoords, width, offsetX);
   unsigned int size = width * sizeof(float);
@@ -82,7 +82,7 @@ void runTest(const int width, const float offsetX) {
 
   bool result = true;
   for (int j = 0; j < width; j++) {
-    float expectedValue = getExpectedValue<addressMode, filterMode>(width, offsetX + j, hData);
+    float expectedValue = getExpectedValue<float, addressMode, filterMode>(width, offsetX + j, hData);
     if (!hipTextureSamplingVerify<float, filterMode>(hOutputData[j], expectedValue)) {
       INFO("Mismatch at " << offsetX + j << ":" << hOutputData[j] <<
            " expected:" << expectedValue);

catch/unit/texture/hipTextureObj2DCheckModes.cc

@@ -35,7 +35,7 @@ __global__ void tex2DKernel(float *outputData, hipTextureObject_t textureObject,
 }
 
 template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool normalizedCoords>
-void runTest(const int width, const int height, const float offsetX, const float offsetY) {
+static void runTest(const int width, const int height, const float offsetX, const float offsetY) {
   //printf("%s(addressMode=%d, filterMode=%d, normalizedCoords=%d, width=%d, height=%d, offsetX=%f, offsetY=%f)\n",
   //     __FUNCTION__, addressMode, filterMode, normalizedCoords, width, height, offsetX, offsetY);
   unsigned int size = width * height * sizeof(float);
@@ -92,7 +92,7 @@ void runTest(const int width, const int height, const float offsetX, const float
   for (int i = 0; i < height; i++) {
     for (int j = 0; j < width; j++) {
       int index = i * width + j;
-      float expectedValue = getExpectedValue<addressMode, filterMode>(width, height,
+      float expectedValue = getExpectedValue<float, addressMode, filterMode>(width, height,
                                                     offsetX + j, offsetY + i, hData);
       if (!hipTextureSamplingVerify<float, filterMode>(hOutputData[index], expectedValue)) {
         INFO("Mismatch at (" << offsetX + j << ", " << offsetY + i << "):" <<
@@ -128,7 +128,7 @@ TEST_CASE("Unit_hipTextureObj2DCheckModes") {
     runTest<hipAddressModeBorder, hipFilterModePoint, false>(256, 256, 12.5, 6.7);
   }
 
-  SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") {
+  SECTION("hipAddressModeClamp, hipFilterModeLinear, regularCoords") {
     runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, 256, -0.4, -0.4);
     runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, 256, 4, 14.6);
   }

catch/unit/texture/hipTextureObj2DCheckSRGBModes.cc

@@ -0,0 +1,242 @@
+/*
+Copyright (c) 2022 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.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <hip_texture_helper.hh>
+
+template<bool normalizedCoords>
+__global__ void tex2DRGBAKernel(float4 *outputData, hipTextureObject_t textureObject,
+                            int width, int height, float offsetX,
+                            float offsetY) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
+  int x = blockIdx.x * blockDim.x + threadIdx.x;
+  int y = blockIdx.y * blockDim.y + threadIdx.y;
+  outputData[y * width + x] = tex2D<float4>(textureObject,
+                                           normalizedCoords ? (x + offsetX) / width : x + offsetX,
+                                           normalizedCoords ? (y + offsetY) / height : y + offsetY);
+#endif
+}
+
+template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool normalizedCoords, bool sRGB = false>
+static void runTest(const int width, const int height, const float offsetX, const float offsetY) {
+  //printf("%s(addressMode=%d, filterMode=%d, normalizedCoords=%d, width=%d, height=%d, offsetX=%f, offsetY=%f)\n",
+  //     __FUNCTION__, addressMode, filterMode, normalizedCoords, width, height, offsetX, offsetY);
+  constexpr float uCharMax = UCHAR_MAX;
+  unsigned int size = width * height * sizeof(uchar4);
+  uchar4 *hData = (uchar4*) malloc(size);
+  memset(hData, 0, size);
+  for (int i = 0; i < height; i++) {
+    for (int j = 0; j < width; j++) {
+      int index = i * width + j;
+      hData[index].x = static_cast<unsigned char>(j);
+      hData[index].y = static_cast<unsigned char>(i);
+      hData[index].z = static_cast<unsigned char>(index);
+      hData[index].w = static_cast<unsigned char>(i + j);
+    }
+  }
+
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc<uchar4>();
+  hipArray *hipArray;
+  HIP_CHECK(hipMallocArray(&hipArray, &channelDesc, width, height));
+
+  HIP_CHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, width * sizeof(uchar4), width * sizeof(uchar4), height, hipMemcpyHostToDevice));
+
+  hipResourceDesc resDesc;
+  memset(&resDesc, 0, sizeof(resDesc));
+  resDesc.resType = hipResourceTypeArray;
+  resDesc.res.array.array = hipArray;
+
+  // Specify texture object parameters
+  hipTextureDesc texDesc;
+  memset(&texDesc, 0, sizeof(texDesc));
+  texDesc.addressMode[0] = addressMode;
+  texDesc.addressMode[1] = addressMode;
+  texDesc.filterMode = filterMode;
+  texDesc.readMode = hipReadModeNormalizedFloat;
+  texDesc.normalizedCoords = normalizedCoords;
+  texDesc.sRGB = sRGB ? 1 : 0;
+
+  // Create texture object
+  hipTextureObject_t textureObject = 0;
+  HIP_CHECK(hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL));
+
+  float4 *dData = nullptr;
+  size = width * height * sizeof(float4);
+  HIP_CHECK(hipMalloc((void**) &dData, size));
+
+  dim3 dimBlock(16, 16, 1);
+  dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y, 1);
+
+  hipLaunchKernelGGL(tex2DRGBAKernel<normalizedCoords>, dimGrid, dimBlock, 0, 0, dData,
+                     textureObject, width, height, offsetX, offsetY);
+
+  HIP_CHECK(hipDeviceSynchronize());
+
+  float4 *hInputData = (float4*) malloc(size);  // CPU expected values
+  float4 *hOutputData = (float4*) malloc(size); // GPU output values
+  memset(hInputData, 0, size);
+  memset(hOutputData, 0, size);
+
+  for (int i = 0; i < height; i++) {
+    for (int j = 0; j < width; j++) {
+      int index = i * width + j;
+      hInputData[index].x = hData[index].x / uCharMax;
+      hInputData[index].y = hData[index].y / uCharMax;
+      hInputData[index].z = hData[index].z / uCharMax;
+      hInputData[index].w = hData[index].w / uCharMax;
+    }
+  }
+  HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost));
+
+  bool result = true;
+
+  for (int i = 0; i < height; i++) {
+    for (int j = 0; j < width; j++) {
+      int index = i * width + j;
+        //printf("(%d, %d): hInputData=(%f, %f, %f, %f), hOutputData=(%f, %f, %f, %f)\n", i, j,
+        //     hInputData[index].x, hInputData[index].y, hInputData[index].z, hInputData[index].w,
+        //     hOutputData[index].x, hOutputData[index].y, hOutputData[index].z, hOutputData[index].w);
+
+      float4 cpuExpected = getExpectedValue<float4, addressMode, filterMode, sRGB>(width, height,
+                                                    offsetX + j, offsetY + i, hInputData);
+      float4 gpuOutput = hOutputData[index];
+      if (sRGB) {
+        // CTS will map to sRGP before comparison, so we do so
+        cpuExpected = hipSRGBMap(cpuExpected);
+        gpuOutput = hipSRGBMap(gpuOutput);
+      }
+      // Convert from [0, 1] back to [0, 255]
+      gpuOutput *= uCharMax;
+      cpuExpected *= uCharMax;
+      if (!hipTextureSamplingVerify<float4, filterMode, sRGB>(gpuOutput, cpuExpected)) {
+        WARN("Mismatch at (" << offsetX + j << ", " << offsetY + i << ") GPU output : " <<
+             gpuOutput.x << ", " <<
+             gpuOutput.y << ", " <<
+             gpuOutput.z << ", " <<
+             gpuOutput.w << ", " <<
+             " CPU expected: " <<
+             cpuExpected.x << ", " <<
+             cpuExpected.y << ", " <<
+             cpuExpected.z << ", " <<
+             cpuExpected.w << "\n");
+        result = false;
+        goto line1;
+      }
+    }
+  }
+
+line1:
+  HIP_CHECK(hipDestroyTextureObject(textureObject));
+  HIP_CHECK(hipFree(dData));
+  HIP_CHECK(hipFreeArray(hipArray));
+  free(hData);
+  free(hOutputData);
+  free(hInputData);
+  REQUIRE(result);
+}
+
+TEST_CASE("Unit_hipTextureObj2DCheckRGBAModes") {
+  CHECK_IMAGE_SUPPORT
+
+  SECTION("RGBA 2D hipAddressModeClamp, hipFilterModePoint, regularCoords") {
+    runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, -3.9, 6.1);
+    runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, 4.4, -7.0);
+  }
+
+  SECTION("RGBA 2D hipAddressModeBorder, hipFilterModePoint, regularCoords") {
+    runTest<hipAddressModeBorder, hipFilterModePoint, false>(256, 256, -8.5, 2.9);
+    runTest<hipAddressModeBorder, hipFilterModePoint, false>(256, 256, 12.5, 6.7);
+  }
+
+  SECTION("RGBA 2D hipAddressModeClamp, hipFilterModeLinear, regularCoords") {
+    runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, 256, -0.4, -0.4);
+    runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, 256, 4, 14.6);
+  }
+
+  SECTION("RGBA 2D hipAddressModeBorder, hipFilterModeLinear, regularCoords") {
+    runTest<hipAddressModeBorder, hipFilterModeLinear, false>(256, 256, -0.4, 0.4);
+    runTest<hipAddressModeBorder, hipFilterModeLinear, false>(256, 256, 12.5, 23.7);
+  }
+
+  SECTION("RGBA 2D hipAddressModeClamp, hipFilterModePoint, normalizedCoords") {
+    runTest<hipAddressModeClamp, hipFilterModePoint, true>(256, 256, -3, 8.9);
+    runTest<hipAddressModeClamp, hipFilterModePoint, true>(256, 256, 4, -0.1);
+  }
+
+  SECTION("RGBA 2D hipAddressModeBorder, hipFilterModePoint, normalizedCoords") {
+    runTest<hipAddressModeBorder, hipFilterModePoint, true>(256, 256, -8.5, 15.9);
+    runTest<hipAddressModeBorder, hipFilterModePoint, true>(256, 256, 12.5, -17.9);
+  }
+
+  SECTION("RGBA 2D hipAddressModeClamp, hipFilterModeLinear, normalizedCoords") {
+    runTest<hipAddressModeClamp, hipFilterModeLinear, true>(256, 256, -3, 5.8);
+    runTest<hipAddressModeClamp, hipFilterModeLinear, true>(256, 256, 4, 9.1);
+  }
+
+  SECTION("RGBA 2D hipAddressModeBorder, hipFilterModeLinear, normalizedCoords") {
+    runTest<hipAddressModeBorder, hipFilterModeLinear, true>(256, 256, -8.5, 6.6);
+    runTest<hipAddressModeBorder, hipFilterModeLinear, true>(256, 256, 12.5, 0.01);
+  }
+}
+
+
+TEST_CASE("Unit_hipTextureObj2DCheckSRGBAModes") {
+  CHECK_IMAGE_SUPPORT
+
+  SECTION("SRGBA 2D hipAddressModeClamp, hipFilterModePoint, regularCoords") {
+    runTest<hipAddressModeClamp, hipFilterModePoint, false, true>(256, 256, -3.9, 6.1);
+    runTest<hipAddressModeClamp, hipFilterModePoint, false, true>(256, 256, 4.4, -7.0);
+  }
+
+  SECTION("SRGBA 2D hipAddressModeBorder, hipFilterModePoint, regularCoords") {
+    runTest<hipAddressModeBorder, hipFilterModePoint, false, true>(256, 256, -8.5, 2.9);
+    runTest<hipAddressModeBorder, hipFilterModePoint, false, true>(256, 256, 12.5, 6.7);
+  }
+
+  SECTION("SRGBA 2D hipAddressModeClamp, hipFilterModeLinear, regularCoords") {
+    runTest<hipAddressModeClamp, hipFilterModeLinear, false, true>(256, 256, -0.4, -0.4);
+    runTest<hipAddressModeClamp, hipFilterModeLinear, false, true>(256, 256, 4, 14.6);
+  }
+
+  SECTION("SRGBA 2D hipAddressModeBorder, hipFilterModeLinear, regularCoords") {
+    runTest<hipAddressModeBorder, hipFilterModeLinear, false, true>(256, 256, -0.4, 0.4);
+    runTest<hipAddressModeBorder, hipFilterModeLinear, false, true>(256, 256, 12.5, 23.7);
+  }
+
+  SECTION("SRGBA 2D hipAddressModeClamp, hipFilterModePoint, normalizedCoords") {
+    runTest<hipAddressModeClamp, hipFilterModePoint, true, true>(256, 256, -3, 8.9);
+    runTest<hipAddressModeClamp, hipFilterModePoint, true, true>(256, 256, 4, -0.1);
+  }
+
+  SECTION("SRGBA 2D hipAddressModeBorder, hipFilterModePoint, normalizedCoords") {
+    runTest<hipAddressModeBorder, hipFilterModePoint, true, true>(256, 256, -8.5, 15.9);
+    runTest<hipAddressModeBorder, hipFilterModePoint, true, true>(256, 256, 12.5, -17.9);
+  }
+
+  SECTION("SRGBA 2D hipAddressModeClamp, hipFilterModeLinear, normalizedCoords") {
+    runTest<hipAddressModeClamp, hipFilterModeLinear, true, true>(256, 256, -3, 5.8);
+    runTest<hipAddressModeClamp, hipFilterModeLinear, true, true>(256, 256, 4, 9.1);
+  }
+
+  SECTION("SRGBA 2D hipAddressModeBorder, hipFilterModeLinear, normalizedCoords") {
+    runTest<hipAddressModeBorder, hipFilterModeLinear, true, true>(256, 256, -8.5, 6.6);
+    runTest<hipAddressModeBorder, hipFilterModeLinear, true, true>(256, 256, 12.5, 0.01);
+  }
+}

catch/unit/texture/hipTextureObj3DCheckModes.cc

@@ -39,7 +39,7 @@ __global__ void tex3DKernel(float *outputData, hipTextureObject_t textureObject,
 }
 
 template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool normalizedCoords>
-void runTest(const int width, const int height, const int depth, const float offsetX, const float offsetY,
+static void runTest(const int width, const int height, const int depth, const float offsetX, const float offsetY,
              const float offsetZ) {
   //printf("%s(addressMode=%d, filterMode=%d, normalizedCoords=%d, width=%d, height=%d, depth=%d, offsetX=%f, offsetY=%f, offsetZ=%f)\n",
   //    __FUNCTION__, addressMode, filterMode, normalizedCoords, width, height,
@@ -124,7 +124,7 @@ void runTest(const int width, const int height, const int depth, const float off
     for (int j = 0; j < height; j++) {
       for (int k = 0; k < width; k++) {
         int index = i * width * height + j * width + k;
-        float expectedValue = getExpectedValue<addressMode, filterMode>(
+        float expectedValue = getExpectedValue<float, addressMode, filterMode>(
             width, height, depth, offsetX + k, offsetY + j, offsetZ + i, hData);
 
         if (!hipTextureSamplingVerify<float, filterMode>(hOutputData[index], expectedValue)) {