diff --git a/projects/hip-tests/catch/include/hip_test_defgroups.hh b/projects/hip-tests/catch/include/hip_test_defgroups.hh index 77e389b5da..4fa9facedb 100644 --- a/projects/hip-tests/catch/include/hip_test_defgroups.hh +++ b/projects/hip-tests/catch/include/hip_test_defgroups.hh @@ -42,3 +42,10 @@ THE SOFTWARE. * This section describes the warp shuffle types & functions of HIP runtime API. * @} */ + +/** + * @defgroup StreamTest Stream Management + * @{ + * This section describes the stream management types & functions of HIP runtime API. + * @} + */ diff --git a/projects/hip-tests/catch/unit/stream/hipStreamCreateWithPriority.cc b/projects/hip-tests/catch/unit/stream/hipStreamCreateWithPriority.cc index 39b4b98061..03b34e4da1 100644 --- a/projects/hip-tests/catch/unit/stream/hipStreamCreateWithPriority.cc +++ b/projects/hip-tests/catch/unit/stream/hipStreamCreateWithPriority.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022-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 @@ -18,38 +18,20 @@ THE SOFTWARE. */ /** -Testcase Scenarios : + * @addtogroup hipStreamCreateWithPriority hipStreamCreateWithPriority + * @{ + * @ingroup StreamTest + * `hipStreamCreateWithPriority (hipStream_t *stream, unsigned int flags, int priority)` - + * begins graph capture on a stream + */ -1)Create streams with default flag for all available priority levels and -queue tasks in each of these streams, perform device synchronize and validate behavior. - -2)Create streams with non-blocking flag for all available priority levels and -queue tasks in each of these streams, perform stream synchronize and validate behavior. - -3)Create streams with default flag for all available priority levels and -queue tasks in each of these streams, perform stream synchronize and validate behavior. - -4)Create streams with non-blocking flag for all available priority levels and -queue tasks in each of these streams, perform device synchronize and validate behavior. - -5)Create a stream for each priority level with default flag, Launch memcpy and kernel -tasks on these streams from multiple threads. Validate all the results. - -6)Create a stream for each priority level with non-blocking flag, Launch memcpy and -kernel tasks on these streams from multiple threads. Validate all the results. - -7) Validate negative scenarios for hipStreamCreateWithPriority api. - -8) Validate stream priorities with event after classifying them as low, medium, high. -*/ - -#include "streamCommon.hh" #include #include #include +#include "streamCommon.hh" // NOLINT -#define MEMCPYSIZE 64*1024*1024 -#define MEMCPYSIZE2 1024*1024 +#define MEMCPYSIZE1 (64*1024*1024) +#define MEMCPYSIZE2 (1024*1024) #define NUMITERS 2 #define GRIDSIZE 1024 #define BLOCKSIZE 256 @@ -87,7 +69,7 @@ void funcTestsForAllPriorityLevelsWrtNullStrm(unsigned int flags, int priority; int priority_low{}; int priority_high{}; - size_t size = MEMCPYSIZE2*sizeof(int); + size_t size = MEMCPYSIZE2 * sizeof(int); // Test is to get the Stream Priority Range HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); @@ -194,7 +176,7 @@ void funcTestsForAllPriorityLevelsWrtNullStrm(unsigned int flags, */ void queueTasksInStreams(std::vector *stream, const int arrsize) { - size_t size = MEMCPYSIZE2*sizeof(int); + size_t size = MEMCPYSIZE2 * sizeof(int); // Allocate memory int **A_d = reinterpret_cast(malloc(arrsize*sizeof(int *))); int **C_d = reinterpret_cast(malloc(arrsize*sizeof(int *))); @@ -211,8 +193,8 @@ void queueTasksInStreams(std::vector *stream, HIPASSERT(A_h[idx] != nullptr); C_h[idx] = reinterpret_cast(malloc(size)); HIPASSERT(C_h[idx] != nullptr); - HIPCHECK(hipMalloc(&A_d[idx], size)); - HIPCHECK(hipMalloc(&C_d[idx], size)); + HIP_CHECK(hipMalloc(&A_d[idx], size)); + HIP_CHECK(hipMalloc(&C_d[idx], size)); } // Initialize host memory constexpr int initVal = 2; @@ -223,20 +205,20 @@ void queueTasksInStreams(std::vector *stream, } // Launch task on each stream for (int idx = 0; idx < arrsize; idx++) { - HIPCHECK(hipMemcpyAsync(A_d[idx], A_h[idx], size, + HIP_CHECK(hipMemcpyAsync(A_d[idx], A_h[idx], size, hipMemcpyHostToDevice, (*stream)[idx])); hipLaunchKernelGGL((HipTest::vector_square), dim3(GRIDSIZE), dim3(BLOCKSIZE), 0, (*stream)[idx], A_d[idx], C_d[idx], MEMCPYSIZE2); HIP_CHECK(hipGetLastError()); - HIPCHECK(hipMemcpyAsync(C_h[idx], C_d[idx], size, + HIP_CHECK(hipMemcpyAsync(C_h[idx], C_d[idx], size, hipMemcpyDeviceToHost, (*stream)[idx])); } bool isPassed = true; // Validate the output of each queue for (int idx = 0; idx < arrsize; idx++) { - HIPCHECK(hipStreamSynchronize((*stream)[idx])); + HIP_CHECK(hipStreamSynchronize((*stream)[idx])); for (int idy = 0; idy < MEMCPYSIZE2; idy++) { if (C_h[idx][idy] != A_h[idx][idy] * A_h[idx][idy]) { UNSCOPED_INFO("Data mismatch at idx:" << idx << " idy:" << idy); @@ -248,8 +230,8 @@ void queueTasksInStreams(std::vector *stream, } // Deallocate memory for (int idx = 0; idx < arrsize; idx++) { - HIPCHECK(hipFree(reinterpret_cast(C_d[idx]))); - HIPCHECK(hipFree(reinterpret_cast(A_d[idx]))); + HIP_CHECK(hipFree(reinterpret_cast(C_d[idx]))); + HIP_CHECK(hipFree(reinterpret_cast(A_d[idx]))); free(C_h[idx]); free(A_h[idx]); } @@ -322,7 +304,7 @@ bool runFuncTestsForAllPriorityLevelsMultThread(unsigned int flags) { template bool validateStreamPrioritiesWithEvents() { - size_t size = NUMITERS*MEMCPYSIZE; + size_t size = NUMITERS * MEMCPYSIZE1; // get the range of priorities available #define OP(x) \ @@ -421,13 +403,13 @@ bool validateStreamPrioritiesWithEvents() { #undef OP // launch kernels repeatedly on each of the prioritiy streams - for (int i = 0; i < static_cast(size); i += MEMCPYSIZE) { + for (int i = 0; i < static_cast(size); i += MEMCPYSIZE1) { int j = i / sizeof(T); #define OP(x) \ if (enable_priority_##x) { \ hipLaunchKernelGGL((memcpy_kernel), dim3(GRIDSIZE), \ dim3(BLOCKSIZE), 0, stream_##x, dst_d_##x + j, src_d_##x + j, \ - (MEMCPYSIZE / sizeof(T))); \ + (MEMCPYSIZE1 / sizeof(T))); \ HIP_CHECK(hipGetLastError()); \ } OP(low) @@ -522,19 +504,330 @@ bool validateStreamPrioritiesWithEvents() { return true; } +#define LOW_PRIORITY_STREAMCOUNT 2 +#define HIGH_PRIORITY_STREAMCOUNT 2 +#define NORMAL_PRIORITY_STREAMCOUNT 2 +template +void TestForMultipleStreamWithPriority(void) { + size_t size = NUMITERS * MEMCPYSIZE1; + // get the range of priorities available + #define OP(x) \ + int priority_##x; \ + bool enable_priority_##x = false; + OP(low) + OP(normal) + OP(high) + #undef OP + HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); + INFO("HIP stream priority range - low: " << priority_low << ",high: " + << priority_high << ",normal: " + << (priority_low + priority_high)/2); + + // Check if priorities are indeed supported + // Enable/disable priorities based on number of available priority levels + enable_priority_low = true; + enable_priority_high = true; + if ((priority_low - priority_high) > 1) { + enable_priority_normal = true; + } + if (enable_priority_normal) { + priority_normal = ((priority_low + priority_high) / 2); + } + // create streams with low priority + hipStream_t stream_low[LOW_PRIORITY_STREAMCOUNT]; + for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_low) { + HIP_CHECK(hipStreamCreateWithPriority(&stream_low[i], + hipStreamDefault, priority_low)); + } + } + // create streams with normal priority + hipStream_t stream_normal[NORMAL_PRIORITY_STREAMCOUNT]; + for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_normal) { + HIP_CHECK(hipStreamCreateWithPriority(&stream_normal[i], + hipStreamDefault, priority_normal)); + } + } + // create streams with high priority + hipStream_t stream_high[HIGH_PRIORITY_STREAMCOUNT]; + for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_high) { + HIP_CHECK(hipStreamCreateWithPriority(&stream_high[i], + hipStreamDefault, priority_high)); + } + } + // allocate and initialise host source and destination buffers for + // low streams + T* src_h_low[LOW_PRIORITY_STREAMCOUNT]; + T* dst_h_low[LOW_PRIORITY_STREAMCOUNT]; + for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_low) { + src_h_low[i] = reinterpret_cast(malloc(size)); + REQUIRE(src_h_low[i] != nullptr); + mem_init(src_h_low[i], (size / sizeof(T))); + dst_h_low[i] = reinterpret_cast(malloc(size)); + REQUIRE(dst_h_low[i] != nullptr); + memset(dst_h_low[i], 0, size); + } + } + // allocate and initialise host source and destination buffers for + // normal streams + T* src_h_normal[NORMAL_PRIORITY_STREAMCOUNT]; + T* dst_h_normal[NORMAL_PRIORITY_STREAMCOUNT]; + for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_normal) { + src_h_normal[i] = reinterpret_cast(malloc(size)); + REQUIRE(src_h_normal[i] != nullptr); + mem_init(src_h_normal[i], (size / sizeof(T))); + dst_h_normal[i] = reinterpret_cast(malloc(size)); + REQUIRE(dst_h_normal[i] != nullptr); + memset(dst_h_normal[i], 0, size); + } + } + // allocate and initialise host source and destination buffers for + // high streams + T* src_h_high[HIGH_PRIORITY_STREAMCOUNT]; + T* dst_h_high[HIGH_PRIORITY_STREAMCOUNT]; + for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_high) { + src_h_high[i] = reinterpret_cast(malloc(size)); + REQUIRE(src_h_high[i] != nullptr); + mem_init(src_h_high[i], (size / sizeof(T))); + dst_h_high[i] = reinterpret_cast(malloc(size)); + REQUIRE(dst_h_high[i] != nullptr); + memset(dst_h_high[i], 0, size); + } + } + // allocate and initialize device source and destination buffers for + // low streams + T* src_d_low[LOW_PRIORITY_STREAMCOUNT]; + T* dst_d_low[LOW_PRIORITY_STREAMCOUNT]; + for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_low) { + HIP_CHECK(hipMalloc(&src_d_low[i], size)); + HIP_CHECK(hipMemcpy(src_d_low[i], src_h_low[i], size, + hipMemcpyHostToDevice)); + HIP_CHECK(hipMalloc(&dst_d_low[i], size)); + } + } + // allocate and initialize device source and destination buffers for + // normal streams + T* src_d_normal[NORMAL_PRIORITY_STREAMCOUNT]; + T* dst_d_normal[NORMAL_PRIORITY_STREAMCOUNT]; + for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_normal) { + HIP_CHECK(hipMalloc(&src_d_normal[i], size)); + HIP_CHECK(hipMemcpy(src_d_normal[i], src_h_normal[i], size, + hipMemcpyHostToDevice)); + HIP_CHECK(hipMalloc(&dst_d_normal[i], size)); + } + } + // allocate and initialize device source and destination buffers for + // high streams + T* src_d_high[HIGH_PRIORITY_STREAMCOUNT]; + T* dst_d_high[HIGH_PRIORITY_STREAMCOUNT]; + for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_high) { + HIP_CHECK(hipMalloc(&src_d_high[i], size)); + HIP_CHECK(hipMemcpy(src_d_high[i], src_h_high[i], size, + hipMemcpyHostToDevice)); + HIP_CHECK(hipMalloc(&dst_d_high[i], size)); + } + } + + hipEvent_t event_start_low[LOW_PRIORITY_STREAMCOUNT]; + hipEvent_t event_end_low[LOW_PRIORITY_STREAMCOUNT]; + // create events for measuring time spent in kernel execution for low + // priority streams + for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_low) { + HIP_CHECK(hipEventCreate(&event_start_low[i])); + HIP_CHECK(hipEventCreate(&event_end_low[i])); + } + } + + hipEvent_t event_start_normal[NORMAL_PRIORITY_STREAMCOUNT]; + hipEvent_t event_end_normal[NORMAL_PRIORITY_STREAMCOUNT]; + // create events for measuring time spent in kernel execution for + // normal priority streams + for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_normal) { + HIP_CHECK(hipEventCreate(&event_start_normal[i])); + HIP_CHECK(hipEventCreate(&event_end_normal[i])); + } + } + + hipEvent_t event_start_high[HIGH_PRIORITY_STREAMCOUNT]; + hipEvent_t event_end_high[HIGH_PRIORITY_STREAMCOUNT]; + // create events for measuring time spent in kernel execution for + // high priority streams + for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_high) { + HIP_CHECK(hipEventCreate(&event_start_high[i])); + HIP_CHECK(hipEventCreate(&event_end_high[i])); + } + } + // record start events for each of the low priority streams + for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_low) { + HIP_CHECK(hipEventRecord(event_start_low[i], stream_low[i])); + } + } + // record start events for each of the normal priority streams + for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_normal) { + HIP_CHECK(hipEventRecord(event_start_normal[i], stream_normal[i])); + } + } + // record start events for each of the high priority streams + for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_high) { + HIP_CHECK(hipEventRecord(event_start_high[i], stream_high[i])); + } + } + // launch kernels repeatedly on each of the low prioritiy stream + for (int k = 0; k < LOW_PRIORITY_STREAMCOUNT; ++k) { + for (int i = 0; i < size; i += MEMCPYSIZE1) { + int j = i / sizeof(T); + if (enable_priority_low) { + hipLaunchKernelGGL((memcpy_kernel), dim3(GRIDSIZE), dim3(BLOCKSIZE), + 0, stream_low[k], dst_d_low[k] + j, src_d_low[k] + j, + (MEMCPYSIZE1 / sizeof(T))); + } + } + } + // launch kernels repeatedly on each of the normal prioritiy stream + for (int k = 0; k < NORMAL_PRIORITY_STREAMCOUNT; ++k) { + for (int i = 0; i < size; i += MEMCPYSIZE1) { + int j = i / sizeof(T); + if (enable_priority_normal) { + hipLaunchKernelGGL((memcpy_kernel), dim3(GRIDSIZE), dim3(BLOCKSIZE), + 0, stream_normal[k], dst_d_normal[k] + j, src_d_normal[k] + j, + (MEMCPYSIZE1 / sizeof(T))); + } + } + } + // launch kernels repeatedly on each of the high prioritiy stream + for (int k = 0; k < HIGH_PRIORITY_STREAMCOUNT; ++k) { + for (int i = 0; i < size; i += MEMCPYSIZE1) { + int j = i / sizeof(T); + if (enable_priority_high) { + hipLaunchKernelGGL((memcpy_kernel), dim3(GRIDSIZE), dim3(BLOCKSIZE), + 0, stream_high[k], dst_d_high[k] + j, src_d_high[k] + j, + (MEMCPYSIZE1 / sizeof(T))); + } + } + } + // record end events for each of the low priority streams + for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_low) { + HIP_CHECK(hipEventRecord(event_end_low[i], stream_low[i])); + } + } + // record end events for each of the normal priority streams + for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_normal) { + HIP_CHECK(hipEventRecord(event_end_normal[i], stream_normal[i])); + } + } + // record end events for each of the high priority streams + for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_high) { + HIP_CHECK(hipEventRecord(event_end_high[i], stream_high[i])); + } + } + // synchronize events for each of the low priority streams + for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_low) { + HIP_CHECK(hipEventSynchronize(event_end_low[i])); + } + } + // synchronize events for each of the normal priority streams + for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_normal) { + HIP_CHECK(hipEventSynchronize(event_end_normal[i])); + } + } + // synchronize events for each of the high priority streams + for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_high) { + HIP_CHECK(hipEventSynchronize(event_end_high[i])); + } + } + // compute time spent for memcpy in each low stream + float time_spent_low[LOW_PRIORITY_STREAMCOUNT]; + for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_low) { + HIP_CHECK(hipEventElapsedTime(&time_spent_low[i], event_start_low[i], + event_end_low[i])); + } + } + // compute time spent for memcpy in each normal stream + float time_spent_normal[NORMAL_PRIORITY_STREAMCOUNT]; + for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_normal) { + HIP_CHECK(hipEventElapsedTime(&time_spent_normal[i], + event_start_normal[i], event_end_normal[i])); + } + } + // compute time spent for memcpy in each high stream + float time_spent_high[HIGH_PRIORITY_STREAMCOUNT]; + for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_high) { + HIP_CHECK(hipEventElapsedTime(&time_spent_high[i], event_start_high[i], + event_end_high[i])); + } + } + // sanity check for low priority streams + for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_low) { + HIP_CHECK(hipMemcpy(dst_h_low[i], dst_d_low[i], size, + hipMemcpyDeviceToHost)); + REQUIRE(memcmp(dst_h_low[i], src_h_low[i], size) == 0); + } + } + // sanity check for normal priority streams + for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_normal) { + HIP_CHECK(hipMemcpy(dst_h_normal[i], dst_d_normal[i], size, + hipMemcpyDeviceToHost)); + REQUIRE(memcmp(dst_h_normal[i], src_h_normal[i], size) == 0); + } + } + // sanity check for high priority streams + for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) { + if (enable_priority_high) { + HIP_CHECK(hipMemcpy(dst_h_high[i], dst_d_high[i], size, + hipMemcpyDeviceToHost)); + REQUIRE(memcmp(dst_h_high[i], src_h_high[i], size) == 0); + } + } +} } // namespace hipStreamCreateWithPriorityTest /** -Tests following scenarios. - 1)Create streams with default flag for all available priority levels and - queue tasks in each of these streams, perform device synchronize and validate behavior. - 2)Create streams with non-blocking flag for all available priority levels and - queue tasks in each of these streams, perform stream synchronize and validate behavior. - 3)Create streams with default flag for all available priority levels and - queue tasks in each of these streams, perform stream synchronize and validate behavior. - 4)Create streams with non-blocking flag for all available priority levels and - queue tasks in each of these streams, perform device synchronize and validate behavior. -*/ + * Test Description + * ------------------------ + * - Create streams with default flag for all available priority levels and + * queue tasks in each of these streams, perform device synchronize and validate + * behavior. + * - Create streams with non-blocking flag for all available priority levels + * and queue tasks in each of these streams, perform stream synchronize and + * validate behavior. + * - Create streams with default flag for all available priority levels and + * queue tasks in each of these streams, perform stream synchronize and validate + * behavior. + * - Create streams with non-blocking flag for all available priority levels + * and queue tasks in each of these streams, perform device synchronize and validate + * behavior. + * Test source + * ------------------------ + * - catch\unit\stream\hipStreamCreateWithPriority.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamCreateWithPriority_FunctionalForAllPriorities") { SECTION("Default flag and device synchronize") { hipStreamCreateWithPriorityTest:: @@ -558,8 +851,16 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_FunctionalForAllPriorities") { } /** - * Create a stream for each priority level with default flag, Launch memcpy and kernel - * tasks on these streams from multiple threads. Validate all the results. + * Test Description + * ------------------------ + * - Create a stream for each priority level with default flag, Launch + * memcpy and kernel tasks on these streams from multiple threads. Validate + * all the results. + * ------------------------ + * - catch\unit\stream\hipStreamCreateWithPriority.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadDefaultflag") { bool TestPassed = true; @@ -569,8 +870,16 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadDefaultflag") { } /** - * Create a stream for each priority level with non-blocking flag, Launch memcpy and - * kernel tasks on these streams from multiple threads. Validate all the results. + * Test Description + * ------------------------ + * - Create a stream for each priority level with non-blocking flag, Launch + * memcpy and kernel tasks on these streams from multiple threads. Validate all + * the results. + * ------------------------ + * - catch\unit\stream\hipStreamCreateWithPriority.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadNonblockingflag") { bool TestPassed = true; @@ -579,13 +888,17 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadNonblockingflag") { REQUIRE(TestPassed); } - /** - * Scenario1: Validates functionality of hipStreamCreateWithPriority when - stream = nullptr. - * Scenario2: Validates functionality of hipStreamCreateWithPriority when - flag = 0xffffffff. -*/ + * Test Description + * ------------------------ + * - Validates functionality of hipStreamCreateWithPriority when stream = nullptr + * - Validates functionality of hipStreamCreateWithPriority when flag = 0xffffffff + * ------------------------ + * - catch\unit\stream\hipStreamCreateWithPriority.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamCreateWithPriority_NegTst") { hipStream_t stream{nullptr}; int priority_low{0}; @@ -601,15 +914,26 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_NegTst") { SECTION("stream = nullptr test") { REQUIRE(hipErrorInvalidValue == - hipStreamCreateWithPriority(nullptr, hipStreamDefault, priority_low)); + hipStreamCreateWithPriority(nullptr, hipStreamDefault, priority_low)); } SECTION("flag value invalid test") { - REQUIRE(hipErrorInvalidValue == hipStreamCreateWithPriority(&stream, 0xffffffff, priority_low)); + REQUIRE(hipErrorInvalidValue == + hipStreamCreateWithPriority(&stream, 0xffffffff, priority_low)); } } -TEST_CASE("Unit_hipStreamCreateWithPriority") { +/** + * Test Description + * ------------------------ + * - Set and Get Priority Value + * ------------------------ + * - catch\unit\stream\hipStreamCreateWithPriority.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipStreamCreateWithPriority_CheckPriorityVal") { int id = GENERATE(range(0, HipTest::getDeviceCount())); HIP_CHECK(hipSetDevice(id)); @@ -619,45 +943,73 @@ TEST_CASE("Unit_hipStreamCreateWithPriority") { hipStream_t stream{nullptr}; SECTION("Setting high priority") { - HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_high)); + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, + priority_high)); REQUIRE(stream != nullptr); REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_high)); } SECTION("Setting low priority") { - HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_low)); + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, + priority_low)); REQUIRE(stream != nullptr); REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_low)); } SECTION("Setting lowest possible priority") { - HIP_CHECK( - hipStreamCreateWithPriority(&stream, hipStreamDefault, std::numeric_limits::max())); + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, + std::numeric_limits::max())); REQUIRE(stream != nullptr); REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_low)); } SECTION("Setting highest possible priority") { - HIP_CHECK( - hipStreamCreateWithPriority(&stream, hipStreamDefault, std::numeric_limits::min())); + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, + std::numeric_limits::min())); REQUIRE(stream != nullptr); REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_high)); } SECTION("Setting flags to hipStreamNonBlocking") { - HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_high)); + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, + priority_high)); REQUIRE(stream != nullptr); - REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_high, hipStreamNonBlocking)); + REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_high, + hipStreamNonBlocking)); } HIP_CHECK(hipStreamDestroy(stream)); } /** - * Validate stream priorities with event after classifying them as low, medium and high. + * Test Description + * ------------------------ + * - Validate stream priorities with event after classifying them as low, + * medium and high. + * ------------------------ + * - catch\unit\stream\hipStreamCreateWithPriority.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamCreateWithPriority_ValidateWithEvents") { bool TestPassed = true; - TestPassed = hipStreamCreateWithPriorityTest::validateStreamPrioritiesWithEvents(); + TestPassed = + hipStreamCreateWithPriorityTest::validateStreamPrioritiesWithEvents(); REQUIRE(TestPassed); } + +/** + * Test Description + * ------------------------ + * - Test that create Multiple streams with priority low, normal and high + * then these streams used to launch the kernel in random seq high,normal,low. + * ------------------------ + * - catch\unit\stream\hipStreamCreateWithPriority.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipStreamCreateWithPriority_TestMultipleStreamWithPriority") { + hipStreamCreateWithPriorityTest::TestForMultipleStreamWithPriority(); +}