Skip to content

Commit e534d2e

Browse files
music-dinomusic-dino
committed
EXSWHTEC-94 - Implement helper classes and functions for memory tests
1 parent c02b15f commit e534d2e

File tree

2 files changed

+246
-0
lines changed

2 files changed

+246
-0
lines changed

‎tests/catch/include/resource_guards.hh‎

Lines changed: 144 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,144 @@
1+
/*
2+
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
3+
Permission is hereby granted, free of charge, to any person obtaining a copy
4+
of this software and associated documentation files (the "Software"), to deal
5+
in the Software without restriction, including without limitation the rights
6+
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
7+
copies of the Software, and to permit persons to whom the Software is
8+
furnished to do so, subject to the following conditions:
9+
The above copyright notice and this permission notice shall be included in
10+
all copies or substantial portions of the Software.
11+
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
12+
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
13+
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
14+
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
15+
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
16+
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
17+
THE SOFTWARE.
18+
*/
19+
20+
#pragma once
21+
22+
#include <hip_test_common.hh>
23+
#include <hip/hip_runtime_api.h>
24+
25+
enum class LinearAllocs {
26+
malloc,
27+
mallocAndRegister,
28+
hipHostMalloc,
29+
hipMalloc,
30+
hipMallocManaged,
31+
};
32+
33+
template <typename T> class LinearAllocGuard {
34+
public:
35+
LinearAllocGuard(const LinearAllocs allocation_type, const size_t size,
36+
const unsigned int flags = 0u)
37+
: allocation_type_{allocation_type} {
38+
switch (allocation_type_) {
39+
case LinearAllocs::malloc:
40+
ptr_ = host_ptr_ = reinterpret_cast<T*>(malloc(size));
41+
break;
42+
case LinearAllocs::mallocAndRegister:
43+
host_ptr_ = reinterpret_cast<T*>(malloc(size));
44+
HIP_CHECK(hipHostRegister(host_ptr_, size, flags));
45+
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&ptr_), host_ptr_, 0u));
46+
break;
47+
case LinearAllocs::hipHostMalloc:
48+
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&ptr_), size, flags));
49+
host_ptr_ = ptr_;
50+
break;
51+
case LinearAllocs::hipMalloc:
52+
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&ptr_), size));
53+
break;
54+
case LinearAllocs::hipMallocManaged:
55+
HIP_CHECK(hipMallocManaged(reinterpret_cast<void**>(&ptr_), size, flags ? flags : 1u));
56+
host_ptr_ = ptr_;
57+
}
58+
}
59+
60+
LinearAllocGuard(const LinearAllocGuard&) = delete;
61+
LinearAllocGuard(LinearAllocGuard&&) = delete;
62+
63+
~LinearAllocGuard() {
64+
// No Catch macros, don't want to possibly throw in the destructor
65+
switch (allocation_type_) {
66+
case LinearAllocs::malloc:
67+
free(ptr_);
68+
break;
69+
case LinearAllocs::mallocAndRegister:
70+
// Cast to void to suppress nodiscard warnings
71+
static_cast<void>(hipHostUnregister(host_ptr_));
72+
free(host_ptr_);
73+
break;
74+
case LinearAllocs::hipHostMalloc:
75+
static_cast<void>(hipHostFree(ptr_));
76+
break;
77+
case LinearAllocs::hipMalloc:
78+
case LinearAllocs::hipMallocManaged:
79+
static_cast<void>(hipFree(ptr_));
80+
}
81+
}
82+
83+
T* ptr() { return ptr_; };
84+
T* const ptr() const { return ptr_; };
85+
T* host_ptr() { return host_ptr_; }
86+
T* const host_ptr() const { return host_ptr(); }
87+
88+
private:
89+
const LinearAllocs allocation_type_;
90+
T* ptr_ = nullptr;
91+
T* host_ptr_ = nullptr;
92+
};
93+
94+
enum class Streams { nullstream, perThread, created };
95+
96+
class StreamGuard {
97+
public:
98+
StreamGuard(const Streams stream_type) : stream_type_{stream_type} {
99+
switch (stream_type_) {
100+
case Streams::nullstream:
101+
stream_ = nullptr;
102+
break;
103+
case Streams::perThread:
104+
stream_ = hipStreamPerThread;
105+
break;
106+
case Streams::created:
107+
HIP_CHECK(hipStreamCreate(&stream_));
108+
}
109+
}
110+
111+
StreamGuard(const StreamGuard&) = delete;
112+
StreamGuard(StreamGuard&&) = delete;
113+
114+
~StreamGuard() {
115+
if (stream_type_ == Streams::created) {
116+
static_cast<void>(hipStreamDestroy(stream_));
117+
}
118+
}
119+
120+
hipStream_t stream() const { return stream_; }
121+
122+
private:
123+
const Streams stream_type_;
124+
hipStream_t stream_;
125+
};
126+
127+
inline unsigned int GenerateLinearAllocationFlagCombinations(const LinearAllocs allocation_type) {
128+
switch (allocation_type) {
129+
case LinearAllocs::mallocAndRegister:
130+
// TODO
131+
return 0;
132+
case LinearAllocs::hipHostMalloc:
133+
return GENERATE(hipHostMallocDefault, hipHostMallocPortable, hipHostMallocMapped,
134+
hipHostMallocWriteCombined);
135+
case LinearAllocs::hipMallocManaged:
136+
// TODO
137+
return 1u;
138+
case LinearAllocs::malloc:
139+
case LinearAllocs::hipMalloc:
140+
return 0u;
141+
default:
142+
assert("Invalid LinearAllocs enumerator");
143+
}
144+
}

‎tests/catch/include/utils.hh‎

Lines changed: 102 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,102 @@
1+
/*
2+
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
3+
Permission is hereby granted, free of charge, to any person obtaining a copy
4+
of this software and associated documentation files (the "Software"), to deal
5+
in the Software without restriction, including without limitation the rights
6+
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
7+
copies of the Software, and to permit persons to whom the Software is
8+
furnished to do so, subject to the following conditions:
9+
The above copyright notice and this permission notice shall be included in
10+
all copies or substantial portions of the Software.
11+
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
12+
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
13+
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
14+
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
15+
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
16+
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
17+
THE SOFTWARE.
18+
*/
19+
20+
#pragma once
21+
22+
#include <chrono>
23+
24+
#include <hip_test_common.hh>
25+
#include <hip/hip_runtime_api.h>
26+
27+
namespace {
28+
inline constexpr size_t kPageSize = 4096;
29+
} // anonymous namespace
30+
31+
template <typename T>
32+
void ArrayMismatch(T* const expected, T* const actual, const size_t num_elements) {
33+
const auto ret = std::mismatch(expected, expected + num_elements, actual);
34+
if (ret.first != expected + num_elements) {
35+
const auto idx = std::distance(expected, ret.first);
36+
INFO("Value mismatch at index: " << idx);
37+
REQUIRE(expected[idx] == actual[idx]);
38+
}
39+
}
40+
41+
template <typename It, typename T> void ArrayFindIfNot(It begin, It end, const T expected_value) {
42+
const auto it = std::find_if_not(
43+
begin, end, [expected_value](const int elem) { return expected_value == elem; });
44+
45+
if (it != end) {
46+
const auto idx = std::distance(begin, it);
47+
INFO("Value mismatch at index " << idx);
48+
REQUIRE(expected_value == *it);
49+
}
50+
}
51+
52+
template <typename T>
53+
void ArrayFindIfNot(T* const array, const T expected_value, const size_t num_elements) {
54+
ArrayFindIfNot(array, array + num_elements, expected_value);
55+
}
56+
57+
template <typename T>
58+
__global__ void VectorIncrement(T* const vec, const T increment_value, size_t N) {
59+
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
60+
size_t stride = blockDim.x * gridDim.x;
61+
62+
for (size_t i = offset; i < N; i += stride) {
63+
vec[i] += increment_value;
64+
}
65+
}
66+
67+
template <typename T> __global__ void VectorSet(T* const vec, const T value, size_t N) {
68+
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
69+
size_t stride = blockDim.x * gridDim.x;
70+
71+
for (size_t i = offset; i < N; i += stride) {
72+
vec[i] = value;
73+
}
74+
}
75+
76+
// Will execute for atleast interval milliseconds
77+
static __global__ void Delay(uint32_t interval, const uint32_t ticks_per_ms) {
78+
while (interval--) {
79+
uint64_t start = clock();
80+
while (clock() - start < ticks_per_ms) {
81+
}
82+
}
83+
}
84+
85+
inline void LaunchDelayKernel(const std::chrono::milliseconds interval, const hipStream_t stream) {
86+
int ticks_per_ms = 0;
87+
// Clock rate is in kHz => number of clock ticks in a millisecond
88+
HIP_CHECK(hipDeviceGetAttribute(&ticks_per_ms, hipDeviceAttributeClockRate, 0));
89+
Delay<<<1, 1, 0, stream>>>(interval.count(), ticks_per_ms);
90+
HIP_CHECK(hipGetLastError());
91+
}
92+
93+
template <typename... Attributes>
94+
inline bool DeviceAttributesSupport(const int device, Attributes... attributes) {
95+
constexpr auto DeviceAttributeSupport = [](const int device,
96+
const hipDeviceAttribute_t attribute) {
97+
int value = 0;
98+
HIP_CHECK(hipDeviceGetAttribute(&value, attribute, device));
99+
return value;
100+
};
101+
return (... && DeviceAttributeSupport(device, attributes));
102+
}

0 commit comments

Comments
 (0)