From 4feb1f9337ae1178f8c1ecc988952871979aedb1 Mon Sep 17 00:00:00 2001 From: Anusha GodavarthySurya Date: Wed, 7 Feb 2024 06:02:53 +0000 Subject: [PATCH] SWDEV-448586 - Added implementation for new API hipStreamBeginCaptureToGraph Change-Id: I1ce802102cef2b66c92d3375f769983841de793f --- .../include/hip/amd_detail/hip_api_trace.hpp | 9 +++- hipamd/include/hip/amd_detail/hip_prof_str.h | 41 ++++++++++++++- hipamd/src/amdhip.def | 3 +- hipamd/src/hip_api_trace.cpp | 31 ++++++----- hipamd/src/hip_graph.cpp | 52 +++++++++++++++++-- hipamd/src/hip_hcc.map.in | 1 + hipamd/src/hip_table_interface.cpp | 9 +++- 7 files changed, 124 insertions(+), 22 deletions(-) diff --git a/hipamd/include/hip/amd_detail/hip_api_trace.hpp b/hipamd/include/hip/amd_detail/hip_api_trace.hpp index 65e9997dc1..4685890336 100644 --- a/hipamd/include/hip/amd_detail/hip_api_trace.hpp +++ b/hipamd/include/hip/amd_detail/hip_api_trace.hpp @@ -61,7 +61,7 @@ // - Reset any of the *_STEP_VERSION defines to zero if the corresponding *_MAJOR_VERSION increases #define HIP_API_TABLE_STEP_VERSION 0 #define HIP_COMPILER_API_TABLE_STEP_VERSION 0 -#define HIP_RUNTIME_API_TABLE_STEP_VERSION 1 +#define HIP_RUNTIME_API_TABLE_STEP_VERSION 2 // HIP API interface typedef hipError_t (*t___hipPopCallConfiguration)(dim3* gridDim, dim3* blockDim, size_t* sharedMem, @@ -947,7 +947,11 @@ typedef hipError_t (*t_hipTexRefGetBorderColor)(float* pBorderColor, typedef hipError_t (*t_hipTexRefGetArray)(hipArray_t* pArray, const textureReference* texRef); typedef hipError_t (*t_hipGetProcAddress)(const char* symbol, void** pfn, int hipVersion, uint64_t flags, hipDriverProcAddressQueryResult* symbolStatus); - +typedef hipError_t (*t_hipStreamBeginCaptureToGraph)(hipStream_t stream, hipGraph_t graph, + const hipGraphNode_t* dependencies, + const hipGraphEdgeData* dependencyData, + size_t numDependencies, + hipStreamCaptureMode mode); // HIP Compiler dispatch table struct HipCompilerDispatchTable { size_t size; @@ -1412,4 +1416,5 @@ struct HipDispatchTable { t_hipTexRefGetBorderColor hipTexRefGetBorderColor_fn; t_hipTexRefGetArray hipTexRefGetArray_fn; t_hipGetProcAddress hipGetProcAddress_fn; + t_hipStreamBeginCaptureToGraph hipStreamBeginCaptureToGraph_fn; }; diff --git a/hipamd/include/hip/amd_detail/hip_prof_str.h b/hipamd/include/hip/amd_detail/hip_prof_str.h index 9a9655577f..55b2679050 100644 --- a/hipamd/include/hip/amd_detail/hip_prof_str.h +++ b/hipamd/include/hip/amd_detail/hip_prof_str.h @@ -408,7 +408,8 @@ enum hip_api_id_t { HIP_API_ID_hipDrvGraphExecMemsetNodeSetParams = 388, HIP_API_ID_hipTexRefGetArray = 389, HIP_API_ID_hipTexRefGetBorderColor = 390, - HIP_API_ID_LAST = 390, + HIP_API_ID_hipStreamBeginCaptureToGraph = 391, + HIP_API_ID_LAST = 391, HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice), HIP_API_ID_hipGetDeviceProperties = HIP_API_ID_CONCAT(HIP_API_ID_,hipGetDeviceProperties), @@ -795,6 +796,7 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipStreamAddCallback: return "hipStreamAddCallback"; case HIP_API_ID_hipStreamAttachMemAsync: return "hipStreamAttachMemAsync"; case HIP_API_ID_hipStreamBeginCapture: return "hipStreamBeginCapture"; + case HIP_API_ID_hipStreamBeginCaptureToGraph: return "hipStreamBeginCaptureToGraph"; case HIP_API_ID_hipStreamCreate: return "hipStreamCreate"; case HIP_API_ID_hipStreamCreateWithFlags: return "hipStreamCreateWithFlags"; case HIP_API_ID_hipStreamCreateWithPriority: return "hipStreamCreateWithPriority"; @@ -1188,6 +1190,7 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipStreamAddCallback", name) == 0) return HIP_API_ID_hipStreamAddCallback; if (strcmp("hipStreamAttachMemAsync", name) == 0) return HIP_API_ID_hipStreamAttachMemAsync; if (strcmp("hipStreamBeginCapture", name) == 0) return HIP_API_ID_hipStreamBeginCapture; + if (strcmp("hipStreamBeginCaptureToGraph", name) == 0) return HIP_API_ID_hipStreamBeginCaptureToGraph; if (strcmp("hipStreamCreate", name) == 0) return HIP_API_ID_hipStreamCreate; if (strcmp("hipStreamCreateWithFlags", name) == 0) return HIP_API_ID_hipStreamCreateWithFlags; if (strcmp("hipStreamCreateWithPriority", name) == 0) return HIP_API_ID_hipStreamCreateWithPriority; @@ -3267,6 +3270,16 @@ typedef struct hip_api_data_s { hipStream_t stream; hipStreamCaptureMode mode; } hipStreamBeginCapture; + struct { + hipStream_t stream; + hipGraph_t graph; + const hipGraphNode_t* dependencies; + hipGraphNode_t dependencies__val; + const hipGraphEdgeData* dependencyData; + hipGraphEdgeData dependencyData__val; + size_t numDependencies; + hipStreamCaptureMode mode; + } hipStreamBeginCaptureToGraph; struct { hipStream_t* stream; hipStream_t stream__val; @@ -5582,6 +5595,15 @@ typedef struct hip_api_data_s { cb_data.args.hipStreamBeginCapture.stream = (hipStream_t)stream; \ cb_data.args.hipStreamBeginCapture.mode = (hipStreamCaptureMode)mode; \ }; +// hipStreamBeginCaptureToGraph[('hipStream_t', 'stream'), ('hipGraph_t', 'graph'), ('const hipGraphNode_t*', 'dependencies'), ('const hipGraphEdgeData*', 'dependencyData'), ('size_t', 'numDependencies'), ('hipStreamCaptureMode', 'mode')] +#define INIT_hipStreamBeginCaptureToGraph_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipStreamBeginCaptureToGraph.stream = (hipStream_t)stream; \ + cb_data.args.hipStreamBeginCaptureToGraph.graph = (hipGraph_t)graph; \ + cb_data.args.hipStreamBeginCaptureToGraph.dependencies = (const hipGraphNode_t*)dependencies; \ + cb_data.args.hipStreamBeginCaptureToGraph.dependencyData = (const hipGraphEdgeData*)dependencyData; \ + cb_data.args.hipStreamBeginCaptureToGraph.numDependencies = (size_t)numDependencies; \ + cb_data.args.hipStreamBeginCaptureToGraph.mode = (hipStreamCaptureMode)mode; \ +}; // hipStreamCreate[('hipStream_t*', 'stream')] #define INIT_hipStreamCreate_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipStreamCreate.stream = (hipStream_t*)stream; \ @@ -7231,6 +7253,11 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipStreamBeginCapture[('hipStream_t', 'stream'), ('hipStreamCaptureMode', 'mode')] case HIP_API_ID_hipStreamBeginCapture: break; +// hipStreamBeginCaptureToGraph[('hipStream_t', 'stream'), ('hipGraph_t', 'graph'), ('const hipGraphNode_t*', 'dependencies'), ('const hipGraphEdgeData*', 'dependencyData'), ('size_t', 'numDependencies'), ('hipStreamCaptureMode', 'mode')] + case HIP_API_ID_hipStreamBeginCaptureToGraph: + if (data->args.hipStreamBeginCaptureToGraph.dependencies) data->args.hipStreamBeginCaptureToGraph.dependencies__val = *(data->args.hipStreamBeginCaptureToGraph.dependencies); + if (data->args.hipStreamBeginCaptureToGraph.dependencyData) data->args.hipStreamBeginCaptureToGraph.dependencyData__val = *(data->args.hipStreamBeginCaptureToGraph.dependencyData); + break; // hipStreamCreate[('hipStream_t*', 'stream')] case HIP_API_ID_hipStreamCreate: if (data->args.hipStreamCreate.stream) data->args.hipStreamCreate.stream__val = *(data->args.hipStreamCreate.stream); @@ -10156,6 +10183,18 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << ", mode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBeginCapture.mode); oss << ")"; break; + case HIP_API_ID_hipStreamBeginCaptureToGraph: + oss << "hipStreamBeginCaptureToGraph("; + oss << "stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBeginCaptureToGraph.stream); + oss << ", graph="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBeginCaptureToGraph.graph); + if (data->args.hipStreamBeginCaptureToGraph.dependencies == NULL) oss << ", dependencies=NULL"; + else { oss << ", dependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBeginCaptureToGraph.dependencies__val); } + if (data->args.hipStreamBeginCaptureToGraph.dependencyData == NULL) oss << ", dependencyData=NULL"; + else { oss << ", dependencyData="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBeginCaptureToGraph.dependencyData__val); } + oss << ", numDependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBeginCaptureToGraph.numDependencies); + oss << ", mode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBeginCaptureToGraph.mode); + oss << ")"; + break; case HIP_API_ID_hipStreamCreate: oss << "hipStreamCreate("; if (data->args.hipStreamCreate.stream == NULL) oss << "stream=NULL"; diff --git a/hipamd/src/amdhip.def b/hipamd/src/amdhip.def index 126c6fe0b8..d9764ab4af 100644 --- a/hipamd/src/amdhip.def +++ b/hipamd/src/amdhip.def @@ -460,4 +460,5 @@ hipGraphExternalSemaphoresWaitNodeGetParams hipGraphExecExternalSemaphoresSignalNodeSetParams hipGraphExecExternalSemaphoresWaitNodeSetParams hipGraphAddNode -hipGraphInstantiateWithParams \ No newline at end of file +hipGraphInstantiateWithParams +hipStreamBeginCaptureToGraph diff --git a/hipamd/src/hip_api_trace.cpp b/hipamd/src/hip_api_trace.cpp index 917e142016..50de6878e6 100644 --- a/hipamd/src/hip_api_trace.cpp +++ b/hipamd/src/hip_api_trace.cpp @@ -763,6 +763,10 @@ hipError_t hipTexRefGetBorderColor(float* pBorderColor, const textureReference* hipError_t hipTexRefGetArray(hipArray_t* pArray, const textureReference* texRef); hipError_t hipGetProcAddress(const char* symbol, void** pfn, int hipVersion, uint64_t flags, hipDriverProcAddressQueryResult* symbolStatus = NULL); +hipError_t hipStreamBeginCaptureToGraph(hipStream_t stream, hipGraph_t graph, + const hipGraphNode_t* dependencies, + const hipGraphEdgeData* dependencyData, + size_t numDependencies, hipStreamCaptureMode mode); } // namespace hip namespace hip { @@ -1237,6 +1241,7 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) { ptrDispatchTable->hipTexRefGetBorderColor_fn = hip::hipTexRefGetBorderColor; ptrDispatchTable->hipTexRefGetArray_fn = hip::hipTexRefGetArray; ptrDispatchTable->hipGetProcAddress_fn = hip::hipGetProcAddress; + ptrDispatchTable->hipStreamBeginCaptureToGraph_fn = hip::hipStreamBeginCaptureToGraph; } #if HIP_ROCPROFILER_REGISTER > 0 @@ -1783,20 +1788,22 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipStreamGetCaptureInfo_v2_spt_fn, 425) HIP_ENFORCE_ABI(HipDispatchTable, hipLaunchHostFunc_spt_fn, 426) HIP_ENFORCE_ABI(HipDispatchTable, hipGetStreamDeviceId_fn, 427) HIP_ENFORCE_ABI(HipDispatchTable, hipDrvGraphAddMemsetNode_fn, 428) -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphAddExternalSemaphoresWaitNode_fn, 429) -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphAddExternalSemaphoresSignalNode_fn, 430) -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExternalSemaphoresSignalNodeSetParams_fn, 431) -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExternalSemaphoresWaitNodeSetParams_fn, 432) -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExternalSemaphoresSignalNodeGetParams_fn, 433) -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExternalSemaphoresWaitNodeGetParams_fn, 434) -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExecExternalSemaphoresSignalNodeSetParams_fn, 435) -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExecExternalSemaphoresWaitNodeSetParams_fn, 436) -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphAddNode_fn, 437) -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphInstantiateWithParams_fn, 438) +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphAddExternalSemaphoresWaitNode_fn, 429); +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphAddExternalSemaphoresSignalNode_fn, 430); +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExternalSemaphoresSignalNodeSetParams_fn, 431); +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExternalSemaphoresWaitNodeSetParams_fn, 432); +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExternalSemaphoresSignalNodeGetParams_fn, 433); +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExternalSemaphoresWaitNodeGetParams_fn, 434); +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExecExternalSemaphoresSignalNodeSetParams_fn, 435); +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExecExternalSemaphoresWaitNodeSetParams_fn, 436); +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphAddNode_fn, 437); +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphInstantiateWithParams_fn, 438); HIP_ENFORCE_ABI(HipDispatchTable, hipExtGetLastError_fn, 439) HIP_ENFORCE_ABI(HipDispatchTable, hipTexRefGetBorderColor_fn, 440) HIP_ENFORCE_ABI(HipDispatchTable, hipTexRefGetArray_fn, 441) HIP_ENFORCE_ABI(HipDispatchTable, hipGetProcAddress_fn, 442) +HIP_ENFORCE_ABI(HipDispatchTable, hipStreamBeginCaptureToGraph_fn, 443); + // if HIP_ENFORCE_ABI entries are added for each new function pointer in the table, the number below // will be +1 of the number in the last HIP_ENFORCE_ABI line. E.g.: @@ -1804,9 +1811,9 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipGetProcAddress_fn, 442) // HIP_ENFORCE_ABI(, , 8) // // HIP_ENFORCE_ABI_VERSIONING(
, 9) <- 8 + 1 = 9 -HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 443) +HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 444) -static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 1, +static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 2, "If you get this error, add new HIP_ENFORCE_ABI(...) code for the new function " "pointers and then update this check so it is true"); #endif diff --git a/hipamd/src/hip_graph.cpp b/hipamd/src/hip_graph.cpp index 5f261ddd79..0fff3323ce 100644 --- a/hipamd/src/hip_graph.cpp +++ b/hipamd/src/hip_graph.cpp @@ -33,7 +33,8 @@ amd::Monitor g_captureStreamsLock{"StreamCaptureGlobalList"}; amd::Monitor g_streamSetLock{"StreamCaptureset"}; std::unordered_set g_allCapturingStreams; hipError_t ihipGraphDebugDotPrint(hipGraph_t graph, const char* path, unsigned int flags); - +hipError_t ihipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t* dependencies, + size_t numDependencies, unsigned int flags); inline hipError_t ihipGraphUpload(hipGraphExec_t graphExec, hipStream_t stream) { if (graphExec == nullptr) { @@ -863,7 +864,8 @@ hipError_t hipThreadExchangeStreamCaptureMode(hipStreamCaptureMode* mode) { HIP_RETURN_DURATION(hipSuccess); } -hipError_t hipStreamBeginCapture_common(hipStream_t stream, hipStreamCaptureMode mode) { +hipError_t hipStreamBeginCapture_common(hipStream_t stream, hipStreamCaptureMode mode, + hipGraph_t graph = nullptr) { if (!hip::isValid(stream)) { return hipErrorContextIsDestroyed; } @@ -879,8 +881,11 @@ hipError_t hipStreamBeginCapture_common(hipStream_t stream, hipStreamCaptureMode if (s->GetCaptureStatus() == hipStreamCaptureStatusActive) { return hipErrorIllegalState; } - - s->SetCaptureGraph(new hip::Graph(s->GetDevice())); + if(graph == nullptr) { + s->SetCaptureGraph(new hip::Graph(s->GetDevice())); + } else { + s->SetCaptureGraph(reinterpret_cast(graph)); + } s->SetCaptureId(); s->SetCaptureMode(mode); s->SetOriginStream(); @@ -903,6 +908,37 @@ hipError_t hipStreamBeginCapture(hipStream_t stream, hipStreamCaptureMode mode) HIP_RETURN_DURATION(hipStreamBeginCapture_common(stream, mode)); } +hipError_t hipStreamBeginCaptureToGraph(hipStream_t stream, hipGraph_t graph, + const hipGraphNode_t* dependencies, + const hipGraphEdgeData* dependencyData, + size_t numDependencies, hipStreamCaptureMode mode) { + HIP_INIT_API(hipStreamBeginCapture, stream, graph, dependencies, dependencyData, numDependencies, + mode); + if (dependencyData != nullptr) { + return hipErrorNotSupported; + } else if (graph == nullptr) { + return hipErrorInvalidValue; + } else if (dependencies == nullptr && numDependencies != 0) { + return hipErrorInvalidValue; + } else if (dependencies != nullptr && numDependencies == 0) { + return hipErrorInvalidValue; + } + hip::Graph* g = reinterpret_cast(graph); + const std::vector nodes = g->GetNodes(); + for (auto& node : nodes) { + g->AddManualNodeDuringCapture(node); + } + hipError_t status = hipStreamBeginCapture_common(stream, mode, graph); + if (status != hipSuccess) { + HIP_RETURN_DURATION(status); + } + if (dependencies != nullptr) { + status = ihipStreamUpdateCaptureDependencies(stream, const_cast(dependencies), + numDependencies, hipStreamSetCaptureDependencies); + } + HIP_RETURN_DURATION(status); +} + hipError_t hipStreamBeginCapture_spt(hipStream_t stream, hipStreamCaptureMode mode) { HIP_INIT_API(hipStreamBeginCapture, stream, mode); PER_THREAD_DEFAULT_STREAM(stream); @@ -1774,7 +1810,7 @@ hipError_t hipStreamGetCaptureInfo_v2_spt(hipStream_t stream, dependencies_out, numDependencies_out)); } -hipError_t hipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t* dependencies, +hipError_t ihipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t* dependencies, size_t numDependencies, unsigned int flags) { HIP_INIT_API(hipStreamUpdateCaptureDependencies, stream, dependencies, numDependencies, flags); if (!hip::isValid(stream)) { @@ -1809,6 +1845,12 @@ hipError_t hipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t HIP_RETURN(hipSuccess); } +hipError_t hipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t* dependencies, + size_t numDependencies, unsigned int flags) { + HIP_INIT_API(hipStreamUpdateCaptureDependencies, stream, dependencies, numDependencies, flags); + HIP_RETURN(ihipStreamUpdateCaptureDependencies(stream, dependencies, numDependencies, flags)); +} + hipError_t hipGraphRemoveDependencies(hipGraph_t graph, const hipGraphNode_t* from, const hipGraphNode_t* to, size_t numDependencies) { HIP_INIT_API(hipGraphRemoveDependencies, graph, from, to, numDependencies); diff --git a/hipamd/src/hip_hcc.map.in b/hipamd/src/hip_hcc.map.in index ca41830d68..0ce6ffd2ab 100644 --- a/hipamd/src/hip_hcc.map.in +++ b/hipamd/src/hip_hcc.map.in @@ -552,6 +552,7 @@ hip_6.1 { global: hipGraphInstantiateWithParams; hipGetProcAddress; + hipStreamBeginCaptureToGraph; local: *; } hip_6.0; diff --git a/hipamd/src/hip_table_interface.cpp b/hipamd/src/hip_table_interface.cpp index 11f6f6e013..b561091687 100644 --- a/hipamd/src/hip_table_interface.cpp +++ b/hipamd/src/hip_table_interface.cpp @@ -1732,4 +1732,11 @@ extern "C" hipError_t hipGetProcAddress(const char* symbol, void** pfn, int hip hipDriverProcAddressQueryResult* symbolStatus = nullptr) { return hip::GetHipDispatchTable()->hipGetProcAddress_fn(symbol, pfn, hipVersion, flags, symbolStatus); -} \ No newline at end of file +} +hipError_t hipStreamBeginCaptureToGraph(hipStream_t stream, hipGraph_t graph, + const hipGraphNode_t* dependencies, + const hipGraphEdgeData* dependencyData, + size_t numDependencies, hipStreamCaptureMode mode) { + return hip::GetHipDispatchTable()->hipStreamBeginCaptureToGraph_fn( + stream, graph, dependencies, dependencyData, numDependencies, mode); +}