AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-448586 - Added implementation for new API hipStreamBeginCaptureToGraph

Browse Source 
Change-Id: I1ce802102cef2b66c92d3375f769983841de793f


[ROCm/clr commit: 4feb1f9337]
This commit is contained in:
Anusha GodavarthySurya
2024-02-07 06:02:53 +00:00
parent 4bb028a49d
commit 41ec66ff67
7 changed files with 124 additions and 22 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp
+7 -2
View File
@@ -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;
};
projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h
+40 -1
View File
@@ -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";
projects/clr/hipamd/src/amdhip.def
+2 -1
View File
@@ -460,4 +460,5 @@ hipGraphExternalSemaphoresWaitNodeGetParams
hipGraphExecExternalSemaphoresSignalNodeSetParams
hipGraphExecExternalSemaphoresWaitNodeSetParams
hipGraphAddNode
hipGraphInstantiateWithParams
hipGraphInstantiateWithParams
hipStreamBeginCaptureToGraph
projects/clr/hipamd/src/hip_api_trace.cpp
+19 -12
View File
@@ -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(<table>, <functor>, 8)
//
// HIP_ENFORCE_ABI_VERSIONING(<table>, 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
projects/clr/hipamd/src/hip_graph.cpp
+47 -5
View File
@@ -33,7 +33,8 @@ amd::Monitor g_captureStreamsLock{"StreamCaptureGlobalList"};
amd::Monitor g_streamSetLock{"StreamCaptureset"};
std::unordered_set<hip::Stream*> 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<hip::Graph*>(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<hip::Graph*>(graph);
const std::vector<Node> 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<hipGraphNode_t*>(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);
projects/clr/hipamd/src/hip_hcc.map.in
+1
View File
@@ -552,6 +552,7 @@ hip_6.1 {
global:
hipGraphInstantiateWithParams;
hipGetProcAddress;
hipStreamBeginCaptureToGraph;
local:
*;
} hip_6.0;
projects/clr/hipamd/src/hip_table_interface.cpp
+8 -1
View File
@@ -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);
}
}
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);
}