From 617f19e6d96977ffa40fa748da200668aa022b4e Mon Sep 17 00:00:00 2001 From: Sarbojit Sarkar Date: Fri, 25 Jun 2021 08:28:59 +0000 Subject: [PATCH] SWDEV-291783 - Catch2 tests for some of Stream APIs Change-Id: I656fc59184863ef142cf89d7541982492f6b8484 --- catch/hipTestMain/CMakeLists.txt | 1 + catch/include/hip_test_common.hh | 6 +- catch/unit/CMakeLists.txt | 1 + catch/unit/stream/CMakeLists.txt | 12 +++ catch/unit/stream/hipMultiStream.cc | 97 +++++++++++++++++++++++ catch/unit/stream/hipStreamCreate.cc | 52 ++++++++++++ catch/unit/stream/hipStreamGetFlags.cc | 38 +++++++++ catch/unit/stream/hipStreamGetPriority.cc | 92 +++++++++++++++++++++ 8 files changed, 298 insertions(+), 1 deletion(-) create mode 100755 catch/unit/stream/CMakeLists.txt create mode 100755 catch/unit/stream/hipMultiStream.cc create mode 100755 catch/unit/stream/hipStreamCreate.cc create mode 100755 catch/unit/stream/hipStreamGetFlags.cc create mode 100755 catch/unit/stream/hipStreamGetPriority.cc diff --git a/catch/hipTestMain/CMakeLists.txt b/catch/hipTestMain/CMakeLists.txt index 2bee9c9adf..37d7d6568d 100644 --- a/catch/hipTestMain/CMakeLists.txt +++ b/catch/hipTestMain/CMakeLists.txt @@ -11,6 +11,7 @@ endif() target_link_libraries(UnitTests PRIVATE UnitDeviceTests MemoryTest + StreamTest stdc++fs) # Add AMD Only Tests diff --git a/catch/include/hip_test_common.hh b/catch/include/hip_test_common.hh index dc2d6980d5..135d220d9a 100644 --- a/catch/include/hip_test_common.hh +++ b/catch/include/hip_test_common.hh @@ -43,5 +43,9 @@ THE SOFTWARE. // Utility Functions namespace HipTest { -int getDeviceCount(); +static inline int getDeviceCount() { + int dev = 0; + HIP_CHECK(hipGetDeviceCount(&dev)); + return dev; +} } diff --git a/catch/unit/CMakeLists.txt b/catch/unit/CMakeLists.txt index 323feaa277..5aca994a4c 100644 --- a/catch/unit/CMakeLists.txt +++ b/catch/unit/CMakeLists.txt @@ -1,4 +1,5 @@ add_subdirectory(memory) add_subdirectory(deviceLib) +add_subdirectory(stream) # Disable Saxpy test temporarily to see if CI Passes # add_subdirectory(rtc) diff --git a/catch/unit/stream/CMakeLists.txt b/catch/unit/stream/CMakeLists.txt new file mode 100755 index 0000000000..4b8e57a112 --- /dev/null +++ b/catch/unit/stream/CMakeLists.txt @@ -0,0 +1,12 @@ +set(TEST_SRC + hipStreamCreate.cc + hipStreamGetFlags.cc + hipStreamGetPriority.cc + hipMultiStream.cc +) + +# Create shared lib of all tests +add_library(StreamTest SHARED EXCLUDE_FROM_ALL ${TEST_SRC}) + +# Add dependency on build_tests to build it on this custom target +add_dependencies(build_tests StreamTest) \ No newline at end of file diff --git a/catch/unit/stream/hipMultiStream.cc b/catch/unit/stream/hipMultiStream.cc new file mode 100755 index 0000000000..c7572dd2c1 --- /dev/null +++ b/catch/unit/stream/hipMultiStream.cc @@ -0,0 +1,97 @@ + +/* +Copyright (c) 2021-Present 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 +constexpr int NN = 1 << 21; +__global__ void kernel_do_nothing(__attribute__((unused))int a) { + // empty kernel +} +__global__ void kernel(float* x, float* y, int n) { + size_t tid{threadIdx.x}; + if (tid < 1) { + for (int i = 0; i < n; i++) { + x[i] = sqrt(powf(3.14159, i)); + } + y[tid] = y[tid] + 1.0f; + } +} +__global__ void nKernel(float* y) { + size_t tid{threadIdx.x}; + y[tid] = y[tid] + 1.0f; +} +TEST_CASE("Unit_hipMultiStream_sameDevice") { + constexpr int num_streams{8}; + hipStream_t streams[num_streams]; + float *data[num_streams], *yd, *xd; + float y{1.0f}, x{1.0f}; + HIP_CHECK(hipMalloc((void**)&yd, sizeof(float))); + HIP_CHECK(hipMalloc((void**)&xd, sizeof(float))); + HIP_CHECK(hipMemcpy(yd, &y, sizeof(float), hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(xd, &x, sizeof(float), hipMemcpyHostToDevice)); + for (int i = 0; i < num_streams; i++) { + HIP_CHECK(hipStreamCreate(&streams[i])); + HIP_CHECK(hipMalloc(&data[i], NN * sizeof(float))); + hipLaunchKernelGGL(kernel, dim3(1), dim3(1), 0, streams[i], data[i], xd, NN); + hipLaunchKernelGGL(HIP_KERNEL_NAME(nKernel), dim3(1), dim3(1), 0, 0, yd); + } + HIP_CHECK(hipMemcpy(&x, xd, sizeof(float), hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(&y, yd, sizeof(float), hipMemcpyDeviceToHost)); + REQUIRE(x == Approx(y)); +} +TEST_CASE("Unit_hipMultiStream_multimeDevice") { + constexpr int nLoops = 100000; + constexpr int nStreams = 2; + std::vector streams(nStreams); + int nGpu = 0; + HIP_CHECK(hipGetDeviceCount(&nGpu)); + if (nGpu < 1) { + std::cout << "info: didn't find any GPU! skipping the test!\n"; + REQUIRE(true); + } + static int device = 0; + HIP_CHECK(hipSetDevice(device)); + hipDeviceProp_t props; + HIP_CHECK(hipGetDeviceProperties(&props, device)); + std::cout << "info: running on bus " + << "0x" << props.pciBusID << " " << props.name << std::endl; + for (int i = 0; i < nStreams; i++) { + HIP_CHECK(hipStreamCreate(&streams[i])); + } + for (int k = 0; k <= nLoops; ++k) { + HIP_CHECK(hipDeviceSynchronize()); + // Launch kernel with default stream + hipLaunchKernelGGL(kernel_do_nothing, dim3(1), dim3(1), 0, 0, 1); + // Launch kernel on all streams + for (int i = 0; i < nStreams; i++) { + hipLaunchKernelGGL(kernel_do_nothing, dim3(1), dim3(1), 0, streams[i], 1); + } + // Sync stream 1 + HIP_CHECK(hipStreamSynchronize(streams[0])); + if (k % 10000 == 0 || k == nLoops) { + std::cout << "Info: Iteration = " << k << std::endl; + } + } + HIP_CHECK(hipDeviceSynchronize()); + // Clean up + for (int i = 0; i < nStreams; i++) { + HIP_CHECK(hipStreamDestroy(streams[i])); + } +} diff --git a/catch/unit/stream/hipStreamCreate.cc b/catch/unit/stream/hipStreamCreate.cc new file mode 100755 index 0000000000..59033ae63b --- /dev/null +++ b/catch/unit/stream/hipStreamCreate.cc @@ -0,0 +1,52 @@ + +/* +Copyright (c) 2021-Present 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 +TEST_CASE("Unit_hipStreamCreate_default") { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); +} +TEST_CASE("Unit_hipStreamCreateWithFlags_Negative") { + hipStream_t stream; + auto status = hipStreamCreateWithFlags(&stream, 0xFF); + REQUIRE(status == hipErrorInvalidValue); + status = hipStreamCreateWithFlags(nullptr, hipStreamDefault); + REQUIRE(status == hipErrorInvalidValue); +} +TEST_CASE("Unit_hipStreamCreateWithFlags") { + hipStream_t stream; + HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); + HIP_CHECK(hipStreamDestroy(stream)); +} +TEST_CASE("Unit_hipStreamCreateWithPriority") { + int priority_low = 0; + int priority_high = 0; + HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); + hipStream_t stream; + + SECTION("Setting high prirority") { + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_high)); + } + SECTION("Setting low priority") { + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_low)); + } + HIP_CHECK(hipStreamDestroy(stream)); +} diff --git a/catch/unit/stream/hipStreamGetFlags.cc b/catch/unit/stream/hipStreamGetFlags.cc new file mode 100755 index 0000000000..824aa9858d --- /dev/null +++ b/catch/unit/stream/hipStreamGetFlags.cc @@ -0,0 +1,38 @@ + +/* +Copyright (c) 2021-Present 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 +TEST_CASE("Unit_hipStreamGetFlags_Negative") { + // Get flags for uninitialized stream + hipStream_t stream; + HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault)); + REQUIRE(hipStreamGetFlags(stream, nullptr) == hipErrorInvalidValue); +} +TEST_CASE("Unit_hipStreamGetFlags") { + hipStream_t stream; + unsigned int flags; + HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault)); + HIP_CHECK(hipStreamGetFlags(stream, &flags)); + REQUIRE(flags == hipStreamDefault); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); + HIP_CHECK(hipStreamGetFlags(stream, &flags)); + REQUIRE(flags == hipStreamNonBlocking); + HIP_CHECK(hipStreamDestroy(stream)); +} diff --git a/catch/unit/stream/hipStreamGetPriority.cc b/catch/unit/stream/hipStreamGetPriority.cc new file mode 100755 index 0000000000..8394ffee84 --- /dev/null +++ b/catch/unit/stream/hipStreamGetPriority.cc @@ -0,0 +1,92 @@ + +/* +Copyright (c) 2021-Present 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 +TEST_CASE("Unit_hipStreamGetPriority_Negative") { + hipStream_t stream = 0; + REQUIRE(hipStreamGetPriority(stream, nullptr) == hipErrorInvalidValue); +} +TEST_CASE("Unit_hipStreamGetPriority_default") { + int priority_low = 0; + int priority_high = 0; + int devID = GENERATE(range(0, HipTest::getDeviceCount())); + HIP_CHECK(hipSetDevice(devID)); + HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + int priority = 0; + HIP_CHECK(hipStreamGetPriority(stream, &priority)); + // valid priority + // Lower the value higher the priority, higher the value lower the priority + REQUIRE(priority_low >= priority); + REQUIRE(priority >= priority_high); + HIP_CHECK(hipStreamDestroy(stream)); +} +TEST_CASE("Unit_hipStreamGetPriority_high") { + int priority_low = 0; + int priority_high = 0; + int devID = GENERATE(range(0, HipTest::getDeviceCount())); + HIP_CHECK(hipSetDevice(devID)); + HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); + hipStream_t stream; + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_high)); + int priority = 0; + HIP_CHECK(hipStreamGetPriority(stream, &priority)); + REQUIRE(priority == priority_high); + HIP_CHECK(hipStreamDestroy(stream)); +} +TEST_CASE("Unit_hipStreamGetPriority_higher") { + int priority_low = 0; + int priority_high = 0; + int devID = GENERATE(range(0, HipTest::getDeviceCount())); + HIP_CHECK(hipSetDevice(devID)); + HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); + hipStream_t stream; + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_high-1)); + int priority = 0; + HIP_CHECK(hipStreamGetPriority(stream, &priority)); + REQUIRE(priority == priority_high); + HIP_CHECK(hipStreamDestroy(stream)); +} +TEST_CASE("Unit_hipStreamGetPriority_low") { + int priority_low = 0; + int priority_high = 0; + int devID = GENERATE(range(0, HipTest::getDeviceCount())); + HIP_CHECK(hipSetDevice(devID)); + HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); + hipStream_t stream; + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_low)); + int priority = 0; + HIP_CHECK(hipStreamGetPriority(stream, &priority)); + REQUIRE(priority_low == priority); + HIP_CHECK(hipStreamDestroy(stream)); +} +TEST_CASE("Unit_hipStreamGetPriority_lower") { + int priority_low = 0; + int priority_high = 0; + int devID = GENERATE(range(0, HipTest::getDeviceCount())); + HIP_CHECK(hipSetDevice(devID)); + HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); + hipStream_t stream; + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_low+1)); + int priority = 0; + HIP_CHECK(hipStreamGetPriority(stream, &priority)); + REQUIRE(priority_low == priority); + HIP_CHECK(hipStreamDestroy(stream)); +}