From f80f85f252b15952b00a13cdbf7ca5e68b7f48e6 Mon Sep 17 00:00:00 2001 From: music-dino <111048524+music-dino@users.noreply.github.com> Date: Mon, 6 Mar 2023 09:13:45 +0100 Subject: [PATCH] EXSWHTEC-193 - Implement new and update existing tests for the hipGraph*MemsetNode family of APIs (#12) - Implement positive and negative tests for hipGraphAddMemsetNode - Implement positive and negative tests for hipGraphMemsetNodeSetParams - Implement positive and negative tests for hipGraphExecMemsetNodeSetParams - Implement negative tests for hipGraphMemsetNodeGetParams - Add doxygen annotations - Add a dynamic section in positive test for memset for different dimensions [ROCm/hip-tests commit: 77a37371b014c74ea3be52cc24eff52ab02b24c8] --- .../config/config_amd_linux_common.json | 9 +- .../config/config_amd_windows_common.json | 15 +- .../graph/graph_memset_node_test_common.hh | 103 ++++++ .../catch/unit/graph/hipGraphAddMemsetNode.cc | 211 ++++++----- .../graph/hipGraphExecMemsetNodeSetParams.cc | 330 +++++++++--------- .../unit/graph/hipGraphMemsetNodeGetParams.cc | 166 +++------ .../unit/graph/hipGraphMemsetNodeSetParams.cc | 263 ++++++-------- 7 files changed, 562 insertions(+), 535 deletions(-) create mode 100644 projects/hip-tests/catch/unit/graph/graph_memset_node_test_common.hh diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json index f1a3283fb6..d14b0971fe 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json @@ -36,6 +36,13 @@ "Unit_hipGraphAddMemcpyNode1D_Negative_Basic", "Unit_hipStreamGetCaptureInfo_Nullstream_CaptureInfo", "intermittent issue: corrupted double-linked list", - "Unit_hipGraphRetainUserObject_Functional_2" + "Unit_hipGraphRetainUserObject_Functional_2", + "Note: Following four tests disabled due to defect - EXSWHTEC-203", + "Unit_hipGraphAddMemsetNode_Positive_Basic - uint16_t", + "Unit_hipGraphAddMemsetNode_Positive_Basic - uint32_t", + "Unit_hipGraphMemsetNodeSetParams_Positive_Basic - uint16_t", + "Unit_hipGraphMemsetNodeSetParams_Positive_Basic - uint32_t", + "Note: Test disabled due to defect - EXSWHTEC-207", + "Unit_hipGraphExecMemsetNodeSetParams_Negative_Updating_Non1D_Node" ] } diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json index 09eadeb673..67c76ef238 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json @@ -56,7 +56,7 @@ "Unit_hipStreamWaitEvent_DifferentStreams", "Unit_hipStreamQuery_WithFinishedWork", "Unit_hipDeviceGetCacheConfig_Positive_Basic", - "Unit_hipDeviceGetCacheConfig_Positive_Basic", + "Unit_hipDeviceGetCacheConfig_Positive_Basic", "Unit_hipDeviceGetCacheConfig_Positive_Threaded", "Unit_hipStreamValue_Wait32_Blocking_Mask_Gte", "Unit_hipStreamValue_Wait32_Blocking_Mask_Eq_1", @@ -115,6 +115,17 @@ "Unit_hipStreamQuery_WithFinishedWork", "Unit_hipLaunchHostFunc_Graph", "Unit_hipLaunchHostFunc_KernelHost", - "Unit_hipStreamSetCaptureDependencies_Positive_Functional" + "Unit_hipStreamSetCaptureDependencies_Positive_Functional", + "Note: Following four tests disabled due to defect - EXSWHTEC-203", + "Unit_hipGraphAddMemsetNode_Positive_Basic - uint16_t", + "Unit_hipGraphAddMemsetNode_Positive_Basic - uint32_t", + "Unit_hipGraphMemsetNodeSetParams_Positive_Basic - uint16_t", + "Unit_hipGraphMemsetNodeSetParams_Positive_Basic - uint32_t", + "Note: Test disabled due to defect - EXSWHTEC-207", + "Unit_hipGraphExecMemsetNodeSetParams_Negative_Updating_Non1D_Node", + "Unit_hipGraphExecMemsetNodeSetParams_Positive_Basic - uint8_t", + "Unit_hipGraphExecMemsetNodeSetParams_Positive_Basic - uint16_t", + "Unit_hipGraphExecMemsetNodeSetParams_Positive_Basic - uint32_t", + "Unit_hipGraphMemsetNodeSetParams_Positive_Basic - uint8_t" ] } diff --git a/projects/hip-tests/catch/unit/graph/graph_memset_node_test_common.hh b/projects/hip-tests/catch/unit/graph/graph_memset_node_test_common.hh new file mode 100644 index 0000000000..f4b957283e --- /dev/null +++ b/projects/hip-tests/catch/unit/graph/graph_memset_node_test_common.hh @@ -0,0 +1,103 @@ +/* +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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +#include +#include +#include + +template void GraphMemsetNodeCommonPositive(F f) { + const size_t width = GENERATE(1, 64, kPageSize / sizeof(T) + 1); + const size_t height = GENERATE(1, 2, 1024); + DYNAMIC_SECTION("Width: " << width << " Height: " << height) { + LinearAllocGuard2D alloc(width, height); + + constexpr T set_value = 42; + hipMemsetParams params = {}; + params.dst = alloc.ptr(); + params.elementSize = sizeof(T); + params.width = width; + params.height = height; + params.pitch = alloc.pitch(); + params.value = set_value; + + HIP_CHECK(f(¶ms)); + + LinearAllocGuard buffer(LinearAllocs::hipHostMalloc, width * sizeof(T) * height); + HIP_CHECK(hipMemcpy2D(buffer.ptr(), width * sizeof(T), alloc.ptr(), alloc.pitch(), + width * sizeof(T), height, hipMemcpyDeviceToHost)); + ArrayFindIfNot(buffer.ptr(), set_value, width * height); + } +} + +template void MemsetCommonNegative(F f, hipMemsetParams params) { + SECTION("pMemsetParams == nullptr") { HIP_CHECK_ERROR(f(nullptr), hipErrorInvalidValue); } + + SECTION("pMemsetParams.dst == nullptr") { + params.dst = nullptr; + HIP_CHECK_ERROR(f(¶ms), hipErrorInvalidValue); + } + + SECTION("pMemsetParams.elementSize != 1, 2, 4") { + params.elementSize = GENERATE(0, 3, 5); + HIP_CHECK_ERROR(f(¶ms), hipErrorInvalidValue); + } + +// Disabled on AMD due to defect - EXSWHTEC-204 +#if HT_NVIDIA + SECTION("pMemsetParams.width == 0") { + params.width = 0; + HIP_CHECK_ERROR(f(¶ms), hipErrorInvalidValue); + } +#endif + + SECTION("pMemsetParams.width > allocation size") { + params.width = params.width + 1000; + HIP_CHECK_ERROR(f(¶ms), hipErrorInvalidValue); + } + + SECTION("pMemsetParams.height == 0") { + params.height = 0; + HIP_CHECK_ERROR(f(¶ms), hipErrorInvalidValue); + } + +// Disabled on AMD due to defect - EXSWHTEC-205 +#if HT_NVIDIA + SECTION("pMemsetParams.pitch < width when height > 1") { + params.width = 2; + params.height = 2; + params.pitch = params.elementSize; + HIP_CHECK_ERROR(f(¶ms), hipErrorInvalidValue); + } +#endif + +// Disabled on AMD due to defect - EXSWHTEC-206 +#if HT_NVIDIA + SECTION("pMemsetParams.pitch * height > allocation size") { + params.width = 2; + params.height = 2; + params.pitch = 3 * params.elementSize; + HIP_CHECK_ERROR(f(¶ms), hipErrorInvalidValue); + } +#endif +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/graph/hipGraphAddMemsetNode.cc b/projects/hip-tests/catch/unit/graph/hipGraphAddMemsetNode.cc index c89d1a6f5c..1080741731 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphAddMemsetNode.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphAddMemsetNode.cc @@ -6,132 +6,121 @@ 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 + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY 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 +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER 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 +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Negative Testcase Scenarios for api hipGraphAddMemsetNode : -1) Pass pGraphNode as nullptr and check if api returns error. -2) Pass pGraphNode as un-initialize object and check. -3) Pass Graph as nullptr and check if api returns error. -4) Pass Graph as empty object(skipping graph creation), api should return error code. -5) Pass pDependencies as nullptr, api should return success. -6) Pass numDependencies is max(size_t) and pDependencies is not valid ptr, api expected to return error code. -7) Pass pDependencies is nullptr, but numDependencies is non-zero, api expected to return error. -8) Pass pMemsetParams as nullptr and check if api returns error code. -9) Pass pMemsetParams as un-initialize object and check if api returns error code. -10) Pass hipMemsetParams::dst as nullptr should return error code. -11) Pass hipMemsetParams::element size other than 1, 2, or 4 and check api should return error code. -12) Pass hipMemsetParams::height as zero and check api should return error code. -Functional Scenarios for api hipGraphAddMemsetNode : -1. Allocate a 2D array using hipMallocPitch. Initialize the allocated memory using hipGraphAddMemsetNode. - Copy the values in device memory to host using hipGraphAddMemcpyNode. Verify the results -2. Allocate a 1D array using hipMallocPitch. Initialize the allocated memory using hipGraphAddMemsetNode. - Copy the values in device memory to host using hipGraphAddMemcpyNode. Verify the results.. -3. Allocate a 2D array using hipMalloc3D. Initialize the allocated memory using hipGraphAddMemsetNode. - Copy the values in device memory to host using hipGraphAddMemcpyNode. Verify the results. -4. Allocate a 1D array using hipMalloc3D. Initialize the allocated memory using hipGraphAddMemsetNode. - Copy the values in device memory to host using hipGraphAddMemcpyNode. Verify the results. -5. Allocate a 1D array using hipMalloc. Initialize the allocated memory using hipGraphAddMemsetNode. - Copy the values in device memory to host using hipGraphAddMemcpyNode. Verify the results. -6. Allocate memory using hipMallocManaged. Initialize the allocated memory using hipGraphAddMemsetNode. - Copy the values in device memory to host using hipGraphAddMemcpyNode. Verify the results. -*/ +#include +#include +#include #include -/** - * Negative Test for API hipGraphAddMemsetNode - */ +#include +#include + +#include "graph_memset_node_test_common.hh" +#include "graph_tests_common.hh" + #define SIZE 1024 static char memSetVal = 'a'; -TEST_CASE("Unit_hipGraphAddMemsetNode_Negative") { - hipError_t ret; - hipGraph_t graph; - hipGraphNode_t memsetNode; - char *devData; - HIP_CHECK(hipMalloc(&devData, 1024)); +/** + * @addtogroup hipGraphAddMemsetNode hipGraphAddMemsetNode + * @{ + * @ingroup GraphTest + * `hipGraphAddMemsetNode(hipGraphNode_t *pGraphNode, hipGraph_t graph, const hipGraphNode_t + * *pDependencies, size_t numDependencies, const hipMemsetParams *pMemsetParams)` - + * Creates a memset node and adds it to a graph + */ + +/** + * Test Description + * ------------------------ + * - Verify that all elements of destination memory are set to the correct value. + * The test is repeated for all valid element sizes(1, 2, 4), and several allocations of different + * height and width, both on host and device. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddMemsetNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_hipGraphAddMemsetNode_Positive_Basic", "", uint8_t, uint16_t, uint32_t) { + const auto f = [](hipMemsetParams* params) { + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + hipGraphNode_t node = nullptr; + HIP_CHECK(hipGraphAddMemsetNode(&node, graph, nullptr, 0, params)); + + hipGraphExec_t graph_exec = nullptr; + HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); + + HIP_CHECK(hipGraphLaunch(graph_exec, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + + HIP_CHECK(hipGraphExecDestroy(graph_exec)); + HIP_CHECK(hipGraphDestroy(graph)); + + return hipSuccess; + }; + + GraphMemsetNodeCommonPositive(f); +} + +/** + * Test Description + * ------------------------ + * - Verify API behaviour with invalid arguments: + * -# pGraphNode is nullptr + * -# graph is nullptr + * -# pDependencies is nullptr when numDependencies is not zero + * -# A node in pDependencies originates from a different graph + * -# numDependencies is invalid + * -# A node is duplicated in pDependencies + * -# pMemsetParams is nullptr + * -# pMemsetParams::dst is nullptr + * -# pMemsetParams::elementSize is different from 1, 2, and 4 + * -# pMemsetParams::width is zero + * -# pMemsetParams::width is larger than the allocated memory region + * -# pMemsetParams::height is zero + * -# pMemsetParams::pitch is less than width when height is more than 1 + * -# pMemsetParams::pitch * pMemsetParams::height is larger than the allocated memory region + * Test source + * ------------------------ + * - unit/graph/hipGraphAddMemsetNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipGraphAddMemsetNode_Negative_Parameters") { + using namespace std::placeholders; + hipGraph_t graph = nullptr; HIP_CHECK(hipGraphCreate(&graph, 0)); - hipMemsetParams memsetParams{}; - memset(&memsetParams, 0, sizeof(memsetParams)); - memsetParams.dst = reinterpret_cast(devData); - memsetParams.value = 0; - memsetParams.pitch = 0; - memsetParams.elementSize = sizeof(char); - memsetParams.width = 1024; - memsetParams.height = 1; + LinearAllocGuard alloc(LinearAllocs::hipMalloc, 4 * sizeof(int)); + hipMemsetParams params = {}; + params.dst = alloc.ptr(); + params.elementSize = sizeof(*alloc.ptr()); + params.width = 1; + params.height = 1; + params.value = 42; + + GraphAddNodeCommonNegativeTests(std::bind(hipGraphAddMemsetNode, _1, _2, _3, _4, ¶ms), graph); + + hipGraphNode_t node = nullptr; + MemsetCommonNegative(std::bind(hipGraphAddMemsetNode, &node, graph, nullptr, 0, _1), params); - SECTION("Pass pGraphNode as nullptr") { - ret = hipGraphAddMemsetNode(nullptr, graph, nullptr, 0, &memsetParams); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass pGraphNode as un-initialize object") { - hipGraphNode_t memsetNode_1; - ret = hipGraphAddMemsetNode(&memsetNode_1, graph, - nullptr, 0, &memsetParams); - REQUIRE(hipSuccess == ret); - } - SECTION("Pass graph as nullptr") { - ret = hipGraphAddMemsetNode(&memsetNode, nullptr, - nullptr, 0, &memsetParams); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass Graph as empty object") { - hipGraph_t graph_1{}; - ret = hipGraphAddMemsetNode(&memsetNode, graph_1, - nullptr, 0, &memsetParams); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass pDependencies as nullptr") { - ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams); - REQUIRE(hipSuccess == ret); - } - SECTION("Pass numDependencies is max and pDependencies is not valid ptr") { - ret = hipGraphAddMemsetNode(&memsetNode, graph, - nullptr, INT_MAX, &memsetParams); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass pDependencies as nullptr, but numDependencies is non-zero") { - ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 9, &memsetParams); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass pMemsetParams as nullptr") { - ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, nullptr); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass pMemsetParams as un-initialize object") { - hipMemsetParams memsetParams1; - ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, - &memsetParams1); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass hipMemsetParams::dst as nullptr") { - memsetParams.dst = nullptr; - ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass hipMemsetParams::element size other than 1, 2, or 4") { - memsetParams.dst = reinterpret_cast(devData); - memsetParams.elementSize = 9; - ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass hipMemsetParams::height as zero") { - memsetParams.elementSize = sizeof(char); - memsetParams.height = 0; - ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams); - REQUIRE(hipErrorInvalidValue == ret); - } - HIP_CHECK(hipFree(devData)); HIP_CHECK(hipGraphDestroy(graph)); } /* diff --git a/projects/hip-tests/catch/unit/graph/hipGraphExecMemsetNodeSetParams.cc b/projects/hip-tests/catch/unit/graph/hipGraphExecMemsetNodeSetParams.cc index 7e51eb296b..72ea020ca8 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphExecMemsetNodeSetParams.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphExecMemsetNodeSetParams.cc @@ -6,196 +6,208 @@ 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 + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY 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 +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER 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 +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : -Functional- -1) Instantiate a graph with memset node, obtain executable graph and update the - hipMemsetParams node params with set. Make sure they are taking effect. -Negative- -1) Pass hGraphExec as nullptr and verify api returns error code. -2) Pass graph node as nullptr and verify api returns error code. -3) Pass different hipGraphNode_t which was not used in graphExec and verify api returns error code. -4) Pass Pass different Graph which was not used in graphExec and verify api returns error code. -5) Pass pNodeParams as nullptr and verify api returns error code. -6) Pass pNodeParams as empty structure object and verify api returns error code. -7) Pass hipMemsetParams::dst as nullptr, api should return error code. -8) Pass hipMemsetParams::element size other than 1, 2, or 4 and check api should return error code. -9) Pass hipMemsetParams::height as zero and check api should return error code. -*/ +#include +#include #include -#include -/* Test verifies hipGraphExecMemsetNodeSetParams API Negative scenarios. +#include "graph_memset_node_test_common.hh" +#include "graph_tests_common.hh" + +/** + * @addtogroup hipGraphExecMemsetNodeSetParams hipGraphExecMemsetNodeSetParams + * @{ + * @ingroup GraphTest + * `hipGraphExecMemsetNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node, const + * hipMemsetParams *pNodeParams)` - + * Sets the parameters for a memset node in the given graphExec */ -TEST_CASE("Unit_hipGraphExecMemsetNodeSetParams_Negative") { - constexpr size_t N = 1024; - constexpr size_t Nbytes = N * sizeof(char); - constexpr size_t val = 0; - char *devData, *hOutputData; - HIP_CHECK(hipMalloc(&devData, Nbytes)); - hOutputData = reinterpret_cast(malloc(Nbytes)); - REQUIRE(hOutputData != nullptr); - memset(hOutputData, 0, Nbytes); - - hipGraph_t graph; - hipError_t ret; - hipGraphExec_t graphExec; - hipStream_t streamForGraph; - hipGraphNode_t memsetNode; +/** + * Test Description + * ------------------------ + * - Verify that node parameters get updated correctly by creating a node with valid but + * incorrect parameters, and then setting them to the correct values in the executable graph. + * The executable graph is run and the results of the memset is verified. + * hipGraphMemsetNodeGetParams is used to verify that node parameters in the graph were not updated, + * which also constitutes a test for said API. + * The test is repeated for all valid element sizes(1, + * 2, 4), and several allocations of different width(height is always 1 because only 1D memset nodes + * can be updated), both on host and device + * Test source + * ------------------------ + * - unit/graph/hipGraphExecMemsetNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_hipGraphExecMemsetNodeSetParams_Positive_Basic", "", uint8_t, uint16_t, + uint32_t) { + const size_t width = GENERATE(1, 64, kPageSize / sizeof(TestType) + 1); + hipGraph_t graph = nullptr; HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipStreamCreate(&streamForGraph)); - hipMemsetParams mParams{}; - memset(&mParams, 0, sizeof(mParams)); - mParams.dst = reinterpret_cast(devData); - mParams.value = val; - mParams.pitch = 0; - mParams.elementSize = sizeof(char); - mParams.width = Nbytes; - mParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &mParams)); + hipGraphNode_t node = nullptr; + LinearAllocGuard initial_alloc(LinearAllocs::hipMalloc, 2 * sizeof(TestType)); - std::vector dependencies; - dependencies.push_back(memsetNode); + hipMemsetParams initial_params = {}; + initial_params.dst = initial_alloc.ptr(); + initial_params.elementSize = sizeof(TestType); + initial_params.width = 2; + initial_params.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&node, graph, nullptr, 0, &initial_params)); - HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + hipGraphExec_t graph_exec = nullptr; + HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); - SECTION("Pass hGraphExec as nullptr") { - ret = hipGraphExecMemsetNodeSetParams(nullptr, memsetNode, &mParams); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass hGraphNode as nullptr") { - ret = hipGraphExecMemsetNodeSetParams(graphExec, nullptr, &mParams); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass different hGraphNode which was not used in graphExec") { - hipGraphNode_t memsetNode1{}; - ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode1, &mParams); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass different Graph which was not used in graphExec") { - hipGraph_t graph1; - HIP_CHECK(hipGraphCreate(&graph1, 0)); - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph1, nullptr, 0, &mParams)); - ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParams); - REQUIRE(hipErrorInvalidValue == ret); - HIP_CHECK(hipGraphDestroy(graph1)); - } - SECTION("Pass pNodeParams as nullptr") { - ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, nullptr); - REQUIRE(hipErrorInvalidValue == ret); - } -#if HT_NVIDIA - SECTION("Pass pNodeParams as empty structure object") { - hipMemsetParams mParmTemp{}; - ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParmTemp); - REQUIRE(hipErrorInvalidValue == ret); - } -#endif - SECTION("Pass hipMemsetParams::dst as nullptr") { - mParams.dst = nullptr; - ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParams); - REQUIRE(hipErrorInvalidValue == ret); - } -#if HT_NVIDIA - SECTION("Pass hipMemsetParams::element size other than 1, 2, or 4") { - mParams.dst = reinterpret_cast(devData); - mParams.elementSize = 9; - ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParams); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass hipMemsetParams::height as zero") { - mParams.elementSize = sizeof(char); - mParams.height = 0; - ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParams); - REQUIRE(hipErrorInvalidValue == ret); - } -#endif + LinearAllocGuard2D alloc(width, 1); + constexpr TestType set_value = 42; + hipMemsetParams params = {}; + params.dst = alloc.ptr(); + params.elementSize = sizeof(TestType); + params.width = width; + params.height = 1; + params.value = set_value; + HIP_CHECK(hipGraphExecMemsetNodeSetParams(graph_exec, node, ¶ms)); - free(hOutputData); - HIP_CHECK(hipFree(devData)); - HIP_CHECK(hipGraphExecDestroy(graphExec)); + hipMemsetParams retrieved_params = {}; + HIP_CHECK(hipGraphMemsetNodeGetParams(node, &retrieved_params)); + REQUIRE(initial_params.dst == retrieved_params.dst); + REQUIRE(initial_params.elementSize == retrieved_params.elementSize); + REQUIRE(initial_params.width == retrieved_params.width); + REQUIRE(initial_params.height == retrieved_params.height); + REQUIRE(initial_params.pitch == retrieved_params.pitch); + REQUIRE(initial_params.value == retrieved_params.value); + + HIP_CHECK(hipGraphLaunch(graph_exec, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + + HIP_CHECK(hipGraphExecDestroy(graph_exec)); HIP_CHECK(hipGraphDestroy(graph)); - HIP_CHECK(hipStreamDestroy(streamForGraph)); + + LinearAllocGuard buffer(LinearAllocs::hipHostMalloc, width * sizeof(TestType)); + HIP_CHECK(hipMemcpy2D(buffer.ptr(), width * sizeof(TestType), alloc.ptr(), alloc.pitch(), + width * sizeof(TestType), 1, hipMemcpyDeviceToHost)); + ArrayFindIfNot(buffer.ptr(), set_value, width); } -/* Test verifies hipGraphExecMemsetNodeSetParams API Functional scenarios. +/** + * Test Description + * ------------------------ + * - Verify API behaviour with invalid arguments: + * -# pGraphExec is nullptr + * -# node is nullptr + * -# pNodeParams is nullptr + * -# pNodeParams::dst is nullptr + * -# pNodeParams::elementSize is different from 1, 2, and 4 + * -# pNodeParams::width is zero + * -# pNodeParams::width is larger than the allocated memory region + * -# pNodeParams::height is zero + * -# pNodeParams::pitch is less than width when height is more than 1 + * -# pNodeParams::pitch * pMemsetParams::height is larger than the allocated memory region + * -# pNodeParams::dst holds a pointer to memory allocated on a device different from the one + * the original dst was allocated on + * Test source + * ------------------------ + * - unit/graph/hipGraphExecMemsetNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ -TEST_CASE("Unit_hipGraphExecMemsetNodeSetParams_Functional") { - constexpr size_t N = 1024; - constexpr size_t Nbytes = N * sizeof(char); - constexpr size_t val = 0; - constexpr size_t updateVal = 2; - char *devData, *devData1, *hOutputData, *hOutputData1; - - HIP_CHECK(hipMalloc(&devData, Nbytes)); - HIP_CHECK(hipMalloc(&devData1, Nbytes)); - hOutputData = reinterpret_cast(malloc(Nbytes)); - REQUIRE(hOutputData != nullptr); - memset(hOutputData, updateVal, Nbytes); - hOutputData1 = reinterpret_cast(malloc(Nbytes)); - REQUIRE(hOutputData1 != nullptr); - memset(hOutputData1, 0, Nbytes); - - hipGraph_t graph; - hipGraphExec_t graphExec; - hipStream_t streamForGraph; - hipGraphNode_t memsetNode; +TEST_CASE("Unit_hipGraphExecMemsetNodeSetParams_Negative_Parameters") { + using namespace std::placeholders; + hipGraph_t graph = nullptr; HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipStreamCreate(&streamForGraph)); - hipMemsetParams memsetParams{}; - memset(&memsetParams, 0, sizeof(memsetParams)); - memsetParams.dst = reinterpret_cast(devData); - memsetParams.value = val; - memsetParams.pitch = 0; - memsetParams.elementSize = sizeof(char); - memsetParams.width = Nbytes; - memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, - &memsetParams)); + LinearAllocGuard alloc(LinearAllocs::hipMalloc, 4 * sizeof(int)); + hipMemsetParams params = {}; + params.dst = alloc.ptr(); + params.elementSize = sizeof(*alloc.ptr()); + params.width = 1; + params.height = 1; + params.value = 42; - std::vector dependencies; - dependencies.push_back(memsetNode); + hipGraphNode_t node = nullptr; + HIP_CHECK(hipGraphAddMemsetNode(&node, graph, nullptr, 0, ¶ms)) - HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + hipGraphExec_t graph_exec = nullptr; + HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); - memset(&memsetParams, 0, sizeof(memsetParams)); - memsetParams.dst = reinterpret_cast(devData1); - memsetParams.value = updateVal; - memsetParams.pitch = 0; - memsetParams.elementSize = sizeof(char); - memsetParams.width = Nbytes; - memsetParams.height = 1; + SECTION("pGraphExec == nullptr") { + HIP_CHECK_ERROR(hipGraphExecMemsetNodeSetParams(nullptr, node, ¶ms), hipErrorInvalidValue); + } - REQUIRE(hipSuccess == hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, - &memsetParams)); - HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); - HIP_CHECK(hipStreamSynchronize(streamForGraph)); + SECTION("node == nullptr") { + HIP_CHECK_ERROR(hipGraphExecMemsetNodeSetParams(graph_exec, nullptr, ¶ms), + hipErrorInvalidValue); + } - HIP_CHECK(hipMemcpy(hOutputData1, devData1, Nbytes, hipMemcpyDeviceToHost)); - HipTest::checkArray(hOutputData, hOutputData1, Nbytes, 1); + MemsetCommonNegative(std::bind(hipGraphExecMemsetNodeSetParams, graph_exec, node, _1), params); - free(hOutputData); - free(hOutputData1); - HIP_CHECK(hipFree(devData)); - HIP_CHECK(hipFree(devData1)); - HIP_CHECK(hipGraphExecDestroy(graphExec)); + SECTION("Changing dst allocation device") { + if (HipTest::getDeviceCount() < 2) { + HipTest::HIP_SKIP_TEST("Test requires two connected GPUs"); + return; + } + HIP_CHECK(hipSetDevice(1)); + LinearAllocGuard new_alloc(LinearAllocs::hipMalloc, 4 * sizeof(int)); + params.dst = new_alloc.ptr(); + HIP_CHECK_ERROR(hipGraphExecMemsetNodeSetParams(graph_exec, node, ¶ms), + hipErrorInvalidValue); + } + + HIP_CHECK(hipGraphExecDestroy(graph_exec)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/** + * Test Description + * ------------------------ + * - Verify that a 2D node cannot be updated + * Test source + * ------------------------ + * - unit/graph/hipGraphExecMemsetNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipGraphExecMemsetNodeSetParams_Negative_Updating_Non1D_Node") { + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + LinearAllocGuard2D alloc(2, 2); + hipMemsetParams params = {}; + params.dst = alloc.ptr(); + params.elementSize = sizeof(*alloc.ptr()); + params.width = 1; + params.height = 2; + params.pitch = alloc.pitch(); + params.value = 42; + + hipGraphNode_t node = nullptr; + HIP_CHECK(hipGraphAddMemsetNode(&node, graph, nullptr, 0, ¶ms)) + + hipGraphExec_t graph_exec = nullptr; + HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); + + params.width = 2; + HIP_CHECK_ERROR(hipGraphExecMemsetNodeSetParams(graph_exec, node, ¶ms), hipErrorInvalidValue); + + HIP_CHECK(hipGraphExecDestroy(graph_exec)); HIP_CHECK(hipGraphDestroy(graph)); - HIP_CHECK(hipStreamDestroy(streamForGraph)); } diff --git a/projects/hip-tests/catch/unit/graph/hipGraphMemsetNodeGetParams.cc b/projects/hip-tests/catch/unit/graph/hipGraphMemsetNodeGetParams.cc index 03cbcaaf2c..b776b9171e 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphMemsetNodeGetParams.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphMemsetNodeGetParams.cc @@ -6,8 +6,10 @@ in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions: + The above copyright notice and this permission notice shall be included in all copies or substantial portions of the Software. + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE @@ -17,125 +19,65 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase 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 Memset node to graph with desired node params. - Verify api fetches the node params mentioned while adding Memset node. -2) Set Memset node params with hipGraphMemsetNodeSetParams, - now get the params and verify both are same. -*/ - +#include #include +#include -/* Test verifies hipGraphMemsetNodeGetParams API Negative scenarios. - */ -TEST_CASE("Unit_hipGraphMemsetNodeGetParams_Negative") { - hipError_t ret; - hipGraph_t graph; - hipGraphNode_t memsetNode; - - HIP_CHECK(hipGraphCreate(&graph, 0)); - - char *devData; - HIP_CHECK(hipMalloc(&devData, 1024)); - hipMemsetParams memsetParams{}; - memset(&memsetParams, 0, sizeof(memsetParams)); - memsetParams.dst = reinterpret_cast(devData); - memsetParams.value = 0; - memsetParams.pitch = 0; - memsetParams.elementSize = sizeof(char); - memsetParams.width = 1024; - memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, - &memsetParams)); - SECTION("Pass node as nullptr") { - ret = hipGraphMemsetNodeGetParams(nullptr, &memsetParams); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass GetNodeParams as nullptr") { - ret = hipGraphMemsetNodeGetParams(memsetNode, nullptr); - REQUIRE(hipErrorInvalidValue == ret); - } - HIP_CHECK(hipFree(devData)); - HIP_CHECK(hipGraphDestroy(graph)); -} - -/* Test verifies hipGraphMemsetNodeGetParams API Functional scenarios. +/** + * @addtogroup hipGraphMemsetNodeGetParams hipGraphMemsetNodeGetParams + * @{ + * @ingroup GraphTest + * `hipGraphMemsetNodeGetParams(hipGraphNode_t node, hipMemsetParams *pNodeParams)` - + * Gets a memset node's parameters + * ________________________ + * Test cases from other APIs: + * - @ref Unit_hipGraphMemsetNodeSetParams_Positive_Basic + * - @ref Unit_hipGraphExecMemsetNodeSetParams_Positive_Basic */ -bool memsetNodeCompare(hipMemsetParams *mNode1, hipMemsetParams *mNode2) { - if (mNode1->dst != mNode2->dst) - return false; - if (mNode1->elementSize != mNode2->elementSize) - return false; - if (mNode1->height != mNode2->height) - return false; - if (mNode1->pitch != mNode2->pitch) - return false; - if (mNode1->value != mNode2->value) - return false; - if (mNode1->width != mNode2->width) - return false; - return true; -} +/** + * Test Description + * ------------------------ + * - Verify API behaviour with invalid arguments: + * -# node is nullptr + * -# pNodeParams is nullptr + * -# node is destroyed + * Test source + * ------------------------ + * - unit/graph/hipGraphMemsetNodeGetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipGraphMemsetNodeGetParams_Negative_Parameters") { + LinearAllocGuard2D alloc(1, 1); + hipMemsetParams params = {}; + params.dst = alloc.ptr(); + params.elementSize = sizeof(int); + params.width = 1; + params.height = 1; -TEST_CASE("Unit_hipGraphMemsetNodeGetParams_Functional") { - constexpr size_t N = 1024; - constexpr size_t Nbytes = N * sizeof(char); - constexpr size_t val = 0; + hipGraph_t graph = nullptr; + hipGraphNode_t node = nullptr; - char *devData; - HIP_CHECK(hipMalloc(&devData, Nbytes)); - - hipGraph_t graph; - hipGraphNode_t memsetNode; - HIP_CHECK(hipGraphCreate(&graph, 0)); - hipStream_t streamForGraph; - HIP_CHECK(hipStreamCreate(&streamForGraph)); - - hipMemsetParams memsetParams{}; - memset(&memsetParams, 0, sizeof(memsetParams)); - memsetParams.dst = reinterpret_cast(devData); - memsetParams.value = val; - memsetParams.pitch = 0; - memsetParams.elementSize = sizeof(char); - memsetParams.width = Nbytes; - memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, - &memsetParams)); - SECTION("Get Memset Param and verify.") { - hipMemsetParams memsetGetParams; - REQUIRE(hipSuccess == hipGraphMemsetNodeGetParams(memsetNode, - &memsetGetParams)); - // Validating the result - REQUIRE(true == memsetNodeCompare(&memsetParams, &memsetGetParams)); + SECTION("node == nullptr") { + HIP_CHECK_ERROR(hipGraphMemsetNodeGetParams(nullptr, ¶ms), hipErrorInvalidValue); } - SECTION("Set memset node params then Get and verify.") { - constexpr size_t updateVal = 2; - char *devData1; - HIP_CHECK(hipMalloc(&devData1, Nbytes)); - memset(&memsetParams, 0, sizeof(memsetParams)); - memsetParams.dst = reinterpret_cast(devData1); - memsetParams.value = updateVal; - memsetParams.pitch = 0; - memsetParams.elementSize = sizeof(char); - memsetParams.width = Nbytes; - memsetParams.height = 1; - hipMemsetParams memsetGetParams; - REQUIRE(hipSuccess == hipGraphMemsetNodeSetParams(memsetNode, - &memsetParams)); - REQUIRE(hipSuccess == hipGraphMemsetNodeGetParams(memsetNode, - &memsetGetParams)); - // Validating the result - REQUIRE(true == memsetNodeCompare(&memsetParams, &memsetGetParams)); - HIP_CHECK(hipFree(devData1)); + SECTION("pNodeParams == nullptr") { + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemsetNode(&node, graph, nullptr, 0, ¶ms)); + HIP_CHECK_ERROR(hipGraphMemsetNodeGetParams(node, nullptr), hipErrorInvalidValue); + HIP_CHECK(hipGraphDestroy(graph)); } - HIP_CHECK(hipFree(devData)); - HIP_CHECK(hipGraphDestroy(graph)); - HIP_CHECK(hipStreamDestroy(streamForGraph)); + +// Disabled on AMD due to defect - EXSWHTEC-208 +#if 0 + SECTION("Node is destroyed") { + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemsetNode(&node, graph, nullptr, 0, ¶ms)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK_ERROR(hipGraphMemsetNodeGetParams(node, ¶ms), hipErrorInvalidValue); + } +#endif } diff --git a/projects/hip-tests/catch/unit/graph/hipGraphMemsetNodeSetParams.cc b/projects/hip-tests/catch/unit/graph/hipGraphMemsetNodeSetParams.cc index fe188d48c8..937eaa06dd 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphMemsetNodeSetParams.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphMemsetNodeSetParams.cc @@ -6,8 +6,10 @@ in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions: + The above copyright notice and this permission notice shall be included in all copies or substantial portions of the Software. + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE @@ -17,169 +19,130 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : -Negative - -1) Pass node as nullptr and verify api returns error code. -2) Pass pNodeParams as nullptr and verify api returns error code. -3) Passing hipMemsetParams::dst as nullptr should return error code. -4) Passing hipMemsetParams::element size other than 1, 2, or 4 and check. -5) Passing hipMemsetParams::height as zero and check api should return error code. -Functional - -1) Add Memset node to graph, update the node params with set and - launch the graph and check the set params are executing properly. -2) Add Memset node to graph, launch graph, then update the Memset node params - with set and launch the graph and check updated params are taking effect. -*/ +#include +#include #include -#include -/* Test verifies hipGraphMemsetNodeSetParams API invalid params scenarios. +#include "graph_memset_node_test_common.hh" +#include "graph_tests_common.hh" + +/** + * @addtogroup hipGraphMemsetNodeSetParams hipGraphMemsetNodeSetParams + * @{ + * @ingroup GraphTest + * `hipGraphMemsetNodeSetParams(hipGraphNode_t node, const hipMemsetParams *pNodeParams)` - + * Sets a memset node's parameters */ -TEST_CASE("Unit_hipGraphMemsetNodeSetParams_InvalidParams") { - hipError_t ret; - hipGraph_t graph; - hipStream_t streamForGraph; - hipGraphNode_t memsetNode; - HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipStreamCreate(&streamForGraph)); - - char *devData; - HIP_CHECK(hipMalloc(&devData, 1024)); - hipMemsetParams memsetParams{}; - memset(&memsetParams, 0, sizeof(memsetParams)); - memsetParams.dst = reinterpret_cast(devData); - memsetParams.value = 0; - memsetParams.pitch = 0; - memsetParams.elementSize = sizeof(char); - memsetParams.width = 1024; - memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, - &memsetParams)); - SECTION("Pass node as nullptr") { - ret = hipGraphMemsetNodeSetParams(nullptr, &memsetParams); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass GetNodeParams as nullptr") { - ret = hipGraphMemsetNodeSetParams(memsetNode, nullptr); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass dest as nullptr") { - memsetParams.dst = nullptr; - ret = hipGraphMemsetNodeSetParams(memsetNode, &memsetParams); - REQUIRE(hipErrorInvalidValue == ret); - } -#if HT_NVIDIA - SECTION("Pass element size other than 1, 2, or 4") { - memsetParams.dst = reinterpret_cast(devData); - memsetParams.elementSize = 9; - ret = hipGraphMemsetNodeSetParams(memsetNode, &memsetParams); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass height as zero or negative") { - memsetParams.elementSize = 2; - memsetParams.height = 0; - ret = hipGraphMemsetNodeSetParams(memsetNode, &memsetParams); - REQUIRE(hipErrorInvalidValue == ret); - } -#endif - HIP_CHECK(hipFree(devData)); - HIP_CHECK(hipGraphDestroy(graph)); - HIP_CHECK(hipStreamDestroy(streamForGraph)); -} - -static void validate_result(char *hData, size_t size, char val) { - // Validating the result - for (size_t i = 0; i < size; i++) { - if (hData[i] != val) { - WARN("Validation failed at- " << i << " hData[i] " << hData[i]); - REQUIRE(false); - } - } -} - -/* Test verifies hipGraphMemsetNodeSetParams API Functional scenarios. +/** + * Test Description + * ------------------------ + * - Verify that node parameters get updated correctly by creating a node with valid but + * incorrect parameters, and then setting them to the correct values after which the graph is + * executed and the results verified. + * The parameters are also verified via hipGraphMemsetNodeGetParams, which also constitutes a test + * for said API. + * The test is repeated for all valid element sizes(1, 2, 4), and several allocations of different + * height and width both on host and device + * Test source + * ------------------------ + * - unit/graph/hipGraphMemsetNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ -TEST_CASE("Unit_hipGraphMemsetNodeSetParams_Functional") { - constexpr size_t N = 1024; - constexpr size_t Nbytes = N * sizeof(char); - constexpr size_t val = 0; - constexpr size_t updateVal = 1; - constexpr size_t updateVal2 = 2; - char *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; - char *A_h{nullptr}, *B_h{nullptr}; +TEMPLATE_TEST_CASE("Unit_hipGraphMemsetNodeSetParams_Positive_Basic", "", uint8_t, uint16_t, + uint32_t) { + const auto f = [](hipMemsetParams* params) { + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); - HipTest::initArrays(&A_d, &B_d, &C_d, - &A_h, &B_h, nullptr, N, false); + hipGraphNode_t node = nullptr; + LinearAllocGuard initial_alloc(LinearAllocs::hipMalloc, 2 * sizeof(TestType)); - hipGraph_t graph; - hipGraphExec_t graphExec; - hipStream_t streamForGraph; - hipGraphNode_t memsetNode; + hipMemsetParams initial_params = {}; + initial_params.dst = initial_alloc.ptr(); + initial_params.elementSize = sizeof(TestType); + initial_params.width = 2; + initial_params.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&node, graph, nullptr, 0, &initial_params)); - HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipStreamCreate(&streamForGraph)); + hipMemsetParams retrieved_params = {}; + HIP_CHECK(hipGraphMemsetNodeGetParams(node, &retrieved_params)); + REQUIRE(initial_params.dst == retrieved_params.dst); + REQUIRE(initial_params.elementSize == retrieved_params.elementSize); + REQUIRE(initial_params.width == retrieved_params.width); + REQUIRE(initial_params.height == retrieved_params.height); + REQUIRE(initial_params.pitch == retrieved_params.pitch); + REQUIRE(initial_params.value == retrieved_params.value); - hipMemsetParams memsetParams{}; - memset(&memsetParams, 0, sizeof(memsetParams)); - memsetParams.dst = reinterpret_cast(C_d); - memsetParams.value = val; - memsetParams.pitch = 0; - memsetParams.elementSize = sizeof(char); - memsetParams.width = Nbytes; - memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, - &memsetParams)); + HIP_CHECK(hipGraphMemsetNodeSetParams(node, params)); + HIP_CHECK(hipGraphMemsetNodeGetParams(node, &retrieved_params)); + REQUIRE(params->dst == retrieved_params.dst); + REQUIRE(params->elementSize == retrieved_params.elementSize); + REQUIRE(params->width == retrieved_params.width); + REQUIRE(params->height == retrieved_params.height); + REQUIRE(params->pitch == retrieved_params.pitch); + REQUIRE(params->value == retrieved_params.value); - std::vector dependencies; - dependencies.push_back(memsetNode); + hipGraphExec_t graph_exec = nullptr; + HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); - SECTION("Update the memsetNode and check") { - memset(&memsetParams, 0, sizeof(memsetParams)); - memsetParams.dst = reinterpret_cast(A_d); - memsetParams.value = updateVal; - memsetParams.pitch = 0; - memsetParams.elementSize = sizeof(char); - memsetParams.width = Nbytes; - memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, dependencies.data(), - dependencies.size(), &memsetParams)); - HIP_CHECK(hipGraphMemsetNodeSetParams(memsetNode, &memsetParams)); - dependencies.push_back(memsetNode); + HIP_CHECK(hipGraphLaunch(graph_exec, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); - HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); - HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); - HIP_CHECK(hipStreamSynchronize(streamForGraph)); + HIP_CHECK(hipGraphExecDestroy(graph_exec)); + HIP_CHECK(hipGraphDestroy(graph)); - HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); - validate_result(A_h, Nbytes, updateVal); - } - SECTION("Update the memsetNode again and check") { - memset(&memsetParams, 0, sizeof(memsetParams)); - memsetParams.dst = reinterpret_cast(B_d); - memsetParams.value = updateVal2; - memsetParams.pitch = 0; - memsetParams.elementSize = sizeof(char); - memsetParams.width = Nbytes; - memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, dependencies.data(), - dependencies.size(), &memsetParams)); - HIP_CHECK(hipGraphMemsetNodeSetParams(memsetNode, &memsetParams)); - dependencies.push_back(memsetNode); + return hipSuccess; + }; - HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); - HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); - HIP_CHECK(hipStreamSynchronize(streamForGraph)); - - HIP_CHECK(hipMemcpy(B_h, B_d, Nbytes, hipMemcpyDeviceToHost)); - validate_result(B_h, Nbytes, updateVal2); - } - - HipTest::freeArrays(A_d, B_d, C_d, - A_h, B_h, nullptr, false); - HIP_CHECK(hipGraphExecDestroy(graphExec)); - HIP_CHECK(hipGraphDestroy(graph)); - HIP_CHECK(hipStreamDestroy(streamForGraph)); + GraphMemsetNodeCommonPositive(f); +} + +/** + * Test Description + * ------------------------ + * - Verify API behaviour with invalid arguments: + * -# node is nullptr + * -# pNodeParams is nullptr + * -# pNodeParams::dst is nullptr + * -# pNodeParams::elementSize is different from 1, 2, and 4 + * -# pNodeParams::width is zero + * -# pNodeParams::width is larger than the allocated memory region + * -# pNodeParams::height is zero + * -# pNodeParams::pitch is less than width when height is more than 1 + * -# pNodeParams::pitch * pMemsetParams::height is larger than the allocated memory region + * Test source + * ------------------------ + * - unit/graph/hipGraphMemsetNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipGraphMemsetNodeSetParams_Negative_Parameters") { + using namespace std::placeholders; + + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + LinearAllocGuard alloc(LinearAllocs::hipMalloc, 4 * sizeof(int)); + hipMemsetParams params = {}; + params.dst = alloc.ptr(); + params.elementSize = sizeof(*alloc.ptr()); + params.width = 1; + params.height = 1; + params.value = 42; + + hipGraphNode_t node = nullptr; + HIP_CHECK(hipGraphAddMemsetNode(&node, graph, nullptr, 0, ¶ms)) + + SECTION("node == nullptr") { + HIP_CHECK_ERROR(hipGraphMemsetNodeSetParams(nullptr, ¶ms), hipErrorInvalidValue); + } + + MemsetCommonNegative(std::bind(hipGraphMemsetNodeSetParams, node, _1), params); + + HIP_CHECK(hipGraphDestroy(graph)); }