diff --git a/projects/hip-tests/catch/unit/CMakeLists.txt b/projects/hip-tests/catch/unit/CMakeLists.txt index 167a984c30..0e8b37c993 100644 --- a/projects/hip-tests/catch/unit/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/CMakeLists.txt @@ -21,6 +21,7 @@ add_subdirectory(memory) add_subdirectory(deviceLib) add_subdirectory(stream) +add_subdirectory(streamperthread) add_subdirectory(event) add_subdirectory(occupancy) add_subdirectory(device) diff --git a/projects/hip-tests/catch/unit/streamperthread/CMakeLists.txt b/projects/hip-tests/catch/unit/streamperthread/CMakeLists.txt new file mode 100644 index 0000000000..51a789d426 --- /dev/null +++ b/projects/hip-tests/catch/unit/streamperthread/CMakeLists.txt @@ -0,0 +1,13 @@ +# Common Tests - Test independent of all platforms +set(TEST_SRC + hipStreamPerThread_Basic.cc + hipStreamPerThread_Event.cc + hipStreamPerThread_MultiThread.cc + hipStreamPerThread_DeviceReset.cc +) + +# Create shared lib of all tests +add_library(StreamPerThreadTest SHARED EXCLUDE_FROM_ALL ${TEST_SRC}) + +# Add dependency on build_tests to build it on this custom target +add_dependencies(build_tests StreamPerThreadTest) diff --git a/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThread_Basic.cc b/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThread_Basic.cc new file mode 100644 index 0000000000..3780631141 --- /dev/null +++ b/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThread_Basic.cc @@ -0,0 +1,133 @@ +/* +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, INNCLUDING 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 ANNY 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 + +#define MEM_SIZE (1024*1024*32) +#define SEED 5 + +constexpr unsigned int MAX_THREAD_CNT = 10; + +__global__ void copy_kernl(int* devPtr) { + for (int i = 0; i < MEM_SIZE; ++i) { + devPtr[i] = (i+1) + SEED; + } +} + +TEST_CASE("Unit_hipStreamPerThread_Basic") { + constexpr int size = sizeof(int) * MEM_SIZE; + int* hostMem = nullptr; + int* devMem = nullptr; + + HIP_CHECK(hipHostMalloc(&hostMem, size)); + HIP_CHECK(hipMalloc(&devMem, size)); + + // Init host mem with different values + for (int i = 0; i < MEM_SIZE; ++i) { + hostMem[i] = i; + } + + /* + hipStreamPerThread is an implicit stream which works independent of null stream. + Null stream synchronize will account hipStreamPerThread into account. + test scenario: Launch kernel + Async mem copy on hipStreamPerThread and call synchronize on null stream. + Result : Null stream synchronize should sync hipStreamPerThread as well + */ + copy_kernl<<<1, 1, 0, hipStreamPerThread>>>(devMem); + + HIP_CHECK(hipMemcpyAsync(hostMem, devMem, size, hipMemcpyDeviceToHost, hipStreamPerThread)); + + HIP_CHECK(hipStreamSynchronize(0)); + + // validate result + for (int i = MEM_SIZE-1; i >= 0; --i) { + CHECK(hostMem[i] == (i+1+SEED)); + } +} + +TEST_CASE("Unit_hipStreamPerThread_StreamQuery") { + std::vector threads(MAX_THREAD_CNT); + + for (auto &th : threads) { + th = std::thread([](){HIP_CHECK(hipStreamQuery(hipStreamPerThread));}); + } + + for (auto& th : threads) { + th.join(); + } + REQUIRE(true); +} + +TEST_CASE("Unit_hipStreamPerThread_StreamSynchronize") { + constexpr unsigned int MAX_THREAD_CNT = 10; + std::vector threads(MAX_THREAD_CNT); + + for (auto &th : threads) { + th = std::thread([](){HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));}); + } + + for (auto& th : threads) { + th.join(); + } + REQUIRE(true); +} + +TEST_CASE("Unit_hipStreamPerThread_StreamGetPriority") { + int priority = 0; + HIP_CHECK(hipStreamGetPriority(hipStreamPerThread, &priority)); +} + +TEST_CASE("Unit_hipStreamPerThread_StreamGetFlags") { + unsigned int flags = 0; + HIP_CHECK(hipStreamGetFlags(hipStreamPerThread, &flags)); +} + +TEST_CASE("Unit_hipStreamPerThread_StreamDestroy") { + hipError_t status = hipStreamDestroy(hipStreamPerThread); + REQUIRE(status != hipSuccess); +} + +TEST_CASE("Unit_hipStreamPerThread_MemcpyAsync") { + unsigned int ele_size = (16 * 1024); // 16KB + int* A_h = nullptr; + int* A_d = nullptr; + + HIP_CHECK(hipHostMalloc(&A_h, ele_size*sizeof(int))); + HIP_CHECK(hipMalloc(&A_d, ele_size * sizeof(int))); + + for (unsigned int i = 0; i < ele_size; ++i) { + A_h[i] = 123; + } + + HIP_CHECK(hipMemcpy(A_d, A_h, ele_size * sizeof(int), hipMemcpyHostToDevice)); + + // Rest host memory + for (unsigned int i = 0; i < ele_size; ++i) { + A_h[i] = 0; + } + + HIP_CHECK(hipMemcpyAsync(A_h, A_d, ele_size * sizeof(int), hipMemcpyDeviceToHost, + hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + + // Verify result + for (unsigned int i = 0; i < ele_size; ++i) { + REQUIRE(A_h[i] == 123); + } +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThread_DeviceReset.cc b/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThread_DeviceReset.cc new file mode 100644 index 0000000000..837446d2a9 --- /dev/null +++ b/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThread_DeviceReset.cc @@ -0,0 +1,89 @@ +/* +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, INNCLUDING 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 ANNY 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 + +/* + hipDeviceReset deletes all active streams including hipStreamPerThread. + Scenario: App calls hipDeviceReset while in other thread some Async operation is in + progress on hipStreamPerThread. + Watch out: hipDeviceRest should be successfull without any crash + */ +static void Copy_to_device() { + unsigned int ele_size = (32 * 1024); // 32KB + int* A_h = nullptr; + int* A_d = nullptr; + + hipError_t status = hipHostMalloc(&A_h, ele_size*sizeof(int)); + if (status != hipSuccess) return; + + status = hipMalloc(&A_d, ele_size * sizeof(int)); + if (status != hipSuccess) return; + + for(unsigned int i = 0; i < ele_size; ++i) { + A_h[i] = 123; + } + hipMemcpyAsync(A_d, A_h, ele_size * sizeof(int), hipMemcpyHostToDevice, + hipStreamPerThread); +} + +TEST_CASE("Unit_hipStreamPerThread_DeviceReset_1") { + constexpr unsigned int MAX_THREAD_CNT = 10; + std::vector threads(MAX_THREAD_CNT); + + for (auto &th : threads) { + th = std::thread(Copy_to_device); + th.detach(); + } + HIP_CHECK(hipDeviceReset()); +} + +/* + hipDeviceReset deletes all active streams including hipStreamPerThread. + Scenario: i) Launch Async task on hipStreamPerThread and waits for it to complete. + ii) Call hipDeviceReset to delete all active stream + iii) Again try to launch Async task on hipStreamPerThread + Watch out: Since hipStreamPerThread is an implicit stream hence even after device reset + it should available to use. + */ +TEST_CASE("Unit_hipStreamPerThread_DeviceReset_2") { + unsigned int ele_size = (32 * 1024); // 32KB + int* A_h = nullptr; + int* A_d = nullptr; + + hipError_t status = hipHostMalloc(&A_h, ele_size*sizeof(int)); + if (status != hipSuccess) return; + status = hipMalloc(&A_d, ele_size * sizeof(int)); + if (status != hipSuccess) return; + + for (unsigned int i = 0; i < ele_size; ++i) { + A_h[i] = 123; + } + hipMemcpyAsync(A_d, A_h, ele_size * sizeof(int), hipMemcpyHostToDevice, + hipStreamPerThread); + hipStreamSynchronize(hipStreamPerThread); + + hipDeviceReset(); + + hipMemcpyAsync(A_d, A_h, ele_size * sizeof(int), hipMemcpyHostToDevice, + hipStreamPerThread); + hipStreamSynchronize(hipStreamPerThread); +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThread_Event.cc b/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThread_Event.cc new file mode 100644 index 0000000000..6fef9914eb --- /dev/null +++ b/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThread_Event.cc @@ -0,0 +1,64 @@ +/* +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, INNCLUDING 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 ANNY 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 + +TEST_CASE("Unit_hipStreamPerThread_EventRecord") { + hipEvent_t event; + HIP_CHECK(hipEventCreate(&event)); + HIP_CHECK(hipEventRecord(event, hipStreamPerThread)); +} + +__global__ void update_even_odd(unsigned int N, int* out) { + for (unsigned int i = 0; i < N; ++i) { + if (i%2 == 0) { + out[i] = 2; + } else { + out[i] = 3; + } + } +} +TEST_CASE("Unit_hipStreamPerThread_EventSynchronize") { + int* A_h = nullptr; + int* A_d = nullptr; + unsigned int size = 1000; + + HIP_CHECK(hipHostMalloc(&A_h, size*sizeof(int))); + HIP_CHECK(hipMalloc(&A_d, size * sizeof(int))); + + hipEvent_t start, end; + HIP_CHECK(hipEventCreate(&start)); + HIP_CHECK(hipEventCreate(&end)); + + HIP_CHECK(hipEventRecord(start, hipStreamPerThread)); + update_even_odd<<<1, 1>>>(size, A_d); + HIP_CHECK(hipEventRecord(end, hipStreamPerThread)); + + HIP_CHECK(hipEventSynchronize(end)); + HIP_CHECK(hipMemcpy(A_h, A_d, size*sizeof(int), hipMemcpyDeviceToHost)); + + // Verify result + for (unsigned int i = 0; i < size; ++i) { + if (i%2 == 0 && A_h[i] != 2) + REQUIRE(false); + else if (i%2 != 0 && A_h[i] != 3) { + REQUIRE(false); + } + } +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThread_MultiThread.cc b/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThread_MultiThread.cc new file mode 100644 index 0000000000..a0ff7ed6be --- /dev/null +++ b/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThread_MultiThread.cc @@ -0,0 +1,55 @@ +/* +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, INNCLUDING 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 ANNY 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 + +static void Copy_to_device() { + unsigned int ele_size = (32 * 1024); // 32KB + int* A_h = nullptr; + int* A_d = nullptr; + + hipHostMalloc(&A_h, ele_size*sizeof(int)); + hipMalloc(&A_d, ele_size * sizeof(int)); + + for (unsigned int i = 0; i < ele_size; ++i) { + A_h[i] = 123; + } + hipMemcpyAsync(A_d, A_h, ele_size * sizeof(int), hipMemcpyHostToDevice, + hipStreamPerThread); +} + +/* +hipStreamPerThread is an implicit stream which gets destroyed once thread is completed. +Scenario : App pushes Async task(s) into hipStreamPerThread and did not wait for it to complete. +Watch out : Incomplete task in hipStreamPerThread should not cause any crash due to thread exit. + */ +TEST_CASE("Unit_hipStreamPerThread_MultiThread") { + constexpr unsigned int MAX_THREAD_CNT = 10; + std::vector threads(MAX_THREAD_CNT); + + for (auto &th : threads) { + th = std::thread(Copy_to_device); + } + + for (auto& th : threads) { + th.detach(); + } +}