diff --git a/catch/include/hip_test_defgroups.hh b/catch/include/hip_test_defgroups.hh index 680dfa8a0..ac935399b 100644 --- a/catch/include/hip_test_defgroups.hh +++ b/catch/include/hip_test_defgroups.hh @@ -131,6 +131,26 @@ THE SOFTWARE. * @} */ +/** + * @defgroup StreamOTest Ordered Memory Allocator + * @{ + * This section describes the tests for Stream Ordered Memory Allocator functions of HIP runtime + * API. + */ + +/** + * @defgroup StreamTest Stream Management + * @{ + * This section describes tests for the stream management functions of HIP runtime API. + * @} + */ + +/** + * @defgroup StreamMTest Stream Memory Operations + * @{ + * This section describes tests for the Stream Memory Wait and Write functions of HIP runtime API. + */ + /** * @defgroup ShflTest warp shuffle function Management * @{ diff --git a/catch/unit/memory/hipMemPoolApi.cc b/catch/unit/memory/hipMemPoolApi.cc index 927899f1b..a39d99083 100644 --- a/catch/unit/memory/hipMemPoolApi.cc +++ b/catch/unit/memory/hipMemPoolApi.cc @@ -17,11 +17,6 @@ THE SOFTWARE. */ -/* Test Case Description: - 1) This testcase verifies the basic scenario - supported on - all devices -*/ - #include #include #include @@ -31,6 +26,12 @@ #include #include +/** + * @addtogroup hipMallocAsync hipMallocAsync + * @{ + * @ingroup StreamOTest + */ + constexpr hipMemPoolProps kPoolProps = { hipMemAllocationTypePinned, hipMemHandleTypeNone, @@ -42,10 +43,25 @@ constexpr hipMemPoolProps kPoolProps = { {0} }; -/* - This testcase verifies HIP Mem Pool API basic scenario - supported on all devices +/** + * @addtogroup hipMallocAsync hipMallocAsync + * @{ + * @ingroup StreamOTest */ + /** + * Test Description + * ------------------------ + * - Allocates memory for the array. + * - Checks basic functionalities. + * Test source + * ------------------------ + * - unit/memory/hipMemPoolApi.cc + * Test requirements + * ------------------------ + * - Runtime supports Memory Pools + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipMemPoolApi_Basic") { int mem_pool_support = 0; HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); @@ -108,6 +124,100 @@ TEST_CASE("Unit_hipMemPoolApi_Basic") { HIP_CHECK(hipStreamDestroy(stream)); } +/** + * Test Description + * ------------------------ + * - Checks that the freed memory is used for allocation again. + * - Launches kernel to create a realistic test case. + * Test source + * ------------------------ + * - unit/memory/hipMemPoolApi.cc + * Test requirements + * ------------------------ + * - Runtime supports Memory Pools + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipMemPoolApi_Default") { + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); + return; + } + + hipMemPool_t mem_pool; + HIP_CHECK(hipDeviceGetDefaultMemPool(&mem_pool, 0)); + + float *A, *B, *C; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + size_t numElements = 8 * 1024 * 1024; + HIP_CHECK(hipMallocAsync(reinterpret_cast(&A), numElements * sizeof(float), stream)); + + numElements = 1024; + HIP_CHECK(hipMallocAsync(reinterpret_cast(&C), numElements * sizeof(float), stream)); + + int blocks = 2; + int clkRate; + + if (IsGfx11()) { + HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0)); + kernel500ms_gfx11<<<32, blocks, 0, stream>>>(A, clkRate); + } else { + HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0)); + + kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate); + } + + hipMemPoolAttr attr; + // Not a real free, since kernel isn't done + HIP_CHECK(hipFreeAsync(reinterpret_cast(A), stream)); + + numElements = 8 * 1024 * 1024; + HIP_CHECK(hipMallocAsync(reinterpret_cast(&B), numElements * sizeof(float), stream)); + // Runtime must reuse the pointer + REQUIRE(A == B); + + // Make a sync before the second kernel launch to make sure memory B isn't gone + HIP_CHECK(hipStreamSynchronize(stream)); + + // Second kernel launch with new memory + if (IsGfx11()) { + kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate); + } else { + kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate); + } + + HIP_CHECK(hipFreeAsync(reinterpret_cast(B), stream)); + + HIP_CHECK(hipStreamSynchronize(stream)); + + std::uint64_t value64 = 0; + attr = hipMemPoolAttrReservedMemCurrent; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64)); + // Make sure the current reserved is at least allocation size of buffer C (4KB) + REQUIRE(sizeof(float) * 1024 <= value64); + + attr = hipMemPoolAttrUsedMemHigh; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64)); + // Make sure the high watermark usage works - the both buffers must be reported + REQUIRE(sizeof(float) * (8 * 1024 * 1024 + 1024) == value64); + + attr = hipMemPoolAttrUsedMemCurrent; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64)); + // Make sure the current usage reports just one buffer, because the above free doesn't hold memory + REQUIRE(sizeof(float) * 1024 == value64); + + HIP_CHECK(hipFreeAsync(reinterpret_cast(C), stream)); + HIP_CHECK(hipStreamDestroy(stream)); +} + +/** + * End doxygen group hipMallocAsync. + * @} + */ + constexpr auto wait_ms = 500; __global__ void kernel500ms(float* hostRes, int clkRate) { @@ -138,6 +248,25 @@ __global__ void kernel500ms_gfx11(float* hostRes, int clkRate) { #endif } +/** + * @addtogroup hipFreeAsync hipFreeAsync + * @{ + * @ingroup StreamOTest + */ + +/** + * Test Description + * ------------------------ + * - Checks if memory usage is different before and after synchronization. + * - Synchronization will force free to execute. + * Test source + * ------------------------ + * - unit/memory/hipMemPoolApi.cc + * Test requirements + * ------------------------ + * - Runtime supports Memory Pools + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipMemPoolApi_BasicAlloc") { int mem_pool_support = 0; HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); @@ -226,6 +355,30 @@ TEST_CASE("Unit_hipMemPoolApi_BasicAlloc") { HIP_CHECK(hipStreamDestroy(stream)); } +/** + * End doxygen group hipFreeAsync. + * @} + */ + +/** + * @addtogroup hipMemPoolTrimTo hipMemPoolTrimTo + * @{ + * @ingroup StreamOTest + */ + +/** + * Test Description + * ------------------------ + * - Check if a trim operation is no-op when memory is still in use. + * - Check that trim works correctly once the memory is not in use. + * Test source + * ------------------------ + * - unit/memory/hipMemPoolApi.cc + * Test requirements + * ------------------------ + * - Runtime supports Memory Pools + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipMemPoolApi_BasicTrim") { int mem_pool_support = 0; HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); @@ -314,6 +467,31 @@ TEST_CASE("Unit_hipMemPoolApi_BasicTrim") { HIP_CHECK(hipStreamDestroy(stream)); } +/** + * End doxygen group hipMemPoolTrimTo. + * @} + */ + +/** + * @addtogroup hipMallocFromPoolAsync hipMallocFromPoolAsync + * @{ + * @ingroup StreamOTest + */ + +/** + * Test Description + * ------------------------ + * - Checks that memory from pool is reused when freed. + * - Allocate the same array after the memory is freed. + * - Verify that the old and new pointers are the same. + * Test source + * ------------------------ + * - unit/memory/hipMemPoolApi.cc + * Test requirements + * ------------------------ + * - Runtime supports Memory Pools + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipMemPoolApi_BasicReuse") { int mem_pool_support = 0; HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); @@ -390,6 +568,21 @@ TEST_CASE("Unit_hipMemPoolApi_BasicReuse") { HIP_CHECK(hipStreamDestroy(stream)); } +/** + * Test Description + * ------------------------ + * - Verifies that an oportunistic flag behaves correctly with allocations. + * -# When oportunistic is disallowed and no reuse + * -# When oportunistic is allowed and reuse + * -# When oportunistic is allowed and no reuse + * Test source + * ------------------------ + * - unit/memory/hipMemPoolApi.cc + * Test requirements + * ------------------------ + * - Runtime supports Memory Pools + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") { int mem_pool_support = 0; HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); @@ -553,78 +746,7 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") { HIP_CHECK(hipStreamDestroy(stream2)); } -TEST_CASE("Unit_hipMemPoolApi_Default") { - int mem_pool_support = 0; - HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); - if (!mem_pool_support) { - SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); - return; - } - - hipMemPool_t mem_pool; - HIP_CHECK(hipDeviceGetDefaultMemPool(&mem_pool, 0)); - - float *A, *B, *C; - hipStream_t stream; - HIP_CHECK(hipStreamCreate(&stream)); - - size_t numElements = 8 * 1024 * 1024; - HIP_CHECK(hipMallocAsync(reinterpret_cast(&A), numElements * sizeof(float), stream)); - - numElements = 1024; - HIP_CHECK(hipMallocAsync(reinterpret_cast(&C), numElements * sizeof(float), stream)); - - int blocks = 2; - int clkRate; - - if (IsGfx11()) { - HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0)); - kernel500ms_gfx11<<<32, blocks, 0, stream>>>(A, clkRate); - } else { - HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0)); - - kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate); - } - - hipMemPoolAttr attr; - // Not a real free, since kernel isn't done - HIP_CHECK(hipFreeAsync(reinterpret_cast(A), stream)); - - numElements = 8 * 1024 * 1024; - HIP_CHECK(hipMallocAsync(reinterpret_cast(&B), numElements * sizeof(float), stream)); - // Runtime must reuse the pointer - REQUIRE(A == B); - - // Make a sync before the second kernel launch to make sure memory B isn't gone - HIP_CHECK(hipStreamSynchronize(stream)); - - // Second kernel launch with new memory - if (IsGfx11()) { - kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate); - } else { - kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate); - } - - HIP_CHECK(hipFreeAsync(reinterpret_cast(B), stream)); - - HIP_CHECK(hipStreamSynchronize(stream)); - - std::uint64_t value64 = 0; - attr = hipMemPoolAttrReservedMemCurrent; - HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64)); - // Make sure the current reserved is at least allocation size of buffer C (4KB) - REQUIRE(sizeof(float) * 1024 <= value64); - - attr = hipMemPoolAttrUsedMemHigh; - HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64)); - // Make sure the high watermark usage works - the both buffers must be reported - REQUIRE(sizeof(float) * (8 * 1024 * 1024 + 1024) == value64); - - attr = hipMemPoolAttrUsedMemCurrent; - HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64)); - // Make sure the current usage reports just one buffer, because the above free doesn't hold memory - REQUIRE(sizeof(float) * 1024 == value64); - - HIP_CHECK(hipFreeAsync(reinterpret_cast(C), stream)); - HIP_CHECK(hipStreamDestroy(stream)); -} +/** + * End doxygen group hipMallocFromPoolAsync. + * @} + */ diff --git a/catch/unit/stream/hipAPIStreamDisable.cc b/catch/unit/stream/hipAPIStreamDisable.cc index e1163471e..85195ad6f 100644 --- a/catch/unit/stream/hipAPIStreamDisable.cc +++ b/catch/unit/stream/hipAPIStreamDisable.cc @@ -19,6 +19,12 @@ THE SOFTWARE. #include #include "hip/math_functions.h" +/** + * @addtogroup hipStreamCreate hipStreamCreate + * @{ + * @ingroup StreamTest + */ + #define NUM_STREAMS 8 namespace hipAPIStreamDisableTest { @@ -41,7 +47,16 @@ __global__ void nKernel(float* y) { } // namespace hipAPIStreamDisableTest /** - * Validate basic multistream functionalities + * Test Description + * ------------------------ + * - Validate basic multistream functionalities. + * - Launch the same kernel for multiple streams. + * Test source + * ------------------------ + * - unit/stream/hipAPIStreamDisable.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamCreate_MultistreamBasicFunctionalities") { hipStream_t streams[NUM_STREAMS]; diff --git a/catch/unit/stream/hipDeviceGetStreamPriorityRange.cc b/catch/unit/stream/hipDeviceGetStreamPriorityRange.cc index eb813f908..eb0a8ac62 100644 --- a/catch/unit/stream/hipDeviceGetStreamPriorityRange.cc +++ b/catch/unit/stream/hipDeviceGetStreamPriorityRange.cc @@ -19,13 +19,28 @@ 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. */ -/* -Testcase Scenarios : -Unit_hipDeviceGetStreamPriorityRange_Default - Check if device stream piority range is valid -*/ #include +/** + * @addtogroup hipDeviceGetStreamPriorityRange hipDeviceGetStreamPriorityRange + * @{ + * @ingroup StreamTest + * `hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPriority)` - + * Returns numerical values that correspond to the least and greatest stream priority. + */ + +/** + * Test Description + * ------------------------ + * - Checks that the low and high priority limits are valid. + * Test source + * ------------------------ + * - unit/stream/hipDeviceGetStreamPriorityRange.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipDeviceGetStreamPriorityRange_Default") { int priority_low = 0; int priority_high = 0; diff --git a/catch/unit/stream/hipMultiStream.cc b/catch/unit/stream/hipMultiStream.cc index cffc9af88..838152c2c 100644 --- a/catch/unit/stream/hipMultiStream.cc +++ b/catch/unit/stream/hipMultiStream.cc @@ -19,6 +19,13 @@ THE SOFTWARE. #include #include #include + +/** + * @addtogroup hipStreamCreate hipStreamCreate + * @{ + * @ingroup StreamTest + */ + constexpr int NN = 1 << 21; __global__ void kernel_do_nothing(__attribute__((unused))int a) { // empty kernel @@ -36,6 +43,18 @@ __global__ void nKernel(float* y) { size_t tid{threadIdx.x}; y[tid] = y[tid] + 1.0f; } + +/** + * Test Description + * ------------------------ + * - Validate creation of multiple streams on the same device. + * Test source + * ------------------------ + * - unit/stream/hipMultiStream.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipMultiStream_sameDevice") { constexpr int num_streams{8}; hipStream_t streams[num_streams]; @@ -60,6 +79,17 @@ TEST_CASE("Unit_hipMultiStream_sameDevice") { REQUIRE(x == Approx(y)); } +/** + * Test Description + * ------------------------ + * - Validate creation of multiple streams on multiple devices. + * Test source + * ------------------------ + * - unit/stream/hipMultiStream.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipMultiStream_multimeDevice") { constexpr int nLoops = 50000; constexpr int nStreams = 2; diff --git a/catch/unit/stream/hipStreamACb_MultiThread.cc b/catch/unit/stream/hipStreamACb_MultiThread.cc index 9d1a780de..129cdb17f 100644 --- a/catch/unit/stream/hipStreamACb_MultiThread.cc +++ b/catch/unit/stream/hipStreamACb_MultiThread.cc @@ -17,16 +17,16 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenario : -Validate behaviour of HIP when multiple hipStreaAddCallback() are called over -multiple Threads. -*/ - #include #include #include +/** + * @addtogroup hipStreamAddCallback hipStreamAddCallback + * @{ + * @ingroup StreamTest + */ + static constexpr size_t N = 4096; static constexpr int numThreads = 1000; static std::atomic Cb_count{0}, Data_mismatch{0}; @@ -89,8 +89,15 @@ void Thread2_func() { } /** - Test multiple hipStreamAddCallback() called over - multiple Threads. + * Test Description + * ------------------------ + * - Add callbacks on the streams from multiple threads. + * Test source + * ------------------------ + * - unit/stream/hipStreamACb_MultiThread.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamAddCallback_MultipleThreads") { float *A_d, *C_d; diff --git a/catch/unit/stream/hipStreamACb_StrmSyncTiming.cc b/catch/unit/stream/hipStreamACb_StrmSyncTiming.cc index a60760496..feee595e5 100644 --- a/catch/unit/stream/hipStreamACb_StrmSyncTiming.cc +++ b/catch/unit/stream/hipStreamACb_StrmSyncTiming.cc @@ -17,18 +17,18 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenario : -Validate behaviour of HIP when multiple hipStreaAddCallback() are called over -multiple Threads. -*/ - #include #include #include #include #include +/** + * @addtogroup hipStreamAddCallback hipStreamAddCallback + * @{ + * @ingroup StreamTest + */ + #ifdef __HIP_PLATFORM_AMD__ #define HIPRT_CB #endif @@ -77,8 +77,15 @@ static void HIPRT_CB Callback1(hipStream_t stream, hipError_t status, void* user } /** - Test multiple hipStreamAddCallback() called over - multiple Threads. + * Test Description + * ------------------------ + * - Add callbacks on the streams utilizing synchronization. + * Test source + * ------------------------ + * - unit/stream/hipStreamACb_MultiThread.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamAddCallback_StrmSyncTiming") { float *A_d, *C_d; diff --git a/catch/unit/stream/hipStreamAddCallback.cc b/catch/unit/stream/hipStreamAddCallback.cc index 567cfa168..a7d8d8ae4 100644 --- a/catch/unit/stream/hipStreamAddCallback.cc +++ b/catch/unit/stream/hipStreamAddCallback.cc @@ -17,18 +17,23 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : - 1) Validates parameter list of hipStreamAddCallback. - 2) Validates hipStreamAddCallback functionality with default stream. - 3) Validates hipStreamAddCallback functionality with defined stream. -*/ - #include #include #include #include +/** + * @addtogroup hipStreamAddCallback hipStreamAddCallback + * @{ + * @ingroup StreamTest + * `hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, + * void* userData, unsigned int flags)` - + * Adds a callback to be called on the host after all currently enqueued + * items in the stream have completed. For each + * hipStreamAddCallback call, a callback will be executed exactly once. + * The callback will block later work in the stream until it is finished. + */ + #define UNUSED(expr) do { (void)(expr); } while (0) #ifdef __HIP_PLATFORM_AMD__ @@ -59,9 +64,8 @@ void HIPRT_CB Callback(hipStream_t stream, hipError_t status, } gcbDone = true; } -/** - * Validates functionality of hipStreamAddCallback with default/created stream. - */ + +// Validates functionality of hipStreamAddCallback with default/created stream. bool testStreamCallbackFunctionality(bool isDefault) { float *A_d, *C_d; size_t Nbytes = NSize * sizeof(float); @@ -115,14 +119,14 @@ bool testStreamCallbackFunctionality(bool isDefault) { free(A_h); return gPassed; } -/** - * Scenario1: Validates if callback = nullptr returns error code for created stream. - * Scenario2: Validates if callback = nullptr returns error code for default stream. - * Scenario3: Validates if flag != 0 returns error code for created stream. - * Scenario4: Validates if flag != 0 returns error code for default stream. - * Scenario5: Validates if userData pointer is passed properly to callback. - * Scenario6: Validates if stream value is passed properly to callback. - */ +/* +Scenario1: Validates if callback = nullptr returns error code for created stream. +Scenario2: Validates if callback = nullptr returns error code for default stream. +Scenario3: Validates if flag != 0 returns error code for created stream. +Scenario4: Validates if flag != 0 returns error code for default stream. +Scenario5: Validates if userData pointer is passed properly to callback. +Scenario6: Validates if stream value is passed properly to callback. +*/ void Callback_ChkUsrdataPtr(hipStream_t stream, hipError_t status, void* userData) { REQUIRE(stream == gstream); @@ -157,9 +161,20 @@ using hipStreaAddCallbackTest::Callback; using hipStreaAddCallbackTest::Callback_ChkUsrdataPtr; using hipStreaAddCallbackTest::Callback_ChkStreamValue; - -/* - * Validates parameter list of hipStreamAddCallback. +/** + * Test Description + * ------------------------ + * - Test that all parameters behave correctly: + * -# When userData pointer is valid + * - Expected output: return `hipSuccess` + * -# When stream is created + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/stream/hipStreamAddCallback.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamAddCallback_ParamTst_Positive") { hipStream_t mystream; @@ -193,8 +208,24 @@ TEST_CASE("Unit_hipStreamAddCallback_ParamTst_Positive") { HIP_CHECK(hipStreamDestroy(mystream)); } -/* - * Negative tests for validation of hipStreamAddCallback parameter list. +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When callback is `nullptr` for non-default stream + * - Expected output: do not return `hipSuccess` + * -# When callback is `nullptr` for default stream + * - Expected output: do not return `hipSuccess` + * -# When flag is non-zero for non-default stream + * - Expected output: do not return `hipSuccess` + * -# When flag is non-zero for default stream + * - Expected output: do not return `hipSuccess` + * Test source + * ------------------------ + * - unit/stream/hipStreamAddCallback.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamAddCallback_ParamTst_Negative") { hipStream_t mystream; @@ -223,8 +254,16 @@ TEST_CASE("Unit_hipStreamAddCallback_ParamTst_Negative") { HIP_CHECK(hipStreamDestroy(mystream)); } -/* - * Validates hipStreamAddCallback functionality with default stream. +/** + * Test Description + * ------------------------ + * - Validates adding callback functionality with default stream. + * Test source + * ------------------------ + * - unit/stream/hipStreamAddCallback.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamAddCallback_WithDefaultStream") { bool TestPassed = true; @@ -232,12 +271,19 @@ TEST_CASE("Unit_hipStreamAddCallback_WithDefaultStream") { REQUIRE(TestPassed); } -/* - * Validates hipStreamAddCallback functionality with defined stream. +/** + * Test Description + * ------------------------ + * - Validates adding callback functionality with defined stream. + * Test source + * ------------------------ + * - unit/stream/hipStreamAddCallback.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamAddCallback_WithCreatedStream") { bool TestPassed = true; TestPassed = testStreamCallbackFunctionality(false); REQUIRE(TestPassed); } - diff --git a/catch/unit/stream/hipStreamCreate.cc b/catch/unit/stream/hipStreamCreate.cc index 21fff252b..8512e62e2 100644 --- a/catch/unit/stream/hipStreamCreate.cc +++ b/catch/unit/stream/hipStreamCreate.cc @@ -19,6 +19,25 @@ THE SOFTWARE. #include "streamCommon.hh" +/** + * @addtogroup hipStreamCreate hipStreamCreate + * @{ + * @ingroup StreamTest + * `hipStreamCreate(hipStream_t* stream)` - + * Create an asynchronous stream. + */ + +/** + * Test Description + * ------------------------ + * - Create valid stream and check its flags and priority correctness. + * Test source + * ------------------------ + * - unit/stream/hipStreamCreate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamCreate_default") { int id = GENERATE(range(0, HipTest::getDeviceCount())); HIP_CHECK(hipSetDevice(id)); @@ -30,6 +49,19 @@ TEST_CASE("Unit_hipStreamCreate_default") { HIP_CHECK(hipStreamDestroy(stream)); } +/** + * Test Description + * ------------------------ + * - Validate handling of invalid arguments: + * -# When output pointer to the stream is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/stream/hipStreamCreate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamCreate_Negative") { REQUIRE(hipErrorInvalidValue == hipStreamCreate(nullptr)); } diff --git a/catch/unit/stream/hipStreamCreateWithFlags.cc b/catch/unit/stream/hipStreamCreateWithFlags.cc index 1ffb9c152..3cb14a038 100644 --- a/catch/unit/stream/hipStreamCreateWithFlags.cc +++ b/catch/unit/stream/hipStreamCreateWithFlags.cc @@ -21,12 +21,45 @@ THE SOFTWARE. #include #include +/** + * @addtogroup hipStreamCreateWithFlags hipStreamCreateWithFlags + * @{ + * @ingroup StreamTest + * `hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags)` - + * Create an asynchronous stream with flags. + */ + namespace hipStreamCreateWithFlagsTests { +/** + * Test Description + * ------------------------ + * - Verifies handling of invalid arguments: + * -# When output pointer to the stream is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/stream/hipStreamCreateWithFlags.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamCreateWithFlags_Negative_NullStream") { HIP_CHECK_ERROR(hipStreamCreateWithFlags(nullptr, hipStreamDefault), hipErrorInvalidValue); } +/** + * Test Description + * ------------------------ + * - Creates stream with invalid flag. + * - Valid flags are 0x0 and 0x1. + * Test source + * ------------------------ + * - unit/stream/hipStreamCreateWithFlags.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamCreateWithFlags_Negative_InvalidFlag") { hipStream_t stream{}; unsigned int flag = 0xFF; @@ -35,7 +68,18 @@ TEST_CASE("Unit_hipStreamCreateWithFlags_Negative_InvalidFlag") { HIP_CHECK_ERROR(hipStreamCreateWithFlags(&stream, flag), hipErrorInvalidValue); } -// create a stream and check the properties are correctly set +/** + * Test Description + * ------------------------ + * - Creates streams with valid flags. + * - Checks that they are created as expected. + * Test source + * ------------------------ + * - unit/stream/hipStreamCreateWithFlags.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamCreateWithFlags_Default") { const unsigned int flagUnderTest = GENERATE(hipStreamDefault, hipStreamNonBlocking); hipStream_t stream{}; @@ -53,9 +97,25 @@ TEST_CASE("Unit_hipStreamCreateWithFlags_Default") { HIP_CHECK(hipStreamDestroy(stream)); } -// a stream will default to blocking the null stream, but will not block the null stream when -// created with hipStreamNonBlocking #if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */ +/** + * Test Description + * ------------------------ + * - Test how stream set as default interacts with created streams. + * -# When null stream is set as default and stream is created as default + * - Created stream is blocking the null stream and vice versa + * -# When null stream is set as default and stream is created as non blocking + * - Created stream is not blocking the null stream and vice versa + * -# When stream per thread is set as default and stream is created with any flag + * - Created stream is not blocking the default stream and vice versa + * Test source + * ------------------------ + * - unit/stream/hipStreamCreateWithFlags.cc + * Test requirements + * ------------------------ + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamCreateWithFlags_DefaultStreamInteraction") { const hipStream_t defaultStream = GENERATE(static_cast(nullptr), hipStreamPerThread); const unsigned int flagUnderTest = GENERATE(hipStreamDefault, hipStreamNonBlocking); diff --git a/catch/unit/stream/hipStreamCreateWithPriority.cc b/catch/unit/stream/hipStreamCreateWithPriority.cc index ee3b723c7..81f7af1e8 100644 --- a/catch/unit/stream/hipStreamCreateWithPriority.cc +++ b/catch/unit/stream/hipStreamCreateWithPriority.cc @@ -59,11 +59,11 @@ __global__ void memcpy_kernel(T* dst, T* src, size_t n) { } } -/** - * Scenario: Create a stream for all available priority levels - * and queue tasks in each of these streams and default stream. - * Validate the calculated results. - */ +/* +Scenario: Create a stream for all available priority levels +and queue tasks in each of these streams and default stream. +Validate the calculated results. +*/ void funcTestsForAllPriorityLevelsWrtNullStrm(unsigned int flags, bool deviceSynchronize) { int priority; @@ -170,10 +170,10 @@ void funcTestsForAllPriorityLevelsWrtNullStrm(unsigned int flags, free(C_h); } -/** - * Scenario: Queue tasks in each of these streams and default stream. - * Validate the calculated results. - */ +/* +Scenario: Queue tasks in each of these streams and default stream. +Validate the calculated results. +*/ void queueTasksInStreams(std::vector *stream, const int arrsize) { size_t size = MEMCPYSIZE2 * sizeof(int); @@ -242,13 +242,13 @@ void queueTasksInStreams(std::vector *stream, g_thTestPassed &= static_cast(isPassed); } -/** - * Scenario: - * Common streams used across multiple threads:Create a stream for each - * priority level (flag = hipStreamDefault/hipStreamNonBlocking) - * and 1 default stream. - * Launch memcpy and kernel tasks on these streams from multiple threads - * (use 16 threads). Validate all the results. +/* +Scenario: +Common streams used across multiple threads:Create a stream for each +priority level (flag = hipStreamDefault/hipStreamNonBlocking) +and 1 default stream. +Launch memcpy and kernel tasks on these streams from multiple threads +(use 16 threads). Validate all the results. */ bool runFuncTestsForAllPriorityLevelsMultThread(unsigned int flags) { bool TestPassed = true; @@ -807,24 +807,29 @@ void TestForMultipleStreamWithPriority(void) { /** * 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 following scenarios: + * -# Default flag and device synchronize + * - Create streams with default flag for all available priority levels + * - Queue tasks in each of these streams + * - Perform device synchronize and validate behavior + * -# Stream non-blocking flag and stream synchronize + * - Create streams with non-blocking flag for all available priority levels + * - Queue tasks in each of these streams + * - Perform stream synchronize and validate behavior + * -# Default flag and stream synchronize + * - Create streams with default flag for all available priority levels + * - Queue tasks in each of these streams + * - Perform stream synchronize and validate behavior + * -# Stream non-blocking flag and device synchronize + * - Create streams with non-blocking flag for all available priority levels + * - Queue tasks in each of these streams + * - Perform device synchronize and validate behavior * Test source * ------------------------ - * - catch\unit\stream\hipStreamCreateWithPriority.cc + * - unit/stream/hipStreamCreateWithPriority.cc * Test requirements * ------------------------ - * - HIP_VERSION >= 5.2 + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamCreateWithPriority_FunctionalForAllPriorities") { SECTION("Default flag and device synchronize") { @@ -851,14 +856,15 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_FunctionalForAllPriorities") { /** * 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. + * - 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 source * ------------------------ - * - catch\unit\stream\hipStreamCreateWithPriority.cc + * - unit/stream/hipStreamCreateWithPriority.cc * Test requirements * ------------------------ - * - HIP_VERSION >= 5.2 + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadDefaultflag") { bool TestPassed = true; @@ -870,14 +876,15 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadDefaultflag") { /** * 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. + * - 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 source * ------------------------ - * - catch\unit\stream\hipStreamCreateWithPriority.cc + * - unit/stream/hipStreamCreateWithPriority.cc * Test requirements * ------------------------ - * - HIP_VERSION >= 5.2 + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadNonblockingflag") { bool TestPassed = true; @@ -889,13 +896,17 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadNonblockingflag") { /** * Test Description * ------------------------ - * - Validates functionality of hipStreamCreateWithPriority when stream = nullptr - * - Validates functionality of hipStreamCreateWithPriority when flag = 0xffffffff + * - Validates handling of invalid arguments: + * -# Output pointer to the stream is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# Flag is invalid (0xFFFFFFFF) + * - Expected output: return `hipErrorInvalidValue` + * Test source * ------------------------ - * - catch\unit\stream\hipStreamCreateWithPriority.cc + * - unit/stream/hipStreamCreateWithPriority.cc * Test requirements * ------------------------ - * - HIP_VERSION >= 5.2 + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamCreateWithPriority_NegTst") { hipStream_t stream{nullptr}; @@ -982,13 +993,15 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_CheckPriorityVal") { /** * Test Description * ------------------------ - * - Validate stream priorities with event after classifying them as low, - * medium and high. + * - Launches lots of kernels on three priority streams: low, normal, high. + * - Validates that the higher priority lower the execution time. + * - Execution time is tracked with recording events on streams. + * Test source * ------------------------ - * - catch\unit\stream\hipStreamCreateWithPriority.cc + * - unit/stream/hipStreamCreateWithPriority.cc * Test requirements * ------------------------ - * - HIP_VERSION >= 5.2 + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamCreateWithPriority_ValidateWithEvents") { bool TestPassed = true; diff --git a/catch/unit/stream/hipStreamDestroy.cc b/catch/unit/stream/hipStreamDestroy.cc index a2c0f287e..235632eb0 100644 --- a/catch/unit/stream/hipStreamDestroy.cc +++ b/catch/unit/stream/hipStreamDestroy.cc @@ -20,14 +20,44 @@ THE SOFTWARE. #include #include -namespace hipStreamDestroyTests { +/** + * @addtogroup hipStreamDestroy hipStreamDestroy + * @{ + * @ingroup StreamTest + * `hipStreamDestroy(hipStream_t stream)` - + * Destroys the specified stream. + */ +namespace hipStreamDestroyTests { +/** + * Test Description + * ------------------------ + * - Validates that the stream can be destroyed without errors. + * Test source + * ------------------------ + * - unit/stream/hipStreamDestroy.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamDestroy_Default") { hipStream_t stream{}; HIP_CHECK(hipStreamCreate(&stream)); HIP_CHECK(hipStreamDestroy(stream)); } +/** + * Test Description + * ------------------------ + * - Tries to destroy already destroyed stream: + * - Expected output: return `hipErrorContextIsDestroyed` + * Test source + * ------------------------ + * - unit/stream/hipStreamDestroy.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamDestroy_Negative_DoubleDestroy") { hipStream_t stream{}; HIP_CHECK(hipStreamCreate(&stream)); @@ -35,6 +65,18 @@ TEST_CASE("Unit_hipStreamDestroy_Negative_DoubleDestroy") { HIP_CHECK_ERROR(hipStreamDestroy(stream), hipErrorContextIsDestroyed); } +/** + * Test Description + * ------------------------ + * - Tries to destroy null stream: + * - Expected output: return `hipErrorInvalidResourceHandle` + * Test source + * ------------------------ + * - unit/stream/hipStreamDestroy.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamDestroy_Negative_NullStream") { HIP_CHECK_ERROR(hipStreamDestroy(nullptr), hipErrorInvalidResourceHandle); } @@ -54,6 +96,19 @@ __global__ void setToOne(int* x, size_t size) { } } +/** + * Test Description + * ------------------------ + * - Create a default stream. + * - Run a simple kernel that finishes quickly. + * - Destroy the created stream successfully. + * Test source + * ------------------------ + * - unit/stream/hipStreamDestroy.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamDestroy_WithFinishedWork") { hipStream_t stream{}; HIP_CHECK(hipStreamCreate(&stream)); @@ -72,6 +127,20 @@ TEST_CASE("Unit_hipStreamDestroy_WithFinishedWork") { // hipStreamDestroy should return immediately then clean up the resources when the stream is empty // of work #if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */ +/** + * Test Description + * ------------------------ + * - Create a default stream. + * - Run a kernel with delay that runs for 500ms. + * - Destroy the created stream. + * Test source + * ------------------------ + * - unit/stream/hipStreamDestroy.cc + * Test requirements + * ------------------------ + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamDestroy_WithPendingWork") { hipStream_t stream{}; diff --git a/catch/unit/stream/hipStreamGetCUMask.cc b/catch/unit/stream/hipStreamGetCUMask.cc index 695388993..8852fe25b 100644 --- a/catch/unit/stream/hipStreamGetCUMask.cc +++ b/catch/unit/stream/hipStreamGetCUMask.cc @@ -17,22 +17,31 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : -1) Test to verify hipExtStreamGetCUMask api returning default CU Mask or global CU Mask. -2) Test to verify hipExtStreamGetCUMask api returns custom mask set. -3) Negative tests for hipExtStreamGetCUMask api. -*/ - #include #include #include +/** + * @addtogroup hipExtStreamGetCUMask hipExtStreamGetCUMask + * @{ + * @ingroup StreamTest + * `hipExtStreamGetCUMask(hipStream_t stream, uint32_t cuMaskSize, uint32_t* cuMask)` - + * Get CU mask associated with an asynchronous stream. + */ /** - * Scenario to verify hipExtStreamGetCUMask api returning default CU Mask or global CU Mask. - * Scenario to verify hipExtStreamGetCUMask api returns custom mask set. + * Test Description + * ------------------------ + * - Verifies that stream can be created with different CU mask values: + * - Verify with default CU mask or global CU mask + * - Verify with custom mask set + * Test source + * ------------------------ + * - unit/stream/hipStreamGetCUMask.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipExtStreamGetCUMask_verifyDefaultAndCustomMask") { constexpr unsigned maxCUPerValue = 32; @@ -175,7 +184,19 @@ TEST_CASE("Unit_hipExtStreamGetCUMask_verifyDefaultAndCustomMask") { } /** - * Negative tests for hipExtStreamGetCUMask. + * Test Description + * ------------------------ + * - Verifies handling of invalid arguments: + * -# When pointer to the CU mask is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When CU mask size is 0 + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/stream/hipStreamGetCUMask.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipExtStreamGetCUMask_Negative") { hipError_t ret; diff --git a/catch/unit/stream/hipStreamGetFlags.cc b/catch/unit/stream/hipStreamGetFlags.cc index 25624b41a..b71ce9320 100644 --- a/catch/unit/stream/hipStreamGetFlags.cc +++ b/catch/unit/stream/hipStreamGetFlags.cc @@ -16,19 +16,27 @@ 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. */ -/** -Testcase Scenarios : -1) Test flag value of stream created with hipStreamCreateWithFlags/ - /hipStreamCreate/hipStreamCreateWithPriority. -2) Negative tests for hipStreamGetFlags api. -3) Test flag value when streams created with CUMask. -*/ #include /** - * @brief Check that hipStreamGetFlags returns the same flags that were used to create the stream. - * + * @addtogroup hipStreamGetFlags hipStreamGetFlags + * @{ + * @ingroup StreamTest + * `hipStreamGetFlags(hipStream_t stream, unsigned int* flags)` - + * Return flags associated with this stream. + */ + +/** + * Test Description + * ------------------------ + * - Checks that the returned flags are the same as the ones used to create streams. + * Test source + * ------------------------ + * - unit/stream/hipStreamGetFlags.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamGetFlags_Basic") { unsigned int expectedFlag = GENERATE(hipStreamDefault, hipStreamNonBlocking); @@ -42,8 +50,19 @@ TEST_CASE("Unit_hipStreamGetFlags_Basic") { } /** - * @brief Negative scenarios for hipStreamGetFlags. - * + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When stream is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When output pointer to flags is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/stream/hipStreamGetFlags.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamGetFlags_Negative") { hipStream_t validStream; @@ -68,7 +87,16 @@ TEST_CASE("Unit_hipStreamGetFlags_Negative") { #if HT_AMD /** - * Test flag value when streams created with CUMask. + * Test Description + * ------------------------ + * - Create stream with CU mask. + * - Check that flags are valid. + * Test source + * ------------------------ + * - unit/stream/hipStreamGetFlags.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamGetFlags_StreamsCreatedWithCUMask") { hipStream_t stream; diff --git a/catch/unit/stream/hipStreamGetPriority.cc b/catch/unit/stream/hipStreamGetPriority.cc index b3a85d30a..04c10a4bd 100644 --- a/catch/unit/stream/hipStreamGetPriority.cc +++ b/catch/unit/stream/hipStreamGetPriority.cc @@ -17,20 +17,38 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/* -Testcase Scenarios : -1) Negative tests for hipStreamGetPriority api. -2) Create stream and check default priority of stream is within range. -3) Create stream with high or low priority and check priority is set as expected. -4) Create stream with higher priority or lower priority for the priority range returned, the stream -priority should be clamped to the priority range. -5) Create stream with CUMask and check priority is returned as expected. -*/ - #include /** - * Create stream and check priority. + * @addtogroup hipStreamGetPriority hipStreamGetPriority + * @{ + * @ingroup StreamTest + * `hipStreamGetPriority(hipStream_t stream, int* priority)` - + * Query the priority of a stream. + */ + +/** + * Test Description + * ------------------------ + * - Checks different valid scenarios: + * -# When stream is `nullptr` + * - Expected output: valid priority + * -# When default priority stream is created + * - Expected output: valid priority + * -# When high priority stream is created + * - Expected output: valid priority + * -# When stream priority is higher than avaliable + * - Expected output: clamped priority to the highest valid one + * -# When low priority stream is created + * - Expected output: valid priority + * -# When stream priority is lower than available + * - Expected output: clamped priority to the lowest valid one + * Test source + * ------------------------ + * - unit/stream/hipStreamGetPriority.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamGetPriority_happy") { int priority_low = 0; @@ -80,16 +98,33 @@ TEST_CASE("Unit_hipStreamGetPriority_happy") { } /** - * both stream and priority passed as nullptr. + * Test Description + * ------------------------ + * - Verifies the case when both stream and priority pointers are `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/stream/hipStreamGetPriority.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamGetPriority_nullptr_nullptr") { auto res = hipStreamGetPriority(nullptr,nullptr); REQUIRE(res == hipErrorInvalidValue); } - /** - * valid stream and priority passed as nullptr. + * Test Description + * ------------------------ + * - Verifies the case when priority pointer is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/stream/hipStreamGetPriority.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamGetPriority_stream_nullptr") { hipStream_t stream = nullptr; @@ -101,9 +136,17 @@ TEST_CASE("Unit_hipStreamGetPriority_stream_nullptr") { HIP_CHECK(hipStreamDestroy(stream)); } - /** - * nullptr stream and valid priority + * Test Description + * ------------------------ + * - Verifies the case when stream pointer is `nullptr` + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/stream/hipStreamGetPriority.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamGetPriority_nullptr_priority") { int priority = -1; @@ -111,7 +154,15 @@ TEST_CASE("Unit_hipStreamGetPriority_nullptr_priority") { } /** - * both stream and priority passed as valid. + * Test Description + * ------------------------ + * - Both stream and priority pointers are valid. + * Test source + * ------------------------ + * - unit/stream/hipStreamGetPriority.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamGetPriority_stream_priority") { int priority = -1; @@ -125,7 +176,16 @@ TEST_CASE("Unit_hipStreamGetPriority_stream_priority") { #if HT_AMD /** - * Create stream with CUMask and check priority is returned as expected. + * Test Description + * ------------------------ + * - Create stream with CU mask and check priority is returned as expected. + * Test source + * ------------------------ + * - unit/stream/hipStreamGetPriority.cc + * Test requirements + * ------------------------ + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamGetPriority_StreamsWithCUMask") { hipStream_t stream{}; diff --git a/catch/unit/stream/hipStreamQuery.cc b/catch/unit/stream/hipStreamQuery.cc index 6a1f306c6..e525c1d67 100644 --- a/catch/unit/stream/hipStreamQuery.cc +++ b/catch/unit/stream/hipStreamQuery.cc @@ -21,9 +21,29 @@ THE SOFTWARE. #include "streamCommon.hh" #include /** - * @brief Check that querying a stream with no work returns hipSuccess - * - **/ + * @addtogroup hipStreamQuery hipStreamQuery + * @{ + * @ingroup StreamTest + * `hipStreamQuery(hipStream_t stream)` - + * Return `hipSuccess` if all of the operations in the specified stream have completed, or + * `hipErrorNotReady` if not. + */ + +/** + * Test Description + * ------------------------ + * - Query a stream with no work: + * -# When the stream is `nullptr` + * - Expected output: return `hipSuccess` + * -# When the stream is created + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/stream/hipStreamQuery.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamQuery_WithNoWork") { hipStream_t stream{nullptr}; @@ -39,9 +59,20 @@ TEST_CASE("Unit_hipStreamQuery_WithNoWork") { } /** - * @brief Check that querying a stream with finished work returns hipSuccess - * - **/ + * Test Description + * ------------------------ + * - Query a stream with finished work: + * -# When the stream is `nullptr` + * - Expected output: return `hipSuccess` + * -# When the stream is created + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/stream/hipStreamQuery.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamQuery_WithFinishedWork") { hipStream_t stream{nullptr}; @@ -63,11 +94,20 @@ TEST_CASE("Unit_hipStreamQuery_WithFinishedWork") { } #if !HT_NVIDIA +// Test removed for Nvidia devices because it returns unexpected error + /** - * @brief Check that submitting work to a destroyed stream sets its status as - * hipErrorContextIsDestroyed - * - * Test removed for Nvidia devices because it returns unexpected error + * Test Description + * ------------------------ + * - Query a stream that has been destroyed previously + * - Expected output: return `hipErrorContextIsDestroyed` + * Test source + * ------------------------ + * - unit/stream/hipStreamQuery.cc + * Test requirements + * ------------------------ + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamQuery_WithDestroyedStream") { hipStream_t stream{nullptr}; @@ -77,10 +117,17 @@ TEST_CASE("Unit_hipStreamQuery_WithDestroyedStream") { } /** - * @brief Check that submitting work to an uninitialized stream sets its status as - * hipErrorContextIsDestroyed - * - * Test removed for Nvidia devices because it returns unexpected error + * Test Description + * ------------------------ + * - Query an uninitialized stream + * - Expected output: return `hipErrorContextIsDestroyed` + * Test source + * ------------------------ + * - unit/stream/hipStreamQuery.cc + * Test requirements + * ------------------------ + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamQuery_WithUninitializedStream") { hipStream_t stream{reinterpret_cast(0xFFFF)}; @@ -89,11 +136,18 @@ TEST_CASE("Unit_hipStreamQuery_WithUninitializedStream") { #endif #if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */ - /** - * @brief Check that submitting work to a stream sets the status of the nullStream to - * hipErrorNotReady - * + * Test Description + * ------------------------ + * - Query a null stream while another stream has submitted work that has not finished yet + * - Expected output: return `hipErrorNotReady` + * Test source + * ------------------------ + * - unit/stream/hipStreamQuery.cc + * Test requirements + * ------------------------ + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamQuery_SubmitWorkOnStreamAndQueryNullStream") { { @@ -110,9 +164,17 @@ TEST_CASE("Unit_hipStreamQuery_SubmitWorkOnStreamAndQueryNullStream") { } /** - * @brief Check that submitting work to the nullStream properly sets its status as - * hipErrorNotReady. - * + * Test Description + * ------------------------ + * - Query a null stream with submitted work that has not finished yet + * - Expected output: return `hipErrorNotReady` + * Test source + * ------------------------ + * - unit/stream/hipStreamQuery.cc + * Test requirements + * ------------------------ + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamQuery_NullStreamQuery") { HIP_CHECK(hipStreamQuery(hip::nullStream)); @@ -123,9 +185,17 @@ TEST_CASE("Unit_hipStreamQuery_NullStreamQuery") { } /** - * @brief Check that querying a stream with pending work returns hipErrorNotReady - * - **/ + * Test Description + * ------------------------ + * - Query a stream with pending work + * - Expected output: return `hipErrorNotReady` + * Test source + * ------------------------ + * - unit/stream/hipStreamQuery.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamQuery_WithPendingWork") { hipStream_t waitingStream{nullptr}; HIP_CHECK(hipStreamCreate(&waitingStream)); diff --git a/catch/unit/stream/hipStreamSynchronize.cc b/catch/unit/stream/hipStreamSynchronize.cc index d1faaa4f0..a7d40cf7f 100644 --- a/catch/unit/stream/hipStreamSynchronize.cc +++ b/catch/unit/stream/hipStreamSynchronize.cc @@ -18,13 +18,29 @@ THE SOFTWARE. */ #include -#include "streamCommon.hh" #include +#include "streamCommon.hh" + +/** + * @addtogroup hipStreamSynchronize hipStreamSynchronize + * @{ + * @ingroup StreamTest + * `hipStreamSynchronize(hipStream_t stream)` - + * Wait for all commands in stream to complete. + */ + namespace hipStreamSynchronizeTest { /** - * @brief Check that hipStreamSynchronize handles empty streams properly. - * + * Test Description + * ------------------------ + * - Synchronize an empty stream. + * Test source + * ------------------------ + * - unit/stream/hipStreamSynchronize.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamSynchronize_EmptyStream") { hipStream_t stream; @@ -34,11 +50,20 @@ TEST_CASE("Unit_hipStreamSynchronize_EmptyStream") { } #if !HT_NVIDIA +// Test removed for Nvidia devices because it returns unexpected error. + /** - * @brief Check that synchronization of uninitialized stream sets its status to - * hipErrorContextIsDestroyed - * - * Test removed for Nvidia devices because it returns unexpected error + * Test Description + * ------------------------ + * - Synchronize an uninitialized stream + * - Expected output: return `hipErrorContextIsDestroyed` + * Test source + * ------------------------ + * - unit/stream/hipStreamSynchronize.cc + * Test requirements + * ------------------------ + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamSynchronize_UninitializedStream") { hipStream_t stream{reinterpret_cast(0xFFFF)}; @@ -49,9 +74,16 @@ TEST_CASE("Unit_hipStreamSynchronize_UninitializedStream") { #if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */ /** - * @brief Check that all work executing in a stream is finished after a call to - * hipStreamSynchronize. - * + * Test Description + * ------------------------ + * - Check that all work executing in a stream is finished after synchronization. + * Test source + * ------------------------ + * - unit/stream/hipStreamSynchronize.cc + * Test requirements + * ------------------------ + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamSynchronize_FinishWork") { const hipStream_t explicitStream = reinterpret_cast(-1); @@ -72,7 +104,15 @@ TEST_CASE("Unit_hipStreamSynchronize_FinishWork") { } /** - * @brief Check that synchronizing the nullStream implicitly synchronizes all executing streams. + * Test Description + * ------------------------ + * - Check that synchronizing the nullStream implicitly synchronizes all executing streams. + * Test source + * ------------------------ + * - unit/stream/hipStreamSynchronize.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamSynchronize_NullStreamSynchronization") { int totalStreams = 10; @@ -108,9 +148,18 @@ TEST_CASE("Unit_hipStreamSynchronize_NullStreamSynchronization") { } /** - * @brief Check that synchronizing one stream does implicitly synchronize other streams. - * Check that submiting work to the nullStream does not affect synchronization of other - * streams. Check that querying the nullStream does not affect synchronization of other streams. + * Test Description + * ------------------------ + * - Check that synchronizing one stream does not synchronize other streams. + * - Check that submiting work to the nullStream does not affect synchronization of other streams. + * - Check that querying the nullStream does not affect synchronization of other streams. + * Test source + * ------------------------ + * - unit/stream/hipStreamSynchronize.cc + * Test requirements + * ------------------------ + * - Platform specific (NVIDIA) + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamSynchronize_SynchronizeStreamAndQueryNullStream") { #if HT_AMD @@ -152,9 +201,17 @@ TEST_CASE("Unit_hipStreamSynchronize_SynchronizeStreamAndQueryNullStream") { } /** - * @brief Check that synchronizing the nullStream also synchronizes the hipStreamPerThread - * special stream. - * + * Test Description + * ------------------------ + * - Check that synchronizing the null stream also synchronizes the + * per thread special stream. + * Test source + * ------------------------ + * - unit/stream/hipStreamSynchronize.cc + * Test requirements + * ------------------------ + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamSynchronize_NullStreamAndStreamPerThread") { LaunchDelayKernel(std::chrono::milliseconds(500), hip::streamPerThread); diff --git a/catch/unit/stream/hipStreamValue.cc b/catch/unit/stream/hipStreamValue.cc index b5844a255..44c084cf9 100644 --- a/catch/unit/stream/hipStreamValue.cc +++ b/catch/unit/stream/hipStreamValue.cc @@ -233,6 +233,34 @@ template struct TestParams { constexpr static PtrType ptrType = ptrTypeValue; }; +/** + * @addtogroup hipStreamWriteValue32 hipStreamWriteValue32 + * @{ + * @ingroup StreamMTest + * `hipStreamWriteValue32(hipStream_t stream, void* ptr, uint32_t value, unsigned int flags)` - + * Enqueues a write command to the stream, write operation is performed after all earlier commands + * on this stream have completed the execution. + * ________________________ + * Test cases from other modules: + * - @ref Unit_hipStreamValue_Negative_InvalidMemory + * - @ref Unit_hipStreamValue_Negative_UninitializedStream + */ + +/** + * Test Description + * ------------------------ + * - Write a value to a GPU visible pointer. + * - Check if write vas valid for memory types: + * -# Registered memory + * -# Device memory + * -# Signal memory + * Test source + * ------------------------ + * - unit/stream/hipStreamValue.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ #if HT_AMD TEMPLATE_TEST_CASE("Unit_hipStreamValue_Write", "", (TestParams), (TestParams), @@ -277,6 +305,11 @@ TEMPLATE_TEST_CASE("Unit_hipStreamValue_Write", "", (TestParams void syncAndCheckData(hipStream_t stream, UIntT* dataPtr, TestPtr signalPtr, size_t offset, TEST_WAIT tc, std::array& events) { @@ -473,7 +506,34 @@ DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("NoMask_Nor", 0xbddbddbdbddbddbd, 0xbddbddbdbddbddb3)) #undef DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64 -// Negative Tests +/** + * @addtogroup hipStreamWaitValue32 hipStreamWaitValue32 + * @{ + * @ingroup StreamMTest + */ + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments for [hipStreamWriteValue32](@ref hipStreamWriteValue32): + * -# When memory pointer is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * - Validates handling of invalid arguments for [hipStreamWriteValue64](@ref hipStreamWriteValue64): + * -# When memory pointer is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * - Validates handling of invalid arguments for [hipStreamWaitValue32](@ref hipStreamWaitValue32): + * -# When memory pointer is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * - Validates handling of invalid arguments for [hipStreamWaitValue64](@ref hipStreamWaitValue64): + * -# When memory pointer is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/stream/hipStreamValue.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamValue_Negative_InvalidMemory") { if (!streamWaitValueSupported()) { HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device."); @@ -505,6 +565,28 @@ TEST_CASE("Unit_hipStreamValue_Negative_InvalidMemory") { HIP_CHECK(hipStreamDestroy(stream)); } +/** + * Test Description + * ------------------------ + * - Validates handling of uninitialized stream for [hipStreamWriteValue32](@ref hipStreamWriteValue32): + * -# When stream is uninitialized + * - Expected output: return `hipErrorContextIsDestroyed` + * - Validates handling of uninitialized stream for [hipStreamWriteValue64](@ref hipStreamWriteValue64): + * -# When stream is uninitialized + * - Expected output: return `hipErrorContextIsDestroyed` + * - Validates handling of uninitialized stream for [hipStreamWaitValue32](@ref hipStreamWaitValue32): + * -# When stream is uninitialized + * - Expected output: return `hipErrorContextIsDestroyed` + * - Validates handling of uninitialized stream for [hipStreamWaitValue64](@ref hipStreamWaitValue64): + * -# When stream is uninitialized + * - Expected output: return `hipErrorContextIsDestroyed` + * Test source + * ------------------------ + * - unit/stream/hipStreamValue.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEMPLATE_TEST_CASE("Unit_hipStreamValue_Negative_UninitializedStream", "", uint32_t, uint64_t) { if (!streamWaitValueSupported()) { HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device."); @@ -538,6 +620,22 @@ TEMPLATE_TEST_CASE("Unit_hipStreamValue_Negative_UninitializedStream", "", uint3 HIP_CHECK(hipHostUnregister(hostPtr.get())); } +/** + * Test Description + * ------------------------ + * - Validates handling of invalid flags for [hipStreamWaitValue32](@ref hipStreamWaitValue32): + * -# When flags are not in valid range of values + * - Expected output: return `hipErrorInvalidValue` + * - Validates handling of invalid flags for [hipStreamWaitValue64](@ref hipStreamWaitValue64): + * -# When flags are not in valid range of values + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/stream/hipStreamValue.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEMPLATE_TEST_CASE("Unit_hipStreamValue_Negative_InvalidFlag", "", uint32_t, uint64_t) { if (!streamWaitValueSupported()) { HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device."); @@ -563,3 +661,8 @@ TEMPLATE_TEST_CASE("Unit_hipStreamValue_Negative_InvalidFlag", "", uint32_t, uin HIP_CHECK(hipHostUnregister(hostPtr.get())); HIP_CHECK(hipStreamDestroy(stream)); } + +/** + * End doxygen group hipStreamWaitValue32. + * @} + */ diff --git a/catch/unit/stream/hipStreamWaitEvent.cc b/catch/unit/stream/hipStreamWaitEvent.cc index 32dbe9983..8aa4c8281 100644 --- a/catch/unit/stream/hipStreamWaitEvent.cc +++ b/catch/unit/stream/hipStreamWaitEvent.cc @@ -16,16 +16,33 @@ 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. */ -/* -Testcase Scenarios : -Unit_hipStreamWaitEvent_Negative - Test unsuccessful hipStreamWaitEvent when either event or flags are invalid -Unit_hipStreamWaitEvent_UninitializedStream_Negative - Test unsuccessful hipStreamWaitEvent when stream is uninitialized -Unit_hipStreamWaitEvent_Default - Test simple waiting for an event with hipStreamWaitEvent api -Unit_hipStreamWaitEvent_DifferentStreams - Test waiting for an event on a different stream with hipStreamWaitEvent api -*/ #include #include + +/** + * @addtogroup hipStreamWaitEvent hipStreamWaitEvent + * @{ + * @ingroup StreamTest + * `hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags)` - + * Make the specified compute stream wait for an event. + */ + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When event handle is `nullptr` + * - Expected output: return `hipErrorInvalidResourceHandle` + * -# When flags are not valid + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/stream/hipStreamWaitEvent.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamWaitEvent_Negative") { enum class StreamTestType { NullStream = 0, StreamPerThread, CreatedStream }; @@ -67,6 +84,19 @@ TEST_CASE("Unit_hipStreamWaitEvent_Negative") { /* Test removed for Nvidia devices because it returns unexpected error */ #if !HT_NVIDIA +/** + * Test Description + * ------------------------ + * - Waits for event on stream that has not been initialized + * - Expected output: return `hipErrorContextIsDestroyed` + * Test source + * ------------------------ + * - unit/stream/hipStreamWaitEvent.cc + * Test requirements + * ------------------------ + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamWaitEvent_UninitializedStream_Negative") { hipStream_t stream{reinterpret_cast(0xFFFF)}; hipEvent_t event{nullptr}; @@ -104,6 +134,17 @@ TEST_CASE("Unit_hipStreamWaitEvent_Default") { HIP_CHECK(hipEventDestroy(waitEvent)); } +/** + * Test Description + * ------------------------ + * - Create multiple dependant kernels and synchronize between them with streams and waiting on events. + * Test source + * ------------------------ + * - unit/stream/hipStreamWaitEvent.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipStreamWaitEvent_DifferentStreams") { hipStream_t blockedStreamA{nullptr}, streamBlockedOnStreamA{nullptr}, unblockingStream{nullptr}; hipEvent_t waitEvent{nullptr}; diff --git a/catch/unit/stream/hipStreamWithCUMask.cc b/catch/unit/stream/hipStreamWithCUMask.cc index 61535ed2c..60fd46c65 100644 --- a/catch/unit/stream/hipStreamWithCUMask.cc +++ b/catch/unit/stream/hipStreamWithCUMask.cc @@ -17,23 +17,6 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : - -1)Validates functionality of hipStreamAddCallback with created stream. - -2)Validates functionality of stream with cu mask. - -3)Create a stream with all CU masks disabled (0x00000000). -Verify that default CU mask is set for the stream. - -4)Size is greater than physical CU number. In this case the extra elements -are ignored and hipExtStreamCreateWithCUMask must return hipSuccess. - -5)Negative Testing of hipExtStreamCreateWithCUMask. -*/ - - #include #include #include @@ -41,6 +24,15 @@ are ignored and hipExtStreamCreateWithCUMask must return hipSuccess. #include #include +/** + * @addtogroup hipExtStreamCreateWithCUMask hipExtStreamCreateWithCUMask + * @{ + * @ingroup StreamTest + * `hipExtStreamCreateWithCUMask(hipStream_t* stream, uint32_t cuMaskSize, + * const uint32_t* cuMask)` - + * Create an asynchronous stream with the specified CU mask. + */ + #define NUM_CU_PARTITIONS 4 #define CONSTANT 1.618f #define SIZE_INBYTES_OF_MB (1024*1024) @@ -121,9 +113,18 @@ using hipExtStreamCreateWithCUMaskTest::createDefaultCUMask; using hipExtStreamCreateWithCUMaskTest::createDisabledCUMask; using hipExtStreamCreateWithCUMaskTest::Callback; - /** - * Scenario: Validates functionality of hipStreamAddCallback with created stream. + * Test Description + * ------------------------ + * - Creates stream with CU mask. + * - Adds callback to the created stream. + * - Successfully destroys the stream. + * Test source + * ------------------------ + * - unit/stream/hipStreamWithCUMask.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_ValidateCallbackFunc") { float *A_d, *C_d; @@ -171,7 +172,15 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_ValidateCallbackFunc") { } /** - * Scenario: Validates functionality of stream with cu mask. + * Test Description + * ------------------------ + * - Creates a stream for each possible CU mask. + * Test source + * ------------------------ + * - unit/stream/hipStreamWithCUMask.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_Functionality") { const int KNumPartition = NUM_CU_PARTITIONS; @@ -304,8 +313,15 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_Functionality") { } /** - * Scenario: Create a stream with all CU masks disabled (0x00000000). - * Verify that default CU mask is set for the stream. + * Test Description + * ------------------------ + * - Verifies that the stream is created with default CU mask if all of the CU's are disabled. + * Test source + * ------------------------ + * - unit/stream/hipStreamWithCUMask.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_AllCUsMasked") { HIP_CHECK(hipSetDevice(0)); @@ -330,7 +346,21 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_AllCUsMasked") { } /** - * Scenario: Negative Testing of hipExtStreamCreateWithCUMask. + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When the output stream pointer is `nullptr` + * - Expected output: do not return `hipSuccess` + * -# When the CU mask size is 0 + * - Expected output: do not return `hipSuccess` + * -# When the CU mask pointer is `nullptr` + * - Expected output: do not return `hipSuccess` + * Test source + * ------------------------ + * - unit/stream/hipStreamWithCUMask.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_NegTst") { std::vector defaultCUMask;