From c062ab6c593f9bc4a2bfccb14c94495de9a01e96 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 21 Apr 2023 16:44:33 +0530 Subject: [PATCH] SWDEV-375580 - rework the test to be more stable, move it to correct folder and fix some multi threaded access patterns (#211) --- catch/unit/memory/CMakeLists.txt | 1 - .../unit/memory/hipMemsetAsyncMultiThread.cc | 243 ------------------ catch/unit/multiThread/CMakeLists.txt | 1 + .../multiThread/hipMemsetAsyncMultiThread.cc | 182 +++++++++++++ 4 files changed, 183 insertions(+), 244 deletions(-) delete mode 100644 catch/unit/memory/hipMemsetAsyncMultiThread.cc create mode 100644 catch/unit/multiThread/hipMemsetAsyncMultiThread.cc diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index 6c528fcc16..dfecc3471f 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -51,7 +51,6 @@ set(TEST_SRC hipMallocManaged_MultiScenario.cc hipMemsetNegative.cc hipMemset.cc - hipMemsetAsyncMultiThread.cc hipMemset3D.cc hipMemset2D.cc hipHostMallocTests.cc diff --git a/catch/unit/memory/hipMemsetAsyncMultiThread.cc b/catch/unit/memory/hipMemsetAsyncMultiThread.cc deleted file mode 100644 index 68deacda92..0000000000 --- a/catch/unit/memory/hipMemsetAsyncMultiThread.cc +++ /dev/null @@ -1,243 +0,0 @@ -/* - * Copyright (c) 2021 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 WARRANNTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - * THE SOFTWARE. - */ - -/* - * Test that validates functionality of hipmemsetAsync apis over multi threads - */ - -#include -#include - - -#define NUM_THREADS 20 -#define ITER 10 -#define N (4*1024*1024) - - -template -class MemSetAsyncMthreadTest { - public: - T *A_h, *A_d, *B_h; - T memSetVal; - size_t Nbytes; - bool testResult = true; - int validateCount = 0; - hipStream_t stream; - - void memAllocate(T memSetValue) { - memSetVal = memSetValue; - Nbytes = N * sizeof(T); - - A_h = reinterpret_cast(malloc(Nbytes)); - HIP_ASSERT(A_h != nullptr); - - HIP_CHECK(hipMalloc(&A_d, Nbytes)); - B_h = reinterpret_cast(malloc(Nbytes)); - HIP_ASSERT(B_h != nullptr); - - HIP_CHECK(hipStreamCreate(&stream)); - } - - void threadCompleteStatus() { - for (int k = 0 ; k < N ; k++) { - if ((A_h[k] == memSetVal) && (B_h[k] == memSetVal)) { - validateCount+= 1; - } - } - } - - bool resultAfterAllIterations() { - memDeallocate(); - testResult = (validateCount == (ITER * N)) ? true: false; - return testResult; - } - - void memDeallocate() { - HIP_CHECK(hipFree(A_d)); - free(A_h); - free(B_h); - HIP_CHECK(hipStreamDestroy(stream)); - } -}; - -template -void queueJobsForhipMemsetAsync(T* A_d, T* A_h, T memSetVal, size_t Nbytes, - hipStream_t stream) { - HIPCHECK(hipMemsetAsync(A_d, memSetVal, N, stream)); - HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream)); -} - -template -void queueJobsForhipMemsetD32Async(T* A_d, T* A_h, T memSetVal, size_t Nbytes, - hipStream_t stream) { - HIPCHECK(hipMemsetD32Async((hipDeviceptr_t)A_d, memSetVal, N, stream)); - HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream)); -} - -template -void queueJobsForhipMemsetD16Async(T* A_d, T* A_h, T memSetVal, size_t Nbytes, - hipStream_t stream) { - HIPCHECK(hipMemsetD16Async((hipDeviceptr_t)A_d, memSetVal, N, stream)); - HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream)); -} - -template -void queueJobsForhipMemsetD8Async(T* A_d, T* A_h, T memSetVal, size_t Nbytes, - hipStream_t stream) { - HIPCHECK(hipMemsetD8Async((hipDeviceptr_t)A_d, memSetVal, N, stream)); - HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream)); -} - -/* Queue hipMemsetAsync jobs on multiple threads and verify they all - * finished on all threads successfully - */ -bool testhipMemsetAsyncWithMultiThread() { - MemSetAsyncMthreadTest obj; - constexpr char memsetval = 0x42; - obj.memAllocate(memsetval); - std::thread t[NUM_THREADS]; - - for (int i = 0 ; i < ITER ; i++) { - for (int k = 0 ; k < NUM_THREADS ; k++) { - if (k%2) { - t[k] = std::thread(queueJobsForhipMemsetAsync, obj.A_d, obj.A_h, - obj.memSetVal, obj.Nbytes, obj.stream); - } else { - t[k] = std::thread(queueJobsForhipMemsetAsync, obj.A_d, obj.B_h, - obj.memSetVal, obj.Nbytes, obj.stream); - } - } - - for (int j = 0 ; j < NUM_THREADS ; j++) { - t[j].join(); - } - - HIP_CHECK(hipStreamSynchronize(obj.stream)); - obj.threadCompleteStatus(); - } - return obj.resultAfterAllIterations(); -} - -bool testhipMemsetD32AsyncWithMultiThread() { - MemSetAsyncMthreadTest obj; - constexpr int memsetD32val = 0xDEADBEEF; - obj.memAllocate(memsetD32val); - std::thread t[NUM_THREADS]; - - for (int i = 0 ; i < ITER ; i++) { - for (int k = 0 ; k < NUM_THREADS ; k++) { - if (k%2) { - t[k] = std::thread(queueJobsForhipMemsetD32Async, obj.A_d, - obj.A_h, obj.memSetVal, obj.Nbytes, obj.stream); - } else { - t[k] = std::thread(queueJobsForhipMemsetD32Async, obj.A_d, - obj.B_h, obj.memSetVal, obj.Nbytes, obj.stream); - } - } - - for (int j = 0 ; j < NUM_THREADS ; j++) { - t[j].join(); - } - - HIP_CHECK(hipStreamSynchronize(obj.stream)); - obj.threadCompleteStatus(); - } - return obj.resultAfterAllIterations(); -} - -bool testhipMemsetD16AsyncWithMultiThread() { - MemSetAsyncMthreadTest obj; - constexpr int16_t memsetD16val = 0xDEAD; - obj.memAllocate(memsetD16val); - std::thread t[NUM_THREADS]; - - for (int i = 0 ; i < ITER ; i++) { - for (int k = 0 ; k < NUM_THREADS ; k++) { - if (k%2) { - t[k] = std::thread(queueJobsForhipMemsetD16Async, obj.A_d, - obj.A_h, obj.memSetVal, obj.Nbytes, obj.stream); - } else { - t[k] = std::thread(queueJobsForhipMemsetD16Async, obj.A_d, - obj.B_h, obj.memSetVal, obj.Nbytes, obj.stream); - } - } - - for (int j = 0 ; j < NUM_THREADS ; j++) { - t[j].join(); - } - - HIP_CHECK(hipStreamSynchronize(obj.stream)); - obj.threadCompleteStatus(); - } - return obj.resultAfterAllIterations(); -} - -bool testhipMemsetD8AsyncWithMultiThread() { - MemSetAsyncMthreadTest obj; - constexpr char memsetD8val = 0xDE; - obj.memAllocate(memsetD8val); - std::thread t[NUM_THREADS]; - - for (int i = 0 ; i < ITER ; i++) { - for (int k = 0 ; k < NUM_THREADS ; k++) { - if (k%2) { - t[k] = std::thread(queueJobsForhipMemsetD8Async, obj.A_d, - obj.A_h, obj.memSetVal, obj.Nbytes, obj.stream); - } else { - t[k] = std::thread(queueJobsForhipMemsetD8Async, obj.A_d, - obj.B_h, obj.memSetVal, obj.Nbytes, obj.stream); - } - } - for (int j = 0 ; j < NUM_THREADS ; j++) { - t[j].join(); - } - - HIP_CHECK(hipStreamSynchronize(obj.stream)); - obj.threadCompleteStatus(); - } - return obj.resultAfterAllIterations(); -} - - -/* - * Test that validates functionality of hipmemsetAsync apis over multi threads - */ -TEST_CASE("Unit_hipMemsetAsync_QueueJobsMultithreaded") { - bool ret; - - SECTION("hipMemsetAsync With MultiThread") { - ret = testhipMemsetAsyncWithMultiThread(); - REQUIRE(ret == true); - } - - SECTION("hipMemsetD32Async With MultiThread") { - ret = testhipMemsetD32AsyncWithMultiThread(); - REQUIRE(ret == true); - } - - SECTION("hipMemsetD16Async With MultiThread") { - ret = testhipMemsetD16AsyncWithMultiThread(); - REQUIRE(ret == true); - } - - SECTION("hipMemsetD8Async With MultiThread") { - ret = testhipMemsetD8AsyncWithMultiThread(); - REQUIRE(ret == true); - } -} diff --git a/catch/unit/multiThread/CMakeLists.txt b/catch/unit/multiThread/CMakeLists.txt index 32abf0f5f3..8c77831a71 100644 --- a/catch/unit/multiThread/CMakeLists.txt +++ b/catch/unit/multiThread/CMakeLists.txt @@ -1,5 +1,6 @@ # Common Tests - Test independent of all platforms set(TEST_SRC + hipMemsetAsyncMultiThread.cc hipMultiThreadDevice.cc hipMultiThreadStreams1.cc hipMultiThreadStreams2.cc diff --git a/catch/unit/multiThread/hipMemsetAsyncMultiThread.cc b/catch/unit/multiThread/hipMemsetAsyncMultiThread.cc new file mode 100644 index 0000000000..0826ad7535 --- /dev/null +++ b/catch/unit/multiThread/hipMemsetAsyncMultiThread.cc @@ -0,0 +1,182 @@ +/* + * Copyright (c) 2023 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 WARRANNTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + */ + +#include +#include + +#include +#include +#include +#include +#include +#include +#include + +enum TestType { SameStream = 0, DifferentStreams }; + +// Func -> Memset/MemsetD[8/16/32], T - type of data to be worked on +// basically each thread memsets value present in input one by one on its +// allocated buffer +template +void threadCall(Func f, hipStream_t stream) { + // Should match hipMemsetAsync or hipMemsetD*Async arguments + static_assert( + (std::is_same:: + value || // hipMemsetAsync + std::is_same::value || // hipMemsetD32Async + std::is_same::value || // hipMemsetD16Async + std::is_same::value) && // hipMemsetD8Async + "Func f should be hipMemsetAsync or hipMemsetD*Async"); + + constexpr bool cast_2_void = + std::is_same::value; + + // Use the unsiged type, since memset concerns with set bit values over a mem + // address + typedef typename std::make_unsigned::type unsigned_t; + + unsigned_t min = 0; + unsigned_t max = std::numeric_limits::max(); + + std::mt19937_64 engine(std::random_device{}()); + auto distribution = std::uniform_int_distribution( + min, max); // this needs to be unsigned because windows does not treats + // char as numeric types + + T *ptr{nullptr}; + constexpr size_t size = 1024; + constexpr size_t iter = 1024; + HIP_CHECK_THREAD(hipMalloc(&ptr, sizeof(T) * size)); + hipEvent_t event{}; + HIP_CHECK_THREAD(hipEventCreate(&event)); + + union overlay_val_t { + T t_val; + unsigned_t u_val; + } overlay_val; + + std::vector dst(size, 0); + for (size_t i = 0; i < iter; i++) { + overlay_val.u_val = static_cast( + distribution(engine)); // generate an unsigned int number + if constexpr (cast_2_void) { + HIP_CHECK_THREAD(f((void *)ptr, overlay_val.t_val, size, stream)); + } else { + HIP_CHECK_THREAD( + f(*(hipDeviceptr_t *)&ptr, overlay_val.t_val, size, stream)); + } + HIP_CHECK_THREAD(hipMemcpyAsync(dst.data(), ptr, size * sizeof(T), + hipMemcpyDeviceToHost, stream)); + HIP_CHECK_THREAD(hipEventRecord(event, stream)); + HIP_CHECK_THREAD( + hipStreamWaitEvent(stream, event, 0)); // wait till memcpy is done + REQUIRE_THREAD(std::all_of(dst.begin(), dst.end(), [&](T v) { + // If this test ever fails, add prints here on mismatch + return v == overlay_val.t_val; + })); + } + + HIP_CHECK_THREAD(hipEventDestroy(event)); + HIP_CHECK_THREAD(hipFree(ptr)); +} + +// Func -> Memset/MemsetD[8/16/32], T - type of data to be worked on +template void launchThreads(Func f, TestType type) { + static_assert(!std::is_pointer::value && "Argument cant be a pointer"); + + // Should match hipMemsetAsync or hipMemsetD*Async arguments + static_assert( + (std::is_same:: + value || // hipMemsetAsync + std::is_same::value || // hipMemsetD32Async + std::is_same::value || // hipMemsetD16Async + std::is_same::value) && // hipMemsetD8Async + "Func f should be hipMemsetAsync or hipMemsetD*Async"); + + const size_t num_threads = + (std::thread::hardware_concurrency() > 8) + ? (((std::thread::hardware_concurrency() / 4) >= 127) + ? 127 + : (std::thread::hardware_concurrency() / 4)) + : 2; // thread count between 2 - 127 + + const size_t num_streams = (type == SameStream) ? 1 : num_threads; + std::vector streams(num_streams, nullptr); + + for (size_t i = 0; i < num_streams; i++) { + HIP_CHECK(hipStreamCreate(&streams[i])); + REQUIRE(streams[i] != nullptr); + } + + std::vector thread_pool; + thread_pool.reserve(num_threads); + auto thread_func = threadCall; + for (size_t i = 0; i < num_threads; i++) { + auto stream = (type == SameStream) ? streams[0] : streams[i]; + thread_pool.emplace_back(std::thread(thread_func, f, stream)); + } + + for (size_t i = 0; i < num_threads; i++) { + thread_pool[i].join(); + } + + HIP_CHECK_THREAD_FINALIZE(); // Make sure all thread have exited properly + + for (size_t i = 0; i < num_streams; i++) { + HIP_CHECK(hipStreamDestroy(streams[i])); + } +} + +TEST_CASE("Unit_hipMemsetAsync_QueueJobsMultithreaded") { + using hipMemsetAsync_t = + hipError_t (*)(void *, int, const size_t, hipStream_t); + using hipMemsetAsyncD8_t = + hipError_t (*)(hipDeviceptr_t, unsigned char, const size_t, hipStream_t); + using hipMemsetAsyncD16_t = + hipError_t (*)(hipDeviceptr_t, unsigned short, const size_t, hipStream_t); + using hipMemsetAsyncD32_t = + hipError_t (*)(hipDeviceptr_t, int, const size_t, hipStream_t); + + launchThreads(hipMemsetAsync, SameStream); + launchThreads(hipMemsetAsync, DifferentStreams); + + launchThreads(hipMemsetD8Async, + SameStream); + launchThreads(hipMemsetD8Async, + DifferentStreams); + launchThreads(hipMemsetD16Async, + SameStream); + launchThreads(hipMemsetD16Async, + DifferentStreams); + launchThreads(hipMemsetD32Async, SameStream); + launchThreads(hipMemsetD32Async, DifferentStreams); +}