diff --git a/hipamd/include/hip/amd_detail/hip_prof_str.h b/hipamd/include/hip/amd_detail/hip_prof_str.h index 39fda98129..592beed9b2 100644 --- a/hipamd/include/hip/amd_detail/hip_prof_str.h +++ b/hipamd/include/hip/amd_detail/hip_prof_str.h @@ -236,14 +236,23 @@ enum hip_api_id_t { HIP_API_ID_hipGetTextureObjectResourceViewDesc = HIP_API_ID_NONE, HIP_API_ID_hipGetTextureObjectTextureDesc = HIP_API_ID_NONE, HIP_API_ID_hipGetTextureReference = HIP_API_ID_NONE, + HIP_API_ID_hipGraphAddDependencies = HIP_API_ID_NONE, HIP_API_ID_hipGraphAddKernelNode = HIP_API_ID_NONE, HIP_API_ID_hipGraphAddMemcpyNode = HIP_API_ID_NONE, HIP_API_ID_hipGraphAddMemsetNode = HIP_API_ID_NONE, HIP_API_ID_hipGraphCreate = HIP_API_ID_NONE, HIP_API_ID_hipGraphDestroy = HIP_API_ID_NONE, HIP_API_ID_hipGraphExecDestroy = HIP_API_ID_NONE, + HIP_API_ID_hipGraphGetNodes = HIP_API_ID_NONE, + HIP_API_ID_hipGraphGetRootNodes = HIP_API_ID_NONE, HIP_API_ID_hipGraphInstantiate = HIP_API_ID_NONE, + HIP_API_ID_hipGraphKernelNodeGetParams = HIP_API_ID_NONE, + HIP_API_ID_hipGraphKernelNodeSetParams = HIP_API_ID_NONE, HIP_API_ID_hipGraphLaunch = HIP_API_ID_NONE, + HIP_API_ID_hipGraphMemcpyNodeGetParams = HIP_API_ID_NONE, + HIP_API_ID_hipGraphMemcpyNodeSetParams = HIP_API_ID_NONE, + HIP_API_ID_hipGraphMemsetNodeGetParams = HIP_API_ID_NONE, + HIP_API_ID_hipGraphMemsetNodeSetParams = HIP_API_ID_NONE, HIP_API_ID_hipGraphicsGLRegisterBuffer = HIP_API_ID_NONE, HIP_API_ID_hipGraphicsMapResources = HIP_API_ID_NONE, HIP_API_ID_hipGraphicsResourceGetMappedPointer = HIP_API_ID_NONE, @@ -3136,6 +3145,8 @@ typedef struct hip_api_data_s { #define INIT_hipGetTextureObjectTextureDesc_CB_ARGS_DATA(cb_data) {}; // hipGetTextureReference() #define INIT_hipGetTextureReference_CB_ARGS_DATA(cb_data) {}; +// hipGraphAddDependencies() +#define INIT_hipGraphAddDependencies_CB_ARGS_DATA(cb_data) {}; // hipGraphAddKernelNode() #define INIT_hipGraphAddKernelNode_CB_ARGS_DATA(cb_data) {}; // hipGraphAddMemcpyNode() @@ -3148,10 +3159,26 @@ typedef struct hip_api_data_s { #define INIT_hipGraphDestroy_CB_ARGS_DATA(cb_data) {}; // hipGraphExecDestroy() #define INIT_hipGraphExecDestroy_CB_ARGS_DATA(cb_data) {}; +// hipGraphGetNodes() +#define INIT_hipGraphGetNodes_CB_ARGS_DATA(cb_data) {}; +// hipGraphGetRootNodes() +#define INIT_hipGraphGetRootNodes_CB_ARGS_DATA(cb_data) {}; // hipGraphInstantiate() #define INIT_hipGraphInstantiate_CB_ARGS_DATA(cb_data) {}; +// hipGraphKernelNodeGetParams() +#define INIT_hipGraphKernelNodeGetParams_CB_ARGS_DATA(cb_data) {}; +// hipGraphKernelNodeSetParams() +#define INIT_hipGraphKernelNodeSetParams_CB_ARGS_DATA(cb_data) {}; // hipGraphLaunch() #define INIT_hipGraphLaunch_CB_ARGS_DATA(cb_data) {}; +// hipGraphMemcpyNodeGetParams() +#define INIT_hipGraphMemcpyNodeGetParams_CB_ARGS_DATA(cb_data) {}; +// hipGraphMemcpyNodeSetParams() +#define INIT_hipGraphMemcpyNodeSetParams_CB_ARGS_DATA(cb_data) {}; +// hipGraphMemsetNodeGetParams() +#define INIT_hipGraphMemsetNodeGetParams_CB_ARGS_DATA(cb_data) {}; +// hipGraphMemsetNodeSetParams() +#define INIT_hipGraphMemsetNodeSetParams_CB_ARGS_DATA(cb_data) {}; // hipGraphicsGLRegisterBuffer() #define INIT_hipGraphicsGLRegisterBuffer_CB_ARGS_DATA(cb_data) {}; // hipGraphicsMapResources() diff --git a/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h b/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h index adddb2f484..1e1fbd13a4 100644 --- a/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -2292,6 +2292,64 @@ inline static hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGr cudaGraphAddMemsetNode(pGraphNode, graph, pDependencies, numDependencies, pMemsetParams)); } +inline static hipError_t hipGraphGetNodes(hipGraph_t graph, hipGraphNode_t* nodes, + size_t* numNodes) { + return hipCUDAErrorTohipError(cudaGraphGetNodes(graph, nodes, numNodes)); +} + +inline static hipError_t hipGraphGetRootNodes(hipGraph_t graph, hipGraphNode_t* pRootNodes, + size_t* pNumRootNodes) { + return hipCUDAErrorTohipError(cudaGraphGetRootNodes(graph, pRootNodes, pNumRootNodes)); +} + +inline static hipError_t hipGraphKernelNodeGetParams(hipGraphNode_t node, + hipKernelNodeParams* pNodeParams) { + return hipCUDAErrorTohipError(cudaGraphKernelNodeGetParams(node, pNodeParams)); +} + +inline static hipError_t hipGraphKernelNodeSetParams(hipGraphNode_t node, + const hipKernelNodeParams* pNodeParams) { + return hipCUDAErrorTohipError(cudaGraphKernelNodeSetParams(node, pNodeParams)); +} + +inline static hipError_t hipGraphMemcpyNodeGetParams(hipGraphNode_t node, + hipMemcpy3DParms* pNodeParams) { + return hipCUDAErrorTohipError(cudaGraphMemcpyNodeGetParams(node, pNodeParams)); +} + +inline static hipError_t hipGraphMemcpyNodeSetParams(hipGraphNode_t node, + const hipMemcpy3DParms* pNodeParams) { + return hipCUDAErrorTohipError(cudaGraphMemcpyNodeSetParams(node, pNodeParams)); +} + +inline static hipError_t hipGraphMemsetNodeGetParams(hipGraphNode_t node, + hipMemsetParams* pNodeParams) { + return hipCUDAErrorTohipError(cudaGraphMemsetNodeGetParams(node, pNodeParams)); +} + +inline static hipError_t hipGraphMemsetNodeSetParams(hipGraphNode_t node, + const hipMemsetParams* pNodeParams) { + return hipCUDAErrorTohipError(cudaGraphMemsetNodeSetParams(node, pNodeParams)); +} + +inline static hipError_t hipGraphExecKernelNodeSetParams(hipGraphExec_t hGraphExec, + hipGraphNode_t node, + const hipKernelNodeParams* pNodeParams) { + return hipCUDAErrorTohipError(cudaGraphExecKernelNodeSetParams(hGraphExec, node, pNodeParams)); +} + +inline static hipError_t hipGraphAddDependencies(hipGraph_t graph, const hipGraphNode_t* from, + const hipGraphNode_t* to, size_t numDependencies) { + return hipCUDAErrorTohipError(cudaGraphAddDependencies(graph, from, to, numDependencies)); +} + +inline static hipError_t hipGraphAddEmptyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, + size_t numDependencies) { + return hipCUDAErrorTohipError( + cudaGraphAddEmptyNode(pGraphNode, graph, pDependencies, numDependencies)); +} + inline static hipError_t hipStreamWriteValue32(hipStream_t stream, void* ptr, int32_t value, unsigned int flags) { if (value < 0) { diff --git a/hipamd/src/amdhip.def b/hipamd/src/amdhip.def index 559a08289d..a0923a5e0d 100644 --- a/hipamd/src/amdhip.def +++ b/hipamd/src/amdhip.def @@ -291,3 +291,14 @@ hipImportExternalSemaphore hipSignalExternalSemaphoresAsync hipWaitExternalSemaphoresAsync hipDestroyExternalSemaphore +hipGraphGetNodes +hipGraphGetRootNodes +hipGraphKernelNodeGetParams +hipGraphKernelNodeSetParams +hipGraphMemcpyNodeGetParams +hipGraphMemcpyNodeSetParams +hipGraphMemsetNodeGetParams +hipGraphMemsetNodeSetParams +hipGraphExecKernelNodeSetParams +hipGraphAddDependencies +hipGraphAddEmptyNode diff --git a/hipamd/src/hip_graph.cpp b/hipamd/src/hip_graph.cpp index 2aa4b896b4..6e8235ec4d 100644 --- a/hipamd/src/hip_graph.cpp +++ b/hipamd/src/hip_graph.cpp @@ -413,3 +413,100 @@ hipError_t hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream) { HIP_INIT_API(hipGraphLaunch, graphExec, stream); HIP_RETURN_DURATION(ihipGraphlaunch(graphExec, stream)); } + +hipError_t hipGraphGetNodes(hipGraph_t graph, hipGraphNode_t* nodes, size_t* numNodes) { + HIP_INIT_API(hipGraphGetNodes, graph, nodes, numNodes); + if (graph == nullptr || numNodes == nullptr) { + *numNodes = graph->GetNodeCount(); + } + if (*numNodes > 0) { + nodes = graph->GetNodes().data(); + } + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphGetRootNodes(hipGraph_t graph, hipGraphNode_t* pRootNodes, + size_t* pNumRootNodes) { + HIP_INIT_API(hipGraphGetRootNodes, graph, pRootNodes, pNumRootNodes); + if (graph == nullptr || pNumRootNodes == nullptr) { + return HIP_RETURN(hipErrorInvalidValue); + } + std::vector rootNodes = graph->GetRootNodes(); + pRootNodes = rootNodes.data(); + *pNumRootNodes = rootNodes.size(); + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphKernelNodeGetParams(hipGraphNode_t node, hipKernelNodeParams* pNodeParams) { + HIP_INIT_API(hipGraphKernelNodeGetParams, node, pNodeParams); + if (node == nullptr || pNodeParams == nullptr) { + return HIP_RETURN(hipErrorInvalidValue); + } + reinterpret_cast(node)->GetParams(pNodeParams); + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphKernelNodeSetParams(hipGraphNode_t node, + const hipKernelNodeParams* pNodeParams) { + HIP_INIT_API(hipGraphKernelNodeSetParams, node, pNodeParams); + if (node == nullptr || pNodeParams == nullptr) { + return HIP_RETURN(hipErrorInvalidValue); + } + reinterpret_cast(node)->SetParams(pNodeParams); + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphMemcpyNodeGetParams(hipGraphNode_t node, hipMemcpy3DParms* pNodeParams) { + HIP_INIT_API(hipGraphMemcpyNodeGetParams, node, pNodeParams); + if (node == nullptr || pNodeParams == nullptr) { + return HIP_RETURN(hipErrorInvalidValue); + } + reinterpret_cast(node)->GetParams(pNodeParams); + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphMemcpyNodeSetParams(hipGraphNode_t node, const hipMemcpy3DParms* pNodeParams) { + HIP_INIT_API(hipGraphMemcpyNodeSetParams, node, pNodeParams); + if (node == nullptr || pNodeParams == nullptr) { + return HIP_RETURN(hipErrorInvalidValue); + } + reinterpret_cast(node)->SetParams(pNodeParams); + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphMemsetNodeGetParams(hipGraphNode_t node, hipMemsetParams* pNodeParams) { + HIP_INIT_API(hipGraphMemsetNodeGetParams, node, pNodeParams); + if (node == nullptr || pNodeParams == nullptr) { + return HIP_RETURN(hipErrorInvalidValue); + } + reinterpret_cast(node)->GetParams(pNodeParams); + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphMemsetNodeSetParams(hipGraphNode_t node, const hipMemsetParams* pNodeParams) { + HIP_INIT_API(hipGraphMemsetNodeSetParams, node, pNodeParams); + if (node == nullptr || pNodeParams == nullptr) { + return HIP_RETURN(hipErrorInvalidValue); + } + reinterpret_cast(node)->SetParams(pNodeParams); + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphAddDependencies(hipGraph_t graph, const hipGraphNode_t* from, + const hipGraphNode_t* to, size_t numDependencies) { + HIP_INIT_API(hipGraphAddDependencies, graph, from, to, numDependencies); + if (graph == nullptr) { + return HIP_RETURN(hipErrorInvalidValue); + } + if (numDependencies == 0) { + HIP_RETURN(hipSuccess); + } else if (from == nullptr || to == nullptr) { + return HIP_RETURN(hipErrorInvalidValue); + } + for (size_t i = 0; i < numDependencies; i++) { + if (graph->AddEdge(from[i], to[i]) != hipSuccess) { + HIP_RETURN(hipErrorInvalidValue); + } + } + HIP_RETURN(hipSuccess); +} diff --git a/hipamd/src/hip_graph_internal.hpp b/hipamd/src/hip_graph_internal.hpp index b70717dc80..7b949eca94 100644 --- a/hipamd/src/hip_graph_internal.hpp +++ b/hipamd/src/hip_graph_internal.hpp @@ -120,7 +120,7 @@ class hipGraphKernelNode : public hipGraphNode { void GetParams(hipKernelNodeParams* params) { std::memcpy(params, pKernelParams_, sizeof(hipKernelNodeParams)); } - void SetParams(hipKernelNodeParams* params) { + void SetParams(const hipKernelNodeParams* params) { std::memcpy(pKernelParams_, params, sizeof(hipKernelNodeParams)); } }; @@ -145,7 +145,7 @@ class hipGraphMemcpyNode : public hipGraphNode { void GetParams(hipMemcpy3DParms* params) { std::memcpy(params, pCopyParams_, sizeof(hipMemcpy3DParms)); } - void SetParams(hipMemcpy3DParms* params) { + void SetParams(const hipMemcpy3DParms* params) { std::memcpy(pCopyParams_, params, sizeof(hipMemcpy3DParms)); } }; @@ -263,7 +263,7 @@ class hipGraphMemsetNode : public hipGraphNode { void GetParams(hipMemsetParams* params) { std::memcpy(params, pMemsetParams_, sizeof(hipMemsetParams)); } - void SetParams(hipMemsetParams* params) { + void SetParams(const hipMemsetParams* params) { std::memcpy(pMemsetParams_, params, sizeof(hipMemsetParams)); } }; diff --git a/hipamd/src/hip_hcc.def.in b/hipamd/src/hip_hcc.def.in index 2778cb13be..ac9b168cb8 100755 --- a/hipamd/src/hip_hcc.def.in +++ b/hipamd/src/hip_hcc.def.in @@ -297,3 +297,14 @@ hipGraphicsMapResources hipGraphicsResourceGetMappedPointer hipGraphicsUnmapResources hipGraphicsUnregisterResource +hipGraphGetNodes +hipGraphGetRootNodes +hipGraphKernelNodeGetParams +hipGraphKernelNodeSetParams +hipGraphMemcpyNodeGetParams +hipGraphMemcpyNodeSetParams +hipGraphMemsetNodeGetParams +hipGraphMemsetNodeSetParams +hipGraphExecKernelNodeSetParams +hipGraphAddDependencies +hipGraphAddEmptyNode diff --git a/hipamd/src/hip_hcc.map.in b/hipamd/src/hip_hcc.map.in index 381be32823..05b69fa777 100755 --- a/hipamd/src/hip_hcc.map.in +++ b/hipamd/src/hip_hcc.map.in @@ -314,3 +314,20 @@ global: local: *; } hip_4.2; + +hip_4.4 { +global: + hipGraphGetNodes; + hipGraphGetRootNodes; + hipGraphKernelNodeGetParams; + hipGraphKernelNodeSetParams; + hipGraphMemcpyNodeGetParams; + hipGraphMemcpyNodeSetParams; + hipGraphMemsetNodeGetParams; + hipGraphMemsetNodeSetParams; + hipGraphExecKernelNodeSetParams; + hipGraphAddDependencies; + hipGraphAddEmptyNode; +local: + *; +}hip_4.3;