EXSWHTEC-317 - Implement tests for 1D texture device functions #366

Change-Id: Iafefada132fb94a1b120d91915e24524e7d766d0

catch/unit/texture/CMakeLists.txt

@@ -51,6 +51,13 @@ set(TEST_SRC
     hipMipmappedArrayCreate.cc
     hipMipmappedArrayDestroy.cc
     hipMipmappedArrayGetLevel.cc
+    tex1Dfetch.cc
+    tex1D.cc
+    tex1DLayered.cc
+    tex1DGrad.cc
+    tex1DLayeredGrad.cc
+    tex1DLayeredLod.cc
+    tex1DLod.cc
 )
 
 if(WIN32)

catch/unit/texture/kernels.hh

@@ -33,6 +33,14 @@ __host__ __device__ inline float GetCoordinate(size_t iteration, size_t N, size_
   return normalized_coords ? x / dim : x;
 }
 
+template <typename TexelType>
+__global__ void tex1DfetchKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj) {
+  const auto tid = cg::this_grid().thread_rank();
+  if (tid >= N) return;
+
+  out[tid] = tex1D<TexelType>(tex_obj, tid);
+}
+
 template <typename TexelType>
 __global__ void tex1DKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj,
                             size_t width, size_t num_subdivisions, bool normalized_coords) {
@@ -43,6 +51,50 @@ __global__ void tex1DKernel(TexelType* const out, size_t N, hipTextureObject_t t
   out[tid] = tex1D<TexelType>(tex_obj, x);
 }
 
+template <typename TexelType>
+__global__ void tex1DLodKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj,
+                               size_t width, size_t num_subdivisions, bool normalized_coords,
+                               float level_of_detail) {
+  const auto tid = cg::this_grid().thread_rank();
+  if (tid >= N) return;
+
+  float x = GetCoordinate(tid, N, width, num_subdivisions, normalized_coords);
+  out[tid] = tex1DLod<TexelType>(tex_obj, x, level_of_detail);
+}
+
+template <typename TexelType>
+__global__ void tex1DLayeredLodKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj,
+                                      size_t width, size_t num_subdivisions, bool normalized_coords,
+                                      int layer, float level_of_detail) {
+  const auto tid = cg::this_grid().thread_rank();
+  if (tid >= N) return;
+
+  float x = GetCoordinate(tid, N, width, num_subdivisions, normalized_coords);
+  out[tid] = tex1DLayeredLod<TexelType>(tex_obj, x, layer, level_of_detail);
+}
+
+template <typename TexelType>
+__global__ void tex1DGradKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj,
+                                size_t width, size_t num_subdivisions, bool normalized_coords,
+                                float dx, float dy) {
+  const auto tid = cg::this_grid().thread_rank();
+  if (tid >= N) return;
+
+  float x = GetCoordinate(tid, N, width, num_subdivisions, normalized_coords);
+  out[tid] = tex1DGrad<TexelType>(tex_obj, x, dx, dy);
+}
+
+template <typename TexelType>
+__global__ void tex1DLayeredGradKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj,
+                                       size_t width, size_t num_subdivisions,
+                                       bool normalized_coords, float dx, float dy, int layer) {
+  const auto tid = cg::this_grid().thread_rank();
+  if (tid >= N) return;
+
+  float x = GetCoordinate(tid, N, width, num_subdivisions, normalized_coords);
+  out[tid] = tex1DLayeredGrad<TexelType>(tex_obj, x, layer, dx, dy);
+}
+
 template <typename TexelType>
 __global__ void tex2DKernel(TexelType* const out, size_t N_x, size_t N_y,
                             hipTextureObject_t tex_obj, size_t width, size_t height,

catch/unit/texture/tex1D.cc

@@ -0,0 +1,140 @@
+/*
+Copyright (c) 2023 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 "kernels.hh"
+#include "test_fixture.hh"
+
+/**
+ * @addtogroup tex1D tex1D
+ * @{
+ * @ingroup TextureTest
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Test texture fetching with `tex1D` and read mode set to `hipReadModeElementType`. The test
+ * is performed with:
+ *      - normalized coordinates
+ *      - non-normalized coordinates
+ *      - Nearest-point sampling
+ *      - Linear filtering
+ *      - All combinations of different addressing modes.
+ * Test source
+ * ------------------------
+ *    - unit/texture/tex1D.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_tex1D_Positive_ReadModeElementType", "", char, unsigned char, short,
+                   unsigned short, int, unsigned int, float) {
+  CHECK_IMAGE_SUPPORT;
+
+  TextureTestParams<TestType> params = {};
+  params.extent = make_hipExtent(1024, 0, 0);
+  params.num_subdivisions = 4;
+  params.GenerateTextureDesc();
+
+  TextureTestFixture<TestType> fixture{params};
+
+  const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
+  tex1DKernel<vec4<TestType>><<<num_blocks, num_threads>>>(
+      fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
+      params.num_subdivisions, params.tex_desc.normalizedCoords);
+
+  fixture.LoadOutput();
+
+  for (auto i = 0u; i < params.NumItersX(); ++i) {
+    float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
+                            params.tex_desc.normalizedCoords);
+
+    INFO("Index: " << i);
+    INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
+    INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
+    INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
+    INFO("x: " << std::fixed << std::setprecision(16) << x);
+
+    auto ref_val = fixture.tex_h.Tex1D(x, params.tex_desc);
+    REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
+    REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
+    REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
+    REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Test texture fetching with `tex1D` and read mode set to `hipReadModeNormalizedFloat`. The
+ * test is performed with:
+ *      - normalized coordinates
+ *      - non-normalized coordinates
+ *      - Nearest-point sampling
+ *      - Linear filtering
+ *      - All combinations of different addressing modes.
+ * Test source
+ * ------------------------
+ *    - unit/texture/tex1D.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_tex1D_Positive_ReadModeNormalizedFloat", "", char, unsigned char, short,
+                   unsigned short) {
+  CHECK_IMAGE_SUPPORT;
+
+  TextureTestParams<TestType> params = {};
+  params.extent = make_hipExtent(1024, 0, 0);
+  params.num_subdivisions = 4;
+  params.GenerateTextureDesc(hipReadModeNormalizedFloat);
+
+  TextureTestFixture<TestType, true> fixture{params};
+
+  const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
+  tex1DKernel<vec4<float>><<<num_blocks, num_threads>>>(
+      fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
+      params.num_subdivisions, params.tex_desc.normalizedCoords);
+
+  fixture.LoadOutput();
+
+  for (auto i = 0u; i < params.NumItersX(); ++i) {
+    float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
+                            params.tex_desc.normalizedCoords);
+
+    INFO("i: " << i);
+    INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
+    INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
+    INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
+    INFO("Filter mode: " << FilteringModeToString(params.tex_desc.filterMode));
+    INFO("x: " << std::fixed << std::setprecision(16) << x);
+
+    auto ref_val =
+        Vec4Map<TestType>(fixture.tex_h.Tex1D(x, params.tex_desc), NormalizeInteger<TestType>);
+    REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
+    REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
+    REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
+    REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
+  }
+}

catch/unit/texture/tex1DGrad.cc

@@ -0,0 +1,140 @@
+/*
+Copyright (c) 2023 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 "kernels.hh"
+#include "test_fixture.hh"
+
+/**
+ * @addtogroup tex1DGrad tex1DGrad
+ * @{
+ * @ingroup TextureTest
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Test texture fetching with `tex1DGrad` and read mode set to `hipReadModeElementType`. The
+ * test is performed with:
+ *      - normalized coordinates
+ *      - non-normalized coordinates
+ *      - Nearest-point sampling
+ *      - Linear filtering
+ *      - All combinations of different addressing modes.
+ * Test source
+ * ------------------------
+ *    - unit/texture/tex1DGrad.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeElementType", "", char, unsigned char, short,
+                   unsigned short, int, unsigned int, float) {
+  CHECK_IMAGE_SUPPORT;
+
+  TextureTestParams<TestType> params = {};
+  params.extent = make_hipExtent(1024, 0, 0);
+  params.num_subdivisions = 4;
+  params.GenerateTextureDesc();
+
+  TextureTestFixture<TestType> fixture{params};
+
+  const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
+  tex1DGradKernel<vec4<TestType>><<<num_blocks, num_threads>>>(
+      fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
+      params.num_subdivisions, params.tex_desc.normalizedCoords, 0.5f, 0.5f);
+
+  fixture.LoadOutput();
+
+  for (auto i = 0u; i < params.NumItersX(); ++i) {
+    float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
+                            params.tex_desc.normalizedCoords);
+
+    INFO("Index: " << i);
+    INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
+    INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
+    INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
+    INFO("x: " << std::fixed << std::setprecision(16) << x);
+
+    auto ref_val = fixture.tex_h.Tex1D(x, params.tex_desc);
+    REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
+    REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
+    REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
+    REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Test texture fetching with `tex1DGrad` and read mode set to `hipReadModeNormalizedFloat`.
+ * The test is performed with:
+ *      - normalized coordinates
+ *      - non-normalized coordinates
+ *      - Nearest-point sampling
+ *      - Linear filtering
+ *      - All combinations of different addressing modes.
+ * Test source
+ * ------------------------
+ *    - unit/texture/tex1DGrad.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeNormalizedFloat", "", char, unsigned char,
+                   short, unsigned short) {
+  CHECK_IMAGE_SUPPORT;
+
+  TextureTestParams<TestType> params = {};
+  params.extent = make_hipExtent(1024, 0, 0);
+  params.num_subdivisions = 4;
+  params.GenerateTextureDesc(hipReadModeNormalizedFloat);
+
+  TextureTestFixture<TestType, true> fixture{params};
+
+  const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
+  tex1DGradKernel<vec4<float>><<<num_blocks, num_threads>>>(
+      fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
+      params.num_subdivisions, params.tex_desc.normalizedCoords, 0.5f, 0.5f);
+
+  fixture.LoadOutput();
+
+  for (auto i = 0u; i < params.NumItersX(); ++i) {
+    float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
+                            params.tex_desc.normalizedCoords);
+
+    INFO("i: " << i);
+    INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
+    INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
+    INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
+    INFO("Filter mode: " << FilteringModeToString(params.tex_desc.filterMode));
+    INFO("x: " << std::fixed << std::setprecision(16) << x);
+
+    auto ref_val =
+        Vec4Map<TestType>(fixture.tex_h.Tex1D(x, params.tex_desc), NormalizeInteger<TestType>);
+    REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
+    REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
+    REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
+    REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
+  }
+}

catch/unit/texture/tex1DLayered.cc

@@ -0,0 +1,149 @@
+/*
+Copyright (c) 2023 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 "kernels.hh"
+#include "test_fixture.hh"
+
+/**
+ * @addtogroup tex1DLayered tex1DLayered
+ * @{
+ * @ingroup TextureTest
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Test texture fetching with `tex1DLayered` and read mode set to `hipReadModeElementType`. The
+ * test is performed with:
+ *      - normalized coordinates
+ *      - non-normalized coordinates
+ *      - Nearest-point sampling
+ *      - Linear filtering
+ *      - All combinations of different addressing modes.
+ * Test source
+ * ------------------------
+ *    - unit/texture/tex1DLayered.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_tex1DLayered_Positive_ReadModeElementType", "", char, unsigned char, short,
+                   unsigned short, int, unsigned int, float) {
+  CHECK_IMAGE_SUPPORT;
+
+  TextureTestParams<TestType> params = {};
+  params.extent = make_hipExtent(1024, 0, 0);
+  params.layers = 2;
+  params.num_subdivisions = 4;
+  params.GenerateTextureDesc();
+
+  TextureTestFixture<TestType> fixture{params};
+
+  const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
+
+  for (auto layer = 0u; layer < params.layers; ++layer) {
+    tex1DLayeredKernel<vec4<TestType>><<<num_blocks, num_threads>>>(
+        fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
+        params.num_subdivisions, params.tex_desc.normalizedCoords, layer);
+
+    fixture.LoadOutput();
+
+    for (auto i = 0u; i < params.NumItersX(); ++i) {
+      float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
+                              params.tex_desc.normalizedCoords);
+
+      INFO("Layer: " << layer);
+      INFO("i: " << i);
+      INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
+      INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
+      INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
+      INFO("x: " << std::fixed << std::setprecision(16) << x);
+
+      const auto ref_val = fixture.tex_h.Tex1DLayered(x, layer, params.tex_desc);
+      REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
+      REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
+      REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
+      REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Test texture fetching with `tex1DLayered` and read mode set to `hipReadModeNormalizedFloat`.
+ * The test is performed with:
+ *      - normalized coordinates
+ *      - non-normalized coordinates
+ *      - Nearest-point sampling
+ *      - Linear filtering
+ *      - All combinations of different addressing modes.
+ * Test source
+ * ------------------------
+ *    - unit/texture/tex1DLayered.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_tex1DLayered_Positive_ReadModeNormalizedFloat", "", char, unsigned char,
+                   short, unsigned short) {
+  CHECK_IMAGE_SUPPORT;
+
+  TextureTestParams<TestType> params = {};
+  params.extent = make_hipExtent(1024, 0, 0);
+  params.layers = 2;
+  params.num_subdivisions = 4;
+  params.GenerateTextureDesc(hipReadModeNormalizedFloat);
+
+  TextureTestFixture<TestType, true> fixture{params};
+
+  const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
+
+  for (auto layer = 0u; layer < params.layers; ++layer) {
+    tex1DLayeredKernel<vec4<float>><<<num_blocks, num_threads>>>(
+        fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
+        params.num_subdivisions, params.tex_desc.normalizedCoords, layer);
+
+    fixture.LoadOutput();
+
+    for (auto i = 0u; i < params.NumItersX(); ++i) {
+      float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
+                              params.tex_desc.normalizedCoords);
+
+      INFO("Layer: " << layer);
+      INFO("Index: " << i);
+      INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
+      INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
+      INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
+      INFO("x: " << std::fixed << std::setprecision(16) << x);
+
+      auto ref_val = Vec4Map<TestType>(fixture.tex_h.Tex1DLayered(x, layer, params.tex_desc),
+                                       NormalizeInteger<TestType>);
+      REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
+      REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
+      REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
+      REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
+    }
+  }
+}

catch/unit/texture/tex1DLayeredGrad.cc

@@ -0,0 +1,140 @@
+/*
+Copyright (c) 2023 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 "kernels.hh"
+#include "test_fixture.hh"
+
+/**
+ * @addtogroup tex1DLayeredGrad tex1DLayeredGrad
+ * @{
+ * @ingroup TextureTest
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Test texture fetching with `tex1DLayeredGrad` and read mode set to `hipReadModeElementType`.
+ * The test is performed with:
+ *      - normalized coordinates
+ *      - non-normalized coordinates
+ *      - Nearest-point sampling
+ *      - Linear filtering
+ *      - All combinations of different addressing modes.
+ * Test source
+ * ------------------------
+ *    - unit/texture/tex1DLayeredGrad.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_tex1DLayeredGrad_Positive_ReadModeElementType", "", char, unsigned char,
+                   short, unsigned short, int, unsigned int, float) {
+  CHECK_IMAGE_SUPPORT;
+
+  TextureTestParams<TestType> params = {};
+  params.extent = make_hipExtent(1024, 0, 0);
+  params.num_subdivisions = 4;
+  params.GenerateTextureDesc();
+
+  TextureTestFixture<TestType> fixture{params};
+
+  const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
+  tex1DLayeredGradKernel<vec4<TestType>><<<num_blocks, num_threads>>>(
+      fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
+      params.num_subdivisions, params.tex_desc.normalizedCoords, 0.5f, 0.5f, 0);
+
+  fixture.LoadOutput();
+
+  for (auto i = 0u; i < params.NumItersX(); ++i) {
+    float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
+                            params.tex_desc.normalizedCoords);
+
+    INFO("Index: " << i);
+    INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
+    INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
+    INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
+    INFO("x: " << std::fixed << std::setprecision(16) << x);
+
+    auto ref_val = fixture.tex_h.Tex1D(x, params.tex_desc);
+    REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
+    REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
+    REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
+    REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Test texture fetching with `tex1DLayeredGrad` and read mode set to
+ * `hipReadModeNormalizedFloat`. The test is performed with:
+ *      - normalized coordinates
+ *      - non-normalized coordinates
+ *      - Nearest-point sampling
+ *      - Linear filtering
+ *      - All combinations of different addressing modes.
+ * Test source
+ * ------------------------
+ *    - unit/texture/tex1DLayeredGrad.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_tex1DLayeredGrad_Positive_ReadModeNormalizedFloat", "", char,
+                   unsigned char, short, unsigned short) {
+  CHECK_IMAGE_SUPPORT;
+
+  TextureTestParams<TestType> params = {};
+  params.extent = make_hipExtent(1024, 0, 0);
+  params.num_subdivisions = 4;
+  params.GenerateTextureDesc(hipReadModeNormalizedFloat);
+
+  TextureTestFixture<TestType, true> fixture{params};
+
+  const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
+  tex1DLayeredGradKernel<vec4<float>><<<num_blocks, num_threads>>>(
+      fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
+      params.num_subdivisions, params.tex_desc.normalizedCoords, 0.5f, 0.5f, 0);
+
+  fixture.LoadOutput();
+
+  for (auto i = 0u; i < params.NumItersX(); ++i) {
+    float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
+                            params.tex_desc.normalizedCoords);
+
+    INFO("i: " << i);
+    INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
+    INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
+    INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
+    INFO("Filter mode: " << FilteringModeToString(params.tex_desc.filterMode));
+    INFO("x: " << std::fixed << std::setprecision(16) << x);
+
+    auto ref_val =
+        Vec4Map<TestType>(fixture.tex_h.Tex1D(x, params.tex_desc), NormalizeInteger<TestType>);
+    REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
+    REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
+    REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
+    REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
+  }
+}

catch/unit/texture/tex1DLayeredLod.cc

@@ -0,0 +1,140 @@
+/*
+Copyright (c) 2023 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 "kernels.hh"
+#include "test_fixture.hh"
+
+/**
+ * @addtogroup tex1DLayeredLod tex1DLayeredLod
+ * @{
+ * @ingroup TextureTest
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Test texture fetching with `tex1DLayeredLod` and read mode set to `hipReadModeElementType`.
+ * The test is performed with:
+ *      - normalized coordinates
+ *      - non-normalized coordinates
+ *      - Nearest-point sampling
+ *      - Linear filtering
+ *      - All combinations of different addressing modes.
+ * Test source
+ * ------------------------
+ *    - unit/texture/tex1DLayeredLod.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_tex1DLayeredLod_Positive_ReadModeElementType", "", char, unsigned char,
+                   short, unsigned short, int, unsigned int, float) {
+  CHECK_IMAGE_SUPPORT;
+
+  TextureTestParams<TestType> params = {};
+  params.extent = make_hipExtent(1024, 0, 0);
+  params.num_subdivisions = 4;
+  params.GenerateTextureDesc();
+
+  TextureTestFixture<TestType> fixture{params};
+
+  const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
+  tex1DLayeredLodKernel<vec4<TestType>><<<num_blocks, num_threads>>>(
+      fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
+      params.num_subdivisions, params.tex_desc.normalizedCoords, 0, 0);
+
+  fixture.LoadOutput();
+
+  for (auto i = 0u; i < params.NumItersX(); ++i) {
+    float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
+                            params.tex_desc.normalizedCoords);
+
+    INFO("Index: " << i);
+    INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
+    INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
+    INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
+    INFO("x: " << std::fixed << std::setprecision(16) << x);
+
+    auto ref_val = fixture.tex_h.Tex1D(x, params.tex_desc);
+    REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
+    REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
+    REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
+    REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Test texture fetching with `tex1DLayeredLod` and read mode set to
+ * `hipReadModeNormalizedFloat`. The test is performed with:
+ *      - normalized coordinates
+ *      - non-normalized coordinates
+ *      - Nearest-point sampling
+ *      - Linear filtering
+ *      - All combinations of different addressing modes.
+ * Test source
+ * ------------------------
+ *    - unit/texture/tex1DLayeredLod.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_tex1DLayeredLod_Positive_ReadModeNormalizedFloat", "", char, unsigned char,
+                   short, unsigned short) {
+  CHECK_IMAGE_SUPPORT;
+
+  TextureTestParams<TestType> params = {};
+  params.extent = make_hipExtent(1024, 0, 0);
+  params.num_subdivisions = 4;
+  params.GenerateTextureDesc(hipReadModeNormalizedFloat);
+
+  TextureTestFixture<TestType, true> fixture{params};
+
+  const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
+  tex1DLayeredLodKernel<vec4<float>><<<num_blocks, num_threads>>>(
+      fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
+      params.num_subdivisions, params.tex_desc.normalizedCoords, 0, 0);
+
+  fixture.LoadOutput();
+
+  for (auto i = 0u; i < params.NumItersX(); ++i) {
+    float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
+                            params.tex_desc.normalizedCoords);
+
+    INFO("i: " << i);
+    INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
+    INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
+    INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
+    INFO("Filter mode: " << FilteringModeToString(params.tex_desc.filterMode));
+    INFO("x: " << std::fixed << std::setprecision(16) << x);
+
+    auto ref_val =
+        Vec4Map<TestType>(fixture.tex_h.Tex1D(x, params.tex_desc), NormalizeInteger<TestType>);
+    REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
+    REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
+    REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
+    REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
+  }
+}

catch/unit/texture/tex1DLod.cc

@@ -0,0 +1,140 @@
+/*
+Copyright (c) 2023 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 "kernels.hh"
+#include "test_fixture.hh"
+
+/**
+ * @addtogroup tex1DLod tex1DLod
+ * @{
+ * @ingroup TextureTest
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Test texture fetching with `tex1DLod` and read mode set to `hipReadModeElementType`. The
+ * test is performed with:
+ *      - normalized coordinates
+ *      - non-normalized coordinates
+ *      - Nearest-point sampling
+ *      - Linear filtering
+ *      - All combinations of different addressing modes.
+ * Test source
+ * ------------------------
+ *    - unit/texture/tex1DLod.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeElementType", "", char, unsigned char, short,
+                   unsigned short, int, unsigned int, float) {
+  CHECK_IMAGE_SUPPORT;
+
+  TextureTestParams<TestType> params = {};
+  params.extent = make_hipExtent(1024, 0, 0);
+  params.num_subdivisions = 4;
+  params.GenerateTextureDesc();
+
+  TextureTestFixture<TestType> fixture{params};
+
+  const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
+  tex1DLodKernel<vec4<TestType>><<<num_blocks, num_threads>>>(
+      fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
+      params.num_subdivisions, params.tex_desc.normalizedCoords, 0);
+
+  fixture.LoadOutput();
+
+  for (auto i = 0u; i < params.NumItersX(); ++i) {
+    float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
+                            params.tex_desc.normalizedCoords);
+
+    INFO("Index: " << i);
+    INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
+    INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
+    INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
+    INFO("x: " << std::fixed << std::setprecision(16) << x);
+
+    auto ref_val = fixture.tex_h.Tex1D(x, params.tex_desc);
+    REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
+    REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
+    REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
+    REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Test texture fetching with `tex1DLod` and read mode set to `hipReadModeNormalizedFloat`. The
+ * test is performed with:
+ *      - normalized coordinates
+ *      - non-normalized coordinates
+ *      - Nearest-point sampling
+ *      - Linear filtering
+ *      - All combinations of different addressing modes.
+ * Test source
+ * ------------------------
+ *    - unit/texture/tex1DLod.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeNormalizedFloat", "", char, unsigned char, short,
+                   unsigned short) {
+  CHECK_IMAGE_SUPPORT;
+
+  TextureTestParams<TestType> params = {};
+  params.extent = make_hipExtent(1024, 0, 0);
+  params.num_subdivisions = 4;
+  params.GenerateTextureDesc(hipReadModeNormalizedFloat);
+
+  TextureTestFixture<TestType, true> fixture{params};
+
+  const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
+  tex1DLodKernel<vec4<float>><<<num_blocks, num_threads>>>(
+      fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
+      params.num_subdivisions, params.tex_desc.normalizedCoords, 0);
+
+  fixture.LoadOutput();
+
+  for (auto i = 0u; i < params.NumItersX(); ++i) {
+    float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
+                            params.tex_desc.normalizedCoords);
+
+    INFO("i: " << i);
+    INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
+    INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
+    INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
+    INFO("Filter mode: " << FilteringModeToString(params.tex_desc.filterMode));
+    INFO("x: " << std::fixed << std::setprecision(16) << x);
+
+    auto ref_val =
+        Vec4Map<TestType>(fixture.tex_h.Tex1D(x, params.tex_desc), NormalizeInteger<TestType>);
+    REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
+    REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
+    REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
+    REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
+  }
+}

catch/unit/texture/tex1Dfetch.cc

@@ -0,0 +1,162 @@
+/*
+Copyright (c) 2023 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 <vector>
+
+#include <hip_test_common.hh>
+#include <resource_guards.hh>
+
+#include "kernels.hh"
+#include "utils.hh"
+#include "vec4.hh"
+
+/**
+ * @addtogroup tex1D tex1D
+ * @{
+ * @ingroup TextureTest
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Test texture fetching with `tex1Dfetch` and read mode set to `hipReadModeElementType`.
+ * Test source
+ * ------------------------
+ *    - unit/texture/tex1Dfetch.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_tex1Dfetch_Positive_ReadModeElementType", "", char, unsigned char, short,
+                   unsigned short, int, unsigned int, float) {
+  CHECK_IMAGE_SUPPORT;
+
+  std::vector<vec4<TestType>> tex_h(1024);
+  for (auto i = 0u; i < tex_h.size(); ++i) {
+    tex_h[i].x = i + 7;
+    tex_h[i].y = i + 7;
+    tex_h[i].z = i + 7;
+    tex_h[i].w = i + 7;
+  }
+
+  const auto alloc_size = tex_h.size() * sizeof(vec4<TestType>);
+  LinearAllocGuard<vec4<TestType>> tex_alloc_d(LinearAllocs::hipMalloc, alloc_size);
+  HIP_CHECK(hipMemcpy(tex_alloc_d.ptr(), tex_h.data(), alloc_size, hipMemcpyHostToDevice));
+
+  hipResourceDesc res_desc;
+  memset(&res_desc, 0, sizeof(res_desc));
+  res_desc.resType = hipResourceTypeLinear;
+  res_desc.res.linear.devPtr = tex_alloc_d.ptr();
+  res_desc.res.linear.desc = hipCreateChannelDesc<vec4<TestType>>();
+  res_desc.res.linear.sizeInBytes = alloc_size;
+
+  hipTextureDesc tex_desc;
+  memset(&tex_desc, 0, sizeof(tex_desc));
+  tex_desc.filterMode = hipFilterModePoint;
+  tex_desc.readMode = hipReadModeElementType;
+  tex_desc.normalizedCoords = false;
+  tex_desc.addressMode[0] = hipAddressModeClamp;
+
+  LinearAllocGuard<vec4<TestType>> out_alloc_d(LinearAllocs::hipMalloc, alloc_size);
+  TextureGuard tex(&res_desc, &tex_desc);
+
+  const auto num_threads = std::min<size_t>(1024, tex_h.size());
+  const auto num_blocks = (tex_h.size() + num_threads - 1) / num_threads;
+  tex1DfetchKernel<vec4<TestType>>
+      <<<num_blocks, num_threads>>>(out_alloc_d.ptr(), tex_h.size(), tex.object());
+
+  std::vector<vec4<TestType>> out_alloc_h(tex_h.size());
+  HIP_CHECK(hipMemcpy(out_alloc_h.data(), out_alloc_d.ptr(), alloc_size, hipMemcpyDeviceToHost));
+  HIP_CHECK(hipDeviceSynchronize());
+
+  for (auto i = 0u; i < out_alloc_h.size(); ++i) {
+    INFO("Index: " << i);
+    const auto ref_val = tex_h[i];
+    REQUIRE(ref_val.x == out_alloc_h[i].x);
+    REQUIRE(ref_val.y == out_alloc_h[i].y);
+    REQUIRE(ref_val.z == out_alloc_h[i].z);
+    REQUIRE(ref_val.w == out_alloc_h[i].w);
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Test texture fetching with `tex1Dfetch` and read mode set to `hipReadModeNormalizedFloat`.
+ * Test source
+ * ------------------------
+ *    - unit/texture/tex1Dfetch.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_tex1Dfetch_Positive_ReadModeNormalizedFloat", "", char, unsigned char,
+                   short, unsigned short) {
+  CHECK_IMAGE_SUPPORT;
+
+  std::vector<vec4<TestType>> tex_h(1024);
+  for (auto i = 0u; i < tex_h.size(); ++i) {
+    tex_h[i].x = i + 7;
+    tex_h[i].y = i + 7;
+    tex_h[i].z = i + 7;
+    tex_h[i].w = i + 7;
+  }
+
+  const auto alloc_size = tex_h.size() * sizeof(vec4<TestType>);
+  LinearAllocGuard<vec4<TestType>> tex_alloc_d(LinearAllocs::hipMalloc, alloc_size);
+  HIP_CHECK(hipMemcpy(tex_alloc_d.ptr(), tex_h.data(), alloc_size, hipMemcpyHostToDevice));
+
+  hipResourceDesc res_desc;
+  memset(&res_desc, 0, sizeof(res_desc));
+  res_desc.resType = hipResourceTypeLinear;
+  res_desc.res.linear.devPtr = tex_alloc_d.ptr();
+  res_desc.res.linear.desc = hipCreateChannelDesc<vec4<TestType>>();
+  res_desc.res.linear.sizeInBytes = alloc_size;
+
+  hipTextureDesc tex_desc;
+  memset(&tex_desc, 0, sizeof(tex_desc));
+  tex_desc.filterMode = hipFilterModePoint;
+  tex_desc.readMode = hipReadModeElementType;
+  tex_desc.normalizedCoords = false;
+  tex_desc.addressMode[0] = hipAddressModeClamp;
+
+  LinearAllocGuard<vec4<float>> out_alloc_d(LinearAllocs::hipMalloc, alloc_size);
+  TextureGuard tex(&res_desc, &tex_desc);
+
+  const auto num_threads = std::min<size_t>(1024, tex_h.size());
+  const auto num_blocks = (tex_h.size() + num_threads - 1) / num_threads;
+  tex1DfetchKernel<vec4<float>>
+      <<<num_blocks, num_threads>>>(out_alloc_d.ptr(), tex_h.size(), tex.object());
+
+  std::vector<vec4<float>> out_alloc_h(tex_h.size());
+  HIP_CHECK(hipMemcpy(out_alloc_h.data(), out_alloc_d.ptr(), alloc_size, hipMemcpyDeviceToHost));
+  HIP_CHECK(hipDeviceSynchronize());
+
+  for (auto i = 0u; i < out_alloc_h.size(); ++i) {
+    INFO("Index: " << i);
+    const auto ref_val = Vec4Map<TestType>(tex_h[i], NormalizeInteger<TestType>);
+    REQUIRE(ref_val.x == out_alloc_h[i].x);
+    REQUIRE(ref_val.y == out_alloc_h[i].y);
+    REQUIRE(ref_val.z == out_alloc_h[i].z);
+    REQUIRE(ref_val.w == out_alloc_h[i].w);
+  }
+}