diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index 36f488f9fe..da44fdae25 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -60,6 +60,7 @@ "Unit_hipMemPoolApi_Opportunistic", "Unit_hipMemPoolApi_Default", "Unit_hipDeviceGetUuid", + "Unit_hipGraphMemcpyNodeSetParamsFromSymbol_Functional", "Unit_hipGraphExecEventWaitNodeSetEvent_Negative", "Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory" ] diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index e39efe7142..a8a37733d4 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -62,6 +62,7 @@ set(TEST_SRC hipStreamIsCapturing.cc hipStreamGetCaptureInfo.cc hipStreamEndCapture.cc + hipGraphMemcpyNodeSetParamsFromSymbol.cc hipGraphExecEventWaitNodeSetEvent.cc ) diff --git a/tests/catch/unit/graph/hipGraphMemcpyNodeSetParamsFromSymbol.cc b/tests/catch/unit/graph/hipGraphMemcpyNodeSetParamsFromSymbol.cc new file mode 100644 index 0000000000..495685f4a0 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphMemcpyNodeSetParamsFromSymbol.cc @@ -0,0 +1,259 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER 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 of hipGraphMemcpyNodeSetParamsFromSymbol API: + +Functional : +1) Allocate global symbol memory, add node to the graph. + Set/Update the new values to the node. Make sure they are taking effect. +2) Allocate const symbol memory, add node to the graph. + Set/Update the new values to the node. Make sure they are taking effect. + +Negative : +1) Pass GraphNode as nullptr and check if api returns error. +2) Pass destination ptr as nullptr, api expected to return error code. +3) Pass source/symbol ptr as nullptr, api expected to return error code. +4) Pass count as zero, api expected to return error code. +5) Pass count more than allocated size for source and destination ptr, api should return error code. +6) Pass offset+count greater than allocated size, api expected to return error code. +7) Pass same symbol pointer as destination ptr and source ptr, api expected to return error code. +8) Pass both destination ptr and source ptr as 2 different symbol ptr, api expected to return error code. +*/ + +#include +#include +#include +#define SIZE 256 + +__device__ int globalIn[SIZE]; +__device__ int globalOut[SIZE]; +__device__ __constant__ int globalConst[SIZE]; + + +/* Test verifies hipGraphMemcpyNodeSetParamsFromSymbol API Negative scenarios. + */ +TEST_CASE("Unit_hipGraphMemcpyNodeSetParamsFromSymbol_Negative") { + constexpr size_t Nbytes = SIZE * sizeof(int); + int *A_d{nullptr}, *B_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, &B_d, nullptr, + &A_h, &B_h, nullptr, SIZE, false); + + hipError_t ret; + hipGraph_t graph; + hipGraphNode_t memcpyToSymbolNode, memcpyFromSymbolNode, memcpyH2D_A; + std::vector dependencies; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyH2D_A); + + // Adding MemcpyNodeToSymbol + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + dependencies.clear(); + dependencies.push_back(memcpyToSymbolNode); + + HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + B_h, + HIP_SYMBOL(globalConst), + Nbytes, 0, + hipMemcpyDeviceToHost)); + SECTION("Pass GraphNode as nullptr") { + ret = hipGraphMemcpyNodeSetParamsFromSymbol(nullptr, B_h, + HIP_SYMBOL(globalConst), + Nbytes, 0, + hipMemcpyDeviceToHost); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass destination ptr as nullptr") { + ret = hipGraphMemcpyNodeSetParamsFromSymbol(memcpyFromSymbolNode, nullptr, + HIP_SYMBOL(globalConst), + Nbytes, 0, + hipMemcpyDeviceToHost); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass source/symbol ptr as nullptr") { + ret = hipGraphMemcpyNodeSetParamsFromSymbol(memcpyFromSymbolNode, B_h, + nullptr, + Nbytes, 0, + hipMemcpyDeviceToHost); + REQUIRE(hipErrorInvalidSymbol == ret); + } + SECTION("Pass count as zero") { + ret = hipGraphMemcpyNodeSetParamsFromSymbol(memcpyFromSymbolNode, B_h, + HIP_SYMBOL(globalConst), + 0, 0, + hipMemcpyDeviceToHost); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass count more than allocated size for source and dstn ptr") { + ret = hipGraphMemcpyNodeSetParamsFromSymbol(memcpyFromSymbolNode, B_h, + HIP_SYMBOL(globalConst), + Nbytes+10, 0, + hipMemcpyDeviceToHost); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass offset non zero so that offset+count > allocated size") { + ret = hipGraphMemcpyNodeSetParamsFromSymbol(memcpyFromSymbolNode, B_h, + HIP_SYMBOL(globalConst), + Nbytes, 10, + hipMemcpyDeviceToHost); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass same symbol pointer as dstn ptr and source ptr") { + ret = hipGraphMemcpyNodeSetParamsFromSymbol(memcpyFromSymbolNode, + HIP_SYMBOL(globalConst), + HIP_SYMBOL(globalConst), + Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass both dstn ptr and source ptr as 2 different symbol ptr") { + ret = hipGraphMemcpyNodeSetParamsFromSymbol(memcpyFromSymbolNode, + HIP_SYMBOL(globalOut), + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + HipTest::freeArrays(A_d, B_d, nullptr, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphDestroy(graph)); +} + +static +void hipGraphMemcpyNodeSetParamsFromSymbol_GlobalMem(bool useConstDeviceVar) { + constexpr size_t Nbytes = SIZE * sizeof(int); + hipGraphNode_t memcpyD2H_B; + int *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, &B_d, &C_d, + &A_h, &B_h, nullptr, SIZE, false); + + hipGraph_t graph; + hipGraphExec_t graphExec; + hipGraphNode_t memcpyToSymbolNode, memcpyFromSymbolNode, memcpyH2D_A; + std::vector dependencies; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyH2D_A); + + if (useConstDeviceVar) { + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalConst), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + } else { + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + } + dependencies.clear(); + dependencies.push_back(memcpyToSymbolNode); + + if (useConstDeviceVar) { + HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + C_d, + HIP_SYMBOL(globalConst), + Nbytes, 0, + hipMemcpyDeviceToDevice)); + } else { + HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + C_d, + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToDevice)); + } + dependencies.clear(); + dependencies.push_back(memcpyFromSymbolNode); + + // Update the node with B_d destination pointer from C_d + if (useConstDeviceVar) { + HIP_CHECK(hipGraphMemcpyNodeSetParamsFromSymbol(memcpyFromSymbolNode, + B_d, + HIP_SYMBOL(globalConst), + Nbytes, 0, + hipMemcpyDeviceToDevice)); + } else { + HIP_CHECK(hipGraphMemcpyNodeSetParamsFromSymbol(memcpyFromSymbolNode, + B_d, + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToDevice)); + } + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_B, graph, dependencies.data(), + dependencies.size(), B_h, B_d, + Nbytes, hipMemcpyDeviceToHost)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + + // Validating the result + for (int i = 0; i < SIZE; i++) { + if (B_h[i] != A_h[i]) { + WARN("Validation failed B_h[i] " << B_h[i] << "A_h[i] " << A_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, B_d, C_d, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/* Test verifies hipGraphMemcpyNodeSetParamsFromSymbol API Functional scenario. + 1) Allocate global symbol memory, add node to the graph. + Set/Update the new values to the node. Make sure they are taking effect. + 2) Allocate const symbol memory, add node to the graph. + Set/Update the new values to the node. Make sure they are taking effect. + */ +TEST_CASE("Unit_hipGraphMemcpyNodeSetParamsFromSymbol_Functional") { + SECTION("Check and update with Global Device Symbol Memory") { + hipGraphMemcpyNodeSetParamsFromSymbol_GlobalMem(false); + } + SECTION("Check and update with Constant Global Device Symbol Memory") { + hipGraphMemcpyNodeSetParamsFromSymbol_GlobalMem(true); + } +} diff --git a/tests/catch/unit/graph/hipStreamBeginCapture.cc b/tests/catch/unit/graph/hipStreamBeginCapture.cc index b1d4dfdd06..af550da759 100644 --- a/tests/catch/unit/graph/hipStreamBeginCapture.cc +++ b/tests/catch/unit/graph/hipStreamBeginCapture.cc @@ -18,7 +18,7 @@ THE SOFTWARE. */ /** -Testcase Scenarios : +Testcase Scenarios : Functional 1) Initiate stream capture with different modes on custom stream. Capture stream sequence and replay the sequence in multiple iterations. 2) End capture and validate that API returns captured graph for @@ -183,3 +183,46 @@ TEST_CASE("Unit_hipStreamBeginCapture_hipStreamPerThread") { HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipFree(C_d)); } + + +/* Test verifies hipStreamBeginCapture API Negative scenarios. + */ + +TEST_CASE("Unit_hipStreamBeginCapture_Negative") { + hipError_t ret; + hipStream_t stream{}; + HIP_CHECK(hipStreamCreate(&stream)); + + SECTION("Stream capture on legacy/null stream returns error code.") { + ret = hipStreamBeginCapture(nullptr, hipStreamCaptureModeGlobal); + REQUIRE(hipErrorStreamCaptureUnsupported == ret); + } + SECTION("Capturing hipStream status with same stream again") { + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + ret = hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal); + REQUIRE(hipErrorIllegalState == ret); + } + SECTION("Creating hipStream with invalid mode") { + ret = hipStreamBeginCapture(stream, hipStreamCaptureMode(-1)); + REQUIRE(hipErrorInvalidValue == ret); + } + HIP_CHECK(hipStreamDestroy(stream)); +} + +TEST_CASE("Unit_hipStreamBeginCapture_Basic") { + hipStream_t s1, s2, s3; + + HIP_CHECK(hipStreamCreate(&s1)); + HIP_CHECK(hipStreamBeginCapture(s1, hipStreamCaptureModeGlobal)); + + HIP_CHECK(hipStreamCreate(&s2)); + HIP_CHECK(hipStreamBeginCapture(s2, hipStreamCaptureModeThreadLocal)); + + HIP_CHECK(hipStreamCreate(&s3)); + HIP_CHECK(hipStreamBeginCapture(s3, hipStreamCaptureModeRelaxed)); + + HIP_CHECK(hipStreamDestroy(s1)); + HIP_CHECK(hipStreamDestroy(s2)); + HIP_CHECK(hipStreamDestroy(s3)); +} +