From e534d2e050446f4e936c0de76fc81b5858782654 Mon Sep 17 00:00:00 2001 From: Dino Music Date: Thu, 6 Oct 2022 15:47:34 +0200 Subject: [PATCH 01/12] 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 +#include + +enum class LinearAllocs { + malloc, + mallocAndRegister, + hipHostMalloc, + hipMalloc, + hipMallocManaged, +}; + +template 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(malloc(size)); + break; + case LinearAllocs::mallocAndRegister: + host_ptr_ = reinterpret_cast(malloc(size)); + HIP_CHECK(hipHostRegister(host_ptr_, size, flags)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&ptr_), host_ptr_, 0u)); + break; + case LinearAllocs::hipHostMalloc: + HIP_CHECK(hipHostMalloc(reinterpret_cast(&ptr_), size, flags)); + host_ptr_ = ptr_; + break; + case LinearAllocs::hipMalloc: + HIP_CHECK(hipMalloc(reinterpret_cast(&ptr_), size)); + break; + case LinearAllocs::hipMallocManaged: + HIP_CHECK(hipMallocManaged(reinterpret_cast(&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(hipHostUnregister(host_ptr_)); + free(host_ptr_); + break; + case LinearAllocs::hipHostMalloc: + static_cast(hipHostFree(ptr_)); + break; + case LinearAllocs::hipMalloc: + case LinearAllocs::hipMallocManaged: + static_cast(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(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 + +#include +#include + +namespace { +inline constexpr size_t kPageSize = 4096; +} // anonymous namespace + +template +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 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 +void ArrayFindIfNot(T* const array, const T expected_value, const size_t num_elements) { + ArrayFindIfNot(array, array + num_elements, expected_value); +} + +template +__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 __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 +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 Date: Thu, 6 Oct 2022 16:19:21 +0200 Subject: [PATCH 02/12] 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 Date: Thu, 6 Oct 2022 12:09:41 -0400 Subject: [PATCH 03/12] 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 Date: Thu, 6 Oct 2022 15:47:34 +0200 Subject: [PATCH 04/12] 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 +#include + +enum class LinearAllocs { + malloc, + mallocAndRegister, + hipHostMalloc, + hipMalloc, + hipMallocManaged, +}; + +template 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(malloc(size)); + break; + case LinearAllocs::mallocAndRegister: + host_ptr_ = reinterpret_cast(malloc(size)); + HIP_CHECK(hipHostRegister(host_ptr_, size, flags)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&ptr_), host_ptr_, 0u)); + break; + case LinearAllocs::hipHostMalloc: + HIP_CHECK(hipHostMalloc(reinterpret_cast(&ptr_), size, flags)); + host_ptr_ = ptr_; + break; + case LinearAllocs::hipMalloc: + HIP_CHECK(hipMalloc(reinterpret_cast(&ptr_), size)); + break; + case LinearAllocs::hipMallocManaged: + HIP_CHECK(hipMallocManaged(reinterpret_cast(&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(hipHostUnregister(host_ptr_)); + free(host_ptr_); + break; + case LinearAllocs::hipHostMalloc: + static_cast(hipHostFree(ptr_)); + break; + case LinearAllocs::hipMalloc: + case LinearAllocs::hipMallocManaged: + static_cast(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(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 + +#include +#include + +namespace { +inline constexpr size_t kPageSize = 4096; +} // anonymous namespace + +template +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 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 +void ArrayFindIfNot(T* const array, const T expected_value, const size_t num_elements) { + ArrayFindIfNot(array, array + num_elements, expected_value); +} + +template +__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 __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 +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 Date: Thu, 6 Oct 2022 16:19:21 +0200 Subject: [PATCH 05/12] 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 Date: Thu, 6 Oct 2022 12:09:41 -0400 Subject: [PATCH 06/12] 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 Date: Thu, 6 Oct 2022 15:47:34 +0200 Subject: [PATCH 07/12] 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 +#include + +enum class LinearAllocs { + malloc, + mallocAndRegister, + hipHostMalloc, + hipMalloc, + hipMallocManaged, +}; + +template 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(malloc(size)); + break; + case LinearAllocs::mallocAndRegister: + host_ptr_ = reinterpret_cast(malloc(size)); + HIP_CHECK(hipHostRegister(host_ptr_, size, flags)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&ptr_), host_ptr_, 0u)); + break; + case LinearAllocs::hipHostMalloc: + HIP_CHECK(hipHostMalloc(reinterpret_cast(&ptr_), size, flags)); + host_ptr_ = ptr_; + break; + case LinearAllocs::hipMalloc: + HIP_CHECK(hipMalloc(reinterpret_cast(&ptr_), size)); + break; + case LinearAllocs::hipMallocManaged: + HIP_CHECK(hipMallocManaged(reinterpret_cast(&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(hipHostUnregister(host_ptr_)); + free(host_ptr_); + break; + case LinearAllocs::hipHostMalloc: + static_cast(hipHostFree(ptr_)); + break; + case LinearAllocs::hipMalloc: + case LinearAllocs::hipMallocManaged: + static_cast(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(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 + +#include +#include + +namespace { +inline constexpr size_t kPageSize = 4096; +} // anonymous namespace + +template +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 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 +void ArrayFindIfNot(T* const array, const T expected_value, const size_t num_elements) { + ArrayFindIfNot(array, array + num_elements, expected_value); +} + +template +__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 __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 +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 Date: Thu, 6 Oct 2022 16:19:21 +0200 Subject: [PATCH 08/12] 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 Date: Thu, 6 Oct 2022 12:09:41 -0400 Subject: [PATCH 09/12] 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 Date: Wed, 12 Oct 2022 10:25:11 +0200 Subject: [PATCH 10/12] 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 class LinearAllocGuard { T* host_ptr_ = nullptr; }; +template class LinearAllocGuardMultiDim { + protected: + LinearAllocGuardMultiDim(hipExtent extent) + : extent_{extent} {} + + ~LinearAllocGuardMultiDim() { + static_cast(hipFree(pitched_ptr_.ptr)); + } + + public: + T* ptr() const { return reinterpret_cast(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 class LinearAllocGuard2D : public LinearAllocGuardMultiDim { + public: + LinearAllocGuard2D(const size_t width_logical, const size_t height) + : LinearAllocGuardMultiDim{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 class LinearAllocGuard3D : public LinearAllocGuardMultiDim { + public: + LinearAllocGuard3D(const size_t width_logical, const size_t height, const size_t depth) + : LinearAllocGuardMultiDim{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 Date: Fri, 14 Oct 2022 19:47:44 +0200 Subject: [PATCH 11/12] 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 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 class LinearAllocGuard3D : public LinearAllocGuardMultiDim HIP_CHECK(hipMalloc3D(&this->pitched_ptr_, this->extent_)); } + LinearAllocGuard3D(const hipExtent extent) : LinearAllocGuardMultiDim(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 +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(ptr) + pitch * height * z; + const auto row = slice + pitch * y; + if (reinterpret_cast(row)[x] != expected_value_generator(x, y, z)) { + INFO("Mismatch at indices: " << x << ", " << y << ", " << z); + REQUIRE(reinterpret_cast(row)[x] == expected_value_generator(x, y, z)); + } + } + } + } +} + +template +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(ptr) + pitch * height * z; + const auto row = slice + pitch * y; + reinterpret_cast(row)[x] = expected_value_generator(x, y, z); + } + } + } +} + template __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 +__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(out) + pitch * h * z; + char* const row = slice + pitch * y; + reinterpret_cast(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 e67bd18a5e5baaf75865c12fb18e3cd813f85b9b Mon Sep 17 00:00:00 2001 From: Dino Music Date: Mon, 17 Oct 2022 10:03:21 -0400 Subject: [PATCH 12/12] EXSWHTEC-100 - Miscellaneous modifications to existing tests --- tests/catch/unit/memory/hipFree.cc | 18 +++---- .../unit/memory/hipHostGetDevicePointer.cc | 30 +++++++++-- tests/catch/unit/memory/hipHostRegister.cc | 7 ++- tests/catch/unit/memory/hipHostUnregister.cc | 6 +++ tests/catch/unit/memory/hipMallocPitch.cc | 51 ++++++------------- .../unit/memory/hipPointerGetAttribute.cc | 5 +- 6 files changed, 61 insertions(+), 56 deletions(-) diff --git a/tests/catch/unit/memory/hipFree.cc b/tests/catch/unit/memory/hipFree.cc index 1248deebc1..018b95b9c3 100644 --- a/tests/catch/unit/memory/hipFree.cc +++ b/tests/catch/unit/memory/hipFree.cc @@ -48,11 +48,10 @@ using namespace std::chrono_literals; const std::chrono::duration delay = 50ms; constexpr size_t numAllocs = 10; -#if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */ -TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncDev", "", char, float, float2, float4) { - TestType* devPtr{}; +TEST_CASE("Unit_hipFreeImplicitSyncDev") { + int* devPtr{}; size_t size_mult = GENERATE(1, 32, 64, 128, 256); - HIP_CHECK(hipMalloc(&devPtr, sizeof(TestType) * size_mult)); + HIP_CHECK(hipMalloc(&devPtr, sizeof(*devPtr) * size_mult)); HipTest::runKernelForDuration(delay); // make sure device is busy @@ -61,11 +60,11 @@ TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncDev", "", char, float, float2, float HIP_CHECK(hipStreamQuery(nullptr)); } -TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncHost", "", char, float, float2, float4) { - TestType* hostPtr{}; +TEST_CASE("Unit_hipFreeImplicitSyncHost") { + int* hostPtr{}; size_t size_mult = GENERATE(1, 32, 64, 128, 256); - HIP_CHECK(hipHostMalloc(&hostPtr, sizeof(TestType) * size_mult)); + HIP_CHECK(hipHostMalloc(&hostPtr, sizeof(*hostPtr) * size_mult)); HipTest::runKernelForDuration(delay); // make sure device is busy @@ -74,7 +73,7 @@ TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncHost", "", char, float, float2, floa HIP_CHECK(hipStreamQuery(nullptr)); } -#if HT_NVIDIA // Meaningless at the moment, since we are not running wait kernel on nvidia. +#if HT_NVIDIA TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncArray", "", char, float, float2, float4) { using vec_info = vector_info; DriverContext ctx; @@ -134,7 +133,6 @@ TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncArray", "", char, float, float2, flo } } -#endif #endif // Freeing a invalid pointer with on device @@ -165,8 +163,6 @@ TEST_CASE("Unit_hipFreeNegativeHost") { #if HT_NVIDIA TEST_CASE("Unit_hipFreeNegativeArray") { DriverContext ctx; - hipArray_t arrayPtr{}; - hiparray cuArrayPtr{}; SECTION("ArrayFree") { HIP_CHECK(hipFreeArray(nullptr)); } SECTION("ArrayDestroy") { diff --git a/tests/catch/unit/memory/hipHostGetDevicePointer.cc b/tests/catch/unit/memory/hipHostGetDevicePointer.cc index 7c3e689e05..7f07468935 100644 --- a/tests/catch/unit/memory/hipHostGetDevicePointer.cc +++ b/tests/catch/unit/memory/hipHostGetDevicePointer.cc @@ -21,11 +21,19 @@ THE SOFTWARE. */ #include +#include TEST_CASE("Unit_hipHostGetDevicePointer_Negative") { int* hPtr{nullptr}; + int* dPtr{nullptr}; HIP_CHECK(hipHostMalloc(&hPtr, sizeof(int))); + if (!DeviceAttributesSupport(0, hipDeviceAttributeCanMapHostMemory)) { + HIP_CHECK_ERROR(hipHostGetDevicePointer(reinterpret_cast(&dPtr), hPtr, 0), + hipErrorNotSupported); + return; + } + SECTION("Nullptr as device") { HIP_CHECK_ERROR(hipHostGetDevicePointer(nullptr, hPtr, 0), hipErrorInvalidValue); } @@ -36,13 +44,29 @@ TEST_CASE("Unit_hipHostGetDevicePointer_Negative") { hipErrorInvalidValue); } - // Not adding check for flags since CUDA spec states that there might be more values added to it + SECTION("Non pinned memory as host") { + int* hPtr = reinterpret_cast(malloc(sizeof(*hPtr))); + HIP_CHECK_ERROR(hipHostGetDevicePointer(reinterpret_cast(&dPtr), hPtr, 0), + hipErrorInvalidValue); + free(hPtr); + } + + SECTION("Flags non-zero") { + HIP_CHECK_ERROR(hipHostGetDevicePointer(reinterpret_cast(&dPtr), hPtr, 1), + hipErrorInvalidValue); + } + HIP_CHECK(hipHostFree(hPtr)); } template __global__ void set(T* ptr, T val) { *ptr = val; } TEST_CASE("Unit_hipHostGetDevicePointer_UseCase") { + if(!DeviceAttributesSupport(0, hipDeviceAttributeCanMapHostMemory)) { + HipTest::HIP_SKIP_TEST("Device does not support mapping host memory"); + return; + } + int* hPtr{nullptr}; HIP_CHECK(hipHostMalloc(&hPtr, sizeof(int))); @@ -71,8 +95,8 @@ TEST_CASE("Unit_hipHostGetDevicePointer_UseCase") { HIP_CHECK(hipDeviceSynchronize()); HIP_CHECK(hipHostUnregister(&res)); - REQUIRE(value == 10); + REQUIRE(res == value); } HIP_CHECK(hipHostFree(hPtr)); -} +} \ No newline at end of file diff --git a/tests/catch/unit/memory/hipHostRegister.cc b/tests/catch/unit/memory/hipHostRegister.cc index f6964db616..5e1b10d234 100644 --- a/tests/catch/unit/memory/hipHostRegister.cc +++ b/tests/catch/unit/memory/hipHostRegister.cc @@ -27,9 +27,10 @@ This testfile verifies the following scenarios of hipHostRegister API 2. hipHostRegister and perform hipMemcpy on it. */ +#include "hip/hip_runtime_api.h" #include #include -#include "hip/hip_runtime_api.h" +#include #define OFFSET 128 static constexpr auto LEN{1024 * 1024}; @@ -63,9 +64,7 @@ void doMemCopy(size_t numElements, int offset, T* A, T* Bh, T* Bd, bool internal HIP_CHECK(hipMemcpy(Bh, Bd, sizeBytes, hipMemcpyDeviceToHost)); // Make sure the copy worked - for (size_t i = 0; i < numElements; i++) { - REQUIRE(Bh[i] == A[i]); - } + ArrayMismatch(A, Bh, numElements); if (internalRegister) { HIP_CHECK(hipHostUnregister(A)); diff --git a/tests/catch/unit/memory/hipHostUnregister.cc b/tests/catch/unit/memory/hipHostUnregister.cc index 69373133d0..ea3d018a33 100644 --- a/tests/catch/unit/memory/hipHostUnregister.cc +++ b/tests/catch/unit/memory/hipHostUnregister.cc @@ -68,6 +68,12 @@ TEST_CASE("Unit_hipHostUnregister_NullPtr") { HIP_CHECK_ERROR(hipHostUnregister(nullptr), hipErrorInvalidValue); } +TEST_CASE("Unit_hipHostUnregister_Ptr_Different_Than_Specified_To_Register") { + std::vector alloc(2); + HIP_CHECK(hipHostRegister(alloc.data(), alloc.size(), 0)); + HIP_CHECK_ERROR(hipHostUnregister(&alloc.data()[1]), hipErrorHostMemoryNotRegistered); +} + TEST_CASE("Unit_hipHostUnregister_NotRegisteredPointer") { auto x = std::unique_ptr(new int); HIP_CHECK_ERROR(hipHostUnregister(x.get()), hipErrorHostMemoryNotRegistered); diff --git a/tests/catch/unit/memory/hipMallocPitch.cc b/tests/catch/unit/memory/hipMallocPitch.cc index 5a20671e14..b84b45087b 100644 --- a/tests/catch/unit/memory/hipMallocPitch.cc +++ b/tests/catch/unit/memory/hipMallocPitch.cc @@ -228,6 +228,21 @@ TEST_CASE("Unit_hipMallocPitch_Negative") { } } +TEST_CASE("Unit_hipMallocPitch_Zero_Dims") { + void* ptr = nullptr; + size_t pitch = 0; + + SECTION("width == 0") { + HIP_CHECK(hipMallocPitch(&ptr, &pitch, 0, 1)); + REQUIRE(ptr == nullptr); + } + + SECTION("height == 0") { + HIP_CHECK(hipMallocPitch(&ptr, &pitch, 1, 0)); + REQUIRE(ptr == nullptr); + } +} + TEST_CASE("Unit_hipMemAllocPitch_Negative") { size_t pitch = 0; hipDeviceptr_t ptr{}; @@ -366,42 +381,7 @@ static void MemoryAllocDiffSizes(int gpu) { static void threadFunc(int gpu) { MemoryAllocDiffSizes(gpu); } -/* - * This testcase verifies the negative scenarios of hipMallocPitch API - */ -#if 0 //TODO: Review, fix and re-enable test -TEST_CASE("Unit_hipMallocPitch_Negative") { - float* A_d; - size_t pitch_A = 0; - size_t width{NUM_W * sizeof(float)}; -#if HT_NVIDIA - SECTION("NullPtr to Pitched Ptr") { - REQUIRE(hipMallocPitch(nullptr, - &pitch_A, width, NUM_H) != hipSuccess); - } - - SECTION("nullptr to pitch") { - REQUIRE(hipMallocPitch(reinterpret_cast(&A_d), - nullptr, width, NUM_H) != hipSuccess); - } -#endif - SECTION("Width 0 in hipMallocPitch") { - REQUIRE(hipMallocPitch(reinterpret_cast(&A_d), - &pitch_A, 0, NUM_H) == hipSuccess); - } - SECTION("Height 0 in hipMallocPitch") { - REQUIRE(hipMallocPitch(reinterpret_cast(&A_d), - &pitch_A, width, 0) == hipSuccess); - } - - SECTION("Max int values") { - REQUIRE(hipMallocPitch(reinterpret_cast(&A_d), - &pitch_A, std::numeric_limits::max(), - std::numeric_limits::max()) != hipSuccess); - } -} -#endif /* * This testcase verifies the basic scenario of * hipMallocPitch API for different datatypes @@ -414,6 +394,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocPitch_Basic", size_t width{NUM_W * sizeof(TestType)}; REQUIRE(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, width, NUM_H) == hipSuccess); + REQUIRE(width <= pitch_A); HIP_CHECK(hipFree(A_d)); } diff --git a/tests/catch/unit/memory/hipPointerGetAttribute.cc b/tests/catch/unit/memory/hipPointerGetAttribute.cc index 393221da11..b0e1e7a5f8 100644 --- a/tests/catch/unit/memory/hipPointerGetAttribute.cc +++ b/tests/catch/unit/memory/hipPointerGetAttribute.cc @@ -316,9 +316,8 @@ TEST_CASE("Unit_hipPointerGetAttribute_Negative") { == hipErrorInvalidValue); } SECTION("Pass invalid attribute") { - hipPointer_attribute attr{HIP_POINTER_ATTRIBUTE_DEVICE_POINTER}; - REQUIRE(hipPointerGetAttribute(&data, attr, - reinterpret_cast(A_h)) == hipErrorInvalidValue); + REQUIRE(hipPointerGetAttribute(&data, static_cast(-1), + reinterpret_cast(A_h)) == hipErrorInvalidValue); } #if HT_AMD SECTION("Pass HIP_POINTER_ATTRIBUTE_IS_GPU_DIRECT_RDMA_CAPABLE"