SWDEV-333188 - Port image tests to catch2 (#2793)

- Port image tests to catch2. - Disable 3D filter test for gfx9a. - Some other improvement - Temporarily disable linear filter tests on Windows Change-Id: Ieeddac9f8c50aa9b6b12ca9d8fd792a51973b8ca [ROCm/hip-tests commit: 0fc518d281]
2022-07-19 19:26:56 +05:30
@@ -203,7 +203,7 @@ static inline int RAND_R(unsigned* rand_seed) {

 inline bool isImageSupported() {
  int imageSupport = 1;
-#ifdef __HIP_PLATFORM_AMD__
+#if HT_AMD
  int device;
  HIP_CHECK(hipGetDevice(&device));
  HIPCHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, device));
@@ -276,8 +276,6 @@ void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerB

 // This must be called in the beginning of image test app's main() to indicate whether image
 // is supported.
-#define checkImageSupport()                                                                        \
-  if (!HipTest::isImageSupported()) {                                                              \
-    printf("Texture is not support on the device. Skipped.\n");                                    \
-    return;                                                                                        \
-  }
+#define CHECK_IMAGE_SUPPORT                                                                \
+  if (!HipTest::isImageSupported())                                                      \
+    { INFO("Texture is not support on the device. Skipped."); return; }
@@ -0,0 +1,227 @@
+#pragma once
+
+#define HIP_SAMPLING_VERIFY_EPSILON     0.00001
+// The internal precision varies by the GPU family and sometimes within the family.
+// Thus the following threshold is subject to change.
+#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) {
+  bool testResult = false;
+  if (fMode == hipFilterModePoint) {
+    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);
+    if (ratio <= HIP_SAMPLING_VERIFY_RELATIVE_THRESHOLD) {
+      testResult = true;
+    } else if (diff <= HIP_SAMPLING_VERIFY_ABSOLUTE_THRESHOLD) {
+      // Some small outputs have big ratio due to float operation difference of ALU and GPU
+      testResult = true;
+    }
+  }
+  return testResult;
+}
+
+// Simulate CTS static AddressingTable sAddressingTable
+template<hipTextureAddressMode addressMode>
+void hipTextureGetAddress(int &value, const int maxValue)
+{
+  switch(addressMode)
+  {
+    case hipAddressModeClamp:
+      value = value < 0 ? 0
+                  : (value > maxValue - 1 ? maxValue - 1 : value);
+      break;
+    case hipAddressModeBorder:
+      value = value < -1 ? -1
+                  : (value > maxValue ? maxValue : value);
+      break;
+    default:
+      break;
+  }
+}
+
+// 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();
+  switch (addressMode) {
+    case hipAddressModeClamp:
+      if (width > 0) {
+        if (height == 0 && depth == 0) {
+          result = data[x];  // 1D
+        } else if (depth == 0) {
+          result = data[y * width + x];  // 2D
+        } else {
+          result = data[z * width * height + y * width + x];  // 3D
+        }
+      }
+      break;
+    case hipAddressModeBorder:
+      if (width > 0) {
+        if (height == 0 && depth == 0) {
+          result = (x >= 0 && x < width) ? data[x] : 0;  // 1D
+        } else if (depth == 0) {
+          result = (x >= 0 && x < width && y >= 0 && y < height) ?
+              data[y * width + x] : 0;  // 2D
+        } else {
+          result = (x >= 0 && x < width && y >= 0 && y < height && z >= 0 && z < depth) ?
+              data[z * width * height + y * width + x] : 0;  // 3D
+        }
+      }
+      break;
+    default:
+      break;
+  }
+  return result;
+}
+
+template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode>
+float getExpectedValue(const int width, float x, const float *data) {
+  float result = std::numeric_limits<float>::lowest();
+  switch (filterMode) {
+    case hipFilterModePoint: {
+      int i1 = static_cast<int>(floor(x));
+      hipTextureGetAddress < addressMode > (i1, width);
+      result = hipTextureGetValue < addressMode > (data, i1, width);
+    }
+      break;
+    case hipFilterModeLinear: {
+      x -= 0.5;
+      int i1 = static_cast<int>(floor(x));
+      int i2 = i1 + 1;
+      float a = x - i1;
+      hipTextureGetAddress < addressMode > (i1, width);
+      hipTextureGetAddress < addressMode > (i2, width);
+
+      float t1 = hipTextureGetValue < addressMode > (data, i1, width);
+      float t2 = hipTextureGetValue < addressMode > (data, i2, width);
+
+      return (1 - a) * t1 + a * t2;
+    }
+      break;
+  }
+  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();
+  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);
+    }
+      break;
+    case hipFilterModeLinear: {
+      x -= 0.5;
+      y -= 0.5;
+
+      int i1 = static_cast<int>(floor(x));
+      int j1 = static_cast<int>(floor(y));
+
+      int i2 = i1 + 1;
+      int j2 = j1 + 1;
+
+      float a = x - i1;
+      float b = y - j1;
+
+      hipTextureGetAddress < addressMode > (i1, width);
+      hipTextureGetAddress < addressMode > (i2, width);
+      hipTextureGetAddress < addressMode > (j1, height);
+      hipTextureGetAddress < addressMode > (j2, height);
+
+      float t11 = hipTextureGetValue < addressMode
+          > (data, i1, width, j1, height);
+      float t21 = hipTextureGetValue < addressMode
+          > (data, i2, width, j1, height);
+      float t12 = hipTextureGetValue < addressMode
+          > (data, i1, width, j2, height);
+      float t22 = hipTextureGetValue < addressMode
+          > (data, i2, width, j2, height);
+
+      result = (1 - a) * (1 - b) * t11 + a * (1 - b) * t21 + (1 - a) * b * t12
+          + a * b * t22;
+    }
+      break;
+  }
+  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();
+  switch (filterMode) {
+    case hipFilterModePoint: {
+      int i1 = static_cast<int>(floor(x));
+      int j1 = static_cast<int>(floor(y));
+      int k1 = static_cast<int>(floor(z));
+
+      hipTextureGetAddress < addressMode > (i1, width);
+      hipTextureGetAddress < addressMode > (j1, height);
+      hipTextureGetAddress < addressMode > (k1, depth);
+
+      result = hipTextureGetValue < addressMode > (data, i1, width, j1, height, k1, depth);
+    }
+      break;
+    case hipFilterModeLinear: {
+      x -= 0.5;
+      y -= 0.5;
+      z -= 0.5;
+
+      int i1 = static_cast<int>(floor(x));
+      int j1 = static_cast<int>(floor(y));
+      int k1 = static_cast<int>(floor(z));
+
+      int i2 = i1 + 1;
+      int j2 = j1 + 1;
+      int k2 = k1 + 1;
+
+      float a = x - i1;
+      float b = y - j1;
+      float c = z - k1;
+
+      hipTextureGetAddress < addressMode > (i1, width);
+      hipTextureGetAddress < addressMode > (i2, width);
+      hipTextureGetAddress < addressMode > (j1, height);
+      hipTextureGetAddress < addressMode > (j2, height);
+      hipTextureGetAddress < addressMode > (k1, depth);
+      hipTextureGetAddress < addressMode > (k2, depth);
+
+      float t111 = hipTextureGetValue < addressMode
+          > (data, i1, width, j1, height, k1, depth);
+      float t211 = hipTextureGetValue < addressMode
+          > (data, i2, width, j1, height, k1, depth);
+      float t121 = hipTextureGetValue < addressMode
+          > (data, i1, width, j2, height, k1, depth);
+      float t112 = hipTextureGetValue < addressMode
+          > (data, i1, width, j1, height, k2, depth);
+      float t122 = hipTextureGetValue < addressMode
+          > (data, i1, width, j2, height, k2, depth);
+      float t212 = hipTextureGetValue < addressMode
+          > (data, i2, width, j1, height, k2, depth);
+      float t221 = hipTextureGetValue < addressMode
+          > (data, i2, width, j2, height, k1, depth);
+      float t222 = hipTextureGetValue < addressMode
+          > (data, i2, width, j2, height, k2, depth);
+
+      result =
+                (1 - a) * (1 - b) * (1 - c) * t111 + a * (1 - b) * (1 - c) * t211 +
+                (1 - a) * b * (1 - c) * t121 + a * b * (1 - c) * t221 +
+                (1 - a) * (1 - b) * c * t112 + a * (1 - b) * c * t212 +
+                (1 - a) * b * c * t122 + a * b * c * t222;
+
+    }
+      break;
+  }
+  return result;
+}
@@ -37,6 +37,12 @@ set(TEST_SRC
    hipGetChanDesc.cc
    hipTexObjPitch.cc
    hipTextureObj1DFetch.cc
+    hipBindTex2DPitch.cc
+    hipBindTexRef1DFetch.cc
+    hipTex1DFetchCheckModes.cc
+    hipTextureObj1DCheckModes.cc
+    hipTextureObj2DCheckModes.cc
+    hipTextureObj3DCheckModes.cc
 )

 hip_add_exe_to_target(NAME TextureTest
@@ -28,15 +28,19 @@ texture<TYPE_t, 2, hipReadModeElementType> tex;

 // texture object is a kernel argument
 static __global__ void texture2dCopyKernel(TYPE_t* dst) {
-    int x = threadIdx.x + blockIdx.x * blockDim.x;
-    int y = threadIdx.y + blockIdx.y * blockDim.y;
-    if ( (x < SIZE_W) && (y < SIZE_H) ) {
-        dst[SIZE_W*y+x] = tex2D(tex, x, y);
-    }
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
+  int x = threadIdx.x + blockIdx.x * blockDim.x;
+  int y = threadIdx.y + blockIdx.y * blockDim.y;
+  if ( (x < SIZE_W) && (y < SIZE_H) ) {
+    dst[SIZE_W*y+x] = tex2D(tex, x, y);
+  }
+#endif
 }


 TEST_CASE("Unit_hipBindTexture2D_Pitch") {
+  CHECK_IMAGE_SUPPORT
+
  TYPE_t* B;
  TYPE_t* A;
  TYPE_t* devPtrB;
@@ -24,14 +24,17 @@ THE SOFTWARE.
 texture<float, 1, hipReadModeElementType> tex;

 static __global__ void kernel(float *out) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
  int x = blockIdx.x * blockDim.x + threadIdx.x;
  if (x < N) {
      out[x] = tex1Dfetch(tex, x);
  }
+#endif
 }

-
 TEST_CASE("Unit_hipBindTexture_tex1DfetchVerification") {
+  CHECK_IMAGE_SUPPORT
+
  float *texBuf;
  float val[N], output[N];
  size_t offset = 0;
@@ -25,7 +25,7 @@ THE SOFTWARE.
 * Validate argument list of texture object api.
 */
 TEST_CASE("Unit_hipCreateTextureObject_ArgValidation") {
-  checkImageSupport();
+  CHECK_IMAGE_SUPPORT

  float *texBuf;
  hipError_t ret;
@@ -23,7 +23,7 @@ THE SOFTWARE.
 * Validates Array Resource texture object with negative/functional tests.
 */
 TEST_CASE("Unit_hipCreateTextureObject_ArrayResource") {
-  checkImageSupport();
+  CHECK_IMAGE_SUPPORT

  hipError_t ret;
  hipResourceDesc resDesc;
@@ -49,7 +49,7 @@ TEST_CASE("Unit_hipCreateTextureObject_ArrayResource") {
 * with negative/functional tests.
 */
 TEST_CASE("Unit_hipCreateTextureObject_MmArrayResource") {
-  checkImageSupport();
+  CHECK_IMAGE_SUPPORT

  hipError_t ret;
  hipResourceDesc resDesc;
@@ -26,7 +26,7 @@ THE SOFTWARE.
 * Validates Linear Resource texture object with negative/functional tests.
 */
 TEST_CASE("Unit_hipCreateTextureObject_LinearResource") {
-  checkImageSupport();
+  CHECK_IMAGE_SUPPORT

  float *texBuf;
  hipError_t ret;
@@ -29,7 +29,7 @@ THE SOFTWARE.
 * Validates Pitch2D Resource texture object with negative and functional tests
 */
 TEST_CASE("Unit_hipCreateTextureObject_Pitch2DResource") {
-  checkImageSupport();
+  CHECK_IMAGE_SUPPORT

  hipError_t ret;
  hipResourceDesc resDesc;
@@ -24,19 +24,11 @@ THE SOFTWARE.


 TEST_CASE("Unit_hipGetChannelDesc_CreateAndGet") {
+  CHECK_IMAGE_SUPPORT
+
  hipChannelFormatDesc chan_test, chan_desc;
  hipArray *hipArray;

-#if HT_AMD
-  int imageSupport{};
-  HIP_CHECK(hipDeviceGetAttribute(&imageSupport,
-                           hipDeviceAttributeImageSupport, 0));
-  if (!imageSupport) {
-    INFO("Texture is not supported on the device. Test is skipped");
-    return;
-  }
-#endif
-
  chan_desc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindSigned);
  HIP_CHECK(hipMallocArray(&hipArray, &chan_desc, C, R, 0));
  HIP_CHECK(hipGetChannelDesc(&chan_test, hipArray));
@@ -142,16 +142,10 @@ static void runTest_hipTextureFilterMode() {
 }

 TEST_CASE("Unit_hipNormalizedFloatValueTex_CheckModes") {
+  CHECK_IMAGE_SUPPORT
+
 #if HT_AMD
-  int imageSupport{};
-  HIP_CHECK(hipDeviceGetAttribute(&imageSupport,
-                           hipDeviceAttributeImageSupport, 0));
-  if (!imageSupport) {
-    INFO("Texture is not supported on the device. Test is skipped");
-    return;
-  }
  hipDeviceProp_t props;
-  HIP_CHECK(hipSetDevice(0));
  HIP_CHECK(hipGetDeviceProperties(&props, 0));
  INFO("Device :: " << props.name);
  INFO("Arch - AMD GPU :: " << props.gcnArch);
@@ -27,13 +27,17 @@ texture<float, hipTextureType2DLayered> tex2DL;

 __global__ void simpleKernelLayeredArray(T* outputData,
                                         int width, int height, int layer) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
  unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
  unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
  outputData[layer * width * height + y * width + x] = tex2DLayered(tex2DL,
                                                              x, y, layer);
+#endif
 }

 TEST_CASE("Unit_hipSimpleTexture2DLayered_Check") {
+  CHECK_IMAGE_SUPPORT
+
  constexpr int SIZE = 512;
  constexpr int num_layers = 5;
  constexpr unsigned int width = SIZE;
@@ -110,15 +110,8 @@ static void runSimpleTexture3D_Check(int width, int height, int depth,
 }

 TEST_CASE("Unit_hipSimpleTexture3D_Check_DataTypes") {
-#if HT_AMD
-  int imageSupport{};
-  HIP_CHECK(hipDeviceGetAttribute(&imageSupport,
-                           hipDeviceAttributeImageSupport, 0));
-  if (!imageSupport) {
-    INFO("Texture is not supported on the device. Test is skipped");
-    return;
-  }
-#endif
+  CHECK_IMAGE_SUPPORT
+
  for ( int i = 1; i < 25; i++ ) {
    runSimpleTexture3D_Check<float>(i, i, i, &texf);
    runSimpleTexture3D_Check<int>(i+1, i, i, &texi);
@@ -23,9 +23,11 @@ THE SOFTWARE.
 #define offset 3

 static __global__ void tex1dKernel(float *val, hipTextureObject_t obj) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
  int k = blockIdx.x * blockDim.x + threadIdx.x;
  if (k < (N - offset))
      val[k] = tex1Dfetch<float>(obj, k+offset);
+#endif
 }


@@ -84,7 +86,8 @@ static void runTest(hipTextureAddressMode addressMode,

  for (int i = 0; i < (N - offset); i++) {
      if (output[i] != val[i + offset]) {
-          INFO("Output not matching at index " << i);
+          INFO("Mismatch found at output[" << i << "]:" << output[i] <<
+               " val[" << i + offset << "]:" << val[i + offset]);
          REQUIRE(false);
      }
  }
@@ -103,16 +106,18 @@ static void runTest(hipTextureAddressMode addressMode,


 TEST_CASE("Unit_tex1Dfetch_CheckModes") {
-    SECTION("hipAddressModeClamp AND hipFilterModePoint") {
-      runTest(hipAddressModeClamp, hipFilterModePoint);
-    }
-    SECTION("hipAddressModeClamp AND hipFilterModeLinear") {
-      runTest(hipAddressModeClamp, hipFilterModeLinear);
-    }
-    SECTION("hipAddressModeWrap AND hipFilterModePoint") {
-      runTest(hipAddressModeWrap, hipFilterModePoint);
-    }
-    SECTION("hipAddressModeWrap AND hipFilterModeLinear") {
-      runTest(hipAddressModeWrap, hipFilterModeLinear);
-    }
+  CHECK_IMAGE_SUPPORT
+
+  SECTION("hipAddressModeClamp AND hipFilterModePoint") {
+    runTest(hipAddressModeClamp, hipFilterModePoint);
+  }
+  SECTION("hipAddressModeClamp AND hipFilterModeLinear") {
+    runTest(hipAddressModeClamp, hipFilterModeLinear);
+  }
+  SECTION("hipAddressModeWrap AND hipFilterModePoint") {
+    runTest(hipAddressModeWrap, hipFilterModePoint);
+  }
+  SECTION("hipAddressModeWrap AND hipFilterModeLinear") {
+    runTest(hipAddressModeWrap, hipFilterModeLinear);
+  }
 }
@@ -38,21 +38,13 @@ static __global__ void texture2dCopyKernel(hipTextureObject_t texObj,

 TEMPLATE_TEST_CASE("Unit_hipTexObjPitch_texture2D", "", float, int,
                    unsigned char, int16_t, char, unsigned int) {
+  CHECK_IMAGE_SUPPORT
+
  TestType* B;
  TestType* A;
  TestType* devPtrB;
  TestType* devPtrA;

-#if HT_AMD
-  int imageSupport{};
-  HIP_CHECK(hipDeviceGetAttribute(&imageSupport,
-                           hipDeviceAttributeImageSupport, 0));
-  if (!imageSupport) {
-    INFO("Texture is not supported on the device. Test is skipped");
-    return;
-  }
-#endif
-
  B = new TestType[SIZE_H*SIZE_W];
  A = new TestType[SIZE_H*SIZE_W];
  for (size_t i=1; i <= (SIZE_H*SIZE_W); i++) {
@@ -121,14 +121,8 @@ static void runMipMapTest(unsigned int width, unsigned int height, unsigned int
 #endif

 TEST_CASE("Unit_hipTextureMipmapObj2D_Check") {
-#if HT_AMD
-  int imageSupport{};
-  HIP_CHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, 0));
-  if (!imageSupport) {
-    INFO("Texture is not supported on the device. Test is skipped");
-    return;
-  }
-#endif
+  CHECK_IMAGE_SUPPORT
+
 #ifdef _WIN32
  for (auto& hw : hw_vector) {
    for (auto& mip : mip_vector) {
@@ -0,0 +1,148 @@
+/*
+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 tex1DKernel(float *outputData, hipTextureObject_t textureObject,
+                            int width, float offsetX) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
+  int x = blockIdx.x * blockDim.x + threadIdx.x;
+  outputData[x] = tex1D<float>(textureObject, normalizedCoords ? (x + offsetX) / width : x + offsetX);
+#endif
+}
+
+template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool normalizedCoords>
+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);
+  float *hData = (float*) malloc(size);
+  memset(hData, 0, size);
+  for (int j = 0; j < width; j++) {
+    hData[j] = j;
+  }
+
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(
+      32, 0, 0, 0, hipChannelFormatKindFloat);
+  hipArray *hipArray;
+  hipMallocArray(&hipArray, &channelDesc, width);
+
+  HIP_CHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, width * sizeof(float), width * sizeof(float), 1, 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.filterMode = filterMode;
+  texDesc.readMode = hipReadModeElementType;
+  texDesc.normalizedCoords = normalizedCoords;
+
+  // Create texture object
+  hipTextureObject_t textureObject = 0;
+  HIP_CHECK(hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL));
+
+  float *dData = nullptr;
+  hipMalloc((void**) &dData, size);
+
+  dim3 dimBlock(16, 1, 1);
+  dim3 dimGrid((width + dimBlock.x - 1)/ dimBlock.x, 1, 1);
+
+  hipLaunchKernelGGL(tex1DKernel<normalizedCoords>, dimGrid, dimBlock, 0, 0, dData,
+                     textureObject, width, offsetX);
+
+  hipDeviceSynchronize();
+
+  float *hOutputData = (float*) malloc(size);
+  memset(hOutputData, 0, size);
+  hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost);
+
+  bool result = true;
+  for (int j = 0; j < width; j++) {
+    float expectedValue = getExpectedValue<addressMode, filterMode>(width, offsetX + j, hData);
+    if (!hipTextureSamplingVerify<float, filterMode>(hOutputData[j], expectedValue)) {
+      INFO("Mismatch at " << offsetX + j << ":" << hOutputData[j] <<
+           " expected:" << expectedValue);
+      result = false;
+      break;
+    }
+  }
+
+  hipDestroyTextureObject(textureObject);
+  hipFree(dData);
+  hipFreeArray(hipArray);
+  free(hData);
+  free(hOutputData);
+  REQUIRE(result);
+}
+
+TEST_CASE("Unit_hipTextureObj1DCheckModes") {
+  CHECK_IMAGE_SUPPORT
+
+#ifdef _WIN32
+  INFO("Unit_hipTextureObj1DCheckModes skipped on Windows");
+  return;
+#endif
+  SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") {
+    runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, -3);
+    runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 4);
+  }
+
+  SECTION("hipAddressModeBorder, hipFilterModePoint, regularCoords") {
+    runTest<hipAddressModeBorder, hipFilterModePoint, false>(256, -8.5);
+    runTest<hipAddressModeBorder, hipFilterModePoint, false>(256, 12.5);
+  }
+
+  SECTION("hipAddressModeClamp, hipFilterModeLinear, regularCoords") {
+    runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, -3);
+    runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, 4);
+  }
+
+  SECTION("hipAddressModeBorder, hipFilterModeLinear, regularCoords") {
+    runTest<hipAddressModeBorder, hipFilterModeLinear, false>(256, -8.5);
+    runTest<hipAddressModeBorder, hipFilterModeLinear, false>(256, 12.5);
+  }
+
+  SECTION("hipAddressModeClamp, hipFilterModePoint, normalizedCoords") {
+    runTest<hipAddressModeClamp, hipFilterModePoint, true>(256, -3);
+    runTest<hipAddressModeClamp, hipFilterModePoint, true>(256, 4);
+  }
+
+  SECTION("hipAddressModeBorder, hipFilterModePoint, normalizedCoords") {
+    runTest<hipAddressModeBorder, hipFilterModePoint, true>(256, -8.5);
+    runTest<hipAddressModeBorder, hipFilterModePoint, true>(256, 12.5);
+  }
+
+  SECTION("hipAddressModeClamp, hipFilterModeLinear, normalizedCoords") {
+    runTest<hipAddressModeClamp, hipFilterModeLinear, true>(256, -3);
+    runTest<hipAddressModeClamp, hipFilterModeLinear, true>(256, 4);
+  }
+
+  SECTION("hipAddressModeBorder, hipFilterModeLinear, normalizedCoords") {
+    runTest<hipAddressModeBorder, hipFilterModeLinear, true>(256, -8.5);
+    runTest<hipAddressModeBorder, hipFilterModeLinear, true>(256, 12.5);
+  }
+}
@@ -22,14 +22,18 @@ THE SOFTWARE.
 #define N 512

 static __global__ void tex1dKernel(float *val, hipTextureObject_t obj) {
-    int k = blockIdx.x * blockDim.x + threadIdx.x;
-    if (k < N) {
-        val[k] = tex1Dfetch<float>(obj, k);
-    }
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
+  int k = blockIdx.x * blockDim.x + threadIdx.x;
+  if (k < N) {
+    val[k] = tex1Dfetch<float>(obj, k);
+  }
+#endif
 }


 TEST_CASE("Unit_hipCreateTextureObject_tex1DfetchVerification") {
+  CHECK_IMAGE_SUPPORT
+
  // Allocating the required buffer on gpu device
  float *texBuf, *texBufOut;
  float val[N], output[N];
@@ -22,22 +22,15 @@ THE SOFTWARE.
 __global__ void tex2DKernel(float* outputData,
                            hipTextureObject_t textureObject, int width) {
 #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<float>(textureObject, x, y);
+  int x = blockIdx.x * blockDim.x + threadIdx.x;
+  int y = blockIdx.y * blockDim.y + threadIdx.y;
+  outputData[y * width + x] = tex2D<float>(textureObject, x, y);
 #endif
 }

 TEST_CASE("Unit_hipTextureObj2D_Check") {
-#if HT_AMD
-  int imageSupport{};
-  HIP_CHECK(hipDeviceGetAttribute(&imageSupport,
-                           hipDeviceAttributeImageSupport, 0));
-  if (!imageSupport) {
-    INFO("Texture is not supported on the device. Test is skipped");
-    return;
-  }
-#endif
+  CHECK_IMAGE_SUPPORT
+
  constexpr int SIZE = 256;
  constexpr unsigned int width = SIZE;
  constexpr unsigned int height = SIZE;
@@ -0,0 +1,160 @@
+/*
+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 tex2DKernel(float *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<float>(textureObject,
+                                           normalizedCoords ? (x + offsetX) / width : x + offsetX,
+                                           normalizedCoords ? (y + offsetY) / height : y + offsetY);
+#endif
+}
+
+template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool normalizedCoords>
+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);
+  float *hData = (float*) 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] = index;
+    }
+  }
+
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(
+      32, 0, 0, 0, hipChannelFormatKindFloat);
+  hipArray *hipArray;
+  hipMallocArray(&hipArray, &channelDesc, width, height);
+
+  HIP_CHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, width * sizeof(float), width * sizeof(float), 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 = hipReadModeElementType;
+  texDesc.normalizedCoords = normalizedCoords;
+
+  // Create texture object
+  hipTextureObject_t textureObject = 0;
+  HIP_CHECK(hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL));
+
+  float *dData = nullptr;
+  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(tex2DKernel<normalizedCoords>, dimGrid, dimBlock, 0, 0, dData,
+                     textureObject, width, height, offsetX, offsetY);
+
+  hipDeviceSynchronize();
+
+  float *hOutputData = (float*) malloc(size);
+  memset(hOutputData, 0, size);
+  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;
+      float expectedValue = getExpectedValue<addressMode, filterMode>(width, height,
+                                                    offsetX + j, offsetY + i, hData);
+      if (!hipTextureSamplingVerify<float, filterMode>(hOutputData[index], expectedValue)) {
+        INFO("Mismatch at (" << offsetX + j << ", " << offsetY + i << "):" <<
+             hOutputData[index] << " expected:" << expectedValue);
+        result = false;
+        goto line1;
+      }
+    }
+  }
+line1:
+  hipDestroyTextureObject(textureObject);
+  hipFree(dData);
+  hipFreeArray(hipArray);
+  free(hData);
+  free(hOutputData);
+  REQUIRE(result);
+}
+
+TEST_CASE("Unit_hipTextureObj2DCheckModes") {
+  CHECK_IMAGE_SUPPORT
+
+#ifdef _WIN32
+  INFO("Unit_hipTextureObj2DCheckModes skipped on Windows");
+  return;
+#endif
+  SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") {
+    runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, -3.9, 6.1);
+    runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, 4.4, -7.0);
+  }
+
+  SECTION("hipAddressModeBorder, hipFilterModePoint, regularCoords") {
+    runTest<hipAddressModeBorder, hipFilterModePoint, false>(256, 256, -8.5, 2.9);
+    runTest<hipAddressModeBorder, hipFilterModePoint, false>(256, 256, 12.5, 6.7);
+  }
+
+  SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") {
+    runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, 256, -0.4, -0.4);
+    runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, 256, 4, 14.6);
+  }
+
+  SECTION("hipAddressModeBorder, hipFilterModeLinear, regularCoords") {
+    runTest<hipAddressModeBorder, hipFilterModeLinear, false>(256, 256, -0.4, 0.4);
+    runTest<hipAddressModeBorder, hipFilterModeLinear, false>(256, 256, 12.5, 23.7);
+  }
+
+  SECTION("hipAddressModeClamp, hipFilterModePoint, normalizedCoords") {
+    runTest<hipAddressModeClamp, hipFilterModePoint, true>(256, 256, -3, 8.9);
+    runTest<hipAddressModeClamp, hipFilterModePoint, true>(256, 256, 4, -0.1);
+  }
+
+  SECTION("hipAddressModeBorder, hipFilterModePoint, normalizedCoords") {
+    runTest<hipAddressModeBorder, hipFilterModePoint, true>(256, 256, -8.5, 15.9);
+    runTest<hipAddressModeBorder, hipFilterModePoint, true>(256, 256, 12.5, -17.9);
+  }
+
+  SECTION("hipAddressModeClamp, hipFilterModeLinear, normalizedCoords") {
+    runTest<hipAddressModeClamp, hipFilterModeLinear, true>(256, 256, -3, 5.8);
+    runTest<hipAddressModeClamp, hipFilterModeLinear, true>(256, 256, 4, 9.1);
+  }
+
+  SECTION("hipAddressModeBorder, hipFilterModeLinear, normalizedCoords") {
+    runTest<hipAddressModeBorder, hipFilterModeLinear, true>(256, 256, -8.5, 6.6);
+    runTest<hipAddressModeBorder, hipFilterModeLinear, true>(256, 256, 12.5, 0.01);
+  }
+}
@@ -0,0 +1,218 @@
+/*
+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>
+
+bool isGfx90a = false;
+
+template<bool normalizedCoords>
+__global__ void tex3DKernel(float *outputData, hipTextureObject_t textureObject,
+                            int width, int height, int depth, float offsetX,
+                            float offsetY, float offsetZ) {
+#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;
+  int z = blockIdx.z * blockDim.z + threadIdx.z;
+  outputData[z * width * height + y * width + x] = tex3D<float>(textureObject,
+                        normalizedCoords ? (x + offsetX) / width : x + offsetX,
+                        normalizedCoords ? (y + offsetY) / height : y + offsetY,
+                        normalizedCoords ? (z + offsetZ) / depth : z + offsetZ);
+#endif
+}
+
+template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool normalizedCoords>
+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,
+  //    depth, offsetX, offsetY, offsetZ);
+  bool result = true;
+  unsigned int size = width * height * depth * sizeof(float);
+  float *hData = (float*) malloc(size);
+  memset(hData, 0, size);
+
+  for (int i = 0; i < depth; i++) {
+    for (int j = 0; j < height; j++) {
+      for (int k = 0; k < width; k++) {
+        int index = i * width * height + j * width + k;
+        hData[index] = index;
+      }
+    }
+  }
+
+  // Allocate array and copy image data
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();
+  hipArray *arr;
+
+  HIP_CHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayDefault));
+  hipMemcpy3DParms myparms;
+  memset(&myparms, 0, sizeof(myparms));
+  myparms.srcPos = make_hipPos(0,0,0);
+  myparms.dstPos = make_hipPos(0,0,0);
+  myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(float), width, height);
+  myparms.dstArray = arr;
+  myparms.extent = make_hipExtent(width, height, depth);
+  myparms.kind = hipMemcpyHostToDevice;
+
+  HIP_CHECK(hipMemcpy3D(&myparms));
+
+  hipResourceDesc resDesc;
+  memset(&resDesc, 0, sizeof(resDesc));
+  resDesc.resType = hipResourceTypeArray;
+  resDesc.res.array.array = arr;
+
+  // Specify texture object parameters
+  hipTextureDesc texDesc;
+  memset(&texDesc, 0, sizeof(texDesc));
+  texDesc.addressMode[0] = addressMode;
+  texDesc.addressMode[1] = addressMode;
+  texDesc.addressMode[2] = addressMode;
+  texDesc.filterMode = filterMode;
+  texDesc.readMode = hipReadModeElementType;
+  texDesc.normalizedCoords = normalizedCoords;
+
+  // Create texture object
+  hipTextureObject_t textureObject = 0;
+  hipError_t res = hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL);
+  if (res != hipSuccess) {
+    hipFreeArray(arr);
+    free(hData);
+    if (res == hipErrorNotSupported && isGfx90a) {
+      printf("gfx90a doesn't support 3D linear filter! Skipped!\n");
+    } else {
+      result = false;
+    }
+    REQUIRE(result);
+    return;
+  }
+
+  float *dData = nullptr;
+  hipMalloc((void**) &dData, size);
+  hipMemset(dData, 0, size);
+  dim3 dimBlock(8, 8, 8); // 512 threads
+  dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y,
+               (depth + dimBlock.z - 1) / dimBlock.z);
+
+  hipLaunchKernelGGL(tex3DKernel<normalizedCoords>, dimGrid, dimBlock, 0, 0, dData,
+                     textureObject, width, height, depth, offsetX, offsetY, offsetZ);
+
+  hipDeviceSynchronize();
+
+  float *hOutputData = (float*) malloc(size);
+  memset(hOutputData, 0, size);
+  hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost);
+
+  for (int i = 0; i < depth; i++) {
+    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>(
+            width, height, depth, offsetX + k, offsetY + j, offsetZ + i, hData);
+
+        if (!hipTextureSamplingVerify<float, filterMode>(hOutputData[index], expectedValue)) {
+          INFO("Mismatch at (" << offsetX + k << ", " << offsetY + j << ", " << offsetZ + i << "):" <<
+               hOutputData[index] << " expected:" << expectedValue);
+          result = false;
+          goto line1;
+        }
+      }
+    }
+  }
+line1:
+  hipDestroyTextureObject(textureObject);
+  free(hOutputData);
+  hipFree(dData);
+  hipFreeArray(arr);
+  free(hData);
+  REQUIRE(result);
+
+}
+
+TEST_CASE("Unit_hipTextureObj3DCheckModes") {
+  CHECK_IMAGE_SUPPORT
+#ifdef _WIN32
+  INFO("Unit_hipTextureObj3DCheckModes skipped on Windows");
+  return;
+#endif
+
+  int device = 0;
+  hipDeviceProp_t props;
+  HIPCHECK(hipGetDeviceProperties(&props, device));
+  if (!strncmp(props.gcnArchName, "gfx90a", strlen("gfx90a"))) {
+    isGfx90a = true;
+  }
+
+  SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") {
+    runTest<hipAddressModeClamp, hipFilterModePoint, false>
+      (256, 256, 256, -3.9, 6.1, 9.5);
+    runTest<hipAddressModeClamp, hipFilterModePoint, false>
+      (256, 256, 256, 4.4, -7.0, 5.3);
+  }
+
+  SECTION("hipAddressModeBorder, hipFilterModePoint, regularCoords") {
+    runTest<hipAddressModeBorder, hipFilterModePoint, false>
+      (256, 256, 256, -8.5, 2.9, 5.8);
+    runTest<hipAddressModeBorder, hipFilterModePoint, false>
+      (256, 256, 256, 12.5, 6.7, 11.4);
+  }
+
+  SECTION("hipAddressModeClamp, hipFilterModeLinear, regularCoords") {
+    runTest<hipAddressModeClamp, hipFilterModeLinear, false>
+      (256, 256, 256, -0.4, -0.4, -0.4);
+    runTest<hipAddressModeClamp, hipFilterModeLinear, false>
+      (256, 256, 256, 4, 14.6, -0.3);
+  }
+
+  SECTION("hipAddressModeBorder, hipFilterModeLinear, regularCoords") {
+    runTest<hipAddressModeBorder, hipFilterModeLinear, false>
+      (256, 256, 256, 6.9, 7.4, 0.4);
+    runTest<hipAddressModeBorder, hipFilterModeLinear, false>
+      (256, 256, 256, 12.5, 23.7, 0.34);
+  }
+
+  SECTION("hipAddressModeClamp, hipFilterModePoint, normalizedCoords") {
+    runTest<hipAddressModeClamp, hipFilterModePoint, true>
+      (256, 256, 256, -3, 8.9, -4);
+    runTest<hipAddressModeClamp, hipFilterModePoint, true>
+      (256, 256, 256, 4, -0.1, 8.2);
+  }
+
+  SECTION("hipAddressModeBorder, hipFilterModePoint, normalizedCoords") {
+    runTest<hipAddressModeBorder, hipFilterModePoint, true>
+      (256, 256, 256, -8.5, 15.9, 0.1);
+    runTest<hipAddressModeBorder, hipFilterModePoint, true>
+      (256, 256, 256, 12.5, -17.9, -0.35);
+  }
+
+  SECTION("hipAddressModeClamp, hipFilterModeLinear, normalizedCoords") {
+    runTest<hipAddressModeClamp, hipFilterModeLinear, true>
+      (256, 256, 256, -3, 5.8, 0.89);
+    runTest<hipAddressModeClamp, hipFilterModeLinear, true>
+      (256, 256, 256, 4, 9.1, 2.08);
+  }
+
+  SECTION("hipAddressModeBorder, hipFilterModeLinear, normalizedCoords") {
+    runTest<hipAddressModeBorder, hipFilterModeLinear, true>
+      (256, 256, 256, -8.5, 6.6, 3.67);
+    runTest<hipAddressModeBorder, hipFilterModeLinear, true>
+      (256, 256, 256, 12.5, 0.01, -9.9);
+  }
+}
@@ -138,14 +138,12 @@ static inline bool isEqual(const T &val0, const T &val1) {
 }

 template<typename T>
-bool runTest(const char *description) {
+bool runTest() {
  const int N = 1024;
  bool testResult = true;
  // Allocating the required buffer on gpu device
  T *texBuf, *texBufOut;
  T val[N], output[N];
-  printf("%s<%s>(): size: %zu, %zu\n", __FUNCTION__, description,
-         sizeof(T), sizeof(decltype(T::x)));

  memset(output, 0, sizeof(output));
  std::srand(std::time(nullptr)); // use current time as seed for random generator
@@ -199,46 +197,45 @@ bool runTest(const char *description) {
  HIP_CHECK(hipFree(texBuf));
  HIP_CHECK(hipFree(texBufOut));

-  printf(": %s\n", testResult ? "succeeded" : "failed");
  REQUIRE(testResult == true);
  return testResult;
 }

 TEST_CASE("Unit_hipTextureFetch_vector") {
-  checkImageSupport();
+  CHECK_IMAGE_SUPPORT

  // test for char
-  runTest<char1>("char1");
-  runTest<char2>("char2");
-  runTest<char4>("char4");
+  runTest<char1>();
+  runTest<char2>();
+  runTest<char4>();

  // test for uchar
-  runTest<uchar1>("uchar1");
-  runTest<uchar2>("uchar2");
-  runTest<uchar4>("uchar4");
+  runTest<uchar1>();
+  runTest<uchar2>();
+  runTest<uchar4>();

  // test for short
-  runTest<short1>("short1");
-  runTest<short2>("short2");
-  runTest<short4>("short4");
+  runTest<short1>();
+  runTest<short2>();
+  runTest<short4>();

  // test for ushort
-  runTest<ushort1>("ushort1");
-  runTest<ushort2>("ushort2");
-  runTest<ushort4>("ushort4");
+  runTest<ushort1>();
+  runTest<ushort2>();
+  runTest<ushort4>();

  // test for int
-  runTest<int1>("int1");
-  runTest<int2>("int2");
-  runTest<int4>("int4");
+  runTest<int1>();
+  runTest<int2>();
+  runTest<int4>();

  // test for unsigned int
-  runTest<uint1>("uint1");
-  runTest<uint2>("uint2");
-  runTest<uint4>("uint4");
+  runTest<uint1>();
+  runTest<uint2>();
+  runTest<uint4>();

  // test for float
-  runTest<float1>("float1");
-  runTest<float2>("float2");
-  runTest<float4>("float4");
+  runTest<float1>();
+  runTest<float2>();
+  runTest<float4>();
 }
@@ -22,12 +22,16 @@ THE SOFTWARE.
 texture<float, 2, hipReadModeElementType> tex;

 __global__ void tex2DKernel(float* outputData, int width) {
+#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(tex, x, y);
+#endif
 }

 TEST_CASE("Unit_hipTextureRef2D_Check") {
+  CHECK_IMAGE_SUPPORT
+
  constexpr int SIZE = 256;
  constexpr unsigned int width = SIZE;
  constexpr unsigned int height = SIZE;