diff --git a/catch/hipTestMain/config/config_nvidia_linux_common.json b/catch/hipTestMain/config/config_nvidia_linux_common.json index 691386f1a1..07b62e9329 100644 --- a/catch/hipTestMain/config/config_nvidia_linux_common.json +++ b/catch/hipTestMain/config/config_nvidia_linux_common.json @@ -21,11 +21,16 @@ "Unit_ChannelDescriptor_Positive_Basic_4D - long4", "=== Below test fails in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/38 ===", "Unit_hipFreeAsync_negative", + "=== Below test fails in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/222 ===", + "Unit_hipStreamCreate_Performance", + "Unit_hipStreamCreate_WithPriorityPerformance_Nonblocking_high", + "Unit_hipStreamCreate_WithPriorityPerformance_Default_low", + "Unit_hipStreamCreate_WithPriorityPerformance_Nonblocking_low", "=== Below test fails in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/92 ===", "Unit_hipGetTexObjectResourceDesc_positive", - "Unit_hipGetTexObjectResourceDesc_Negative_Parameters", - "Unit_hipGetTexObjectTextureDesc_positive", - "Unit_hipGetTexObjectTextureDesc_Negative_Parameters", - "Unit_hipTexObjectDestroy_positive" + "Unit_hipGetTexObjectResourceDesc_Negative_Parameters", + "Unit_hipGetTexObjectTextureDesc_positive", + "Unit_hipGetTexObjectTextureDesc_Negative_Parameters", + "Unit_hipTexObjectDestroy_positive" ] } diff --git a/catch/unit/stream/CMakeLists.txt b/catch/unit/stream/CMakeLists.txt index 267b00d38a..5c0f349e27 100644 --- a/catch/unit/stream/CMakeLists.txt +++ b/catch/unit/stream/CMakeLists.txt @@ -17,6 +17,7 @@ set(TEST_SRC hipStreamACb_StrmSyncTiming.cc hipLaunchHostFunc.cc hipStreamGetDevice.cc + hipStreamCreatePerformance.cc ) if(HIP_PLATFORM MATCHES "amd") diff --git a/catch/unit/stream/hipStreamCreatePerformance.cc b/catch/unit/stream/hipStreamCreatePerformance.cc new file mode 100644 index 0000000000..999a4fe7bf --- /dev/null +++ b/catch/unit/stream/hipStreamCreatePerformance.cc @@ -0,0 +1,544 @@ +/* +Copyright (c) 2023 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 + +/** + * @addtogroup hipStreamCreate hipStreamCreate + * @{ + * @ingroup StreamTest + * `hipError_t hipStreamCreate(hipStream_t* stream)` - + * Create a new asynchronous stream. + */ + +/** + * Test Description + * ------------------------ + * - Test case to verify hipStreamCreate performance by recording below sets of time. + * create 4 set of streams and record that time taken as b1 + * create another 4 set of streams before destroying earlier streams and record time taken as b2 + * destroy streams and then create 4 streams again and record time taken as b3 + * and verify if the condition b1 > b2 > b3 is satisfied. + + * Test source + * ------------------------ + * - catch/unit/stream/hipStreamCreatePerformance.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ + +TEST_CASE("Unit_hipStreamCreate_Performance") { + HIP_CHECK(hipSetDevice(0)); // just to initialise HIP runtime. + // create stream + hipStream_t streamb1[4]; + // record time for batch1 stream creation + auto Startb1 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreate(&streamb1[k])); + } + auto Stopb1 = std::chrono::high_resolution_clock::now(); + double performb1 = std::chrono::duration(Stopb1 - Startb1).count(); // NOLINT + printf("Stream create performance for batch1 is %lf\n", performb1); + + // record time for batch2 stream creation before + // destroying already created streams + hipStream_t streamb2[4]; + auto Startb2 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreate(&streamb2[k])); + } + auto Stopb2 = std::chrono::high_resolution_clock::now(); + double performb2 = std::chrono::duration(Stopb2 - Startb2).count(); // NOLINT + printf("Stream create performance for batch2 is %lf\n", performb2); + // destroy batch 1 streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb1[k])); + } + // destroy batch 2 streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb2[k])); + } + // record time for batch3 stream creation after stream destroy + hipStream_t streamb3[4]; + auto Startb3 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreate(&streamb3[k])); + } + auto Stopb3 = std::chrono::high_resolution_clock::now(); + double performb3 = std::chrono::duration(Stopb3 - Startb3).count(); // NOLINT + printf("Stream create performance for batch3 is %lf\n", performb3); + + // destroy streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb3[k])); + } + REQUIRE(performb1 > performb2); + REQUIRE(performb2 > performb3); +} + +/** + * End doxygen group hipStreamCreate. + * @} + */ + +/** + * @addtogroup hipStreamCreateWithFlags hipStreamCreateWithFlags + * @{ + * @ingroup StreamTest + * `hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags)` - + * Create a new asynchronous stream. + */ + +/** + * Test Description + * ------------------------ + * - Test case to verify hipStreamCreateWithFlags performance with + * hipStreamNonBlocking flagby recording below sets of time. + * create 4 set of streams and record that time taken as b1 + * create another 4 set of streams before destroying earlier streams + * and record time taken as b2 + * destroy streams and then create 4 streams again and record time taken as b3 + * and verify if the condition b1 > b2 > b3 is satisfied. + + * Test source + * ------------------------ + * - catch/unit/stream/hipStreamCreatePerformance.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ + +TEST_CASE("Unit_hipStreamCreate_WithFlagsPerformance_Nonblocking") { + HIP_CHECK(hipSetDevice(0)); // just to initialise HIP runtime. + // create stream + hipStream_t streamb1[4]; + // record time for batch1 stream creation + auto Startb1 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreateWithFlags(&streamb1[k], hipStreamNonBlocking)); + } + auto Stopb1 = std::chrono::high_resolution_clock::now(); + double performb1 = std::chrono::duration(Stopb1 - Startb1).count(); // NOLINT + printf("Stream create performance for batch1 is %lf\n", performb1); + + // record time for batch2 stream creation before + // destroying already created streams + hipStream_t streamb2[4]; + auto Startb2 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreateWithFlags(&streamb2[k], hipStreamNonBlocking)); + } + auto Stopb2 = std::chrono::high_resolution_clock::now(); + double performb2 = std::chrono::duration(Stopb2 - Startb2).count(); // NOLINT + printf("Stream create performance for batch2 is %lf\n", performb2); + // destroy batch 1 streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb1[k])); + } + // destroy batch 2 streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb2[k])); + } + // record time for batch3 stream creation after stream destroy + hipStream_t streamb3[4]; + auto Startb3 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreateWithFlags(&streamb3[k], hipStreamNonBlocking)); + } + auto Stopb3 = std::chrono::high_resolution_clock::now(); + double performb3 = std::chrono::duration(Stopb3 - Startb3).count(); // NOLINT + printf("Stream create performance for batch3 is %lf\n", performb3); + + // destroy streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb3[k])); + } + REQUIRE(performb1 > performb2); + REQUIRE(performb2 > performb3); +} + +/** + * Test Description + * ------------------------ + * - Test case to verify hipStreamCreateWithFlags performance + * with hipStreamDefault flagby recording below sets of time. + * create 4 set of streams and record that time taken as b1 + * create another 4 set of streams before destroying earlier streams + * and record time taken as b2 + * destroy streams and then create 4 streams again and record time taken as b3 + * and verify if the condition b1 > b2 > b3 is satisfied. + + * Test source + * ------------------------ + * - catch/unit/stream/hipStreamCreatePerformance.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ + +TEST_CASE("Unit_hipStreamCreate_WithFlagsPerformance_Default") { + HIP_CHECK(hipSetDevice(0)); // just to initialise HIP runtime. + // create stream + hipStream_t streamb1[4]; + // record time for batch1 stream creation + auto Startb1 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreateWithFlags(&streamb1[k], hipStreamDefault)); + } + auto Stopb1 = std::chrono::high_resolution_clock::now(); + double performb1 = std::chrono::duration(Stopb1 - Startb1).count(); // NOLINT + printf("Stream create performance for batch1 is %lf\n", performb1); + + // record time for batch2 stream creation before + // destroying already created streams + hipStream_t streamb2[4]; + auto Startb2 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreateWithFlags(&streamb2[k], hipStreamDefault)); + } + auto Stopb2 = std::chrono::high_resolution_clock::now(); + double performb2 = std::chrono::duration(Stopb2 - Startb2).count(); // NOLINT + printf("Stream create performance for batch2 is %lf\n", performb2); + // destroy batch 1 streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb1[k])); + } + // destroy batch 2 streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb2[k])); + } + // record time for batch3 stream creation after stream destroy + hipStream_t streamb3[4]; + auto Startb3 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreateWithFlags(&streamb3[k], hipStreamDefault)); + } + auto Stopb3 = std::chrono::high_resolution_clock::now(); + double performb3 = std::chrono::duration(Stopb3 - Startb3).count(); // NOLINT + printf("Stream create performance for batch3 is %lf\n", performb3); + + // destroy streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb3[k])); + } + REQUIRE(performb1 > performb2); + REQUIRE(performb2 > performb3); +} +/** + * End doxygen group hipStreamCreateWithFlags. + * @} + */ + +/** + * @addtogroup hipStreamCreateWithPriority hipStreamCreateWithPriority + * @{ + * @ingroup StreamTest + * `hipError_t hipStreamCreateWithPriority(hipStream_t* stream, unsigned int flags, int priority)` - + * Create a new asynchronous stream. + */ + +/** + * Test Description + * ------------------------ + * - Test case to verify hipStreamCreateWithPriority performance + * with hipStreamNonBlocking flag along with low priority + * by recording below sets of time. + * create 4 set of streams and record that time taken as b1 + * create another 4 set of streams before destroying earlier streams + * and record time taken as b2 + * destroy streams and then create 4 streams again and record time taken as b3 + * and verify if the condition b1 > b2 > b3 is satisfied. + + * Test source + * ------------------------ + * - catch/unit/stream/hipStreamCreatePerformance.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ + +TEST_CASE("Unit_hipStreamCreate_WithPriorityPerformance_Nonblocking_low") { + int priority_low, priority_high; + HIP_CHECK(hipSetDevice(0)); // just to initialise HIP runtime. + HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); + // create stream + hipStream_t streamb1[4]; + // record time for batch1 stream creation + auto Startb1 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreateWithPriority(&streamb1[k], hipStreamNonBlocking, priority_low)); + } + auto Stopb1 = std::chrono::high_resolution_clock::now(); + double performb1 = std::chrono::duration(Stopb1 - Startb1).count(); // NOLINT + printf("Stream create performance for batch1 is %lf\n", performb1); + + // record time for batch2 stream creation before + // destroying already created streams + hipStream_t streamb2[4]; + auto Startb2 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreateWithPriority(&streamb2[k], hipStreamNonBlocking, priority_low)); + } + auto Stopb2 = std::chrono::high_resolution_clock::now(); + double performb2 = std::chrono::duration(Stopb2 - Startb2).count(); // NOLINT + printf("Stream create performance for batch2 is %lf\n", performb2); + // destroy batch 1 streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb1[k])); + } + // destroy batch 2 streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb2[k])); + } + // record time for batch3 stream creation after stream destroy + hipStream_t streamb3[4]; + auto Startb3 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreateWithPriority(&streamb3[k], hipStreamNonBlocking, priority_low)); + } + auto Stopb3 = std::chrono::high_resolution_clock::now(); + double performb3 = std::chrono::duration(Stopb3 - Startb3).count(); // NOLINT + printf("Stream create performance for batch3 is %lf\n", performb3); + + // destroy streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb3[k])); + } + REQUIRE(performb1 > performb2); + REQUIRE(performb2 > performb3); +} + +/** + * Test Description + * ------------------------ + * - Test case to verify hipStreamCreateWithPriority performance + * with hipStreamNonBlocking flag along with high priority + * by recording below sets of time. + * create 4 set of streams and record that time taken as b1 + * create another 4 set of streams before destroying earlier streams + * and record time taken as b2 + * destroy streams and then create 4 streams again and record time taken as b3 + * and verify if the condition b1 > b2 > b3 is satisfied. + + * Test source + * ------------------------ + * - catch/unit/stream/hipStreamCreatePerformance.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ + +TEST_CASE("Unit_hipStreamCreate_WithPriorityPerformance_Nonblocking_high") { + int priority_low, priority_high; + HIP_CHECK(hipSetDevice(0)); // just to initialise HIP runtime. + HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); + // create stream + hipStream_t streamb1[4]; + // record time for batch1 stream creation + auto Startb1 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreateWithPriority(&streamb1[k], hipStreamNonBlocking, priority_high)); + } + auto Stopb1 = std::chrono::high_resolution_clock::now(); + double performb1 = std::chrono::duration(Stopb1 - Startb1).count(); // NOLINT + printf("Stream create performance for batch1 is %lf\n", performb1); + + // record time for batch2 stream creation before + // destroying already created streams + hipStream_t streamb2[4]; + auto Startb2 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreateWithPriority(&streamb2[k], hipStreamNonBlocking, priority_high)); + } + auto Stopb2 = std::chrono::high_resolution_clock::now(); + double performb2 = std::chrono::duration(Stopb2 - Startb2).count(); // NOLINT + printf("Stream create performance for batch2 is %lf\n", performb2); + // destroy batch 1 streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb1[k])); + } + // destroy batch 2 streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb2[k])); + } + // record time for batch3 stream creation after stream destroy + hipStream_t streamb3[4]; + auto Startb3 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreateWithPriority(&streamb3[k], hipStreamNonBlocking, priority_high)); + } + auto Stopb3 = std::chrono::high_resolution_clock::now(); + double performb3 = std::chrono::duration(Stopb3 - Startb3).count(); // NOLINT + printf("Stream create performance for batch3 is %lf\n", performb3); + + // destroy streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb3[k])); + } + REQUIRE(performb1 > performb2); + REQUIRE(performb2 > performb3); +} + +/** + * Test Description + * ------------------------ + * - Test case to verify hipStreamCreateWithPriority performance + * with hipStreamDefault flag along with low priority by recording below sets of time. + * create 4 set of streams and record that time taken as b1 + * create another 4 set of streams before destroying earlier streams + * and record time taken as b2 + * destroy streams and then create 4 streams again and record time taken as b3 + * and verify if the condition b1 > b2 > b3 is satisfied. + + * Test source + * ------------------------ + * - catch/unit/stream/hipStreamCreatePerformance.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ + +TEST_CASE("Unit_hipStreamCreate_WithPriorityPerformance_Default_low") { + int priority_low, priority_high; + HIP_CHECK(hipSetDevice(0)); // just to initialise HIP runtime. + HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); + // create stream + hipStream_t streamb1[4]; + // record time for batch1 stream creation + auto Startb1 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreateWithPriority(&streamb1[k], hipStreamDefault, priority_low)); + } + auto Stopb1 = std::chrono::high_resolution_clock::now(); + double performb1 = std::chrono::duration(Stopb1 - Startb1).count(); // NOLINT + printf("Stream create performance for batch1 is %lf\n", performb1); + + // record time for batch2 stream creation before + // destroying already created streams + hipStream_t streamb2[4]; + auto Startb2 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreateWithPriority(&streamb2[k], hipStreamDefault, priority_low)); + } + auto Stopb2 = std::chrono::high_resolution_clock::now(); + double performb2 = std::chrono::duration(Stopb2 - Startb2).count(); // NOLINT + printf("Stream create performance for batch2 is %lf\n", performb2); + // destroy batch 1 streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb1[k])); + } + // destroy batch 2 streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb2[k])); + } + // record time for batch3 stream creation after stream destroy + hipStream_t streamb3[4]; + auto Startb3 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreateWithPriority(&streamb3[k], hipStreamDefault, priority_low)); + } + auto Stopb3 = std::chrono::high_resolution_clock::now(); + double performb3 = std::chrono::duration(Stopb3 - Startb3).count(); // NOLINT + printf("Stream create performance for batch3 is %lf\n", performb3); + + // destroy streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb3[k])); + } + REQUIRE(performb1 > performb2); + REQUIRE(performb2 > performb3); +} + +/** + * Test Description + * ------------------------ + * - Test case to verify hipStreamCreateWithPriority performance + * with hipStreamDefault flag along with high priority by recording below sets of time. + * create 4 set of streams and record that time taken as b1 + * create another 4 set of streams before destroying earlier streams + * and record time taken as b2 + * destroy streams and then create 4 streams again and record time taken as b3 + * and verify if the condition b1 > b2 > b3 is satisfied. + + * Test source + * ------------------------ + * - catch/unit/stream/hipStreamCreatePerformance.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ + +TEST_CASE("Unit_hipStreamCreate_WithPriorityPerformance_Default_high") { + int priority_low = 1, priority_high = 2; + HIP_CHECK(hipSetDevice(0)); // just to initialise HIP runtime. + HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); + // create stream + hipStream_t streamb1[4]; + // record time for batch1 stream creation + auto Startb1 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreateWithPriority(&streamb1[k], hipStreamDefault, priority_high)); + } + auto Stopb1 = std::chrono::high_resolution_clock::now(); + double performb1 = std::chrono::duration(Stopb1 - Startb1).count(); // NOLINT + printf("Stream create performance for batch1 is %lf\n", performb1); + + // record time for batch2 stream creation before + // destroying already created streams + hipStream_t streamb2[4]; + auto Startb2 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreateWithPriority(&streamb2[k], hipStreamDefault, priority_high)); + } + auto Stopb2 = std::chrono::high_resolution_clock::now(); + double performb2 = std::chrono::duration(Stopb2 - Startb2).count(); // NOLINT + printf("Stream create performance for batch2 is %lf\n", performb2); + // destroy batch 1 streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb1[k])); + } + // destroy batch 2 streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb2[k])); + } + // record time for batch3 stream creation after stream destroy + hipStream_t streamb3[4]; + auto Startb3 = std::chrono::high_resolution_clock::now(); + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamCreateWithPriority(&streamb3[k], hipStreamDefault, priority_high)); + } + auto Stopb3 = std::chrono::high_resolution_clock::now(); + double performb3 = std::chrono::duration(Stopb3 - Startb3).count(); // NOLINT + printf("Stream create performance for batch3 is %lf\n", performb3); + + // destroy streams + for (int k = 0; k < 4; k++) { + HIP_CHECK(hipStreamDestroy(streamb3[k])); + } + REQUIRE(performb1 > performb2); + REQUIRE(performb2 > performb3); +} +/** + * End doxygen group hipStreamCreateWithPriority. + * @} + */