SWDEV-321872 - [catch2][dtest] Texture tests migrated to catch2 (#2536)

Change-Id: I4e29001812e2310680c6c24028667eb560194e87

[ROCm/hip-tests commit: 3f4fb28faf]

projects/hip-tests/catch/CMakeLists.txt

@@ -85,7 +85,7 @@ set(ADD_SCRIPT_PATH ${CMAKE_CURRENT_BINARY_DIR}/script/CatchAddTests.cmake)
 
 
 if(HIP_PLATFORM MATCHES "amd" AND HIP_COMPILER MATCHES "clang")
-    add_compile_options(-Wall -Wextra -pedantic -Werror)
+    add_compile_options(-Wall -Wextra -pedantic -Werror -Wno-deprecated)
 endif()
 
 cmake_policy(PUSH)

projects/hip-tests/catch/unit/texture/CMakeLists.txt

@@ -25,6 +25,12 @@ set(TEST_SRC
     hipCreateTextureObject_Pitch2D.cc
     hipCreateTextureObject_Array.cc
     hipTextureObjFetchVector.cc
+    hipNormalizedFloatValueTex.cc
+    hipTextureObj2D.cc
+    hipSimpleTexture3D.cc
+    hipTextureRef2D.cc
+    hipSimpleTexture2DLayered.cc
+    hipTextureMipmapObj2D.cc
 )
 
 hip_add_exe_to_target(NAME TextureTest

projects/hip-tests/catch/unit/texture/hipNormalizedFloatValueTex.cc

@@ -0,0 +1,166 @@
+/*
+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>
+
+#define SIZE          10
+#define EPSILON       0.00001
+#define THRESH_HOLD   0.01  // For filter mode
+
+static float getNormalizedValue(const float value,
+                                const hipChannelFormatDesc& desc) {
+  if ((desc.x == 8) && (desc.f == hipChannelFormatKindSigned))
+    return (value / SCHAR_MAX);
+  if ((desc.x == 8) && (desc.f == hipChannelFormatKindUnsigned))
+    return (value / UCHAR_MAX);
+  if ((desc.x == 16) && (desc.f == hipChannelFormatKindSigned))
+    return (value / SHRT_MAX);
+  if ((desc.x == 16) && (desc.f == hipChannelFormatKindUnsigned))
+    return (value / USHRT_MAX);
+  return value;
+}
+
+texture<char, hipTextureType1D, hipReadModeNormalizedFloat>            texc;
+texture<unsigned char, hipTextureType1D, hipReadModeNormalizedFloat>   texuc;
+
+template<typename T>
+__global__ void normalizedValTextureTest(unsigned int numElements,
+                                         float* pDst) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
+  unsigned int elementID = hipThreadIdx_x;
+  if (elementID >= numElements)
+    return;
+  float coord = elementID/static_cast<float>(numElements);
+  if (std::is_same<T, char>::value)
+    pDst[elementID] = tex1D(texc, coord);
+  else if (std::is_same<T, unsigned char>::value)
+    pDst[elementID] = tex1D(texuc, coord);
+#endif
+}
+
+static void textureVerifyFilterModePoint(float *hOutputData,
+                                         float *expected, int size) {
+  for (int i = 0; i < size; i++) {
+    if ((hOutputData[i] == expected[i])
+        || (i >= 1 && hOutputData[i] == expected[i - 1]) ||  // round down
+        (i < (size - 1) && hOutputData[i] == expected[i + 1])) {  // round up
+      continue;
+    }
+      INFO("Mismatch at output[" << i << "]:" << hOutputData[i] <<
+           " expected[" << i << "]:" << expected[i]);
+    if (i >= 1) {
+      INFO(", expected[" << i - 1 << "]:" << expected[i - 1]);
+    }
+    if (i < (size - 1)) {
+      INFO(", expected[" << i + 1 << "]:" << expected[i + 1]);
+    }
+    REQUIRE(false);
+  }
+}
+
+static void textureVerifyFilterModeLinear(float *hOutputData,
+                                          float *expected,  int size) {
+  for (int i = 0; i < size; i++) {
+    float mean = (fabs(expected[i]) + fabs(hOutputData[i])) / 2;
+    float ratio = fabs(expected[i] - hOutputData[i]) / (mean + EPSILON);
+    if (ratio > THRESH_HOLD) {
+      INFO("Mismatch found at output[" << i << "]:" << hOutputData[i] <<
+           " expected[" << i << "]:" << expected[i] << ", ratio:" << ratio);
+      REQUIRE(false);
+    }
+  }
+}
+
+template<hipTextureFilterMode fMode = hipFilterModePoint>
+static void textureVerify(float *hOutputData, float *expected, size_t size) {
+  if (fMode == hipFilterModePoint) {
+    textureVerifyFilterModePoint(hOutputData, expected, size);
+  } else if (fMode == hipFilterModeLinear) {
+    textureVerifyFilterModeLinear(hOutputData, expected, size);
+  }
+}
+
+template<typename T, hipTextureFilterMode fMode = hipFilterModePoint>
+static void textureTest(texture<T, hipTextureType1D,
+                        hipReadModeNormalizedFloat> *tex) {
+  hipChannelFormatDesc desc = hipCreateChannelDesc<T>();
+  hipArray_t dData;
+  HIP_CHECK(hipMallocArray(&dData, &desc, SIZE, 1, hipArrayDefault));
+
+  T hData[] = {65, 66, 67, 68, 69, 70, 71, 72, 73, 74};
+  HIP_CHECK(hipMemcpy2DToArray(dData, 0, 0, hData, sizeof(T) * SIZE,
+            sizeof(T) * SIZE, 1, hipMemcpyHostToDevice));
+
+  tex->normalized = true;
+  tex->channelDesc = desc;
+  tex->filterMode = fMode;
+  HIP_CHECK(hipBindTextureToArray(tex, dData, &desc));
+
+  float *dOutputData = NULL;
+  HIP_CHECK(hipMalloc(&dOutputData, sizeof(float) * SIZE));
+  REQUIRE(dOutputData != nullptr);
+
+  hipLaunchKernelGGL(normalizedValTextureTest<T>, dim3(1, 1, 1),
+                     dim3(SIZE, 1, 1), 0, 0, SIZE, dOutputData);
+
+  float *hOutputData = new float[SIZE];
+  REQUIRE(hOutputData != nullptr);
+  HIP_CHECK(hipMemcpy(hOutputData, dOutputData, (sizeof(float) * SIZE),
+                     hipMemcpyDeviceToHost));
+
+  float expected[SIZE];
+  for (int i = 0; i < SIZE; i++) {
+    expected[i] = getNormalizedValue(static_cast<float>(hData[i]), desc);
+  }
+  textureVerify<fMode>(hOutputData, expected, SIZE);
+
+  HIP_CHECK(hipFreeArray(dData));
+  HIP_CHECK(hipFree(dOutputData));
+  delete [] hOutputData;
+}
+
+template<hipTextureFilterMode fMode = hipFilterModePoint>
+static void runTest_hipTextureFilterMode() {
+  textureTest<char, fMode>(&texc);
+  textureTest<unsigned char, fMode>(&texuc);
+}
+
+TEST_CASE("Unit_hipNormalizedFloatValueTex_CheckModes") {
+#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);
+#endif
+
+  SECTION("hipNormalizedFloatValueTexture for hipFilterModePoint") {
+    runTest_hipTextureFilterMode<hipFilterModePoint>();
+  }
+  SECTION("hipNormalizedFloatValueTexture for hipFilterModeLinear") {
+    runTest_hipTextureFilterMode<hipFilterModeLinear>();
+  }
+}

projects/hip-tests/catch/unit/texture/hipSimpleTexture2DLayered.cc

@@ -0,0 +1,103 @@
+/*
+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>
+
+typedef float T;
+
+// Texture reference for 2D Layered texture
+texture<float, hipTextureType2DLayered> tex2DL;
+
+__global__ void simpleKernelLayeredArray(T* outputData,
+                                         int width, int height, int layer) {
+  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);
+}
+
+TEST_CASE("Unit_hipSimpleTexture2DLayered_Check") {
+  constexpr int SIZE = 512;
+  constexpr int num_layers = 5;
+  constexpr unsigned int width = SIZE;
+  constexpr unsigned int height = SIZE;
+  constexpr unsigned int size = width * height * num_layers * sizeof(T);
+
+  T* hData = reinterpret_cast<T*>(malloc(size));
+  REQUIRE(hData != nullptr);
+  memset(hData, 0, size);
+
+  for (unsigned int layer = 0; layer < num_layers; layer++) {
+    for (int i = 0; i < static_cast<int>(width * height); i++) {
+      hData[layer * width * height + i] = i;
+    }
+  }
+  hipChannelFormatDesc channelDesc;
+  // Allocate array and copy image data
+  channelDesc = hipCreateChannelDesc(sizeof(T)*8, 0, 0, 0,
+                                     hipChannelFormatKindFloat);
+  hipArray *arr;
+
+  HIP_CHECK(hipMalloc3DArray(&arr, &channelDesc,
+               make_hipExtent(width, height, num_layers), hipArrayLayered));
+  hipMemcpy3DParms myparms{};
+  myparms.srcPos = make_hipPos(0, 0, 0);
+  myparms.dstPos = make_hipPos(0, 0, 0);
+  myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
+  myparms.dstArray = arr;
+  myparms.extent = make_hipExtent(width , height, num_layers);
+  // myparms.kind = hipMemcpyHostToDevice;
+  HIP_CHECK(hipMemcpy3D(&myparms));
+
+  // set texture parameters
+  tex2DL.addressMode[0] = hipAddressModeWrap;
+  tex2DL.addressMode[1] = hipAddressModeWrap;
+  tex2DL.filterMode = hipFilterModePoint;
+  tex2DL.normalized = false;
+
+  // Bind the array to the texture
+  HIP_CHECK(hipBindTextureToArray(tex2DL, arr, channelDesc));
+
+  // Allocate device memory for result
+  T* dData = nullptr;
+  HIP_CHECK(hipMalloc(&dData, size));
+  REQUIRE(dData != nullptr);
+
+  dim3 dimBlock(8, 8, 1);
+  dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
+  for (unsigned int layer = 0; layer < num_layers; layer++)
+    hipLaunchKernelGGL(simpleKernelLayeredArray, dimGrid, dimBlock, 0, 0,
+                       dData, width, height, layer);
+  HIP_CHECK(hipDeviceSynchronize());
+
+  // Allocate mem for the result on host side
+  T *hOutputData = reinterpret_cast<T*>(malloc(size));
+  REQUIRE(hOutputData != nullptr);
+  memset(hOutputData, 0,  size);
+
+  // copy result from device to host
+  HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost));
+  HipTest::checkArray(hData, hOutputData, width, height, num_layers);
+
+  HIP_CHECK(hipFree(dData));
+  HIP_CHECK(hipFreeArray(arr));
+  free(hData);
+  free(hOutputData);
+}

projects/hip-tests/catch/unit/texture/hipSimpleTexture3D.cc

@@ -0,0 +1,127 @@
+/*
+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>
+
+// Texture reference for 3D texture
+texture<float, hipTextureType3D, hipReadModeElementType> texf;
+texture<int, hipTextureType3D, hipReadModeElementType>   texi;
+texture<char, hipTextureType3D, hipReadModeElementType>  texc;
+
+template <typename T>
+__global__ void simpleKernel3DArray(T* outputData, int width,
+                                    int height, int depth) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
+  for (int i = 0; i < depth; i++) {
+    for (int j = 0; j < height; j++) {
+      for (int k = 0; k < width; k++) {
+        if (std::is_same<T, float>::value)
+          outputData[i*width*height + j*width + k] = tex3D(texf, k, j, i);
+        else if (std::is_same<T, int>::value)
+          outputData[i*width*height + j*width + k] = tex3D(texi, k, j, i);
+        else if (std::is_same<T, char>::value)
+          outputData[i*width*height + j*width + k] = tex3D(texc, k, j, i);
+      }
+    }
+  }
+#endif
+}
+
+template <typename T>
+static void runSimpleTexture3D_Check(int width, int height, int depth,
+            texture<T, hipTextureType3D, hipReadModeElementType> *tex) {
+  unsigned int size = width * height * depth * sizeof(T);
+  T* hData = reinterpret_cast<T*>(malloc(size));
+  REQUIRE(hData != nullptr);
+  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++) {
+        hData[i*width*height + j*width +k] = i*width*height + j*width + k;
+      }
+    }
+  }
+
+  // Allocate array and copy image data
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
+  hipArray *arr;
+
+  HIP_CHECK(hipMalloc3DArray(&arr, &channelDesc,
+            make_hipExtent(width, height, depth), hipArrayDefault));
+  hipMemcpy3DParms myparms{};
+  myparms.srcPos = make_hipPos(0, 0, 0);
+  myparms.dstPos = make_hipPos(0, 0, 0);
+  myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
+  myparms.dstArray = arr;
+  myparms.extent = make_hipExtent(width, height, depth);
+  myparms.kind = hipMemcpyHostToDevice;
+
+  HIP_CHECK(hipMemcpy3D(&myparms));
+
+  // set texture parameters
+  tex->addressMode[0] = hipAddressModeWrap;
+  tex->addressMode[1] = hipAddressModeWrap;
+  tex->filterMode = hipFilterModePoint;
+  tex->normalized = false;
+
+  // Bind the array to the texture
+  HIP_CHECK(hipBindTextureToArray(*tex, arr, channelDesc));
+
+  // Allocate device memory for result
+  T* dData = nullptr;
+  HIP_CHECK(hipMalloc(&dData, size));
+  REQUIRE(dData != nullptr);
+
+  hipLaunchKernelGGL(simpleKernel3DArray, dim3(1, 1, 1), dim3(1, 1, 1),
+                     0, 0, dData, width, height, depth);
+  HIP_CHECK(hipDeviceSynchronize());
+
+  // Allocate mem for the result on host side
+  T *hOutputData = reinterpret_cast<T*>(malloc(size));
+  REQUIRE(hOutputData != nullptr);
+  memset(hOutputData, 0,  size);
+
+  // copy result from device to host
+  HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost));
+  HipTest::checkArray(hData, hOutputData, width, height, depth);
+
+  HIP_CHECK(hipFree(dData));
+  HIP_CHECK(hipFreeArray(arr));
+  free(hData);
+  free(hOutputData);
+}
+
+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
+  for ( int i = 1; i < 25; i++ ) {
+    runSimpleTexture3D_Check<float>(i, i, i, &texf);
+    runSimpleTexture3D_Check<int>(i+1, i, i, &texi);
+    runSimpleTexture3D_Check<char>(i, i+1, i, &texc);
+  }
+}

projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc

@@ -0,0 +1,153 @@
+/*
+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>
+
+// Height Width Vector
+std::vector<unsigned int> hw_vector = {2048, 1024, 512, 256, 64};
+std::vector<unsigned int> mip_vector = {8, 4, 2, 1};
+
+__global__ void tex2DKernel(float* outputData,
+                            hipTextureObject_t textureObject,
+                            int width, float level) {
+#ifndef __gfx90a__
+#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] = tex2DLod<float>(textureObject, x, y, level);
+#endif
+#endif
+}
+
+#ifdef _WIN32  // MipMap is currently supported only on windows
+static void runMipMapTest(unsigned int width, unsigned int height,
+                          unsigned int mipmap_level) {
+  INFO("Width: " << width << "Height: " << height << "mip: " << mipmap_level);
+
+  // Create new width & height to be tested
+  unsigned int orig_width = width;
+  unsigned int orig_height = height;
+  unsigned int i, j;
+  width /= pow(2, mipmap_level);
+  height /= pow(2, mipmap_level);
+  unsigned int size = width * height * sizeof(float);
+
+  float* hData = reinterpret_cast<float*>(malloc(size));
+  REQUIRE(hData != nullptr);
+  memset(hData, 0, size);
+  for (i = 0; i < height; i++) {
+    for (j = 0; j < width; j++) {
+      hData[i * width + j] = i * width + j;
+    }
+  }
+
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, 0,
+                                            hipChannelFormatKindFloat);
+  HIP_ARRAY3D_DESCRIPTOR mipmapped_array_desc;
+  memset(&mipmapped_array_desc, 0x00, sizeof(HIP_ARRAY3D_DESCRIPTOR));
+  mipmapped_array_desc.Width = orig_width;
+  mipmapped_array_desc.Height = orig_height;
+  mipmapped_array_desc.Depth = 0;
+  mipmapped_array_desc.Format = HIP_AD_FORMAT_FLOAT;
+  mipmapped_array_desc.NumChannels = ((channelDesc.x != 0) +
+      (channelDesc.y != 0) + (channelDesc.z != 0) + (channelDesc.w != 0));
+  mipmapped_array_desc.Flags = 0;
+
+  hipMipmappedArray* mip_array_ptr;
+  HIP_CHECK(hipMipmappedArrayCreate(&mip_array_ptr, &mipmapped_array_desc,
+                                    2 * mipmap_level));
+
+  hipArray *hipArray = nullptr;
+  HIP_CHECK(hipMipmappedArrayGetLevel(&hipArray, mip_array_ptr, mipmap_level));
+  HIP_CHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size,
+                             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] = hipAddressModeWrap;
+  texDesc.addressMode[1] = hipAddressModeWrap;
+  texDesc.filterMode = hipFilterModePoint;
+  texDesc.readMode = hipReadModeElementType;
+  texDesc.normalizedCoords = 0;
+
+  // Create texture object
+  hipTextureObject_t textureObject = 0;
+  HIP_CHECK(hipCreateTextureObject(&textureObject, &resDesc,
+                                   &texDesc, nullptr));
+
+  float* dData = nullptr;
+  HIP_CHECK(hipMalloc(&dData, size));
+  REQUIRE(dData != nullptr);
+
+  dim3 dimBlock(16, 16, 1);
+  dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
+
+  hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData,
+                     textureObject, width, (2 * mipmap_level));
+  hipDeviceSynchronize();
+
+  float* hOutputData = reinterpret_cast<float*>(malloc(size));
+  REQUIRE(hOutputData != nullptr);
+  memset(hOutputData, 0, size);
+  HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost));
+
+  for (i = 0; i < height; i++) {
+    for (j = 0; j < width; j++) {
+      if (hData[i * width + j] != hOutputData[i * width + j]) {
+        INFO("Difference found at [ " << i << j << " ]: " <<
+              hData[i * width + j] << hOutputData[i * width + j]);
+        REQUIRE(false);
+      }
+    }
+  }
+  HIP_CHECK(hipDestroyTextureObject(textureObject));
+  HIP_CHECK(hipFree(dData));
+  HIP_CHECK(hipFreeArray(hipArray));
+  free(hData);
+}
+#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
+#ifdef _WIN32
+  for (auto& hw : hw_vector) {
+    for (auto& mip : mip_vector) {
+      if ((hw / static_cast<int>(pow(2, (mip * 2)))) > 0) {
+        runMipMapTest(hw, hw, mip);
+      }
+    }
+  }
+#else
+  SUCCEED("Mipmaps are Supported only on windows, skipping the test.");
+#endif
+}

projects/hip-tests/catch/unit/texture/hipTextureObj2D.cc

@@ -0,0 +1,113 @@
+/*
+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>
+
+__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);
+#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
+  constexpr int SIZE = 256;
+  constexpr unsigned int width = SIZE;
+  constexpr unsigned int height = SIZE;
+  constexpr unsigned int size = width * height * sizeof(float);
+  unsigned int i, j;
+
+  float* dData = nullptr;
+  HIP_CHECK(hipMalloc(&dData, size));
+  REQUIRE(dData != nullptr);
+
+  float* hOutputData = reinterpret_cast<float*>(malloc(size));
+  REQUIRE(hOutputData != nullptr);
+  memset(hOutputData, 0, size);
+
+  float* hData = reinterpret_cast<float*>(malloc(size));
+  REQUIRE(hData != nullptr);
+  memset(hData, 0, size);
+  for (i = 0; i < height; i++) {
+    for (j = 0; j < width; j++) {
+      hData[i * width + j] = i * width + j;
+    }
+  }
+
+  hipChannelFormatDesc channelDesc =
+     hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat);
+  hipArray* hipArray;
+  HIP_CHECK(hipMallocArray(&hipArray, &channelDesc, width, height));
+  HIP_CHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size,
+                             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] = hipAddressModeWrap;
+  texDesc.addressMode[1] = hipAddressModeWrap;
+  texDesc.filterMode = hipFilterModePoint;
+  texDesc.readMode = hipReadModeElementType;
+  texDesc.normalizedCoords = 0;
+
+  // Create texture object
+  hipTextureObject_t textureObject = 0;
+  HIP_CHECK(hipCreateTextureObject(&textureObject, &resDesc,
+                                   &texDesc, nullptr));
+
+  dim3 dimBlock(16, 16, 1);
+  dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
+
+  hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock),
+                     0, 0, dData, textureObject, width);
+
+  HIP_CHECK(hipDeviceSynchronize());
+  HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost));
+
+  for (i = 0; i < height; i++) {
+    for (j = 0; j < width; j++) {
+      if (hData[i * width + j] != hOutputData[i * width + j]) {
+        INFO("Difference found at [ " << i << j << " ]: " <<
+              hData[i * width + j] << hOutputData[i * width + j]);
+        REQUIRE(false);
+      }
+    }
+  }
+
+  HIP_CHECK(hipDestroyTextureObject(textureObject));
+  HIP_CHECK(hipFree(dData));
+  HIP_CHECK(hipFreeArray(hipArray));
+  free(hData);
+}

projects/hip-tests/catch/unit/texture/hipTextureRef2D.cc

@@ -0,0 +1,87 @@
+/*
+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>
+
+texture<float, 2, hipReadModeElementType> tex;
+
+__global__ void tex2DKernel(float* outputData, int width) {
+  int x = blockIdx.x * blockDim.x + threadIdx.x;
+  int y = blockIdx.y * blockDim.y + threadIdx.y;
+  outputData[y * width + x] = tex2D(tex, x, y);
+}
+
+TEST_CASE("Unit_hipTextureRef2D_Check") {
+  constexpr int SIZE = 256;
+  constexpr unsigned int width = SIZE;
+  constexpr unsigned int height = SIZE;
+  constexpr unsigned int size = width * height * sizeof(float);
+  unsigned int i, j;
+
+  float* hData = reinterpret_cast<float*>(malloc(size));
+  REQUIRE(hData != nullptr);
+  memset(hData, 0, size);
+  for (i = 0; i < height; i++) {
+    for (j = 0; j < width; j++) {
+      hData[i * width + j] = i * width + j;
+    }
+  }
+
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, 0,
+                                           hipChannelFormatKindFloat);
+  hipArray* hipArray;
+  HIP_CHECK(hipMallocArray(&hipArray, &channelDesc, width, height));
+  HIP_CHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size,
+                             hipMemcpyHostToDevice));
+
+  tex.addressMode[0] = hipAddressModeWrap;
+  tex.addressMode[1] = hipAddressModeWrap;
+  tex.filterMode = hipFilterModePoint;
+  tex.normalized = 0;
+
+  HIP_CHECK(hipBindTextureToArray(tex, hipArray, channelDesc));
+
+  float* dData = nullptr;
+  HIP_CHECK(hipMalloc(&dData, size));
+  REQUIRE(dData != nullptr);
+
+  dim3 dimBlock(16, 16, 1);
+  dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
+  hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0,
+                     dData, width);
+  hipDeviceSynchronize();
+
+  float* hOutputData = reinterpret_cast<float*>(malloc(size));
+  REQUIRE(hOutputData != nullptr);
+  memset(hOutputData, 0, size);
+  HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost));
+
+  for (i = 0; i < height; i++) {
+    for (j = 0; j < width; j++) {
+      if (hData[i * width + j] != hOutputData[i * width + j]) {
+        INFO("Difference found at [ " << i << j << " ]: " <<
+              hData[i * width + j] << hOutputData[i * width + j]);
+        REQUIRE(false);
+      }
+    }
+  }
+  HIP_CHECK(hipUnbindTexture(tex));
+  HIP_CHECK(hipFree(dData));
+  HIP_CHECK(hipFreeArray(hipArray));
+}