EXSWHTEC-94 - Implement helper classes and functions for memory tests (#2978)

- Implement helper classes and functions for memory tests - Remove c++14 standard constraint on memory tests - Remove GenerateLinearAllocationFlagCombinations until finished [ROCm/hip-tests commit: 77bc96131b]
2022-10-12 05:00:51 +02:00
@@ -0,0 +1,125 @@
+/*
+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_;
+};
@@ -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));
+}
@@ -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)