From 987262809eb56b0d086fba233da00407918e595e Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Mon, 11 Jul 2022 10:07:22 +0100 Subject: [PATCH] EXSWCPHIPT-37 - Add test for hipStreamAttachMemAsync (#2558) [ROCm/hip commit: 2b1794e5d24e24076a8d88abd4d1cfa4b32e850f] --- .../tests/catch/unit/stream/CMakeLists.txt | 4 + .../unit/stream/hipStreamAttachMemAsync.cc | 143 ++++++++++++++++++ 2 files changed, 147 insertions(+) create mode 100644 projects/hip/tests/catch/unit/stream/hipStreamAttachMemAsync.cc diff --git a/projects/hip/tests/catch/unit/stream/CMakeLists.txt b/projects/hip/tests/catch/unit/stream/CMakeLists.txt index 932e76f718..229df157c1 100644 --- a/projects/hip/tests/catch/unit/stream/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/stream/CMakeLists.txt @@ -25,9 +25,13 @@ set(TEST_SRC hipStreamCreateWithFlags.cc hipStreamCreateWithPriority.cc hipAPIStreamDisable.cc + # hipStreamAttachMemAsync.cc # Disabling it on nvidia due to issue in function definition of hipStreamAttachMemAsync + # Fixing would break ABI, to be re-enabled when the fix is made. streamCommon.cc hipStreamValue.cc ) + + set_source_files_properties(hipStreamAttachMemAsync.cc PROPERTIES COMPILE_FLAGS -std=c++17) endif() hip_add_exe_to_target(NAME StreamTest diff --git a/projects/hip/tests/catch/unit/stream/hipStreamAttachMemAsync.cc b/projects/hip/tests/catch/unit/stream/hipStreamAttachMemAsync.cc new file mode 100644 index 0000000000..e02754dfbb --- /dev/null +++ b/projects/hip/tests/catch/unit/stream/hipStreamAttachMemAsync.cc @@ -0,0 +1,143 @@ +/* +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. +*/ + +// TODO Enable it after hipStreamAttachMemAsync is feature complete on HIP + +#include +#include + +__device__ __managed__ int var = 0; + +enum class StreamAttachTestType { NullStream = 0, StreamPerThread, CreatedStream }; + +TEST_CASE("Unit_hipStreamAttachMemAsync_Negative") { + hipStream_t stream{nullptr}; + + auto streamType = + GENERATE(StreamAttachTestType::NullStream, StreamAttachTestType::StreamPerThread, + StreamAttachTestType::CreatedStream); + + if (streamType == StreamAttachTestType::StreamPerThread) { + stream = hipStreamPerThread; + } else if (streamType == StreamAttachTestType::CreatedStream) { + HIP_CHECK(hipStreamCreate(&stream)); + REQUIRE(stream != nullptr); + } + + SECTION("Invalid Resource Handle") { + int definitelyNotAManagedVariable = 0; + HIP_CHECK_ERROR( + hipStreamAttachMemAsync(stream, reinterpret_cast(&definitelyNotAManagedVariable), + sizeof(int), hipMemAttachSingle), + hipErrorInvalidValue); + } + + SECTION("Invalid devptr") { + HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream, nullptr, sizeof(int), hipMemAttachSingle), + hipErrorInvalidValue); + } + + SECTION("Invalid Resource Size") { + HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream, reinterpret_cast(&var), sizeof(int) - 1, + hipMemAttachSingle), + hipErrorInvalidValue); + } + + SECTION("Invalid Flags") { + HIP_CHECK_ERROR( + hipStreamAttachMemAsync(stream, reinterpret_cast(&var), sizeof(int) - 1, + hipMemAttachSingle | hipMemAttachHost | hipMemAttachGlobal), + hipErrorInvalidValue); + } + + if (streamType == StreamAttachTestType::CreatedStream) { + HIP_CHECK(hipStreamDestroy(stream)); + } +} + +__global__ void kernel(int* ptr, size_t size) { + auto i = threadIdx.x; + if (i < size) { + ptr[i] = 1024; + } +} + +constexpr size_t size = 1024; +__device__ __managed__ int m_memory[size]; + +TEST_CASE("Unit_hipStreamAttachMemAsync_UseCase") { + hipStream_t stream{nullptr}; + + auto streamType = + GENERATE(StreamAttachTestType::NullStream, StreamAttachTestType::StreamPerThread, + StreamAttachTestType::CreatedStream); + + if (streamType == StreamAttachTestType::CreatedStream) { + HIP_CHECK(hipStreamCreate(&stream)); + REQUIRE(stream != nullptr); + } + + SECTION("Size zero is valid") { + int* d_memory{nullptr}; + HIP_CHECK(hipMallocManaged(&d_memory, sizeof(int) * size, hipMemAttachHost)); + HIP_CHECK( + hipStreamAttachMemAsync(stream, reinterpret_cast(d_memory), 0, hipMemAttachHost)); + HIP_CHECK(hipStreamSynchronize(stream)); // Wait for command to complete + HIP_CHECK(hipFree(d_memory)); + } + + SECTION("Access from device and host") { + int* d_memory{nullptr}; + + HIP_CHECK(hipMallocManaged(&d_memory, sizeof(int) * size, hipMemAttachHost)); + HIP_CHECK(hipMemset(d_memory, 0, sizeof(int) * size)); + HIP_CHECK( + hipStreamAttachMemAsync(stream, reinterpret_cast(d_memory), 0, hipMemAttachHost)); + HIP_CHECK(hipStreamSynchronize(stream)); // Wait for the command to complete + + kernel<<<1, size, 0, stream>>>(d_memory, size); + HIP_CHECK(hipStreamSynchronize(stream)); // Wait for the kernel to complete + + auto ptr = std::make_unique(size); + std::copy(d_memory, d_memory + size, ptr.get()); + + HIP_CHECK(hipFree(d_memory)); + + REQUIRE(std::all_of(ptr.get(), ptr.get() + size, [](int n) { return n == size; })); + } + + SECTION("Access ManagedMemory") { + HIP_CHECK(hipMemset(m_memory, 0, sizeof(int) * size)); + HIP_CHECK( + hipStreamAttachMemAsync(stream, reinterpret_cast(m_memory), 0, hipMemAttachHost)); + HIP_CHECK(hipStreamSynchronize(stream)); // Wait for the command to complete + + kernel<<<1, size, 0, stream>>>(m_memory, size); + HIP_CHECK(hipStreamSynchronize(stream)); // Wait for the kernel to complete + + auto ptr = std::make_unique(size); + std::copy(m_memory, m_memory + size, ptr.get()); + + REQUIRE(std::all_of(ptr.get(), ptr.get() + size, [](int n) { return n == size; })); + } + + if (streamType == StreamAttachTestType::CreatedStream) { + HIP_CHECK(hipStreamDestroy(stream)); + } +} \ No newline at end of file