diff --git a/projects/hip-tests/catch/include/memcpy1d_tests_common.hh b/projects/hip-tests/catch/include/memcpy1d_tests_common.hh index 7eae09ab8b..e63394f148 100644 --- a/projects/hip-tests/catch/include/memcpy1d_tests_common.hh +++ b/projects/hip-tests/catch/include/memcpy1d_tests_common.hh @@ -247,12 +247,19 @@ void MemcpySyncBehaviorCheck(F memcpy_func, const bool should_sync, } template -void MemcpyHtoDSyncBehavior(F memcpy_func, const bool should_sync, - const hipStream_t kernel_stream = nullptr) { - using LA = LinearAllocs; - const auto host_alloc_type = GENERATE(LA::malloc, LA::hipHostMalloc); - LinearAllocGuard host_alloc(host_alloc_type, kPageSize); - LinearAllocGuard device_alloc(LA::hipMalloc, kPageSize); +void MemcpyHPageabletoDSyncBehavior(F memcpy_func, const bool should_sync, + const hipStream_t kernel_stream = nullptr) { + LinearAllocGuard host_alloc(LinearAllocs::malloc, kPageSize); + LinearAllocGuard device_alloc(LinearAllocs::hipMalloc, kPageSize); + MemcpySyncBehaviorCheck(std::bind(memcpy_func, device_alloc.ptr(), host_alloc.ptr(), kPageSize), + should_sync, kernel_stream); +} + +template +void MemcpyHPinnedtoDSyncBehavior(F memcpy_func, const bool should_sync, + const hipStream_t kernel_stream = nullptr) { + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, kPageSize); + LinearAllocGuard device_alloc(LinearAllocs::hipMalloc, kPageSize); MemcpySyncBehaviorCheck(std::bind(memcpy_func, device_alloc.ptr(), host_alloc.ptr(), kPageSize), should_sync, kernel_stream); } @@ -288,8 +295,9 @@ template void MemcpyHtoHSyncBehavior(F memcpy_func, const bool should_sync, const hipStream_t kernel_stream = nullptr) { using LA = LinearAllocs; - const auto src_alloc_type = GENERATE(LA::malloc, LA::hipHostMalloc); - const auto dst_alloc_type = GENERATE(LA::malloc, LA::hipHostMalloc); + const auto [src_alloc_type, dst_alloc_type] = GENERATE( + std::make_tuple(LA::malloc, LA::hipHostMalloc), + std::make_tuple(LA::hipHostMalloc, LA::malloc), std::make_tuple(LA::malloc, LA::malloc)); LinearAllocGuard src_alloc(src_alloc_type, kPageSize); LinearAllocGuard dst_alloc(dst_alloc_type, kPageSize); @@ -297,6 +305,15 @@ void MemcpyHtoHSyncBehavior(F memcpy_func, const bool should_sync, should_sync, kernel_stream); } +template +void MemcpyHPinnedtoHPinnedSyncBehavior(F memcpy_func, const bool should_sync, + const hipStream_t kernel_stream = nullptr) { + LinearAllocGuard src_alloc(LinearAllocs::hipHostMalloc, kPageSize); + LinearAllocGuard dst_alloc(LinearAllocs::hipHostMalloc, kPageSize); + MemcpySyncBehaviorCheck(std::bind(memcpy_func, dst_alloc.ptr(), src_alloc.ptr(), kPageSize), + should_sync, kernel_stream); +} + // Common negative tests template void MemcpyCommonNegativeTests(F f, void* dst, void* src, size_t count) { SECTION("dst == nullptr") { HIP_CHECK_ERROR(f(nullptr, src, count), hipErrorInvalidValue); } @@ -316,4 +333,4 @@ void MemcpyWithDirectionCommonNegativeTests(F f, void* dst, void* src, size_t co hipErrorInvalidMemcpyDirection); } #endif -} \ No newline at end of file +} diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpyAsync.cc b/projects/hip-tests/catch/unit/memory/hipMemcpyAsync.cc index a57449e84f..0b9ec7fa81 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpyAsync.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemcpyAsync.cc @@ -38,9 +38,14 @@ TEST_CASE("Unit_hipMemcpyAsync_Positive_Synchronization_Behavior") { using namespace std::placeholders; HIP_CHECK(hipDeviceSynchronize()); - SECTION("Host memory to device memory") { - MemcpyHtoDSyncBehavior(std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyHostToDevice, nullptr), - true); + SECTION("Host pageable memory to device memory") { + MemcpyHPageabletoDSyncBehavior( + std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyHostToDevice, nullptr), true); + } + + SECTION("Host pinned memory to device memory") { + MemcpyHPinnedtoDSyncBehavior( + std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyHostToDevice, nullptr), false); } SECTION("Device memory to pageable host memory") { @@ -66,6 +71,9 @@ TEST_CASE("Unit_hipMemcpyAsync_Positive_Synchronization_Behavior") { SECTION("Host memory to host memory") { MemcpyHtoHSyncBehavior(std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyHostToHost, nullptr), true); + + MemcpyHPinnedtoHPinnedSyncBehavior( + std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyHostToHost, nullptr), false); } } diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpyAsync_derivatives.cc b/projects/hip-tests/catch/unit/memory/hipMemcpyAsync_derivatives.cc index 7fed6121fe..b2583d1ccc 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpyAsync_derivatives.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemcpyAsync_derivatives.cc @@ -98,7 +98,7 @@ TEST_CASE("Unit_hipMemcpyHtoDAsync_Positive_Synchronization_Behavior") { "Nvidia"); return; #endif - MemcpyHtoDSyncBehavior( + MemcpyHPinnedtoDSyncBehavior( [](void* dst, void* src, size_t count) { return hipMemcpyHtoDAsync(reinterpret_cast(dst), src, count, nullptr); }, diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpyWithStream.cc b/projects/hip-tests/catch/unit/memory/hipMemcpyWithStream.cc index 24e11c25f3..d8f9b5bc0f 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpyWithStream.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemcpyWithStream.cc @@ -38,7 +38,7 @@ TEST_CASE("Unit_hipMemcpy_Positive_Synchronization_Behavior") { // For transfers from pinned host memory to device memory, the function is synchronous with // respect to the host SECTION("Host memory to device memory") { - MemcpyHtoDSyncBehavior(std::bind(hipMemcpy, _1, _2, _3, hipMemcpyHostToDevice), true); + MemcpyHPageabletoDSyncBehavior(std::bind(hipMemcpy, _1, _2, _3, hipMemcpyHostToDevice), true); } // For transfers from device to either pageable or pinned host memory, the function returns only @@ -65,6 +65,8 @@ TEST_CASE("Unit_hipMemcpy_Positive_Synchronization_Behavior") { // respect to the host SECTION("Host memory to host memory") { MemcpyHtoHSyncBehavior(std::bind(hipMemcpy, _1, _2, _3, hipMemcpyHostToHost), true); + MemcpyHPinnedtoHPinnedSyncBehavior( + std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyHostToHost, nullptr), true); } } diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpy_derivatives.cc b/projects/hip-tests/catch/unit/memory/hipMemcpy_derivatives.cc index 0375602508..e3ff93e8fb 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpy_derivatives.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemcpy_derivatives.cc @@ -60,7 +60,7 @@ TEST_CASE("Unit_hipMemcpyHtoD_Positive_Basic") { } TEST_CASE("Unit_hipMemcpyHtoD_Positive_Synchronization_Behavior") { - MemcpyHtoDSyncBehavior( + MemcpyHPageabletoDSyncBehavior( [](void* dst, void* src, size_t count) { return hipMemcpyHtoD(reinterpret_cast(dst), src, count); }, @@ -116,4 +116,4 @@ TEST_CASE("Unit_hipMemcpyDtoD_Negative_Parameters") { reinterpret_cast(src), count); }, dst_alloc.ptr(), src_alloc.ptr(), kPageSize); -} \ No newline at end of file +}