diff --git a/projects/hip-tests/catch/unit/vulkan_interop/CMakeLists.txt b/projects/hip-tests/catch/unit/vulkan_interop/CMakeLists.txt index a0c39ebb0b..35212887f5 100644 --- a/projects/hip-tests/catch/unit/vulkan_interop/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/vulkan_interop/CMakeLists.txt @@ -8,6 +8,10 @@ set(TEST_SRC hipSignalExternalSemaphoresAsync.cc hipImportExternalSemaphore.cc hipDestroyExternalSemaphore.cc + hipGraphAddExternalSemaphoresSignalNode.cc + hipGraphExternalSemaphoresSignalNodeSetParams.cc + hipGraphExternalSemaphoresSignalNodeGetParams.cc + hipGraphExecExternalSemaphoresSignalNodeSetParams.cc ) if(WIN32) diff --git a/projects/hip-tests/catch/unit/vulkan_interop/graph_tests_common.hh b/projects/hip-tests/catch/unit/vulkan_interop/graph_tests_common.hh new file mode 100644 index 0000000000..bb28ec5ea5 --- /dev/null +++ b/projects/hip-tests/catch/unit/vulkan_interop/graph_tests_common.hh @@ -0,0 +1,76 @@ +/* +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 GraphAddNodeCommonNegativeTests(F f, hipGraph_t graph) { + hipGraphNode_t node = nullptr; + SECTION("graph == nullptr") { + HIP_CHECK_ERROR(f(&node, nullptr, nullptr, 0), hipErrorInvalidValue); + } + + SECTION("node == nullptr") { + HIP_CHECK_ERROR(f(nullptr, graph, nullptr, 0), hipErrorInvalidValue); + } + + SECTION("dependencies == nullptr with size != 0") { + HIP_CHECK_ERROR(f(&node, graph, nullptr, 1), hipErrorInvalidValue); + } + +// Disabled on AMD due to defect - EXSWHTEC-202 +#if HT_NVIDIA + SECTION("Node in dependency is from different graph") { + hipGraph_t other_graph = nullptr; + HIP_CHECK(hipGraphCreate(&other_graph, 0)); + hipGraphNode_t other_node = nullptr; + HIP_CHECK(hipGraphAddEmptyNode(&other_node, other_graph, nullptr, 0)); + hipGraphNode_t node = nullptr; + HIP_CHECK(hipGraphAddEmptyNode(&node, graph, nullptr, 0)); + HIP_CHECK_ERROR(f(&node, graph, &other_node, 1), hipErrorInvalidValue); + HIP_CHECK(hipGraphDestroy(other_graph)); + } +#endif + + SECTION("Invalid numNodes") { + hipGraphNode_t dep_node = nullptr; + HIP_CHECK(hipGraphAddEmptyNode(&dep_node, graph, nullptr, 0)); + HIP_CHECK_ERROR(f(&node, graph, &dep_node, 2), hipErrorInvalidValue); + } + +// Disabled on AMD due to defect - EXSWHTEC-201 +#if HT_NVIDIA + SECTION("Duplicate node in dependencies") { + hipGraphNode_t dep_node = nullptr; + // Need to create two nodes to avoid overlap with Invalid numNodes case + // First one is left dangling as the graph will be destroyed after the section anyway + HIP_CHECK(hipGraphAddEmptyNode(&dep_node, graph, nullptr, 0)); + HIP_CHECK(hipGraphAddEmptyNode(&dep_node, graph, nullptr, 0)); + hipGraphNode_t deps[] = {dep_node, dep_node}; + HIP_CHECK_ERROR(f(&node, graph, deps, 2), hipErrorInvalidValue); + } +#endif +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/vulkan_interop/hipGraphAddExternalSemaphoresSignalNode.cc b/projects/hip-tests/catch/unit/vulkan_interop/hipGraphAddExternalSemaphoresSignalNode.cc new file mode 100644 index 0000000000..932f99cb41 --- /dev/null +++ b/projects/hip-tests/catch/unit/vulkan_interop/hipGraphAddExternalSemaphoresSignalNode.cc @@ -0,0 +1,135 @@ +/* +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 "signal_semaphore_common.hh" +#include "graph_tests_common.hh" + +/** + * @addtogroup hipGraphAddExternalSemaphoresSignalNode hipGraphAddExternalSemaphoresSignalNode + * @{ + * @ingroup GraphTest + * `hipGraphAddExternalSemaphoresSignalNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, const + * hipGraphNode_t* pDependencies, size_t numDependencies, const + * hipExternalSemaphoreSignalNodeParams* nodeParams);` - Creates a external semaphor signal 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/hipGraphAddExternalSemaphoresSignalNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphAddExternalSemaphoresSignalNode_Positive_Basic") { + SignalExternalSemaphoreCommon(GraphExtSemaphoreSignalWrapper<>); +} + +// 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/hipGraphAddExternalSemaphoresSignalNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphAddExternalSemaphoresSignalNode_Vulkan_Positive_Timeline_Semaphore") { + SignalExternalTimelineSemaphoreCommon(GraphExtSemaphoreSignalWrapper<>); +} + +/** + * 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/hipGraphAddExternalSemaphoresSignalNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphAddExternalSemaphoresSignalNode_Vulkan_Positive_Multiple_Semaphores") { + SignalExternalMultipleSemaphoresCommon(GraphExtSemaphoreSignalWrapper<>); +} +#endif + + +/** + * Test Description + * ------------------------ + * - Test to verify hipGraphAddExternalSemaphoresSignalNode 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/hipGraphAddExternalSemaphoresSignalNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphAddExternalSemaphoresSignalNode_Vulkan_Negative_Parameters") { + using namespace std::placeholders; + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + VulkanTest vkt(enable_validation); + hipExternalSemaphoreSignalParams signal_params = {}; + signal_params.params.fence.value = 1; + auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + + hipExternalSemaphoreSignalNodeParams node_params = {}; + node_params.extSemArray = &hip_ext_semaphore; + node_params.paramsArray = &signal_params; + node_params.numExtSems = 1; + + GraphAddNodeCommonNegativeTests( + std::bind(hipGraphAddExternalSemaphoresSignalNode, _1, _2, _3, _4, &node_params), graph); + + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + HIP_CHECK(hipGraphDestroy(graph)); +} diff --git a/projects/hip-tests/catch/unit/vulkan_interop/hipGraphExecExternalSemaphoresSignalNodeSetParams.cc b/projects/hip-tests/catch/unit/vulkan_interop/hipGraphExecExternalSemaphoresSignalNodeSetParams.cc new file mode 100644 index 0000000000..939d95beb4 --- /dev/null +++ b/projects/hip-tests/catch/unit/vulkan_interop/hipGraphExecExternalSemaphoresSignalNodeSetParams.cc @@ -0,0 +1,193 @@ +/* +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 "signal_semaphore_common.hh" + +/** + * @addtogroup hipGraphExecExternalSemaphoresSignalNodeSetParams + * hipGraphExecExternalSemaphoresSignalNodeSetParams + * @{ + * @ingroup GraphTest + * `hipGraphExecExternalSemaphoresSignalNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t + * hNode, const hipExternalSemaphoreSignalNodeParams* nodeParams)` - Updates node parameters in the + * external semaphore signal node in the given graphExec. + */ + +static hipError_t GraphExecSemaphoreSetParamsSignalWrapper( + hipExternalSemaphore_t* extSemArray, hipExternalSemaphoreSignalParams* paramsArray, + unsigned int numExtSems, hipStream_t stream) { + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t node = nullptr; + + hipExternalSemaphoreSignalNodeParams node_params = {}; + node_params.extSemArray = extSemArray; + node_params.paramsArray = paramsArray; + node_params.numExtSems = numExtSems; + + hipExternalSemaphoreSignalParams* signal_params = + new hipExternalSemaphoreSignalParams[numExtSems]; + for (unsigned int i = 0; i < numExtSems; i++) { + signal_params[i].params.fence.value = 10 + i; + } + + hipExternalSemaphoreSignalNodeParams initial_params = {}; + initial_params.extSemArray = extSemArray; + initial_params.paramsArray = signal_params; + initial_params.numExtSems = numExtSems; + + HIP_CHECK(hipGraphAddExternalSemaphoresSignalNode(&node, graph, nullptr, 0, &initial_params)); + + hipGraphExec_t graph_exec = nullptr; + HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); + + HIP_CHECK(hipGraphExecExternalSemaphoresSignalNodeSetParams(graph_exec, node, &node_params)); + + hipExternalSemaphoreSignalNodeParams retrieved_params = {0}; + HIP_CHECK(hipGraphExternalSemaphoresSignalNodeGetParams(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)); + delete[] signal_params; + + 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 signals the external binary semaphore and + * operation finishes successfully. + * Test source + * ------------------------ + * - unit/vulkan_interop/hipGraphExecExternalSemaphoresSignalNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphExecExternalSemaphoresSignalNodeSetParams_Positive_Basic") { + SignalExternalSemaphoreCommon(GraphExecSemaphoreSetParamsSignalWrapper); +} + +// 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 signals the external timeline semaphore and + * operation finishes successfully. + * Test source + * ------------------------ + * - unit/vulkan_interop/hipGraphExecExternalSemaphoresSignalNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE( + "Unit_hipGraphExecExternalSemaphoresSignalNodeSetParams_Vulkan_Positive_Timeline_Semaphore") { + SignalExternalTimelineSemaphoreCommon(GraphExecSemaphoreSetParamsSignalWrapper); +} + +/** + * 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 signals the external binary semaphores and + * operation finishes successfully. + * Test source + * ------------------------ + * - unit/vulkan_interop/hipGraphExecExternalSemaphoresSignalNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE( + "Unit_hipGraphExecExternalSemaphoresSignalNodeSetParams_Vulkan_Positive_Multiple_Semaphores") { + SignalExternalMultipleSemaphoresCommon(GraphExecSemaphoreSetParamsSignalWrapper); +} +#endif + +/** + * Test Description + * ------------------------ + * - Test to verify hipGraphExecExternalSemaphoresSignalNodeSetParams behavior with invalid + * arguments: + * -# Nullptr graphexec + * -# Nullptr graph node + * -# Nullptr params + * Test source + * ------------------------ + * - /unit/vulkan_interop/hipGraphExecExternalSemaphoresSignalNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphExecExternalSemaphoresSignalNodeSetParams_Vulkan_Negative_Parameters") { + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + VulkanTest vkt(enable_validation); + hipExternalSemaphoreSignalParams signal_params = {}; + signal_params.params.fence.value = 1; + auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + + hipExternalSemaphoreSignalNodeParams node_params = {}; + node_params.extSemArray = &hip_ext_semaphore; + node_params.paramsArray = &signal_params; + node_params.numExtSems = 1; + + hipGraphNode_t node = nullptr; + HIP_CHECK(hipGraphAddExternalSemaphoresSignalNode(&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(hipGraphExecExternalSemaphoresSignalNodeSetParams(nullptr, node, &node_params), + hipErrorInvalidValue); + } + + SECTION("node == nullptr") { + HIP_CHECK_ERROR( + hipGraphExecExternalSemaphoresSignalNodeSetParams(graph_exec, nullptr, &node_params), + hipErrorInvalidValue); + } + + SECTION("params == nullptr") { + HIP_CHECK_ERROR(hipGraphExecExternalSemaphoresSignalNodeSetParams(graph_exec, node, nullptr), + hipErrorInvalidValue); + } + + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + HIP_CHECK(hipGraphExecDestroy(graph_exec)); + HIP_CHECK(hipGraphDestroy(graph)); +} diff --git a/projects/hip-tests/catch/unit/vulkan_interop/hipGraphExternalSemaphoresSignalNodeGetParams.cc b/projects/hip-tests/catch/unit/vulkan_interop/hipGraphExternalSemaphoresSignalNodeGetParams.cc new file mode 100644 index 0000000000..6f6c3c2787 --- /dev/null +++ b/projects/hip-tests/catch/unit/vulkan_interop/hipGraphExternalSemaphoresSignalNodeGetParams.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 "signal_semaphore_common.hh" + +/** + * @addtogroup hipGraphExternalSemaphoresSignalNodeGetParams + * hipGraphExternalSemaphoresSignalNodeGetParams + * @{ + * @ingroup GraphTest + * `hipGraphExternalSemaphoresSignalNodeGetParams(hipGraphNode_t hNode, + * hipExternalSemaphoreSignalNodeParams* params_out)` - Returns external semaphore signal node + * params. + * ________________________ + * Test cases from other APIs: + * - @ref Unit_hipGraphExternalSemaphoresSignalNodeSetParams_Positive_Basic + * - @ref Unit_hipGraphExternalSemaphoresSignalNodeSetParams_Vulkan_Positive_Timeline_Semaphore + * - @ref Unit_hipGraphExternalSemaphoresSignalNodeSetParams_Vulkan_Positive_Multiple_Semaphores + */ + +/** + * Test Description + * ------------------------ + * - Test to verify hipGraphExternalSemaphoresSignalNodeGetParams behavior with invalid + * arguments: + * -# Nullptr graph node + * -# Nullptr params + * -# Node is destroyed + * Test source + * ------------------------ + * - /unit/vulkan_interop/hipGraphExternalSemaphoresSignalNodeGetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphExternalSemaphoresSignalNodeGetParams_Negative_Parameters") { + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + VulkanTest vkt(enable_validation); + hipExternalSemaphoreSignalParams signal_params = {}; + signal_params.params.fence.value = 1; + auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + + hipExternalSemaphoreSignalNodeParams node_params = {}; + node_params.extSemArray = &hip_ext_semaphore; + node_params.paramsArray = &signal_params; + node_params.numExtSems = 1; + hipExternalSemaphoreSignalNodeParams retrieved_params; + + hipGraphNode_t node = nullptr; + HIP_CHECK(hipGraphAddExternalSemaphoresSignalNode(&node, graph, nullptr, 0, &node_params)); + + SECTION("node == nullptr") { + HIP_CHECK_ERROR(hipGraphExternalSemaphoresSignalNodeGetParams(nullptr, &retrieved_params), + hipErrorInvalidValue); + } + + SECTION("params_out == nullptr") { + HIP_CHECK_ERROR(hipGraphExternalSemaphoresSignalNodeGetParams(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( + hipGraphAddExternalSemaphoresSignalNode(&node_temp, graph_temp, nullptr, 0, &node_params)); + HIP_CHECK(hipGraphDestroy(graph_temp)); + HIP_CHECK_ERROR(hipGraphExternalSemaphoresSignalNodeGetParams(node_temp, &retrieved_params), + hipErrorInvalidValue); + } +#endif +} diff --git a/projects/hip-tests/catch/unit/vulkan_interop/hipGraphExternalSemaphoresSignalNodeSetParams.cc b/projects/hip-tests/catch/unit/vulkan_interop/hipGraphExternalSemaphoresSignalNodeSetParams.cc new file mode 100644 index 0000000000..8f964d966d --- /dev/null +++ b/projects/hip-tests/catch/unit/vulkan_interop/hipGraphExternalSemaphoresSignalNodeSetParams.cc @@ -0,0 +1,137 @@ +/* +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 "signal_semaphore_common.hh" + +/** + * @addtogroup hipGraphExternalSemaphoresSignalNodeSetParams + * hipGraphExternalSemaphoresSignalNodeSetParams + * @{ + * @ingroup GraphTest + * `hipGraphExternalSemaphoresSignalNodeSetParams(hipGraphNode_t hNode, const + * hipExternalSemaphoreSignalNodeParams* nodeParams)` - Updates node parameters in the external + * semaphore signal 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 signals the external binary semaphore and operation finishes + * successfully. + * Test source + * ------------------------ + * - unit/vulkan_interop/hipGraphExternalSemaphoresSignalNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphExternalSemaphoresSignalNodeSetParams_Positive_Basic") { + SignalExternalSemaphoreCommon(GraphExtSemaphoreSignalWrapper); +} + +// 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 signals the external timeline semaphore and operation finishes + * successfully. + * Test source + * ------------------------ + * - unit/vulkan_interop/hipGraphExternalSemaphoresSignalNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphExternalSemaphoresSignalNodeSetParams_Vulkan_Positive_Timeline_Semaphore") { + SignalExternalTimelineSemaphoreCommon(GraphExtSemaphoreSignalWrapper); +} + +/** + * 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 signals the external binary semaphores and operation finishes + * successfully. + * Test source + * ------------------------ + * - unit/vulkan_interop/hipGraphExternalSemaphoresSignalNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE( + "Unit_hipGraphExternalSemaphoresSignalNodeSetParams_Vulkan_Positive_Multiple_Semaphores") { + SignalExternalMultipleSemaphoresCommon(GraphExtSemaphoreSignalWrapper); +} +#endif + +/** + * Test Description + * ------------------------ + * - Test to verify hipGraphExternalSemaphoresSignalNodeSetParams behavior with invalid + * arguments: + * -# Nullptr graph node + * -# Nullptr params + * Test source + * ------------------------ + * - /unit/vulkan_interop/hipGraphExternalSemaphoresSignalNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphExternalSemaphoresSignalNodeSetParams_Vulkan_Negative_Parameters") { + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + VulkanTest vkt(enable_validation); + hipExternalSemaphoreSignalParams signal_params = {}; + signal_params.params.fence.value = 1; + auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + + hipExternalSemaphoreSignalNodeParams node_params = {}; + node_params.extSemArray = &hip_ext_semaphore; + node_params.paramsArray = &signal_params; + node_params.numExtSems = 1; + + SECTION("node == nullptr") { + HIP_CHECK_ERROR(hipGraphExternalSemaphoresSignalNodeSetParams(nullptr, &node_params), + hipErrorInvalidValue); + } + + hipGraphNode_t node = nullptr; + HIP_CHECK(hipGraphAddExternalSemaphoresSignalNode(&node, graph, nullptr, 0, &node_params)); + + SECTION("params == nullptr") { + HIP_CHECK_ERROR(hipGraphExternalSemaphoresSignalNodeSetParams(node, nullptr), + hipErrorInvalidValue); + } + + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + HIP_CHECK(hipGraphDestroy(graph)); +} diff --git a/projects/hip-tests/catch/unit/vulkan_interop/hipSignalExternalSemaphoresAsync.cc b/projects/hip-tests/catch/unit/vulkan_interop/hipSignalExternalSemaphoresAsync.cc index 4485a3ad5f..64ae1e3637 100644 --- a/projects/hip-tests/catch/unit/vulkan_interop/hipSignalExternalSemaphoresAsync.cc +++ b/projects/hip-tests/catch/unit/vulkan_interop/hipSignalExternalSemaphoresAsync.cc @@ -20,152 +20,20 @@ THE SOFTWARE. */ #include "vulkan_test.hh" - -constexpr bool enable_validation = false; +#include "signal_semaphore_common.hh" TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_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)); - hipExternalSemaphoreSignalParams signal_params = {}; - signal_params.params.fence.value = 0; - HIP_CHECK(hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, &signal_params, 1, nullptr)); - HIP_CHECK(hipDeviceSynchronize()); - VkSubmitInfo submit_info = {}; - submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; - submit_info.commandBufferCount = 1; - submit_info.pCommandBuffers = &command_buffer; - VkSemaphore waitSemaphores[] = {semaphore}; - // VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT or VK_PIPELINE_STAGE_TRANSFER_BIT can work - VkPipelineStageFlags waitStages[] = {VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT}; - submit_info.waitSemaphoreCount = 1; - submit_info.pWaitSemaphores = waitSemaphores; - submit_info.pWaitDstStageMask = waitStages; - const auto fence = vkt.CreateFence(); - VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence)); - REQUIRE(vkGetFenceStatus(vkt.GetDevice(), fence) == VK_NOT_READY); - PollStream(nullptr, hipSuccess); - VK_CHECK_RESULT( - vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/)); - HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + SignalExternalSemaphoreCommon(hipSignalExternalSemaphoresAsync); } // Timeline semaphores unsupported on AMD #if HT_NVIDIA TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Positive_Timeline_Semaphore") { - VulkanTest vkt(enable_validation); - constexpr uint64_t signal_value = 2; - - 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)); - - hipExternalSemaphoreSignalParams signal_params = {}; - signal_params.params.fence.value = signal_value; - - HIP_CHECK(hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, &signal_params, 1, nullptr)); - PollStream(nullptr, hipSuccess); - - uint64_t sem_value = 0u; - VK_CHECK_RESULT(vkGetSemaphoreCounterValue(vkt.GetDevice(), semaphore, &sem_value)); - - REQUIRE(2 == sem_value); - - HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + SignalExternalTimelineSemaphoreCommon(hipSignalExternalSemaphoresAsync); } TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Positive_Multiple_Semaphores") { - 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); - -#if HT_AMD - constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_BINARY; -#else - constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_TIMELINE; -#endif - - 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)); - - uint64_t wait_values[] = {1, 0}; - VkTimelineSemaphoreSubmitInfo timeline_submit_info = {}; - timeline_submit_info.sType = VK_STRUCTURE_TYPE_TIMELINE_SEMAPHORE_SUBMIT_INFO; - timeline_submit_info.waitSemaphoreValueCount = 2; - timeline_submit_info.pWaitSemaphoreValues = wait_values; - - VkSemaphore wait_semaphores[] = {timeline_semaphore, binary_semaphore}; - VkSubmitInfo submit_info = {}; - submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; - submit_info.commandBufferCount = 1; - submit_info.pCommandBuffers = &command_buffer; - submit_info.waitSemaphoreCount = 2; - submit_info.pWaitSemaphores = wait_semaphores; - submit_info.pNext = - second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? &timeline_submit_info : nullptr; - const auto fence = vkt.CreateFence(); - VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence)); - - REQUIRE(vkGetFenceStatus(vkt.GetDevice(), fence) == VK_NOT_READY); - - hipExternalSemaphoreSignalParams binary_signal_params = {}; - binary_signal_params.params.fence.value = 0; - hipExternalSemaphoreSignalParams timeline_signal_params = {}; - timeline_signal_params.params.fence.value = - second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? 2 : 0; - hipExternalSemaphore_t ext_semaphores[] = {hip_binary_ext_semaphore, hip_timeline_ext_semaphore}; - hipExternalSemaphoreSignalParams signal_params[] = {binary_signal_params, - timeline_signal_params}; - HIP_CHECK(hipSignalExternalSemaphoresAsync(ext_semaphores, signal_params, 2, nullptr)); - - VK_CHECK_RESULT( - vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/)); - - HIP_CHECK(hipDestroyExternalSemaphore(hip_binary_ext_semaphore)); - HIP_CHECK(hipDestroyExternalSemaphore(hip_timeline_ext_semaphore)); + SignalExternalMultipleSemaphoresCommon(hipSignalExternalSemaphoresAsync); } #endif diff --git a/projects/hip-tests/catch/unit/vulkan_interop/signal_semaphore_common.hh b/projects/hip-tests/catch/unit/vulkan_interop/signal_semaphore_common.hh new file mode 100644 index 0000000000..2c8b2f1c86 --- /dev/null +++ b/projects/hip-tests/catch/unit/vulkan_interop/signal_semaphore_common.hh @@ -0,0 +1,236 @@ +/* +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 SignalExternalSemaphoreCommon(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)); + hipExternalSemaphoreSignalParams signal_params = {}; + signal_params.params.fence.value = 0; + HIP_CHECK(f(&hip_ext_semaphore, &signal_params, 1, nullptr)); + HIP_CHECK(hipDeviceSynchronize()); + VkSubmitInfo submit_info = {}; + submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + submit_info.commandBufferCount = 1; + submit_info.pCommandBuffers = &command_buffer; + VkSemaphore waitSemaphores[] = {semaphore}; + // VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT or VK_PIPELINE_STAGE_TRANSFER_BIT can work + VkPipelineStageFlags waitStages[] = {VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT}; + submit_info.waitSemaphoreCount = 1; + submit_info.pWaitSemaphores = waitSemaphores; + submit_info.pWaitDstStageMask = waitStages; + const auto fence = vkt.CreateFence(); + VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence)); + REQUIRE(vkGetFenceStatus(vkt.GetDevice(), fence) == VK_NOT_READY); + PollStream(nullptr, hipSuccess); + VK_CHECK_RESULT( + vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/)); + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); +} + +#if HT_NVIDIA +template void SignalExternalTimelineSemaphoreCommon(F f) { + VulkanTest vkt(enable_validation); + constexpr uint64_t signal_value = 2; + + 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)); + + hipExternalSemaphoreSignalParams signal_params = {}; + signal_params.params.fence.value = signal_value; + + HIP_CHECK(f(&hip_ext_semaphore, &signal_params, 1, nullptr)); + PollStream(nullptr, hipSuccess); + + uint64_t sem_value = 0u; + VK_CHECK_RESULT(vkGetSemaphoreCounterValue(vkt.GetDevice(), semaphore, &sem_value)); + + REQUIRE(2 == sem_value); + + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); +} + +template void SignalExternalMultipleSemaphoresCommon(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); + +#if HT_AMD + constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_BINARY; +#else + constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_TIMELINE; +#endif + + 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)); + + uint64_t wait_values[] = {1, 0}; + VkTimelineSemaphoreSubmitInfo timeline_submit_info = {}; + timeline_submit_info.sType = VK_STRUCTURE_TYPE_TIMELINE_SEMAPHORE_SUBMIT_INFO; + timeline_submit_info.waitSemaphoreValueCount = 2; + timeline_submit_info.pWaitSemaphoreValues = wait_values; + + VkSemaphore wait_semaphores[] = {timeline_semaphore, binary_semaphore}; + VkSubmitInfo submit_info = {}; + submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + submit_info.commandBufferCount = 1; + submit_info.pCommandBuffers = &command_buffer; + submit_info.waitSemaphoreCount = 2; + submit_info.pWaitSemaphores = wait_semaphores; + submit_info.pNext = + second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? &timeline_submit_info : nullptr; + const auto fence = vkt.CreateFence(); + VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence)); + + REQUIRE(vkGetFenceStatus(vkt.GetDevice(), fence) == VK_NOT_READY); + + hipExternalSemaphoreSignalParams binary_signal_params = {}; + binary_signal_params.params.fence.value = 0; + hipExternalSemaphoreSignalParams timeline_signal_params = {}; + timeline_signal_params.params.fence.value = + second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? 2 : 0; + hipExternalSemaphore_t ext_semaphores[] = {hip_binary_ext_semaphore, hip_timeline_ext_semaphore}; + hipExternalSemaphoreSignalParams signal_params[] = {binary_signal_params, timeline_signal_params}; + HIP_CHECK(f(ext_semaphores, signal_params, 2, nullptr)); + + VK_CHECK_RESULT( + vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/)); + + HIP_CHECK(hipDestroyExternalSemaphore(hip_binary_ext_semaphore)); + HIP_CHECK(hipDestroyExternalSemaphore(hip_timeline_ext_semaphore)); +} +#endif + +static inline bool operator==(const hipExternalSemaphoreSignalNodeParams& lhs, + const hipExternalSemaphoreSignalNodeParams& 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 GraphExtSemaphoreSignalWrapper(hipExternalSemaphore_t* extSemArray, + hipExternalSemaphoreSignalParams* paramsArray, + unsigned int numExtSems, hipStream_t stream) { + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t node = nullptr; + hipExternalSemaphoreSignalNodeParams retrieved_params = {}; + memset(&retrieved_params, 0, sizeof(retrieved_params)); + + hipExternalSemaphoreSignalNodeParams node_params = {}; + node_params.extSemArray = extSemArray; + node_params.paramsArray = paramsArray; + node_params.numExtSems = numExtSems; + + if constexpr (set_params) { + hipExternalSemaphoreSignalParams* signal_params = + new hipExternalSemaphoreSignalParams[numExtSems]; + for (unsigned int i = 0; i < numExtSems; i++) { + signal_params[i].params.fence.value = 10 + i; + } + + hipExternalSemaphoreSignalNodeParams initial_params = {}; + initial_params.extSemArray = extSemArray; + initial_params.paramsArray = signal_params; + initial_params.numExtSems = numExtSems; + + HIP_CHECK(hipGraphAddExternalSemaphoresSignalNode(&node, graph, nullptr, 0, &initial_params)); + + HIP_CHECK(hipGraphExternalSemaphoresSignalNodeGetParams(node, &retrieved_params)); + REQUIRE(initial_params == retrieved_params); + HIP_CHECK(hipGraphExternalSemaphoresSignalNodeSetParams(node, &node_params)); + + delete[] signal_params; + } else { + HIP_CHECK(hipGraphAddExternalSemaphoresSignalNode(&node, graph, nullptr, 0, &node_params)); + } + + HIP_CHECK(hipGraphExternalSemaphoresSignalNodeGetParams(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; +}