diff --git a/catch/unit/vulkan_interop/CMakeLists.txt b/catch/unit/vulkan_interop/CMakeLists.txt index 35212887f5..3728fd2bf9 100644 --- a/catch/unit/vulkan_interop/CMakeLists.txt +++ b/catch/unit/vulkan_interop/CMakeLists.txt @@ -12,6 +12,10 @@ set(TEST_SRC hipGraphExternalSemaphoresSignalNodeSetParams.cc hipGraphExternalSemaphoresSignalNodeGetParams.cc hipGraphExecExternalSemaphoresSignalNodeSetParams.cc + hipGraphAddExternalSemaphoresWaitNode.cc + hipGraphExternalSemaphoresWaitNodeSetParams.cc + hipGraphExternalSemaphoresWaitNodeGetParams.cc + hipGraphExecExternalSemaphoresWaitNodeSetParams.cc ) if(WIN32) diff --git a/catch/unit/vulkan_interop/hipGraphAddExternalSemaphoresWaitNode.cc b/catch/unit/vulkan_interop/hipGraphAddExternalSemaphoresWaitNode.cc new file mode 100644 index 0000000000..a64eec5f6f --- /dev/null +++ b/catch/unit/vulkan_interop/hipGraphAddExternalSemaphoresWaitNode.cc @@ -0,0 +1,134 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +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. +*/ + +#include + +#include "vulkan_test.hh" +#include "wait_semaphore_common.hh" +#include "graph_tests_common.hh" + +/** + * @addtogroup hipGraphAddExternalSemaphoresWaitNode hipGraphAddExternalSemaphoresWaitNode + * @{ + * @ingroup GraphTest + * `hipGraphAddExternalSemaphoresWaitNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, const + * hipGraphNode_t* pDependencies, size_t numDependencies, const hipExternalSemaphoreWaitNodeParams* + * nodeParams)` - Creates a external semaphor wait node and adds it to a graph. + */ + +/** + * Test Description + * ------------------------ + * - Creates two host visible Vulkan buffers. + * - Adds a buffer copy command which will copy from one buffer to another. + * - Creates an external Vulkan binary semaphore. + * - Creates a Vulkan fence and signals semaphore asynchronously. + * - Waits for the operation to finish successfully. + * Test source + * ------------------------ + * - unit/vulkan_interop/hipGraphAddExternalSemaphoresWaitNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ + +TEST_CASE("Unit_hipGraphAddExternalSemaphoresWaitNode_Positive_Basic") { + WaitExternalSemaphoreCommon(GraphExtSemaphoreWaitWrapper<>); +} + +// Timeline semaphores unsupported on AMD +#if HT_NVIDIA + +/** + * Test Description + * ------------------------ + * - Creates an external Vulkan timeline semaphore. + * - Imports the semaphore and signals. + * - Waits for the operation to finish successfully. + * Test source + * ------------------------ + * - unit/vulkan_interop/hipGraphAddExternalSemaphoresWaitNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphAddExternalSemaphoresWaitNode_Vulkan_Positive_Timeline_Semaphore") { + WaitExternalTimelineSemaphoreCommon(GraphExtSemaphoreWaitWrapper<>); +} +#endif + +/** + * Test Description + * ------------------------ + * - Creates two host visible Vulkan buffers. + * - Adds a buffer copy command which will copy from one buffer to another. + * - Creates multiple external Vulkan binary semaphores. + * - Createas a Vulkan fence and signals semaphores. + * - Waits for the operations to finish successfully. + * Test source + * ------------------------ + * - unit/vulkan_interop/hipGraphAddExternalSemaphoresWaitNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphAddExternalSemaphoresWaitNode_Vulkan_Positive_Multiple_Semaphores") { + WaitExternalMultipleSemaphoresCommon(GraphExtSemaphoreWaitWrapper<>); +} + +/** + * Test Description + * ------------------------ + * - Test to verify hipGraphAddExternalSemaphoresWaitNode behavior with invalid arguments: + * -# Nullptr graph + * -# Nullptr graph node + * -# Invalid numDependencies for null list of dependencies + * -# Node in dependency is from different graph + * -# Invalid numNodes + * -# Duplicate node in dependencies + * Test source + * ------------------------ + * - /unit/vulkan_interop/hipGraphAddExternalSemaphoresWaitNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphAddExternalSemaphoresWaitNode_Vulkan_Negative_Parameters") { + using namespace std::placeholders; + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + VulkanTest vkt(enable_validation); + hipExternalSemaphoreWaitParams wait_params = {}; + wait_params.params.fence.value = 1; + auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + + hipExternalSemaphoreWaitNodeParams node_params = {}; + node_params.extSemArray = &hip_ext_semaphore; + node_params.paramsArray = &wait_params; + node_params.numExtSems = 1; + + GraphAddNodeCommonNegativeTests( + std::bind(hipGraphAddExternalSemaphoresWaitNode, _1, _2, _3, _4, &node_params), graph); + + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + HIP_CHECK(hipGraphDestroy(graph)); +} diff --git a/catch/unit/vulkan_interop/hipGraphExecExternalSemaphoresWaitNodeSetParams.cc b/catch/unit/vulkan_interop/hipGraphExecExternalSemaphoresWaitNodeSetParams.cc new file mode 100644 index 0000000000..6d03be3d40 --- /dev/null +++ b/catch/unit/vulkan_interop/hipGraphExecExternalSemaphoresWaitNodeSetParams.cc @@ -0,0 +1,189 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +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. +*/ + +#include "vulkan_test.hh" +#include "wait_semaphore_common.hh" + +/** + * @addtogroup hipGraphExecExternalSemaphoresWaitNodeSetParams + * hipGraphExecExternalSemaphoresWaitNodeSetParams + * @{ + * @ingroup GraphTest + * `hipGraphExecExternalSemaphoresWaitNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t hNode, + * const hipExternalSemaphoreWaitNodeParams* nodeParams)` - Updates node parameters in the external + * semaphore wait node in the given graphExec. + */ + +static hipError_t GraphExecSemaphoreSetParamsWaitWrapper( + hipExternalSemaphore_t* extSemArray, hipExternalSemaphoreWaitParams* paramsArray, + unsigned int numExtSems, hipStream_t stream) { + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t node = nullptr; + + hipExternalSemaphoreWaitNodeParams node_params = {}; + node_params.extSemArray = extSemArray; + node_params.paramsArray = paramsArray; + node_params.numExtSems = numExtSems; + + hipExternalSemaphoreWaitParams wait_params[numExtSems]; + for (unsigned int i = 0; i < numExtSems; i++) { + wait_params[i].flags = 0; + wait_params[i].params.fence.value = 10 + i; + } + + hipExternalSemaphoreWaitNodeParams initial_params = {}; + initial_params.extSemArray = extSemArray; + initial_params.paramsArray = wait_params; + initial_params.numExtSems = numExtSems; + + HIP_CHECK(hipGraphAddExternalSemaphoresWaitNode(&node, graph, nullptr, 0, &initial_params)); + + hipGraphExec_t graph_exec = nullptr; + HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); + + HIP_CHECK(hipGraphExecExternalSemaphoresWaitNodeSetParams(graph_exec, node, &node_params)); + + hipExternalSemaphoreWaitNodeParams retrieved_params = {0}; + HIP_CHECK(hipGraphExternalSemaphoresWaitNodeGetParams(node, &retrieved_params)); + REQUIRE(initial_params == retrieved_params); + + HIP_CHECK(hipGraphLaunch(graph_exec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + HIP_CHECK(hipGraphExecDestroy(graph_exec)); + HIP_CHECK(hipGraphDestroy(graph)); + + return hipSuccess; +} + +/** + * 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 + * graph is run and it is verified that the graph node waits for the external binary semaphore and + * operation finishes successfully. + * Test source + * ------------------------ + * - unit/vulkan_interop/hipGraphExecExternalSemaphoresWaitNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphExecExternalSemaphoresWaitNodeSetParams_Positive_Basic") { + WaitExternalSemaphoreCommon(GraphExecSemaphoreSetParamsWaitWrapper); +} + +// Timeline semaphores unsupported on AMD +#if HT_NVIDIA + +/** + * 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 + * graph is run and it is verified that the graph node waits for the external timeline semaphore and + * operation finishes successfully. + * Test source + * ------------------------ + * - unit/vulkan_interop/hipGraphExecExternalSemaphoresWaitNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE( + "Unit_hipGraphExecExternalSemaphoresWaitNodeSetParams_Vulkan_Positive_Timeline_Semaphore") { + WaitExternalTimelineSemaphoreCommon(GraphExecSemaphoreSetParamsWaitWrapper); +} +#endif + +/** + * 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 + * graph is run and it is verified that the graph node waits for the external binary semaphores and + * operation finishes successfully. + * Test source + * ------------------------ + * - unit/vulkan_interop/hipGraphExecExternalSemaphoresWaitNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE( + "Unit_hipGraphExecExternalSemaphoresWaitNodeSetParams_Vulkan_Positive_Multiple_Semaphores") { + WaitExternalMultipleSemaphoresCommon(GraphExecSemaphoreSetParamsWaitWrapper); +} + +/** + * Test Description + * ------------------------ + * - Test to verify hipGraphExecExternalSemaphoresWaitNodeSetParams behavior with invalid + * arguments: + * -# Nullptr graphexec + * -# Nullptr graph node + * -# Nullptr params + * Test source + * ------------------------ + * - /unit/vulkan_interop/hipGraphExecExternalSemaphoresWaitNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphExecExternalSemaphoresWaitNodeSetParams_Vulkan_Negative_Parameters") { + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + VulkanTest vkt(enable_validation); + hipExternalSemaphoreWaitParams wait_params = {}; + wait_params.params.fence.value = 1; + auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + + hipExternalSemaphoreWaitNodeParams node_params = {}; + node_params.extSemArray = &hip_ext_semaphore; + node_params.paramsArray = &wait_params; + node_params.numExtSems = 1; + + hipGraphNode_t node = nullptr; + HIP_CHECK(hipGraphAddExternalSemaphoresWaitNode(&node, graph, nullptr, 0, &node_params)); + + hipGraphExec_t graph_exec = nullptr; + HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); + + SECTION("pGraphExec == nullptr") { + HIP_CHECK_ERROR(hipGraphExecExternalSemaphoresWaitNodeSetParams(nullptr, node, &node_params)); + } + + SECTION("node == nullptr") { + HIP_CHECK_ERROR( + hipGraphExecExternalSemaphoresWaitNodeSetParams(graph_exec, nullptr, &node_params)); + } + + SECTION("params == nullptr") { + HIP_CHECK_ERROR(hipGraphExecExternalSemaphoresWaitNodeSetParams(graph_exec, node, nullptr)); + } + + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + HIP_CHECK(hipGraphExecDestroy(graph_exec)); + HIP_CHECK(hipGraphDestroy(graph)); +} diff --git a/catch/unit/vulkan_interop/hipGraphExternalSemaphoresWaitNodeGetParams.cc b/catch/unit/vulkan_interop/hipGraphExternalSemaphoresWaitNodeGetParams.cc new file mode 100644 index 0000000000..b6c0034c04 --- /dev/null +++ b/catch/unit/vulkan_interop/hipGraphExternalSemaphoresWaitNodeGetParams.cc @@ -0,0 +1,96 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +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. +*/ + +#include "vulkan_test.hh" +#include "wait_semaphore_common.hh" + +/** + * @addtogroup hipGraphExternalSemaphoresWaitNodeGetParams + * hipGraphExternalSemaphoresWaitNodeGetParams + * @{ + * @ingroup GraphTest + * `hipGraphExternalSemaphoresWaitNodeGetParams(hipGraphNode_t hNode, + * hipExternalSemaphoreWaitNodeParams* params_out)` - Returns external semaphore wait node params. + * ________________________ + * Test cases from other APIs: + * - @ref Unit_hipGraphExternalSemaphoresWaitNodeSetParams_Positive_Basic + * - @ref Unit_hipGraphExternalSemaphoresWaitNodeSetParams_Vulkan_Positive_Timeline_Semaphore + * - @ref Unit_hipGraphExternalSemaphoresWaitNodeSetParams_Vulkan_Positive_Multiple_Semaphores + */ + + +/** + * Test Description + * ------------------------ + * - Test to verify hipGraphExternalSemaphoresWaitNodeGetParams behavior with invalid + * arguments: + * -# Nullptr graph node + * -# Nullptr params + * -# Node is destroyed + * Test source + * ------------------------ + * - /unit/vulkan_interop/hipGraphExternalSemaphoresWaitNodeGetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphExternalSemaphoresWaitNodeGetParams_Negative_Parameters") { + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + VulkanTest vkt(enable_validation); + hipExternalSemaphoreWaitParams wait_params = {}; + wait_params.params.fence.value = 1; + auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + + hipExternalSemaphoreWaitNodeParams node_params = {}; + node_params.extSemArray = &hip_ext_semaphore; + node_params.paramsArray = &wait_params; + node_params.numExtSems = 1; + hipExternalSemaphoreWaitNodeParams retrieved_params; + + hipGraphNode_t node = nullptr; + HIP_CHECK(hipGraphAddExternalSemaphoresWaitNode(&node, graph, nullptr, 0, &node_params)); + + SECTION("node == nullptr") { + HIP_CHECK_ERROR(hipGraphExternalSemaphoresWaitNodeGetParams(nullptr, &retrieved_params), + hipErrorInvalidValue); + } + + SECTION("params_out == nullptr") { + HIP_CHECK_ERROR(hipGraphExternalSemaphoresWaitNodeGetParams(node, nullptr), + hipErrorInvalidValue); + } + +// Disabled on AMD due to defect - EXSWHTEC-208 +#if HT_NVIDIA + SECTION("Node is destroyed") { + hipGraph_t graph_temp = nullptr; + HIP_CHECK(hipGraphCreate(&graph_temp, 0)); + hipGraphNode_t node_temp = nullptr; + HIP_CHECK( + hipGraphAddExternalSemaphoresWaitNode(&node_temp, graph_temp, nullptr, 0, &node_params)); + HIP_CHECK(hipGraphDestroy(graph_temp)); + HIP_CHECK_ERROR(hipGraphExternalSemaphoresWaitNodeGetParams(node_temp, &retrieved_params), + hipErrorInvalidValue); + } +#endif +} diff --git a/catch/unit/vulkan_interop/hipGraphExternalSemaphoresWaitNodeSetParams.cc b/catch/unit/vulkan_interop/hipGraphExternalSemaphoresWaitNodeSetParams.cc new file mode 100644 index 0000000000..c21c810ac1 --- /dev/null +++ b/catch/unit/vulkan_interop/hipGraphExternalSemaphoresWaitNodeSetParams.cc @@ -0,0 +1,136 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +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. +*/ + +#include "vulkan_test.hh" +#include "wait_semaphore_common.hh" + +/** + * @addtogroup hipGraphExternalSemaphoresWaitNodeSetParams + * hipGraphExternalSemaphoresWaitNodeSetParams + * @{ + * @ingroup GraphTest + * `hipGraphExternalSemaphoresWaitNodeSetParams(hipGraphNode_t hNode, const + * hipExternalSemaphoreWaitNodeParams* nodeParams)` - Updates node parameters in the external + * semaphore wait node. + */ + +/** + * Test Description + * ------------------------ + * - Verify that node parameters get updated correctly by creating a node with valid but + * incorrect parameters, and the setting them to the correct values. The graph is run and it is + * verified that the graph node waits for the external binary semaphore and operation finishes + * successfully. + * Test source + * ------------------------ + * - unit/vulkan_interop/hipGraphExternalSemaphoresWaitNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphExternalSemaphoresWaitNodeSetParams_Positive_Basic") { + WaitExternalSemaphoreCommon(GraphExtSemaphoreWaitWrapper); +} + +// Timeline semaphores unsupported on AMD +#if HT_NVIDIA + +/** + * Test Description + * ------------------------ + * - Verify that node parameters get updated correctly by creating a node with valid but + * incorrect parameters, and the setting them to the correct values. The graph is run and it is + * verified that the graph node waits for the external timeline semaphore and operation finishes + * successfully. + * Test source + * ------------------------ + * - unit/vulkan_interop/hipGraphExternalSemaphoresWaitNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphExternalSemaphoresWaitNodeSetParams_Vulkan_Positive_Timeline_Semaphore") { + WaitExternalTimelineSemaphoreCommon(GraphExtSemaphoreWaitWrapper); +} +#endif + +/** + * Test Description + * ------------------------ + * - Verify that node parameters get updated correctly by creating a node with valid but + * incorrect parameters, and the setting them to the correct values. The graph is run and it is + * verified that the graph node waits for the external binary semaphores and operation finishes + * successfully. + * Test source + * ------------------------ + * - unit/vulkan_interop/hipGraphExternalSemaphoresWaitNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphExternalSemaphoresWaitNodeSetParams_Vulkan_Positive_Multiple_Semaphores") { + WaitExternalMultipleSemaphoresCommon(GraphExtSemaphoreWaitWrapper); +} + +/** + * Test Description + * ------------------------ + * - Test to verify hipGraphExternalSemaphoresWaitNodeSetParams behavior with invalid + * arguments: + * -# Nullptr graph node + * -# Nullptr params + * Test source + * ------------------------ + * - /unit/vulkan_interop/hipGraphExternalSemaphoresWaitNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphExternalSemaphoresWaitNodeSetParams_Vulkan_Negative_Parameters") { + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + VulkanTest vkt(enable_validation); + hipExternalSemaphoreWaitParams wait_params = {}; + wait_params.params.fence.value = 1; + auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + + hipExternalSemaphoreWaitNodeParams node_params = {}; + node_params.extSemArray = &hip_ext_semaphore; + node_params.paramsArray = &wait_params; + node_params.numExtSems = 1; + + SECTION("node == nullptr") { + HIP_CHECK_ERROR(hipGraphExternalSemaphoresWaitNodeSetParams(nullptr, &node_params), + hipErrorInvalidValue); + } + + hipGraphNode_t node = nullptr; + HIP_CHECK(hipGraphAddExternalSemaphoresWaitNode(&node, graph, nullptr, 0, &node_params)); + + SECTION("params == nullptr") { + HIP_CHECK_ERROR(hipGraphExternalSemaphoresWaitNodeSetParams(node, nullptr), + hipErrorInvalidValue); + } + + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + HIP_CHECK(hipGraphDestroy(graph)); +} diff --git a/catch/unit/vulkan_interop/hipWaitExternalSemaphoresAsync.cc b/catch/unit/vulkan_interop/hipWaitExternalSemaphoresAsync.cc index edebebe52a..ee8a175b6f 100644 --- a/catch/unit/vulkan_interop/hipWaitExternalSemaphoresAsync.cc +++ b/catch/unit/vulkan_interop/hipWaitExternalSemaphoresAsync.cc @@ -20,178 +20,21 @@ THE SOFTWARE. */ #include "vulkan_test.hh" - -constexpr bool enable_validation = false; +#include "wait_semaphore_common.hh" TEST_CASE("Unit_hipWaitExternalSemaphoresAsync_Vulkan_Positive_Binary_Semaphore") { - VulkanTest vkt(enable_validation); - - constexpr uint32_t count = 1; - const auto src_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT); - const auto dst_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT); - - const auto command_buffer = vkt.GetCommandBuffer(); - - VkCommandBufferBeginInfo begin_info = {}; - begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; - begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; - VK_CHECK_RESULT(vkBeginCommandBuffer(command_buffer, &begin_info)); - VkBufferCopy buffer_copy = {}; - buffer_copy.size = count * sizeof(*src_storage.host_ptr); - vkCmdCopyBuffer(command_buffer, src_storage.buffer, dst_storage.buffer, 1, &buffer_copy); - VK_CHECK_RESULT(vkEndCommandBuffer(command_buffer)); - - const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY); - const auto hip_sem_handle_desc = - vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_BINARY); - - hipExternalSemaphore_t hip_ext_semaphore; - HIP_CHECK(hipImportExternalSemaphore(&hip_ext_semaphore, &hip_sem_handle_desc)); - - hipExternalSemaphoreWaitParams hip_ext_semaphore_wait_params = {}; - hip_ext_semaphore_wait_params.flags = 0; - hip_ext_semaphore_wait_params.params.fence.value = 0; - HIP_CHECK(hipWaitExternalSemaphoresAsync(&hip_ext_semaphore, &hip_ext_semaphore_wait_params, 1, - nullptr)); - PollStream(nullptr, hipErrorNotReady); - - VkSubmitInfo submit_info = {}; - submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; - submit_info.commandBufferCount = 1; - submit_info.pCommandBuffers = &command_buffer; - submit_info.signalSemaphoreCount = 1; - submit_info.pSignalSemaphores = &semaphore; - - *src_storage.host_ptr = 42; - - const auto fence = vkt.CreateFence(); - VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence)); - VK_CHECK_RESULT( - vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/)); - - PollStream(nullptr, hipSuccess); - - REQUIRE(42 == *dst_storage.host_ptr); - - HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + WaitExternalSemaphoreCommon(hipWaitExternalSemaphoresAsync); } // Timeline semaphores unsupported on AMD #if HT_NVIDIA TEST_CASE("Unit_hipWaitExternalSemaphoresAsync_Vulkan_Positive_Timeline_Semaphore") { - VulkanTest vkt(enable_validation); - - const auto [wait_value, signal_value] = - GENERATE(std::make_pair(2, 2), std::make_pair(2, 3), std::make_pair(3, 2)); - INFO("Wait value: " << wait_value << ", signal value: " << signal_value); - - const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_TIMELINE); - const auto hip_sem_handle_desc = - vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_TIMELINE); - hipExternalSemaphore_t hip_ext_semaphore; - HIP_CHECK(hipImportExternalSemaphore(&hip_ext_semaphore, &hip_sem_handle_desc)); - - hipExternalSemaphoreWaitParams hip_ext_semaphore_wait_params = {}; - hip_ext_semaphore_wait_params.flags = 0; - hip_ext_semaphore_wait_params.params.fence.value = wait_value; - HIP_CHECK(hipWaitExternalSemaphoresAsync(&hip_ext_semaphore, &hip_ext_semaphore_wait_params, 1, - nullptr)); - PollStream(nullptr, hipErrorNotReady); - - VkSemaphoreSignalInfo signal_info = {}; - signal_info.sType = VK_STRUCTURE_TYPE_SEMAPHORE_SIGNAL_INFO; - signal_info.semaphore = semaphore; - signal_info.value = signal_value; - VK_CHECK_RESULT(vkSignalSemaphore(vkt.GetDevice(), &signal_info)); - if (wait_value > signal_value) { - PollStream(nullptr, hipErrorNotReady); - signal_info.value = wait_value; - VK_CHECK_RESULT(vkSignalSemaphore(vkt.GetDevice(), &signal_info)); - } - PollStream(nullptr, hipSuccess); - - HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + WaitExternalTimelineSemaphoreCommon(hipWaitExternalSemaphoresAsync); } #endif TEST_CASE("Unit_hipWaitExternalSemaphoresAsync_Vulkan_Positive_Multiple_Semaphores") { - VulkanTest vkt(enable_validation); - -#if HT_AMD - constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_BINARY; -#else - constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_TIMELINE; -#endif - - constexpr uint32_t count = 1; - const auto src_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT); - const auto dst_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT); - - const auto command_buffer = vkt.GetCommandBuffer(); - - VkCommandBufferBeginInfo begin_info = {}; - begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; - begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; - VK_CHECK_RESULT(vkBeginCommandBuffer(command_buffer, &begin_info)); - VkBufferCopy buffer_copy = {}; - buffer_copy.size = count * sizeof(*src_storage.host_ptr); - vkCmdCopyBuffer(command_buffer, src_storage.buffer, dst_storage.buffer, 1, &buffer_copy); - VK_CHECK_RESULT(vkEndCommandBuffer(command_buffer)); - - const auto binary_semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY); - const auto hip_binary_sem_handle_desc = - vkt.BuildSemaphoreDescriptor(binary_semaphore, VK_SEMAPHORE_TYPE_BINARY); - hipExternalSemaphore_t hip_binary_ext_semaphore; - HIP_CHECK(hipImportExternalSemaphore(&hip_binary_ext_semaphore, &hip_binary_sem_handle_desc)); - - const auto timeline_semaphore = vkt.CreateExternalSemaphore(second_semaphore_type); - const auto hip_timeline_sem_handle_desc = - vkt.BuildSemaphoreDescriptor(timeline_semaphore, second_semaphore_type); - hipExternalSemaphore_t hip_timeline_ext_semaphore; - HIP_CHECK(hipImportExternalSemaphore(&hip_timeline_ext_semaphore, &hip_timeline_sem_handle_desc)); - - hipExternalSemaphoreWaitParams binary_semaphore_wait_params = {}; - binary_semaphore_wait_params.params.fence.value = 0; - - hipExternalSemaphoreWaitParams timeline_semaphore_wait_params = {}; - timeline_semaphore_wait_params.params.fence.value = - second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? 1 : 0; - - hipExternalSemaphore_t ext_semaphores[] = {hip_binary_ext_semaphore, hip_timeline_ext_semaphore}; - hipExternalSemaphoreWaitParams wait_params[] = {binary_semaphore_wait_params, - timeline_semaphore_wait_params}; - HIP_CHECK(hipWaitExternalSemaphoresAsync(ext_semaphores, wait_params, 2, nullptr)); - - PollStream(nullptr, hipErrorNotReady); - - if (second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE) { - VkSemaphoreSignalInfo signal_info = {}; - signal_info.sType = VK_STRUCTURE_TYPE_SEMAPHORE_SIGNAL_INFO; - signal_info.semaphore = timeline_semaphore; - signal_info.value = 1; - VK_CHECK_RESULT(vkSignalSemaphore(vkt.GetDevice(), &signal_info)); - - PollStream(nullptr, hipErrorNotReady); - } - - VkSubmitInfo submit_info = {}; - VkSemaphore signal_semaphores[] = {binary_semaphore, timeline_semaphore}; - submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; - submit_info.commandBufferCount = 1; - submit_info.pCommandBuffers = &command_buffer; - submit_info.signalSemaphoreCount = second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? 1 : 2; - submit_info.pSignalSemaphores = - second_semaphore_type == VK_SEMAPHORE_TYPE_MAX_ENUM ? &binary_semaphore : signal_semaphores; - - const auto fence = vkt.CreateFence(); - VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence)); - VK_CHECK_RESULT( - vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/)); - - PollStream(nullptr, hipSuccess); - - HIP_CHECK(hipDestroyExternalSemaphore(hip_timeline_ext_semaphore)); - HIP_CHECK(hipDestroyExternalSemaphore(hip_binary_ext_semaphore)); + WaitExternalMultipleSemaphoresCommon(hipWaitExternalSemaphoresAsync); } TEST_CASE("Unit_hipWaitExternalSemaphoresAsync_Vulkan_Negative_Parameters") { diff --git a/catch/unit/vulkan_interop/wait_semaphore_common.hh b/catch/unit/vulkan_interop/wait_semaphore_common.hh new file mode 100644 index 0000000000..e590a6d54a --- /dev/null +++ b/catch/unit/vulkan_interop/wait_semaphore_common.hh @@ -0,0 +1,263 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +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 + +constexpr bool enable_validation = false; + +template void WaitExternalSemaphoreCommon(F f) { + VulkanTest vkt(enable_validation); + + constexpr uint32_t count = 1; + const auto src_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT); + const auto dst_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT); + + const auto command_buffer = vkt.GetCommandBuffer(); + + VkCommandBufferBeginInfo begin_info = {}; + begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; + begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; + VK_CHECK_RESULT(vkBeginCommandBuffer(command_buffer, &begin_info)); + VkBufferCopy buffer_copy = {}; + buffer_copy.size = count * sizeof(*src_storage.host_ptr); + vkCmdCopyBuffer(command_buffer, src_storage.buffer, dst_storage.buffer, 1, &buffer_copy); + VK_CHECK_RESULT(vkEndCommandBuffer(command_buffer)); + + const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY); + const auto hip_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_BINARY); + + hipExternalSemaphore_t hip_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_ext_semaphore, &hip_sem_handle_desc)); + + hipExternalSemaphoreWaitParams hip_ext_semaphore_wait_params = {}; + hip_ext_semaphore_wait_params.flags = 0; + hip_ext_semaphore_wait_params.params.fence.value = 0; + HIP_CHECK(f(&hip_ext_semaphore, &hip_ext_semaphore_wait_params, 1, nullptr)); + PollStream(nullptr, hipErrorNotReady); + + VkSubmitInfo submit_info = {}; + submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + submit_info.commandBufferCount = 1; + submit_info.pCommandBuffers = &command_buffer; + submit_info.signalSemaphoreCount = 1; + submit_info.pSignalSemaphores = &semaphore; + + *src_storage.host_ptr = 42; + + const auto fence = vkt.CreateFence(); + VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence)); + VK_CHECK_RESULT( + vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/)); + + PollStream(nullptr, hipSuccess); + + REQUIRE(42 == *dst_storage.host_ptr); + + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); +} + +#if HT_NVIDIA +template void WaitExternalTimelineSemaphoreCommon(F f) { + VulkanTest vkt(enable_validation); + + const auto [wait_value, signal_value] = + GENERATE(std::make_pair(2, 2), std::make_pair(2, 3), std::make_pair(3, 2)); + INFO("Wait value: " << wait_value << ", signal value: " << signal_value); + + const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_TIMELINE); + const auto hip_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_TIMELINE); + hipExternalSemaphore_t hip_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_ext_semaphore, &hip_sem_handle_desc)); + + hipExternalSemaphoreWaitParams hip_ext_semaphore_wait_params = {}; + hip_ext_semaphore_wait_params.flags = 0; + hip_ext_semaphore_wait_params.params.fence.value = wait_value; + HIP_CHECK(f(&hip_ext_semaphore, &hip_ext_semaphore_wait_params, 1, nullptr)); + PollStream(nullptr, hipErrorNotReady); + + VkSemaphoreSignalInfo signal_info = {}; + signal_info.sType = VK_STRUCTURE_TYPE_SEMAPHORE_SIGNAL_INFO; + signal_info.semaphore = semaphore; + signal_info.value = signal_value; + VK_CHECK_RESULT(vkSignalSemaphore(vkt.GetDevice(), &signal_info)); + if (wait_value > signal_value) { + PollStream(nullptr, hipErrorNotReady); + signal_info.value = wait_value; + VK_CHECK_RESULT(vkSignalSemaphore(vkt.GetDevice(), &signal_info)); + } + PollStream(nullptr, hipSuccess); + + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); +} +#endif + +template void WaitExternalMultipleSemaphoresCommon(F f) { + VulkanTest vkt(enable_validation); + +#if HT_AMD + constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_BINARY; +#else + constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_TIMELINE; +#endif + + constexpr uint32_t count = 1; + const auto src_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT); + const auto dst_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT); + + const auto command_buffer = vkt.GetCommandBuffer(); + + VkCommandBufferBeginInfo begin_info = {}; + begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; + begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; + VK_CHECK_RESULT(vkBeginCommandBuffer(command_buffer, &begin_info)); + VkBufferCopy buffer_copy = {}; + buffer_copy.size = count * sizeof(*src_storage.host_ptr); + vkCmdCopyBuffer(command_buffer, src_storage.buffer, dst_storage.buffer, 1, &buffer_copy); + VK_CHECK_RESULT(vkEndCommandBuffer(command_buffer)); + + const auto binary_semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY); + const auto hip_binary_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(binary_semaphore, VK_SEMAPHORE_TYPE_BINARY); + hipExternalSemaphore_t hip_binary_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_binary_ext_semaphore, &hip_binary_sem_handle_desc)); + + const auto timeline_semaphore = vkt.CreateExternalSemaphore(second_semaphore_type); + const auto hip_timeline_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(timeline_semaphore, second_semaphore_type); + hipExternalSemaphore_t hip_timeline_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_timeline_ext_semaphore, &hip_timeline_sem_handle_desc)); + + hipExternalSemaphoreWaitParams binary_semaphore_wait_params = {}; + binary_semaphore_wait_params.params.fence.value = 0; + + hipExternalSemaphoreWaitParams timeline_semaphore_wait_params = {}; + timeline_semaphore_wait_params.params.fence.value = + second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? 1 : 0; + + hipExternalSemaphore_t ext_semaphores[] = {hip_binary_ext_semaphore, hip_timeline_ext_semaphore}; + hipExternalSemaphoreWaitParams wait_params[] = {binary_semaphore_wait_params, + timeline_semaphore_wait_params}; + HIP_CHECK(f(ext_semaphores, wait_params, 2, nullptr)); + + PollStream(nullptr, hipErrorNotReady); + + if (second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE) { + VkSemaphoreSignalInfo signal_info = {}; + signal_info.sType = VK_STRUCTURE_TYPE_SEMAPHORE_SIGNAL_INFO; + signal_info.semaphore = timeline_semaphore; + signal_info.value = 1; + VK_CHECK_RESULT(vkSignalSemaphore(vkt.GetDevice(), &signal_info)); + + PollStream(nullptr, hipErrorNotReady); + } + + VkSubmitInfo submit_info = {}; + VkSemaphore signal_semaphores[] = {binary_semaphore, timeline_semaphore}; + submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + submit_info.commandBufferCount = 1; + submit_info.pCommandBuffers = &command_buffer; + submit_info.signalSemaphoreCount = second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? 1 : 2; + submit_info.pSignalSemaphores = + second_semaphore_type == VK_SEMAPHORE_TYPE_MAX_ENUM ? &binary_semaphore : signal_semaphores; + + const auto fence = vkt.CreateFence(); + VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence)); + VK_CHECK_RESULT( + vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/)); + + PollStream(nullptr, hipSuccess); + + HIP_CHECK(hipDestroyExternalSemaphore(hip_timeline_ext_semaphore)); + HIP_CHECK(hipDestroyExternalSemaphore(hip_binary_ext_semaphore)); +} + +static inline bool operator==(const hipExternalSemaphoreWaitNodeParams& lhs, + const hipExternalSemaphoreWaitNodeParams& rhs) { + bool equal = true; + if (lhs.numExtSems != rhs.numExtSems) { + return false; + } + for (unsigned int i = 0; i < lhs.numExtSems; i++) { + if ((lhs.extSemArray[i] != rhs.extSemArray[i]) || + (lhs.paramsArray[i].params.fence.value != rhs.paramsArray[i].params.fence.value)) { + equal = false; + break; + } + } + return equal; +} + +template +hipError_t GraphExtSemaphoreWaitWrapper(hipExternalSemaphore_t* extSemArray, + hipExternalSemaphoreWaitParams* paramsArray, + unsigned int numExtSems, hipStream_t stream) { + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t node = nullptr; + hipExternalSemaphoreWaitNodeParams retrieved_params = {}; + memset(&retrieved_params, 0, sizeof(retrieved_params)); + + hipExternalSemaphoreWaitNodeParams node_params = {}; + node_params.extSemArray = extSemArray; + node_params.paramsArray = paramsArray; + node_params.numExtSems = numExtSems; + + if constexpr (set_params) { + hipExternalSemaphoreWaitParams* wait_params = new hipExternalSemaphoreWaitParams[numExtSems]; + for (unsigned int i = 0; i < numExtSems; i++) { + wait_params[i].flags = 0; + wait_params[i].params.fence.value = 10 + i; + } + + hipExternalSemaphoreWaitNodeParams initial_params = {}; + initial_params.extSemArray = extSemArray; + initial_params.paramsArray = wait_params; + initial_params.numExtSems = numExtSems; + + HIP_CHECK(hipGraphAddExternalSemaphoresWaitNode(&node, graph, nullptr, 0, &initial_params)); + + HIP_CHECK(hipGraphExternalSemaphoresWaitNodeGetParams(node, &retrieved_params)); + REQUIRE(initial_params == retrieved_params); + HIP_CHECK(hipGraphExternalSemaphoresWaitNodeSetParams(node, &node_params)); + + delete[] wait_params; + } else { + HIP_CHECK(hipGraphAddExternalSemaphoresWaitNode(&node, graph, nullptr, 0, &node_params)); + } + + HIP_CHECK(hipGraphExternalSemaphoresWaitNodeGetParams(node, &retrieved_params)); + REQUIRE(node_params == retrieved_params); + + hipGraphExec_t graph_exec = nullptr; + HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); + + HIP_CHECK(hipGraphLaunch(graph_exec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + HIP_CHECK(hipGraphExecDestroy(graph_exec)); + HIP_CHECK(hipGraphDestroy(graph)); + + return hipSuccess; +}