diff --git a/catch/unit/graph/hipGraph.cc b/catch/unit/graph/hipGraph.cc index 59c552c75..bd0f16dc4 100644 --- a/catch/unit/graph/hipGraph.cc +++ b/catch/unit/graph/hipGraph.cc @@ -17,16 +17,14 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : - 1) Execution Without HIPGraphs : Regular procedure of using stream with async api calls. - 2) Manual HIPGraph : Manual procedure of adding nodes to graphs and mapping dependencies. - 3) HIPGraphs Using StreamCapture : Capturing sequence of operations in stream and launching - graph with the nodes automatically added. -*/ - #include +/** + * @addtogroup hipGraphCreate hipGraphCreate + * @{ + * @ingroup GraphTest + */ + #define THREADS_PER_BLOCK 512 #define GRAPH_LAUNCH_ITERATIONS 1000 @@ -65,9 +63,7 @@ static void init_input(float* a, size_t size) { } } -/** - * Regular procedure of using stream with async api calls - */ +// Regular procedure of using stream with async api calls static void hipWithoutGraphs(float* inputVec_h, float* inputVec_d, double* outputVec_d, double* result_d, size_t inputSize, size_t numOfBlocks) { hipStream_t stream1, stream2, stream3; @@ -116,10 +112,10 @@ static void hipWithoutGraphs(float* inputVec_h, float* inputVec_d, REQUIRE(result_h_cpu == result_h); } -/** - * Capturing sequence of operations in stream and launching graph - * with the nodes automatically added. - */ +/* + Capturing sequence of operations in stream and launching graph + with the nodes automatically added. +*/ static void hipGraphsUsingStreamCapture(float* inputVec_h, float* inputVec_d, double* outputVec_d, double* result_d, size_t inputSize, size_t numOfBlocks) { @@ -197,9 +193,7 @@ static void hipGraphsUsingStreamCapture(float* inputVec_h, float* inputVec_d, REQUIRE(result_h_cpu == result_h); } -/** - * Manual procedure of adding nodes to graphs and mapping dependencies. - */ +// Manual procedure of adding nodes to graphs and mapping dependencies. static void hipGraphsManual(float* inputVec_h, float* inputVec_d, double* outputVec_d, double* result_d, size_t inputSize, size_t numOfBlocks) { @@ -304,8 +298,18 @@ static void hipGraphsManual(float* inputVec_h, float* inputVec_d, } /** - * Tests basic functionality of hipGraph APIs by - * Execution Without HIPGraphs, Manual HIPGraph, HIPGraphs Using StreamCapture. + * Test Description + * ------------------------ + * - Validates several basic Graph functionalities: + * -# Execution without Graphs: regular procedure of using stream with async API calls. + * -# Manual Graphs: manual procedure of adding nodes to graphs and adding dependencies. + * -# Graphs using Stream Capture: capturing sequence of operations in stream and launching. + * Test source + * ------------------------ + * - unit/graph/hipGraph.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraph_BasicFunctional") { constexpr size_t size = 1 << 12; diff --git a/catch/unit/graph/hipGraphAddChildGraphNode.cc b/catch/unit/graph/hipGraphAddChildGraphNode.cc index 9c8bf4c09..da8450b16 100644 --- a/catch/unit/graph/hipGraphAddChildGraphNode.cc +++ b/catch/unit/graph/hipGraphAddChildGraphNode.cc @@ -17,65 +17,41 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios of hipGraphAddChildGraphNode API: - -Functional: -1. Create child graph as root node and execute the main graph. -2. Create multiple child graph nodes and check the behaviour. -3. Clone the child graph node, Add new nodes and execute the cloned graph. -4. Create child graph, add it to main graph and execute child graph. -5. Pass original graph as child graph and execute the org graph. -6. This test case verifies nested graph functionality. Parent graph - containing child graph, which in turn, contains another child graph. - Execute the graph in loop taking random input data and Validate the - output in each iteration. -7. This test case verifies clones the nested graph created in scenario6. - Execute the cloned graph in loop taking random input data and Validate - the output in each iteration. -8. Verify if an empty graph can be added as child node. -9. Create the nested graph of scenario6 and update the property of add kernel - node (innermost graph) with subtract kernel functionality. Clone the graph. - Execute both the updated graph. -10. The updated nested graph in 9 is cloned and the cloned graph is then - executed and the result is validated. -11. Create the nested graph of 6 and update the block size and grid size - property of add kernel node. -12. Create the nested graph of 6 and delete the add kernel node - (innermost graph) and add a subtract kernel node. -13. The updated nested graph in 12 is cloned and the cloned graph is then - executed and the result is validated. -14. Create the nested graph of 6 and delete the add kernel node - (innermost graph), add a child graph that contains an event record node, - a subtract kernel node followed by another event record node. Clone the - graph. Execute both the original and cloned graph. -15. The updated nested graph in 14 is cloned and the cloned graph is then - executed and the result is validated. -16. Create one nested graph per GPU context. Execute all the created graphs - in their respective GPUs and validate the output. -17. Functional Test to use child node as barrier to wait for multiple nodes. - This test uses child nodes to resolve dependencies between graphs. 4 - graphs are created. Graph1 contains 3 independent memcpy h2d nodes, graph2 - contains 3 independent kernel nodes and graph3 contains 3 independent - memcpy d2h nodes. Graph1, graph2 and graph3 are added as child nodes in - graph4. Graph4 is validated for functionality. - -Negative: -1. Pass nullptr to graph node -2. Pass nullptr to graph -3. Pass invalid number of numDepdencies -4. Pass nullptr to child graph -*/ - #include #include #include #define TEST_LOOP_SIZE 50 -/* -This testcase verifies the negative scenarios of -hipGraphAddChildGraphNode API -*/ + +/** + * @addtogroup hipGraphAddChildGraphNode hipGraphAddChildGraphNode + * @{ + * @ingroup GraphTest + * `hipGraphAddChildGraphNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + * const hipGraphNode_t* pDependencies, size_t numDependencies, + * hipGraph_t childGraph)` - + * Creates a child graph node and adds it to a graph. + */ + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When graph node pointer is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When graph handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When child graph handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When dependencies pointer is `nullptr` and number is not valid + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphAddChildGraphNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphAddChildGraphNode_Negative") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); @@ -93,7 +69,7 @@ TEST_CASE("Unit_hipGraphAddChildGraphNode_Negative") { 0, A_h, B_d, Nbytes, hipMemcpyDeviceToHost)); - SECTION("Pass nullptr to graph noe") { + SECTION("Pass nullptr to graph node") { REQUIRE(hipGraphAddChildGraphNode(nullptr, graph, nullptr, 0, childgraph1) == hipErrorInvalidValue); @@ -111,18 +87,26 @@ TEST_CASE("Unit_hipGraphAddChildGraphNode_Negative") { == hipErrorInvalidValue); } - SECTION("Pass invalid depdencies") { + SECTION("Pass invalid dependencies") { REQUIRE(hipGraphAddChildGraphNode(&childGraphNode1, graph, nullptr, 10, childgraph1) == hipErrorInvalidValue); } } -/* -This testcase verifies the following scenario -Creates the graph, add the graph as a child node -and verify the number of the nodes in the original graph -*/ +/** + * Test Description + * ------------------------ + * - Creates the graph. + * - Adds the graph as a child node. + * - Verifies the number of the nodes in the original graph. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddChildGraphNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphAddChildGraphNode_OrgGraphAsChildGraph") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); @@ -160,11 +144,19 @@ TEST_CASE("Unit_hipGraphAddChildGraphNode_OrgGraphAsChildGraph") { HIP_CHECK(hipStreamDestroy(streamForGraph)); } -/* -This testcase verifies the following scenario -Create graph, Add child nodes to the graph and execute only the -child graph node and verify the behaviour -*/ +/** + * Test Description + * ------------------------ + * - Creates a graph. + * - Adds child nodes to the graph. + * - Executes only the child graphs and verifies the behaviour. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddChildGraphNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphAddChildGraphNode_ExecuteChildGraph") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); @@ -217,11 +209,19 @@ TEST_CASE("Unit_hipGraphAddChildGraphNode_ExecuteChildGraph") { HIP_CHECK(hipStreamDestroy(streamForGraph)); } -/* -This testcase verifies the following scenario -creates graph, Add child nodes to graph, clone the graph and execute -the cloned graph -*/ +/** + * Test Description + * ------------------------ + * - Creates the graphs. + * - Adds child nodes to graph. + * - Clones the graphs and executes the cloned graph. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddChildGraphNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphAddChildGraphNode_CloneChildGraph") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); @@ -271,11 +271,18 @@ TEST_CASE("Unit_hipGraphAddChildGraphNode_CloneChildGraph") { HIP_CHECK(hipStreamDestroy(streamForGraph)); } -/* -This testcase verifies the following scenario -Create graph, add multiple child nodes and validates the -behaviour -*/ +/** + * Test Description + * ------------------------ + * - Creates the graph. + * - Adds multiple child nodes and validates the behaviour. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddChildGraphNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphAddChildGraphNode_MultipleChildNodes") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); @@ -344,9 +351,18 @@ TEST_CASE("Unit_hipGraphAddChildGraphNode_MultipleChildNodes") { HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipStreamDestroy(streamForGraph)); } + /** - This testcase verifies hipGraphAddChildGraphNode functionality - where root node is the child node. + * Test Description + * ------------------------ + * - Create child graph as root node. + * - Execute the main graph. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddChildGraphNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphAddChildGraphNode_SingleChildNode") { constexpr size_t N = 1024; diff --git a/catch/unit/graph/hipGraphAddEmptyNode.cc b/catch/unit/graph/hipGraphAddEmptyNode.cc index 23d477f59..4b0871fa9 100644 --- a/catch/unit/graph/hipGraphAddEmptyNode.cc +++ b/catch/unit/graph/hipGraphAddEmptyNode.cc @@ -17,27 +17,34 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : - 1) Create and add empty node to graph and verify addition is successful. - 2) Negative Scenarios - 3) Functional Test to use empty node as barrier to wait for multiple nodes. - In a graph add 3 independent memcpy_h2d nodes, add an empty node with - dependencies on the 3 memcpy_h2d nodes, add 3 independent kernel nodes, - add another empty node with dependencies on the 3 kernel nodes and - add 3 independent memcpy_d2h nodes with dependencies on previous empty - node. Execute the graph and validate the results. -*/ -#include #include #include #define TEST_LOOP_SIZE 50 /** - * Functional Test to add empty node with dependencies + * @addtogroup hipGraphAddEmptyNode hipGraphAddEmptyNode + * @{ + * @ingroup GraphTest + * `hipGraphAddEmptyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + * const hipGraphNode_t* pDependencies, size_t numDependencies)` - + * Creates an empty node and adds it to a graph. + */ + +/** + * Test Description + * ------------------------ + * - Creates an empty node. + * - Adds empty node to the graph. + * - Verifies that the addition is successful. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddEmptyNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphAddEmptyNode_Functional") { - char *pOutBuff_d{}; + char* pOutBuff_d{}; constexpr size_t size = 1024; hipGraph_t graph{}; hipGraphNode_t memsetNode{}, emptyNode{}; @@ -52,13 +59,11 @@ TEST_CASE("Unit_hipGraphAddEmptyNode_Functional") { memsetParams.width = size * sizeof(char); memsetParams.height = 1; HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, - &memsetParams)); + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams)); dependencies.push_back(memsetNode); // Create emptyNode and add it to graph with dependency - HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, graph, dependencies.data(), - dependencies.size())); + HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, graph, dependencies.data(), dependencies.size())); REQUIRE(emptyNode != nullptr); HIP_CHECK(hipFree(pOutBuff_d)); @@ -66,10 +71,24 @@ TEST_CASE("Unit_hipGraphAddEmptyNode_Functional") { } /** - * Negative Scenarios hipGraphAddEmptyNode + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When empty graph node pointer is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When graph is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When dependencies are `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphAddEmptyNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphAddEmptyNode_NegTest") { - char *pOutBuff_d{}; + char* pOutBuff_d{}; constexpr size_t size = 1024; hipGraph_t graph; hipGraphNode_t memsetNode{}, emptyNode{}; @@ -84,23 +103,22 @@ TEST_CASE("Unit_hipGraphAddEmptyNode_NegTest") { memsetParams.width = size * sizeof(char); memsetParams.height = 1; HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, - &memsetParams)); + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams)); dependencies.push_back(memsetNode); // pGraphNode is nullptr SECTION("Null Empty Graph Node") { - REQUIRE(hipErrorInvalidValue == hipGraphAddEmptyNode(nullptr, graph, - dependencies.data(), dependencies.size())); + REQUIRE(hipErrorInvalidValue == + hipGraphAddEmptyNode(nullptr, graph, dependencies.data(), dependencies.size())); } // graph is nullptr SECTION("Null Graph") { - REQUIRE(hipErrorInvalidValue == hipGraphAddEmptyNode(&emptyNode, nullptr, - dependencies.data(), dependencies.size())); + REQUIRE(hipErrorInvalidValue == + hipGraphAddEmptyNode(&emptyNode, nullptr, dependencies.data(), dependencies.size())); } // pDependencies is nullptr - SECTION("Dependencies is null") { - REQUIRE(hipErrorInvalidValue == hipGraphAddEmptyNode(&emptyNode, graph, - nullptr, dependencies.size())); + SECTION("Null Dependencies") { + REQUIRE(hipErrorInvalidValue == + hipGraphAddEmptyNode(&emptyNode, graph, nullptr, dependencies.size())); } HIP_CHECK(hipFree(pOutBuff_d)); @@ -108,7 +126,7 @@ TEST_CASE("Unit_hipGraphAddEmptyNode_NegTest") { } // Function to fill input data -static void fillRandInpData(int *A1_h, int *A2_h, int *A3_h, size_t N) { +static void fillRandInpData(int* A1_h, int* A2_h, int* A3_h, size_t N) { unsigned int seed = time(nullptr); for (size_t i = 0; i < N; i++) { A1_h[i] = (HipTest::RAND_R(&seed) & 0xFF); @@ -118,9 +136,9 @@ static void fillRandInpData(int *A1_h, int *A2_h, int *A3_h, size_t N) { } // Function to validate result -static void validateOutData(int *A1_h, int *A2_h, size_t N) { +static void validateOutData(int* A1_h, int* A2_h, size_t N) { for (size_t i = 0; i < N; i++) { - int result = (A1_h[i]*A1_h[i]); + int result = (A1_h[i] * A1_h[i]); REQUIRE(result == A2_h[i]); } } @@ -131,106 +149,98 @@ TEST_CASE("Unit_hipGraphAddEmptyNode_BarrierFunc") { size_t size = 1024; constexpr auto blocksPerCU = 6; constexpr auto threadsPerBlock = 256; - unsigned blocks = HipTest::setNumBlocks(blocksPerCU, - threadsPerBlock, size); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, size); hipGraph_t graph; std::vector nodeDependencies; HIP_CHECK(hipGraphCreate(&graph, 0)); - int *inputVec_d1{nullptr}, *inputVec_h1{nullptr}, *outputVec_h1{nullptr}, - *outputVec_d1{nullptr}; - int *inputVec_d2{nullptr}, *inputVec_h2{nullptr}, *outputVec_h2{nullptr}, - *outputVec_d2{nullptr}; - int *inputVec_d3{nullptr}, *inputVec_h3{nullptr}, *outputVec_h3{nullptr}, - *outputVec_d3{nullptr}; + int *inputVec_d1{nullptr}, *inputVec_h1{nullptr}, *outputVec_h1{nullptr}, *outputVec_d1{nullptr}; + int *inputVec_d2{nullptr}, *inputVec_h2{nullptr}, *outputVec_h2{nullptr}, *outputVec_d2{nullptr}; + int *inputVec_d3{nullptr}, *inputVec_h3{nullptr}, *outputVec_h3{nullptr}, *outputVec_d3{nullptr}; // host and device allocation - HipTest::initArrays(&inputVec_d1, &outputVec_d1, nullptr, - &inputVec_h1, &outputVec_h1, nullptr, size, false); - HipTest::initArrays(&inputVec_d2, &outputVec_d2, nullptr, - &inputVec_h2, &outputVec_h2, nullptr, size, false); - HipTest::initArrays(&inputVec_d3, &outputVec_d3, nullptr, - &inputVec_h3, &outputVec_h3, nullptr, size, false); + HipTest::initArrays(&inputVec_d1, &outputVec_d1, nullptr, &inputVec_h1, &outputVec_h1, + nullptr, size, false); + HipTest::initArrays(&inputVec_d2, &outputVec_d2, nullptr, &inputVec_h2, &outputVec_h2, + nullptr, size, false); + HipTest::initArrays(&inputVec_d3, &outputVec_d3, nullptr, &inputVec_h3, &outputVec_h3, + nullptr, size, false); // add nodes to graph hipGraphNode_t memcpyH2D_1, memcpyH2D_2, memcpyH2D_3; hipGraphNode_t vecSqr1, vecSqr2, vecSqr3; hipGraphNode_t memcpyD2H_1, memcpyD2H_2, memcpyD2H_3; hipGraphNode_t emptyNode1, emptyNode2; // Create memcpy h2d nodes - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_1, graph, nullptr, - 0, inputVec_d1, inputVec_h1, (sizeof(int)*size), hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_2, graph, nullptr, - 0, inputVec_d2, inputVec_h2, (sizeof(int)*size), hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_3, graph, nullptr, - 0, inputVec_d3, inputVec_h3, (sizeof(int)*size), hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_1, graph, nullptr, 0, inputVec_d1, inputVec_h1, + (sizeof(int) * size), hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_2, graph, nullptr, 0, inputVec_d2, inputVec_h2, + (sizeof(int) * size), hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_3, graph, nullptr, 0, inputVec_d3, inputVec_h3, + (sizeof(int) * size), hipMemcpyHostToDevice)); // Create dependency list nodeDependencies.push_back(memcpyH2D_1); nodeDependencies.push_back(memcpyH2D_2); nodeDependencies.push_back(memcpyH2D_3); // Create emptyNode and add it to graph with dependency - HIP_CHECK(hipGraphAddEmptyNode(&emptyNode1, graph, nodeDependencies.data(), - nodeDependencies.size())); + HIP_CHECK( + hipGraphAddEmptyNode(&emptyNode1, graph, nodeDependencies.data(), nodeDependencies.size())); nodeDependencies.clear(); nodeDependencies.push_back(emptyNode1); // Creating kernel nodes hipKernelNodeParams kerNodeParams1{}, kerNodeParams2{}, kerNodeParams3{}; void* kernelArgs1[] = {reinterpret_cast(&inputVec_d1), - reinterpret_cast(&outputVec_d1), - reinterpret_cast(&size)}; + reinterpret_cast(&outputVec_d1), reinterpret_cast(&size)}; kerNodeParams1.func = reinterpret_cast(HipTest::vector_square); kerNodeParams1.gridDim = dim3(blocks); kerNodeParams1.blockDim = dim3(threadsPerBlock); kerNodeParams1.sharedMemBytes = 0; kerNodeParams1.kernelParams = reinterpret_cast(kernelArgs1); kerNodeParams1.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&vecSqr1, graph, nodeDependencies.data(), - nodeDependencies.size(), &kerNodeParams1)); + HIP_CHECK(hipGraphAddKernelNode(&vecSqr1, graph, nodeDependencies.data(), nodeDependencies.size(), + &kerNodeParams1)); void* kernelArgs2[] = {reinterpret_cast(&inputVec_d2), - reinterpret_cast(&outputVec_d2), - reinterpret_cast(&size)}; + reinterpret_cast(&outputVec_d2), reinterpret_cast(&size)}; kerNodeParams2.func = reinterpret_cast(HipTest::vector_square); kerNodeParams2.gridDim = dim3(blocks); kerNodeParams2.blockDim = dim3(threadsPerBlock); kerNodeParams2.sharedMemBytes = 0; kerNodeParams2.kernelParams = reinterpret_cast(kernelArgs2); kerNodeParams2.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&vecSqr2, graph, nodeDependencies.data(), - nodeDependencies.size(), &kerNodeParams2)); + HIP_CHECK(hipGraphAddKernelNode(&vecSqr2, graph, nodeDependencies.data(), nodeDependencies.size(), + &kerNodeParams2)); void* kernelArgs3[] = {reinterpret_cast(&inputVec_d3), - reinterpret_cast(&outputVec_d3), - reinterpret_cast(&size)}; + reinterpret_cast(&outputVec_d3), reinterpret_cast(&size)}; kerNodeParams3.func = reinterpret_cast(HipTest::vector_square); kerNodeParams3.gridDim = dim3(blocks); kerNodeParams3.blockDim = dim3(threadsPerBlock); kerNodeParams3.sharedMemBytes = 0; kerNodeParams3.kernelParams = reinterpret_cast(kernelArgs3); kerNodeParams3.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&vecSqr3, graph, nodeDependencies.data(), - nodeDependencies.size(), &kerNodeParams3)); + HIP_CHECK(hipGraphAddKernelNode(&vecSqr3, graph, nodeDependencies.data(), nodeDependencies.size(), + &kerNodeParams3)); nodeDependencies.clear(); nodeDependencies.push_back(vecSqr1); nodeDependencies.push_back(vecSqr2); nodeDependencies.push_back(vecSqr3); // Create emptyNode and add it to graph with dependency - HIP_CHECK(hipGraphAddEmptyNode(&emptyNode2, graph, nodeDependencies.data(), - nodeDependencies.size())); + HIP_CHECK( + hipGraphAddEmptyNode(&emptyNode2, graph, nodeDependencies.data(), nodeDependencies.size())); nodeDependencies.clear(); nodeDependencies.push_back(emptyNode2); // Create memcpy d2h nodes - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_1, graph, - nodeDependencies.data(), nodeDependencies.size(), outputVec_h1, - outputVec_d1, (sizeof(int)*size), hipMemcpyDeviceToHost)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_2, graph, - nodeDependencies.data(), nodeDependencies.size(), outputVec_h2, - outputVec_d2, (sizeof(int)*size), hipMemcpyDeviceToHost)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_3, graph, - nodeDependencies.data(), nodeDependencies.size(), outputVec_h3, - outputVec_d3, (sizeof(int)*size), hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_1, graph, nodeDependencies.data(), + nodeDependencies.size(), outputVec_h1, outputVec_d1, + (sizeof(int) * size), hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_2, graph, nodeDependencies.data(), + nodeDependencies.size(), outputVec_h2, outputVec_d2, + (sizeof(int) * size), hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_3, graph, nodeDependencies.data(), + nodeDependencies.size(), outputVec_h3, outputVec_d3, + (sizeof(int) * size), hipMemcpyDeviceToHost)); nodeDependencies.clear(); // Create executable graph hipStream_t streamForGraph; hipGraphExec_t graphExec{nullptr}; HIP_CHECK(hipStreamCreate(&streamForGraph)); - HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, - nullptr, 0)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); // Execute graph for (int iter = 0; iter < TEST_LOOP_SIZE; iter++) { fillRandInpData(inputVec_h1, inputVec_h2, inputVec_h3, size); @@ -243,11 +253,11 @@ TEST_CASE("Unit_hipGraphAddEmptyNode_BarrierFunc") { HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipStreamDestroy(streamForGraph)); // Free - HipTest::freeArrays(inputVec_d1, outputVec_d1, nullptr, - inputVec_h1, outputVec_h1, nullptr, false); - HipTest::freeArrays(inputVec_d2, outputVec_d2, nullptr, - inputVec_h2, outputVec_h2, nullptr, false); - HipTest::freeArrays(inputVec_d3, outputVec_d3, nullptr, - inputVec_h3, outputVec_h3, nullptr, false); + HipTest::freeArrays(inputVec_d1, outputVec_d1, nullptr, inputVec_h1, outputVec_h1, nullptr, + false); + HipTest::freeArrays(inputVec_d2, outputVec_d2, nullptr, inputVec_h2, outputVec_h2, nullptr, + false); + HipTest::freeArrays(inputVec_d3, outputVec_d3, nullptr, inputVec_h3, outputVec_h3, nullptr, + false); HIP_CHECK(hipGraphDestroy(graph)); } diff --git a/catch/unit/graph/hipGraphAddEventRecordNode.cc b/catch/unit/graph/hipGraphAddEventRecordNode.cc index 4f5b3b0da..223f80a3e 100644 --- a/catch/unit/graph/hipGraphAddEventRecordNode.cc +++ b/catch/unit/graph/hipGraphAddEventRecordNode.cc @@ -17,41 +17,6 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : - 1) Simple Scenario: Create an event node and add it to graph. -Instantiate and Launch the Graph. Wait for the event to complete. -The operation must succeed without any failures. - 2) Add different kinds of nodes to graph and add dependencies to nodes. -Create an event record node at the end. Instantiate and Launch the Graph. -Wait for the event to complete. Verify the results. Event is created using -hipEventCreate. - 3) Add different kinds of nodes to graph and add dependencies to nodes. -Create event record nodes at the beginning and end. Instantiate and Launch -the Graph. Wait for the event to complete. Verify the results. Also verify -the elapsed time. Events are created using hipEventCreate. - 4) Add different kinds of nodes to graph and add dependencies to nodes. -Create an event record node at the end. Instantiate and Launch graph. -Wait for the event to complete. Verify the results. Event is created -using hipEventCreateWithFlags (for different flag values). - 5) Create event record node at the beginning with -flag = hipEventDisableTiming, a memset node and event record nodes at the -end. Instantiate and Launch the Graph. Wait for the event to complete. -Verify that hipEventElapsedTime() returns error. - 6) Validate scenario 2 by running the graph multiple times in a loop -(100 times) after instantiation. - 7) Validate that no error is reported when numDeps <= dependencies length - 8) Negative Scenarios - - Output node is a nullptr. - - Input graph is a nullptr. - - Input dependencies is a nullptr. - - Node in dependency is from different graph - - Invalid numNodes - - Duplicate node in dependencies - - Input event is a nullptr. - - Input graph is uninitialized. - - Input event is uninitialized. -*/ #include #include @@ -61,8 +26,26 @@ Verify that hipEventElapsedTime() returns error. #include "graph_tests_common.hh" /** - * Scenario 1: Create s simple graph with just one event record - * node and instantiate and launch the graph. + * @addtogroup hipGraphAddEventRecordNode hipGraphAddEventRecordNode + * @{ + * @ingroup GraphTest + * `hipGraphAddEventRecordNode(hipGraphNode_t* pGraphNode, + * hipGraph_t graph, const hipGraphNode_t* pDependencies, + * size_t numDependencies, hipEvent_t event)` - + * Creates an event record node and adds it to a graph. + */ + +/** + * Test Description + * ------------------------ + * - Creates a simple graph with just one event record node. + * - Instantiates the graph and launches it without errors. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddEventRecordNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_Simple") { hipGraph_t graph; @@ -198,7 +181,19 @@ static void validateAddEventRecordNode(bool measureTime, bool withFlags, int nst } /** - * Scenario 2: Validate event record nodes created without flags. + * Test Description + * ------------------------ + * - Add different kinds of nodes to graph and add dependencies to nodes. + * - Create an event record node at the end. + * - Instantiate and launch the graph. + * - Wait for the event to complete. + * - Verify the results. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddEventRecordNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags") { // Create events without flags using hipEventCreate and @@ -207,7 +202,20 @@ TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags") { } /** - * Scenario 3: Validate elapsed time between 2 recorded events. + * Test Description + * ------------------------ + * - Add different kinds of nodes to graph and add dependencies to nodes. + * - Create event record nodes at the beginning and end. + * - Instantiate and launch the graph. + * - Wait for the event to complete. + * - Verify the results. + * - Also verify the elapsed time. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddEventRecordNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_ElapsedTime") { // Create events without flags using hipEventCreate and @@ -216,8 +224,22 @@ TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_ElapsedTime") { } /** - * Scenario 4: Validate event record nodes created with different - * event flags. + * Test Description + * ------------------------ + * - Add different kinds of nodes to graph and add dependencies to nodes. + * - Create an event record nodes with flags at the end. + * -# When flag is `hipEventDefault` + * -# When flag is `hipEventBlockingSync` + * -# When flag is `hipEventDisableTiming` + * - Instantiate and launch graph. + * - Wait for the event to complete. + * - Verify the results. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddEventRecordNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_WithFlags") { // Create events with different flags using hipEventCreate and @@ -236,15 +258,36 @@ TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_WithFlags") { } /** - * Scenario 5: Validate hipGraphAddEventRecordNode by executing graph - * 100 times in a loop. + * Test Description + * ------------------------ + * - Validate scenario @ref Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags + * by running the graph multiple times in a loop + * (100 times) after instantiation. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddEventRecordNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphAddEventRecordNode_MultipleRun") { validateAddEventRecordNode(false, false, 100); } /** - * Scenario 6: Validate hipGraphAddEventRecordNode with time disabled events. + * Test Description + * ------------------------ + * - Create event record node at the beginning with flag `hipEventDisableTiming`. + * - Add a memset node and event record nodes at the end. + * - Instantiate and launch the graph. + * - Wait for the event to complete. + * - Verify that elapsed time returns error. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddEventRecordNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_TimingDisabled") { constexpr size_t Nbytes = 1024; @@ -297,7 +340,21 @@ TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_TimingDisabled") { } /** - * Scenario 7: Positive parameter tests + * Test Description + * ------------------------ + * - Validate several positive scenarios: + * -# When number of dependencies is zero, and dependencies are `nullptr` + * - Expected output: returned number for dependencies count is equal to 0 + * -# When number of dependencies is less than total lenght + * - Expected output: returned number of dependencies count equal to 1 + * -# When number of dependencies is equal to the total length + * - Expected output: returned number of depencencies is equal to the total length + * Test source + * ------------------------ + * - unit/graph/hipGraphAddEventRecordNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphAddEventRecordNode_Positive_Parameters") { hipGraph_t graph; @@ -336,7 +393,35 @@ TEST_CASE("Unit_hipGraphAddEventRecordNode_Positive_Parameters") { } /** - * Scenario 8: All negative tests + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When graph handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When node dependencies are `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When dependencies are not `nullptr` and the size is not zero + * - Expected output: return `hipErrorInvalidValue` + * -# When node in dependency is from different graph + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidValue` + * -# When number of nodes is not valid (0) + * - Expected output: return `hipErrorInvalidValue` + * -# When duplicate node in dependencies + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidValue` + * -# When node event handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When graph is not initialized + * - Expected output: return `hipErrorInvalidValue` + * -# When event is not initialized + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphAddEventRecordNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphAddEventRecordNode_Negative") { using namespace std::placeholders; diff --git a/catch/unit/graph/hipGraphAddEventWaitNode.cc b/catch/unit/graph/hipGraphAddEventWaitNode.cc index 89e7e03d7..69a1ac65a 100644 --- a/catch/unit/graph/hipGraphAddEventWaitNode.cc +++ b/catch/unit/graph/hipGraphAddEventWaitNode.cc @@ -17,33 +17,6 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : - 1) Simple Scenario: Create an event record node and then create an event -wait node using the same event and add it to graph. Instantiate and Launch -the Graph. Wait for the graph to complete. The operation must succeed without -any failures. - 2) Create a graph 1 with memcpyh2d, event record node (event A), kernel1 -and memcpyd2h nodes. Create a graph 2 with Event Wait (event A) , kernel2 -and memcpyd2h nodes. Instantiate and launch graph1 on stream1 and graph2 on -stream2. Wait for both graph1 and graph2 to complete. Validate the result of -both graphs. - 3) Execute graph1 and graph2 in scenario 2 multiple times in a loop -(100 times). - 4) Execute scenario 2 with stream1 = stream2. - 5) Repeat scenario 2 for different event flags. - 6) Validate that no error is reported when numDeps <= dependencies length - 7) Negative Scenarios - - Pass input node parameter as nullptr. - - Pass input graph parameter as nullptr. - - Pass input dependency parameter as nullptr. - - Node in dependency is from different graph - - Invalid numNodes - - Duplicate node in dependencies - - Pass input event parameter as nullptr. - - Pass uninitialized input graph parameter. - - Pass uninitialized input event parameter. -*/ #include #include @@ -53,7 +26,28 @@ both graphs. #include "graph_tests_common.hh" /** - * Scenario 1 + * @addtogroup hipGraphAddEventWaitNode hipGraphAddEventWaitNode + * @{ + * @ingroup GraphTest + * `hipGraphAddEventWaitNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + * const hipGraphNode_t* pDependencies, size_t numDependencies, hipEvent_t event)` - + * Creates an event wait node and adds it to a graph. + */ + +/** + * Test Description + * ------------------------ + * - Create an event record node. + * - Create an event wait node using the same event and add it to graph. + * - Instantiate and launch the Graph. + * - Wait for the graph to complete. + * - The operation must succeed without any failures. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddEventWaitNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphAddEventWaitNode_Functional_Simple") { hipGraph_t graph; @@ -79,9 +73,7 @@ TEST_CASE("Unit_hipGraphAddEventWaitNode_Functional_Simple") { HIP_CHECK(hipStreamDestroy(streamForGraph)); } -/** - * Local Function - */ +// Local Function static void validate_hipGraphAddEventWaitNode_internodedep(int test, int nstep, unsigned flag = hipEventDefault) { constexpr size_t N = 1024; @@ -214,28 +206,70 @@ static void validate_hipGraphAddEventWaitNode_internodedep(int test, int nstep, } /** - * Scenario 2 + * Test Description + * ------------------------ + * - Create a graph 1 with memcpyh2d, event record node, kernel1 + * and memcpyd2h nodes. + * - Create a graph 2 with Event Wait, kernel2 and memcpyd2h nodes. + * - Instantiate and launch graph1 on stream1 and graph2 on stream2. + * - Wait for both graph1 and graph2 to complete. + * - Validate the result of both graphs. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddEventWaitNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphAddEventWaitNode_MultGraphMultStrmDependency") { validate_hipGraphAddEventWaitNode_internodedep(0, 1); } /** - * Scenario 3 + * Test Description + * ------------------------ + * - Execute graph1 and graph2 in scenario @ref Unit_hipGraphAddEventWaitNode_MultGraphMultStrmDependency + * multiple times in a loop (100 times). + * Test source + * ------------------------ + * - unit/graph/hipGraphAddEventWaitNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphAddEventWaitNode_MultipleRun") { validate_hipGraphAddEventWaitNode_internodedep(0, 100); } /** - * Scenario 4 + * Test Description + * ------------------------ + * - Execute scenario @ref Unit_hipGraphAddEventWaitNode_MultGraphMultStrmDependency + * with stream1 = stream2. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddEventWaitNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphAddEventWaitNode_MultGraphOneStrmDependency") { validate_hipGraphAddEventWaitNode_internodedep(1, 1); } /** - * Scenario 5 + * Test Description + * ------------------------ + * - Repeat scenario @ref Unit_hipGraphAddEventWaitNode_MultGraphMultStrmDependency + * for different event flags. + * -# When flag is `hipEventBlockingSync` + * -# When flag is `hipEventDisableTiming` + * Test source + * ------------------------ + * - unit/graph/hipGraphAddEventWaitNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphAddEventWaitNode_differentFlags") { SECTION("flag = hipEventBlockingSync") { @@ -247,7 +281,21 @@ TEST_CASE("Unit_hipGraphAddEventWaitNode_differentFlags") { } /** - * Scenario 6: Positive parameter tests + * Test Description + * ------------------------ + * - Validate that no error is reported for different scenarios: + * -# When number of dependencies is zero and dependencies are not `nullptr` + * - Expected output: return dependencies number is zero + * -# When number of dependencies is less than total length + * - Expected output: return dependencies number less than total length + * -# When number of dependencies is equal to the total length + * - Expected output: return dependencies number equal to the total length + * Test source + * ------------------------ + * - unit/graph/hipGraphAddEventWaitNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphAddEventWaitNode_Positive_Parameters") { hipGraph_t graph; @@ -286,7 +334,35 @@ TEST_CASE("Unit_hipGraphAddEventWaitNode_Positive_Parameters") { } /** - * Scenario 7 + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When graph handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When node dependencies are `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When dependencies are not `nullptr` and the size is not zero + * - Expected output: return `hipErrorInvalidValue` + * -# When node in dependency is from different graph + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidValue` + * -# When number of nodes is not valid (0) + * - Expected output: return `hipErrorInvalidValue` + * -# When duplicate node in dependencies + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidValue` + * -# When node event handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When graph is not initialized + * - Expected output: return `hipErrorInvalidValue` + * -# When event is not initialized + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphAddEventWaitNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphAddEventWaitNode_Negative") { using namespace std::placeholders; diff --git a/catch/unit/graph/hipGraphAddHostNode.cc b/catch/unit/graph/hipGraphAddHostNode.cc index 971ac187c..b4f64adc5 100644 --- a/catch/unit/graph/hipGraphAddHostNode.cc +++ b/catch/unit/graph/hipGraphAddHostNode.cc @@ -19,27 +19,19 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Test Case Scenarios of hipGraphAddHostNode API: - -Functional: -1. Creates graph, Adds HostNode which updates the variable and validates the result -2. Create graph, Add Graph nodes and clones the graph. Add Host node to the cloned graph - and validate the result -3. Creates graph which performs the square of number in the kernel function and the result - is validated in the callback function of hipGraphAddHostNode API - -Negative: - -1) Pass pGraphNode as nullptr and verify api doesn’t crash, returns error code. -2) Pass graph as nullptr and verify api doesn’t crash, returns error code. -3) Pass pNodeParams as nullptr and verify api doesn’t crash, returns error code. -4) Pass hipHostNodeParams::hipHostFn_t as nullptr and verify api doesn't crash, returns error code. -*/ - #include #include +/** + * @addtogroup hipGraphAddHostNode hipGraphAddHostNode + * @{ + * @ingroup GraphTest + * `hipGraphAddHostNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + * const hipGraphNode_t* pDependencies, size_t numDependencies, + * const hipHostNodeParams* pNodeParams)` - + * Creates a host execution node and adds it to a graph. + */ + #define SIZE 1024 static int* B_h; @@ -73,10 +65,29 @@ static void vectorsquare_callback(void* ptr) { } } -/* -This test case verifies the negative scenarios of -hipGraphAddHostNode API -*/ +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When pointer to the graph node is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When graph handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When number of dependencies is not valid + * - Expected output: return `hipErrorInvalidValue` + * -# When list of dependencies is valid but number of dependencies is not valid + * - Expected output: return `hipErrorInvalidValue` + * -# When pointer to host params is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When host params host function data member is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphAddHostNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphAddHostNode_Negative") { constexpr size_t N = 1024; hipGraph_t graph; @@ -130,11 +141,20 @@ TEST_CASE("Unit_hipGraphAddHostNode_Negative") { HIP_CHECK(hipGraphDestroy(graph)); } -/* -This test case verifies hipGraphAddHostNode API in cloned graph -Creates graph, Add graph nodes and clone the graph -Add HostNode to the cloned graph and validate the result -*/ +/** + * Test Description + * ------------------------ + * - Verifies API behaviour in cloned graph. + * - Creates graph and adds nodes. + * - Clones graph and adds host node to the cloned graph. + * - Validates the result. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddHostNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphAddHostNode_ClonedGraphWithHostNode") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); @@ -198,6 +218,19 @@ This test case verifies the square of number by creating graph, Add kernel node which does the square of number and the result is validated by hipGraphAddHostNode API */ +/** + * Test Description + * ------------------------ + * - Verifies square of number by creating graph. + * - Adds kernel node which does the square of number. + * - Result is validated. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddHostNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphAddHostNode_VectorSquare") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); @@ -256,6 +289,19 @@ Create graph, calls the host function and updates the parameters in the callback function and validates it. */ +/** + * Test Description + * ------------------------ + * - Creates graph and calls the host function. + * - Updates the prameters in the callback function. + * - Validates results. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddHostNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphAddHostNode_BasicFunc") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); diff --git a/catch/unit/graph/hipGraphChildGraphNodeGetGraph.cc b/catch/unit/graph/hipGraphChildGraphNodeGetGraph.cc index 4a975dd3c..4dcf2ba7e 100644 --- a/catch/unit/graph/hipGraphChildGraphNodeGetGraph.cc +++ b/catch/unit/graph/hipGraphChildGraphNodeGetGraph.cc @@ -16,29 +16,32 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios of hipGraphChildGraphNodeGetGraph API: - -Functional Scenarios: -1. Get the child graph node from the original graph and execute it - -Negative Scenarios: -1. Pass nullptr to graph -2. Pass nullptr to graphnode -3. Pass uninitialized graph node -4. Pass orginial graph node instead of child graph node -**/ - #include #include #include -/* -This testcase verifies the following scenario -Create graph, add multiple child nodes and gets the -graph of one of the child nodes using hipGraphChildGraphNodeGetGraph API -executes it and validates the results -*/ +/** + * @addtogroup hipGraphChildGraphNodeGetGraph hipGraphChildGraphNodeGetGraph + * @{ + * @ingroup GraphTest + * `hipGraphChildGraphNodeGetGraph(hipGraphNode_t node, hipGraph_t* pGraph)` - + * Gets a handle to the embedded graph of a child graph node. + */ + +/** + * Test Description + * ------------------------ + * - Creates the graph. + * - Adds multiple child nodes. + * - Gets the graph of one of the child nodes. + * - Executes it and validates the results. + * Test source + * ------------------------ + * - unit/graph/hipGraphChildGraphNodeGetGraph.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphChildGraphNodeGetGraph_Functional") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); @@ -118,6 +121,25 @@ TEST_CASE("Unit_hipGraphChildGraphNodeGetGraph_Functional") { This testcase verifies the negative scenarios of hipGraphChildGraphNodeGetGraph API */ +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When child node handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When output pointer to the graph is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When passing parent instead of child graph node + * - Expected output: return `hipErrorInvalidValue` + * -# When child node is not initialized + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphChildGraphNodeGetGraph.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphChildGraphNodeGetGraph_Negative") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); diff --git a/catch/unit/graph/hipGraphClone.cc b/catch/unit/graph/hipGraphClone.cc index 75b83c82a..fac0bb6d4 100644 --- a/catch/unit/graph/hipGraphClone.cc +++ b/catch/unit/graph/hipGraphClone.cc @@ -17,32 +17,35 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/* -Testcase Scenarios of hipGraphClone API: - -Negative: - -1. Pass nullptr to cloned graph -2. pass nullptr to original graph - -Functional: - -1. Clone the graph,Instantiate and execute the cloned graph -2. Clone the graph and modify the original graph and ensure that the - cloned graph is not modified -3. Create graph on one GPU device and clone it from peer GPU device -4. Create graph in one thread and clone it from multiple threads. -*/ - #include #include #include -#define NUM_THREADS 10 +/** + * @addtogroup hipGraphClone hipGraphClone + * @{ + * @ingroup GraphTest + * `hipGraphClone(hipGraph_t* pGraphClone, hipGraph_t originalGraph)` - + * Clones a graph. + */ -/* This test covers the negative scenarios of - hipGraphClone API */ +#define NUM_THREADS 10 +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When output pointer to the cloned graph is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When original graph handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphClone.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphClone_Negative") { SECTION("Passing nullptr to Cloned graph") { hipGraph_t graph; @@ -56,10 +59,11 @@ TEST_CASE("Unit_hipGraphClone_Negative") { REQUIRE(hipGraphClone(&clonedGraph, nullptr) == hipErrorInvalidValue); } } + /* -This function creates the graph with dependencies -then performs device context change and clones the cloned graph -Executes the cloned graph and validates the result + This function creates the graph with dependencies + then performs device context change and clones the cloned graph + Executes the cloned graph and validates the result */ void hipGraphClone_DeviceContextChange() { constexpr size_t N = 1024; @@ -97,13 +101,14 @@ void hipGraphClone_DeviceContextChange() { HIP_CHECK(hipGraphDestroy(clonedgraph)); HIP_CHECK(hipStreamDestroy(streamForGraph)); } + /* -This function does the following -1. Creates the graph with multiple dependencies - clones the graph and validates the result. -2. Creates the graph, clones the graph and modifies - the existing graph and execute the cloned graph - to ensure that cloned graph is not modified + This function does the following + 1. Creates the graph with multiple dependencies + clones the graph and validates the result. + 2. Creates the graph, clones the graph and modifies + the existing graph and execute the cloned graph + to ensure that cloned graph is not modified */ void hipGraphClone_Func(bool ModifyOrigGraph = false) { constexpr size_t N = 1024; @@ -232,13 +237,20 @@ void hipGraphClone_Func(bool ModifyOrigGraph = false) { HIP_CHECK(hipStreamDestroy(streamForGraph)); } -/* -This testcase verifies following scenarios -1. Clones the graph and verify the result -2. Clones the graph, Modify the original graph and - validate the result of the cloned graph -3. Device context change for cloned graph -*/ +/** + * Test Description + * ------------------------ + * - Clones the graph and verifies the result. + * - Clones the graph, modifies the original graph and validates + * the result of the cloned graph. + * - Changes device context for cloned graph. + * Test source + * ------------------------ + * - unit/graph/hipGraphClone.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphClone_Functional") { SECTION("hipGraphClone Basic Functionality") { hipGraphClone_Func(); @@ -264,13 +276,23 @@ TEST_CASE("Unit_hipGraphClone_Functional") { } } -/* -This testcase creates the graph with dependencies -then creates multiple threads and clones the graph -in each thread and executes the cloned graph -hipGraphClone is failing in CUDA in multi threaded -scenario so excluded for nvidia -*/ +/** + * Test Description + * ------------------------ + * - Creates the graph with dependencies. + * - Creates multiple threads. + * - Clones the graph in eahc thread. + * - Executes the cloned graph. + * - This scenario is failing in CUDA. + * Test source + * ------------------------ + * - unit/graph/hipGraphClone.cc + * Test requirements + * ------------------------ + * - Multi-threaded + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 + */ #if HT_AMD TEST_CASE("Unit_hipGraphClone_MultiThreaded") { constexpr size_t N = 1024; diff --git a/catch/unit/graph/hipGraphDestroyNode.cc b/catch/unit/graph/hipGraphDestroyNode.cc index a6a028e60..bd82890d0 100644 --- a/catch/unit/graph/hipGraphDestroyNode.cc +++ b/catch/unit/graph/hipGraphDestroyNode.cc @@ -17,59 +17,56 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/* -Testcase Scenarios of hipGraphDestroyNode API: - -Negative :: - -1) Pass nullptr to graph node - -Functional :: - -1) Create Node and destroy the node -2) Create graph with dependencies and destroy one of the dependency node - before executing the graph. -3) Create a graph with N nodes and (N-1) dependencies between them as shown - below. Start destroying the nodes in iteration from left. In each iteration - verify the number of nodes and dependencies using hipGraphGetNodes and - hipGraphGetEdges. - Node1-->Node2-->Node3->...................->NodeN -4) Create a graph with N nodes and (N-1) dependencies between them as shown - above. Clone the graph. Start destroying the nodes in iteration from left - in the cloned graph. In each iteration verify the number of nodes and - dependencies using hipGraphGetNodes and hipGraphGetEdges. Once all nodes - in the cloned graph are deleted, verify the number of nodes in the original - graph are intact. -5) Create a graph1 with N nodes and (N-1) dependencies between them as shown - above. Create another empty graph0. Add graph1 as child node to graph0. - Delete the child node in graph0. Verify that the nodes in graph1 are still - intact after deleting the child node using hipGraphGetNodes and hipGraphGetEdges. -*/ - #include #include #include #define NUM_OF_DUMMY_NODES 8 -static __global__ void dummyKernel() { - return; -} +static __global__ void dummyKernel() { return; } + +/** + * @addtogroup hipGraphDestroyNode hipGraphDestroyNode + * @{ + * @ingroup GraphTest + * `hipGraphDestroyNode(hipGraphNode_t node)` - + * Remove a node from the graph. + */ -/* This test covers the negative scenarios of - hipGraphDestroyNode API */ +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When graph handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphDestroyNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphDestroyNode_Negative") { SECTION("Passing nullptr to graph Node") { REQUIRE(hipGraphDestroyNode(nullptr) == hipErrorInvalidValue); } } -/* This test covers the basic functionality of - hipGraphDestroyNode API where we create and destroy - the node -*/ +/** + * Test Description + * ------------------------ + * - Creates an empty graph. + * - Checks that it is not `nullptr`. + * - Destroys the graph successfully. + * Test source + * ------------------------ + * - unit/graph/hipGraphDestroyNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphDestroyNode_BasicFunctionality") { - char *pOutBuff_d{}; + char* pOutBuff_d{}; constexpr size_t size = 1024; hipGraph_t graph{}; hipGraphNode_t memsetNode{}; @@ -83,17 +80,23 @@ TEST_CASE("Unit_hipGraphDestroyNode_BasicFunctionality") { memsetParams.width = size * sizeof(char); memsetParams.height = 1; HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, - &memsetParams)); + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams)); REQUIRE(hipGraphDestroyNode(memsetNode) == hipSuccess); HIP_CHECK(hipFree(pOutBuff_d)); } -/* -This testcase verifies the following scenario where -graph is created with dependencies and one of the dependency is -destroyed before execute the graph -*/ +/** + * Test Description + * ------------------------ + * - Create graph with dependencies and destroy one of the dependency node before executing the + * graph. + * Test source + * ------------------------ + * - unit/graph/hipGraphDestroyNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphDestroyNode_DestroyDependencyNode") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); @@ -113,32 +116,29 @@ TEST_CASE("Unit_hipGraphDestroyNode_DestroyDependencyNode") { HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B2Copies, graph, nullptr, - 0, B_d, C_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B2Copies, graph, nullptr, 0, B_d, C_h, Nbytes, + hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, B_h, C_d, - Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, B_h, C_d, Nbytes, + hipMemcpyDeviceToHost)); - void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; - kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); kernelNodeParams.gridDim = dim3(blocks); kernelNodeParams.blockDim = dim3(threadsPerBlock); kernelNodeParams.sharedMemBytes = 0; kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); kernelNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, - &kernelNodeParams)); + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, &kernelNodeParams)); // Create dependencies HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1)); HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1)); - HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B2Copies, &kernel_vecAdd, - 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B2Copies, &kernel_vecAdd, 1)); HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1)); // Destroy one of the dependency node @@ -159,8 +159,17 @@ TEST_CASE("Unit_hipGraphDestroyNode_DestroyDependencyNode") { } /** - * Functional Test to test hipGraphDestroyNode using hipGraphGetNodes - * and hipGraphGetEdges APIs. + * Test Description + * ------------------------ + * - Create a graph with N nodes and (N-1) dependencies between them as shown below. Start + * destroying the nodes in iteration from left. In each iteration verify the number of nodes and + * dependencies using hipGraphGetNodes and hipGraphGetEdges. + * Test source + * ------------------------ + * - unit/graph/hipGraphDestroyNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep") { hipGraph_t graph; @@ -170,19 +179,17 @@ TEST_CASE("Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep") { // Create graph with no dependencies for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) { void* kernelArgs[] = {nullptr}; - kernelNodeParams[i].func = reinterpret_cast(dummyKernel); + kernelNodeParams[i].func = reinterpret_cast(dummyKernel); kernelNodeParams[i].gridDim = dim3(1); kernelNodeParams[i].blockDim = dim3(1); kernelNodeParams[i].sharedMemBytes = 0; kernelNodeParams[i].kernelParams = reinterpret_cast(kernelArgs); kernelNodeParams[i].extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph, nullptr, - 0, &kernelNodeParams[i])); + HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph, nullptr, 0, &kernelNodeParams[i])); } // Create dependencies between nodes for (int i = 1; i < NUM_OF_DUMMY_NODES; i++) { - HIP_CHECK(hipGraphAddDependencies(graph, &kernelnode[i-1], - &kernelnode[i], 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernelnode[i - 1], &kernelnode[i], 1)); } // Start destroying nodes from 0 size_t numOfNodes = 0, numOfDep = 0; @@ -194,15 +201,26 @@ TEST_CASE("Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep") { HIP_CHECK(hipGraphGetEdges(graph, nullptr, nullptr, &numOfDep)); REQUIRE(numOfDep == (NUM_OF_DUMMY_NODES - i - 2)); } - HIP_CHECK(hipGraphDestroyNode(kernelnode[NUM_OF_DUMMY_NODES-1])); + HIP_CHECK(hipGraphDestroyNode(kernelnode[NUM_OF_DUMMY_NODES - 1])); HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numOfNodes)); REQUIRE(numOfNodes == 0); HIP_CHECK(hipGraphDestroy(graph)); } /** - * Functional Test to test hipGraphDestroyNode using hipGraphGetNodes - * and hipGraphGetEdges APIs on a cloned graph + * Test Description + * ------------------------ + * - Create a graph with N nodes and (N-1) dependencies between them as shown above. Clone the + * graph. Start destroying the nodes in iteration from left in the cloned graph. In each iteration + * verify the number of nodes and dependencies using hipGraphGetNodes and hipGraphGetEdges. Once all + * nodes in the cloned graph are deleted, verify the number of nodes in the original graph are + * intact. + * Test source + * ------------------------ + * - unit/graph/hipGraphDestroyNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ClonedGrph") { hipGraph_t graph, clonedgraph; @@ -213,19 +231,17 @@ TEST_CASE("Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ClonedGrph") { // Create graph with no dependencies for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) { void* kernelArgs[] = {nullptr}; - kernelNodeParams[i].func = reinterpret_cast(dummyKernel); + kernelNodeParams[i].func = reinterpret_cast(dummyKernel); kernelNodeParams[i].gridDim = dim3(1); kernelNodeParams[i].blockDim = dim3(1); kernelNodeParams[i].sharedMemBytes = 0; kernelNodeParams[i].kernelParams = reinterpret_cast(kernelArgs); kernelNodeParams[i].extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph, nullptr, - 0, &kernelNodeParams[i])); + HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph, nullptr, 0, &kernelNodeParams[i])); } // Create dependencies between nodes for (int i = 1; i < NUM_OF_DUMMY_NODES; i++) { - HIP_CHECK(hipGraphAddDependencies(graph, &kernelnode[i-1], - &kernelnode[i], 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernelnode[i - 1], &kernelnode[i], 1)); } HIP_CHECK(hipGraphClone(&clonedgraph, graph)); // Start destroying nodes from 0 and validate number of nodes in @@ -250,8 +266,18 @@ TEST_CASE("Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ClonedGrph") { } /** - * Functional Test to test hipGraphDestroyNode on child node using - * hipGraphGetNodes and hipGraphGetEdges APIs on a cloned graph. + * Test Description + * ------------------------ + * - Create a graph1 with N nodes and (N-1) dependencies between them as shown above. Create + * another empty graph0. Add graph1 as child node to graph0. Delete the child node in graph0. Verify + * that the nodes in graph1 are still intact after deleting the child node using hipGraphGetNodes + * and hipGraphGetEdges. + * Test source + * ------------------------ + * - unit/graph/hipGraphDestroyNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ChldNode") { hipGraph_t graph0, graph1; @@ -262,23 +288,20 @@ TEST_CASE("Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ChldNode") { // Create graph with no dependencies for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) { void* kernelArgs[] = {nullptr}; - kernelNodeParams[i].func = reinterpret_cast(dummyKernel); + kernelNodeParams[i].func = reinterpret_cast(dummyKernel); kernelNodeParams[i].gridDim = dim3(1); kernelNodeParams[i].blockDim = dim3(1); kernelNodeParams[i].sharedMemBytes = 0; kernelNodeParams[i].kernelParams = reinterpret_cast(kernelArgs); kernelNodeParams[i].extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph0, nullptr, - 0, &kernelNodeParams[i])); + HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph0, nullptr, 0, &kernelNodeParams[i])); } // Create dependencies between nodes for (int i = 1; i < NUM_OF_DUMMY_NODES; i++) { - HIP_CHECK(hipGraphAddDependencies(graph0, &kernelnode[i-1], - &kernelnode[i], 1)); + HIP_CHECK(hipGraphAddDependencies(graph0, &kernelnode[i - 1], &kernelnode[i], 1)); } // Create child node and add it to graph1 - HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph1, - nullptr, 0, graph0)); + HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph1, nullptr, 0, graph0)); // delete the child node from graph1 HIP_CHECK(hipGraphDestroyNode(childGraphNode)); // Start destroying nodes from 0 diff --git a/catch/unit/graph/hipGraphEventRecordNodeGetEvent.cc b/catch/unit/graph/hipGraphEventRecordNodeGetEvent.cc index 49fab0012..8957b4796 100644 --- a/catch/unit/graph/hipGraphEventRecordNodeGetEvent.cc +++ b/catch/unit/graph/hipGraphEventRecordNodeGetEvent.cc @@ -17,26 +17,19 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : - 1) Validate that the event returned by hipGraphEventRecordNodeGetEvent matches -with the event set in hipGraphAddEventRecordNode. - 2) Negative Scenarios - - Input node is a nullptr. - - Output event is a nullptr. - - Input node is an empty node. - - Input node is a memset node. - - Input node is event wait node - - Input node is an uninitialized node. -*/ - #include #include #include /** - * Local Function to set and get event record node property. + * @addtogroup hipGraphEventRecordNodeGetEvent hipGraphEventRecordNodeGetEvent + * @{ + * @ingroup GraphTest + * `hipGraphEventRecordNodeGetEvent(hipGraphNode_t node, hipEvent_t* event_out)` - + * Returns the event associated with an event record node. */ + +// Local Function to set and get event record node property. static void validateEventRecordNodeGetEvent(unsigned flag) { hipGraph_t graph; HIP_CHECK(hipGraphCreate(&graph, 0)); @@ -53,7 +46,18 @@ static void validateEventRecordNodeGetEvent(unsigned flag) { } /** - * Scenario: Validate scenario 1 for different event flags. + * Test Description + * ------------------------ + * - Validate that the event returned matches with the event that is previously set. + * -# When flag is `hipEventDefault` + * -# When flag is `hipEventBlockingSync` + * -# When flag is `hipEventDisableTiming` + * Test source + * ------------------------ + * - unit/graph/hipGraphEventRecordNodeGetEvent.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphEventRecordNodeGetEvent_Functional") { // Create event nodes with different flags and validate with @@ -72,7 +76,27 @@ TEST_CASE("Unit_hipGraphEventRecordNodeGetEvent_Functional") { } /** - * Scenario 2: Negative tests. + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When node handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When output pointer to the event is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When input node is empty node + * - Expected output: return `hipErrorInvalidValue` + * -# When input node is memset node + * - Expected output: return `hipErrorInvalidValue` + * -# When input node is a wait node + * - Expected output: return `hipErrorInvalidValue` + * -# When input node is not initialized node + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphEventRecordNodeGetEvent.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphEventRecordNodeGetEvent_Negative") { hipGraph_t graph; diff --git a/catch/unit/graph/hipGraphEventRecordNodeSetEvent.cc b/catch/unit/graph/hipGraphEventRecordNodeSetEvent.cc index 93bb0418f..0cebb676f 100644 --- a/catch/unit/graph/hipGraphEventRecordNodeSetEvent.cc +++ b/catch/unit/graph/hipGraphEventRecordNodeSetEvent.cc @@ -17,32 +17,12 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : - 1) Set a different type of event using hipGraphEventRecordNodeSetEvent and - validate using hipGraphEventRecordNodeGetEvent. - 2) Add different kinds of nodes to graph and add dependencies to nodes. - Create an event record node at the end with Default flag. Set a different - type of event using hipGraphEventRecordNodeSetEvent. Instantiate and - Launch graph. Wait for the event to complete. Verify the results. - 3) Negative Scenarios - - Input node parameter is nullptr. - - Input event parameter is nullptr. - - Empty node is passed as input node. - - Memset node is passed as input node. - - Event wait node is passed as input node. - - Input node is an uninitialized node. - - Input event is an uninitialized event. -*/ - #include #include #include -/** - * Local Function: Set Get test - */ +// Local Function: Set Get test static void validateEventRecordNodeSetEvent(unsigned flag) { hipGraph_t graph; HIP_CHECK(hipGraphCreate(&graph, 0)); @@ -63,9 +43,7 @@ static void validateEventRecordNodeSetEvent(unsigned flag) { HIP_CHECK(hipEventDestroy(event2)); } -/** - * Local Function - */ +// Local Function static void setEventWaitNode() { hipGraph_t graph; HIP_CHECK(hipGraphCreate(&graph, 0)); @@ -84,7 +62,20 @@ static void setEventWaitNode() { } /** - * Scenario 2: Validate Change of event property in event record node. + * Test Description + * ------------------------ + * - Add different kinds of nodes to graph and add dependencies to nodes. + * - Create an event record node at the end with Default flag. + * - Set a different type of event. + * - Instantiate and launch graph. + * - Wait for the event to complete. + * - Verify the results. + * Test source + * ------------------------ + * - unit/graph/hipGraphEventRecordNodeSetEvent.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphEventRecordNodeSetEvent_SetEventProperty") { hipGraph_t graph; @@ -153,7 +144,18 @@ TEST_CASE("Unit_hipGraphEventRecordNodeSetEvent_SetEventProperty") { } /** - * Scenario 1: Validate Set Get test for all Event flags + * Test Description + * ------------------------ + * - Set a different type of event and validate using get. + * -# When flag is `hipEventDefault` + * -# When flag is `hipEventBlockingSync` + * -# When flag is `hipEventDisableTiming` + * Test source + * ------------------------ + * - unit/graph/hipGraphEventRecordNodeSetEvent.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphEventRecordNodeSetEvent_SetGet") { SECTION("Flag = hipEventDefault") { @@ -170,7 +172,29 @@ TEST_CASE("Unit_hipGraphEventRecordNodeSetEvent_SetGet") { } /** - * Scenario 3: Negative Tests + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When node handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When event handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When input node is empty node + * - Expected output: return `hipErrorInvalidValue` + * -# When input node is memset node + * - Expected output: return `hipErrorInvalidValue` + * -# When input node is wait node + * - Expected output: return `hipErrorInvalidValue` + * -# When input node is not initialized + * - Expected output: return `hipErrorInvalidValue` + * -# When event is not initialized + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphEventRecordNodeSetEvent.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphEventRecordNodeSetEvent_Negative") { hipGraph_t graph; diff --git a/catch/unit/graph/hipGraphEventWaitNodeGetEvent.cc b/catch/unit/graph/hipGraphEventWaitNodeGetEvent.cc index 74bbf87c8..a9bf65ef3 100644 --- a/catch/unit/graph/hipGraphEventWaitNodeGetEvent.cc +++ b/catch/unit/graph/hipGraphEventWaitNodeGetEvent.cc @@ -17,27 +17,20 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : - 1) Validate that the event returned by hipGraphEventWaitNodeGetEvent matches -with the event set in hipGraphAddEventWaitNode. - 2) Negative Scenarios - - Input node parameter is passed as nullptr. - - Output event parameter is passed as nullptr. - - Input node parameter is an empty node. - - Input node parameter is a memset node. - - Input node parameter is a event record node. - - Input node parameter is an uninitialized node. -*/ - #include #include #include /** - * Local Function + * @addtogroup hipGraphEventWaitNodeGetEvent hipGraphEventWaitNodeGetEvent + * @{ + * @ingroup GraphTest + * `hipGraphEventWaitNodeGetEvent(hipGraphNode_t node, hipEvent_t* event_out)` - + * Returns the event associated with an event wait node. */ + +// Local Function static void validateEventWaitNodeGetEvent(unsigned flag) { hipGraph_t graph; HIP_CHECK(hipGraphCreate(&graph, 0)); @@ -54,7 +47,18 @@ static void validateEventWaitNodeGetEvent(unsigned flag) { } /** - * Scenario 1 + * Test Description + * ------------------------ + * - Validate that the event returned matches with the event that is set previously. + * -# When flag is `hipEventDefault` + * -# When flag is `hipEventBlockingSync` + * -# When flag is `hipEventDisableTiming` + * Test source + * ------------------------ + * - unit/graph/hipGraphEventWaitNodeGetEvent.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphEventWaitNodeGetEvent_Functional") { // Create event nodes with different flags and validate with @@ -73,7 +77,27 @@ TEST_CASE("Unit_hipGraphEventWaitNodeGetEvent_Functional") { } /** - * Scenario 2 + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When node handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When output pointer to the event is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When input node is empty node + * - Expected output: return `hipErrorInvalidValue` + * -# When input node is memset node + * - Expected output: return `hipErrorInvalidValue` + * -# When input node is a record node + * - Expected output: return `hipErrorInvalidValue` + * -# When input node is not initialized node + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphEventWaitNodeGetEvent.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphEventWaitNodeGetEvent_Negative") { hipGraph_t graph; diff --git a/catch/unit/graph/hipGraphEventWaitNodeSetEvent.cc b/catch/unit/graph/hipGraphEventWaitNodeSetEvent.cc index 6c3132262..d6fc9a74c 100644 --- a/catch/unit/graph/hipGraphEventWaitNodeSetEvent.cc +++ b/catch/unit/graph/hipGraphEventWaitNodeSetEvent.cc @@ -17,34 +17,19 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : - 1) Set a different type of event using hipGraphEventWaitNodeSetEvent and - validate using hipGraphEventWaitNodeGetEvent. - 2) Create a graph1 with memset (Value 1) node, event record node (event A) - and memset (Value 2) node and event record node (event B). Create a - graph2 with Event Wait (event A ) node and memcpyd2h. Instantiate - graph1 on stream1 and graph2 on stream2. Set the Event Wait node event - to B using hipGraphEventWaitNodeSetEvent. Launch graphs. Wait for the - event to complete. Verify the results. - 3) Negative Scenarios - - Input node parameter is passed as nullptr. - - Input event parameter is passed as nullptr. - - Input node is an empty node. - - Input node is a memset node. - - Input node is a event record node. - - Input node is an uninitialized node. - - Input event is an uninitialized node. -*/ - #include #include #include - /** - * Local Function + * @addtogroup hipGraphEventWaitNodeSetEvent hipGraphEventWaitNodeSetEvent + * @{ + * @ingroup GraphTest + * `hipGraphEventWaitNodeSetEvent(hipGraphNode_t node, hipEvent_t event)` - + * Sets an event wait node's event. */ + +// Local Function static void validateEventWaitNodeSetEvent(unsigned flag) { hipGraph_t graph; HIP_CHECK(hipGraphCreate(&graph, 0)); @@ -65,9 +50,7 @@ static void validateEventWaitNodeSetEvent(unsigned flag) { HIP_CHECK(hipEventDestroy(event2)); } -/** - * Local Function - */ +// Local function static void setEventRecordNode() { hipGraph_t graph; HIP_CHECK(hipGraphCreate(&graph, 0)); @@ -86,7 +69,21 @@ static void setEventRecordNode() { } /** - * Scenario 2 + * Test Description + * ------------------------ + * - Create a graph1 with memset node, event record node + * and memset node and event record node. + * - Create a graph2 with Event Wait node and memcpyd2h. + * - Instantiate graph1 on stream1 and graph2 on stream2. + * - Set the Event Wait node event to graph2. + * - Launch graphs and wait for the event to complete. + * - Verify the results. + * Test source + * ------------------------ + * - unit/graph/hipGraphEventWaitNodeSetEvent.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphEventWaitNodeSetEvent_SetProp") { constexpr size_t N = 512; @@ -215,8 +212,20 @@ TEST_CASE("Unit_hipGraphEventWaitNodeSetEvent_SetProp") { HIP_CHECK(hipStreamDestroy(streamForGraph1)); HIP_CHECK(hipStreamDestroy(streamForGraph2)); } + /** - * Scenario 1 + * Test Description + * ------------------------ + * - Validate that the event returned matches with the event that is set previously. + * -# When flag is `hipEventDefault` + * -# When flag is `hipEventBlockingSync` + * -# When flag is `hipEventDisableTiming` + * Test source + * ------------------------ + * - unit/graph/hipGraphEventWaitNodeSetEvent.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphEventWaitNodeSetEvent_SetGet") { SECTION("Flag = hipEventDefault") { @@ -233,7 +242,29 @@ TEST_CASE("Unit_hipGraphEventWaitNodeSetEvent_SetGet") { } /** - * Scenario 3 + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When node handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When output pointer to the event is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When input node is empty node + * - Expected output: return `hipErrorInvalidValue` + * -# When input node is memset node + * - Expected output: return `hipErrorInvalidValue` + * -# When input node is a record node + * - Expected output: return `hipErrorInvalidValue` + * -# When input node is not initialized node + * - Expected output: return `hipErrorInvalidValue` + * -# When event is not initialized + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphEventWaitNodeSetEvent.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphEventWaitNodeSetEvent_Negative") { hipGraph_t graph; diff --git a/catch/unit/graph/hipGraphExecChildGraphNodeSetParams.cc b/catch/unit/graph/hipGraphExecChildGraphNodeSetParams.cc index ab7d3ac9a..cc0489cbf 100644 --- a/catch/unit/graph/hipGraphExecChildGraphNodeSetParams.cc +++ b/catch/unit/graph/hipGraphExecChildGraphNodeSetParams.cc @@ -16,30 +16,40 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** - Testcase Scenarios of hipGraphExecChildGraphNodeSetParams API: - - Functional Scenarios: - 1. Create child graph, Instantiate the graph and update the child graph - using hipGraphExecChildGraphNodeSetParams API - 2. Create child graph with topology, Instantiate the graph - and update the child graph - using hipGraphExecChildGraphNodeSetParams API - - Negative Scenarios: - 1. Pass nullptr to child graph - 2. Pass nullptr to graphnode - 3. Pass nullptr to graphExec - 4. Pass uninitialized graph node - 5. Pass orginial graph node instead of child graph node - 6. Change topology of child graph node - **/ - #include #include #include +/** + * @addtogroup hipGraphExecChildGraphNodeSetParams hipGraphExecChildGraphNodeSetParams + * @{ + * @ingroup GraphTest + * `hipGraphExecChildGraphNodeSetParams(hipGraphExec_t hGraphExec, + * hipGraphNode_t node, hipGraph_t childGraph)` - + * Updates node parameters in the child graph node in the given graphExec. + */ +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When executable graph handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When child graph node handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When child graph handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When parent graph is passed instead of child graph + * - Expected output: do not return `hipSuccess` + * -# When updating the child graph topology + * - Expected output: do not return `hipSuccess` + * Test source + * ------------------------ + * - unit/graph/hipGraphExecChildGraphNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphExecChildGraphNodeSetParams_Negative") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); @@ -130,13 +140,21 @@ TEST_CASE("Unit_hipGraphExecChildGraphNodeSetParams_Negative") { HIP_CHECK(hipStreamDestroy(streamForGraph)); } -/* - This testcase verifies the following scenario - Create graph, add child node to graph, Instantiate the graph - and update the child graph node with a new graph - using hipGraphExecChildGraphNodeSetParams API - and execute it - */ +/** + * Test Description + * ------------------------ + * - Creates a graph. + * - Adds the child node to the graph. + * - Instantiates the graph. + * - Updates the child graph node with a new graph. + * - Executes the graph. + * Test source + * ------------------------ + * - unit/graph/hipGraphExecChildGraphNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphExecChildGraphNodeSetParams_BasicFunc") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); @@ -208,14 +226,22 @@ TEST_CASE("Unit_hipGraphExecChildGraphNodeSetParams_BasicFunc") { HIP_CHECK(hipStreamDestroy(streamForGraph)); } -/* - This testcase verifies the following scenario - Create graph, Create child graph with a topology and - add child node to graph, Instantiate the graph - and update the child graph node with a new graph - using hipGraphExecChildGraphNodeSetParams API - and execute it - */ +/** + * Test Description + * ------------------------ + * - Creates a graph. + * - Creates child graph with a topology. + * - Adds child node to graph. + * - Instantiate the graph. + * - Updates the child graph node with a new graph. + * - Executes the graph. + * Test source + * ------------------------ + * - unit/graph/hipGraphExecChildGraphNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphExecChildGraphNodeSetParams_ChildTopology") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); diff --git a/catch/unit/graph/hipGraphExecEventRecordNodeSetEvent.cc b/catch/unit/graph/hipGraphExecEventRecordNodeSetEvent.cc index e820dc1cd..35081e084 100644 --- a/catch/unit/graph/hipGraphExecEventRecordNodeSetEvent.cc +++ b/catch/unit/graph/hipGraphExecEventRecordNodeSetEvent.cc @@ -17,53 +17,43 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : - 1) Create a graph with event record nodes as follows: - event_record_start_node(event1) --> MemcpyH2DNode --> kernel --> - MemcpyD2HNode --> event_record_stop_node(event2).Instantiate the graph. - Set a different event 'event3' in event_record_stop_node using - hipGraphExecEventRecordNodeSetEvent. Launch the graph. Verify the - hipGraphExecEventRecordNodeSetEvent functionality by measuring the - time difference between event2 & event1 and between event3 and - event1. - 2) Scenario to verify that hipGraphExecEventRecordNodeSetEvent does not - impact the graph and changes only the executable graph. - Create an event record node with event1 and add it to graph. Instantiate - the graph to create an executable graph. Change the event in the - executable graph to event2. Verify that the event record node still - contains event1. - 3) Scenario to verify that hipGraphExecEventRecordNodeSetEvent can set event - created on different device. Create an event record node with event1 and add it to graph. - Instantiate the graph to create an executable graph. Call the API to change the event in the - executable graph to event2 which has been created on different device. Verify that graph can be - launched and no error is reported. - 4) Negative Scenarios - - Input executable graph is a nullptr. - - Input node is a nullptr. - - Input event to set is a nullptr. - - Input executable graph is uninitialized. - - Input node is uninitialized. - - Input event is uninitialized. - - Event record node does not exist in graph. - - Input node is a memset node. - - Input node is a event wait node. -*/ - #include #include #include /** - * Kernel Functions to copy. + * @addtogroup hipGraphExecEventRecordNodeSetEvent hipGraphExecEventRecordNodeSetEvent + * @{ + * @ingroup GraphTest + * `hipGraphExecEventRecordNodeSetEvent(hipGraphExec_t hGraphExec, + * hipGraphNode_t hNode, hipEvent_t event)` - + * Sets the event for an event record node in the given graphExec. */ + +// Kernel Functions to copy. static __global__ void copy_ker_func(int* a, int* b, size_t N) { int tx = blockIdx.x * blockDim.x + threadIdx.x; if (tx < N) b[tx] = a[tx]; } /** - * Scenario 1: Functional scenario (See description Above) + * Test Description + * ------------------------ + * - Create a graph with event record nodes as follows: + * -# Record event on start node. + * -# Launch kernel. + * -# Record event on stop node. + * - Instantiate the graph. + * - Set a different event for stop node. + * - Launch the graph. + * - Verify the set functionality by measuring the time difference + * between start and stop event. + * Test source + * ------------------------ + * - unit/graph/hipGraphExecEventRecordNodeSetEvent.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_Functional") { constexpr size_t gridSize = 512; @@ -159,8 +149,20 @@ TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_Functional") { } /** - * Scenario 2: This test verifies that changes to executable graph does - * not impact the original graph. + * Test Description + * ------------------------ + * - Scenario to verify that event set does not + * impact the graph and changes only the executable graph. + * - Create an event record node with event and add it to graph. + * - Instantiate the graph to create an executable graph. + * - Change the event in the executable graph to the new event. + * - Verify that the event record node still contains first event. + * Test source + * ------------------------ + * - unit/graph/hipGraphExecEventRecordNodeSetEvent.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_VerifyEventNotChanged") { hipGraph_t graph; @@ -184,8 +186,22 @@ TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_VerifyEventNotChanged") { } /** - * Scenario 3: This test verifies event in node of the executable graph can be changed to event on - * different device + * Test Description + * ------------------------ + * - Scenario to verify that hipGraphExecEventRecordNodeSetEvent can set event created on different + * device. + * - Create an event record node with event1 and add it to graph. + * - Instantiate the graph to create an executable graph. + * - Call the API to change the event in the executable graph to the event which has been created + * on different device. + * - Verify that graph can be launched and no error is reported. + * Test source + * ------------------------ + * - unit/graph/hipGraphExecEventRecordNodeSetEvent.cc + * Test requirements + * ------------------------ + * - Multi-device + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_Positive_DifferentDevices") { const auto device_count = HipTest::getDeviceCount(); @@ -225,7 +241,33 @@ TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_Positive_DifferentDevices") } /** - * Scenario 4: Negative Parameter Tests + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When executable graph handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When node handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When pointer to the event is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When executable graph is not initialized + * - Expected output: return `hipErrorInvalidValue` + * -# When node is not initialized + * - Expected output: return `hipErrorInvalidValue` + * -# When event is not initialized + * - Expected output: return `hipErrorInvalidValue` + * -# When event record node does not exist + * - Expected output: return `hipErrorInvalidValue` + * -# When node is memset node + * - Expected output: return `hipErrorInvalidValue` + * -# When node is event wait node + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphExecEventRecordNodeSetEvent.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_Negative") { hipGraph_t graph; diff --git a/catch/unit/graph/hipGraphExecEventWaitNodeSetEvent.cc b/catch/unit/graph/hipGraphExecEventWaitNodeSetEvent.cc index 529f8df3c..120521ec0 100644 --- a/catch/unit/graph/hipGraphExecEventWaitNodeSetEvent.cc +++ b/catch/unit/graph/hipGraphExecEventWaitNodeSetEvent.cc @@ -17,43 +17,20 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : - 1) Create a graph1 with event record nodes as follows: - MemcpyH2DNode --> kernel1(x*x) --> event_record_node(event1) - Instantiate graph1. - Create a graph2 with event record nodes as follows: - event_wait_node(event1) --> MemcpyD2HNode.Instantiate graph2. - Change the event in event_record_node in graph1 to event2 using - hipGraphExecEventRecordNodeSetEvent. - Change the event in event_wait_node in graph2 to event2 using - hipGraphExecEventWaitNodeSetEvent.Execute graph1 and then graph2. - Verify the result matches with x*x. - 2) Scenario to verify that hipGraphExecEventWaitNodeSetEvent does not - impact the graph and changes only the executable graph. - Create an event wait node with event1 and add it to graph. Instantiate - the graph to create an executable graph. Change the event in the - executable graph to event2. Verify that the event wait node still - contains event1. - 3) Negative Scenarios - - Input executable graph is nullptr. - - Input node is nullptr. - - Input set event is nullptr. - - Input executable graph is uninitialized. - - Input node is uninitialized. - - Input set event is uninitialized. - - Graph does not contain event wait node. - - Pass memset node as input node. - - Pass event record node as input node. -*/ - #include #include #include /** - * Kernel Functions to perform square and introduce delay in device. + * @addtogroup hipGraphExecEventWaitNodeSetEvent hipGraphExecEventWaitNodeSetEvent + * @{ + * @ingroup GraphTest + * `hipGraphExecEventWaitNodeSetEvent(hipGraphExec_t hGraphExec, + * hipGraphNode_t hNode, hipEvent_t event)` - + * Sets the event for an event record node in the given graphExec. */ + +// Kernel Functions to perform square and introduce delay in device. static __global__ void sqr_ker_func(int* a, int* b, size_t N, int clockrate, size_t delayMs) { int tx = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; if (tx < N) b[tx] = a[tx] * a[tx]; @@ -75,7 +52,25 @@ static __global__ void sqr_ker_func_gfx11(int* a, int* b, size_t N, int clockrat } /** - * Scenario 1: Test to validate setting different events in executable graph. + * Test Description + * ------------------------ + * - Create a graph with event record nodes as follows: + * -# Launch kernel. + * -# Record event. + * - Instantiate graph. + * - Create another graph with event record nodes as follows: + * -# Create wait event node on event + * - Instantiate the second graph. + * - Change the event in first graph to the second event + * - Change the event in second graph to the second event + * - Execute both graphs. + * - Verify the kernel results match. + * Test source + * ------------------------ + * - unit/graph/hipGraphExecEventWaitNodeSetEvent.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory") { constexpr size_t gridSize = 64; @@ -199,8 +194,20 @@ TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory") { } /** - * Scenario 2: Test to validate setting a different event in an executable - * graph does not impact the original graph and nodes. + * Test Description + * ------------------------ + * - Scenario to verify that event set does not + * impact the graph and changes only the executable graph. + * - Create an event wait node with first event and add it to graph. + * - Instantiate the graph to create an executable graph. + * - Change the event in the executable graph to the second event. + * - Verify that the event wait node still contains first event. + * Test source + * ------------------------ + * - unit/graph/hipGraphExecEventWaitNodeSetEvent.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_VerifyEventNotChanged") { hipGraph_t graph; @@ -224,7 +231,33 @@ TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_VerifyEventNotChanged") { } /** - * Scenario 3: Negative and Parameter Tests. + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When executable graph handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When node handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When pointer to the event is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When executable graph is not initialized + * - Expected output: return `hipErrorInvalidValue` + * -# When node is not initialized + * - Expected output: return `hipErrorInvalidValue` + * -# When event is not initialized + * - Expected output: return `hipErrorInvalidValue` + * -# When event wait node does not exist + * - Expected output: return `hipErrorInvalidValue` + * -# When node is memset node + * - Expected output: return `hipErrorInvalidValue` + * -# When node is record wait node + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphExecEventWaitNodeSetEvent.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_Negative") { hipGraph_t graph; diff --git a/catch/unit/graph/hipGraphExecHostNodeSetParams.cc b/catch/unit/graph/hipGraphExecHostNodeSetParams.cc index 6341e73e0..48916b73b 100644 --- a/catch/unit/graph/hipGraphExecHostNodeSetParams.cc +++ b/catch/unit/graph/hipGraphExecHostNodeSetParams.cc @@ -19,29 +19,18 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Test Case Scenarios of hipGraphExecHostNodeSetParams API: - -Functional: -1. Creates graph, Adds HostNode, update hostNode params using hipGraphExecHostNodeSetParams API - and validates the result -2. Create graph, Add Graph nodes and clones the graph. Add Host node to the cloned graph, update - hostNode params using hipGraphExecHostNodeSetParams API and validate the result - -Negative: - -1) Pass hGraphExec as nullptr and verify api doesn't crash, returns error code. -2) Pass node as nullptr and verify api doesn't crash, returns error code. -3) Pass pNodeParams as nullptr and verify api doesn't crash, returns error code. -3) Pass hipHostNodeParams::hipHostFn_t as nullptr and verify api doesn't crash, returns error code. -4) Pass uninitialized host params and verify api doesn't crash, returns error code. -5) Pass uninitialized graph and verify api doesn't crash, returns error code. -6) Pass nullptr to host func and verify api doesn't crash, returns error code. -*/ - #include #include +/** + * @addtogroup hipGraphExecHostNodeSetParams hipGraphExecHostNodeSetParams + * @{ + * @ingroup GraphTest + * `hipGraphExecHostNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node, + * const hipHostNodeParams* pNodeParams)` - + * Sets the parameters for a host node in the given graphExec. + */ + #define SIZE 1024 void callbackfunc(void* A_h) { @@ -58,10 +47,32 @@ void callbackfunc_setparams(void* B_h) { } } -/* -This test case verifies the negative scenarios of -hipGraphExecHostNodeSetParams API -*/ +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When graph exec handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When host node params pointer is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When graph handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When host params function data member is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When host params are not initialized + * - Expected output: return `hipErrorInvalidValue` + * -# When node is not a host node + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidValue` + * -# When node is not instantiated + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphExecHostNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphExecHostNodeSetParams_Negative") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); @@ -129,7 +140,7 @@ TEST_CASE("Unit_hipGraphExecHostNodeSetParams_Negative") { hipErrorInvalidValue); } -#if HT_NVIDIA // segfaults on AMD +#if HT_NVIDIA // segfaults on AMD SECTION("node is not a host node") { HIP_CHECK_ERROR(hipGraphExecHostNodeSetParams(graphExec, empty_node, &sethostParams), hipErrorInvalidValue); @@ -148,12 +159,21 @@ TEST_CASE("Unit_hipGraphExecHostNodeSetParams_Negative") { HIP_CHECK(hipStreamDestroy(streamForGraph)); } -/* -This test case verifies hipGraphExecHostNodeSetParams API in cloned graph -Creates graph, Add graph nodes and clone the graph -Add HostNode to the cloned graph,update the host params using -hipGraphExecHostNodeSetParams API and validates the result -*/ +/** + * Test Description + * ------------------------ + * - Verifies API behaviour in a cloned graph. + * - Creates graph and adds graph nodes. + * - Clones the graph and adds host node to the cloned graph. + * - Update the host params using set API. + * - Validates the result. + * Test source + * ------------------------ + * - unit/graph/hipGraphExecHostNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphExecHostNodeSetParams_ClonedGraphWithHostNode") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); @@ -210,12 +230,19 @@ TEST_CASE("Unit_hipGraphExecHostNodeSetParams_ClonedGraphWithHostNode") { HIP_CHECK(hipStreamDestroy(streamForGraph)); } -/* -This test case verifies the following scenario -Create graph, Adds host node to the graph, -updates the host params using hipGraphExecHostNodeSetParams API -and validates the result -*/ +/** + * Test Description + * ------------------------ + * - Creates graph and adds host node to the graph. + * - Updates the host params using set functionality. + * - Validates the result. + * Test source + * ------------------------ + * - unit/graph/hipGraphExecHostNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphExecHostNodeSetParams_BasicFunc") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); diff --git a/catch/unit/graph/hipGraphExecKernelNodeSetParams.cc b/catch/unit/graph/hipGraphExecKernelNodeSetParams.cc index 524e38b46..50aa2e0c2 100644 --- a/catch/unit/graph/hipGraphExecKernelNodeSetParams.cc +++ b/catch/unit/graph/hipGraphExecKernelNodeSetParams.cc @@ -19,25 +19,47 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Test Case Scenarios : -Negative - -1) Pass hGraphExec as nullptr and verify api returns error code. -2) Pass node as nullptr and verify api returns error code. -3) Pass NodeParams as un-initialized structure object and verify api returns error code. -4) Pass pNodeParams as nullptr and verify api returns error code. -5) Pass NodeParams:func data member as nullptr and verify api returns error code. -Functional - -1) Instantiate a graph with kernel node, obtain executable graph and update - the kernel node params with set and check it is taking effect. -*/ - #include #include #include /** - * Negative Test for API hipGraphExecKernelNodeSetParams + * @addtogroup hipGraphExecKernelNodeSetParams hipGraphExecKernelNodeSetParams + * @{ + * @ingroup GraphTest + * `hipGraphExecKernelNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node, + * const hipKernelNodeParams* pNodeParams)` - + * Sets the parameters for a kernel node in the given graphExec. + */ + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When graph exec handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When graph node handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When node params pointer is `nullptr` + * - Platform specific (AMD) + * - Expected output: return `hipErrorInvalidValue` + * -# When node params func data member is `nullptr` + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidDeviceFunction` + * -# When node params kernel params data member is `nullptr` + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidValue` + * -# When node is not a kernel node + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidValue` + * -# When node is not instantiated + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphExecKernelNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphExecKernelNodeSetParams_Negative") { constexpr size_t N = 1024; @@ -126,7 +148,18 @@ TEST_CASE("Unit_hipGraphExecKernelNodeSetParams_Negative") { } /** - * Functional Test for API Exec Kernel Params + * Test Description + * ------------------------ + * - Instantiates a graph with kernel node. + * - Obtains executable graph. + * - Updates the kernel node params with set. + * - Checks its taking effect. + * Test source + * ------------------------ + * - unit/graph/hipGraphExecKernelNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphExecKernelNodeSetParams_Functional") { constexpr size_t N = 1024; diff --git a/catch/unit/graph/hipGraphHostNodeGetParams.cc b/catch/unit/graph/hipGraphHostNodeGetParams.cc index ffb04cc17..2a9d8f5d5 100644 --- a/catch/unit/graph/hipGraphHostNodeGetParams.cc +++ b/catch/unit/graph/hipGraphHostNodeGetParams.cc @@ -19,26 +19,17 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Test Case Scenarios of hipGraphHostNodeGetParams API: - -Functional Scenarios: -1) Create a graph, add Host node to graph with desired node params. Verify api fetches the node -params which were mentioned while adding the host node. 2) Set host node params with -hipGraphHostNodeSetParams, now get the params and verify both are same. 3) Create graph, Add Graph -nodes and clones the graph. Add Host node to the cloned graph, update hostNode params using -hipGraphHostNodeSetParams API now get the params and verify both are same - -Negative Scenarios: - -1) Pass pGraphNode as nullptr and verify api doesn’t crash, returns error code. -2) Pass pNodeParams as nullptr and verify api doesn’t crash, returns error code. -3) Pass uninitialized graph node -*/ - #include #include +/** + * @addtogroup hipGraphHostNodeGetParams hipGraphHostNodeGetParams + * @{ + * @ingroup GraphTest + * `hipGraphHostNodeGetParams(hipGraphNode_t node, hipHostNodeParams* pNodeParams)` - + * Returns a host node's parameters. + */ + #define SIZE 1024 static void callbackfunc(void* A_h) { @@ -55,10 +46,24 @@ static void callbackfunc_setparams(void* B_h) { } } -/* -This test case verifies the negative scenarios of -hipGraphHostNodeGetParams API -*/ +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When graph node handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When output pointer to the node params is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When node is not a host node + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphHostNodeGetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphHostNodeGetParams_Negative") { constexpr size_t N = 1024; hipGraph_t graph; @@ -83,7 +88,7 @@ TEST_CASE("Unit_hipGraphHostNodeGetParams_Negative") { HIP_CHECK_ERROR(hipGraphHostNodeGetParams(hostNode, nullptr), hipErrorInvalidValue); } -#if HT_NVIDIA // segfaults on AMD +#if HT_NVIDIA // segfaults on AMD SECTION("node is not a host node") { hipGraphNode_t empty_node; HIP_CHECK(hipGraphAddEmptyNode(&empty_node, graph, nullptr, 0)); @@ -95,13 +100,21 @@ TEST_CASE("Unit_hipGraphHostNodeGetParams_Negative") { HIP_CHECK(hipGraphDestroy(graph)); } -/* -This test case verifies hipGraphHostNodeGetParams API in cloned graph -Creates graph, Add graph nodes and clone the graph -Add HostNode to the cloned graph, update hostNode using hipGraphHostNodeSetParams, -then get the host node params using hipGraphHostNodeGetParams API and -compare it. -*/ +/** + * Test Description + * ------------------------ + * - Verifies API behaviour in cloned graph. + * - Creates graph and adds graph nodes. + * - Clones the graph and adds host node to the cloned graph. + * - Updates host node by setting its params. + * - Gets the host node params and compare them with the ones that are set. + * Test source + * ------------------------ + * - unit/graph/hipGraphHostNodeGetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphHostNodeGetParams_ClonedGraphWithHostNode") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); @@ -241,11 +254,18 @@ void hipGraphHostNodeGetParams_func(bool setparams) { HIP_CHECK(hipStreamDestroy(streamForGraph)); } -/* -This test case verifies hipGraphHostNodeGetParams API by -adding host node to graph and gets the host params and -validates it -*/ +/** + * Test Description + * ------------------------ + * - Adds host node to the graph. + * - Gets host params and validates them. + * Test source + * ------------------------ + * - unit/graph/hipGraphHostNodeGetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphHostNodeGetParams_BasicFunc") { hipGraphHostNodeGetParams_func(false); } /* @@ -254,4 +274,17 @@ adding host node to graph, updates host node params using hipGraphHostNodeSetParams and gets the host params validates it */ +/** + * Test Description + * ------------------------ + * - Adds host node to the graph. + * - Updates host node params using set function. + * - Gets params and validates them. + * Test source + * ------------------------ + * - unit/graph/hipGraphHostNodeGetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphHostNodeGetParams_SetParams") { hipGraphHostNodeGetParams_func(true); } diff --git a/catch/unit/graph/hipGraphHostNodeSetParams.cc b/catch/unit/graph/hipGraphHostNodeSetParams.cc index d9b275b37..1800efaba 100644 --- a/catch/unit/graph/hipGraphHostNodeSetParams.cc +++ b/catch/unit/graph/hipGraphHostNodeSetParams.cc @@ -19,26 +19,17 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Test Case Scenarios of hipGraphHostNodeSetParams API: - -Functional: -1. Creates graph, Adds HostNode, update hostNode params using hipGraphHostNodeSetParams API - and validates the result -2. Create graph, Add Graph nodes and clones the graph. Add Host node to the cloned graph, update - hostNode params using hipGraphHostNodeSetParams API and validate the result - -Negative: - -1) Pass pGraphNode as nullptr and verify api doesn’t crash, returns error code. -2) Pass pNodeParams as nullptr and verify api doesn’t crash, returns error code. -3) Pass hipHostNodeParams::hipHostFn_t as nullptr and verify api doesn't crash, returns error code. -4) Pass uninitialized host params -*/ - #include #include +/** + * @addtogroup hipGraphHostNodeSetParams hipGraphHostNodeSetParams + * @{ + * @ingroup GraphTest + * `hipGraphHostNodeSetParams(hipGraphNode_t node, const hipHostNodeParams* pNodeParams)` - + * Sets a host node's parameters. + */ + #define SIZE 1024 static void callbackfunc(void* A_h) { @@ -55,10 +46,28 @@ static void callbackfunc_setparams(void* B_h) { } } -/* -This test case verifies the negative scenarios of -hipGraphHostNodeSetParams API -*/ +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When graph node handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When pointer to the host params is `nullptr` + * - Expected output: return `hipErrorInvaldValue` + * -# When host params host function data member is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When host params are not initialized + * - Expected output: return `hipErrorInvalidValue` + * -# When node is not host node + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphHostNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphHostNodeSetParams_Negative") { constexpr size_t N = 1024; hipGraph_t graph; @@ -92,7 +101,7 @@ TEST_CASE("Unit_hipGraphHostNodeSetParams_Negative") { HIP_CHECK_ERROR(hipGraphHostNodeSetParams(hostNode, &unintParams), hipErrorInvalidValue); } -#if HT_NVIDIA // segfaults on AMD +#if HT_NVIDIA // segfaults on AMD SECTION("node is not a host node") { hipGraphNode_t empty_node; HIP_CHECK(hipGraphAddEmptyNode(&empty_node, graph, nullptr, 0)); @@ -104,12 +113,21 @@ TEST_CASE("Unit_hipGraphHostNodeSetParams_Negative") { HIP_CHECK(hipGraphDestroy(graph)); } -/* -This test case verifies hipGraphHostNodeSetParams API in cloned graph -Creates graph, Add graph nodes and clone the graph -Add HostNode to the cloned graph,update the host params using -hipGraphHostNodeSetParams API and validates the result -*/ +/** + * Test Description + * ------------------------ + * - Verifes API functionality in the cloned graph. + * - Creates graph and adds graph nodes. + * - Clones the graph and adds host node to the cloned graph. + * - Updates the host params using set functionality. + * - Validates the results. + * Test source + * ------------------------ + * - unit/graph/hipGraphHostNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphHostNodeSetParams_ClonedGraphWithHostNode") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); @@ -163,12 +181,19 @@ TEST_CASE("Unit_hipGraphHostNodeSetParams_ClonedGraphWithHostNode") { HIP_CHECK(hipStreamDestroy(streamForGraph)); } -/* -This test case verifies the following scenario -Create graph, Adds host node to the graph, -updates the host params using hipGraphHostNodeSetParams API -and validates the result -*/ +/** + * Test Description + * ------------------------ + * - Creates graph and adds host node to the graph. + * - Updates the host params using set functionality. + * - Validates the result. + * Test source + * ------------------------ + * - unit/graph/hipGraphHostNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphHostNodeSetParams_BasicFunc") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); diff --git a/catch/unit/graph/hipGraphInstantiate.cc b/catch/unit/graph/hipGraphInstantiate.cc index bb9ff15ab..6d1203189 100644 --- a/catch/unit/graph/hipGraphInstantiate.cc +++ b/catch/unit/graph/hipGraphInstantiate.cc @@ -17,42 +17,43 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : -Functional - -1) Create a graph and then used it for hipGraphInstantiate without adding any node to graph. -2.a) Create an invalid graph as shown below: - ---> empty node ---> empty node ---> - ^ | - | | - -------------------------------- - Try instantiating the graph. Instantiating the graph should result in error. -2.b) Create a more complex cyclic graph. Try instantiating the graph. Instantiating the graph - should result in error. -2.c) Create a graph with child node. The graph in the child node is a cyclical graph. - Try instantiating the graph. Instantiating the graph should result in error. -3.a) Create a graph with redundant dependencies. Instantiate and execute the graph and validate - the output. -3.b) Create a graph. Instantiate the graph multiple times. Execute all the instantiated graphs - and validate the output. Destroy the instantiated graphs at the end. -3.c) Create a graph. Instantiate the graph multiple times. In loop, execute an instantiated - graph, validate the output and destroy the current instantiated graph. - -Negative - -1) Pass pGraphExec as null ptr and verify that api returns error code and doesn’t crash. -2) Pass graph as null/invalid ptr and check if api returns error. -3) Pass pGraphExec as un-initilize object and verify that api returns error code and doesn’t crash. -4) Pass Graph as un-initilize and verify that api returns error code and doesn’t crash. -*/ - #include #include #include #define NUM_OF_INSTANCES 10 -/* Test verifies hipGraphInstantiate API Negative scenarios. + +/** + * @addtogroup hipGraphInstantiate hipGraphInstantiate + * @{ + * @ingroup GraphTest + * `hipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph, + * hipGraphNode_t* pErrorNode, char* pLogBuffer, size_t bufferSize)` - + * Creates an executable graph from a graph. + * ________________________ + * Test cases from other modules: + * - @ref Unit_hipGraph_BasicFunctional */ + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When output pointer to the executable graph is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When graph handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When graph is not initialized + * - Expected output: return `hipErrorInvalidValue` + * -# When executable graph is not initialized + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphInstantiate_Negative") { hipError_t ret; hipGraphExec_t gExec{}; @@ -79,8 +80,17 @@ TEST_CASE("Unit_hipGraphInstantiate_Negative") { HIP_CHECK(hipGraphDestroy(graph)); } -/* Test verifies hipGraphInstantiate Basic scenarios. -Create a graph and then used it for hipGraphInstantiate without adding any node to graph. +/** + * Test Description + * ------------------------ + * - Creates a graph. + * - Instantiates an executable graph. + * Test source + * ------------------------ + * - unit/graph/hipGraphInstantiate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphInstantiate_Basic") { hipGraph_t graph; @@ -94,9 +104,20 @@ TEST_CASE("Unit_hipGraphInstantiate_Basic") { HIP_CHECK(hipGraphDestroy(graph)); } #if HT_NVIDIA -/* Test Functional Scenario 2.a, 2.b, 2.c with hipGraphInstantiate and -hipGraphInstantiateWithFlags. -*/ + +/** + * Test Description + * ------------------------ + * - Validates following testcase scenarios, where instantiating the graph should result in error: + * -# Create a cyclic graph with two empty nodes + * -# Create a more complex cyclic graph + * -# Create a graph with child node. The graph in the child node is a cyclical graph + * Test source + * ------------------------ + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphInstantiate_InvalidCyclicGraph") { hipGraph_t graph; hipGraphExec_t graphExec; @@ -112,22 +133,19 @@ TEST_CASE("Unit_hipGraphInstantiate_InvalidCyclicGraph") { HIP_CHECK(hipGraphAddDependencies(graph, &emptyNode1, &emptyNode2, 1)); HIP_CHECK(hipGraphAddDependencies(graph, &emptyNode2, &emptyNode1, 1)); // Detect the error during instantiation - REQUIRE(hipErrorInvalidValue == - hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); - REQUIRE(hipErrorInvalidValue == - hipGraphInstantiateWithFlags(&graphExec, graph, 0)); + REQUIRE(hipErrorInvalidValue == hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + REQUIRE(hipErrorInvalidValue == hipGraphInstantiateWithFlags(&graphExec, graph, 0)); // Clone the illegal graph HIP_CHECK(hipGraphClone(&clonedgraph, graph)); // Try instantiating the cloned graph REQUIRE(hipErrorInvalidValue == - hipGraphInstantiate(&graphExec, clonedgraph, nullptr, nullptr, 0)); - REQUIRE(hipErrorInvalidValue == - hipGraphInstantiateWithFlags(&graphExec, clonedgraph, 0)); + hipGraphInstantiate(&graphExec, clonedgraph, nullptr, nullptr, 0)); + REQUIRE(hipErrorInvalidValue == hipGraphInstantiateWithFlags(&graphExec, clonedgraph, 0)); } SECTION("A More Complex Cyclic Graph") { - hipGraphNode_t emptyNode1, emptyNode2, emptyNode3, emptyNode4, - emptyNode5, emptyNode6, emptyNode7; + hipGraphNode_t emptyNode1, emptyNode2, emptyNode3, emptyNode4, emptyNode5, emptyNode6, + emptyNode7; hipGraph_t clonedgraph; // Create emptyNode and add it to graph with dependency HIP_CHECK(hipGraphAddEmptyNode(&emptyNode1, graph, nullptr, 0)); @@ -147,24 +165,21 @@ TEST_CASE("Unit_hipGraphInstantiate_InvalidCyclicGraph") { HIP_CHECK(hipGraphAddDependencies(graph, &emptyNode6, &emptyNode7, 1)); HIP_CHECK(hipGraphAddDependencies(graph, &emptyNode4, &emptyNode1, 1)); // Detect the error during instantiation - REQUIRE(hipErrorInvalidValue == - hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); - REQUIRE(hipErrorInvalidValue == - hipGraphInstantiateWithFlags(&graphExec, graph, 0)); + REQUIRE(hipErrorInvalidValue == hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + REQUIRE(hipErrorInvalidValue == hipGraphInstantiateWithFlags(&graphExec, graph, 0)); // Clone the illegal graph HIP_CHECK(hipGraphClone(&clonedgraph, graph)); // Try instantiating the cloned graph REQUIRE(hipErrorInvalidValue == - hipGraphInstantiate(&graphExec, clonedgraph, nullptr, nullptr, 0)); - REQUIRE(hipErrorInvalidValue == - hipGraphInstantiateWithFlags(&graphExec, clonedgraph, 0)); + hipGraphInstantiate(&graphExec, clonedgraph, nullptr, nullptr, 0)); + REQUIRE(hipErrorInvalidValue == hipGraphInstantiateWithFlags(&graphExec, clonedgraph, 0)); } SECTION("A Cyclic Graph as Child Node") { hipGraph_t childgraph; HIP_CHECK(hipGraphCreate(&childgraph, 0)); - hipGraphNode_t emptyNode1, emptyNode2, emptyNode3, emptyNode4, - emptyNode5, emptyNode6, childNode; + hipGraphNode_t emptyNode1, emptyNode2, emptyNode3, emptyNode4, emptyNode5, emptyNode6, + childNode; // Create emptyNode and add it to graph with dependency HIP_CHECK(hipGraphAddEmptyNode(&emptyNode1, graph, nullptr, 0)); HIP_CHECK(hipGraphAddEmptyNode(&emptyNode2, graph, nullptr, 0)); @@ -173,30 +188,24 @@ TEST_CASE("Unit_hipGraphInstantiate_InvalidCyclicGraph") { HIP_CHECK(hipGraphAddEmptyNode(&emptyNode4, childgraph, nullptr, 0)); HIP_CHECK(hipGraphAddEmptyNode(&emptyNode5, childgraph, nullptr, 0)); HIP_CHECK(hipGraphAddEmptyNode(&emptyNode6, childgraph, nullptr, 0)); - HIP_CHECK(hipGraphAddDependencies(childgraph, &emptyNode3, - &emptyNode4, 1)); - HIP_CHECK(hipGraphAddDependencies(childgraph, &emptyNode4, - &emptyNode5, 1)); - HIP_CHECK(hipGraphAddDependencies(childgraph, &emptyNode5, - &emptyNode6, 1)); + HIP_CHECK(hipGraphAddDependencies(childgraph, &emptyNode3, &emptyNode4, 1)); + HIP_CHECK(hipGraphAddDependencies(childgraph, &emptyNode4, &emptyNode5, 1)); + HIP_CHECK(hipGraphAddDependencies(childgraph, &emptyNode5, &emptyNode6, 1)); // Illegal dependency - HIP_CHECK(hipGraphAddDependencies(childgraph, &emptyNode5, - &emptyNode4, 1)); - HIP_CHECK(hipGraphAddChildGraphNode(&childNode, graph, nullptr, - 0, childgraph)); + HIP_CHECK(hipGraphAddDependencies(childgraph, &emptyNode5, &emptyNode4, 1)); + HIP_CHECK(hipGraphAddChildGraphNode(&childNode, graph, nullptr, 0, childgraph)); HIP_CHECK(hipGraphAddDependencies(graph, &emptyNode1, &childNode, 1)); HIP_CHECK(hipGraphAddDependencies(graph, &childNode, &emptyNode2, 1)); // Detect the error during instantiation - REQUIRE(hipErrorInvalidValue == - hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); - REQUIRE(hipErrorInvalidValue == - hipGraphInstantiateWithFlags(&graphExec, graph, 0)); + REQUIRE(hipErrorInvalidValue == hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + REQUIRE(hipErrorInvalidValue == hipGraphInstantiateWithFlags(&graphExec, graph, 0)); } HIP_CHECK(hipGraphDestroy(graph)); } #endif + /* Local function to initialize input data. -*/ + */ static void init_input(int* a, size_t size) { unsigned int seed = time(nullptr); for (size_t i = 0; i < size; i++) { @@ -204,8 +213,22 @@ static void init_input(int* a, size_t size) { } } -/* Test Functional Scenario 3.a, 3.b and 3.c. -*/ +/** + * Test Description + * ------------------------ + * - Validates following testcase scenarios:: + * -# Create a graph with redundant dependencies. Instantiate and execute the graph and validate + * the output + * -# Create a graph. Instantiate the graph multiple times. Execute all the instantiated graphs + * and validate the output. Destroy the instantiated graphs at the end + * -# Create a graph. Instantiate the graph multiple times. In loop, execute an instantiated + * graph, validate the output and destroy the current instantiated graph + * Test source + * ------------------------ + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphInstantiate_functionalScenarios") { hipGraph_t graph; hipGraphExec_t graphExec[NUM_OF_INSTANCES]; @@ -214,36 +237,29 @@ TEST_CASE("Unit_hipGraphInstantiate_functionalScenarios") { constexpr size_t size = 1024; constexpr auto blocksPerCU = 6; constexpr auto threadsPerBlock = 256; - unsigned blocks = HipTest::setNumBlocks(blocksPerCU, - threadsPerBlock, size); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, size); hipGraphNode_t memcpyh2d, kernelNode, memcpyd2h; - int *inputVec_d{nullptr}, *inputVec_h{nullptr}, *outputVec_h{nullptr}, - *outputVec_d{nullptr}; + int *inputVec_d{nullptr}, *inputVec_h{nullptr}, *outputVec_h{nullptr}, *outputVec_d{nullptr}; // host and device allocation - HipTest::initArrays(&inputVec_d, &outputVec_d, nullptr, - &inputVec_h, &outputVec_h, nullptr, size, false); + HipTest::initArrays(&inputVec_d, &outputVec_d, nullptr, &inputVec_h, &outputVec_h, nullptr, + size, false); // Create graph - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyh2d, graph, nullptr, 0, - inputVec_d, inputVec_h, sizeof(int) * size, - hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyd2h, graph, nullptr, 0, - outputVec_h, outputVec_d, sizeof(int) * size, - hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyh2d, graph, nullptr, 0, inputVec_d, inputVec_h, + sizeof(int) * size, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyd2h, graph, nullptr, 0, outputVec_h, outputVec_d, + sizeof(int) * size, hipMemcpyDeviceToHost)); hipKernelNodeParams kernelNodeParams{}; size_t N = size; void* kernelArgs[3] = {reinterpret_cast(&inputVec_d), - reinterpret_cast(&outputVec_d), - reinterpret_cast(&N)}; - kernelNodeParams.func = - reinterpret_cast(HipTest::vector_square); + reinterpret_cast(&outputVec_d), reinterpret_cast(&N)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vector_square); kernelNodeParams.gridDim = dim3(blocks, 1, 1); kernelNodeParams.blockDim = dim3(threadsPerBlock, 1, 1); kernelNodeParams.sharedMemBytes = 0; kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs); kernelNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, nullptr, 0, - &kernelNodeParams)); + HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, nullptr, 0, &kernelNodeParams)); HIP_CHECK(hipGraphAddDependencies(graph, &memcpyh2d, &kernelNode, 1)); HIP_CHECK(hipGraphAddDependencies(graph, &kernelNode, &memcpyd2h, 1)); @@ -259,7 +275,7 @@ TEST_CASE("Unit_hipGraphInstantiate_functionalScenarios") { HIP_CHECK(hipGraphLaunch(graphExec[0], stream)); HIP_CHECK(hipStreamSynchronize(stream)); for (size_t i = 0; i < size; i++) { - REQUIRE(outputVec_h[i] == (inputVec_h[i]*inputVec_h[i])); + REQUIRE(outputVec_h[i] == (inputVec_h[i] * inputVec_h[i])); } HIP_CHECK(hipStreamDestroy(stream)); HIP_CHECK(hipGraphExecDestroy(graphExec[0])); @@ -268,8 +284,7 @@ TEST_CASE("Unit_hipGraphInstantiate_functionalScenarios") { SECTION("Creating Multiple Instances Graph") { // Create Executable Graphs for (int i = 0; i < NUM_OF_INSTANCES; i++) { - HIP_CHECK(hipGraphInstantiate(&graphExec[i], graph, - nullptr, nullptr, 0)); + HIP_CHECK(hipGraphInstantiate(&graphExec[i], graph, nullptr, nullptr, 0)); REQUIRE(graphExec[i] != nullptr); } // Execute all the instances of the graph @@ -279,8 +294,8 @@ TEST_CASE("Unit_hipGraphInstantiate_functionalScenarios") { for (int i = 0; i < NUM_OF_INSTANCES; i++) { HIP_CHECK(hipGraphLaunch(graphExec[i], stream)); HIP_CHECK(hipStreamSynchronize(stream)); - for (size_t ii = 0; ii< size; ii++) { - REQUIRE(outputVec_h[ii] == (inputVec_h[ii]*inputVec_h[ii])); + for (size_t ii = 0; ii < size; ii++) { + REQUIRE(outputVec_h[ii] == (inputVec_h[ii] * inputVec_h[ii])); } } HIP_CHECK(hipStreamDestroy(stream)); @@ -292,8 +307,7 @@ TEST_CASE("Unit_hipGraphInstantiate_functionalScenarios") { SECTION("Creating Multiple Instances Graph and Destroying After Use") { // Create Executable Graphs for (int i = 0; i < NUM_OF_INSTANCES; i++) { - HIP_CHECK(hipGraphInstantiate(&graphExec[i], graph, - nullptr, nullptr, 0)); + HIP_CHECK(hipGraphInstantiate(&graphExec[i], graph, nullptr, nullptr, 0)); REQUIRE(graphExec[i] != nullptr); } // Execute all the instances of the graph @@ -303,15 +317,15 @@ TEST_CASE("Unit_hipGraphInstantiate_functionalScenarios") { for (int i = 0; i < NUM_OF_INSTANCES; i++) { HIP_CHECK(hipGraphLaunch(graphExec[i], stream)); HIP_CHECK(hipStreamSynchronize(stream)); - for (size_t ii = 0; ii< size; ii++) { - REQUIRE(outputVec_h[ii] == (inputVec_h[ii]*inputVec_h[ii])); + for (size_t ii = 0; ii < size; ii++) { + REQUIRE(outputVec_h[ii] == (inputVec_h[ii] * inputVec_h[ii])); } HIP_CHECK(hipGraphExecDestroy(graphExec[i])); } HIP_CHECK(hipStreamDestroy(stream)); } // Free - HipTest::freeArrays(inputVec_d, outputVec_d, nullptr, - inputVec_h, outputVec_h, nullptr, false); + HipTest::freeArrays(inputVec_d, outputVec_d, nullptr, inputVec_h, outputVec_h, nullptr, + false); HIP_CHECK(hipGraphDestroy(graph)); } diff --git a/catch/unit/graph/hipGraphInstantiateWithFlags.cc b/catch/unit/graph/hipGraphInstantiateWithFlags.cc index 9571962ef..7a0791348 100644 --- a/catch/unit/graph/hipGraphInstantiateWithFlags.cc +++ b/catch/unit/graph/hipGraphInstantiateWithFlags.cc @@ -17,42 +17,48 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/* -hipGraphInstantiateWithFlags(hipGraphExec_t* pGraphExec, hipGraph_t graph, unsigned long long flags); -Testcase Scenarios of hipGraphInstantiateWithFlags API: -Negative: -1) Pass nullptr to pGraphExec -2) Pass nullptr to graph -4) Pass invalid flag -Functional: -1) Create dependencies graph and instantiate the graph -2) Create graph in one GPU device and instantiate, launch in peer GPU device -3) Create stream capture graph and instantite the graph -4) Create stream capture graph in one GPU device and instantite the graph launch - in peer GPU device -Mapping is missing for NVIDIA platform hence skipping the testcases -*/ - - #include #include #include +/** + * @addtogroup hipGraphInstantiateWithFlags hipGraphInstantiateWithFlags + * @{ + * @ingroup GraphTest + * `hipGraphInstantiateWithFlags(hipGraphExec_t* pGraphExec, + * hipGraph_t graph, unsigned long long flags)` - + * Creates an executable graph from a graph. + */ + constexpr size_t N = 1000000; -/* This test covers the negative scenarios of - hipGraphInstantiateWithFlags API */ + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When output pointer to the executable graph is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When graph handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When flag is not valid + * - Expected output: do not return `hipSuccess` + * Test source + * ------------------------ + * - unit/graph/hipGraphInstantiateWithFlags.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphInstantiateWithFlags_Negative") { SECTION("Passing nullptr pGraphExec") { hipGraph_t graph; HIP_CHECK(hipGraphCreate(&graph, 0)); - REQUIRE(hipGraphInstantiateWithFlags(nullptr, - graph, 0) == hipErrorInvalidValue); + REQUIRE(hipGraphInstantiateWithFlags(nullptr, graph, 0) == hipErrorInvalidValue); } SECTION("Passing nullptr to graph") { hipGraphExec_t graphExec; - REQUIRE(hipGraphInstantiateWithFlags(&graphExec, - nullptr, 0) == hipErrorInvalidValue); + REQUIRE(hipGraphInstantiateWithFlags(&graphExec, nullptr, 0) == hipErrorInvalidValue); } SECTION("Passing Invalid flag") { @@ -63,9 +69,9 @@ TEST_CASE("Unit_hipGraphInstantiateWithFlags_Negative") { } } /* -This function verifies the following scenarios -1. Creates dependency graph, Instantiates the graph with flags and verifies it -2. Creates graph on one GPU-1 device and instantiates the graph on peer GPU device + This function verifies the following scenarios + 1. Creates dependency graph, Instantiates the graph with flags and verifies it + 2. Creates graph on one GPU-1 device and instantiates the graph on peer GPU device */ void GraphInstantiateWithFlags_DependencyGraph(bool ctxt_change = false) { constexpr size_t N = 1024; @@ -96,8 +102,7 @@ void GraphInstantiateWithFlags_DependencyGraph(bool ctxt_change = false) { memsetParams.elementSize = sizeof(char); memsetParams.width = Nbytes; memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, - &memsetParams)); + HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, &memsetParams)); memset(&memsetParams, 0, sizeof(memsetParams)); memsetParams.dst = reinterpret_cast(B_d); @@ -106,37 +111,33 @@ void GraphInstantiateWithFlags_DependencyGraph(bool ctxt_change = false) { memsetParams.elementSize = sizeof(char); memsetParams.width = Nbytes; memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memset_B, graph, nullptr, 0, - &memsetParams)); + HIP_CHECK(hipGraphAddMemsetNode(&memset_B, graph, nullptr, 0, &memsetParams)); - void* kernelArgs1[] = {&C_d, &memsetVal, reinterpret_cast(&NElem)}; - kernelNodeParams.func = - reinterpret_cast(HipTest::memsetReverse); + void* kernelArgs1[] = {&C_d, &memsetVal, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::memsetReverse); kernelNodeParams.gridDim = dim3(blocks); kernelNodeParams.blockDim = dim3(threadsPerBlock); kernelNodeParams.sharedMemBytes = 0; kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs1); kernelNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&memsetKer_C, graph, nullptr, 0, - &kernelNodeParams)); + HIP_CHECK(hipGraphAddKernelNode(&memsetKer_C, graph, nullptr, 0, &kernelNodeParams)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d, - Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d, Nbytes, + hipMemcpyDeviceToHost)); - void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; - kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); kernelNodeParams.gridDim = dim3(blocks); kernelNodeParams.blockDim = dim3(threadsPerBlock); kernelNodeParams.sharedMemBytes = 0; kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); kernelNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, - &kernelNodeParams)); + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, &kernelNodeParams)); // Create dependencies HIP_CHECK(hipGraphAddDependencies(graph, &memset_A, &memcpyH2D_A, 1)); @@ -163,9 +164,9 @@ void GraphInstantiateWithFlags_DependencyGraph(bool ctxt_change = false) { } /* -This function verifies the following scenarios -1. Creates stream capture graph, Instantiates the graph with flags and verifies it -2. Creates graph on one GPU-1 device and instantiates the graph on peer GPU device + This function verifies the following scenarios + 1. Creates stream capture graph, Instantiates the graph with flags and verifies it + 2. Creates graph on one GPU-1 device and instantiates the graph on peer GPU device */ void GraphInstantiateWithFlags_StreamCapture(bool deviceContextChg = false) { float *A_d, *C_d; @@ -182,7 +183,7 @@ void GraphInstantiateWithFlags_StreamCapture(bool deviceContextChg = false) { // Fill with Phi + i for (size_t i = 0; i < N; i++) { - A_h[i] = 1.618f + i; + A_h[i] = 1.618f + i; } HIP_CHECK(hipMalloc(&A_d, Nbytes)); HIP_CHECK(hipMalloc(&C_d, Nbytes)); @@ -199,8 +200,8 @@ void GraphInstantiateWithFlags_StreamCapture(bool deviceContextChg = false) { HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream)); - hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), - dim3(threadsPerBlock), 0, stream, A_d, C_d, N); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), dim3(threadsPerBlock), 0, stream, A_d, + C_d, N); HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream)); HIP_CHECK(hipStreamEndCapture(stream, &graph)); @@ -236,20 +237,38 @@ void GraphInstantiateWithFlags_StreamCapture(bool deviceContextChg = false) { HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipFree(C_d)); } -/* -This testcase verifies hipGraphInstantiateWithFlags API -by creating dependency graph and instantiate, launching and verifying -the result -*/ + +/** + * Test Description + * ------------------------ + * - Creates dependency graph. + * - Instantiates it and launches it. + * - Verifies the results. + * Test source + * ------------------------ + * - unit/graph/hipGraphInstantiateWithFlags.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphInstantiateWithFlags_DependencyGraph") { GraphInstantiateWithFlags_DependencyGraph(); } -/* -This testcase verifies hipGraphInstantiateWithFlags API -by creating dependency graph on GPU-0 and instantiate, launching and verifying -the result on GPU-1 -*/ +/** + * Test Description + * ------------------------ + * - Creates dependency graph on the device 0. + * - Instantiates it and launches it. + * - Verifies the result on the device 1. + * Test source + * ------------------------ + * - unit/graph/hipGraphInstantiateWithFlags.cc + * Test requirements + * ------------------------ + * - Multi-device + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphInstantiateWithFlags_DependencyGraphDeviceCtxtChg") { int numDevices = 0; int canAccessPeer = 0; @@ -266,11 +285,19 @@ TEST_CASE("Unit_hipGraphInstantiateWithFlags_DependencyGraphDeviceCtxtChg") { } } -/* -This testcase verifies hipGraphInstantiateWithFlags API -by creating capture graph and instantiate, launching and verifying -the result -*/ +/** + * Test Description + * ------------------------ + * - Creates capture graph. + * - Instantiates it and launches it. + * - Verifies the results. + * Test source + * ------------------------ + * - unit/graph/hipGraphInstantiateWithFlags.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphInstantiateWithFlags_StreamCapture") { int numDevices = 0; int canAccessPeer = 0; @@ -287,11 +314,20 @@ TEST_CASE("Unit_hipGraphInstantiateWithFlags_StreamCapture") { } } -/* -This testcase verifies hipGraphInstantiateWithFlags API -by creating capture graph on GPU-0 and instantiate, launching and verifying -the result on GPU-1 -*/ +/** + * Test Description + * ------------------------ + * - Creates capture graph on the device 0. + * - Instantiates it and launches it. + * - Verifies the result on the device 1. + * Test source + * ------------------------ + * - unit/graph/hipGraphInstantiateWithFlags.cc + * Test requirements + * ------------------------ + * - Multi-device + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphInstantiateWithFlags_StreamCaptureDeviceContextChg") { int numDevices = 0; int canAccessPeer = 0; @@ -336,14 +372,13 @@ TEST_CASE("Unit_hipGraphInstantiateWithFlags_FlagAutoFreeOnLaunch_check") { allocParam.poolProps.location.id = 0; allocParam.poolProps.location.type = hipMemLocationTypeDevice; - HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, - 0, &allocParam)); + HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, 0, &allocParam)); REQUIRE(allocParam.dptr != nullptr); - int *A_d = reinterpret_cast(allocParam.dptr); + int* A_d = reinterpret_cast(allocParam.dptr); // Instantiate with Flag and launch the graph - HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph, - hipGraphInstantiateFlagAutoFreeOnLaunch)); + HIP_CHECK( + hipGraphInstantiateWithFlags(&graphExec, graph, hipGraphInstantiateFlagAutoFreeOnLaunch)); HIP_CHECK(hipGraphLaunch(graphExec, stream)); HIP_CHECK(hipStreamSynchronize(stream)); diff --git a/catch/unit/graph/hipGraphKernelNodeGetParams.cc b/catch/unit/graph/hipGraphKernelNodeGetParams.cc index 810576865..f364d8340 100644 --- a/catch/unit/graph/hipGraphKernelNodeGetParams.cc +++ b/catch/unit/graph/hipGraphKernelNodeGetParams.cc @@ -19,24 +19,39 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Test Case Scenarios : -Negative - -1) Pass node as nullptr and verify api returns error code. -2) Pass pNodeParams as nullptr and verify api returns error code. -Functional - -1) Create a graph, add kernel node to graph with desired kernel node params. - Verify api fetches the node params mentioned while adding kernel node. -2) Set kernel node params with hipGraphKernelNodeSetParams, - now get the params and verify both are same. -*/ - #include #include +/** + * @addtogroup hipGraphKernelNodeGetParams hipGraphKernelNodeGetParams + * @{ + * @ingroup GraphTest + * `hipGraphKernelNodeGetParams(hipGraphNode_t node, hipKernelNodeParams* pNodeParams)` - + * Gets kernel node's parameters. + * ________________________ + * Test cases from other modules: + * - @ref Unit_hipGraphKernelNodeGetSetParams_Functional + */ + #define THREADS_PER_BLOCK 512 -/* Test verifies hipGraphKernelNodeGetParams API Negative scenarios. +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When node handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When output pointer to the kernel params is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When node is not kernel node + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphKernelNodeGetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphKernelNodeGetParams_Negative") { constexpr int N = 1024; @@ -104,7 +119,23 @@ static bool node_compare(hipKernelNodeParams* kNode1, hipKernelNodeParams* kNode return true; } -/* Test verifies hipGraphKernelNodeGetParams API Functional scenarios. +/** + * Test Description + * ------------------------ + * - Validates API funcitonality in different scenarios: + * -# Verifying returned kernel params + * - Create new graph node with desired params. + * - Get params and compare them with the desired params. + * -# Set kernel params, get them and verify + * - Create new desired params. + * - Set new params to the existing graph. + * - Get params and compare them with the desired params. + * Test source + * ------------------------ + * - unit/graph/hipGraphKernelNodeGetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphKernelNodeGetParams_Functional") { constexpr int N = 1024; diff --git a/catch/unit/graph/hipGraphKernelNodeSetParams.cc b/catch/unit/graph/hipGraphKernelNodeSetParams.cc index eb561c05e..f03fce4fb 100644 --- a/catch/unit/graph/hipGraphKernelNodeSetParams.cc +++ b/catch/unit/graph/hipGraphKernelNodeSetParams.cc @@ -19,24 +19,41 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Test Case Scenarios : -Negative - -1) Pass node as nullptr and verify api returns error code. -2) Pass pNodeParams as nullptr and verify api returns error code. -Functional - -1) Add kernel node to graph with certain kernel params, now update the kernel - node params with set and check taking effect after launching graph. -2) Add kernel node to graph with certain kernel params, now get kernel node parameters - with hipGraphKernelNodeGetParams, then update the kernel node params with - hipGraphKernelNodeSetParams, finally check taking effect after launching graph. -*/ - #include #include #include -/* Test verifies hipGraphKernelNodeSetParams API Negative scenarios. +/** + * @addtogroup hipGraphKernelNodeSetParams hipGraphKernelNodeSetParams + * @{ + * @ingroup GraphTest + * `hipGraphKernelNodeSetParams(hipGraphNode_t node, const hipKernelNodeParams* pNodeParams)` - + * Sets a kernel node's parameters. + */ + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When node handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When pointer to the node params is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When node params func data member is `nullptr` + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidDeviceFunction` + * -# When node params kernel params data member is `nullptr` + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidValue` + * -# When node is not a kernel node + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphKernelNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphKernelNodeSetParams_Negative") { constexpr int N = 1024; @@ -99,7 +116,20 @@ TEST_CASE("Unit_hipGraphKernelNodeSetParams_Negative") { } /** - * Functional Test for API Set Kernel Params + * Test Description + * ------------------------ + * - Add kernel node to graph with certain kernel params. + * - Update the kernel node params with set. + * - Check taking effect after launching graph. + * - Add kernel note to graph with certain kernel params. + * - Get kernel node params, update kernel node params. + * - Check taking effect after launching graph. + * Test source + * ------------------------ + * - unit/graph/hipGraphKernelNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphKernelNodeSetParams_Functional") { constexpr size_t N = 1024; @@ -264,6 +294,17 @@ class GraphKernelNodeGetSetParam { } }; +/** + * Test Description + * ------------------------ + * - Basic functional test that sets and gets node params. + * Test source + * ------------------------ + * - unit/graph/hipGraphKernelNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphKernelNodeGetSetParams_Functional") { hipGraph_t* graph; hipStream_t streamForGraph; diff --git a/catch/unit/graph/hipGraphNodeFindInClone.cc b/catch/unit/graph/hipGraphNodeFindInClone.cc index 174c0ebcb..61a02adf7 100644 --- a/catch/unit/graph/hipGraphNodeFindInClone.cc +++ b/catch/unit/graph/hipGraphNodeFindInClone.cc @@ -17,40 +17,48 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/* -Testcase Scenarios of hipGraphNodeFindInClone API: - -Negative: - -1) Pass nullptr to graph node -2) pass nullptr to original graph node -3) pass nullptr to clonedGraph -4) Pass original graph in place of the cloned graph -5) Pass invalid originalNode -6) Destroy the graph node in the original graph - and try to get the deleted graph node - from the cloned graph -7) Clone the graph,Add node to Original graph - and try to find the original node in the cloned graph - - -Functional: - -1) Get the graph node from the cloned graph corresponding to the original node -2) Create and clone the graph, modify the original graph and clone the graph again, - then try to find the newly added graph node from the cloned graph - -*/ - #include #include #include #include - -/* This test covers the negative scenarios of - hipGraphNodeFindInClone API */ - +/** + * @addtogroup hipGraphNodeFindInClone hipGraphNodeFindInClone + * @{ + * @ingroup GraphTest + * `hipGraphNodeFindInClone(hipGraphNode_t* pNode, + * hipGraphNode_t originalNode, hipGraph_t clonedGraph)` - + * Finds a cloned version of a node. + */ + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When cloned graph handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When original graph node handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When output pointer to the graph node is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When graph is not cloned + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidValue` + * -# When graph node that is to be found is destroyed + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidValue` + * -# When original node is not initialized + * - Expected output: return `hipErrorInvalidValue` + * -# When find node in the cloned graph which is present only + * in the original graph + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphNodeFindInClone.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphNodeFindInClone_Negative") { hipGraph_t graph; hipGraph_t clonedgraph; @@ -231,6 +239,19 @@ void hipGraphNodeFindInClone_Func(bool ModifyOrigGraph = false) { HIP_CHECK(hipGraphDestroy(clonedgraph)); } +/** + * Test Description + * ------------------------ + * - Validates basic functionalities: + * -# Finds the node from the original graph in the cloned graph. + * -# Finds the node from the modified original graph in the cloned graph. + * Test source + * ------------------------ + * - unit/graph/hipGraphNodeFindInClone.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphNodeFindInClone_Functional") { SECTION("hipGraphNodeFindInClone Basic Functionality") { hipGraphNodeFindInClone_Func(); diff --git a/catch/unit/graph/hipGraphNodeGetType.cc b/catch/unit/graph/hipGraphNodeGetType.cc index b373fc3f5..7c2146d45 100644 --- a/catch/unit/graph/hipGraphNodeGetType.cc +++ b/catch/unit/graph/hipGraphNodeGetType.cc @@ -16,37 +16,6 @@ LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : -Functional :: -1) Pass different types of valid nodes and get the corresponding type of the node -2) Add Graph node, destroy the graph node and add the same node of different type. - hipGraphNodeGetTye should return the new type -3) Add graph node, overwrite new type to the same graph node and trigger - hipGraphNodeGetType which should return the updated type. -4) Create a graph with different types of nodes. Clone the graph. Verify the types - of each of these nodes in the cloned graph using hipGraphNodeGetType. -5) Create a graph with different types of nodes. Pass the graph to a thread. In the - thread, verify node types of all the nodes in the graph -6) Create a graph with different types of nodes say X. Create graph Y with few nodes - and X as child graph. Now verify each of nodes of Y including the nodes inside - child graph using hipGraphNodeGetType() -7) Create a graph with different types of nodes along with dependencies between - nodes. Clone the graph. Verify the types of each of these nodes in the cloned - graph using hipGraphNodeGetType. -8) Create a graph with different types of nodes along with dependencies between - nodes. Pass the graph to thread and verify each type of node in the graph - -9) Create a graph with different types of nodes say X with dependencies between - nodes. Create graph Y with few nodes with dependencies and X as child graph. - Now verify each of nodes of Y including the nodes inside child graph using - hipGraphNodeGetType() - -Negative :: -1) Pass nullptr to graph node -2) Pass nullptr to node type -3) Pass invalid to graph node -*/ #include #include #include @@ -55,13 +24,36 @@ Negative :: #define SIZE 1024 __device__ __constant__ static int globalConst[SIZE]; -static void callbackfunc(void *A_h) { - int *A = reinterpret_cast(A_h); +static void callbackfunc(void* A_h) { + int* A = reinterpret_cast(A_h); for (int i = 0; i < SIZE; i++) { A[i] = i; } } +/** + * @addtogroup hipGraphNodeGetType hipGraphNodeGetType + * @{ + * @ingroup GraphTest + * `hipGraphNodeGetType(hipGraphNode_t node, hipGraphNodeType* pType)` - + * Returns a node's type. + */ + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When graph node handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When node is not initialized + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphNodeGetType.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphNodeGetType_Negative") { SECTION("Pass nullptr to graph node") { hipGraphNodeType nodeType; @@ -72,7 +64,7 @@ TEST_CASE("Unit_hipGraphNodeGetType_Negative") { hipGraphNode_t memcpyNode; hipGraph_t graph; HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipGraphAddEmptyNode(&memcpyNode, graph, nullptr , 0)); + HIP_CHECK(hipGraphAddEmptyNode(&memcpyNode, graph, nullptr, 0)); REQUIRE(hipGraphNodeGetType(memcpyNode, nullptr) == hipErrorInvalidValue); } @@ -83,6 +75,19 @@ TEST_CASE("Unit_hipGraphNodeGetType_Negative") { } } +/** + * Test Description + * ------------------------ + * - Validates different functionalitie: + * -# When the node is deleted and different node is assigned. + * -# When the graph node is overridden. + * Test source + * ------------------------ + * - unit/graph/hipGraphNodeGetType.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphNodeGetType_Functional") { constexpr size_t N = 1024; hipGraphNodeType nodeType; @@ -99,30 +104,40 @@ TEST_CASE("Unit_hipGraphNodeGetType_Functional") { SECTION("Delete Node and Assign different Node Type") { HIP_CHECK(hipStreamWaitEvent(stream, event, 0)); - HIP_CHECK(hipGraphAddEventWaitNode(&waiteventNode, graph, nullptr, 0, - event)); + HIP_CHECK(hipGraphAddEventWaitNode(&waiteventNode, graph, nullptr, 0, event)); HIP_CHECK(hipGraphNodeGetType(waiteventNode, &nodeType)); - HIP_CHECK(hipGraphAddEmptyNode(&waiteventNode, graph, nullptr , 0)); + HIP_CHECK(hipGraphAddEmptyNode(&waiteventNode, graph, nullptr, 0)); HIP_CHECK(hipGraphNodeGetType(waiteventNode, &nodeType)); REQUIRE(nodeType == hipGraphNodeTypeEmpty); } SECTION("Override the graph node and get Node Type") { HIP_CHECK(hipStreamWaitEvent(stream, event, 0)); - HIP_CHECK(hipGraphAddEventWaitNode(&waiteventNode, graph, nullptr, 0, - event)); + HIP_CHECK(hipGraphAddEventWaitNode(&waiteventNode, graph, nullptr, 0, event)); HIP_CHECK(hipGraphNodeGetType(waiteventNode, &nodeType)); - HIP_CHECK(hipGraphAddEmptyNode(&waiteventNode, graph, nullptr , 0)); + HIP_CHECK(hipGraphAddEmptyNode(&waiteventNode, graph, nullptr, 0)); HIP_CHECK(hipGraphNodeGetType(waiteventNode, &nodeType)); REQUIRE(nodeType == hipGraphNodeTypeEmpty); } HIP_CHECK(hipStreamDestroy(stream)); HIP_CHECK(hipEventDestroy(event)); } + /** - * Functional Test for hipGraphNodeGetType API - * Create a graph and add nodes to it. - * Verify each node type added. + * Test Description + * ------------------------ + * - Gets different types of graph nodes: + * -# When node is `hipGraphNodeTypeMemcpy` + * -# When node is `hipGraphNodeTypeKernel` + * -# When node is `hipGraphNodeTypeEmpty` + * -# When node is `hipGraphNodeTypeWaitEvent` + * -# When node is `hipGraphNodeTypeEventRecord` + * Test source + * ------------------------ + * - unit/graph/hipGraphNodeGetType.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); @@ -142,9 +157,9 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeType") { SECTION("Get Memcpy node NodeType from Symbol") { hipGraphNode_t memcpyFromSymbolNode; - HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, - nullptr, 0, B_d, HIP_SYMBOL(globalConst), Nbytes, 0, - hipMemcpyDeviceToDevice)); + HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, nullptr, 0, B_d, + HIP_SYMBOL(globalConst), Nbytes, 0, + hipMemcpyDeviceToDevice)); // Verify node type HIP_CHECK(hipGraphNodeGetType(memcpyFromSymbolNode, &nodeType)); REQUIRE(nodeType == hipGraphNodeTypeMemcpy); @@ -152,9 +167,9 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeType") { SECTION("Get Memcpy node NodeType to Symbol") { hipGraphNode_t memcpyToSymbolNode; - HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, - nullptr, 0, HIP_SYMBOL(globalConst), A_d, Nbytes, 0, - hipMemcpyDeviceToDevice)); + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, nullptr, 0, + HIP_SYMBOL(globalConst), A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); // Verify node type HIP_CHECK(hipGraphNodeGetType(memcpyToSymbolNode, &nodeType)); REQUIRE(nodeType == hipGraphNodeTypeMemcpy); @@ -166,9 +181,7 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeType") { hostParams.fn = callbackfunc; hostParams.userData = A_h; // Create a host node - HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, - nullptr, - 0, &hostParams)); + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, nullptr, 0, &hostParams)); // Verify node type HIP_CHECK(hipGraphNodeGetType(hostNode, &nodeType)); REQUIRE(nodeType == hipGraphNodeTypeHost); @@ -180,44 +193,42 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeType") { // Create child graph HIP_CHECK(hipGraphCreate(&childgraph, 0)); // Add child graph node - HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph, - nullptr, 0, childgraph)) + HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph, nullptr, 0, childgraph)) // Verify node type HIP_CHECK(hipGraphNodeGetType(childGraphNode, &nodeType)); REQUIRE(nodeType == hipGraphNodeTypeGraph); } SECTION("Get Memcpy NodeType") { - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); HIP_CHECK(hipGraphNodeGetType(memcpyNode, &nodeType)); // temp disable it until correct node is set // REQUIRE(nodeType == hipGraphNodeTypeMemcpy); - HIP_CHECK(hipGraphAddEmptyNode(&memcpyNode, graph, nullptr , 0)); + HIP_CHECK(hipGraphAddEmptyNode(&memcpyNode, graph, nullptr, 0)); HIP_CHECK(hipGraphNodeGetType(memcpyNode, &nodeType)); REQUIRE(nodeType == hipGraphNodeTypeEmpty); } SECTION("Get Kernel NodeType") { hipKernelNodeParams kernelNodeParams{}; - void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; - kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); kernelNodeParams.gridDim = dim3(blocks); kernelNodeParams.blockDim = dim3(threadsPerBlock); kernelNodeParams.sharedMemBytes = 0; kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs); kernelNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, nullptr, - 0, &kernelNodeParams)); + HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, nullptr, 0, &kernelNodeParams)); HIP_CHECK(hipGraphNodeGetType(kernelNode, &nodeType)); REQUIRE(nodeType == hipGraphNodeTypeKernel); } SECTION("Get Empty NodeType") { hipGraphNode_t emptyNode; - HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, graph, nullptr , 0)); + HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, graph, nullptr, 0)); HIP_CHECK(hipGraphNodeGetType(emptyNode, &nodeType)); REQUIRE(nodeType == hipGraphNodeTypeEmpty); } @@ -232,8 +243,7 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeType") { memsetParams.elementSize = sizeof(char); memsetParams.width = Nbytes; memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, - &memsetParams)); + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams)); HIP_CHECK(hipGraphNodeGetType(memsetNode, &nodeType)); REQUIRE(nodeType == hipGraphNodeTypeMemset); } @@ -245,8 +255,7 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeType") { hipStream_t stream; HIP_CHECK(hipStreamCreate(&stream)); HIP_CHECK(hipStreamWaitEvent(stream, event, 0)); - HIP_CHECK(hipGraphAddEventWaitNode(&waiteventNode, graph, nullptr, 0, - event)); + HIP_CHECK(hipGraphAddEventWaitNode(&waiteventNode, graph, nullptr, 0, event)); HIP_CHECK(hipGraphNodeGetType(waiteventNode, &nodeType)); REQUIRE(nodeType == hipGraphNodeTypeWaitEvent); HIP_CHECK(hipStreamDestroy(stream)); @@ -260,8 +269,7 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeType") { hipStream_t stream; HIP_CHECK(hipStreamCreate(&stream)); HIP_CHECK(hipEventRecord(event, stream)); - HIP_CHECK(hipGraphAddEventRecordNode(&recordeventNode, graph, nullptr, 0, - event)); + HIP_CHECK(hipGraphAddEventRecordNode(&recordeventNode, graph, nullptr, 0, event)); HIP_CHECK(hipGraphNodeGetType(recordeventNode, &nodeType)); REQUIRE(nodeType == hipGraphNodeTypeEventRecord); HIP_CHECK(hipStreamDestroy(stream)); @@ -273,13 +281,12 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeType") { } // Function to verify node Type -static void ChkNodeType(hipGraph_t graph, - const std::map *nodeTypeToQuery) { +static void ChkNodeType(hipGraph_t graph, const std::map* nodeTypeToQuery) { size_t numNodes{}; hipGraphNodeType nodeType; HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numNodes)); int numBytes = sizeof(hipGraphNode_t) * numNodes; - hipGraphNode_t* nodes = reinterpret_cast(malloc(numBytes)); + hipGraphNode_t* nodes = reinterpret_cast(malloc(numBytes)); REQUIRE(nodes != nullptr); HIP_CHECK(hipGraphGetNodes(graph, nodes, &numNodes)); std::map cntNode; @@ -288,8 +295,7 @@ static void ChkNodeType(hipGraph_t graph, cntNode[nodeType] += 1; } std::map::iterator iter; - std::map::const_iterator - iter1 = nodeTypeToQuery->begin(); + std::map::const_iterator iter1 = nodeTypeToQuery->begin(); for (iter = cntNode.begin(); iter != cntNode.end(); iter++) { REQUIRE(iter->first == iter1->first); REQUIRE(iter->second == iter1->second); @@ -301,8 +307,7 @@ static void ChkNodeType(hipGraph_t graph, free(nodes); } // Thread Function -static void thread_func(hipGraph_t graph, - std::map *numNode) { +static void thread_func(hipGraph_t graph, std::map* numNode) { ChkNodeType(graph, numNode); } /* @@ -310,7 +315,7 @@ static void thread_func(hipGraph_t graph, * of each of these nodes in the cloned graph using hipGraphNodeGetType. * 2.Create a graph with different types of nodes. Pass the graph to a thread. In the * thread, verify node types of all the nodes in the graph -*/ + */ TEST_CASE("Unit_hipGraphNodeGetType_NodeTypeOfClonedGraph_NodeTypeInThread") { hipGraph_t graph, childGraph, clonedGraph; int *A_d, *B_d, *C_d; @@ -321,55 +326,48 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeTypeOfClonedGraph_NodeTypeInThread") { unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); HIP_CHECK(hipGraphCreate(&graph, 0)); - hipGraphNode_t memcpyNode, kernelNode, hostNode, childGraphNode, emptyNode, - memsetNode, waiteventNode, recordeventNode; - int numMemcpy{}, numKernel{}, numMemset{}, numHost{}, numWaitEvent{}, - numEventRecord{}, numEmpty{}, numChild{}; + hipGraphNode_t memcpyNode, kernelNode, hostNode, childGraphNode, emptyNode, memsetNode, + waiteventNode, recordeventNode; + int numMemcpy{}, numKernel{}, numMemset{}, numHost{}, numWaitEvent{}, numEventRecord{}, + numEmpty{}, numChild{}; // Host Node hipHostNodeParams hostParams = {0, 0}; hostParams.fn = callbackfunc; hostParams.userData = A_h; - HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, - nullptr, 0, &hostParams)); + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, nullptr, 0, &hostParams)); numHost++; // Host NOde - HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, - nullptr, 0, &hostParams)); + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, nullptr, 0, &hostParams)); numHost++; // Host Node - HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, - nullptr, 0, &hostParams)); + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, nullptr, 0, &hostParams)); numHost++; // MemCpy Node - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); numMemcpy++; // Host Node - HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, - nullptr, 0, &hostParams)); + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, nullptr, 0, &hostParams)); numHost++; // Kernal Node hipKernelNodeParams kernelNodeParams{}; - void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; - kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); kernelNodeParams.gridDim = dim3(blocks); kernelNodeParams.blockDim = dim3(threadsPerBlock); kernelNodeParams.sharedMemBytes = 0; kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs); kernelNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, nullptr, - 0, &kernelNodeParams)); + HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, nullptr, 0, &kernelNodeParams)); numKernel++; // Child Node HIP_CHECK(hipGraphCreate(&childGraph, 0)); - HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph, - nullptr, 0, childGraph)); + HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph, nullptr, 0, childGraph)); numChild++; // Child Node - HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph, - nullptr, 0, childGraph)); + HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph, nullptr, 0, childGraph)); numChild++; // memSet Node hipMemsetParams memsetParams{}; @@ -380,8 +378,7 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeTypeOfClonedGraph_NodeTypeInThread") { memsetParams.elementSize = sizeof(char); memsetParams.width = Nbytes; memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, - &memsetParams)); + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams)); numMemset++; // WaitEvent Node hipEvent_t event1, event2; @@ -389,21 +386,19 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeTypeOfClonedGraph_NodeTypeInThread") { hipStream_t stream1, stream2; HIP_CHECK(hipStreamCreate(&stream1)); HIP_CHECK(hipStreamWaitEvent(stream1, event1, 0)); - HIP_CHECK(hipGraphAddEventWaitNode(&waiteventNode, graph, nullptr, 0, - event1)); + HIP_CHECK(hipGraphAddEventWaitNode(&waiteventNode, graph, nullptr, 0, event1)); numWaitEvent++; // Empty Node - HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, graph, nullptr , 0)); + HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, graph, nullptr, 0)); numEmpty++; // Empty Node - HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, graph, nullptr , 0)); + HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, graph, nullptr, 0)); numEmpty++; // Event Record Node HIP_CHECK(hipEventCreate(&event2)); HIP_CHECK(hipStreamCreate(&stream2)); HIP_CHECK(hipEventRecord(event2, stream2)); - HIP_CHECK(hipGraphAddEventRecordNode(&recordeventNode, graph, nullptr, 0, - event2)); + HIP_CHECK(hipGraphAddEventRecordNode(&recordeventNode, graph, nullptr, 0, event2)); numEventRecord++; numNode[hipGraphNodeTypeHost] = numHost; @@ -438,7 +433,7 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeTypeOfClonedGraph_NodeTypeInThread") { * Create a graph with different types of nodes say X. Create graph Y with * few nodes and X as child graph. Now verify each of nodes of Y including * the nodes inside child graph using hipGraphNodeGetType() -*/ + */ TEST_CASE("Unit_hipGraphNodeGetType_NodeTypeOfChildGraph") { hipGraph_t graph, childGraph, getGraph; int *A_d, *B_d, *C_d; @@ -448,10 +443,10 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeTypeOfChildGraph") { unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); HIP_CHECK(hipGraphCreate(&graph, 0)); - hipGraphNode_t memcpyNode, kernelNode, hostNode, childGraphNode, emptyNode, - memsetNode, waiteventNode, recordeventNode; - int numMemcpy{}, numKernel{}, numMemset{}, numHost{}, numWaitEvent{}, - numEventRecord{}, numEmpty{}, numChild{}; + hipGraphNode_t memcpyNode, kernelNode, hostNode, childGraphNode, emptyNode, memsetNode, + waiteventNode, recordeventNode; + int numMemcpy{}, numKernel{}, numMemset{}, numHost{}, numWaitEvent{}, numEventRecord{}, + numEmpty{}, numChild{}; std::map numNodeParent; std::map numNodeChild; @@ -467,8 +462,7 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeTypeOfChildGraph") { memsetParams.elementSize = sizeof(char); memsetParams.width = Nbytes; memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, childGraph, nullptr, 0, - &memsetParams)); + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, childGraph, nullptr, 0, &memsetParams)); numMemset++; // Add WaitEvent Node to child graph @@ -477,55 +471,50 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeTypeOfChildGraph") { hipStream_t stream1, stream2; HIP_CHECK(hipStreamCreate(&stream1)); HIP_CHECK(hipStreamWaitEvent(stream1, event1, 0)); - HIP_CHECK(hipGraphAddEventWaitNode(&waiteventNode, childGraph, nullptr, 0, - event1)); + HIP_CHECK(hipGraphAddEventWaitNode(&waiteventNode, childGraph, nullptr, 0, event1)); numWaitEvent++; // Add Empty Node to child graph - HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, childGraph, nullptr , 0)); + HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, childGraph, nullptr, 0)); numEmpty++; // Empty Node - HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, childGraph, nullptr , 0)); + HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, childGraph, nullptr, 0)); numEmpty++; // Add Event Record Node to child graph HIP_CHECK(hipEventCreate(&event2)); HIP_CHECK(hipStreamCreate(&stream2)); HIP_CHECK(hipEventRecord(event2, stream2)); - HIP_CHECK(hipGraphAddEventRecordNode(&recordeventNode, childGraph, nullptr, 0, - event2)); + HIP_CHECK(hipGraphAddEventRecordNode(&recordeventNode, childGraph, nullptr, 0, event2)); numEventRecord++; // Add Host Node to parent graph hipHostNodeParams hostParams = {0, 0}; hostParams.fn = callbackfunc; hostParams.userData = A_h; - HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, nullptr, - 0, &hostParams)); + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, nullptr, 0, &hostParams)); numHost++; // Add MemCpy Node to parent graph - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); numMemcpy++; // Add MemCpy Node to parent graph - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); numMemcpy++; // Add Kernal Node to parent graph hipKernelNodeParams kernelNodeParams{}; - void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; - kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); kernelNodeParams.gridDim = dim3(blocks); kernelNodeParams.blockDim = dim3(threadsPerBlock); kernelNodeParams.sharedMemBytes = 0; kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs); kernelNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, nullptr, - 0, &kernelNodeParams)); + HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, nullptr, 0, &kernelNodeParams)); numKernel++; // Add child node to the parent graph - HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph, - nullptr, 0, childGraph)) + HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph, nullptr, 0, childGraph)) numChild++; numNodeParent[hipGraphNodeTypeHost] = numHost; @@ -551,26 +540,22 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeTypeOfChildGraph") { HIP_CHECK(hipGraphDestroy(childGraph)); HIP_CHECK(hipGraphDestroy(graph)); } -enum graphType { - Parent, - Child }; +enum graphType { Parent, Child }; // Function to verify node Type static void ChkNodeTypeWithDependency(hipGraph_t graph, enum graphType Type) { size_t numNodes{}; hipGraphNodeType nodeType; - hipGraphNodeType Arr[] = {hipGraphNodeTypeHost, hipGraphNodeTypeMemcpy, + hipGraphNodeType Arr[] = {hipGraphNodeTypeHost, hipGraphNodeTypeMemcpy, hipGraphNodeTypeKernel, hipGraphNodeTypeGraph, hipGraphNodeTypeMemset, hipGraphNodeTypeWaitEvent, - hipGraphNodeTypeEmpty, hipGraphNodeTypeEventRecord}; + hipGraphNodeTypeEmpty, hipGraphNodeTypeEventRecord}; - hipGraphNodeType childArr[] = {hipGraphNodeTypeMemset, - hipGraphNodeTypeWaitEvent, - hipGraphNodeTypeEmpty, - hipGraphNodeTypeEventRecord}; + hipGraphNodeType childArr[] = {hipGraphNodeTypeMemset, hipGraphNodeTypeWaitEvent, + hipGraphNodeTypeEmpty, hipGraphNodeTypeEventRecord}; HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numNodes)); int numBytes = sizeof(hipGraphNode_t) * numNodes; - hipGraphNode_t* nodes = reinterpret_cast(malloc(numBytes)); + hipGraphNode_t* nodes = reinterpret_cast(malloc(numBytes)); REQUIRE(nodes != nullptr); HIP_CHECK(hipGraphGetNodes(graph, nodes, &numNodes)); @@ -604,37 +589,34 @@ TEST_CASE("Unit_hipGraphNodeGetType_ClonedGraph_InThread_WithDependencies") { unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); HIP_CHECK(hipGraphCreate(&graph, 0)); - hipGraphNode_t memcpyNode, kernelNode, hostNode, childGraphNode, emptyNode, - memsetNode, waiteventNode, recordeventNode; + hipGraphNode_t memcpyNode, kernelNode, hostNode, childGraphNode, emptyNode, memsetNode, + waiteventNode, recordeventNode; // Host Node hipHostNodeParams hostParams = {0, 0}; hostParams.fn = callbackfunc; hostParams.userData = A_h; - HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, - nullptr, 0, &hostParams)); + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, nullptr, 0, &hostParams)); // MemCpy Node - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); HIP_CHECK(hipGraphAddDependencies(graph, &hostNode, &memcpyNode, 1)); // Kernal Node hipKernelNodeParams kernelNodeParams{}; - void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; - kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); kernelNodeParams.gridDim = dim3(blocks); kernelNodeParams.blockDim = dim3(threadsPerBlock); kernelNodeParams.sharedMemBytes = 0; kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs); kernelNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, nullptr, - 0, &kernelNodeParams)); + HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, nullptr, 0, &kernelNodeParams)); HIP_CHECK(hipGraphAddDependencies(graph, &memcpyNode, &kernelNode, 1)); // Child Node HIP_CHECK(hipGraphCreate(&childGraph, 0)); - HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph, - nullptr, 0, childGraph)); + HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph, nullptr, 0, childGraph)); HIP_CHECK(hipGraphAddDependencies(graph, &kernelNode, &childGraphNode, 1)); // memSet Node @@ -646,8 +628,7 @@ TEST_CASE("Unit_hipGraphNodeGetType_ClonedGraph_InThread_WithDependencies") { memsetParams.elementSize = sizeof(char); memsetParams.width = Nbytes; memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, - &memsetParams)); + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams)); HIP_CHECK(hipGraphAddDependencies(graph, &childGraphNode, &memsetNode, 1)); // WaitEvent Node @@ -656,20 +637,18 @@ TEST_CASE("Unit_hipGraphNodeGetType_ClonedGraph_InThread_WithDependencies") { hipStream_t stream1, stream2; HIP_CHECK(hipStreamCreate(&stream1)); HIP_CHECK(hipStreamWaitEvent(stream1, event1, 0)); - HIP_CHECK(hipGraphAddEventWaitNode(&waiteventNode, graph, nullptr, 0, - event1)); + HIP_CHECK(hipGraphAddEventWaitNode(&waiteventNode, graph, nullptr, 0, event1)); HIP_CHECK(hipGraphAddDependencies(graph, &memsetNode, &waiteventNode, 1)); // Empty Node - HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, graph, nullptr , 0)); + HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, graph, nullptr, 0)); HIP_CHECK(hipGraphAddDependencies(graph, &waiteventNode, &emptyNode, 1)); // Event Record Node HIP_CHECK(hipEventCreate(&event2)); HIP_CHECK(hipStreamCreate(&stream2)); HIP_CHECK(hipEventRecord(event2, stream2)); - HIP_CHECK(hipGraphAddEventRecordNode(&recordeventNode, graph, nullptr, 0, - event2)); + HIP_CHECK(hipGraphAddEventRecordNode(&recordeventNode, graph, nullptr, 0, event2)); HIP_CHECK(hipGraphAddDependencies(graph, &emptyNode, &recordeventNode, 1)); // Clone the graph @@ -695,7 +674,7 @@ TEST_CASE("Unit_hipGraphNodeGetType_ClonedGraph_InThread_WithDependencies") { * nodes. Create graph Y with few nodes with dependencies and X as child graph. * Now verify each of nodes of Y including the nodes inside child graph using * hipGraphNodeGetType() -*/ + */ TEST_CASE("Unit_hipGraphNodeGetType_NodeTypeOfChildGraph_WithDependency") { hipGraph_t graph, childGraph, getGraph; int *A_d, *B_d, *C_d; @@ -705,8 +684,8 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeTypeOfChildGraph_WithDependency") { unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); HIP_CHECK(hipGraphCreate(&graph, 0)); - hipGraphNode_t memcpyNode, kernelNode, hostNode, childGraphNode, emptyNode, - memsetNode, waiteventNode, recordeventNode; + hipGraphNode_t memcpyNode, kernelNode, hostNode, childGraphNode, emptyNode, memsetNode, + waiteventNode, recordeventNode; // Create a child graph HIP_CHECK(hipGraphCreate(&childGraph, 0)); @@ -720,8 +699,7 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeTypeOfChildGraph_WithDependency") { memsetParams.elementSize = sizeof(char); memsetParams.width = Nbytes; memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, childGraph, nullptr, 0, - &memsetParams)); + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, childGraph, nullptr, 0, &memsetParams)); // Add WaitEvent Node to child graph hipEvent_t event1, event2; @@ -729,64 +707,53 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeTypeOfChildGraph_WithDependency") { hipStream_t stream1, stream2; HIP_CHECK(hipStreamCreate(&stream1)); HIP_CHECK(hipStreamWaitEvent(stream1, event1, 0)); - HIP_CHECK(hipGraphAddEventWaitNode(&waiteventNode, childGraph, nullptr, 0, - event1)); - HIP_CHECK(hipGraphAddDependencies(childGraph, &memsetNode, - &waiteventNode, 1)); + HIP_CHECK(hipGraphAddEventWaitNode(&waiteventNode, childGraph, nullptr, 0, event1)); + HIP_CHECK(hipGraphAddDependencies(childGraph, &memsetNode, &waiteventNode, 1)); // Add Empty Node to child graph - HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, childGraph, nullptr , 0)); + HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, childGraph, nullptr, 0)); HIP_CHECK(hipGraphAddDependencies(childGraph, &waiteventNode, &emptyNode, 1)); // Add Event Record Node to child graph HIP_CHECK(hipEventCreate(&event2)); HIP_CHECK(hipStreamCreate(&stream2)); HIP_CHECK(hipEventRecord(event2, stream2)); - HIP_CHECK(hipGraphAddEventRecordNode(&recordeventNode, childGraph, nullptr, 0, - event2)); - HIP_CHECK(hipGraphAddDependencies(childGraph, &emptyNode, - &recordeventNode, 1)); + HIP_CHECK(hipGraphAddEventRecordNode(&recordeventNode, childGraph, nullptr, 0, event2)); + HIP_CHECK(hipGraphAddDependencies(childGraph, &emptyNode, &recordeventNode, 1)); // Add Host Node to parent graph hipHostNodeParams hostParams = {0, 0}; hostParams.fn = callbackfunc; hostParams.userData = A_h; - HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, nullptr, - 0, &hostParams)); + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, nullptr, 0, &hostParams)); // Add MemCpy Node to parent graph - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); HIP_CHECK(hipGraphAddDependencies(graph, &hostNode, &memcpyNode, 1)); // Add Kernal Node to parent graph hipKernelNodeParams kernelNodeParams{}; - void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; - kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); kernelNodeParams.gridDim = dim3(blocks); kernelNodeParams.blockDim = dim3(threadsPerBlock); kernelNodeParams.sharedMemBytes = 0; kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs); kernelNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, nullptr, - 0, &kernelNodeParams)); + HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, nullptr, 0, &kernelNodeParams)); HIP_CHECK(hipGraphAddDependencies(graph, &memcpyNode, &kernelNode, 1)); // Add child node to the parent graph - HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph, - nullptr, 0, childGraph)); + HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph, nullptr, 0, childGraph)); HIP_CHECK(hipGraphAddDependencies(graph, &kernelNode, &childGraphNode, 1)); // Check Node Type of graph - SECTION("Graph node Type verification") { - ChkNodeTypeWithDependency(graph, Parent); - } + SECTION("Graph node Type verification") { ChkNodeTypeWithDependency(graph, Parent); } // Get the child graph from graph HIP_CHECK(hipGraphChildGraphNodeGetGraph(childGraphNode, &getGraph)); - SECTION("Child Graph node Type verification") { - ChkNodeTypeWithDependency(getGraph, Child); - } + SECTION("Child Graph node Type verification") { ChkNodeTypeWithDependency(getGraph, Child); } HIP_CHECK(hipStreamDestroy(stream1)); HIP_CHECK(hipEventDestroy(event1)); diff --git a/catch/unit/graph/hipGraphReleaseUserObject.cc b/catch/unit/graph/hipGraphReleaseUserObject.cc index 613d8478b..d6a1630da 100644 --- a/catch/unit/graph/hipGraphReleaseUserObject.cc +++ b/catch/unit/graph/hipGraphReleaseUserObject.cc @@ -25,11 +25,34 @@ THE SOFTWARE. #include "user_object_common.hh" /** - * Negative Test for API - hipGraphReleaseUserObject - 1) Pass graph as nullptr - 2) Pass User Object as nullptr - 3) Pass initialRefcount as 0 - 4) Pass initialRefcount as INT_MAX + * @addtogroup hipGraphReleaseUserObject hipGraphReleaseUserObject + * @{ + * @ingroup GraphTest + * `hipGraphReleaseUserObject(hipGraph_t graph, hipUserObject_t object, unsigned int count + * __dparm(1))` - Release user object from graphs. + * ________________________ + * Test cases from other modules: + * - @ref Unit_hipGraphRetainUserObject_Functional_2 + */ + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When graph handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When used object handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When count is zero + * - Expected output: return `hipErrorInvalidValue` + * -# When count is INT_MAX + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/graph/hipGraphReleaseUserObject.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphReleaseUserObject_Negative") { hipGraph_t graph; diff --git a/catch/unit/graph/hipGraphRetainUserObject.cc b/catch/unit/graph/hipGraphRetainUserObject.cc index ebbe322d1..9ef40ad15 100644 --- a/catch/unit/graph/hipGraphRetainUserObject.cc +++ b/catch/unit/graph/hipGraphRetainUserObject.cc @@ -27,7 +27,12 @@ THE SOFTWARE. #include "user_object_common.hh" /** - * Functional Test for API - hipGraphRetainUserObject + * @addtogroup hipGraphRetainUserObject hipGraphRetainUserObject + * @{ + * @ingroup GraphTest + * `hipGraphRetainUserObject(hipGraph_t graph, hipUserObject_t object, + * unsigned int count __dparm(1), unsigned int flags __dparm(0))` - + * Retain user object for graphs. */ /* 1) Create GraphUserObject and retain it by calling hipGraphRetainUserObject @@ -47,6 +52,27 @@ static void hipGraphRetainUserObject_Functional_1(void* object, void destroyObj( HIP_CHECK(hipGraphDestroy(graph)); } +/** + * Test Description + * ------------------------ + * - Create user object successfully. + * - Release it with no errors. + * - Perform action for different objects: + * -# When object is int + * - Expected output: return `hipSuccess` + * -# When object is float + * - Expected output: return `hipSuccess` + * -# When object is class instance + * - Expected output: return `hipSuccess` + * -# When object is struct instance + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/graph/hipGraphRetainUserObject.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphRetainUserObject_Functional_1") { SECTION("Called with int Object") { int* object = new int(); @@ -72,6 +98,19 @@ TEST_CASE("Unit_hipGraphRetainUserObject_Functional_1") { /* 2) Create UserObject and GraphUserObject and retain using custom reference count and release it by calling hipGraphReleaseUserObject with count. */ +/** + * Test Description + * ------------------------ + * - Create user object and graph user object. + * - Retain graph using custom reference count. + * - Release it by calling release function with count. + * Test source + * ------------------------ + * - unit/graph/hipGraphRetainUserObject.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphRetainUserObject_Functional_2") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); @@ -143,13 +182,27 @@ TEST_CASE("Unit_hipGraphRetainUserObject_Functional_2") { } /** - * Negative Test for API - hipGraphRetainUserObject - 1) Pass graph as nullptr - 2) Pass User Object as nullptr - 3) Pass initialRefcount as 0 - 4) Pass initialRefcount as INT_MAX - 5) Pass flag as 0 - 6) Pass flag as INT_MAX + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When graph handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When user object handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When count is zero + * - Expected output: return `hipErrorInvalidValue` + * -# When count is INT_MAX + * - Expected output: return `hipSuccess` + * -# When flag is zero + * - Expected output: return `hipSuccess` + * -# When flag is INT_MAX + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipGraphRetainUserObject.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraphRetainUserObject_Negative") { hipGraph_t graph; @@ -187,6 +240,25 @@ TEST_CASE("Unit_hipGraphRetainUserObject_Negative") { HIP_CHECK(hipGraphDestroy(graph)); } +/** + * Test Description + * ------------------------ + * - Create user object from float. + * - Retain graph object with reference count 2. + * - Expected output: return `hipSuccess` + * - Release graph object with reference count greater than 2. + * - Expected output: return `hipSuccess` + * - Retain graph object with reference count 8. + * - Expected output: return `hipSuccess` + * - Release graph object with reference count 1. + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/graph/hipGraphRetainUserObject.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphRetainUserObject_Negative_Basic") { hipGraph_t graph; HIP_CHECK(hipGraphCreate(&graph, 0)); @@ -215,6 +287,25 @@ TEST_CASE("Unit_hipGraphRetainUserObject_Negative_Basic") { HIP_CHECK(hipGraphDestroy(graph)); } +/** + * Test Description + * ------------------------ + * - Create user object from `nullptr`. + * - Retain graph object with reference count 2. + * - Expected output: return `hipSuccess` + * - Release graph object with reference count greater than 2. + * - Expected output: return `hipSuccess` + * - Retain graph object with reference count 8. + * - Expected output: return `hipSuccess` + * - Release graph object with reference count 1. + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/graph/hipGraphRetainUserObject.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGraphRetainUserObject_Negative_Null_Object") { hipGraph_t graph; HIP_CHECK(hipGraphCreate(&graph, 0)); diff --git a/catch/unit/graph/hipSimpleGraphWithKernel.cc b/catch/unit/graph/hipSimpleGraphWithKernel.cc index 08f20cbe4..725dd0a18 100644 --- a/catch/unit/graph/hipSimpleGraphWithKernel.cc +++ b/catch/unit/graph/hipSimpleGraphWithKernel.cc @@ -17,13 +17,14 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** - Simple test to demonstrate usage of graph. - Compares implementation with and without using graphs. -*/ - #include +/** + * @addtogroup hipGraphInstantiate hipGraphInstantiate + * @{ + * @ingroup GraphTest + */ + #define N 1024 * 1024 #define NSTEP 1000 #define NKERNEL 25 @@ -146,7 +147,17 @@ static void hipTestWithoutGraph() { } /** - * Simple test to demonstrate usage of graph. + * Test Description + * ------------------------ + * - Executes a simple kernel: + * -# When graph is executed + * -# When there is no graph + * Test source + * ------------------------ + * - unit/graph/hipSimpleGraphWithKernel.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGraph_SimpleGraphWithKernel") { // Sections run test with and without graph. diff --git a/catch/unit/graph/hipUserObjectCreate.cc b/catch/unit/graph/hipUserObjectCreate.cc index 2a22e62d9..830173315 100644 --- a/catch/unit/graph/hipUserObjectCreate.cc +++ b/catch/unit/graph/hipUserObjectCreate.cc @@ -23,17 +23,14 @@ THE SOFTWARE. #include "user_object_common.hh" -/* - Functional Test for API - hipUserObjectCreate -1) Call hipUserObjectCreate once and release it by calling hipUserObjectRelease -2) Call hipUserObjectCreate refCount as X and release it by calling - hipUserObjectRelease with same refCount. -3) Call hipUserObjectCreate, retain it by calling hipUserObjectRetain - and release it by calling hipUserObjectRelease twice. -4) Call hipUserObjectCreate with refCount as X, retain it by calling - hipUserObjectRetain with count as Y and release it by calling - hipUserObjectRelease with count as X+Y. -*/ +/** + * @addtogroup hipUserObjectCreate hipUserObjectCreate + * @{ + * @ingroup GraphTest + * `hipUserObjectCreate(hipUserObject_t* object_out, void* ptr, hipHostFn_t destroy, + * unsigned int initialRefcount, unsigned int flags)` - + * Create an instance of userObject to manage lifetime of a resource. + */ /* 1) Call hipUserObjectCreate once and release it by calling hipUserObjectRelease */ @@ -44,6 +41,23 @@ static void hipUserObjectCreate_Functional_1(void* object, void destroyObj(void* HIP_CHECK(hipUserObjectRelease(hObject)); } +/** + * Test Description + * ------------------------ + * - Creates user object from different types with ref count 1: + * -# When object is int + * -# When object is float + * -# When object is a class instance + * -# When structure is a structure instance + * - Checks that user object is not `nullptr`. + * - Releases user object with ref count 1. + * Test source + * ------------------------ + * - unit/graph/hipUserObjectCreate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipUserObjectCreate_Functional_1") { SECTION("Called with int Object") { int* object = new int(); @@ -78,6 +92,23 @@ static void hipUserObjectCreate_Functional_2(void* object, void destroyObj(void* HIP_CHECK(hipUserObjectRelease(hObject, refCount)); } +/** + * Test Description + * ------------------------ + * - Creates user object from different types with ref count greater than 1: + * -# When object is int + * -# When object is float + * -# When object is a class instance + * -# When structure is a structure instance + * - Checks that user object is not `nullptr`. + * - Releases user object with ref count greater than 1. + * Test source + * ------------------------ + * - unit/graph/hipUserObjectCreate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipUserObjectCreate_Functional_2") { SECTION("Called with int Object") { int* object = new int(); @@ -112,6 +143,24 @@ static void hipUserObjectCreate_Functional_3(void* object, void destroyObj(void* HIP_CHECK(hipUserObjectRelease(hObject)); } +/** + * Test Description + * ------------------------ + * - Creates user object from different types with ref count 1: + * -# When object is int + * -# When object is float + * -# When object is a class instance + * -# When structure is a structure instance + * - Checks that user object is not `nullptr`. + * - Retains user object once. + * - Releases user object twice with ref count 1. + * Test source + * ------------------------ + * - unit/graph/hipUserObjectCreate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipUserObjectCreate_Functional_3") { SECTION("Called with int Object") { int* object = new int(); @@ -149,6 +198,24 @@ static void hipUserObjectCreate_Functional_4(void* object, void destroyObj(void* HIP_CHECK(hipUserObjectRelease(hObject, refCount + refCountRetain)); } +/** + * Test Description + * ------------------------ + * - Creates user object from different types with ref count 5: + * -# When object is int + * -# When object is float + * -# When object is a class instance + * -# When structure is a structure instance + * - Checks that user object is not `nullptr`. + * - Retains user object with ref count 8. + * - Releases user object with ref count 13. + * Test source + * ------------------------ + * - unit/graph/hipUserObjectCreate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipUserObjectCreate_Functional_4") { SECTION("Called with int Object") { int* object = new int(); @@ -173,13 +240,27 @@ TEST_CASE("Unit_hipUserObjectCreate_Functional_4") { } /** - * Negative Test for API - hipUserObjectCreate - 1) Pass User Object as nullptr - 2) Pass object as nullptr - 3) Pass Callback function as nullptr - 4) Pass initialRefcount as 0 - 5) Pass initialRefcount as INT_MAX - 6) Pass flag other than hipUserObjectNoDestructorSync + * Test Description + * ------------------------ + * - Validates handling of invalid argument: + * -# When output pointer to the user object is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When object is `nullptr` + * - Expected output: return `hipSuccess` + * -# When destroy callback function handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When ref count is zero + * - Expected output: return `hipErrorInvalidValue` + * -# When ref count is INT_MAX + * - Expected output: return `hipSuccess` + * -# When flag is not valid + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/graph/hipUserObjectCreate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipUserObjectCreate_Negative") { int* object = new int(); @@ -215,6 +296,24 @@ TEST_CASE("Unit_hipUserObjectCreate_Negative") { } } +/** + * Test Description + * ------------------------ + * - Creates new int object. + * - Expected output: return `hipSuccess` + * - Creates user object with 2 references. + * - Expected output: return `hipSuccess` + * - Releases more user objects (4) than created (2). + * - Expected output: return `hipSuccess` + * - Retains reference to a removed user object. + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/graph/hipUserObjectCreate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipUserObj_Negative_Test") { int* object = new int(); REQUIRE(object != nullptr); diff --git a/catch/unit/graph/hipUserObjectRelease.cc b/catch/unit/graph/hipUserObjectRelease.cc index bb73fdeac..55391e838 100644 --- a/catch/unit/graph/hipUserObjectRelease.cc +++ b/catch/unit/graph/hipUserObjectRelease.cc @@ -25,10 +25,35 @@ THE SOFTWARE. #include "user_object_common.hh" /** - * Negative Test for API - hipUserObjectRelease - 1) Pass User Object as nullptr - 2) Pass initialRefcount as 0 - 3) Pass initialRefcount as INT_MAX + * @addtogroup hipUserObjectRelease hipUserObjectRelease + * @{ + * @ingroup GraphTest + * `hipUserObjectRelease(hipUserObject_t object, unsigned int count __dparm(1))` - + * Release number of references to resource. + * ________________________ + * Test cases from other modules: + * - @ref Unit_hipUserObjectCreate_Functional_1 + * - @ref Unit_hipUserObjectCreate_Functional_2 + * - @ref Unit_hipUserObjectCreate_Functional_3 + * - @ref Unit_hipUserObjectCreate_Functional_4 + */ + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When user object handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When ref count is zero + * - Expected output: return `hipErrorInvalidValue` + * -# When ref count is INT_MAX + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/graph/hipUserObjectRelease.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipUserObjectRelease_Negative") { int* object = new int(); diff --git a/catch/unit/graph/hipUserObjectRetain.cc b/catch/unit/graph/hipUserObjectRetain.cc index 895c89848..13dbcaa2c 100644 --- a/catch/unit/graph/hipUserObjectRetain.cc +++ b/catch/unit/graph/hipUserObjectRetain.cc @@ -25,10 +25,33 @@ THE SOFTWARE. #include "user_object_common.hh" /** - * Negative Test for API - hipUserObjectRetain - 1) Pass User Object as nullptr - 2) Pass initialRefcount as 0 - 3) Pass initialRefcount as INT_MAX + * @addtogroup hipUserObjectRetain hipUserObjectRetain + * @{ + * @ingroup GraphTest + * `hipUserObjectRetain(hipUserObject_t object, unsigned int count __dparm(1))` - + * Retain number of references to resource. + * ________________________ + * Test cases from other modules: + * - @ref Unit_hipUserObjectCreate_Functional_3 + * - @ref Unit_hipUserObjectCreate_Functional_4 + */ + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When user object handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When ref count is zero + * - Expected output: return `hipErrorInvalidValue` + * -# When ref count is INT_MAX + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/graph/hipUserObjectRetain.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipUserObjectRetain_Negative") { int* object = new int();