From 41844f2202152ff05edb17d6d748da7021e8fee6 Mon Sep 17 00:00:00 2001
From: Dino Music <dino.music@htecgroup.com>
Date: Fri, 30 Sep 2022 02:38:41 -0400
Subject: [PATCH 01/21] Cit -m RAII guards for memory allocations and streams,
 define some commonly useful utility functions and kernels

---
 tests/catch/include/resource_guards.hh | 124 +++++++++++++++++++++++++
 tests/catch/include/utils.hh           |  87 +++++++++++++++++
 2 files changed, 211 insertions(+)
 create mode 100644 tests/catch/include/resource_guards.hh
 create mode 100644 tests/catch/include/utils.hh

diff --git a/tests/catch/include/resource_guards.hh b/tests/catch/include/resource_guards.hh
new file mode 100644
index 0000000000..293fd9d493
--- /dev/null
+++ b/tests/catch/include/resource_guards.hh
@@ -0,0 +1,124 @@
+/*
+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.
+*/
+
+#pragma once
+
+#include <hip_test_common.hh>
+#include <hip/hip_runtime_api.h>
+
+enum class LinearAllocs {
+  malloc,
+  mallocAndRegister,
+  hipHostMalloc,
+  hipMalloc,
+  hipMallocManaged,
+};
+
+template <typename T> class LinearAllocGuard {
+ public:
+  LinearAllocGuard(const LinearAllocs allocation_type, const size_t size,
+                   const unsigned int flags = 0u)
+      : allocation_type_{allocation_type} {
+    switch (allocation_type_) {
+      case LinearAllocs::malloc:
+        ptr_ = host_ptr_ = reinterpret_cast<T*>(malloc(size));
+        break;
+      case LinearAllocs::mallocAndRegister:
+        host_ptr_ = reinterpret_cast<T*>(malloc(size));
+        HIP_CHECK(hipHostRegister(host_ptr_, size, flags));
+        HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&ptr_), host_ptr_, 0u));
+        break;
+      case LinearAllocs::hipHostMalloc:
+        HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&ptr_), size, flags));
+        host_ptr_ = ptr_;
+        break;
+      case LinearAllocs::hipMalloc:
+        HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&ptr_), size));
+        break;
+      case LinearAllocs::hipMallocManaged:
+        HIP_CHECK(hipMallocManaged(reinterpret_cast<void**>(&ptr_), size, flags ? flags : 1u));
+        host_ptr_ = ptr_;
+    }
+  }
+
+  LinearAllocGuard(const LinearAllocGuard&) = delete;
+  LinearAllocGuard(LinearAllocGuard&&) = delete;
+
+  ~LinearAllocGuard() {
+    // No Catch macros, don't want to possibly throw in the destructor
+    switch (allocation_type_) {
+      case LinearAllocs::malloc:
+        free(ptr_);
+        break;
+      case LinearAllocs::mallocAndRegister:
+        hipHostUnregister(host_ptr_);
+        free(host_ptr_);
+        break;
+      case LinearAllocs::hipHostMalloc:
+        hipHostFree(ptr_);
+        break;
+      case LinearAllocs::hipMalloc:
+      case LinearAllocs::hipMallocManaged:
+        hipFree(ptr_);
+    }
+  }
+
+  T* ptr() { return ptr_; };
+  T* const ptr() const { return ptr_; };
+  T* host_ptr() { return host_ptr_; }
+  T* const host_ptr() const { return host_ptr(); }
+
+ private:
+  const LinearAllocs allocation_type_;
+  T* ptr_ = nullptr;
+  T* host_ptr_ = nullptr;
+};
+
+enum class Streams { nullstream, perThread, created };
+
+class StreamGuard {
+ public:
+  StreamGuard(const Streams stream_type) : stream_type_{stream_type} {
+    switch (stream_type_) {
+      case Streams::nullstream:
+        stream_ = nullptr;
+        break;
+      case Streams::perThread:
+        stream_ = hipStreamPerThread;
+        break;
+      case Streams::created:
+        HIP_CHECK(hipStreamCreate(&stream_));
+    }
+  }
+
+  StreamGuard(const StreamGuard&) = delete;
+  StreamGuard(StreamGuard&&) = delete;
+
+  ~StreamGuard() {
+    if (stream_type_ == Streams::created) {
+      hipStreamDestroy(stream_);
+    }
+  }
+
+  hipStream_t stream() const { return stream_; }
+
+ private:
+  const Streams stream_type_;
+  hipStream_t stream_;
+};
\ No newline at end of file
diff --git a/tests/catch/include/utils.hh b/tests/catch/include/utils.hh
new file mode 100644
index 0000000000..614159eda7
--- /dev/null
+++ b/tests/catch/include/utils.hh
@@ -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.
+*/
+
+#pragma once
+
+#include <chrono>
+
+#include <hip_test_common.hh>
+#include <hip/hip_runtime_api.h>
+
+namespace {
+inline constexpr size_t kPageSize = 4096;
+}  // anonymous namespace
+
+template <typename T>
+void MemcpyArrayCompare(T* const expected, T* const actual, const size_t num_elements) {
+  const auto ret = std::mismatch(expected, expected + num_elements, actual);
+  if (ret.first != expected + num_elements) {
+    const auto idx = std::distance(expected, ret.first);
+    INFO("Value mismatch at index: " << idx);
+    REQUIRE(expected[idx] == actual[idx]);
+  }
+}
+
+template <typename T>
+void ArrayFindIfNot(T* const array, const T expected_value, const size_t num_elements) {
+  const auto it = std::find_if_not(array, array + num_elements, [expected_value](const int elem) {
+    return expected_value == elem;
+  });
+
+  if (it != array + num_elements) {
+    const auto idx = std::distance(array, it);
+    INFO("Value mismatch at index " << idx);
+    REQUIRE(expected_value == array[idx]);
+  }
+}
+
+template <typename T>
+__global__ void VectorIncrement(T* const vec, const T increment_value, size_t N) {
+  size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
+  size_t stride = blockDim.x * gridDim.x;
+
+  for (size_t i = offset; i < N; i += stride) {
+    vec[i] += increment_value;
+  }
+}
+
+template <typename T> __global__ void VectorSet(T* const vec, const T value, size_t N) {
+  size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
+  size_t stride = blockDim.x * gridDim.x;
+
+  for (size_t i = offset; i < N; i += stride) {
+    vec[i] = value;
+  }
+}
+
+// Will execute for atleast interval milliseconds
+static __global__ void Delay(uint32_t interval, const uint32_t ticks_per_ms) {
+  while (interval--) {
+    uint64_t start = clock();
+    while (clock() - start < ticks_per_ms) {
+    }
+  }
+}
+
+inline void LaunchDelayKernel(const std::chrono::milliseconds interval, const hipStream_t stream) {
+  int ticks_per_ms = 0;
+  // Clock rate is in kHz => number of clock ticks in a millisecond
+  HIP_CHECK(hipDeviceGetAttribute(&ticks_per_ms, hipDeviceAttributeClockRate, 0));
+  Delay<<<1, 1, 0, stream>>>(interval.count(), ticks_per_ms);
+}
\ No newline at end of file

From 858da0e1ae643b9c8f0347e98ad4ffacd6b5ccbe Mon Sep 17 00:00:00 2001
From: Dino Music <dino.music@htecgroup.com>
Date: Mon, 3 Oct 2022 05:22:42 -0400
Subject: [PATCH 02/21] Implement helper function for generating allocation
 flags

---
 tests/catch/include/resource_guards.hh | 21 ++++++++++++++++++++-
 1 file changed, 20 insertions(+), 1 deletion(-)

diff --git a/tests/catch/include/resource_guards.hh b/tests/catch/include/resource_guards.hh
index 293fd9d493..9f50ea443a 100644
--- a/tests/catch/include/resource_guards.hh
+++ b/tests/catch/include/resource_guards.hh
@@ -121,4 +121,23 @@ class StreamGuard {
  private:
   const Streams stream_type_;
   hipStream_t stream_;
-};
\ No newline at end of file
+};
+
+inline unsigned int GenerateLinearAllocationFlagCombinations(const LinearAllocs allocation_type) {
+  switch (allocation_type) {
+    case LinearAllocs::mallocAndRegister:
+      // TODO
+      return 0;
+    case LinearAllocs::hipHostMalloc:
+      return GENERATE(hipHostMallocDefault, hipHostMallocPortable, hipHostMallocMapped,
+                      hipHostMallocWriteCombined);
+    case LinearAllocs::hipMallocManaged:
+      // TODO
+      return 1u;
+    case LinearAllocs::malloc:
+    case LinearAllocs::hipMalloc:
+      return 0u;
+    default:
+      assert("Invalid LinearAllocs enumerator");
+  }
+}
\ No newline at end of file

From 899b91f978a97f2cd4a082d71d8517bd74461ad9 Mon Sep 17 00:00:00 2001
From: Dino Music <dino.music@htecgroup.com>
Date: Tue, 4 Oct 2022 08:27:42 +0200
Subject: [PATCH 03/21] Implement helper function DeviceAttributesSupport to
 check if a device supports any number of attributes

---
 tests/catch/include/utils.hh | 13 ++++++++++++-
 1 file changed, 12 insertions(+), 1 deletion(-)

diff --git a/tests/catch/include/utils.hh b/tests/catch/include/utils.hh
index 614159eda7..1448c4f768 100644
--- a/tests/catch/include/utils.hh
+++ b/tests/catch/include/utils.hh
@@ -29,7 +29,7 @@ inline constexpr size_t kPageSize = 4096;
 }  // anonymous namespace
 
 template <typename T>
-void MemcpyArrayCompare(T* const expected, T* const actual, const size_t num_elements) {
+void ArrayMismatch(T* const expected, T* const actual, const size_t num_elements) {
   const auto ret = std::mismatch(expected, expected + num_elements, actual);
   if (ret.first != expected + num_elements) {
     const auto idx = std::distance(expected, ret.first);
@@ -84,4 +84,15 @@ inline void LaunchDelayKernel(const std::chrono::milliseconds interval, const hi
   // Clock rate is in kHz => number of clock ticks in a millisecond
   HIP_CHECK(hipDeviceGetAttribute(&ticks_per_ms, hipDeviceAttributeClockRate, 0));
   Delay<<<1, 1, 0, stream>>>(interval.count(), ticks_per_ms);
+}
+
+template <typename... Attributes>
+inline bool DeviceAttributesSupport(const int device, Attributes... attributes) {
+  constexpr auto DeviceAttributeSupport = [](const int device,
+                                             const hipDeviceAttribute_t attribute) {
+    int value = 0;
+    HIP_CHECK(hipDeviceGetAttribute(&value, attribute, device));
+    return value;
+  };
+  return (... && DeviceAttributeSupport(device, attributes));
 }
\ No newline at end of file

From 8accf0c5e8fe9faa22e07c0160b27868637bab24 Mon Sep 17 00:00:00 2001
From: Dino Music <dino.music@htecgroup.com>
Date: Tue, 4 Oct 2022 14:32:09 +0200
Subject: [PATCH 04/21] EXSWHTEC-70 - Reimplement tests for hipMemPrefetchAsync
 - Reimplement Unit_hipMemPrefetchAsync - Reimplement negative tests

---
 .../catch/unit/memory/hipMemPrefetchAsync.cc  | 164 ++++++++----------
 1 file changed, 72 insertions(+), 92 deletions(-)

diff --git a/tests/catch/unit/memory/hipMemPrefetchAsync.cc b/tests/catch/unit/memory/hipMemPrefetchAsync.cc
index 17ef618b77..5f613bd826 100644
--- a/tests/catch/unit/memory/hipMemPrefetchAsync.cc
+++ b/tests/catch/unit/memory/hipMemPrefetchAsync.cc
@@ -1,13 +1,15 @@
 /*
-Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
+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
@@ -17,9 +19,27 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 
+#include <vector>
+
 #include <hip_test_common.hh>
-// Kernel function
-__global__ void MemPrftchAsyncKernel(int* C_d, const int* A_d, size_t N) {
+#include <hip/hip_runtime_api.h>
+#include <utils.hh>
+#include <resource_guards.hh>
+
+std::vector<int> GetDevicesWithPrefetchSupport() {
+  const auto device_count = HipTest::getDeviceCount();
+  std::vector<int> supported_devices;
+  supported_devices.reserve(device_count + 1);
+  for (int i = 0; i < device_count; ++i) {
+    if (DeviceAttributesSupport(i, hipDeviceAttributeManagedMemory,
+                                hipDeviceAttributeConcurrentManagedAccess)) {
+      supported_devices.push_back(i);
+    }
+  }
+  return supported_devices;
+}
+
+__global__ void MemPrefetchAsyncKernel(int* C_d, const int* A_d, size_t N) {
   size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
   size_t stride = blockDim.x * gridDim.x;
   for (size_t i = offset; i < N; i += stride) {
@@ -27,98 +47,58 @@ __global__ void MemPrftchAsyncKernel(int* C_d, const int* A_d, size_t N) {
   }
 }
 
+TEST_CASE("Unit_hipMemPrefetchAsync_Basic") {
+  const auto supported_devices = GetDevicesWithPrefetchSupport();
+  if (supported_devices.empty()) {
+    HipTest::HIP_SKIP_TEST("Test need at least one device with managed memory support");
+  }
 
-static int HmmAttrPrint() {
-  int managed = 0;
-  INFO("The following are the attribute values related to HMM for"
-         " device 0:\n");
-  HIP_CHECK(hipDeviceGetAttribute(&managed,
-              hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
-  INFO("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
-  HIP_CHECK(hipDeviceGetAttribute(&managed,
-                                 hipDeviceAttributeConcurrentManagedAccess, 0));
-  INFO("hipDeviceAttributeConcurrentManagedAccess: " << managed);
-  HIP_CHECK(hipDeviceGetAttribute(&managed,
-                                 hipDeviceAttributePageableMemoryAccess, 0));
-  INFO("hipDeviceAttributePageableMemoryAccess: " << managed);
-  HIP_CHECK(hipDeviceGetAttribute(&managed,
-              hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
-  INFO("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:"
-         << managed);
-
-  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
-                                  0));
-  INFO("hipDeviceAttributeManagedMemory: " << managed);
-  return managed;
-}
+  LinearAllocGuard<int> alloc1(LinearAllocs::hipMallocManaged, kPageSize);
+  const auto count = kPageSize / sizeof(*alloc1.ptr());
+  constexpr auto fill_value = 42;
+  std::fill_n(alloc1.ptr(), count, fill_value);
 
-/*
-  Test Description: This test prefetches the memory to each of the available
-  devices and launch kernel followed by result verification
-  At the end the memory is prefetched to Host and kernel is launched followed
-  by result verification.
-*/
-
-TEST_CASE("Unit_hipMemPrefetchAsync") {
-  int MangdMem = HmmAttrPrint();
-  if (MangdMem == 1) {
-    bool IfTestPassed = true;
-    int A_CONST = 123, MEM_SIZE = (8192 * sizeof(int));
-    int *devPtr1 = NULL, *devPtr2 = NULL, NumDevs = 0, flag = 0;
-    hipStream_t strm;
-    HIP_CHECK(hipMallocManaged(&devPtr1, MEM_SIZE));
-    HIP_CHECK(hipMallocManaged(&devPtr2, MEM_SIZE));
-    HIP_CHECK(hipGetDeviceCount(&NumDevs));
-    // Initializing the memory
-    for (uint32_t k = 0; k < (MEM_SIZE/sizeof(int)); ++k) {
-      devPtr1[k] = A_CONST;
-      devPtr2[k] = 0;
-    }
+  for (const auto device : supported_devices) {
+    HIP_CHECK(hipSetDevice(device));
+    LinearAllocGuard<int> alloc2(LinearAllocs::hipMallocManaged, kPageSize);
+    StreamGuard sg(Streams::created);
+    HIP_CHECK(hipMemPrefetchAsync(alloc1.ptr(), kPageSize, device, sg.stream()));
+    MemPrefetchAsyncKernel<<<count / 1024 + 1, 1024, 0, sg.stream()>>>(alloc2.ptr(), alloc1.ptr(),
+                                                                       count);
+    HIP_CHECK(hipStreamSynchronize(sg.stream()));
+    ArrayFindIfNot(alloc1.ptr(), fill_value, count);
+    ArrayFindIfNot(alloc2.ptr(), fill_value * fill_value, count);
+  }
 
+  HIP_CHECK(hipMemPrefetchAsync(alloc1.ptr(), kPageSize, hipCpuDeviceId));
+  HIP_CHECK(hipStreamSynchronize(nullptr));
+  ArrayFindIfNot(alloc1.ptr(), fill_value, count);
+}
 
-    for (int i = 0; i < NumDevs;  ++i) {
-      HIP_CHECK(hipSetDevice(i));
-      HIP_CHECK(hipStreamCreate(&strm));
-      HIP_CHECK(hipMemPrefetchAsync(devPtr1, MEM_SIZE, i, strm));
-      HIP_CHECK(hipStreamSynchronize(strm));
-      MemPrftchAsyncKernel<<<32, (MEM_SIZE/sizeof(int)/32)>>>(devPtr2, devPtr1,
-                                                       MEM_SIZE/sizeof(int));
-      for (uint32_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) {
-        if (devPtr1[m] != (A_CONST * A_CONST)) {
-          flag = 1;
-        }
-      }
-      HIP_CHECK(hipStreamDestroy(strm));
-      if (!flag) {
-        INFO("Test failed for device: " << i);
-        IfTestPassed = false;
-        flag = 0;
-      }
-    }
-    // The memory will be prefetched from last gpu in the system to the host
-    // memory and kernel is launched followed by result verification.
-    HIP_CHECK(hipStreamCreate(&strm));
-    HIP_CHECK(hipMemPrefetchAsync(devPtr1, MEM_SIZE, hipCpuDeviceId, strm));
-    HIP_CHECK(hipStreamSynchronize(strm));
-    MemPrftchAsyncKernel<<<32, (MEM_SIZE/sizeof(int)/32)>>>(devPtr2, devPtr1,
-                                                     MEM_SIZE/sizeof(int));
-    for (uint32_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) {
-      if (devPtr1[m] != (A_CONST * A_CONST)) {
-        flag = 1;
-      }
-    }
-    HIP_CHECK(hipStreamDestroy(strm));
-    if (!flag) {
-      INFO("Failed to prefetch the memory to System space.\n");
-      IfTestPassed = false;
-      flag = 0;
-    }
+TEST_CASE("Unit_hipMemPrefetchAsync_Negative_Parameters") {
+  auto supported_devices = GetDevicesWithPrefetchSupport();
+  if (supported_devices.empty()) {
+    HipTest::HIP_SKIP_TEST("Test need at least one device with managed memory support");
+  }
+  supported_devices.push_back(hipCpuDeviceId);
+  const auto device = GENERATE_COPY(from_range(supported_devices));
 
-    HIP_CHECK(hipFree(devPtr1));
-    HIP_CHECK(hipFree(devPtr2));
-    REQUIRE(IfTestPassed);
-  } else {
-    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
-           "attribute. Hence skipping the testing with Pass result.\n");
+  LinearAllocGuard<void> alloc(LinearAllocs::hipMallocManaged, kPageSize);
+  SECTION("count == 0") {
+    HIP_CHECK_ERROR(hipMemPrefetchAsync(alloc.ptr(), 0, device), hipErrorInvalidValue);
   }
-}
+  SECTION("count larger than allocation size") {
+    HIP_CHECK_ERROR(hipMemPrefetchAsync(alloc.ptr(), kPageSize + 1, device), hipErrorInvalidValue);
+  }
+  SECTION("Invalid device") {
+    HIP_CHECK_ERROR(hipMemPrefetchAsync(alloc.ptr(), kPageSize, hipInvalidDeviceId),
+                    hipErrorInvalidDevice);
+  }
+  SECTION("Invalid stream") {
+    hipStream_t stream;
+    HIP_CHECK(hipStreamCreate(&stream));
+    HIP_CHECK(hipStreamDestroy(stream));
+    HIP_CHECK_ERROR(hipMemPrefetchAsync(alloc.ptr(), kPageSize, device, stream),
+                    hipErrorContextIsDestroyed);
+  }
+}
\ No newline at end of file

From 94aa6b4f61c904f03cf31cab4e16469c89a5a1b8 Mon Sep 17 00:00:00 2001
From: Dino Music <dino.music@htecgroup.com>
Date: Tue, 4 Oct 2022 15:59:49 +0200
Subject: [PATCH 05/21] EXSWHTEC-70 - Reimplement tests for hipMemPrefetchAsync
 - Implement additional negative tests - Implement tests for synchronization
 and rounding behavior

---
 tests/catch/unit/memory/CMakeLists.txt        |  3 +-
 .../catch/unit/memory/hipMemPrefetchAsync.cc  | 55 +++++++++++++++++++
 2 files changed, 56 insertions(+), 2 deletions(-)

diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt
index f24c63ad8c..54e7708556 100644
--- a/tests/catch/unit/memory/CMakeLists.txt
+++ b/tests/catch/unit/memory/CMakeLists.txt
@@ -183,5 +183,4 @@ endif()
 
 hip_add_exe_to_target(NAME MemoryTest
                       TEST_SRC ${TEST_SRC}
-                      TEST_TARGET_NAME build_tests
-                      COMPILE_OPTIONS -std=c++14)
+                      TEST_TARGET_NAME build_tests)
diff --git a/tests/catch/unit/memory/hipMemPrefetchAsync.cc b/tests/catch/unit/memory/hipMemPrefetchAsync.cc
index 5f613bd826..c7dfb0f123 100644
--- a/tests/catch/unit/memory/hipMemPrefetchAsync.cc
+++ b/tests/catch/unit/memory/hipMemPrefetchAsync.cc
@@ -65,6 +65,7 @@ TEST_CASE("Unit_hipMemPrefetchAsync_Basic") {
     HIP_CHECK(hipMemPrefetchAsync(alloc1.ptr(), kPageSize, device, sg.stream()));
     MemPrefetchAsyncKernel<<<count / 1024 + 1, 1024, 0, sg.stream()>>>(alloc2.ptr(), alloc1.ptr(),
                                                                        count);
+    HIP_CHECK(hipGetLastError());
     HIP_CHECK(hipStreamSynchronize(sg.stream()));
     ArrayFindIfNot(alloc1.ptr(), fill_value, count);
     ArrayFindIfNot(alloc2.ptr(), fill_value * fill_value, count);
@@ -75,6 +76,53 @@ TEST_CASE("Unit_hipMemPrefetchAsync_Basic") {
   ArrayFindIfNot(alloc1.ptr(), fill_value, count);
 }
 
+TEST_CASE("Unit_hipMemPrefetchAsync_Sync_Behavior") {
+  const auto supported_devices = GetDevicesWithPrefetchSupport();
+  if (supported_devices.empty()) {
+    HipTest::HIP_SKIP_TEST("Test need at least one device with managed memory support");
+  }
+  const auto device = supported_devices.front();
+  const auto stream_type = GENERATE(Streams::nullstream, Streams::perThread, Streams::created);
+
+  StreamGuard sg(stream_type);
+  LinearAllocGuard<void> alloc(LinearAllocs::hipMallocManaged, kPageSize);
+  LaunchDelayKernel(std::chrono::milliseconds{100}, sg.stream());
+  HIP_CHECK(hipMemPrefetchAsync(alloc.ptr(), kPageSize, device, sg.stream()));
+  HIP_CHECK_ERROR(hipStreamQuery(sg.stream()), hipErrorNotReady);
+  HIP_CHECK(hipStreamSynchronize(sg.stream()));
+}
+
+TEST_CASE("Unit_hipMemPrefetchAsync_Rounding_Behavior") {
+  auto supported_devices = GetDevicesWithPrefetchSupport();
+  if (supported_devices.empty()) {
+    HipTest::HIP_SKIP_TEST("Test need at least one device with managed memory support");
+  }
+  const auto device = supported_devices.front();
+  LinearAllocGuard<uint8_t> alloc(LinearAllocs::hipMallocManaged, 3 * kPageSize);
+  REQUIRE_FALSE(reinterpret_cast<intptr_t>(alloc.ptr()) % kPageSize);
+  const auto [offset, width] =
+      GENERATE_COPY(std::make_pair(kPageSize / 4, kPageSize / 2),   // Withing page
+                    std::make_pair(kPageSize / 2, kPageSize),       // Across page border
+                    std::make_pair(kPageSize / 2, kPageSize * 2));  // Across two page borders
+  HIP_CHECK(hipMemPrefetchAsync(alloc.ptr() + offset, width, device));
+  HIP_CHECK(hipStreamSynchronize(nullptr));
+  constexpr auto RoundDown = [](const intptr_t a, const intptr_t n) { return a - a % n; };
+  constexpr auto RoundUp = [RoundDown](const intptr_t a, const intptr_t n) {
+    return RoundDown(a + n - 1, n);
+  };
+  const auto base = alloc.ptr();
+  const auto rounded_up = RoundUp(offset + width, kPageSize);
+  unsigned int attribute = 0;
+  HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute),
+                                    hipMemRangeAttributeLastPrefetchLocation,
+                                    reinterpret_cast<void*>(base), rounded_up));
+  REQUIRE(device == attribute);
+  HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute),
+                                    hipMemRangeAttributeLastPrefetchLocation, alloc.ptr(),
+                                    3 * kPageSize));
+  REQUIRE((rounded_up == 3 * kPageSize ? device : hipInvalidDeviceId) == attribute);
+}
+
 TEST_CASE("Unit_hipMemPrefetchAsync_Negative_Parameters") {
   auto supported_devices = GetDevicesWithPrefetchSupport();
   if (supported_devices.empty()) {
@@ -84,6 +132,13 @@ TEST_CASE("Unit_hipMemPrefetchAsync_Negative_Parameters") {
   const auto device = GENERATE_COPY(from_range(supported_devices));
 
   LinearAllocGuard<void> alloc(LinearAllocs::hipMallocManaged, kPageSize);
+  SECTION("dev_ptr == nullptr") {
+    HIP_CHECK_ERROR(hipMemPrefetchAsync(nullptr, kPageSize, device), hipErrorInvalidValue);
+  }
+  SECTION("dev_ptr points to non-managed memory") {
+    LinearAllocGuard<void> alloc(LinearAllocs::hipMalloc, kPageSize);
+    HIP_CHECK_ERROR(hipMemPrefetchAsync(alloc.ptr(), kPageSize, device), hipErrorInvalidValue);
+  }
   SECTION("count == 0") {
     HIP_CHECK_ERROR(hipMemPrefetchAsync(alloc.ptr(), 0, device), hipErrorInvalidValue);
   }

From e534d2e050446f4e936c0de76fc81b5858782654 Mon Sep 17 00:00:00 2001
From: Dino Music <dino.music@htecgroup.com>
Date: Thu, 6 Oct 2022 15:47:34 +0200
Subject: [PATCH 06/21] EXSWHTEC-94 - Implement helper classes and functions
 for memory tests

---
 tests/catch/include/resource_guards.hh | 144 +++++++++++++++++++++++++
 tests/catch/include/utils.hh           | 102 ++++++++++++++++++
 2 files changed, 246 insertions(+)
 create mode 100644 tests/catch/include/resource_guards.hh
 create mode 100644 tests/catch/include/utils.hh

diff --git a/tests/catch/include/resource_guards.hh b/tests/catch/include/resource_guards.hh
new file mode 100644
index 0000000000..f8d1688312
--- /dev/null
+++ b/tests/catch/include/resource_guards.hh
@@ -0,0 +1,144 @@
+/*
+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.
+*/
+
+#pragma once
+
+#include <hip_test_common.hh>
+#include <hip/hip_runtime_api.h>
+
+enum class LinearAllocs {
+  malloc,
+  mallocAndRegister,
+  hipHostMalloc,
+  hipMalloc,
+  hipMallocManaged,
+};
+
+template <typename T> class LinearAllocGuard {
+ public:
+  LinearAllocGuard(const LinearAllocs allocation_type, const size_t size,
+                   const unsigned int flags = 0u)
+      : allocation_type_{allocation_type} {
+    switch (allocation_type_) {
+      case LinearAllocs::malloc:
+        ptr_ = host_ptr_ = reinterpret_cast<T*>(malloc(size));
+        break;
+      case LinearAllocs::mallocAndRegister:
+        host_ptr_ = reinterpret_cast<T*>(malloc(size));
+        HIP_CHECK(hipHostRegister(host_ptr_, size, flags));
+        HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&ptr_), host_ptr_, 0u));
+        break;
+      case LinearAllocs::hipHostMalloc:
+        HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&ptr_), size, flags));
+        host_ptr_ = ptr_;
+        break;
+      case LinearAllocs::hipMalloc:
+        HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&ptr_), size));
+        break;
+      case LinearAllocs::hipMallocManaged:
+        HIP_CHECK(hipMallocManaged(reinterpret_cast<void**>(&ptr_), size, flags ? flags : 1u));
+        host_ptr_ = ptr_;
+    }
+  }
+
+  LinearAllocGuard(const LinearAllocGuard&) = delete;
+  LinearAllocGuard(LinearAllocGuard&&) = delete;
+
+  ~LinearAllocGuard() {
+    // No Catch macros, don't want to possibly throw in the destructor
+    switch (allocation_type_) {
+      case LinearAllocs::malloc:
+        free(ptr_);
+        break;
+      case LinearAllocs::mallocAndRegister:
+        // Cast to void to suppress nodiscard warnings
+        static_cast<void>(hipHostUnregister(host_ptr_));
+        free(host_ptr_);
+        break;
+      case LinearAllocs::hipHostMalloc:
+        static_cast<void>(hipHostFree(ptr_));
+        break;
+      case LinearAllocs::hipMalloc:
+      case LinearAllocs::hipMallocManaged:
+        static_cast<void>(hipFree(ptr_));
+    }
+  }
+
+  T* ptr() { return ptr_; };
+  T* const ptr() const { return ptr_; };
+  T* host_ptr() { return host_ptr_; }
+  T* const host_ptr() const { return host_ptr(); }
+
+ private:
+  const LinearAllocs allocation_type_;
+  T* ptr_ = nullptr;
+  T* host_ptr_ = nullptr;
+};
+
+enum class Streams { nullstream, perThread, created };
+
+class StreamGuard {
+ public:
+  StreamGuard(const Streams stream_type) : stream_type_{stream_type} {
+    switch (stream_type_) {
+      case Streams::nullstream:
+        stream_ = nullptr;
+        break;
+      case Streams::perThread:
+        stream_ = hipStreamPerThread;
+        break;
+      case Streams::created:
+        HIP_CHECK(hipStreamCreate(&stream_));
+    }
+  }
+
+  StreamGuard(const StreamGuard&) = delete;
+  StreamGuard(StreamGuard&&) = delete;
+
+  ~StreamGuard() {
+    if (stream_type_ == Streams::created) {
+      static_cast<void>(hipStreamDestroy(stream_));
+    }
+  }
+
+  hipStream_t stream() const { return stream_; }
+
+ private:
+  const Streams stream_type_;
+  hipStream_t stream_;
+};
+
+inline unsigned int GenerateLinearAllocationFlagCombinations(const LinearAllocs allocation_type) {
+  switch (allocation_type) {
+    case LinearAllocs::mallocAndRegister:
+      // TODO
+      return 0;
+    case LinearAllocs::hipHostMalloc:
+      return GENERATE(hipHostMallocDefault, hipHostMallocPortable, hipHostMallocMapped,
+                      hipHostMallocWriteCombined);
+    case LinearAllocs::hipMallocManaged:
+      // TODO
+      return 1u;
+    case LinearAllocs::malloc:
+    case LinearAllocs::hipMalloc:
+      return 0u;
+    default:
+      assert("Invalid LinearAllocs enumerator");
+  }
+}
\ No newline at end of file
diff --git a/tests/catch/include/utils.hh b/tests/catch/include/utils.hh
new file mode 100644
index 0000000000..9edffc6f7c
--- /dev/null
+++ b/tests/catch/include/utils.hh
@@ -0,0 +1,102 @@
+/*
+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.
+*/
+
+#pragma once
+
+#include <chrono>
+
+#include <hip_test_common.hh>
+#include <hip/hip_runtime_api.h>
+
+namespace {
+inline constexpr size_t kPageSize = 4096;
+}  // anonymous namespace
+
+template <typename T>
+void ArrayMismatch(T* const expected, T* const actual, const size_t num_elements) {
+  const auto ret = std::mismatch(expected, expected + num_elements, actual);
+  if (ret.first != expected + num_elements) {
+    const auto idx = std::distance(expected, ret.first);
+    INFO("Value mismatch at index: " << idx);
+    REQUIRE(expected[idx] == actual[idx]);
+  }
+}
+
+template <typename It, typename T> void ArrayFindIfNot(It begin, It end, const T expected_value) {
+  const auto it = std::find_if_not(
+      begin, end, [expected_value](const int elem) { return expected_value == elem; });
+
+  if (it != end) {
+    const auto idx = std::distance(begin, it);
+    INFO("Value mismatch at index " << idx);
+    REQUIRE(expected_value == *it);
+  }
+}
+
+template <typename T>
+void ArrayFindIfNot(T* const array, const T expected_value, const size_t num_elements) {
+  ArrayFindIfNot(array, array + num_elements, expected_value);
+}
+
+template <typename T>
+__global__ void VectorIncrement(T* const vec, const T increment_value, size_t N) {
+  size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
+  size_t stride = blockDim.x * gridDim.x;
+
+  for (size_t i = offset; i < N; i += stride) {
+    vec[i] += increment_value;
+  }
+}
+
+template <typename T> __global__ void VectorSet(T* const vec, const T value, size_t N) {
+  size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
+  size_t stride = blockDim.x * gridDim.x;
+
+  for (size_t i = offset; i < N; i += stride) {
+    vec[i] = value;
+  }
+}
+
+// Will execute for atleast interval milliseconds
+static __global__ void Delay(uint32_t interval, const uint32_t ticks_per_ms) {
+  while (interval--) {
+    uint64_t start = clock();
+    while (clock() - start < ticks_per_ms) {
+    }
+  }
+}
+
+inline void LaunchDelayKernel(const std::chrono::milliseconds interval, const hipStream_t stream) {
+  int ticks_per_ms = 0;
+  // Clock rate is in kHz => number of clock ticks in a millisecond
+  HIP_CHECK(hipDeviceGetAttribute(&ticks_per_ms, hipDeviceAttributeClockRate, 0));
+  Delay<<<1, 1, 0, stream>>>(interval.count(), ticks_per_ms);
+  HIP_CHECK(hipGetLastError());
+}
+
+template <typename... Attributes>
+inline bool DeviceAttributesSupport(const int device, Attributes... attributes) {
+  constexpr auto DeviceAttributeSupport = [](const int device,
+                                             const hipDeviceAttribute_t attribute) {
+    int value = 0;
+    HIP_CHECK(hipDeviceGetAttribute(&value, attribute, device));
+    return value;
+  };
+  return (... && DeviceAttributeSupport(device, attributes));
+}
\ No newline at end of file

From 09ce86ac14450d712725a7e10377072c8b9050f0 Mon Sep 17 00:00:00 2001
From: Dino Music <dino.music@htecgroup.com>
Date: Thu, 6 Oct 2022 16:19:21 +0200
Subject: [PATCH 07/21] EXSWHTEC-94 - Remove c++14 standard constraint on
 memory tests

---
 tests/catch/unit/memory/CMakeLists.txt | 3 +--
 1 file changed, 1 insertion(+), 2 deletions(-)

diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt
index f24c63ad8c..54e7708556 100644
--- a/tests/catch/unit/memory/CMakeLists.txt
+++ b/tests/catch/unit/memory/CMakeLists.txt
@@ -183,5 +183,4 @@ endif()
 
 hip_add_exe_to_target(NAME MemoryTest
                       TEST_SRC ${TEST_SRC}
-                      TEST_TARGET_NAME build_tests
-                      COMPILE_OPTIONS -std=c++14)
+                      TEST_TARGET_NAME build_tests)

From 48b337f4363897a27649960dd9a3523190ccac1b Mon Sep 17 00:00:00 2001
From: Dino Music <dino.music@htecgroup.com>
Date: Thu, 6 Oct 2022 12:09:41 -0400
Subject: [PATCH 08/21] EXSWHTEC-94 - Remove
 GenerateLinearAllocationFlagCombinations until finished

---
 tests/catch/include/resource_guards.hh | 21 +--------------------
 1 file changed, 1 insertion(+), 20 deletions(-)

diff --git a/tests/catch/include/resource_guards.hh b/tests/catch/include/resource_guards.hh
index f8d1688312..7e6179c81a 100644
--- a/tests/catch/include/resource_guards.hh
+++ b/tests/catch/include/resource_guards.hh
@@ -122,23 +122,4 @@ class StreamGuard {
  private:
   const Streams stream_type_;
   hipStream_t stream_;
-};
-
-inline unsigned int GenerateLinearAllocationFlagCombinations(const LinearAllocs allocation_type) {
-  switch (allocation_type) {
-    case LinearAllocs::mallocAndRegister:
-      // TODO
-      return 0;
-    case LinearAllocs::hipHostMalloc:
-      return GENERATE(hipHostMallocDefault, hipHostMallocPortable, hipHostMallocMapped,
-                      hipHostMallocWriteCombined);
-    case LinearAllocs::hipMallocManaged:
-      // TODO
-      return 1u;
-    case LinearAllocs::malloc:
-    case LinearAllocs::hipMalloc:
-      return 0u;
-    default:
-      assert("Invalid LinearAllocs enumerator");
-  }
-}
\ No newline at end of file
+};
\ No newline at end of file

From 715cf30e7715bea49d4137b4284556a10df7e9c1 Mon Sep 17 00:00:00 2001
From: Dino Music <dino.music@htecgroup.com>
Date: Thu, 6 Oct 2022 15:47:34 +0200
Subject: [PATCH 09/21] EXSWHTEC-94 - Implement helper classes and functions
 for memory tests

---
 tests/catch/include/resource_guards.hh | 144 +++++++++++++++++++++++++
 tests/catch/include/utils.hh           | 102 ++++++++++++++++++
 2 files changed, 246 insertions(+)
 create mode 100644 tests/catch/include/resource_guards.hh
 create mode 100644 tests/catch/include/utils.hh

diff --git a/tests/catch/include/resource_guards.hh b/tests/catch/include/resource_guards.hh
new file mode 100644
index 0000000000..f8d1688312
--- /dev/null
+++ b/tests/catch/include/resource_guards.hh
@@ -0,0 +1,144 @@
+/*
+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.
+*/
+
+#pragma once
+
+#include <hip_test_common.hh>
+#include <hip/hip_runtime_api.h>
+
+enum class LinearAllocs {
+  malloc,
+  mallocAndRegister,
+  hipHostMalloc,
+  hipMalloc,
+  hipMallocManaged,
+};
+
+template <typename T> class LinearAllocGuard {
+ public:
+  LinearAllocGuard(const LinearAllocs allocation_type, const size_t size,
+                   const unsigned int flags = 0u)
+      : allocation_type_{allocation_type} {
+    switch (allocation_type_) {
+      case LinearAllocs::malloc:
+        ptr_ = host_ptr_ = reinterpret_cast<T*>(malloc(size));
+        break;
+      case LinearAllocs::mallocAndRegister:
+        host_ptr_ = reinterpret_cast<T*>(malloc(size));
+        HIP_CHECK(hipHostRegister(host_ptr_, size, flags));
+        HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&ptr_), host_ptr_, 0u));
+        break;
+      case LinearAllocs::hipHostMalloc:
+        HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&ptr_), size, flags));
+        host_ptr_ = ptr_;
+        break;
+      case LinearAllocs::hipMalloc:
+        HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&ptr_), size));
+        break;
+      case LinearAllocs::hipMallocManaged:
+        HIP_CHECK(hipMallocManaged(reinterpret_cast<void**>(&ptr_), size, flags ? flags : 1u));
+        host_ptr_ = ptr_;
+    }
+  }
+
+  LinearAllocGuard(const LinearAllocGuard&) = delete;
+  LinearAllocGuard(LinearAllocGuard&&) = delete;
+
+  ~LinearAllocGuard() {
+    // No Catch macros, don't want to possibly throw in the destructor
+    switch (allocation_type_) {
+      case LinearAllocs::malloc:
+        free(ptr_);
+        break;
+      case LinearAllocs::mallocAndRegister:
+        // Cast to void to suppress nodiscard warnings
+        static_cast<void>(hipHostUnregister(host_ptr_));
+        free(host_ptr_);
+        break;
+      case LinearAllocs::hipHostMalloc:
+        static_cast<void>(hipHostFree(ptr_));
+        break;
+      case LinearAllocs::hipMalloc:
+      case LinearAllocs::hipMallocManaged:
+        static_cast<void>(hipFree(ptr_));
+    }
+  }
+
+  T* ptr() { return ptr_; };
+  T* const ptr() const { return ptr_; };
+  T* host_ptr() { return host_ptr_; }
+  T* const host_ptr() const { return host_ptr(); }
+
+ private:
+  const LinearAllocs allocation_type_;
+  T* ptr_ = nullptr;
+  T* host_ptr_ = nullptr;
+};
+
+enum class Streams { nullstream, perThread, created };
+
+class StreamGuard {
+ public:
+  StreamGuard(const Streams stream_type) : stream_type_{stream_type} {
+    switch (stream_type_) {
+      case Streams::nullstream:
+        stream_ = nullptr;
+        break;
+      case Streams::perThread:
+        stream_ = hipStreamPerThread;
+        break;
+      case Streams::created:
+        HIP_CHECK(hipStreamCreate(&stream_));
+    }
+  }
+
+  StreamGuard(const StreamGuard&) = delete;
+  StreamGuard(StreamGuard&&) = delete;
+
+  ~StreamGuard() {
+    if (stream_type_ == Streams::created) {
+      static_cast<void>(hipStreamDestroy(stream_));
+    }
+  }
+
+  hipStream_t stream() const { return stream_; }
+
+ private:
+  const Streams stream_type_;
+  hipStream_t stream_;
+};
+
+inline unsigned int GenerateLinearAllocationFlagCombinations(const LinearAllocs allocation_type) {
+  switch (allocation_type) {
+    case LinearAllocs::mallocAndRegister:
+      // TODO
+      return 0;
+    case LinearAllocs::hipHostMalloc:
+      return GENERATE(hipHostMallocDefault, hipHostMallocPortable, hipHostMallocMapped,
+                      hipHostMallocWriteCombined);
+    case LinearAllocs::hipMallocManaged:
+      // TODO
+      return 1u;
+    case LinearAllocs::malloc:
+    case LinearAllocs::hipMalloc:
+      return 0u;
+    default:
+      assert("Invalid LinearAllocs enumerator");
+  }
+}
\ No newline at end of file
diff --git a/tests/catch/include/utils.hh b/tests/catch/include/utils.hh
new file mode 100644
index 0000000000..9edffc6f7c
--- /dev/null
+++ b/tests/catch/include/utils.hh
@@ -0,0 +1,102 @@
+/*
+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.
+*/
+
+#pragma once
+
+#include <chrono>
+
+#include <hip_test_common.hh>
+#include <hip/hip_runtime_api.h>
+
+namespace {
+inline constexpr size_t kPageSize = 4096;
+}  // anonymous namespace
+
+template <typename T>
+void ArrayMismatch(T* const expected, T* const actual, const size_t num_elements) {
+  const auto ret = std::mismatch(expected, expected + num_elements, actual);
+  if (ret.first != expected + num_elements) {
+    const auto idx = std::distance(expected, ret.first);
+    INFO("Value mismatch at index: " << idx);
+    REQUIRE(expected[idx] == actual[idx]);
+  }
+}
+
+template <typename It, typename T> void ArrayFindIfNot(It begin, It end, const T expected_value) {
+  const auto it = std::find_if_not(
+      begin, end, [expected_value](const int elem) { return expected_value == elem; });
+
+  if (it != end) {
+    const auto idx = std::distance(begin, it);
+    INFO("Value mismatch at index " << idx);
+    REQUIRE(expected_value == *it);
+  }
+}
+
+template <typename T>
+void ArrayFindIfNot(T* const array, const T expected_value, const size_t num_elements) {
+  ArrayFindIfNot(array, array + num_elements, expected_value);
+}
+
+template <typename T>
+__global__ void VectorIncrement(T* const vec, const T increment_value, size_t N) {
+  size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
+  size_t stride = blockDim.x * gridDim.x;
+
+  for (size_t i = offset; i < N; i += stride) {
+    vec[i] += increment_value;
+  }
+}
+
+template <typename T> __global__ void VectorSet(T* const vec, const T value, size_t N) {
+  size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
+  size_t stride = blockDim.x * gridDim.x;
+
+  for (size_t i = offset; i < N; i += stride) {
+    vec[i] = value;
+  }
+}
+
+// Will execute for atleast interval milliseconds
+static __global__ void Delay(uint32_t interval, const uint32_t ticks_per_ms) {
+  while (interval--) {
+    uint64_t start = clock();
+    while (clock() - start < ticks_per_ms) {
+    }
+  }
+}
+
+inline void LaunchDelayKernel(const std::chrono::milliseconds interval, const hipStream_t stream) {
+  int ticks_per_ms = 0;
+  // Clock rate is in kHz => number of clock ticks in a millisecond
+  HIP_CHECK(hipDeviceGetAttribute(&ticks_per_ms, hipDeviceAttributeClockRate, 0));
+  Delay<<<1, 1, 0, stream>>>(interval.count(), ticks_per_ms);
+  HIP_CHECK(hipGetLastError());
+}
+
+template <typename... Attributes>
+inline bool DeviceAttributesSupport(const int device, Attributes... attributes) {
+  constexpr auto DeviceAttributeSupport = [](const int device,
+                                             const hipDeviceAttribute_t attribute) {
+    int value = 0;
+    HIP_CHECK(hipDeviceGetAttribute(&value, attribute, device));
+    return value;
+  };
+  return (... && DeviceAttributeSupport(device, attributes));
+}
\ No newline at end of file

From a74fe2197e7d11e477ac7a4a35a4adc04c52a19f Mon Sep 17 00:00:00 2001
From: Dino Music <dino.music@htecgroup.com>
Date: Thu, 6 Oct 2022 16:19:21 +0200
Subject: [PATCH 10/21] EXSWHTEC-94 - Remove c++14 standard constraint on
 memory tests

---
 tests/catch/unit/memory/CMakeLists.txt | 3 +--
 1 file changed, 1 insertion(+), 2 deletions(-)

diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt
index f24c63ad8c..54e7708556 100644
--- a/tests/catch/unit/memory/CMakeLists.txt
+++ b/tests/catch/unit/memory/CMakeLists.txt
@@ -183,5 +183,4 @@ endif()
 
 hip_add_exe_to_target(NAME MemoryTest
                       TEST_SRC ${TEST_SRC}
-                      TEST_TARGET_NAME build_tests
-                      COMPILE_OPTIONS -std=c++14)
+                      TEST_TARGET_NAME build_tests)

From 350958e5e6bfa2efa722fefcbff09ac0e4e35f9a Mon Sep 17 00:00:00 2001
From: Dino Music <dino.music@htecgroup.com>
Date: Thu, 6 Oct 2022 12:09:41 -0400
Subject: [PATCH 11/21] EXSWHTEC-94 - Remove
 GenerateLinearAllocationFlagCombinations until finished

---
 tests/catch/include/resource_guards.hh | 21 +--------------------
 1 file changed, 1 insertion(+), 20 deletions(-)

diff --git a/tests/catch/include/resource_guards.hh b/tests/catch/include/resource_guards.hh
index f8d1688312..7e6179c81a 100644
--- a/tests/catch/include/resource_guards.hh
+++ b/tests/catch/include/resource_guards.hh
@@ -122,23 +122,4 @@ class StreamGuard {
  private:
   const Streams stream_type_;
   hipStream_t stream_;
-};
-
-inline unsigned int GenerateLinearAllocationFlagCombinations(const LinearAllocs allocation_type) {
-  switch (allocation_type) {
-    case LinearAllocs::mallocAndRegister:
-      // TODO
-      return 0;
-    case LinearAllocs::hipHostMalloc:
-      return GENERATE(hipHostMallocDefault, hipHostMallocPortable, hipHostMallocMapped,
-                      hipHostMallocWriteCombined);
-    case LinearAllocs::hipMallocManaged:
-      // TODO
-      return 1u;
-    case LinearAllocs::malloc:
-    case LinearAllocs::hipMalloc:
-      return 0u;
-    default:
-      assert("Invalid LinearAllocs enumerator");
-  }
-}
\ No newline at end of file
+};
\ No newline at end of file

From 691d00ed3c7d4c7f34cbe1ea77bb28be34fd7239 Mon Sep 17 00:00:00 2001
From: Dino Music <dino.music@htecgroup.com>
Date: Thu, 6 Oct 2022 15:47:34 +0200
Subject: [PATCH 12/21] EXSWHTEC-94 - Implement helper classes and functions
 for memory tests

---
 tests/catch/include/resource_guards.hh | 144 +++++++++++++++++++++++++
 tests/catch/include/utils.hh           | 102 ++++++++++++++++++
 2 files changed, 246 insertions(+)
 create mode 100644 tests/catch/include/resource_guards.hh
 create mode 100644 tests/catch/include/utils.hh

diff --git a/tests/catch/include/resource_guards.hh b/tests/catch/include/resource_guards.hh
new file mode 100644
index 0000000000..f8d1688312
--- /dev/null
+++ b/tests/catch/include/resource_guards.hh
@@ -0,0 +1,144 @@
+/*
+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.
+*/
+
+#pragma once
+
+#include <hip_test_common.hh>
+#include <hip/hip_runtime_api.h>
+
+enum class LinearAllocs {
+  malloc,
+  mallocAndRegister,
+  hipHostMalloc,
+  hipMalloc,
+  hipMallocManaged,
+};
+
+template <typename T> class LinearAllocGuard {
+ public:
+  LinearAllocGuard(const LinearAllocs allocation_type, const size_t size,
+                   const unsigned int flags = 0u)
+      : allocation_type_{allocation_type} {
+    switch (allocation_type_) {
+      case LinearAllocs::malloc:
+        ptr_ = host_ptr_ = reinterpret_cast<T*>(malloc(size));
+        break;
+      case LinearAllocs::mallocAndRegister:
+        host_ptr_ = reinterpret_cast<T*>(malloc(size));
+        HIP_CHECK(hipHostRegister(host_ptr_, size, flags));
+        HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&ptr_), host_ptr_, 0u));
+        break;
+      case LinearAllocs::hipHostMalloc:
+        HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&ptr_), size, flags));
+        host_ptr_ = ptr_;
+        break;
+      case LinearAllocs::hipMalloc:
+        HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&ptr_), size));
+        break;
+      case LinearAllocs::hipMallocManaged:
+        HIP_CHECK(hipMallocManaged(reinterpret_cast<void**>(&ptr_), size, flags ? flags : 1u));
+        host_ptr_ = ptr_;
+    }
+  }
+
+  LinearAllocGuard(const LinearAllocGuard&) = delete;
+  LinearAllocGuard(LinearAllocGuard&&) = delete;
+
+  ~LinearAllocGuard() {
+    // No Catch macros, don't want to possibly throw in the destructor
+    switch (allocation_type_) {
+      case LinearAllocs::malloc:
+        free(ptr_);
+        break;
+      case LinearAllocs::mallocAndRegister:
+        // Cast to void to suppress nodiscard warnings
+        static_cast<void>(hipHostUnregister(host_ptr_));
+        free(host_ptr_);
+        break;
+      case LinearAllocs::hipHostMalloc:
+        static_cast<void>(hipHostFree(ptr_));
+        break;
+      case LinearAllocs::hipMalloc:
+      case LinearAllocs::hipMallocManaged:
+        static_cast<void>(hipFree(ptr_));
+    }
+  }
+
+  T* ptr() { return ptr_; };
+  T* const ptr() const { return ptr_; };
+  T* host_ptr() { return host_ptr_; }
+  T* const host_ptr() const { return host_ptr(); }
+
+ private:
+  const LinearAllocs allocation_type_;
+  T* ptr_ = nullptr;
+  T* host_ptr_ = nullptr;
+};
+
+enum class Streams { nullstream, perThread, created };
+
+class StreamGuard {
+ public:
+  StreamGuard(const Streams stream_type) : stream_type_{stream_type} {
+    switch (stream_type_) {
+      case Streams::nullstream:
+        stream_ = nullptr;
+        break;
+      case Streams::perThread:
+        stream_ = hipStreamPerThread;
+        break;
+      case Streams::created:
+        HIP_CHECK(hipStreamCreate(&stream_));
+    }
+  }
+
+  StreamGuard(const StreamGuard&) = delete;
+  StreamGuard(StreamGuard&&) = delete;
+
+  ~StreamGuard() {
+    if (stream_type_ == Streams::created) {
+      static_cast<void>(hipStreamDestroy(stream_));
+    }
+  }
+
+  hipStream_t stream() const { return stream_; }
+
+ private:
+  const Streams stream_type_;
+  hipStream_t stream_;
+};
+
+inline unsigned int GenerateLinearAllocationFlagCombinations(const LinearAllocs allocation_type) {
+  switch (allocation_type) {
+    case LinearAllocs::mallocAndRegister:
+      // TODO
+      return 0;
+    case LinearAllocs::hipHostMalloc:
+      return GENERATE(hipHostMallocDefault, hipHostMallocPortable, hipHostMallocMapped,
+                      hipHostMallocWriteCombined);
+    case LinearAllocs::hipMallocManaged:
+      // TODO
+      return 1u;
+    case LinearAllocs::malloc:
+    case LinearAllocs::hipMalloc:
+      return 0u;
+    default:
+      assert("Invalid LinearAllocs enumerator");
+  }
+}
\ No newline at end of file
diff --git a/tests/catch/include/utils.hh b/tests/catch/include/utils.hh
new file mode 100644
index 0000000000..9edffc6f7c
--- /dev/null
+++ b/tests/catch/include/utils.hh
@@ -0,0 +1,102 @@
+/*
+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.
+*/
+
+#pragma once
+
+#include <chrono>
+
+#include <hip_test_common.hh>
+#include <hip/hip_runtime_api.h>
+
+namespace {
+inline constexpr size_t kPageSize = 4096;
+}  // anonymous namespace
+
+template <typename T>
+void ArrayMismatch(T* const expected, T* const actual, const size_t num_elements) {
+  const auto ret = std::mismatch(expected, expected + num_elements, actual);
+  if (ret.first != expected + num_elements) {
+    const auto idx = std::distance(expected, ret.first);
+    INFO("Value mismatch at index: " << idx);
+    REQUIRE(expected[idx] == actual[idx]);
+  }
+}
+
+template <typename It, typename T> void ArrayFindIfNot(It begin, It end, const T expected_value) {
+  const auto it = std::find_if_not(
+      begin, end, [expected_value](const int elem) { return expected_value == elem; });
+
+  if (it != end) {
+    const auto idx = std::distance(begin, it);
+    INFO("Value mismatch at index " << idx);
+    REQUIRE(expected_value == *it);
+  }
+}
+
+template <typename T>
+void ArrayFindIfNot(T* const array, const T expected_value, const size_t num_elements) {
+  ArrayFindIfNot(array, array + num_elements, expected_value);
+}
+
+template <typename T>
+__global__ void VectorIncrement(T* const vec, const T increment_value, size_t N) {
+  size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
+  size_t stride = blockDim.x * gridDim.x;
+
+  for (size_t i = offset; i < N; i += stride) {
+    vec[i] += increment_value;
+  }
+}
+
+template <typename T> __global__ void VectorSet(T* const vec, const T value, size_t N) {
+  size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
+  size_t stride = blockDim.x * gridDim.x;
+
+  for (size_t i = offset; i < N; i += stride) {
+    vec[i] = value;
+  }
+}
+
+// Will execute for atleast interval milliseconds
+static __global__ void Delay(uint32_t interval, const uint32_t ticks_per_ms) {
+  while (interval--) {
+    uint64_t start = clock();
+    while (clock() - start < ticks_per_ms) {
+    }
+  }
+}
+
+inline void LaunchDelayKernel(const std::chrono::milliseconds interval, const hipStream_t stream) {
+  int ticks_per_ms = 0;
+  // Clock rate is in kHz => number of clock ticks in a millisecond
+  HIP_CHECK(hipDeviceGetAttribute(&ticks_per_ms, hipDeviceAttributeClockRate, 0));
+  Delay<<<1, 1, 0, stream>>>(interval.count(), ticks_per_ms);
+  HIP_CHECK(hipGetLastError());
+}
+
+template <typename... Attributes>
+inline bool DeviceAttributesSupport(const int device, Attributes... attributes) {
+  constexpr auto DeviceAttributeSupport = [](const int device,
+                                             const hipDeviceAttribute_t attribute) {
+    int value = 0;
+    HIP_CHECK(hipDeviceGetAttribute(&value, attribute, device));
+    return value;
+  };
+  return (... && DeviceAttributeSupport(device, attributes));
+}
\ No newline at end of file

From 7bdf52f994ff486405fc72012cf98a5963a90442 Mon Sep 17 00:00:00 2001
From: Dino Music <dino.music@htecgroup.com>
Date: Thu, 6 Oct 2022 16:19:21 +0200
Subject: [PATCH 13/21] EXSWHTEC-94 - Remove c++14 standard constraint on
 memory tests

---
 tests/catch/unit/memory/CMakeLists.txt | 3 +--
 1 file changed, 1 insertion(+), 2 deletions(-)

diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt
index f24c63ad8c..54e7708556 100644
--- a/tests/catch/unit/memory/CMakeLists.txt
+++ b/tests/catch/unit/memory/CMakeLists.txt
@@ -183,5 +183,4 @@ endif()
 
 hip_add_exe_to_target(NAME MemoryTest
                       TEST_SRC ${TEST_SRC}
-                      TEST_TARGET_NAME build_tests
-                      COMPILE_OPTIONS -std=c++14)
+                      TEST_TARGET_NAME build_tests)

From 1185c3973331e76e1f5d23e7864b4bff89a4c922 Mon Sep 17 00:00:00 2001
From: Dino Music <dino.music@htecgroup.com>
Date: Thu, 6 Oct 2022 12:09:41 -0400
Subject: [PATCH 14/21] EXSWHTEC-94 - Remove
 GenerateLinearAllocationFlagCombinations until finished

---
 tests/catch/include/resource_guards.hh | 21 +--------------------
 1 file changed, 1 insertion(+), 20 deletions(-)

diff --git a/tests/catch/include/resource_guards.hh b/tests/catch/include/resource_guards.hh
index f8d1688312..7e6179c81a 100644
--- a/tests/catch/include/resource_guards.hh
+++ b/tests/catch/include/resource_guards.hh
@@ -122,23 +122,4 @@ class StreamGuard {
  private:
   const Streams stream_type_;
   hipStream_t stream_;
-};
-
-inline unsigned int GenerateLinearAllocationFlagCombinations(const LinearAllocs allocation_type) {
-  switch (allocation_type) {
-    case LinearAllocs::mallocAndRegister:
-      // TODO
-      return 0;
-    case LinearAllocs::hipHostMalloc:
-      return GENERATE(hipHostMallocDefault, hipHostMallocPortable, hipHostMallocMapped,
-                      hipHostMallocWriteCombined);
-    case LinearAllocs::hipMallocManaged:
-      // TODO
-      return 1u;
-    case LinearAllocs::malloc:
-    case LinearAllocs::hipMalloc:
-      return 0u;
-    default:
-      assert("Invalid LinearAllocs enumerator");
-  }
-}
\ No newline at end of file
+};
\ No newline at end of file

From 8911eb7cd62641f50c1555f19118f764ff3ae7bf Mon Sep 17 00:00:00 2001
From: Mirza Halilcevic <mirza.halilcevic@htecgroup.com>
Date: Wed, 12 Oct 2022 10:25:11 +0200
Subject: [PATCH 15/21] EXSWHTEC-94 - Implement resource guards for
 hipMallocPitch and 3D allocations.

---
 tests/catch/include/resource_guards.hh | 55 ++++++++++++++++++++++++++
 1 file changed, 55 insertions(+)

diff --git a/tests/catch/include/resource_guards.hh b/tests/catch/include/resource_guards.hh
index 7e6179c81a..0db1276f15 100644
--- a/tests/catch/include/resource_guards.hh
+++ b/tests/catch/include/resource_guards.hh
@@ -91,6 +91,61 @@ template <typename T> class LinearAllocGuard {
   T* host_ptr_ = nullptr;
 };
 
+template <typename T> class LinearAllocGuardMultiDim {
+  protected:
+  LinearAllocGuardMultiDim(hipExtent extent)
+  : extent_{extent} {}
+
+  ~LinearAllocGuardMultiDim() {
+    static_cast<void>(hipFree(pitched_ptr_.ptr));
+  }
+  
+  public:
+  T* ptr() const { return reinterpret_cast<T*>(pitched_ptr_.ptr); };
+
+  size_t pitch() const { return pitched_ptr_.pitch; }
+
+  hipExtent extent() const { return extent_; }
+
+  hipPitchedPtr pitched_ptr() const { return pitched_ptr_; }
+
+  size_t width() const { return extent_.width; }
+
+  size_t width_logical() const { return extent_.width / sizeof(T); }
+
+  size_t height() const { return extent_.height; }
+
+  public:
+  hipPitchedPtr pitched_ptr_;
+  const hipExtent extent_;
+};
+
+template <typename T> class LinearAllocGuard2D : public LinearAllocGuardMultiDim<T> {
+  public:
+  LinearAllocGuard2D(const size_t width_logical, const size_t height) 
+  : LinearAllocGuardMultiDim<T>{make_hipExtent(width_logical * sizeof(T), height, 1)}
+  {
+    HIP_CHECK(hipMallocPitch(&this->pitched_ptr_.ptr, &this->pitched_ptr_.pitch, this->extent_.width, this->extent_.height));
+  }
+
+  LinearAllocGuard2D(const LinearAllocGuard2D&) = delete;
+  LinearAllocGuard2D(LinearAllocGuard2D&&) = delete;
+};
+
+template <typename T> class LinearAllocGuard3D : public LinearAllocGuardMultiDim<T> {
+  public:
+  LinearAllocGuard3D(const size_t width_logical, const size_t height, const size_t depth)
+  : LinearAllocGuardMultiDim<T>{make_hipExtent(width_logical * sizeof(T), height, depth)}
+  {
+    HIP_CHECK(hipMalloc3D(&this->pitched_ptr_, this->extent_));
+  }
+
+  LinearAllocGuard3D(const LinearAllocGuard3D&) = delete;
+  LinearAllocGuard3D(LinearAllocGuard3D&&) = delete;
+
+  size_t depth() const { return this->extent_.depth; }
+};
+
 enum class Streams { nullstream, perThread, created };
 
 class StreamGuard {

From 76c8e3104c5881e6469b9229d2226011e175e4b7 Mon Sep 17 00:00:00 2001
From: Dino Music <dino.music@htecgroup.com>
Date: Fri, 14 Oct 2022 19:47:44 +0200
Subject: [PATCH 16/21] EXSWHTEC-94 - Add resource guards for 2D and 3D
 allocations and utils for handling pitched memory

---
 tests/catch/include/resource_guards.hh | 10 +++---
 tests/catch/include/utils.hh           | 43 ++++++++++++++++++++++++++
 2 files changed, 49 insertions(+), 4 deletions(-)

diff --git a/tests/catch/include/resource_guards.hh b/tests/catch/include/resource_guards.hh
index 0db1276f15..b3ef7813f7 100644
--- a/tests/catch/include/resource_guards.hh
+++ b/tests/catch/include/resource_guards.hh
@@ -80,10 +80,8 @@ template <typename T> class LinearAllocGuard {
     }
   }
 
-  T* ptr() { return ptr_; };
-  T* const ptr() const { return ptr_; };
-  T* host_ptr() { return host_ptr_; }
-  T* const host_ptr() const { return host_ptr(); }
+  T* ptr() const { return ptr_; };
+  T* host_ptr() const { return host_ptr_; }
 
  private:
   const LinearAllocs allocation_type_;
@@ -140,6 +138,10 @@ template <typename T> class LinearAllocGuard3D : public LinearAllocGuardMultiDim
     HIP_CHECK(hipMalloc3D(&this->pitched_ptr_, this->extent_));
   }
 
+  LinearAllocGuard3D(const hipExtent extent) : LinearAllocGuardMultiDim<T>(extent) {
+    HIP_CHECK(hipMalloc3D(&this->pitched_ptr_, this->extent_));
+  }
+
   LinearAllocGuard3D(const LinearAllocGuard3D&) = delete;
   LinearAllocGuard3D(LinearAllocGuard3D&&) = delete;
 
diff --git a/tests/catch/include/utils.hh b/tests/catch/include/utils.hh
index 9edffc6f7c..05eecea79f 100644
--- a/tests/catch/include/utils.hh
+++ b/tests/catch/include/utils.hh
@@ -54,6 +54,37 @@ void ArrayFindIfNot(T* const array, const T expected_value, const size_t num_ele
   ArrayFindIfNot(array, array + num_elements, expected_value);
 }
 
+template <typename T, typename F>
+void PitchedMemoryVerify(T* const ptr, const size_t pitch, const size_t width, const size_t height,
+                         const size_t depth, F expected_value_generator) {
+  for (int z = 0; z < depth; ++z) {
+    for (int y = 0; y < height; ++y) {
+      for (int x = 0; x < width; ++x) {
+        const auto slice = reinterpret_cast<uint8_t*>(ptr) + pitch * height * z;
+        const auto row = slice + pitch * y;
+        if (reinterpret_cast<T*>(row)[x] != expected_value_generator(x, y, z)) {
+          INFO("Mismatch at indices: " << x << ", " << y << ", " << z);
+          REQUIRE(reinterpret_cast<T*>(row)[x] == expected_value_generator(x, y, z));
+        }
+      }
+    }
+  }
+}
+
+template <typename T, typename F>
+void PitchedMemorySet(T* const ptr, const size_t pitch, const size_t width, const size_t height,
+                      const size_t depth, F expected_value_generator) {
+  for (int z = 0; z < depth; ++z) {
+    for (int y = 0; y < height; ++y) {
+      for (int x = 0; x < width; ++x) {
+        const auto slice = reinterpret_cast<uint8_t*>(ptr) + pitch * height * z;
+        const auto row = slice + pitch * y;
+        reinterpret_cast<T*>(row)[x] = expected_value_generator(x, y, z);
+      }
+    }
+  }
+}
+
 template <typename T>
 __global__ void VectorIncrement(T* const vec, const T increment_value, size_t N) {
   size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
@@ -82,6 +113,18 @@ static __global__ void Delay(uint32_t interval, const uint32_t ticks_per_ms) {
   }
 }
 
+template <typename T>
+__global__ void Iota(T* const out, size_t pitch, size_t w, size_t h, size_t d) {
+  const auto x = blockIdx.x * blockDim.x + threadIdx.x;
+  const auto y = blockIdx.y * blockDim.y + threadIdx.y;
+  const auto z = blockIdx.z * blockDim.z + threadIdx.z;
+  if (x < w && y < h && z < d) {
+    char* const slice = reinterpret_cast<char*>(out) + pitch * h * z;
+    char* const row = slice + pitch * y;
+    reinterpret_cast<T*>(row)[x] = z * w * h + y * w + x;
+  }
+}
+
 inline void LaunchDelayKernel(const std::chrono::milliseconds interval, const hipStream_t stream) {
   int ticks_per_ms = 0;
   // Clock rate is in kHz => number of clock ticks in a millisecond

From 7734178657b4b586219695d617ede5f8e5a98777 Mon Sep 17 00:00:00 2001
From: Mirza Halilcevic <mirza.halilcevic@htecgroup.com>
Date: Tue, 18 Oct 2022 11:44:28 +0200
Subject: [PATCH 17/21] EXSWHTEC-94 - Implement resource guards for arrays.

---
 tests/catch/include/resource_guards.hh | 82 ++++++++++++++++++++------
 1 file changed, 65 insertions(+), 17 deletions(-)

diff --git a/tests/catch/include/resource_guards.hh b/tests/catch/include/resource_guards.hh
index b3ef7813f7..a9c7512a3d 100644
--- a/tests/catch/include/resource_guards.hh
+++ b/tests/catch/include/resource_guards.hh
@@ -19,6 +19,7 @@ THE SOFTWARE.
 
 #pragma once
 
+#include <hip_array_common.hh>
 #include <hip_test_common.hh>
 #include <hip/hip_runtime_api.h>
 
@@ -90,15 +91,12 @@ template <typename T> class LinearAllocGuard {
 };
 
 template <typename T> class LinearAllocGuardMultiDim {
-  protected:
-  LinearAllocGuardMultiDim(hipExtent extent)
-  : extent_{extent} {}
+ protected:
+  LinearAllocGuardMultiDim(hipExtent extent) : extent_{extent} {}
 
-  ~LinearAllocGuardMultiDim() {
-    static_cast<void>(hipFree(pitched_ptr_.ptr));
-  }
-  
-  public:
+  ~LinearAllocGuardMultiDim() { static_cast<void>(hipFree(pitched_ptr_.ptr)); }
+
+ public:
   T* ptr() const { return reinterpret_cast<T*>(pitched_ptr_.ptr); };
 
   size_t pitch() const { return pitched_ptr_.pitch; }
@@ -113,17 +111,17 @@ template <typename T> class LinearAllocGuardMultiDim {
 
   size_t height() const { return extent_.height; }
 
-  public:
+ public:
   hipPitchedPtr pitched_ptr_;
   const hipExtent extent_;
 };
 
 template <typename T> class LinearAllocGuard2D : public LinearAllocGuardMultiDim<T> {
-  public:
-  LinearAllocGuard2D(const size_t width_logical, const size_t height) 
-  : LinearAllocGuardMultiDim<T>{make_hipExtent(width_logical * sizeof(T), height, 1)}
-  {
-    HIP_CHECK(hipMallocPitch(&this->pitched_ptr_.ptr, &this->pitched_ptr_.pitch, this->extent_.width, this->extent_.height));
+ public:
+  LinearAllocGuard2D(const size_t width_logical, const size_t height)
+      : LinearAllocGuardMultiDim<T>{make_hipExtent(width_logical * sizeof(T), height, 1)} {
+    HIP_CHECK(hipMallocPitch(&this->pitched_ptr_.ptr, &this->pitched_ptr_.pitch,
+                             this->extent_.width, this->extent_.height));
   }
 
   LinearAllocGuard2D(const LinearAllocGuard2D&) = delete;
@@ -131,10 +129,9 @@ template <typename T> class LinearAllocGuard2D : public LinearAllocGuardMultiDim
 };
 
 template <typename T> class LinearAllocGuard3D : public LinearAllocGuardMultiDim<T> {
-  public:
+ public:
   LinearAllocGuard3D(const size_t width_logical, const size_t height, const size_t depth)
-  : LinearAllocGuardMultiDim<T>{make_hipExtent(width_logical * sizeof(T), height, depth)}
-  {
+      : LinearAllocGuardMultiDim<T>{make_hipExtent(width_logical * sizeof(T), height, depth)} {
     HIP_CHECK(hipMalloc3D(&this->pitched_ptr_, this->extent_));
   }
 
@@ -148,6 +145,57 @@ template <typename T> class LinearAllocGuard3D : public LinearAllocGuardMultiDim
   size_t depth() const { return this->extent_.depth; }
 };
 
+template <typename T> class ArrayAllocGuard {
+ public:
+  // extent should contain logical width
+  ArrayAllocGuard(const hipExtent extent, const unsigned int flags = 0u) : extent_{extent} {
+    hipChannelFormatDesc desc = hipCreateChannelDesc<T>();
+    HIP_CHECK(hipMalloc3DArray(&ptr_, &desc, extent_, flags));
+  }
+
+  ~ArrayAllocGuard() { static_cast<void>(hipFreeArray(ptr_)); }
+
+  ArrayAllocGuard(const ArrayAllocGuard&) = delete;
+  ArrayAllocGuard(ArrayAllocGuard&&) = delete;
+
+  hipArray_t ptr() const { return ptr_; }
+
+  hipExtent extent() const { return extent_; }
+
+ private:
+  hipArray_t ptr_ = nullptr;
+  const hipExtent extent_;
+};
+
+template <typename T> class DrvArrayAllocGuard {
+ public:
+  // extent should contain width in bytes
+  DrvArrayAllocGuard(const hipExtent extent, const unsigned int flags = 0u) : extent_{extent} {
+    HIP_ARRAY3D_DESCRIPTOR desc{};
+    using vec_info = vector_info<T>;
+    desc.Format = vec_info::format;
+    desc.NumChannels = vec_info::size;
+    desc.Width = extent_.width / sizeof(T);
+    desc.Height = extent_.height;
+    desc.Depth = extent_.depth;
+    desc.Flags = flags;
+    HIP_CHECK(hipArray3DCreate(&ptr_, &desc));
+  }
+
+  ~DrvArrayAllocGuard() { static_cast<void>(hipArrayDestroy(ptr_)); }
+
+  DrvArrayAllocGuard(const DrvArrayAllocGuard&) = delete;
+  DrvArrayAllocGuard(DrvArrayAllocGuard&&) = delete;
+
+  hiparray ptr() const { return ptr_; }
+
+  hipExtent extent() const { return extent_; }
+
+ private:
+  hiparray ptr_ = nullptr;
+  const hipExtent extent_;
+};
+
 enum class Streams { nullstream, perThread, created };
 
 class StreamGuard {

From 1fd1cb0cdcde42e67f97c3aedf246b2061328132 Mon Sep 17 00:00:00 2001
From: Mirza Halilcevic <mirza.halilcevic@htecgroup.com>
Date: Tue, 18 Oct 2022 13:46:02 +0200
Subject: [PATCH 18/21] EXSWHTEC-94 - Add hip_array_common.hh.

---
 tests/catch/include/hip_array_common.hh | 84 +++++++++++++++++++++++++
 1 file changed, 84 insertions(+)
 create mode 100644 tests/catch/include/hip_array_common.hh

diff --git a/tests/catch/include/hip_array_common.hh b/tests/catch/include/hip_array_common.hh
new file mode 100644
index 0000000000..fd6f094f8d
--- /dev/null
+++ b/tests/catch/include/hip_array_common.hh
@@ -0,0 +1,84 @@
+/*
+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.
+*/
+
+#pragma once
+
+#include <hip_test_common.hh>
+
+template <class T, size_t N, hipArray_Format Format> struct type_and_size_and_format {
+  using type = T;
+  static constexpr size_t size = N;
+  static constexpr hipArray_Format format = Format;
+};
+
+// Create a map of type to scalar type, vector size and scalar type format enum.
+// This is useful for creating simpler function that depend on the vector size.
+template <typename T> struct vector_info;
+template <>
+struct vector_info<int> : type_and_size_and_format<int, 1, HIP_AD_FORMAT_SIGNED_INT32> {};
+template <> struct vector_info<float> : type_and_size_and_format<float, 1, HIP_AD_FORMAT_FLOAT> {};
+template <>
+struct vector_info<short> : type_and_size_and_format<short, 1, HIP_AD_FORMAT_SIGNED_INT16> {};
+template <>
+struct vector_info<char> : type_and_size_and_format<char, 1, HIP_AD_FORMAT_SIGNED_INT8> {};
+template <>
+struct vector_info<unsigned int>
+    : type_and_size_and_format<unsigned int, 1, HIP_AD_FORMAT_UNSIGNED_INT32> {};
+template <>
+struct vector_info<unsigned short>
+    : type_and_size_and_format<unsigned short, 1, HIP_AD_FORMAT_UNSIGNED_INT16> {};
+template <>
+struct vector_info<unsigned char>
+    : type_and_size_and_format<unsigned char, 1, HIP_AD_FORMAT_UNSIGNED_INT8> {};
+
+template <>
+struct vector_info<int2> : type_and_size_and_format<int, 2, HIP_AD_FORMAT_SIGNED_INT32> {};
+template <> struct vector_info<float2> : type_and_size_and_format<float, 2, HIP_AD_FORMAT_FLOAT> {};
+template <>
+struct vector_info<short2> : type_and_size_and_format<short, 2, HIP_AD_FORMAT_SIGNED_INT16> {};
+template <>
+struct vector_info<char2> : type_and_size_and_format<char, 2, HIP_AD_FORMAT_SIGNED_INT8> {};
+template <>
+struct vector_info<uint2>
+    : type_and_size_and_format<unsigned int, 2, HIP_AD_FORMAT_UNSIGNED_INT32> {};
+template <>
+struct vector_info<ushort2>
+    : type_and_size_and_format<unsigned short, 2, HIP_AD_FORMAT_UNSIGNED_INT16> {};
+template <>
+struct vector_info<uchar2>
+    : type_and_size_and_format<unsigned char, 2, HIP_AD_FORMAT_UNSIGNED_INT8> {};
+
+template <>
+struct vector_info<int4> : type_and_size_and_format<int, 4, HIP_AD_FORMAT_SIGNED_INT32> {};
+template <> struct vector_info<float4> : type_and_size_and_format<float, 4, HIP_AD_FORMAT_FLOAT> {};
+template <>
+struct vector_info<short4> : type_and_size_and_format<short, 4, HIP_AD_FORMAT_SIGNED_INT16> {};
+template <>
+struct vector_info<char4> : type_and_size_and_format<char, 4, HIP_AD_FORMAT_SIGNED_INT8> {};
+template <>
+struct vector_info<uint4>
+    : type_and_size_and_format<unsigned int, 4, HIP_AD_FORMAT_UNSIGNED_INT32> {};
+template <>
+struct vector_info<ushort4>
+    : type_and_size_and_format<unsigned short, 4, HIP_AD_FORMAT_UNSIGNED_INT16> {};
+template <>
+struct vector_info<uchar4>
+    : type_and_size_and_format<unsigned char, 4, HIP_AD_FORMAT_UNSIGNED_INT8> {};
\ No newline at end of file

From fc3a10712fd51b28185e596ca2e013d8db79d454 Mon Sep 17 00:00:00 2001
From: Mirza Halilcevic <mirza.halilcevic@htecgroup.com>
Date: Tue, 18 Oct 2022 14:51:44 +0200
Subject: [PATCH 19/21] EXSWHTEC-94 - Remove redundancies between
 hip_array_common.hh and hipArrayCommon.hh.

---
 tests/catch/unit/memory/hipArray3DCreate.cc |  1 +
 tests/catch/unit/memory/hipArrayCommon.hh   | 60 ---------------------
 tests/catch/unit/memory/hipArrayCreate.cc   |  1 +
 tests/catch/unit/memory/hipFree.cc          |  1 +
 tests/catch/unit/memory/hipMallocArray.cc   |  1 +
 5 files changed, 4 insertions(+), 60 deletions(-)

diff --git a/tests/catch/unit/memory/hipArray3DCreate.cc b/tests/catch/unit/memory/hipArray3DCreate.cc
index 973868eded..4cf189611b 100644
--- a/tests/catch/unit/memory/hipArray3DCreate.cc
+++ b/tests/catch/unit/memory/hipArray3DCreate.cc
@@ -20,6 +20,7 @@ THE SOFTWARE.
 #include <limits>
 #include "DriverContext.hh"
 #include "hipArrayCommon.hh"
+#include "hip_array_common.hh"
 #include "hip_test_common.hh"
 
 namespace {
diff --git a/tests/catch/unit/memory/hipArrayCommon.hh b/tests/catch/unit/memory/hipArrayCommon.hh
index b40014b490..b0beeb3126 100644
--- a/tests/catch/unit/memory/hipArrayCommon.hh
+++ b/tests/catch/unit/memory/hipArrayCommon.hh
@@ -26,66 +26,6 @@ THE SOFTWARE.
 
 constexpr size_t BlockSize = 16;
 
-template <class T, size_t N, hipArray_Format Format> struct type_and_size_and_format {
-  using type = T;
-  static constexpr size_t size = N;
-  static constexpr hipArray_Format format = Format;
-};
-
-// Create a map of type to scalar type, vector size and scalar type format enum.
-// This is useful for creating simpler function that depend on the vector size.
-template <typename T> struct vector_info;
-template <>
-struct vector_info<int> : type_and_size_and_format<int, 1, HIP_AD_FORMAT_SIGNED_INT32> {};
-template <> struct vector_info<float> : type_and_size_and_format<float, 1, HIP_AD_FORMAT_FLOAT> {};
-template <>
-struct vector_info<short> : type_and_size_and_format<short, 1, HIP_AD_FORMAT_SIGNED_INT16> {};
-template <>
-struct vector_info<char> : type_and_size_and_format<char, 1, HIP_AD_FORMAT_SIGNED_INT8> {};
-template <>
-struct vector_info<unsigned int>
-    : type_and_size_and_format<unsigned int, 1, HIP_AD_FORMAT_UNSIGNED_INT32> {};
-template <>
-struct vector_info<unsigned short>
-    : type_and_size_and_format<unsigned short, 1, HIP_AD_FORMAT_UNSIGNED_INT16> {};
-template <>
-struct vector_info<unsigned char>
-    : type_and_size_and_format<unsigned char, 1, HIP_AD_FORMAT_UNSIGNED_INT8> {};
-
-template <>
-struct vector_info<int2> : type_and_size_and_format<int, 2, HIP_AD_FORMAT_SIGNED_INT32> {};
-template <> struct vector_info<float2> : type_and_size_and_format<float, 2, HIP_AD_FORMAT_FLOAT> {};
-template <>
-struct vector_info<short2> : type_and_size_and_format<short, 2, HIP_AD_FORMAT_SIGNED_INT16> {};
-template <>
-struct vector_info<char2> : type_and_size_and_format<char, 2, HIP_AD_FORMAT_SIGNED_INT8> {};
-template <>
-struct vector_info<uint2>
-    : type_and_size_and_format<unsigned int, 2, HIP_AD_FORMAT_UNSIGNED_INT32> {};
-template <>
-struct vector_info<ushort2>
-    : type_and_size_and_format<unsigned short, 2, HIP_AD_FORMAT_UNSIGNED_INT16> {};
-template <>
-struct vector_info<uchar2>
-    : type_and_size_and_format<unsigned char, 2, HIP_AD_FORMAT_UNSIGNED_INT8> {};
-
-template <>
-struct vector_info<int4> : type_and_size_and_format<int, 4, HIP_AD_FORMAT_SIGNED_INT32> {};
-template <> struct vector_info<float4> : type_and_size_and_format<float, 4, HIP_AD_FORMAT_FLOAT> {};
-template <>
-struct vector_info<short4> : type_and_size_and_format<short, 4, HIP_AD_FORMAT_SIGNED_INT16> {};
-template <>
-struct vector_info<char4> : type_and_size_and_format<char, 4, HIP_AD_FORMAT_SIGNED_INT8> {};
-template <>
-struct vector_info<uint4>
-    : type_and_size_and_format<unsigned int, 4, HIP_AD_FORMAT_UNSIGNED_INT32> {};
-template <>
-struct vector_info<ushort4>
-    : type_and_size_and_format<unsigned short, 4, HIP_AD_FORMAT_UNSIGNED_INT16> {};
-template <>
-struct vector_info<uchar4>
-    : type_and_size_and_format<unsigned char, 4, HIP_AD_FORMAT_UNSIGNED_INT8> {};
-
 // read from a texture using normalized coordinates
 constexpr size_t ChannelToRead = 1;
 template <typename T>
diff --git a/tests/catch/unit/memory/hipArrayCreate.cc b/tests/catch/unit/memory/hipArrayCreate.cc
index 6cc535593a..70a8636922 100644
--- a/tests/catch/unit/memory/hipArrayCreate.cc
+++ b/tests/catch/unit/memory/hipArrayCreate.cc
@@ -27,6 +27,7 @@ hipArrayCreate API test scenarios
 #include <array>
 #include <numeric>
 #include <hip_test_common.hh>
+#include <hip_array_common.hh>
 #include "hipArrayCommon.hh"
 #include "DriverContext.hh"
 
diff --git a/tests/catch/unit/memory/hipFree.cc b/tests/catch/unit/memory/hipFree.cc
index 1248deebc1..b29854271c 100644
--- a/tests/catch/unit/memory/hipFree.cc
+++ b/tests/catch/unit/memory/hipFree.cc
@@ -22,6 +22,7 @@ THE SOFTWARE.
 
 
 #include <hip_test_common.hh>
+#include <hip_array_common.hh>
 #include "hipArrayCommon.hh"
 #include "DriverContext.hh"
 
diff --git a/tests/catch/unit/memory/hipMallocArray.cc b/tests/catch/unit/memory/hipMallocArray.cc
index b6c4939b1e..530eb11077 100644
--- a/tests/catch/unit/memory/hipMallocArray.cc
+++ b/tests/catch/unit/memory/hipMallocArray.cc
@@ -26,6 +26,7 @@ hipMallocArray API test scenarios
 */
 
 #include <hip_test_common.hh>
+#include <hip_array_common.hh>
 #include <limits>
 #include <numeric>
 #include "hipArrayCommon.hh"

From 4e2fe8d999f45d3645a7714b5eff908eeaee64e3 Mon Sep 17 00:00:00 2001
From: Dino Music <dino.music@htecgroup.com>
Date: Fri, 4 Nov 2022 12:11:48 +0100
Subject: [PATCH 20/21] EXSWHTEC-94 - Fix loop counter types in
 PitchedMemoryVerify and PitchedMemorySet

---
 tests/catch/include/utils.hh | 12 ++++++------
 1 file changed, 6 insertions(+), 6 deletions(-)

diff --git a/tests/catch/include/utils.hh b/tests/catch/include/utils.hh
index 05eecea79f..bbab2322fe 100644
--- a/tests/catch/include/utils.hh
+++ b/tests/catch/include/utils.hh
@@ -57,9 +57,9 @@ void ArrayFindIfNot(T* const array, const T expected_value, const size_t num_ele
 template <typename T, typename F>
 void PitchedMemoryVerify(T* const ptr, const size_t pitch, const size_t width, const size_t height,
                          const size_t depth, F expected_value_generator) {
-  for (int z = 0; z < depth; ++z) {
-    for (int y = 0; y < height; ++y) {
-      for (int x = 0; x < width; ++x) {
+  for (size_t z = 0; z < depth; ++z) {
+    for (size_t y = 0; y < height; ++y) {
+      for (size_t x = 0; x < width; ++x) {
         const auto slice = reinterpret_cast<uint8_t*>(ptr) + pitch * height * z;
         const auto row = slice + pitch * y;
         if (reinterpret_cast<T*>(row)[x] != expected_value_generator(x, y, z)) {
@@ -74,9 +74,9 @@ void PitchedMemoryVerify(T* const ptr, const size_t pitch, const size_t width, c
 template <typename T, typename F>
 void PitchedMemorySet(T* const ptr, const size_t pitch, const size_t width, const size_t height,
                       const size_t depth, F expected_value_generator) {
-  for (int z = 0; z < depth; ++z) {
-    for (int y = 0; y < height; ++y) {
-      for (int x = 0; x < width; ++x) {
+  for (size_t z = 0; z < depth; ++z) {
+    for (size_t y = 0; y < height; ++y) {
+      for (size_t x = 0; x < width; ++x) {
         const auto slice = reinterpret_cast<uint8_t*>(ptr) + pitch * height * z;
         const auto row = slice + pitch * y;
         reinterpret_cast<T*>(row)[x] = expected_value_generator(x, y, z);

From bd7e72011704d791a2f30f6e1765406ed60948e8 Mon Sep 17 00:00:00 2001
From: Dino Music <dino.music@htecgroup.com>
Date: Fri, 4 Nov 2022 12:48:07 +0100
Subject: [PATCH 21/21] Disable dev_ptr points to non-managed memory section
 due to defect

---
 tests/catch/unit/memory/hipMemPrefetchAsync.cc | 7 +++++++
 1 file changed, 7 insertions(+)

diff --git a/tests/catch/unit/memory/hipMemPrefetchAsync.cc b/tests/catch/unit/memory/hipMemPrefetchAsync.cc
index c7dfb0f123..92dc8c9a4b 100644
--- a/tests/catch/unit/memory/hipMemPrefetchAsync.cc
+++ b/tests/catch/unit/memory/hipMemPrefetchAsync.cc
@@ -135,20 +135,27 @@ TEST_CASE("Unit_hipMemPrefetchAsync_Negative_Parameters") {
   SECTION("dev_ptr == nullptr") {
     HIP_CHECK_ERROR(hipMemPrefetchAsync(nullptr, kPageSize, device), hipErrorInvalidValue);
   }
+
+#if HT_NVIDIA
   SECTION("dev_ptr points to non-managed memory") {
     LinearAllocGuard<void> alloc(LinearAllocs::hipMalloc, kPageSize);
     HIP_CHECK_ERROR(hipMemPrefetchAsync(alloc.ptr(), kPageSize, device), hipErrorInvalidValue);
   }
+#endif
+
   SECTION("count == 0") {
     HIP_CHECK_ERROR(hipMemPrefetchAsync(alloc.ptr(), 0, device), hipErrorInvalidValue);
   }
+
   SECTION("count larger than allocation size") {
     HIP_CHECK_ERROR(hipMemPrefetchAsync(alloc.ptr(), kPageSize + 1, device), hipErrorInvalidValue);
   }
+
   SECTION("Invalid device") {
     HIP_CHECK_ERROR(hipMemPrefetchAsync(alloc.ptr(), kPageSize, hipInvalidDeviceId),
                     hipErrorInvalidDevice);
   }
+
   SECTION("Invalid stream") {
     hipStream_t stream;
     HIP_CHECK(hipStreamCreate(&stream));