diff --git a/catch/include/resource_guards.hh b/catch/include/resource_guards.hh
index 20c1a20ee5..c2f32e39f5 100644
--- a/catch/include/resource_guards.hh
+++ b/catch/include/resource_guards.hh
@@ -216,6 +216,42 @@ template <typename T> class ArrayAllocGuard {
   const hipExtent extent_;
 };
 
+template <typename T> class MipmappedArrayAllocGuard {
+ public:
+  // extent should contain logical width
+  MipmappedArrayAllocGuard(const hipExtent extent, const unsigned int levels,
+                           const unsigned int flags)
+      : extent_{extent}, levels_{levels} {
+    hipChannelFormatDesc desc = hipCreateChannelDesc<T>();
+    HIP_CHECK(hipMallocMipmappedArray(&ptr_, &desc, extent_, levels_, flags));
+  }
+
+  MipmappedArrayAllocGuard(const hipExtent extent, const unsigned int flags = 0u)
+      : MipmappedArrayAllocGuard{extent, 1, flags} {}
+
+  ~MipmappedArrayAllocGuard() { static_cast<void>(hipFreeMipmappedArray(ptr_)); }
+
+  MipmappedArrayAllocGuard(const MipmappedArrayAllocGuard&) = delete;
+  MipmappedArrayAllocGuard(MipmappedArrayAllocGuard&&) = delete;
+
+  hipMipmappedArray_t ptr() const { return ptr_; }
+
+  hipArray_t GetLevel(unsigned int level) {
+    hipArray_t ret;
+    HIP_CHECK(hipGetMipmappedArrayLevel(&ret, ptr_, level));
+    return ret;
+  }
+
+  hipExtent extent() const { return extent_; }
+
+  unsigned int levels() const { return levels_; }
+
+ private:
+  hipMipmappedArray_t ptr_ = nullptr;
+  const hipExtent extent_;
+  const unsigned int levels_;
+};
+
 template <typename T> class DrvArrayAllocGuard {
  public:
   // extent should contain width in bytes
diff --git a/catch/unit/texture/kernels.hh b/catch/unit/texture/kernels.hh
index f9f7a6a41e..2e0bf256cd 100644
--- a/catch/unit/texture/kernels.hh
+++ b/catch/unit/texture/kernels.hh
@@ -87,7 +87,7 @@ __global__ void tex1DGradKernel(TexelType* const out, size_t N, hipTextureObject
 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) {
+                                       bool normalized_coords, int layer, float dx, float dy) {
   const auto tid = cg::this_grid().thread_rank();
   if (tid >= N) return;
 
diff --git a/catch/unit/texture/test_fixture.hh b/catch/unit/texture/test_fixture.hh
index 28ca9ee2df..222c4deb9e 100644
--- a/catch/unit/texture/test_fixture.hh
+++ b/catch/unit/texture/test_fixture.hh
@@ -99,19 +99,25 @@ template <typename TestType> struct TextureTestParams {
     tex_desc.addressMode[0] = address_mode_x;
     if (extent.height) tex_desc.addressMode[1] = address_mode_y;
     if (extent.depth) tex_desc.addressMode[2] = address_mode_z;
+
+    tex_desc.mipmapFilterMode = tex_desc.filterMode;
   }
 };
 
-template <typename TestType, bool normalized_read = false> struct TextureTestFixture {
+template <typename TestType, bool normalized_read = false, bool mipmap = false>
+struct TextureTestFixture {
   using VecType = vec4<TestType>;
   using OutType = std::conditional_t<normalized_read, vec4<float>, VecType>;
+  template <typename T>
+  using ArrayAllocGuardType =
+      std::conditional_t<mipmap, MipmappedArrayAllocGuard<T>, ArrayAllocGuard<T>>;
 
   TextureTestParams<TestType> params;
   hipResourceDesc res_desc;
 
   LinearAllocGuard<VecType> host_alloc;
   TextureReference<VecType> tex_h;
-  ArrayAllocGuard<VecType> tex_alloc_d;
+  ArrayAllocGuardType<VecType> tex_alloc_d;
   TextureGuard tex;
   LinearAllocGuard<OutType> out_alloc_d;
   std::vector<OutType> out_alloc_h;
@@ -131,9 +137,13 @@ template <typename TestType, bool normalized_read = false> struct TextureTestFix
       SetVec4<TestType>(host_alloc.ptr()[i], i + test_value_offset);
     }
 
-    hipMemcpy3DParms memcpy_params;
+    hipMemcpy3DParms memcpy_params = {};
     memset(&memcpy_params, 0 sizeof(hipMemcpy3DParms));
-    memcpy_params.dstArray = tex_alloc_d.ptr();
+    if constexpr (mipmap) {
+      memcpy_params.dstArray = tex_alloc_d.GetLevel(0);
+    } else {
+      memcpy_params.dstArray = tex_alloc_d.ptr();
+    }
     memcpy_params.extent = params.LayeredExtent();
     memcpy_params.extent.height = memcpy_params.extent.height ?: 1;
     memcpy_params.extent.depth = memcpy_params.extent.depth ?: 1;
@@ -143,8 +153,13 @@ template <typename TestType, bool normalized_read = false> struct TextureTestFix
     HIP_CHECK(hipMemcpy3D(&memcpy_params));
 
     memset(&res_desc, 0, sizeof(res_desc));
-    res_desc.resType = hipResourceTypeArray;
-    res_desc.res.array.array = tex_alloc_d.ptr();
+    if constexpr (mipmap) {
+      res_desc.resType = hipResourceTypeMipmappedArray;
+      res_desc.res.mipmap.mipmap = tex_alloc_d.ptr();
+    } else {
+      res_desc.resType = hipResourceTypeArray;
+      res_desc.res.array.array = tex_alloc_d.ptr();
+    }
     return &res_desc;
   }
 
diff --git a/catch/unit/texture/tex1D.cc b/catch/unit/texture/tex1D.cc
index de92fca2ee..17f360d17f 100644
--- a/catch/unit/texture/tex1D.cc
+++ b/catch/unit/texture/tex1D.cc
@@ -127,7 +127,6 @@ TEMPLATE_TEST_CASE("Unit_tex1D_Positive_ReadModeNormalizedFloat", "", char, unsi
     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 =
diff --git a/catch/unit/texture/tex1DGrad.cc b/catch/unit/texture/tex1DGrad.cc
index 5b893d0954..1c006571b1 100644
--- a/catch/unit/texture/tex1DGrad.cc
+++ b/catch/unit/texture/tex1DGrad.cc
@@ -46,7 +46,7 @@ THE SOFTWARE.
  *    - unit/texture/tex1DGrad.cc
  * Test requirements
  * ------------------------
- *    - HIP_VERSION >= 5.2
+ *    - HIP_VERSION >= 5.7
  */
 TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeElementType", "", char, unsigned char, short,
                    unsigned short, int, unsigned int, float) {
@@ -57,7 +57,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeElementType", "", char, unsi
   params.num_subdivisions = 4;
   params.GenerateTextureDesc();
 
-  TextureTestFixture<TestType> fixture{params};
+  TextureTestFixture<TestType, false, true> fixture{params};
 
   const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
   tex1DGradKernel<vec4<TestType>><<<num_blocks, num_threads>>>(
@@ -99,7 +99,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeElementType", "", char, unsi
  *    - unit/texture/tex1DGrad.cc
  * Test requirements
  * ------------------------
- *    - HIP_VERSION >= 5.2
+ *    - HIP_VERSION >= 5.7
  */
 TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeNormalizedFloat", "", char, unsigned char,
                    short, unsigned short) {
@@ -110,7 +110,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeNormalizedFloat", "", char,
   params.num_subdivisions = 4;
   params.GenerateTextureDesc(hipReadModeNormalizedFloat);
 
-  TextureTestFixture<TestType, true> fixture{params};
+  TextureTestFixture<TestType, true, true> fixture{params};
 
   const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
   tex1DGradKernel<vec4<float>><<<num_blocks, num_threads>>>(
@@ -127,7 +127,6 @@ TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeNormalizedFloat", "", char,
     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 =
diff --git a/catch/unit/texture/tex1DLayeredGrad.cc b/catch/unit/texture/tex1DLayeredGrad.cc
index 07c9734619..6115c939b8 100644
--- a/catch/unit/texture/tex1DLayeredGrad.cc
+++ b/catch/unit/texture/tex1DLayeredGrad.cc
@@ -46,7 +46,7 @@ THE SOFTWARE.
  *    - unit/texture/tex1DLayeredGrad.cc
  * Test requirements
  * ------------------------
- *    - HIP_VERSION >= 5.2
+ *    - HIP_VERSION >= 5.7
  */
 TEMPLATE_TEST_CASE("Unit_tex1DLayeredGrad_Positive_ReadModeElementType", "", char, unsigned char,
                    short, unsigned short, int, unsigned int, float) {
@@ -54,33 +54,38 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredGrad_Positive_ReadModeElementType", "", cha
 
   TextureTestParams<TestType> params = {};
   params.extent = make_hipExtent(1024, 0, 0);
+  params.layers = 2;
   params.num_subdivisions = 4;
   params.GenerateTextureDesc();
 
-  TextureTestFixture<TestType> fixture{params};
+  TextureTestFixture<TestType, false, true> 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 layer = 0u; layer < params.layers; ++layer) {
+    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, layer, 0.5f, 0.5f);
 
-  for (auto i = 0u; i < params.NumItersX(); ++i) {
-    float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
-                            params.tex_desc.normalizedCoords);
+    fixture.LoadOutput();
 
-    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);
+    for (auto i = 0u; i < params.NumItersX(); ++i) {
+      float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
+                              params.tex_desc.normalizedCoords);
 
-    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);
+      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 = 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);
+    }
   }
 }
 
@@ -99,7 +104,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredGrad_Positive_ReadModeElementType", "", cha
  *    - unit/texture/tex1DLayeredGrad.cc
  * Test requirements
  * ------------------------
- *    - HIP_VERSION >= 5.2
+ *    - HIP_VERSION >= 5.7
  */
 TEMPLATE_TEST_CASE("Unit_tex1DLayeredGrad_Positive_ReadModeNormalizedFloat", "", char,
                    unsigned char, short, unsigned short) {
@@ -107,34 +112,39 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredGrad_Positive_ReadModeNormalizedFloat", "",
 
   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};
+  TextureTestFixture<TestType, true, 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 layer = 0u; layer < params.layers; ++layer) {
+    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, layer, 0.5f, 0.5f);
 
-  for (auto i = 0u; i < params.NumItersX(); ++i) {
-    float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
-                            params.tex_desc.normalizedCoords);
+    fixture.LoadOutput();
 
-    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);
+    for (auto i = 0u; i < params.NumItersX(); ++i) {
+      float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
+                              params.tex_desc.normalizedCoords);
 
-    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);
+      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("Filter mode: " << FilteringModeToString(params.tex_desc.filterMode));
+      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);
+    }
   }
 }
\ No newline at end of file
diff --git a/catch/unit/texture/tex1DLayeredLod.cc b/catch/unit/texture/tex1DLayeredLod.cc
index 874c99c85a..a39d664502 100644
--- a/catch/unit/texture/tex1DLayeredLod.cc
+++ b/catch/unit/texture/tex1DLayeredLod.cc
@@ -46,7 +46,7 @@ THE SOFTWARE.
  *    - unit/texture/tex1DLayeredLod.cc
  * Test requirements
  * ------------------------
- *    - HIP_VERSION >= 5.2
+ *    - HIP_VERSION >= 5.7
  */
 TEMPLATE_TEST_CASE("Unit_tex1DLayeredLod_Positive_ReadModeElementType", "", char, unsigned char,
                    short, unsigned short, int, unsigned int, float) {
@@ -54,33 +54,38 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredLod_Positive_ReadModeElementType", "", char
 
   TextureTestParams<TestType> params = {};
   params.extent = make_hipExtent(1024, 0, 0);
+  params.layers = 2;
   params.num_subdivisions = 4;
   params.GenerateTextureDesc();
 
-  TextureTestFixture<TestType> fixture{params};
+  TextureTestFixture<TestType, false, true> 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 layer = 0u; layer < params.layers; ++layer) {
+    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, layer, 0);
 
-  for (auto i = 0u; i < params.NumItersX(); ++i) {
-    float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
-                            params.tex_desc.normalizedCoords);
+    fixture.LoadOutput();
 
-    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);
+    for (auto i = 0u; i < params.NumItersX(); ++i) {
+      float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
+                              params.tex_desc.normalizedCoords);
 
-    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);
+      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 = 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);
+    }
   }
 }
 
@@ -99,7 +104,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredLod_Positive_ReadModeElementType", "", char
  *    - unit/texture/tex1DLayeredLod.cc
  * Test requirements
  * ------------------------
- *    - HIP_VERSION >= 5.2
+ *    - HIP_VERSION >= 5.7
  */
 TEMPLATE_TEST_CASE("Unit_tex1DLayeredLod_Positive_ReadModeNormalizedFloat", "", char, unsigned char,
                    short, unsigned short) {
@@ -107,34 +112,38 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredLod_Positive_ReadModeNormalizedFloat", "",
 
   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};
+  TextureTestFixture<TestType, true, 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 layer = 0u; layer < params.layers; ++layer) {
+    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, layer, 0);
 
-  for (auto i = 0u; i < params.NumItersX(); ++i) {
-    float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
-                            params.tex_desc.normalizedCoords);
+    fixture.LoadOutput();
 
-    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);
+    for (auto i = 0u; i < params.NumItersX(); ++i) {
+      float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
+                              params.tex_desc.normalizedCoords);
 
-    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);
+      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);
+
+      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);
+    }
   }
 }
\ No newline at end of file
diff --git a/catch/unit/texture/tex1DLod.cc b/catch/unit/texture/tex1DLod.cc
index ceee1211b3..e38ed60745 100644
--- a/catch/unit/texture/tex1DLod.cc
+++ b/catch/unit/texture/tex1DLod.cc
@@ -46,7 +46,7 @@ THE SOFTWARE.
  *    - unit/texture/tex1DLod.cc
  * Test requirements
  * ------------------------
- *    - HIP_VERSION >= 5.2
+ *    - HIP_VERSION >= 5.7
  */
 TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeElementType", "", char, unsigned char, short,
                    unsigned short, int, unsigned int, float) {
@@ -57,7 +57,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeElementType", "", char, unsig
   params.num_subdivisions = 4;
   params.GenerateTextureDesc();
 
-  TextureTestFixture<TestType> fixture{params};
+  TextureTestFixture<TestType, false, true> fixture{params};
 
   const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
   tex1DLodKernel<vec4<TestType>><<<num_blocks, num_threads>>>(
@@ -99,7 +99,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeElementType", "", char, unsig
  *    - unit/texture/tex1DLod.cc
  * Test requirements
  * ------------------------
- *    - HIP_VERSION >= 5.2
+ *    - HIP_VERSION >= 5.7
  */
 TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeNormalizedFloat", "", char, unsigned char, short,
                    unsigned short) {
@@ -110,7 +110,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeNormalizedFloat", "", char, u
   params.num_subdivisions = 4;
   params.GenerateTextureDesc(hipReadModeNormalizedFloat);
 
-  TextureTestFixture<TestType, true> fixture{params};
+  TextureTestFixture<TestType, true, true> fixture{params};
 
   const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
   tex1DLodKernel<vec4<float>><<<num_blocks, num_threads>>>(
@@ -127,7 +127,6 @@ TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeNormalizedFloat", "", char, u
     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 =