From 41844f2202152ff05edb17d6d748da7021e8fee6 Mon Sep 17 00:00:00 2001 From: Dino Music Date: Fri, 30 Sep 2022 02:38:41 -0400 Subject: [PATCH 01/15] 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 +#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: + 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 + +#include +#include + +namespace { +inline constexpr size_t kPageSize = 4096; +} // anonymous namespace + +template +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 +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 +__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); +} \ No newline at end of file From 858da0e1ae643b9c8f0347e98ad4ffacd6b5ccbe Mon Sep 17 00:00:00 2001 From: Dino Music Date: Mon, 3 Oct 2022 05:22:42 -0400 Subject: [PATCH 02/15] 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 421488251e0bb702de624943193803f9c4902bae Mon Sep 17 00:00:00 2001 From: Dino Music Date: Mon, 3 Oct 2022 10:13:44 -0400 Subject: [PATCH 03/15] EXSWHTEC-75 - Implement tests for hipMemcpyAsync and derivatives - Implement tests for hipMemcpyAsync - Implement tests for hipMemcpyHtoDAsync - Implement tests for hipMemcpyDtoHAsync - Implement tests for hipMemcpyDtoDAsync --- tests/catch/unit/memory/CMakeLists.txt | 4 + tests/catch/unit/memory/hipMemcpyAsync.cc | 452 ++++-------------- .../unit/memory/hipMemcpyAsync_derivatives.cc | 161 +++++++ tests/catch/unit/memory/hipMemcpyAsync_old.cc | 406 ++++++++++++++++ 4 files changed, 676 insertions(+), 347 deletions(-) create mode 100644 tests/catch/unit/memory/hipMemcpyAsync_derivatives.cc create mode 100644 tests/catch/unit/memory/hipMemcpyAsync_old.cc diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index f24c63ad8c..cd3e433d5a 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -76,6 +76,8 @@ set(TEST_SRC hipHostMalloc.cc hipMemcpy.cc hipMemcpyAsync.cc + hipMemcpyAsync_derivatives.cc + hipMemcpyAsync_old.cc hipMemsetFunctional.cc hipMallocPitch.cc hipMallocArray.cc @@ -152,6 +154,8 @@ set(TEST_SRC hipHostMalloc.cc hipMemcpy.cc hipMemcpyAsync.cc + hipMemcpyAsync_derivatives.cc + hipMemcpyAsync_old.cc hipMemsetFunctional.cc hipMallocPitch.cc hipMallocArray.cc diff --git a/tests/catch/unit/memory/hipMemcpyAsync.cc b/tests/catch/unit/memory/hipMemcpyAsync.cc index b9798f963e..7539f456e0 100644 --- a/tests/catch/unit/memory/hipMemcpyAsync.cc +++ b/tests/catch/unit/memory/hipMemcpyAsync.cc @@ -1,13 +1,15 @@ /* -Copyright (c) 2022 - present 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,390 +19,146 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/* -This testcase verifies the following scenarios -1. hipMemcpyAsync with kernel launch -2. H2D-D2D-D2H-H2PinnMem and device context change scenarios -3. This test launches multiple threads which uses same stream to deploy kernel - and also launch hipMemcpyAsync() api. This test case is simulate the scenario - reported in SWDEV-181598. -*/ +#include "linear_memcpy_tests_common.hh" #include -#include -#include -#include - -#define NUM_THREADS 16 - -static constexpr auto NUM_ELM{1024 * 1024}; - +#include +#include +#include +TEST_CASE("Unit_hipMemcpyAsync_Basic") { + using namespace std::placeholders; + const auto stream_type = GENERATE(Streams::nullstream, Streams::perThread, Streams::created); + const StreamGuard stream_guard(stream_type); + const hipStream_t stream = stream_guard.stream(); -static constexpr size_t N_ELMTS{32 * 1024}; -std::atomic Thread_count { 0 }; -static unsigned blocksPerCU{6}; // to hide latency -static unsigned threadsPerBlock{256}; - -template -void Thread_func(T *A_d, T *B_d, T* C_d, T* C_h, size_t Nbytes, - hipStream_t mystream) { - unsigned blocks = HipTest::setNumBlocks(blocksPerCU, - threadsPerBlock, N_ELMTS); - hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), - dim3(threadsPerBlock), 0, - mystream, A_d, C_d, N_ELMTS); - HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream)); - // The following two MemcpyAsync calls are for sole - // purpose of loading stream with multiple async calls - HIP_CHECK(hipMemcpyAsync(B_d, A_d, Nbytes, - hipMemcpyDeviceToDevice, mystream)); - HIP_CHECK(hipMemcpyAsync(B_d, A_d, Nbytes, - hipMemcpyDeviceToDevice, mystream)); - Thread_count++; + MemcpyWithDirectionCommonTests(std::bind(hipMemcpyAsync, _1, _2, _3, _4, stream), true); } -template -void Thread_func_MultiStream() { - int Data_mismatch = 0; - T *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; - T *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; - size_t Nbytes = N_ELMTS * sizeof(T); - unsigned blocks = HipTest::setNumBlocks(blocksPerCU, - threadsPerBlock, N_ELMTS); - - HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N_ELMTS, false); - hipStream_t mystream; - HIP_CHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking)); - HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream)); - hipLaunchKernelGGL((HipTest::vector_square), dim3(blocks), - dim3(threadsPerBlock), 0, - mystream, A_d, C_d, N_ELMTS); - HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream)); - // The following hipMemcpyAsync() is called only to - // load stream with multiple Async calls - HIP_CHECK(hipMemcpyAsync(B_d, A_d, Nbytes, - hipMemcpyDeviceToDevice, mystream)); - Thread_count++; - HIP_CHECK(hipStreamSynchronize(mystream)); - HIP_CHECK(hipStreamDestroy(mystream)); - // Verifying result of the kernel computation - for (size_t i = 0; i < N_ELMTS; i++) { - if (C_h[i] != A_h[i] * A_h[i]) { - Data_mismatch++; - } +TEST_CASE("Unit_hipMemcpyAsync_Synchronization_Behavior") { + using namespace std::placeholders; + HIP_CHECK(hipDeviceSynchronize()); + + SECTION("Host memory to device memory") { + // This behavior differs on NVIDIA and AMD, on AMD the hipMemcpy calls is synchronous with + // respect to the host +#if HT_AMD + HipTest::HIP_SKIP_TEST( + "EXSWCPHIPT-127 - MemcpyAsync from host to device memory behavior differs on AMD and " + "Nvidia"); + return; +#endif + MemcpyHtoDSyncBehavior(std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyHostToDevice, nullptr), + false); } - // Releasing resources - HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); - REQUIRE(Data_mismatch == 0); -} - -/* -This testcase verifies hipMemcpyAsync API -Initializes device variables -Launches kernel and performs the sum of device variables -copies the result to host variable and validates the result. -*/ -TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_KernelLaunch", "", int, float, - double) { - size_t Nbytes = NUM_ELM * sizeof(TestType); - TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; - TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; - HIP_CHECK(hipSetDevice(0)); - hipStream_t stream; - hipStreamCreate(&stream); - - HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, NUM_ELM, false); - - HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); - HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - - hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), 0, 0, - static_cast(A_d), - static_cast(B_d), C_d, NUM_ELM); - - HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_CHECK(hipStreamDestroy(stream)); - - HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM); - - HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); -} -/* -This testcase verifies the following scenarios -1. H2H,H2PinMem and PinnedMem2Host -2. H2D-D2D-D2H in same GPU -3. Pinned Host Memory to device variables in same GPU -4. Device context change -5. H2D-D2D-D2H peer GPU -*/ -TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_H2H-H2D-D2H-H2PinMem", "", char, int, - float, double) { - TestType *A_d{nullptr}, *B_d{nullptr}; - TestType *A_h{nullptr}, *B_h{nullptr}; - TestType *A_Ph{nullptr}, *B_Ph{nullptr}; - HIP_CHECK(hipSetDevice(0)); - hipStream_t stream; - hipStreamCreate(&stream); - HipTest::initArrays(&A_d, &B_d, nullptr, - &A_h, &B_h, nullptr, - NUM_ELM*sizeof(TestType)); - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_Ph, &B_Ph, nullptr, - NUM_ELM*sizeof(TestType), true); - - SECTION("H2H, H2PinMem and PinMem2H") { - HIP_CHECK(hipMemcpyAsync(B_h, A_h, NUM_ELM*sizeof(TestType), - hipMemcpyHostToHost, stream)); - HIP_CHECK(hipMemcpyAsync(A_Ph, B_h, NUM_ELM*sizeof(TestType), - hipMemcpyHostToHost, stream)); - HIP_CHECK(hipMemcpyAsync(B_Ph, A_Ph, NUM_ELM*sizeof(TestType), - hipMemcpyHostToHost, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - HipTest::checkTest(A_h, B_Ph, NUM_ELM); + SECTION("Device memory to pageable host memory") { + MemcpyDtoHPageableSyncBehavior( + std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyDeviceToHost, nullptr), true); } - SECTION("H2D-D2D-D2H-SameGPU") { - HIP_CHECK(hipMemcpyAsync(A_d, A_h, NUM_ELM*sizeof(TestType), - hipMemcpyHostToDevice, stream)); - HIP_CHECK(hipMemcpyAsync(B_d, A_d, NUM_ELM*sizeof(TestType), - hipMemcpyDeviceToDevice, stream)); - HIP_CHECK(hipMemcpyAsync(B_h, B_d, NUM_ELM*sizeof(TestType), - hipMemcpyDeviceToHost, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - HipTest::checkTest(A_h, B_h, NUM_ELM); + SECTION("Device memory to pinned host memory") { + MemcpyDtoHPinnedSyncBehavior( + std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyDeviceToHost, nullptr), false); } - SECTION("pH2D-D2D-D2pH-SameGPU") { - HIP_CHECK(hipMemcpyAsync(A_d, A_Ph, NUM_ELM*sizeof(TestType), - hipMemcpyHostToDevice, stream)); - HIP_CHECK(hipMemcpyAsync(B_d, A_d, NUM_ELM*sizeof(TestType), - hipMemcpyDeviceToDevice, stream)); - HIP_CHECK(hipMemcpyAsync(B_Ph, B_d, NUM_ELM*sizeof(TestType), - hipMemcpyDeviceToHost, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - HipTest::checkTest(A_Ph, B_Ph, NUM_ELM); - } - SECTION("H2D-D2D-D2H-DeviceContextChange") { - int deviceCount = 0; - HIP_CHECK(hipGetDeviceCount(&deviceCount)); - if (deviceCount < 2) { - SUCCEED("deviceCount less then 2"); - } else { - int canAccessPeer = 0; - HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); - if (canAccessPeer) { - HIP_CHECK(hipSetDevice(1)); - HIP_CHECK(hipMemcpyAsync(A_d, A_h, NUM_ELM*sizeof(TestType), - hipMemcpyHostToDevice, stream)); - HIP_CHECK(hipMemcpyAsync(B_d, A_d, NUM_ELM*sizeof(TestType), - hipMemcpyDeviceToDevice, stream)); - HIP_CHECK(hipMemcpyAsync(B_h, B_d, NUM_ELM*sizeof(TestType), - hipMemcpyDeviceToHost, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - HipTest::checkTest(A_h, B_h, NUM_ELM); - - } else { - SUCCEED("P2P capability is not present"); - } - } + SECTION("Device memory to device memory") { + MemcpyDtoDSyncBehavior(std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyDeviceToDevice, nullptr), + false); } - SECTION("H2D-D2D-D2H-PeerGPU") { - int deviceCount = 0; - HIP_CHECK(hipGetDeviceCount(&deviceCount)); - if (deviceCount < 2) { - SUCCEED("deviceCount less then 2"); - } else { - int canAccessPeer = 0; - HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); - if (canAccessPeer) { - HIP_CHECK(hipSetDevice(1)); - TestType *C_d{nullptr}; - HipTest::initArrays(nullptr, nullptr, &C_d, - nullptr, nullptr, nullptr, - NUM_ELM*sizeof(TestType)); - HIP_CHECK(hipMemcpyAsync(A_d, A_h, NUM_ELM*sizeof(TestType), - hipMemcpyHostToDevice, stream)); - HIP_CHECK(hipMemcpyAsync(C_d, A_d, NUM_ELM*sizeof(TestType), - hipMemcpyDeviceToDevice, stream)); - HIP_CHECK(hipMemcpyAsync(B_h, C_d, NUM_ELM*sizeof(TestType), - hipMemcpyDeviceToHost, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - HipTest::checkTest(A_h, B_h, NUM_ELM); - HIP_CHECK(hipFree(C_d)); - - } else { - SUCCEED("P2P capability is not present"); - } - } + SECTION("Host memory to host memory") { + MemcpyHtoHSyncBehavior(std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyHostToHost, nullptr), + true); } - - HIP_CHECK(hipStreamDestroy(stream)); - - HipTest::freeArrays(A_d, B_d, nullptr, A_h, B_h, nullptr, false); - HipTest::freeArrays(nullptr, nullptr, nullptr, A_Ph, - B_Ph, nullptr, true); } -// This test launches multiple threads which uses same stream to deploy kernel -// and also launch hipMemcpyAsync() api. This test case is simulate the scenario -// reported in SWDEV-181598 +TEST_CASE("Unit_hipMemcpyAsync_Negative_Parameters") { + using namespace std::placeholders; -TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_hipMultiMemcpyMultiThread", "", - int, float, double) { - size_t Nbytes = N_ELMTS * sizeof(TestType); + SECTION("Host to device") { + LinearAllocGuard device_alloc(LinearAllocs::hipMalloc, kPageSize); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, kPageSize); - int Data_mismatch = 0; - hipStream_t mystream; - TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; - TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + MemcpyCommonNegativeTests(std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyHostToDevice, nullptr), + device_alloc.ptr(), host_alloc.ptr(), kPageSize); - HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N_ELMTS, false); - - HIP_CHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking)); - HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream)); + SECTION("Invalid MemcpyKind") { + HIP_CHECK_ERROR(hipMemcpyAsync(device_alloc.ptr(), host_alloc.ptr(), kPageSize, + static_cast(-1), nullptr), + hipErrorInvalidMemcpyDirection); + } - std::thread T[NUM_THREADS]; - for (int i = 0; i < NUM_THREADS; i++) { - T[i] = std::thread(Thread_func, A_d, B_d, C_d, - C_h, Nbytes, mystream); + SECTION("Invalid stream") { + hipStream_t stream; + HIP_CHECK_ERROR(hipMemcpyAsync(device_alloc.ptr(), host_alloc.ptr(), kPageSize, + hipMemcpyHostToDevice, stream), + hipErrorInvalidValue); + } } - // Wait until all the threads finish their execution - for (int i = 0; i < NUM_THREADS; i++) { - T[i].join(); - } + SECTION("Device to host") { + LinearAllocGuard device_alloc(LinearAllocs::hipMalloc, kPageSize); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, kPageSize); - HIP_CHECK(hipStreamSynchronize(mystream)); - HIP_CHECK(hipStreamDestroy(mystream)); + MemcpyCommonNegativeTests(std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyDeviceToHost, nullptr), + host_alloc.ptr(), device_alloc.ptr(), kPageSize); - // Verifying the result of the kernel computation - for (size_t i = 0; i < N_ELMTS; i++) { - if (C_h[i] != A_h[i] * A_h[i]) { - Data_mismatch++; + SECTION("Invalid MemcpyKind") { + HIP_CHECK_ERROR(hipMemcpyAsync(host_alloc.ptr(), device_alloc.ptr(), kPageSize, + static_cast(-1), nullptr), + hipErrorInvalidMemcpyDirection); } - } - REQUIRE(Thread_count.load() == NUM_THREADS); - REQUIRE(Data_mismatch == 0); - HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); - Thread_count.exchange(0); -} -TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_hipMultiMemcpyMultiThreadMultiStream", - "", int, float, double) { - std::thread T[NUM_THREADS]; - for (int i = 0; i < NUM_THREADS; i++) { - T[i] = std::thread(Thread_func_MultiStream); + SECTION("Invalid stream") { + hipStream_t stream; + HIP_CHECK_ERROR(hipMemcpyAsync(host_alloc.ptr(), device_alloc.ptr(), kPageSize, + hipMemcpyDeviceToHost, stream), + hipErrorInvalidValue); + } } - // Wait until all the threads finish their execution - for (int i = 0; i < NUM_THREADS; i++) { - T[i].join(); - } + SECTION("Host to host") { + LinearAllocGuard src_alloc(LinearAllocs::hipHostMalloc, kPageSize); + LinearAllocGuard dst_alloc(LinearAllocs::hipHostMalloc, kPageSize); - REQUIRE(Thread_count.load() == NUM_THREADS); - Thread_count.exchange(0); -} + MemcpyCommonNegativeTests(std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyHostToHost, nullptr), + dst_alloc.ptr(), src_alloc.ptr(), kPageSize); -/* -This testcase verifies hipMemcpy API with pinnedMemory and hostRegister -along with kernel launches -*/ - -TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_PinnedRegMemWithKernelLaunch", - "", int, float, double) { - int numDevices = 0; - HIP_CHECK(hipGetDeviceCount(&numDevices)); - if (numDevices < 2) { - SUCCEED("No of devices are less than 2"); - } else { - // 1 refers to pinned Memory - // 2 refers to register Memory - int MallocPinType = GENERATE(0, 1); - size_t Nbytes = NUM_ELM * sizeof(TestType); - unsigned blocks = HipTest::setNumBlocks(blocksPerCU, - threadsPerBlock, NUM_ELM); - - TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; - TestType *X_d{nullptr}, *Y_d{nullptr}, *Z_d{nullptr}; - TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; - if (MallocPinType) { - HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, NUM_ELM, true); - } else { - A_h = reinterpret_cast(malloc(Nbytes)); - HIP_CHECK(hipHostRegister(A_h, Nbytes, hipHostRegisterDefault)); - B_h = reinterpret_cast(malloc(Nbytes)); - HIP_CHECK(hipHostRegister(B_h, Nbytes, hipHostRegisterDefault)); - C_h = reinterpret_cast(malloc(Nbytes)); - HIP_CHECK(hipHostRegister(C_h, Nbytes, hipHostRegisterDefault)); - HipTest::initArrays(&A_d, &B_d, &C_d, nullptr, nullptr, - nullptr, NUM_ELM, false); - HipTest::setDefaultData(NUM_ELM, A_h, B_h, C_h); + SECTION("Invalid MemcpyKind") { + HIP_CHECK_ERROR(hipMemcpyAsync(dst_alloc.ptr(), src_alloc.ptr(), kPageSize, + static_cast(-1), nullptr), + hipErrorInvalidMemcpyDirection); } - HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); - - hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), - 0, 0, static_cast(A_d), - static_cast(B_d), C_d, NUM_ELM); - - HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); - HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM); - - unsigned int seed = time(0); - HIP_CHECK(hipSetDevice(HipTest::RAND_R(&seed) % (numDevices-1)+1)); - - int device; - HIP_CHECK(hipGetDevice(&device)); - INFO("hipMemcpy is set to happen between device 0 and device " - << device); - HipTest::initArrays(&X_d, &Y_d, &Z_d, nullptr, - nullptr, nullptr, NUM_ELM, false); - hipStream_t gpu1Stream; - HIP_CHECK(hipStreamCreate(&gpu1Stream)); - - for (int j = 0; j < NUM_ELM; j++) { - A_h[j] = 0; - B_h[j] = 0; - C_h[j] = 0; + SECTION("Invalid stream") { + hipStream_t stream; + HIP_CHECK_ERROR( + hipMemcpyAsync(dst_alloc.ptr(), src_alloc.ptr(), kPageSize, hipMemcpyHostToHost, stream), + hipErrorInvalidValue); } + } - hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost); - hipMemcpyAsync(X_d, A_h, Nbytes, hipMemcpyHostToDevice, gpu1Stream); - hipMemcpy(B_h, B_d, Nbytes, hipMemcpyDeviceToHost); - hipMemcpyAsync(Y_d, B_h, Nbytes, hipMemcpyHostToDevice, gpu1Stream); - - hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), - 0, 0, static_cast(X_d), - static_cast(Y_d), Z_d, NUM_ELM); + SECTION("Device to device") { + LinearAllocGuard src_alloc(LinearAllocs::hipMalloc, kPageSize); + LinearAllocGuard dst_alloc(LinearAllocs::hipMalloc, kPageSize); - HIP_CHECK(hipMemcpyAsync(C_h, Z_d, Nbytes, - hipMemcpyDeviceToHost, gpu1Stream)); - HIP_CHECK(hipStreamSynchronize(gpu1Stream)); + MemcpyCommonNegativeTests( + std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyDeviceToDevice, nullptr), dst_alloc.ptr(), + src_alloc.ptr(), kPageSize); - HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM); + SECTION("Invalid MemcpyKind") { + HIP_CHECK_ERROR(hipMemcpyAsync(src_alloc.ptr(), dst_alloc.ptr(), kPageSize, + static_cast(-1), nullptr), + hipErrorInvalidMemcpyDirection); + } - if (MallocPinType) { - HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, true); - } else { - HIP_CHECK(hipHostUnregister(A_h)); - free(A_h); - HIP_CHECK(hipHostUnregister(B_h)); - free(B_h); - HIP_CHECK(hipHostUnregister(C_h)); - free(C_h); - HipTest::freeArrays(A_d, B_d, C_d, nullptr, - nullptr, nullptr, false); + SECTION("Invalid stream") { + hipStream_t stream; + HIP_CHECK_ERROR(hipMemcpyAsync(dst_alloc.ptr(), src_alloc.ptr(), kPageSize, + hipMemcpyDeviceToDevice, stream), + hipErrorInvalidValue); } - HipTest::freeArrays(X_d, Y_d, Z_d, nullptr, - nullptr, nullptr, false); - HIP_CHECK(hipStreamDestroy(gpu1Stream)); } -} - +} \ No newline at end of file diff --git a/tests/catch/unit/memory/hipMemcpyAsync_derivatives.cc b/tests/catch/unit/memory/hipMemcpyAsync_derivatives.cc new file mode 100644 index 0000000000..19f2c1e3cf --- /dev/null +++ b/tests/catch/unit/memory/hipMemcpyAsync_derivatives.cc @@ -0,0 +1,161 @@ +/* +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. +*/ + +#include "linear_memcpy_tests_common.hh" + +#include +#include +#include +#include + +TEST_CASE("Unit_hipMemcpyDtoHAsync_Basic") { + const auto stream_type = GENERATE(Streams::nullstream, Streams::perThread, Streams::created); + const StreamGuard stream_guard(stream_type); + + const auto f = [stream = stream_guard.stream()](void* dst, void* src, size_t count) { + return hipMemcpyDtoHAsync(dst, reinterpret_cast(src), count, stream); + }; + MemcpyDeviceToHostShell(f, true, stream_guard.stream()); +} + +TEST_CASE("Unit_hipMemcpyDtoHAsync_Synchronization_Behavior") { + HIP_CHECK(hipDeviceSynchronize()); + + SECTION("Device memory to pageable host memory") { + MemcpyDtoHPageableSyncBehavior( + [](void* dst, void* src, size_t count) { + return hipMemcpyDtoHAsync(dst, reinterpret_cast(src), count, nullptr); + }, + true); + } + + SECTION("Device memory to pinned host memory") { + MemcpyDtoHPinnedSyncBehavior( + [](void* dst, void* src, size_t count) { + return hipMemcpyDtoHAsync(dst, reinterpret_cast(src), count, nullptr); + }, + false); + } +} + +TEST_CASE("Unit_hipMemcpyDtoHAsync_Negative_Parameters") { + using namespace std::placeholders; + LinearAllocGuard device_alloc(LinearAllocs::hipMalloc, kPageSize); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, kPageSize); + + MemcpyCommonNegativeTests( + [](void* dst, void* src, size_t count) { + return hipMemcpyDtoHAsync(dst, reinterpret_cast(src), count, nullptr); + }, + host_alloc.ptr(), device_alloc.ptr(), kPageSize); + + SECTION("Invalid stream") { + hipStream_t stream; + HIP_CHECK_ERROR(hipMemcpyDtoHAsync(host_alloc.ptr(), device_alloc.ptr(), kPageSize, stream), + hipErrorInvalidValue); + } +} + +TEST_CASE("Unit_hipMemcpyHtoDAsync_Basic") { + const auto stream_type = GENERATE(Streams::nullstream, Streams::perThread, Streams::created); + const StreamGuard stream_guard(stream_type); + + const auto f = [stream = stream_guard.stream()](void* dst, void* src, size_t count) { + return hipMemcpyHtoDAsync(reinterpret_cast(dst), src, count, stream); + }; + MemcpyHostToDeviceShell(f, true, stream_guard.stream()); +} + +TEST_CASE("Unit_hipMemcpyHtoDAsync_Synchronization_Behavior") { + // This behavior differs on NVIDIA and AMD, on AMD the hipMemcpy calls is synchronous with + // respect to the host +#if HT_AMD + HipTest::HIP_SKIP_TEST( + "EXSWCPHIPT-127 - MemcpyAsync from host to device memory behavior differs on AMD and " + "Nvidia"); + return; +#endif + MemcpyHtoDSyncBehavior( + [](void* dst, void* src, size_t count) { + return hipMemcpyHtoDAsync(reinterpret_cast(dst), src, count, nullptr); + }, + false); +} + +TEST_CASE("Unit_hipMemcpyHtoDAsync_Negative_Parameters") { + using namespace std::placeholders; + LinearAllocGuard device_alloc(LinearAllocs::hipMalloc, kPageSize); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, kPageSize); + + MemcpyCommonNegativeTests( + [](void* dst, void* src, size_t count) { + return hipMemcpyHtoDAsync(reinterpret_cast(dst), src, count, nullptr); + }, + device_alloc.ptr(), host_alloc.ptr(), kPageSize); + + SECTION("Invalid stream") { + hipStream_t stream; + HIP_CHECK_ERROR(hipMemcpyHtoDAsync(device_alloc.ptr(), host_alloc.ptr(), kPageSize, stream), + hipErrorInvalidValue); + } +} + +TEST_CASE("Unit_hipMemcpyDtoDAsync_Basic") { + const auto stream_type = GENERATE(Streams::nullstream, Streams::perThread, Streams::created); + const StreamGuard stream_guard(stream_type); + + SECTION("Device to device") { + MemcpyDeviceToDeviceShell( + [stream = stream_guard.stream()](void* dst, void* src, size_t count) { + return hipMemcpyDtoDAsync(reinterpret_cast(dst), + reinterpret_cast(src), count, stream); + }, + true); + } +} + +TEST_CASE("Unit_hipMemcpyDtoDAsync_Synchronization_Behavior") { + MemcpyDtoDSyncBehavior( + [](void* dst, void* src, size_t count) { + return hipMemcpyDtoDAsync(reinterpret_cast(dst), + reinterpret_cast(src), count, nullptr); + }, + false); +} + +TEST_CASE("Unit_hipMemcpyDtoDAsync_Negative_Parameters") { + using namespace std::placeholders; + LinearAllocGuard src_alloc(LinearAllocs::hipMalloc, kPageSize); + LinearAllocGuard dst_alloc(LinearAllocs::hipMalloc, kPageSize); + + MemcpyCommonNegativeTests( + [](void* dst, void* src, size_t count) { + return hipMemcpyDtoDAsync(reinterpret_cast(dst), + reinterpret_cast(src), count, nullptr); + }, + dst_alloc.ptr(), src_alloc.ptr(), kPageSize); + + SECTION("Invalid stream") { + hipStream_t stream; + HIP_CHECK_ERROR(hipMemcpyDtoDAsync(dst_alloc.ptr(), src_alloc.ptr(), kPageSize, stream), + hipErrorInvalidValue); + } +} \ No newline at end of file diff --git a/tests/catch/unit/memory/hipMemcpyAsync_old.cc b/tests/catch/unit/memory/hipMemcpyAsync_old.cc new file mode 100644 index 0000000000..b9798f963e --- /dev/null +++ b/tests/catch/unit/memory/hipMemcpyAsync_old.cc @@ -0,0 +1,406 @@ +/* +Copyright (c) 2022 - present 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. +*/ + +/* +This testcase verifies the following scenarios +1. hipMemcpyAsync with kernel launch +2. H2D-D2D-D2H-H2PinnMem and device context change scenarios +3. This test launches multiple threads which uses same stream to deploy kernel + and also launch hipMemcpyAsync() api. This test case is simulate the scenario + reported in SWDEV-181598. +*/ + +#include +#include +#include +#include + +#define NUM_THREADS 16 + +static constexpr auto NUM_ELM{1024 * 1024}; + + + +static constexpr size_t N_ELMTS{32 * 1024}; +std::atomic Thread_count { 0 }; +static unsigned blocksPerCU{6}; // to hide latency +static unsigned threadsPerBlock{256}; + +template +void Thread_func(T *A_d, T *B_d, T* C_d, T* C_h, size_t Nbytes, + hipStream_t mystream) { + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, + threadsPerBlock, N_ELMTS); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, + mystream, A_d, C_d, N_ELMTS); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream)); + // The following two MemcpyAsync calls are for sole + // purpose of loading stream with multiple async calls + HIP_CHECK(hipMemcpyAsync(B_d, A_d, Nbytes, + hipMemcpyDeviceToDevice, mystream)); + HIP_CHECK(hipMemcpyAsync(B_d, A_d, Nbytes, + hipMemcpyDeviceToDevice, mystream)); + Thread_count++; +} + +template +void Thread_func_MultiStream() { + int Data_mismatch = 0; + T *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + T *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + size_t Nbytes = N_ELMTS * sizeof(T); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, + threadsPerBlock, N_ELMTS); + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N_ELMTS, false); + hipStream_t mystream; + HIP_CHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream)); + hipLaunchKernelGGL((HipTest::vector_square), dim3(blocks), + dim3(threadsPerBlock), 0, + mystream, A_d, C_d, N_ELMTS); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream)); + // The following hipMemcpyAsync() is called only to + // load stream with multiple Async calls + HIP_CHECK(hipMemcpyAsync(B_d, A_d, Nbytes, + hipMemcpyDeviceToDevice, mystream)); + Thread_count++; + + HIP_CHECK(hipStreamSynchronize(mystream)); + HIP_CHECK(hipStreamDestroy(mystream)); + // Verifying result of the kernel computation + for (size_t i = 0; i < N_ELMTS; i++) { + if (C_h[i] != A_h[i] * A_h[i]) { + Data_mismatch++; + } + } + // Releasing resources + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + REQUIRE(Data_mismatch == 0); +} + +/* +This testcase verifies hipMemcpyAsync API +Initializes device variables +Launches kernel and performs the sum of device variables +copies the result to host variable and validates the result. +*/ +TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_KernelLaunch", "", int, float, + double) { + size_t Nbytes = NUM_ELM * sizeof(TestType); + + TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + HIP_CHECK(hipSetDevice(0)); + hipStream_t stream; + hipStreamCreate(&stream); + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, NUM_ELM, false); + + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), 0, 0, + static_cast(A_d), + static_cast(B_d), C_d, NUM_ELM); + + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamDestroy(stream)); + + HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); +} +/* +This testcase verifies the following scenarios +1. H2H,H2PinMem and PinnedMem2Host +2. H2D-D2D-D2H in same GPU +3. Pinned Host Memory to device variables in same GPU +4. Device context change +5. H2D-D2D-D2H peer GPU +*/ +TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_H2H-H2D-D2H-H2PinMem", "", char, int, + float, double) { + TestType *A_d{nullptr}, *B_d{nullptr}; + TestType *A_h{nullptr}, *B_h{nullptr}; + TestType *A_Ph{nullptr}, *B_Ph{nullptr}; + HIP_CHECK(hipSetDevice(0)); + hipStream_t stream; + hipStreamCreate(&stream); + HipTest::initArrays(&A_d, &B_d, nullptr, + &A_h, &B_h, nullptr, + NUM_ELM*sizeof(TestType)); + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_Ph, &B_Ph, nullptr, + NUM_ELM*sizeof(TestType), true); + + SECTION("H2H, H2PinMem and PinMem2H") { + HIP_CHECK(hipMemcpyAsync(B_h, A_h, NUM_ELM*sizeof(TestType), + hipMemcpyHostToHost, stream)); + HIP_CHECK(hipMemcpyAsync(A_Ph, B_h, NUM_ELM*sizeof(TestType), + hipMemcpyHostToHost, stream)); + HIP_CHECK(hipMemcpyAsync(B_Ph, A_Ph, NUM_ELM*sizeof(TestType), + hipMemcpyHostToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HipTest::checkTest(A_h, B_Ph, NUM_ELM); + } + + SECTION("H2D-D2D-D2H-SameGPU") { + HIP_CHECK(hipMemcpyAsync(A_d, A_h, NUM_ELM*sizeof(TestType), + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(B_d, A_d, NUM_ELM*sizeof(TestType), + hipMemcpyDeviceToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(B_h, B_d, NUM_ELM*sizeof(TestType), + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + } + + SECTION("pH2D-D2D-D2pH-SameGPU") { + HIP_CHECK(hipMemcpyAsync(A_d, A_Ph, NUM_ELM*sizeof(TestType), + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(B_d, A_d, NUM_ELM*sizeof(TestType), + hipMemcpyDeviceToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(B_Ph, B_d, NUM_ELM*sizeof(TestType), + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HipTest::checkTest(A_Ph, B_Ph, NUM_ELM); + } + SECTION("H2D-D2D-D2H-DeviceContextChange") { + int deviceCount = 0; + HIP_CHECK(hipGetDeviceCount(&deviceCount)); + if (deviceCount < 2) { + SUCCEED("deviceCount less then 2"); + } else { + int canAccessPeer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, NUM_ELM*sizeof(TestType), + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(B_d, A_d, NUM_ELM*sizeof(TestType), + hipMemcpyDeviceToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(B_h, B_d, NUM_ELM*sizeof(TestType), + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + + } else { + SUCCEED("P2P capability is not present"); + } + } + } + + SECTION("H2D-D2D-D2H-PeerGPU") { + int deviceCount = 0; + HIP_CHECK(hipGetDeviceCount(&deviceCount)); + if (deviceCount < 2) { + SUCCEED("deviceCount less then 2"); + } else { + int canAccessPeer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + HIP_CHECK(hipSetDevice(1)); + TestType *C_d{nullptr}; + HipTest::initArrays(nullptr, nullptr, &C_d, + nullptr, nullptr, nullptr, + NUM_ELM*sizeof(TestType)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, NUM_ELM*sizeof(TestType), + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(C_d, A_d, NUM_ELM*sizeof(TestType), + hipMemcpyDeviceToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(B_h, C_d, NUM_ELM*sizeof(TestType), + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + HIP_CHECK(hipFree(C_d)); + + } else { + SUCCEED("P2P capability is not present"); + } + } + } + + HIP_CHECK(hipStreamDestroy(stream)); + + HipTest::freeArrays(A_d, B_d, nullptr, A_h, B_h, nullptr, false); + HipTest::freeArrays(nullptr, nullptr, nullptr, A_Ph, + B_Ph, nullptr, true); +} + +// This test launches multiple threads which uses same stream to deploy kernel +// and also launch hipMemcpyAsync() api. This test case is simulate the scenario +// reported in SWDEV-181598 + +TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_hipMultiMemcpyMultiThread", "", + int, float, double) { + size_t Nbytes = N_ELMTS * sizeof(TestType); + + int Data_mismatch = 0; + hipStream_t mystream; + TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N_ELMTS, false); + + HIP_CHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream)); + + std::thread T[NUM_THREADS]; + for (int i = 0; i < NUM_THREADS; i++) { + T[i] = std::thread(Thread_func, A_d, B_d, C_d, + C_h, Nbytes, mystream); + } + + // Wait until all the threads finish their execution + for (int i = 0; i < NUM_THREADS; i++) { + T[i].join(); + } + + HIP_CHECK(hipStreamSynchronize(mystream)); + HIP_CHECK(hipStreamDestroy(mystream)); + + // Verifying the result of the kernel computation + for (size_t i = 0; i < N_ELMTS; i++) { + if (C_h[i] != A_h[i] * A_h[i]) { + Data_mismatch++; + } + } + REQUIRE(Thread_count.load() == NUM_THREADS); + REQUIRE(Data_mismatch == 0); + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + Thread_count.exchange(0); +} + +TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_hipMultiMemcpyMultiThreadMultiStream", + "", int, float, double) { + std::thread T[NUM_THREADS]; + for (int i = 0; i < NUM_THREADS; i++) { + T[i] = std::thread(Thread_func_MultiStream); + } + + // Wait until all the threads finish their execution + for (int i = 0; i < NUM_THREADS; i++) { + T[i].join(); + } + + REQUIRE(Thread_count.load() == NUM_THREADS); + Thread_count.exchange(0); +} + +/* +This testcase verifies hipMemcpy API with pinnedMemory and hostRegister +along with kernel launches +*/ + +TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_PinnedRegMemWithKernelLaunch", + "", int, float, double) { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices < 2) { + SUCCEED("No of devices are less than 2"); + } else { + // 1 refers to pinned Memory + // 2 refers to register Memory + int MallocPinType = GENERATE(0, 1); + size_t Nbytes = NUM_ELM * sizeof(TestType); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, + threadsPerBlock, NUM_ELM); + + TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + TestType *X_d{nullptr}, *Y_d{nullptr}, *Z_d{nullptr}; + TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + if (MallocPinType) { + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, NUM_ELM, true); + } else { + A_h = reinterpret_cast(malloc(Nbytes)); + HIP_CHECK(hipHostRegister(A_h, Nbytes, hipHostRegisterDefault)); + B_h = reinterpret_cast(malloc(Nbytes)); + HIP_CHECK(hipHostRegister(B_h, Nbytes, hipHostRegisterDefault)); + C_h = reinterpret_cast(malloc(Nbytes)); + HIP_CHECK(hipHostRegister(C_h, Nbytes, hipHostRegisterDefault)); + HipTest::initArrays(&A_d, &B_d, &C_d, nullptr, nullptr, + nullptr, NUM_ELM, false); + HipTest::setDefaultData(NUM_ELM, A_h, B_h, C_h); + } + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), + 0, 0, static_cast(A_d), + static_cast(B_d), C_d, NUM_ELM); + + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM); + + unsigned int seed = time(0); + HIP_CHECK(hipSetDevice(HipTest::RAND_R(&seed) % (numDevices-1)+1)); + + int device; + HIP_CHECK(hipGetDevice(&device)); + INFO("hipMemcpy is set to happen between device 0 and device " + << device); + HipTest::initArrays(&X_d, &Y_d, &Z_d, nullptr, + nullptr, nullptr, NUM_ELM, false); + + hipStream_t gpu1Stream; + HIP_CHECK(hipStreamCreate(&gpu1Stream)); + + for (int j = 0; j < NUM_ELM; j++) { + A_h[j] = 0; + B_h[j] = 0; + C_h[j] = 0; + } + + hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost); + hipMemcpyAsync(X_d, A_h, Nbytes, hipMemcpyHostToDevice, gpu1Stream); + hipMemcpy(B_h, B_d, Nbytes, hipMemcpyDeviceToHost); + hipMemcpyAsync(Y_d, B_h, Nbytes, hipMemcpyHostToDevice, gpu1Stream); + + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), + 0, 0, static_cast(X_d), + static_cast(Y_d), Z_d, NUM_ELM); + + HIP_CHECK(hipMemcpyAsync(C_h, Z_d, Nbytes, + hipMemcpyDeviceToHost, gpu1Stream)); + HIP_CHECK(hipStreamSynchronize(gpu1Stream)); + + HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM); + + if (MallocPinType) { + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, true); + } else { + HIP_CHECK(hipHostUnregister(A_h)); + free(A_h); + HIP_CHECK(hipHostUnregister(B_h)); + free(B_h); + HIP_CHECK(hipHostUnregister(C_h)); + free(C_h); + HipTest::freeArrays(A_d, B_d, C_d, nullptr, + nullptr, nullptr, false); + } + HipTest::freeArrays(X_d, Y_d, Z_d, nullptr, + nullptr, nullptr, false); + HIP_CHECK(hipStreamDestroy(gpu1Stream)); + } +} + From e534d2e050446f4e936c0de76fc81b5858782654 Mon Sep 17 00:00:00 2001 From: Dino Music Date: Thu, 6 Oct 2022 15:47:34 +0200 Subject: [PATCH 04/15] 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 05/15] 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 06/15] 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 07/15] 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 08/15] 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 09/15] 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 10/15] 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 11/15] 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 12/15] 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 13/15] 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 e1e12091185a79cb1047351f3fb1600246838b0c Mon Sep 17 00:00:00 2001 From: Dino Music Date: Fri, 14 Oct 2022 17:01:25 +0200 Subject: [PATCH 14/15] EXSWHTEC-75 - Update invalid stream negative tests and bring in line with changes to memcpy1d_tests_common --- tests/catch/unit/memory/hipMemcpyAsync.cc | 34 +++++----- .../unit/memory/hipMemcpyAsync_derivatives.cc | 62 +++++++++++-------- 2 files changed, 53 insertions(+), 43 deletions(-) diff --git a/tests/catch/unit/memory/hipMemcpyAsync.cc b/tests/catch/unit/memory/hipMemcpyAsync.cc index 7539f456e0..75a695f37e 100644 --- a/tests/catch/unit/memory/hipMemcpyAsync.cc +++ b/tests/catch/unit/memory/hipMemcpyAsync.cc @@ -19,24 +19,24 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include "linear_memcpy_tests_common.hh" +#include "memcpy1d_tests_common.hh" #include #include #include #include -TEST_CASE("Unit_hipMemcpyAsync_Basic") { +TEST_CASE("Unit_hipMemcpyAsync_Positive_Basic") { using namespace std::placeholders; const auto stream_type = GENERATE(Streams::nullstream, Streams::perThread, Streams::created); const StreamGuard stream_guard(stream_type); const hipStream_t stream = stream_guard.stream(); - MemcpyWithDirectionCommonTests(std::bind(hipMemcpyAsync, _1, _2, _3, _4, stream), true); + MemcpyWithDirectionCommonTests(std::bind(hipMemcpyAsync, _1, _2, _3, _4, stream)); } -TEST_CASE("Unit_hipMemcpyAsync_Synchronization_Behavior") { +TEST_CASE("Unit_hipMemcpyAsync_Positive_Synchronization_Behavior") { using namespace std::placeholders; HIP_CHECK(hipDeviceSynchronize()); @@ -76,6 +76,10 @@ TEST_CASE("Unit_hipMemcpyAsync_Synchronization_Behavior") { TEST_CASE("Unit_hipMemcpyAsync_Negative_Parameters") { using namespace std::placeholders; + constexpr auto InvalidStream = [] { + StreamGuard sg(Streams::created); + return sg.stream(); + }; SECTION("Host to device") { LinearAllocGuard device_alloc(LinearAllocs::hipMalloc, kPageSize); @@ -91,10 +95,9 @@ TEST_CASE("Unit_hipMemcpyAsync_Negative_Parameters") { } SECTION("Invalid stream") { - hipStream_t stream; HIP_CHECK_ERROR(hipMemcpyAsync(device_alloc.ptr(), host_alloc.ptr(), kPageSize, - hipMemcpyHostToDevice, stream), - hipErrorInvalidValue); + hipMemcpyHostToDevice, InvalidStream()), + hipErrorContextIsDestroyed); } } @@ -112,10 +115,9 @@ TEST_CASE("Unit_hipMemcpyAsync_Negative_Parameters") { } SECTION("Invalid stream") { - hipStream_t stream; HIP_CHECK_ERROR(hipMemcpyAsync(host_alloc.ptr(), device_alloc.ptr(), kPageSize, - hipMemcpyDeviceToHost, stream), - hipErrorInvalidValue); + hipMemcpyDeviceToHost, InvalidStream()), + hipErrorContextIsDestroyed); } } @@ -133,10 +135,9 @@ TEST_CASE("Unit_hipMemcpyAsync_Negative_Parameters") { } SECTION("Invalid stream") { - hipStream_t stream; - HIP_CHECK_ERROR( - hipMemcpyAsync(dst_alloc.ptr(), src_alloc.ptr(), kPageSize, hipMemcpyHostToHost, stream), - hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpyAsync(dst_alloc.ptr(), src_alloc.ptr(), kPageSize, + hipMemcpyHostToHost, InvalidStream()), + hipErrorContextIsDestroyed); } } @@ -155,10 +156,9 @@ TEST_CASE("Unit_hipMemcpyAsync_Negative_Parameters") { } SECTION("Invalid stream") { - hipStream_t stream; HIP_CHECK_ERROR(hipMemcpyAsync(dst_alloc.ptr(), src_alloc.ptr(), kPageSize, - hipMemcpyDeviceToDevice, stream), - hipErrorInvalidValue); + hipMemcpyDeviceToDevice, InvalidStream()), + hipErrorContextIsDestroyed); } } } \ No newline at end of file diff --git a/tests/catch/unit/memory/hipMemcpyAsync_derivatives.cc b/tests/catch/unit/memory/hipMemcpyAsync_derivatives.cc index 3a67f0e457..cfb23f705c 100644 --- a/tests/catch/unit/memory/hipMemcpyAsync_derivatives.cc +++ b/tests/catch/unit/memory/hipMemcpyAsync_derivatives.cc @@ -19,24 +19,29 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include "linear_memcpy_tests_common.hh" +#include "memcpy1d_tests_common.hh" #include #include #include #include -TEST_CASE("Unit_hipMemcpyDtoHAsync_Basic") { +static hipStream_t InvalidStream() { + StreamGuard sg(Streams::created); + return sg.stream(); +} + +TEST_CASE("Unit_hipMemcpyDtoHAsync_Positive_Basic") { const auto stream_type = GENERATE(Streams::nullstream, Streams::perThread, Streams::created); const StreamGuard stream_guard(stream_type); const auto f = [stream = stream_guard.stream()](void* dst, void* src, size_t count) { return hipMemcpyDtoHAsync(dst, reinterpret_cast(src), count, stream); }; - MemcpyDeviceToHostShell(f, true, stream_guard.stream()); + MemcpyDeviceToHostShell(f, stream_guard.stream()); } -TEST_CASE("Unit_hipMemcpyDtoHAsync_Synchronization_Behavior") { +TEST_CASE("Unit_hipMemcpyDtoHAsync_Positive_Synchronization_Behavior") { HIP_CHECK(hipDeviceSynchronize()); SECTION("Device memory to pageable host memory") { @@ -68,25 +73,24 @@ TEST_CASE("Unit_hipMemcpyDtoHAsync_Negative_Parameters") { host_alloc.ptr(), device_alloc.ptr(), kPageSize); SECTION("Invalid stream") { - hipStream_t stream; HIP_CHECK_ERROR( hipMemcpyDtoHAsync(host_alloc.ptr(), reinterpret_cast(device_alloc.ptr()), - kPageSize, stream), - hipErrorInvalidValue); + kPageSize, InvalidStream()), + hipErrorContextIsDestroyed); } } -TEST_CASE("Unit_hipMemcpyHtoDAsync_Basic") { +TEST_CASE("Unit_hipMemcpyHtoDAsync_Positive_Basic") { const auto stream_type = GENERATE(Streams::nullstream, Streams::perThread, Streams::created); const StreamGuard stream_guard(stream_type); const auto f = [stream = stream_guard.stream()](void* dst, void* src, size_t count) { return hipMemcpyHtoDAsync(reinterpret_cast(dst), src, count, stream); }; - MemcpyHostToDeviceShell(f, true, stream_guard.stream()); + MemcpyHostToDeviceShell(f, stream_guard.stream()); } -TEST_CASE("Unit_hipMemcpyHtoDAsync_Synchronization_Behavior") { +TEST_CASE("Unit_hipMemcpyHtoDAsync_Positive_Synchronization_Behavior") { // This behavior differs on NVIDIA and AMD, on AMD the hipMemcpy calls is synchronous with // respect to the host #if HT_AMD @@ -114,28 +118,35 @@ TEST_CASE("Unit_hipMemcpyHtoDAsync_Negative_Parameters") { device_alloc.ptr(), host_alloc.ptr(), kPageSize); SECTION("Invalid stream") { - hipStream_t stream; HIP_CHECK_ERROR(hipMemcpyHtoDAsync(reinterpret_cast(device_alloc.ptr()), - host_alloc.ptr(), kPageSize, stream), - hipErrorInvalidValue); + host_alloc.ptr(), kPageSize, InvalidStream()), + hipErrorContextIsDestroyed); } } -TEST_CASE("Unit_hipMemcpyDtoDAsync_Basic") { +TEST_CASE("Unit_hipMemcpyDtoDAsync_Positive_Basic") { const auto stream_type = GENERATE(Streams::nullstream, Streams::perThread, Streams::created); const StreamGuard stream_guard(stream_type); SECTION("Device to device") { - MemcpyDeviceToDeviceShell( - [stream = stream_guard.stream()](void* dst, void* src, size_t count) { - return hipMemcpyDtoDAsync(reinterpret_cast(dst), - reinterpret_cast(src), count, stream); - }, - true); + SECTION("Peer access enabled") { + MemcpyDeviceToDeviceShell( + [stream = stream_guard.stream()](void* dst, void* src, size_t count) { + return hipMemcpyDtoDAsync(reinterpret_cast(dst), + reinterpret_cast(src), count, stream); + }); + } + SECTION("Peer access disabled") { + MemcpyDeviceToDeviceShell( + [stream = stream_guard.stream()](void* dst, void* src, size_t count) { + return hipMemcpyDtoDAsync(reinterpret_cast(dst), + reinterpret_cast(src), count, stream); + }); + } } } -TEST_CASE("Unit_hipMemcpyDtoDAsync_Synchronization_Behavior") { +TEST_CASE("Unit_hipMemcpyDtoDAsync_Positive_Synchronization_Behavior") { MemcpyDtoDSyncBehavior( [](void* dst, void* src, size_t count) { return hipMemcpyDtoDAsync(reinterpret_cast(dst), @@ -157,10 +168,9 @@ TEST_CASE("Unit_hipMemcpyDtoDAsync_Negative_Parameters") { dst_alloc.ptr(), src_alloc.ptr(), kPageSize); SECTION("Invalid stream") { - hipStream_t stream; - HIP_CHECK_ERROR( - hipMemcpyDtoDAsync(reinterpret_cast(dst_alloc.ptr()), - reinterpret_cast(src_alloc.ptr()), kPageSize, stream), - hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpyDtoDAsync(reinterpret_cast(dst_alloc.ptr()), + reinterpret_cast(src_alloc.ptr()), kPageSize, + InvalidStream()), + hipErrorContextIsDestroyed); } } \ No newline at end of file From 76c8e3104c5881e6469b9229d2226011e175e4b7 Mon Sep 17 00:00:00 2001 From: Dino Music Date: Fri, 14 Oct 2022 19:47:44 +0200 Subject: [PATCH 15/15] 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