EXSWHTEC-372 - Implement tests for the hipGraph*ExternalSemaphoresSignalNode APIs #450
Change-Id: I6122be35eb4cbea8ecde5642ac436b2f3b4c3a24
[ROCm/hip-tests commit: dec1be580e]
This commit is contained in:
Nives Vukovic
committad av
Rakesh Roy
Rakesh Roy
förälder
d439c6b986
incheckning
46cbb02cb2
@@ -8,6 +8,10 @@ set(TEST_SRC
|
||||
hipSignalExternalSemaphoresAsync.cc
|
||||
hipImportExternalSemaphore.cc
|
||||
hipDestroyExternalSemaphore.cc
|
||||
hipGraphAddExternalSemaphoresSignalNode.cc
|
||||
hipGraphExternalSemaphoresSignalNodeSetParams.cc
|
||||
hipGraphExternalSemaphoresSignalNodeGetParams.cc
|
||||
hipGraphExecExternalSemaphoresSignalNodeSetParams.cc
|
||||
)
|
||||
|
||||
if(WIN32)
|
||||
|
||||
@@ -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 <stddef.h>
|
||||
|
||||
#include <hip/hip_runtime_api.h>
|
||||
#include <resource_guards.hh>
|
||||
|
||||
template <typename F> 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
|
||||
}
|
||||
+135
@@ -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 <functional>
|
||||
|
||||
#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));
|
||||
}
|
||||
+193
@@ -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));
|
||||
}
|
||||
+96
@@ -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
|
||||
}
|
||||
+137
@@ -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<true>);
|
||||
}
|
||||
|
||||
// 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<true>);
|
||||
}
|
||||
|
||||
/**
|
||||
* 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<true>);
|
||||
}
|
||||
#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));
|
||||
}
|
||||
@@ -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<int>(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT);
|
||||
const auto dst_storage = vkt.CreateMappedStorage<int>(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<int>(count,
|
||||
VK_BUFFER_USAGE_TRANSFER_SRC_BIT);
|
||||
const auto dst_storage = vkt.CreateMappedStorage<int>(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
|
||||
|
||||
|
||||
@@ -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 <vector>
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
|
||||
constexpr bool enable_validation = false;
|
||||
|
||||
template <typename F> void SignalExternalSemaphoreCommon(F f) {
|
||||
VulkanTest vkt(enable_validation);
|
||||
|
||||
constexpr uint32_t count = 1;
|
||||
const auto src_storage = vkt.CreateMappedStorage<int>(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT);
|
||||
const auto dst_storage = vkt.CreateMappedStorage<int>(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 <typename F> 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 <typename F> void SignalExternalMultipleSemaphoresCommon(F f) {
|
||||
VulkanTest vkt(enable_validation);
|
||||
|
||||
constexpr uint32_t count = 1;
|
||||
const auto src_storage = vkt.CreateMappedStorage<int>(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT);
|
||||
const auto dst_storage = vkt.CreateMappedStorage<int>(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 <bool set_params = false>
|
||||
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;
|
||||
}
|
||||
Referens i nytt ärende
Block a user