From a5db63739c4725d019ca2e23165382b3aebba50a Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Fri, 10 Jun 2022 16:47:34 +0100 Subject: [PATCH] Adding/Moving tests for hipMemcpy[From/To]Symbol[Async] (#2638) [ROCm/hip commit: 54b1822167e7058ca194a3e70faf13c0dcd96665] --- .../tests/catch/unit/memory/CMakeLists.txt | 6 - .../catch/unit/memory/hipMemcpyFromSymbol.cc | 179 ++++++++++++++++-- .../unit/memory/hipMemcpyFromSymbolAsync.cc | 47 ----- .../catch/unit/memory/hipMemcpyToSymbol.cc | 43 ----- .../unit/memory/hipMemcpyToSymbolAsync.cc | 47 ----- 5 files changed, 163 insertions(+), 159 deletions(-) delete mode 100644 projects/hip/tests/catch/unit/memory/hipMemcpyFromSymbolAsync.cc delete mode 100644 projects/hip/tests/catch/unit/memory/hipMemcpyToSymbol.cc delete mode 100644 projects/hip/tests/catch/unit/memory/hipMemcpyToSymbolAsync.cc diff --git a/projects/hip/tests/catch/unit/memory/CMakeLists.txt b/projects/hip/tests/catch/unit/memory/CMakeLists.txt index d36f0a1b85..70298c246e 100644 --- a/projects/hip/tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/memory/CMakeLists.txt @@ -58,9 +58,6 @@ set(TEST_SRC hipMallocManaged.cc hipMemRangeGetAttribute.cc hipMemcpyFromSymbol.cc - hipMemcpyFromSymbolAsync.cc - hipMemcpyToSymbol.cc - hipMemcpyToSymbolAsync.cc hipPtrGetAttribute.cc hipMemcpyPeer.cc hipMemcpyPeerAsync.cc @@ -123,9 +120,6 @@ set(TEST_SRC hipMallocManaged.cc hipMemRangeGetAttribute.cc hipMemcpyFromSymbol.cc - hipMemcpyFromSymbolAsync.cc - hipMemcpyToSymbol.cc - hipMemcpyToSymbolAsync.cc hipPtrGetAttribute.cc hipMemcpyPeer.cc hipMemcpyPeerAsync.cc diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpyFromSymbol.cc b/projects/hip/tests/catch/unit/memory/hipMemcpyFromSymbol.cc index 347b4589a1..ff38ddd060 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpyFromSymbol.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpyFromSymbol.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 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 @@ -18,26 +18,173 @@ THE SOFTWARE. */ #include -#define SIZE 1024 -/* Test verifies hipMemcpyFromSymbol API Negative scenarios. +__device__ int devSymbol[10]; + +/* Test verifies hipMemcpy[From/To]Symbol[Async] API Negative scenarios. */ -TEST_CASE("Unit_hipMemcpyFromSymbol_Negative") { - void *Sd; - char S[SIZE]="This is not a device symbol"; - - HIP_CHECK(hipMalloc(&Sd, SIZE)); - - SECTION("Passing void pointer") { - REQUIRE(hipSuccess != hipMemcpyFromSymbol(S, HIP_SYMBOL(Sd), - SIZE, 0, hipMemcpyDeviceToHost)); +TEST_CASE("Unit_hipMemcpyFromToSymbol_Negative") { + SECTION("Invalid Src Ptr") { + int result{0}; + HIP_CHECK_ERROR( + hipMemcpyFromSymbol(nullptr, HIP_SYMBOL(devSymbol), sizeof(int), 0, hipMemcpyDeviceToHost), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpyToSymbol(nullptr, &result, sizeof(int), 0, hipMemcpyHostToDevice), + hipErrorInvalidSymbol); + HIP_CHECK_ERROR( + hipMemcpyToSymbolAsync(nullptr, &result, sizeof(int), 0, hipMemcpyHostToDevice, nullptr), + hipErrorInvalidSymbol); + HIP_CHECK_ERROR(hipMemcpyFromSymbolAsync(nullptr, HIP_SYMBOL(devSymbol), sizeof(int), 0, + hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidValue); } - SECTION("Passing NULL Pointer") { - REQUIRE(hipSuccess != hipMemcpyFromSymbol(S, nullptr, - SIZE, 0, hipMemcpyDeviceToHost)); + SECTION("Invalid Dst Ptr") { + int result{0}; + HIP_CHECK_ERROR(hipMemcpyFromSymbol(&result, nullptr, sizeof(int), 0, hipMemcpyDeviceToHost), + hipErrorInvalidSymbol); + HIP_CHECK_ERROR( + hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), nullptr, sizeof(int), 0, hipMemcpyHostToDevice), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), nullptr, sizeof(int), 0, + hipMemcpyHostToDevice, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpyFromSymbolAsync(&result, nullptr, sizeof(int), 0, hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidSymbol); } - HIP_CHECK(hipFree(Sd)); + SECTION("Invalid Size") { + int result{0}; + HIP_CHECK_ERROR(hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int) * 100, 0, + hipMemcpyDeviceToHost), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), &result, sizeof(int) * 100, 0, + hipMemcpyHostToDevice), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), &result, sizeof(int) * 100, 0, + hipMemcpyHostToDevice, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpyFromSymbolAsync(&result, HIP_SYMBOL(devSymbol), sizeof(int) * 100, 0, + hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidValue); + } + + SECTION("Invalid Offset") { + int result{0}; + HIP_CHECK_ERROR(hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int), 300, + hipMemcpyDeviceToHost), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), &result, sizeof(int), 300, hipMemcpyHostToDevice), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), &result, sizeof(int), 300, + hipMemcpyHostToDevice, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpyFromSymbolAsync(&result, HIP_SYMBOL(devSymbol), sizeof(int), 300, + hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidValue); + } + + SECTION("Invalid Direction") { + int result{0}; + HIP_CHECK_ERROR( + hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int), 0, hipMemcpyHostToDevice), + hipErrorInvalidMemcpyDirection); + HIP_CHECK_ERROR( + hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), &result, sizeof(int), 0, hipMemcpyDeviceToHost), + hipErrorInvalidMemcpyDirection); + HIP_CHECK_ERROR(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), &result, sizeof(int), 0, + hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidMemcpyDirection); + HIP_CHECK_ERROR(hipMemcpyFromSymbolAsync(&result, HIP_SYMBOL(devSymbol), sizeof(int), 0, + hipMemcpyHostToDevice, nullptr), + hipErrorInvalidMemcpyDirection); + } +} + +/* + * Test Verifies hipMemcpyToSymbol/hipMemcpyFromSymbol and Async Variants for simple use case + * For single valuea To and From Symbol + * For Array Values To and From Symbol + * For Array Values with offset To and From Symbol + * For Sync and Async Variants*/ +TEST_CASE("Unit_hipMemcpyToFromSymbol_SyncAndAsync") { + enum StreamTestType { NullStream = 0, StreamPerThread, CreatedStream, NoStream }; + + /* Test type NoStream - Use Sync variants, else use async variants */ + auto streamType = GENERATE(StreamTestType::NoStream, StreamTestType::NullStream, + StreamTestType::StreamPerThread, StreamTestType::CreatedStream); + + hipStream_t stream{nullptr}; + + if (streamType == StreamTestType::StreamPerThread) { + stream = hipStreamPerThread; + } else if (streamType == StreamTestType::CreatedStream) { + HIP_CHECK(hipStreamCreate(&stream)); + } + INFO("Stream :: " << streamType); + + SECTION("Singular Value") { + int set{42}; + int result{0}; + if (streamType == StreamTestType::NoStream) { + HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), &set, sizeof(int))); + HIP_CHECK(hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int))); + } else { + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), &set, sizeof(int), 0, + hipMemcpyHostToDevice, stream)); + + HIP_CHECK(hipMemcpyFromSymbolAsync(&result, HIP_SYMBOL(devSymbol), sizeof(int), 0, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + REQUIRE(result == set); + } + + SECTION("Array Values") { + constexpr size_t size{10}; + int set[size] = {4, 2, 4, 2, 4, 2, 4, 2, 4, 2}; + int result[size] = {0}; + if (streamType == StreamTestType::NoStream) { + HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), set, sizeof(int) * size)); + HIP_CHECK(hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int) * size)); + } else { + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), set, sizeof(int) * size, 0, + hipMemcpyHostToDevice, stream)); + + HIP_CHECK(hipMemcpyFromSymbolAsync(&result, HIP_SYMBOL(devSymbol), sizeof(int) * size, 0, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + for (size_t i = 0; i < size; i++) { + REQUIRE(result[i] == set[i]); + } + } + + SECTION("Offset'ed Values") { + constexpr size_t size{10}; + constexpr size_t offset = 5 * sizeof(int); + int set[size] = {9, 9, 9, 9, 9, 2, 4, 2, 4, 2}; + int result[size] = {0}; + if (streamType == StreamTestType::NoStream) { + HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), set, offset)); + HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), set + 5, offset, offset)); + HIP_CHECK(hipMemcpyFromSymbol(result, HIP_SYMBOL(devSymbol), sizeof(int) * size)); + } else { + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), set, offset, 0, hipMemcpyHostToDevice, + stream)); + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), set + 5, offset, offset, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyFromSymbolAsync(result, HIP_SYMBOL(devSymbol), offset, 0, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipMemcpyFromSymbolAsync(result + 5, HIP_SYMBOL(devSymbol), offset, offset, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + for (size_t i = 0; i < size; i++) { + REQUIRE(result[i] == set[i]); + } + } } diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpyFromSymbolAsync.cc b/projects/hip/tests/catch/unit/memory/hipMemcpyFromSymbolAsync.cc deleted file mode 100644 index ae6ebc3e82..0000000000 --- a/projects/hip/tests/catch/unit/memory/hipMemcpyFromSymbolAsync.cc +++ /dev/null @@ -1,47 +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 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 -#define SIZE 1024 - -/* Test verifies hipMemcpyFromSymbolAsync API Negative scenarios. - */ - -TEST_CASE("Unit_hipMemcpyFromSymbolAsync_Negative") { - void *Sd; - char S[SIZE]="This is not a device symbol"; - - HIP_CHECK(hipMalloc(&Sd, SIZE)); - - hipStream_t stream; - HIP_CHECK(hipStreamCreate(&stream)); - - SECTION("Passing void pointer") { - REQUIRE(hipSuccess != hipMemcpyFromSymbolAsync(S, HIP_SYMBOL(Sd), - SIZE, 0, hipMemcpyDeviceToHost, stream)); - } - - SECTION("Passing NULL pointer") { - REQUIRE(hipSuccess != hipMemcpyFromSymbolAsync(S, nullptr, - SIZE, 0, hipMemcpyDeviceToHost, stream)); - } - - HIP_CHECK(hipStreamDestroy(stream)); - HIP_CHECK(hipFree(Sd)); -} diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpyToSymbol.cc b/projects/hip/tests/catch/unit/memory/hipMemcpyToSymbol.cc deleted file mode 100644 index 23e5768c3c..0000000000 --- a/projects/hip/tests/catch/unit/memory/hipMemcpyToSymbol.cc +++ /dev/null @@ -1,43 +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 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 -#define SIZE 1024 - -/* Test verifies hipMemcpyToSymbol API Negative scenarios. - */ - -TEST_CASE("Unit_hipMemcpyToSymbol_Negative") { - void *Sd; - char S[SIZE]="This is not a device symbol"; - - HIP_CHECK(hipMalloc(&Sd, SIZE)); - - SECTION("Passing void pointer") { - REQUIRE(hipSuccess != hipMemcpyToSymbol(HIP_SYMBOL(Sd), S, - SIZE, 0, hipMemcpyDeviceToHost)); - } - - SECTION("Passing NULL Pointer") { - REQUIRE(hipSuccess != hipMemcpyToSymbol(nullptr, S, - SIZE, 0, hipMemcpyDeviceToHost)); - } - - HIP_CHECK(hipFree(Sd)); -} diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpyToSymbolAsync.cc b/projects/hip/tests/catch/unit/memory/hipMemcpyToSymbolAsync.cc deleted file mode 100644 index 42b82128f9..0000000000 --- a/projects/hip/tests/catch/unit/memory/hipMemcpyToSymbolAsync.cc +++ /dev/null @@ -1,47 +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 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 -#define SIZE 1024 - -/* Test verifies hipMemcpyToSymbolAsync API Negative scenarios. - */ - -TEST_CASE("Unit_hipMemcpyToSymbolAsync_Negative") { - void *Sd; - char S[SIZE]="This is not a device symbol"; - - HIP_CHECK(hipMalloc(&Sd, SIZE)); - - hipStream_t stream; - HIP_CHECK(hipStreamCreate(&stream)); - - SECTION("Passing void pointer") { - REQUIRE(hipSuccess != hipMemcpyToSymbolAsync(HIP_SYMBOL(Sd), S, - SIZE, 0, hipMemcpyHostToDevice, stream)); - } - - SECTION("Passing NULL pointer") { - REQUIRE(hipSuccess != hipMemcpyToSymbolAsync(nullptr, S, - SIZE, 0, hipMemcpyHostToDevice, stream)); - } - - HIP_CHECK(hipStreamDestroy(stream)); - HIP_CHECK(hipFree(Sd)); -}