diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index 32ea04d3c1..0fbb7b7ce5 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -112,6 +112,7 @@ set(TEST_SRC hipMemsetAsync.cc hipMemAdvise.cc hipMemRangeGetAttributes.cc + hipStreamAttachMemAsync.cc hipMemRangeGetAttributes_old.cc hipMemGetAddressRange.cc ) @@ -203,6 +204,7 @@ set(TEST_SRC hipMemRangeGetAttributes.cc hipMemRangeGetAttributes_old.cc hipGetSymbolSizeAddress.cc + hipStreamAttachMemAsync.cc hipMemGetAddressRange.cc ) endif() diff --git a/catch/unit/memory/hipStreamAttachMemAsync.cc b/catch/unit/memory/hipStreamAttachMemAsync.cc new file mode 100644 index 0000000000..05a73e5cc7 --- /dev/null +++ b/catch/unit/memory/hipStreamAttachMemAsync.cc @@ -0,0 +1,224 @@ +/* +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 +#include +#include +#include +#include + +TEST_CASE("Unit_hipStreamAttachMemAsync_Positive_Basic") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) { + HipTest::HIP_SKIP_TEST("Managed memory is not supported"); + return; + } + + StreamGuard stream(Streams::created); + LinearAllocGuard managed(LinearAllocs::hipMallocManaged, kPageSize, + hipMemAttachHost); + + HIP_CHECK(hipStreamAttachMemAsync(stream.stream(), managed.ptr(), 0)); + HIP_CHECK(hipStreamSynchronize(stream.stream())); +} + +TEST_CASE("Unit_hipStreamAttachMemAsync_Positive_Pageable") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) { + HipTest::HIP_SKIP_TEST("Managed memory is not supported"); + return; + } + + if (!DeviceAttributesSupport(0, hipDeviceAttributePageableMemoryAccess)) { + HipTest::HIP_SKIP_TEST("Pageable memory access is not supported"); + return; + } + + StreamGuard stream(Streams::created); + LinearAllocGuard pageable(LinearAllocs::malloc, kPageSize); + + HIP_CHECK(hipStreamAttachMemAsync(stream.stream(), pageable.ptr(), kPageSize)); + HIP_CHECK(hipStreamSynchronize(stream.stream())); +} + +// CUDA docs: +// If the cudaMemAttachGlobal flag is specified, the memory can be accessed by any stream on any +// device. +TEST_CASE("Unit_hipStreamAttachMemAsync_Positive_AttachGlobal") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) { + HipTest::HIP_SKIP_TEST("Managed memory is not supported"); + return; + } + + const auto device_count = HipTest::getDeviceCount(); + const auto stream_count = device_count < 2 ? 8 : device_count; + + std::vector> streams; + streams.reserve(stream_count); + for (int i = 0; i < stream_count; ++i) { + if (device_count > 1) { + HIP_CHECK(hipSetDevice(i)); + } + streams.push_back(std::make_unique(Streams::created)); + } + + LinearAllocGuard managed_global(LinearAllocs::hipMallocManaged, sizeof(int) * stream_count, + hipMemAttachHost); + + HIP_CHECK(hipStreamAttachMemAsync( + nullptr, reinterpret_cast(managed_global.ptr()), 0, hipMemAttachGlobal)); + HIP_CHECK(hipStreamSynchronize(nullptr)); + + for (int i = 0; i < stream_count; ++i) { + HipTest::launchKernel(Set, 1, 1, 0, streams.at(i)->stream(), managed_global.ptr() + i, i); + } + + for (auto&& stream : streams) { + HIP_CHECK(hipStreamSynchronize(stream->stream())); + } + + for (int i = 0; i < stream_count; ++i) { + REQUIRE(managed_global.ptr()[i] == i); + } +} + +// CUDA docs: +// If the cudaMemAttachHost flag is specified, the program makes a guarantee that it won't access +// the memory on the device from any stream on a device that has a zero value for the device +// attribute cudaDevAttrConcurrentManagedAccess. +TEST_CASE("Unit_hipStreamAttachMemAsync_Positive_AttachHost") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) { + HipTest::HIP_SKIP_TEST("Managed memory is not supported"); + return; + } + + if (DeviceAttributesSupport(0, hipDeviceAttributeConcurrentManagedAccess)) { + HipTest::HIP_SKIP_TEST("Device supports concurrent managed access"); + return; + } + + StreamGuard stream(Streams::created); + LinearAllocGuard managed_global(LinearAllocs::hipMallocManaged, sizeof(int)); + LinearAllocGuard managed_host(LinearAllocs::hipMallocManaged, sizeof(int)); + + HIP_CHECK(hipStreamAttachMemAsync( + stream.stream(), reinterpret_cast(managed_host.ptr()), 0, hipMemAttachHost)); + HIP_CHECK(hipStreamSynchronize(stream.stream())); + + HipTest::launchKernel(Set, 1, 1, 0, stream.stream(), managed_global.ptr(), 32); + *managed_host.ptr() = 64; + HIP_CHECK(hipStreamSynchronize(stream.stream())); + + REQUIRE(*managed_global.ptr() == 32); + REQUIRE(*managed_host.ptr() == 64); +} + +// CUDA docs: +// If the cudaMemAttachSingle flag is specified and stream is associated with a device that has a +// zero value for the device attribute cudaDevAttrConcurrentManagedAccess, the program makes a +// guarantee that it will only access the memory on the device from stream. +TEST_CASE("Unit_hipStreamAttachMemAsync_Positive_AttachSingle") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) { + HipTest::HIP_SKIP_TEST("Managed memory is not supported"); + return; + } + + if (DeviceAttributesSupport(0, hipDeviceAttributeConcurrentManagedAccess)) { + HipTest::HIP_SKIP_TEST("Device supports concurrent managed access"); + return; + } + + StreamGuard stream1(Streams::created); + StreamGuard stream2(Streams::created); + + LinearAllocGuard managed_global(LinearAllocs::hipMallocManaged, sizeof(int)); + LinearAllocGuard managed_single(LinearAllocs::hipMallocManaged, sizeof(int), + hipMemAttachHost); + + HIP_CHECK(hipStreamAttachMemAsync(stream1.stream(), + reinterpret_cast(managed_single.ptr()), 0, + hipMemAttachSingle)); + HIP_CHECK(hipStreamSynchronize(stream1.stream())); + + HipTest::launchKernel(Set, 1, 1, 0, stream1.stream(), managed_single.ptr(), 64); + HIP_CHECK(hipStreamSynchronize(stream1.stream())); + + HipTest::launchKernel(Set, 1, 1, 0, stream2.stream(), managed_global.ptr(), 32); + + REQUIRE(*managed_single.ptr() == 64); + *managed_single.ptr() = 128; + + HIP_CHECK(hipStreamSynchronize(stream2.stream())); + + REQUIRE(*managed_global.ptr() == 32); + REQUIRE(*managed_single.ptr() == 128); +} + +TEST_CASE("Unit_hipStreamAttachMemAsync_Negative_Parameters") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) { + HipTest::HIP_SKIP_TEST("Managed memory is not supported"); + return; + } + + StreamGuard stream(Streams::created); + LinearAllocGuard managed(LinearAllocs::hipMallocManaged, kPageSize, + hipMemAttachHost); + + SECTION("invalid stream") { + HIP_CHECK(hipStreamDestroy(stream.stream())); + HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream.stream(), managed.ptr()), + hipErrorContextIsDestroyed); + } + + SECTION("dev_ptr == nullptr") { + HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream.stream(), nullptr), hipErrorInvalidValue); + } + + SECTION("length is not zero nor entire allocation size") { + HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream.stream(), managed.ptr(), kPageSize / 2), + hipErrorInvalidValue); + } + + SECTION("invalid flags") { + HIP_CHECK_ERROR( + hipStreamAttachMemAsync(stream.stream(), managed.ptr(), 0, + hipMemAttachGlobal | hipMemAttachHost | hipMemAttachSingle), + hipErrorInvalidValue); + } + + SECTION("attach single to nullstream") { + HIP_CHECK_ERROR(hipStreamAttachMemAsync(nullptr, managed.ptr(), 0, hipMemAttachSingle), + hipErrorInvalidValue); + } + + LinearAllocGuard pageable(LinearAllocs::malloc, kPageSize); + + if (!DeviceAttributesSupport(0, hipDeviceAttributePageableMemoryAccess)) { + SECTION("dev_ptr is pageable memory") { + HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream.stream(), pageable.ptr(), kPageSize), + hipErrorInvalidValue); + } + } else { + SECTION("length is zero for pageable memory") { + HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream.stream(), pageable.ptr(), 0), + hipErrorInvalidValue); + } + } +} \ No newline at end of file diff --git a/catch/unit/stream/CMakeLists.txt b/catch/unit/stream/CMakeLists.txt index 548bb2e52c..38e651e5b7 100644 --- a/catch/unit/stream/CMakeLists.txt +++ b/catch/unit/stream/CMakeLists.txt @@ -32,8 +32,8 @@ set(TEST_SRC hipStreamCreateWithPriority.cc hipStreamDestroy.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. + # hipStreamAttachMemAsync_old.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 hipStreamSynchronize.cc @@ -42,7 +42,7 @@ set(TEST_SRC hipStreamACb_StrmSyncTiming.cc ) - # set_source_files_properties(hipStreamAttachMemAsync.cc PROPERTIES COMPILE_FLAGS -std=c++17) + # set_source_files_properties(hipStreamAttachMemAsync_old.cc PROPERTIES COMPILE_FLAGS -std=c++17) endif() hip_add_exe_to_target(NAME StreamTest diff --git a/catch/unit/stream/hipStreamAttachMemAsync_old.cc b/catch/unit/stream/hipStreamAttachMemAsync_old.cc new file mode 100644 index 0000000000..e02754dfbb --- /dev/null +++ b/catch/unit/stream/hipStreamAttachMemAsync_old.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