AI/rocm-systems
2
0
Fork 0
Code Issues Pull-Requests Actions Pakete Projekte Releases Wiki Aktivität

SWDEV-421025 - hipGraphInstantiateWithParams API changes

Quellcode durchsuchen 
Change-Id: Ib07d4dd1698220b68ed27f91d58d3bd315a8804c


[ROCm/clr commit: 1a3901fa49]
Dieser Commit ist enthalten in:
Rahul Manocha
2024-01-15 21:19:00 +00:00
committet von Maneesh Gupta
Ursprung 0aadfd0a6f
Commit 7bccc07147
7 geänderte Dateien mit 245 neuen und 29 gelöschten Zeilen
Patch-Datei herunterladen Vergleichs-Datei herunterladen Alle Dateien ausklappen Alle Dateien einklappen
projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp
+4 -1
Datei anzeigen
@@ -917,7 +917,7 @@ typedef hipError_t (*t_hipGraphAddExternalSemaphoresWaitNode)(hipGraphNode_t* pG
const hipExternalSemaphoreWaitNodeParams* nodeParams);
typedef hipError_t (*t_hipGraphAddExternalSemaphoresSignalNode)(hipGraphNode_t* pGraphNode,
hipGraph_t graph, const hipGraphNode_t* pDependencies,
size_t numDependencies,
size_t numDependencies,
const hipExternalSemaphoreSignalNodeParams* nodeParams);
typedef hipError_t (*t_hipGraphExternalSemaphoresSignalNodeSetParams)(hipGraphNode_t hNode,
const hipExternalSemaphoreSignalNodeParams* nodeParams);
@@ -936,6 +936,8 @@ typedef hipError_t (*t_hipGraphExecExternalSemaphoresWaitNodeSetParams)(hipGraph
typedef hipError_t (*t_hipGraphAddNode)(hipGraphNode_t *pGraphNode, hipGraph_t graph,
const hipGraphNode_t *pDependencies, size_t numDependencies,
hipGraphNodeParams *nodeParams);
typedef hipError_t (*t_hipGraphInstantiateWithParams)(hipGraphExec_t* pGraphExec, hipGraph_t graph,
hipGraphInstantiateParams* instantiateParams);
typedef hipError_t (*t_hipExtGetLastError)();
// HIP Compiler dispatch table
@@ -1397,5 +1399,6 @@ struct HipDispatchTable {
t_hipGraphExecExternalSemaphoresSignalNodeSetParams hipGraphExecExternalSemaphoresSignalNodeSetParams_fn;
t_hipGraphExecExternalSemaphoresWaitNodeSetParams hipGraphExecExternalSemaphoresWaitNodeSetParams_fn;
t_hipGraphAddNode hipGraphAddNode_fn;
t_hipGraphInstantiateWithParams hipGraphInstantiateWithParams_fn;
t_hipExtGetLastError hipExtGetLastError_fn;
};
projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h
+168 -19
Datei anzeigen
@@ -398,7 +398,12 @@ enum hip_api_id_t {
HIP_API_ID_hipGraphExternalSemaphoresWaitNodeSetParams = 378,
HIP_API_ID_hipExtGetLastError = 379,
HIP_API_ID_hipGraphAddNode = 380,
HIP_API_ID_LAST = 380,
HIP_API_ID_hipGetProcAddress = 381,
HIP_API_ID_hipGraphExecGetFlags = 382,
HIP_API_ID_hipGraphExecNodeSetParams = 383,
HIP_API_ID_hipGraphInstantiateWithParams = 384,
HIP_API_ID_hipGraphNodeSetParams = 385,
HIP_API_ID_LAST = 385,
HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice),
HIP_API_ID_hipGetDeviceProperties = HIP_API_ID_CONCAT(HIP_API_ID_,hipGetDeviceProperties),
@@ -523,6 +528,7 @@ static inline const char* hip_api_name(const uint32_t id) {
case HIP_API_ID_hipEventQuery: return "hipEventQuery";
case HIP_API_ID_hipEventRecord: return "hipEventRecord";
case HIP_API_ID_hipEventSynchronize: return "hipEventSynchronize";
case HIP_API_ID_hipExtGetLastError: return "hipExtGetLastError";
case HIP_API_ID_hipExtGetLinkTypeAndHopCount: return "hipExtGetLinkTypeAndHopCount";
case HIP_API_ID_hipExtLaunchKernel: return "hipExtLaunchKernel";
case HIP_API_ID_hipExtLaunchMultiKernelMultiDevice: return "hipExtLaunchMultiKernelMultiDevice";
@@ -552,6 +558,7 @@ static inline const char* hip_api_name(const uint32_t id) {
case HIP_API_ID_hipGetErrorString: return "hipGetErrorString";
case HIP_API_ID_hipGetLastError: return "hipGetLastError";
case HIP_API_ID_hipGetMipmappedArrayLevel: return "hipGetMipmappedArrayLevel";
case HIP_API_ID_hipGetProcAddress: return "hipGetProcAddress";
case HIP_API_ID_hipGetSymbolAddress: return "hipGetSymbolAddress";
case HIP_API_ID_hipGetSymbolSize: return "hipGetSymbolSize";
case HIP_API_ID_hipGraphAddChildGraphNode: return "hipGraphAddChildGraphNode";
@@ -587,6 +594,7 @@ static inline const char* hip_api_name(const uint32_t id) {
case HIP_API_ID_hipGraphExecEventWaitNodeSetEvent: return "hipGraphExecEventWaitNodeSetEvent";
case HIP_API_ID_hipGraphExecExternalSemaphoresSignalNodeSetParams: return "hipGraphExecExternalSemaphoresSignalNodeSetParams";
case HIP_API_ID_hipGraphExecExternalSemaphoresWaitNodeSetParams: return "hipGraphExecExternalSemaphoresWaitNodeSetParams";
case HIP_API_ID_hipGraphExecGetFlags: return "hipGraphExecGetFlags";
case HIP_API_ID_hipGraphExecHostNodeSetParams: return "hipGraphExecHostNodeSetParams";
case HIP_API_ID_hipGraphExecKernelNodeSetParams: return "hipGraphExecKernelNodeSetParams";
case HIP_API_ID_hipGraphExecMemcpyNodeSetParams: return "hipGraphExecMemcpyNodeSetParams";
@@ -594,6 +602,7 @@ static inline const char* hip_api_name(const uint32_t id) {
case HIP_API_ID_hipGraphExecMemcpyNodeSetParamsFromSymbol: return "hipGraphExecMemcpyNodeSetParamsFromSymbol";
case HIP_API_ID_hipGraphExecMemcpyNodeSetParamsToSymbol: return "hipGraphExecMemcpyNodeSetParamsToSymbol";
case HIP_API_ID_hipGraphExecMemsetNodeSetParams: return "hipGraphExecMemsetNodeSetParams";
case HIP_API_ID_hipGraphExecNodeSetParams: return "hipGraphExecNodeSetParams";
case HIP_API_ID_hipGraphExecUpdate: return "hipGraphExecUpdate";
case HIP_API_ID_hipGraphExternalSemaphoresSignalNodeGetParams: return "hipGraphExternalSemaphoresSignalNodeGetParams";
case HIP_API_ID_hipGraphExternalSemaphoresSignalNodeSetParams: return "hipGraphExternalSemaphoresSignalNodeSetParams";
@@ -606,6 +615,7 @@ static inline const char* hip_api_name(const uint32_t id) {
case HIP_API_ID_hipGraphHostNodeSetParams: return "hipGraphHostNodeSetParams";
case HIP_API_ID_hipGraphInstantiate: return "hipGraphInstantiate";
case HIP_API_ID_hipGraphInstantiateWithFlags: return "hipGraphInstantiateWithFlags";
case HIP_API_ID_hipGraphInstantiateWithParams: return "hipGraphInstantiateWithParams";
case HIP_API_ID_hipGraphKernelNodeCopyAttributes: return "hipGraphKernelNodeCopyAttributes";
case HIP_API_ID_hipGraphKernelNodeGetAttribute: return "hipGraphKernelNodeGetAttribute";
case HIP_API_ID_hipGraphKernelNodeGetParams: return "hipGraphKernelNodeGetParams";
@@ -627,6 +637,7 @@ static inline const char* hip_api_name(const uint32_t id) {
case HIP_API_ID_hipGraphNodeGetEnabled: return "hipGraphNodeGetEnabled";
case HIP_API_ID_hipGraphNodeGetType: return "hipGraphNodeGetType";
case HIP_API_ID_hipGraphNodeSetEnabled: return "hipGraphNodeSetEnabled";
case HIP_API_ID_hipGraphNodeSetParams: return "hipGraphNodeSetParams";
case HIP_API_ID_hipGraphReleaseUserObject: return "hipGraphReleaseUserObject";
case HIP_API_ID_hipGraphRemoveDependencies: return "hipGraphRemoveDependencies";
case HIP_API_ID_hipGraphRetainUserObject: return "hipGraphRetainUserObject";
@@ -819,7 +830,6 @@ static inline const char* hip_api_name(const uint32_t id) {
case HIP_API_ID_hipUserObjectRelease: return "hipUserObjectRelease";
case HIP_API_ID_hipUserObjectRetain: return "hipUserObjectRetain";
case HIP_API_ID_hipWaitExternalSemaphoresAsync: return "hipWaitExternalSemaphoresAsync";
case HIP_API_ID_hipExtGetLastError: return "hipExtGetLastError";
};
return "unknown";
};
@@ -906,6 +916,7 @@ static inline uint32_t hipApiIdByName(const char* name) {
if (strcmp("hipEventQuery", name) == 0) return HIP_API_ID_hipEventQuery;
if (strcmp("hipEventRecord", name) == 0) return HIP_API_ID_hipEventRecord;
if (strcmp("hipEventSynchronize", name) == 0) return HIP_API_ID_hipEventSynchronize;
if (strcmp("hipExtGetLastError", name) == 0) return HIP_API_ID_hipExtGetLastError;
if (strcmp("hipExtGetLinkTypeAndHopCount", name) == 0) return HIP_API_ID_hipExtGetLinkTypeAndHopCount;
if (strcmp("hipExtLaunchKernel", name) == 0) return HIP_API_ID_hipExtLaunchKernel;
if (strcmp("hipExtLaunchMultiKernelMultiDevice", name) == 0) return HIP_API_ID_hipExtLaunchMultiKernelMultiDevice;
@@ -935,6 +946,7 @@ static inline uint32_t hipApiIdByName(const char* name) {
if (strcmp("hipGetErrorString", name) == 0) return HIP_API_ID_hipGetErrorString;
if (strcmp("hipGetLastError", name) == 0) return HIP_API_ID_hipGetLastError;
if (strcmp("hipGetMipmappedArrayLevel", name) == 0) return HIP_API_ID_hipGetMipmappedArrayLevel;
if (strcmp("hipGetProcAddress", name) == 0) return HIP_API_ID_hipGetProcAddress;
if (strcmp("hipGetSymbolAddress", name) == 0) return HIP_API_ID_hipGetSymbolAddress;
if (strcmp("hipGetSymbolSize", name) == 0) return HIP_API_ID_hipGetSymbolSize;
if (strcmp("hipGraphAddChildGraphNode", name) == 0) return HIP_API_ID_hipGraphAddChildGraphNode;
@@ -970,6 +982,7 @@ static inline uint32_t hipApiIdByName(const char* name) {
if (strcmp("hipGraphExecEventWaitNodeSetEvent", name) == 0) return HIP_API_ID_hipGraphExecEventWaitNodeSetEvent;
if (strcmp("hipGraphExecExternalSemaphoresSignalNodeSetParams", name) == 0) return HIP_API_ID_hipGraphExecExternalSemaphoresSignalNodeSetParams;
if (strcmp("hipGraphExecExternalSemaphoresWaitNodeSetParams", name) == 0) return HIP_API_ID_hipGraphExecExternalSemaphoresWaitNodeSetParams;
if (strcmp("hipGraphExecGetFlags", name) == 0) return HIP_API_ID_hipGraphExecGetFlags;
if (strcmp("hipGraphExecHostNodeSetParams", name) == 0) return HIP_API_ID_hipGraphExecHostNodeSetParams;
if (strcmp("hipGraphExecKernelNodeSetParams", name) == 0) return HIP_API_ID_hipGraphExecKernelNodeSetParams;
if (strcmp("hipGraphExecMemcpyNodeSetParams", name) == 0) return HIP_API_ID_hipGraphExecMemcpyNodeSetParams;
@@ -977,6 +990,7 @@ static inline uint32_t hipApiIdByName(const char* name) {
if (strcmp("hipGraphExecMemcpyNodeSetParamsFromSymbol", name) == 0) return HIP_API_ID_hipGraphExecMemcpyNodeSetParamsFromSymbol;
if (strcmp("hipGraphExecMemcpyNodeSetParamsToSymbol", name) == 0) return HIP_API_ID_hipGraphExecMemcpyNodeSetParamsToSymbol;
if (strcmp("hipGraphExecMemsetNodeSetParams", name) == 0) return HIP_API_ID_hipGraphExecMemsetNodeSetParams;
if (strcmp("hipGraphExecNodeSetParams", name) == 0) return HIP_API_ID_hipGraphExecNodeSetParams;
if (strcmp("hipGraphExecUpdate", name) == 0) return HIP_API_ID_hipGraphExecUpdate;
if (strcmp("hipGraphExternalSemaphoresSignalNodeGetParams", name) == 0) return HIP_API_ID_hipGraphExternalSemaphoresSignalNodeGetParams;
if (strcmp("hipGraphExternalSemaphoresSignalNodeSetParams", name) == 0) return HIP_API_ID_hipGraphExternalSemaphoresSignalNodeSetParams;
@@ -989,6 +1003,7 @@ static inline uint32_t hipApiIdByName(const char* name) {
if (strcmp("hipGraphHostNodeSetParams", name) == 0) return HIP_API_ID_hipGraphHostNodeSetParams;
if (strcmp("hipGraphInstantiate", name) == 0) return HIP_API_ID_hipGraphInstantiate;
if (strcmp("hipGraphInstantiateWithFlags", name) == 0) return HIP_API_ID_hipGraphInstantiateWithFlags;
if (strcmp("hipGraphInstantiateWithParams", name) == 0) return HIP_API_ID_hipGraphInstantiateWithParams;
if (strcmp("hipGraphKernelNodeCopyAttributes", name) == 0) return HIP_API_ID_hipGraphKernelNodeCopyAttributes;
if (strcmp("hipGraphKernelNodeGetAttribute", name) == 0) return HIP_API_ID_hipGraphKernelNodeGetAttribute;
if (strcmp("hipGraphKernelNodeGetParams", name) == 0) return HIP_API_ID_hipGraphKernelNodeGetParams;
@@ -1010,6 +1025,7 @@ static inline uint32_t hipApiIdByName(const char* name) {
if (strcmp("hipGraphNodeGetEnabled", name) == 0) return HIP_API_ID_hipGraphNodeGetEnabled;
if (strcmp("hipGraphNodeGetType", name) == 0) return HIP_API_ID_hipGraphNodeGetType;
if (strcmp("hipGraphNodeSetEnabled", name) == 0) return HIP_API_ID_hipGraphNodeSetEnabled;
if (strcmp("hipGraphNodeSetParams", name) == 0) return HIP_API_ID_hipGraphNodeSetParams;
if (strcmp("hipGraphReleaseUserObject", name) == 0) return HIP_API_ID_hipGraphReleaseUserObject;
if (strcmp("hipGraphRemoveDependencies", name) == 0) return HIP_API_ID_hipGraphRemoveDependencies;
if (strcmp("hipGraphRetainUserObject", name) == 0) return HIP_API_ID_hipGraphRetainUserObject;
@@ -1202,7 +1218,6 @@ static inline uint32_t hipApiIdByName(const char* name) {
if (strcmp("hipUserObjectRelease", name) == 0) return HIP_API_ID_hipUserObjectRelease;
if (strcmp("hipUserObjectRetain", name) == 0) return HIP_API_ID_hipUserObjectRetain;
if (strcmp("hipWaitExternalSemaphoresAsync", name) == 0) return HIP_API_ID_hipWaitExternalSemaphoresAsync;
if (strcmp("hipExtGetLastError", name) == 0) return HIP_API_ID_hipExtGetLastError;
return HIP_API_ID_NONE;
}
@@ -1738,6 +1753,16 @@ typedef struct hip_api_data_s {
hipMipmappedArray_const_t mipmappedArray;
unsigned int level;
} hipGetMipmappedArrayLevel;
struct {
const char* symbol;
char symbol__val;
void** pfn;
void* pfn__val;
int hipVersion;
uint64_t flags;
hipDriverProcAddressQueryResult* symbolStatus;
hipDriverProcAddressQueryResult symbolStatus__val;
} hipGetProcAddress;
struct {
void** devPtr;
void* devPtr__val;
@@ -1993,6 +2018,11 @@ typedef struct hip_api_data_s {
const hipExternalSemaphoreWaitNodeParams* nodeParams;
hipExternalSemaphoreWaitNodeParams nodeParams__val;
} hipGraphExecExternalSemaphoresWaitNodeSetParams;
struct {
hipGraphExec_t graphExec;
unsigned long long* flags;
unsigned long long flags__val;
} hipGraphExecGetFlags;
struct {
hipGraphExec_t hGraphExec;
hipGraphNode_t node;
@@ -2043,6 +2073,12 @@ typedef struct hip_api_data_s {
const hipMemsetParams* pNodeParams;
hipMemsetParams pNodeParams__val;
} hipGraphExecMemsetNodeSetParams;
struct {
hipGraphExec_t graphExec;
hipGraphNode_t node;
hipGraphNodeParams* nodeParams;
hipGraphNodeParams nodeParams__val;
} hipGraphExecNodeSetParams;
struct {
hipGraphExec_t hGraphExec;
hipGraph_t hGraph;
@@ -2120,6 +2156,13 @@ typedef struct hip_api_data_s {
hipGraph_t graph;
unsigned long long flags;
} hipGraphInstantiateWithFlags;
struct {
hipGraphExec_t* pGraphExec;
hipGraphExec_t pGraphExec__val;
hipGraph_t graph;
hipGraphInstantiateParams* instantiateParams;
hipGraphInstantiateParams instantiateParams__val;
} hipGraphInstantiateWithParams;
struct {
hipGraphNode_t hSrc;
hipGraphNode_t hDst;
@@ -2238,6 +2281,11 @@ typedef struct hip_api_data_s {
hipGraphNode_t hNode;
unsigned int isEnabled;
} hipGraphNodeSetEnabled;
struct {
hipGraphNode_t node;
hipGraphNodeParams* nodeParams;
hipGraphNodeParams nodeParams__val;
} hipGraphNodeSetParams;
struct {
hipGraph_t graph;
hipUserObject_t object;
@@ -3984,6 +4032,9 @@ typedef struct hip_api_data_s {
cb_data.args.hipGetMipmappedArrayLevel.mipmappedArray = (hipMipmappedArray_const_t)mipmappedArray; \
cb_data.args.hipGetMipmappedArrayLevel.level = (unsigned int)level; \
};
// hipGetProcAddress[('const char*', 'symbol'), ('void**', 'pfn'), ('int', 'hipVersion'), ('uint64_t', 'flags'), ('hipDriverProcAddressQueryResult*', 'symbolStatus')]
#define INIT_hipGetProcAddress_CB_ARGS_DATA(cb_data) { \
};
// hipGetSymbolAddress[('void**', 'devPtr'), ('const void*', 'symbol')]
#define INIT_hipGetSymbolAddress_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipGetSymbolAddress.devPtr = (void**)devPtr; \
@@ -4034,9 +4085,19 @@ typedef struct hip_api_data_s {
};
// hipGraphAddExternalSemaphoresSignalNode[('hipGraphNode_t*', 'pGraphNode'), ('hipGraph_t', 'graph'), ('const hipGraphNode_t*', 'pDependencies'), ('size_t', 'numDependencies'), ('const hipExternalSemaphoreSignalNodeParams*', 'nodeParams')]
#define INIT_hipGraphAddExternalSemaphoresSignalNode_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipGraphAddExternalSemaphoresSignalNode.pGraphNode = (hipGraphNode_t*)pGraphNode; \
cb_data.args.hipGraphAddExternalSemaphoresSignalNode.graph = (hipGraph_t)graph; \
cb_data.args.hipGraphAddExternalSemaphoresSignalNode.pDependencies = (const hipGraphNode_t*)pDependencies; \
cb_data.args.hipGraphAddExternalSemaphoresSignalNode.numDependencies = (size_t)numDependencies; \
cb_data.args.hipGraphAddExternalSemaphoresSignalNode.nodeParams = (const hipExternalSemaphoreSignalNodeParams*)nodeParams; \
};
// hipGraphAddExternalSemaphoresWaitNode[('hipGraphNode_t*', 'pGraphNode'), ('hipGraph_t', 'graph'), ('const hipGraphNode_t*', 'pDependencies'), ('size_t', 'numDependencies'), ('const hipExternalSemaphoreWaitNodeParams*', 'nodeParams')]
#define INIT_hipGraphAddExternalSemaphoresWaitNode_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipGraphAddExternalSemaphoresWaitNode.pGraphNode = (hipGraphNode_t*)pGraphNode; \
cb_data.args.hipGraphAddExternalSemaphoresWaitNode.graph = (hipGraph_t)graph; \
cb_data.args.hipGraphAddExternalSemaphoresWaitNode.pDependencies = (const hipGraphNode_t*)pDependencies; \
cb_data.args.hipGraphAddExternalSemaphoresWaitNode.numDependencies = (size_t)numDependencies; \
cb_data.args.hipGraphAddExternalSemaphoresWaitNode.nodeParams = (const hipExternalSemaphoreWaitNodeParams*)nodeParams; \
};
// hipGraphAddHostNode[('hipGraphNode_t*', 'pGraphNode'), ('hipGraph_t', 'graph'), ('const hipGraphNode_t*', 'pDependencies'), ('size_t', 'numDependencies'), ('const hipHostNodeParams*', 'pNodeParams')]
#define INIT_hipGraphAddHostNode_CB_ARGS_DATA(cb_data) { \
@@ -4202,9 +4263,18 @@ typedef struct hip_api_data_s {
};
// hipGraphExecExternalSemaphoresSignalNodeSetParams[('hipGraphExec_t', 'hGraphExec'), ('hipGraphNode_t', 'hNode'), ('const hipExternalSemaphoreSignalNodeParams*', 'nodeParams')]
#define INIT_hipGraphExecExternalSemaphoresSignalNodeSetParams_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipGraphExecExternalSemaphoresSignalNodeSetParams.hGraphExec = (hipGraphExec_t)hGraphExec; \
cb_data.args.hipGraphExecExternalSemaphoresSignalNodeSetParams.hNode = (hipGraphNode_t)hNode; \
cb_data.args.hipGraphExecExternalSemaphoresSignalNodeSetParams.nodeParams = (const hipExternalSemaphoreSignalNodeParams*)nodeParams; \
};
// hipGraphExecExternalSemaphoresWaitNodeSetParams[('hipGraphExec_t', 'hGraphExec'), ('hipGraphNode_t', 'hNode'), ('const hipExternalSemaphoreWaitNodeParams*', 'nodeParams')]
#define INIT_hipGraphExecExternalSemaphoresWaitNodeSetParams_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipGraphExecExternalSemaphoresWaitNodeSetParams.hGraphExec = (hipGraphExec_t)hGraphExec; \
cb_data.args.hipGraphExecExternalSemaphoresWaitNodeSetParams.hNode = (hipGraphNode_t)hNode; \
cb_data.args.hipGraphExecExternalSemaphoresWaitNodeSetParams.nodeParams = (const hipExternalSemaphoreWaitNodeParams*)nodeParams; \
};
// hipGraphExecGetFlags[('hipGraphExec_t', 'graphExec'), ('unsigned long long*', 'flags')]
#define INIT_hipGraphExecGetFlags_CB_ARGS_DATA(cb_data) { \
};
// hipGraphExecHostNodeSetParams[('hipGraphExec_t', 'hGraphExec'), ('hipGraphNode_t', 'node'), ('const hipHostNodeParams*', 'pNodeParams')]
#define INIT_hipGraphExecHostNodeSetParams_CB_ARGS_DATA(cb_data) { \
@@ -4259,6 +4329,9 @@ typedef struct hip_api_data_s {
cb_data.args.hipGraphExecMemsetNodeSetParams.node = (hipGraphNode_t)node; \
cb_data.args.hipGraphExecMemsetNodeSetParams.pNodeParams = (const hipMemsetParams*)pNodeParams; \
};
// hipGraphExecNodeSetParams[('hipGraphExec_t', 'graphExec'), ('hipGraphNode_t', 'node'), ('hipGraphNodeParams*', 'nodeParams')]
#define INIT_hipGraphExecNodeSetParams_CB_ARGS_DATA(cb_data) { \
};
// hipGraphExecUpdate[('hipGraphExec_t', 'hGraphExec'), ('hipGraph_t', 'hGraph'), ('hipGraphNode_t*', 'hErrorNode_out'), ('hipGraphExecUpdateResult*', 'updateResult_out')]
#define INIT_hipGraphExecUpdate_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipGraphExecUpdate.hGraphExec = (hipGraphExec_t)hGraphExec; \
@@ -4268,15 +4341,23 @@ typedef struct hip_api_data_s {
};
// hipGraphExternalSemaphoresSignalNodeGetParams[('hipGraphNode_t', 'hNode'), ('hipExternalSemaphoreSignalNodeParams*', 'params_out')]
#define INIT_hipGraphExternalSemaphoresSignalNodeGetParams_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipGraphExternalSemaphoresSignalNodeGetParams.hNode = (hipGraphNode_t)hNode; \
cb_data.args.hipGraphExternalSemaphoresSignalNodeGetParams.params_out = (hipExternalSemaphoreSignalNodeParams*)params_out; \
};
// hipGraphExternalSemaphoresSignalNodeSetParams[('hipGraphNode_t', 'hNode'), ('const hipExternalSemaphoreSignalNodeParams*', 'nodeParams')]
#define INIT_hipGraphExternalSemaphoresSignalNodeSetParams_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipGraphExternalSemaphoresSignalNodeSetParams.hNode = (hipGraphNode_t)hNode; \
cb_data.args.hipGraphExternalSemaphoresSignalNodeSetParams.nodeParams = (const hipExternalSemaphoreSignalNodeParams*)nodeParams; \
};
// hipGraphExternalSemaphoresWaitNodeGetParams[('hipGraphNode_t', 'hNode'), ('hipExternalSemaphoreWaitNodeParams*', 'params_out')]
#define INIT_hipGraphExternalSemaphoresWaitNodeGetParams_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipGraphExternalSemaphoresWaitNodeGetParams.hNode = (hipGraphNode_t)hNode; \
cb_data.args.hipGraphExternalSemaphoresWaitNodeGetParams.params_out = (hipExternalSemaphoreWaitNodeParams*)params_out; \
};
// hipGraphExternalSemaphoresWaitNodeSetParams[('hipGraphNode_t', 'hNode'), ('const hipExternalSemaphoreWaitNodeParams*', 'nodeParams')]
#define INIT_hipGraphExternalSemaphoresWaitNodeSetParams_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipGraphExternalSemaphoresWaitNodeSetParams.hNode = (hipGraphNode_t)hNode; \
cb_data.args.hipGraphExternalSemaphoresWaitNodeSetParams.nodeParams = (const hipExternalSemaphoreWaitNodeParams*)nodeParams; \
};
// hipGraphGetEdges[('hipGraph_t', 'graph'), ('hipGraphNode_t*', 'from'), ('hipGraphNode_t*', 'to'), ('size_t*', 'numEdges')]
#define INIT_hipGraphGetEdges_CB_ARGS_DATA(cb_data) { \
@@ -4321,6 +4402,12 @@ typedef struct hip_api_data_s {
cb_data.args.hipGraphInstantiateWithFlags.graph = (hipGraph_t)graph; \
cb_data.args.hipGraphInstantiateWithFlags.flags = (unsigned long long)flags; \
};
// hipGraphInstantiateWithParams[('hipGraphExec_t*', 'pGraphExec'), ('hipGraph_t', 'graph'), ('hipGraphInstantiateParams*', 'instantiateParams')]
#define INIT_hipGraphInstantiateWithParams_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipGraphInstantiateWithParams.pGraphExec = (hipGraphExec_t*)pGraphExec; \
cb_data.args.hipGraphInstantiateWithParams.graph = (hipGraph_t)graph; \
cb_data.args.hipGraphInstantiateWithParams.instantiateParams = (hipGraphInstantiateParams*)instantiateParams; \
};
// hipGraphKernelNodeCopyAttributes[('hipGraphNode_t', 'hSrc'), ('hipGraphNode_t', 'hDst')]
#define INIT_hipGraphKernelNodeCopyAttributes_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipGraphKernelNodeCopyAttributes.hSrc = (hipGraphNode_t)hSrc; \
@@ -4444,6 +4531,9 @@ typedef struct hip_api_data_s {
cb_data.args.hipGraphNodeSetEnabled.hNode = (hipGraphNode_t)hNode; \
cb_data.args.hipGraphNodeSetEnabled.isEnabled = (unsigned int)isEnabled; \
};
// hipGraphNodeSetParams[('hipGraphNode_t', 'node'), ('hipGraphNodeParams*', 'nodeParams')]
#define INIT_hipGraphNodeSetParams_CB_ARGS_DATA(cb_data) { \
};
// hipGraphReleaseUserObject[('hipGraph_t', 'graph'), ('hipUserObject_t', 'object'), ('unsigned int', 'count')]
#define INIT_hipGraphReleaseUserObject_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipGraphReleaseUserObject.graph = (hipGraph_t)graph; \
@@ -4630,12 +4720,6 @@ typedef struct hip_api_data_s {
};
// hipLaunchKernel[('const void*', 'function_address'), ('dim3', 'numBlocks'), ('dim3', 'dimBlocks'), ('void**', 'args'), ('size_t', 'sharedMemBytes'), ('hipStream_t', 'stream')]
#define INIT_hipLaunchKernel_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipLaunchKernel.function_address = (const void*)hostFunction; \
cb_data.args.hipLaunchKernel.numBlocks = (dim3)gridDim; \
cb_data.args.hipLaunchKernel.dimBlocks = (dim3)blockDim; \
cb_data.args.hipLaunchKernel.args = (void**)args; \
cb_data.args.hipLaunchKernel.sharedMemBytes = (size_t)sharedMemBytes; \
cb_data.args.hipLaunchKernel.stream = (hipStream_t)stream; \
};
// hipMalloc[('void**', 'ptr'), ('size_t', 'size')]
#define INIT_hipMalloc_CB_ARGS_DATA(cb_data) { \
@@ -5569,6 +5653,8 @@ typedef struct hip_api_data_s {
};
// hipTexRefGetMipMappedArray[('hipMipmappedArray_t*', 'pArray'), ('const textureReference*', 'texRef')]
#define INIT_hipTexRefGetMipMappedArray_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipTexRefGetMipMappedArray.pArray = (hipMipmappedArray_t*)pArray; \
cb_data.args.hipTexRefGetMipMappedArray.texRef = (const textureReference*)texRef; \
};
// hipTexRefGetMipmapLevelBias[('float*', 'pbias'), ('const textureReference*', 'texRef')]
#define INIT_hipTexRefGetMipmapLevelBias_CB_ARGS_DATA(cb_data) { \
@@ -5668,9 +5754,6 @@ typedef struct hip_api_data_s {
cb_data.args.hipWaitExternalSemaphoresAsync.numExtSems = (unsigned int)numExtSems; \
cb_data.args.hipWaitExternalSemaphoresAsync.stream = (hipStream_t)stream; \
};
// hipExtGetLastError[]
#define INIT_hipExtGetLastError_CB_ARGS_DATA(cb_data) { \
};
#define INIT_CB_ARGS_DATA(cb_id, cb_data) INIT_##cb_id##_CB_ARGS_DATA(cb_data)
// Macros for non-public API primitives
@@ -6050,6 +6133,9 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
// hipEventSynchronize[('hipEvent_t', 'event')]
case HIP_API_ID_hipEventSynchronize:
break;
// hipExtGetLastError[]
case HIP_API_ID_hipExtGetLastError:
break;
// hipExtGetLinkTypeAndHopCount[('int', 'device1'), ('int', 'device2'), ('unsigned int*', 'linktype'), ('unsigned int*', 'hopcount')]
case HIP_API_ID_hipExtGetLinkTypeAndHopCount:
if (data->args.hipExtGetLinkTypeAndHopCount.linktype) data->args.hipExtGetLinkTypeAndHopCount.linktype__val = *(data->args.hipExtGetLinkTypeAndHopCount.linktype);
@@ -6158,13 +6244,16 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
// hipGetLastError[]
case HIP_API_ID_hipGetLastError:
break;
// hipExtGetLastError[]
case HIP_API_ID_hipExtGetLastError:
break;
// hipGetMipmappedArrayLevel[('hipArray_t*', 'levelArray'), ('hipMipmappedArray_const_t', 'mipmappedArray'), ('unsigned int', 'level')]
case HIP_API_ID_hipGetMipmappedArrayLevel:
if (data->args.hipGetMipmappedArrayLevel.levelArray) data->args.hipGetMipmappedArrayLevel.levelArray__val = *(data->args.hipGetMipmappedArrayLevel.levelArray);
break;
// hipGetProcAddress[('const char*', 'symbol'), ('void**', 'pfn'), ('int', 'hipVersion'), ('uint64_t', 'flags'), ('hipDriverProcAddressQueryResult*', 'symbolStatus')]
case HIP_API_ID_hipGetProcAddress:
if (data->args.hipGetProcAddress.symbol) data->args.hipGetProcAddress.symbol__val = *(data->args.hipGetProcAddress.symbol);
if (data->args.hipGetProcAddress.pfn) data->args.hipGetProcAddress.pfn__val = *(data->args.hipGetProcAddress.pfn);
if (data->args.hipGetProcAddress.symbolStatus) data->args.hipGetProcAddress.symbolStatus__val = *(data->args.hipGetProcAddress.symbolStatus);
break;
// hipGetSymbolAddress[('void**', 'devPtr'), ('const void*', 'symbol')]
case HIP_API_ID_hipGetSymbolAddress:
if (data->args.hipGetSymbolAddress.devPtr) data->args.hipGetSymbolAddress.devPtr__val = *(data->args.hipGetSymbolAddress.devPtr);
@@ -6322,6 +6411,10 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
case HIP_API_ID_hipGraphExecExternalSemaphoresWaitNodeSetParams:
if (data->args.hipGraphExecExternalSemaphoresWaitNodeSetParams.nodeParams) data->args.hipGraphExecExternalSemaphoresWaitNodeSetParams.nodeParams__val = *(data->args.hipGraphExecExternalSemaphoresWaitNodeSetParams.nodeParams);
break;
// hipGraphExecGetFlags[('hipGraphExec_t', 'graphExec'), ('unsigned long long*', 'flags')]
case HIP_API_ID_hipGraphExecGetFlags:
if (data->args.hipGraphExecGetFlags.flags) data->args.hipGraphExecGetFlags.flags__val = *(data->args.hipGraphExecGetFlags.flags);
break;
// hipGraphExecHostNodeSetParams[('hipGraphExec_t', 'hGraphExec'), ('hipGraphNode_t', 'node'), ('const hipHostNodeParams*', 'pNodeParams')]
case HIP_API_ID_hipGraphExecHostNodeSetParams:
if (data->args.hipGraphExecHostNodeSetParams.pNodeParams) data->args.hipGraphExecHostNodeSetParams.pNodeParams__val = *(data->args.hipGraphExecHostNodeSetParams.pNodeParams);
@@ -6347,6 +6440,10 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
case HIP_API_ID_hipGraphExecMemsetNodeSetParams:
if (data->args.hipGraphExecMemsetNodeSetParams.pNodeParams) data->args.hipGraphExecMemsetNodeSetParams.pNodeParams__val = *(data->args.hipGraphExecMemsetNodeSetParams.pNodeParams);
break;
// hipGraphExecNodeSetParams[('hipGraphExec_t', 'graphExec'), ('hipGraphNode_t', 'node'), ('hipGraphNodeParams*', 'nodeParams')]
case HIP_API_ID_hipGraphExecNodeSetParams:
if (data->args.hipGraphExecNodeSetParams.nodeParams) data->args.hipGraphExecNodeSetParams.nodeParams__val = *(data->args.hipGraphExecNodeSetParams.nodeParams);
break;
// hipGraphExecUpdate[('hipGraphExec_t', 'hGraphExec'), ('hipGraph_t', 'hGraph'), ('hipGraphNode_t*', 'hErrorNode_out'), ('hipGraphExecUpdateResult*', 'updateResult_out')]
case HIP_API_ID_hipGraphExecUpdate:
if (data->args.hipGraphExecUpdate.hErrorNode_out) data->args.hipGraphExecUpdate.hErrorNode_out__val = *(data->args.hipGraphExecUpdate.hErrorNode_out);
@@ -6402,6 +6499,11 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
case HIP_API_ID_hipGraphInstantiateWithFlags:
if (data->args.hipGraphInstantiateWithFlags.pGraphExec) data->args.hipGraphInstantiateWithFlags.pGraphExec__val = *(data->args.hipGraphInstantiateWithFlags.pGraphExec);
break;
// hipGraphInstantiateWithParams[('hipGraphExec_t*', 'pGraphExec'), ('hipGraph_t', 'graph'), ('hipGraphInstantiateParams*', 'instantiateParams')]
case HIP_API_ID_hipGraphInstantiateWithParams:
if (data->args.hipGraphInstantiateWithParams.pGraphExec) data->args.hipGraphInstantiateWithParams.pGraphExec__val = *(data->args.hipGraphInstantiateWithParams.pGraphExec);
if (data->args.hipGraphInstantiateWithParams.instantiateParams) data->args.hipGraphInstantiateWithParams.instantiateParams__val = *(data->args.hipGraphInstantiateWithParams.instantiateParams);
break;
// hipGraphKernelNodeCopyAttributes[('hipGraphNode_t', 'hSrc'), ('hipGraphNode_t', 'hDst')]
case HIP_API_ID_hipGraphKernelNodeCopyAttributes:
break;
@@ -6481,6 +6583,10 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
// hipGraphNodeSetEnabled[('hipGraphExec_t', 'hGraphExec'), ('hipGraphNode_t', 'hNode'), ('unsigned int', 'isEnabled')]
case HIP_API_ID_hipGraphNodeSetEnabled:
break;
// hipGraphNodeSetParams[('hipGraphNode_t', 'node'), ('hipGraphNodeParams*', 'nodeParams')]
case HIP_API_ID_hipGraphNodeSetParams:
if (data->args.hipGraphNodeSetParams.nodeParams) data->args.hipGraphNodeSetParams.nodeParams__val = *(data->args.hipGraphNodeSetParams.nodeParams);
break;
// hipGraphReleaseUserObject[('hipGraph_t', 'graph'), ('hipUserObject_t', 'object'), ('unsigned int', 'count')]
case HIP_API_ID_hipGraphReleaseUserObject:
break;
@@ -7760,6 +7866,10 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
oss << "event="; roctracer::hip_support::detail::operator<<(oss, data->args.hipEventSynchronize.event);
oss << ")";
break;
case HIP_API_ID_hipExtGetLastError:
oss << "hipExtGetLastError(";
oss << ")";
break;
case HIP_API_ID_hipExtGetLinkTypeAndHopCount:
oss << "hipExtGetLinkTypeAndHopCount(";
oss << "device1="; roctracer::hip_support::detail::operator<<(oss, data->args.hipExtGetLinkTypeAndHopCount.device1);
@@ -7972,10 +8082,6 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
oss << "hipGetLastError(";
oss << ")";
break;
case HIP_API_ID_hipExtGetLastError:
oss << "hipExtGetLastError(";
oss << ")";
break;
case HIP_API_ID_hipGetMipmappedArrayLevel:
oss << "hipGetMipmappedArrayLevel(";
if (data->args.hipGetMipmappedArrayLevel.levelArray == NULL) oss << "levelArray=NULL";
@@ -7984,6 +8090,18 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
oss << ", level="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetMipmappedArrayLevel.level);
oss << ")";
break;
case HIP_API_ID_hipGetProcAddress:
oss << "hipGetProcAddress(";
if (data->args.hipGetProcAddress.symbol == NULL) oss << "symbol=NULL";
else { oss << "symbol="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetProcAddress.symbol__val); }
if (data->args.hipGetProcAddress.pfn == NULL) oss << ", pfn=NULL";
else { oss << ", pfn="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetProcAddress.pfn__val); }
oss << ", hipVersion="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetProcAddress.hipVersion);
oss << ", flags="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetProcAddress.flags);
if (data->args.hipGetProcAddress.symbolStatus == NULL) oss << ", symbolStatus=NULL";
else { oss << ", symbolStatus="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetProcAddress.symbolStatus__val); }
oss << ")";
break;
case HIP_API_ID_hipGetSymbolAddress:
oss << "hipGetSymbolAddress(";
if (data->args.hipGetSymbolAddress.devPtr == NULL) oss << "devPtr=NULL";
@@ -8309,6 +8427,13 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
else { oss << ", nodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecExternalSemaphoresWaitNodeSetParams.nodeParams__val); }
oss << ")";
break;
case HIP_API_ID_hipGraphExecGetFlags:
oss << "hipGraphExecGetFlags(";
oss << "graphExec="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecGetFlags.graphExec);
if (data->args.hipGraphExecGetFlags.flags == NULL) oss << ", flags=NULL";
else { oss << ", flags="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecGetFlags.flags__val); }
oss << ")";
break;
case HIP_API_ID_hipGraphExecHostNodeSetParams:
oss << "hipGraphExecHostNodeSetParams(";
oss << "hGraphExec="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecHostNodeSetParams.hGraphExec);
@@ -8373,6 +8498,14 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
else { oss << ", pNodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecMemsetNodeSetParams.pNodeParams__val); }
oss << ")";
break;
case HIP_API_ID_hipGraphExecNodeSetParams:
oss << "hipGraphExecNodeSetParams(";
oss << "graphExec="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecNodeSetParams.graphExec);
oss << ", node="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecNodeSetParams.node);
if (data->args.hipGraphExecNodeSetParams.nodeParams == NULL) oss << ", nodeParams=NULL";
else { oss << ", nodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecNodeSetParams.nodeParams__val); }
oss << ")";
break;
case HIP_API_ID_hipGraphExecUpdate:
oss << "hipGraphExecUpdate(";
oss << "hGraphExec="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecUpdate.hGraphExec);
@@ -8474,6 +8607,15 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
oss << ", flags="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphInstantiateWithFlags.flags);
oss << ")";
break;
case HIP_API_ID_hipGraphInstantiateWithParams:
oss << "hipGraphInstantiateWithParams(";
if (data->args.hipGraphInstantiateWithParams.pGraphExec == NULL) oss << "pGraphExec=NULL";
else { oss << "pGraphExec="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphInstantiateWithParams.pGraphExec__val); }
oss << ", graph="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphInstantiateWithParams.graph);
if (data->args.hipGraphInstantiateWithParams.instantiateParams == NULL) oss << ", instantiateParams=NULL";
else { oss << ", instantiateParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphInstantiateWithParams.instantiateParams__val); }
oss << ")";
break;
case HIP_API_ID_hipGraphKernelNodeCopyAttributes:
oss << "hipGraphKernelNodeCopyAttributes(";
oss << "hSrc="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphKernelNodeCopyAttributes.hSrc);
@@ -8634,6 +8776,13 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
oss << ", isEnabled="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphNodeSetEnabled.isEnabled);
oss << ")";
break;
case HIP_API_ID_hipGraphNodeSetParams:
oss << "hipGraphNodeSetParams(";
oss << "node="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphNodeSetParams.node);
if (data->args.hipGraphNodeSetParams.nodeParams == NULL) oss << ", nodeParams=NULL";
else { oss << ", nodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphNodeSetParams.nodeParams__val); }
oss << ")";
break;
case HIP_API_ID_hipGraphReleaseUserObject:
oss << "hipGraphReleaseUserObject(";
oss << "graph="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphReleaseUserObject.graph);
projects/clr/hipamd/src/amdhip.def
+1
Datei anzeigen
@@ -459,3 +459,4 @@ hipGraphExternalSemaphoresWaitNodeGetParams
hipGraphExecExternalSemaphoresSignalNodeSetParams
hipGraphExecExternalSemaphoresWaitNodeSetParams
hipGraphAddNode
hipGraphInstantiateWithParams
projects/clr/hipamd/src/hip_api_trace.cpp
+14 -1
Datei anzeigen
@@ -290,6 +290,8 @@ hipError_t hipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph,
hipGraphNode_t* pErrorNode, char* pLogBuffer, size_t bufferSize);
hipError_t hipGraphInstantiateWithFlags(hipGraphExec_t* pGraphExec, hipGraph_t graph,
unsigned long long flags);
hipError_t hipGraphInstantiateWithParams(hipGraphExec_t* pGraphExec, hipGraph_t graph,
hipGraphInstantiateParams* instantiateParams);
hipError_t hipGraphKernelNodeCopyAttributes(hipGraphNode_t hSrc, hipGraphNode_t hDst);
hipError_t hipGraphKernelNodeGetAttribute(hipGraphNode_t hNode, hipKernelNodeAttrID attr,
hipKernelNodeAttrValue* value);
@@ -942,6 +944,7 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) {
ptrDispatchTable->hipGraphHostNodeSetParams_fn = hip::hipGraphHostNodeSetParams;
ptrDispatchTable->hipGraphInstantiate_fn = hip::hipGraphInstantiate;
ptrDispatchTable->hipGraphInstantiateWithFlags_fn = hip::hipGraphInstantiateWithFlags;
ptrDispatchTable->hipGraphInstantiateWithParams_fn = hip::hipGraphInstantiateWithParams;
ptrDispatchTable->hipGraphKernelNodeCopyAttributes_fn = hip::hipGraphKernelNodeCopyAttributes;
ptrDispatchTable->hipGraphKernelNodeGetAttribute_fn = hip::hipGraphKernelNodeGetAttribute;
ptrDispatchTable->hipGraphKernelNodeGetParams_fn = hip::hipGraphKernelNodeGetParams;
@@ -1310,7 +1313,7 @@ const HipCompilerDispatchTable* GetHipCompilerDispatchTable() {
#define HIP_ENFORCE_ABI(TABLE, ENTRY, NUM) \
static_assert(offsetof(TABLE, ENTRY) == (sizeof(size_t) + (NUM * sizeof(void*))), \
"ABI break for " #TABLE "." #ENTRY \
". Only add new function pointers to end of struct and do not rearrange them");
". Only add new function pointers to end of struct and do not rearrange them " );
// These ensure that function pointers are not re-ordered
HIP_ENFORCE_ABI(HipCompilerDispatchTable, __hipPopCallConfiguration_fn, 0)
@@ -1758,6 +1761,16 @@ 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);
static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 0,
"If you get this error, add new HIP_ENFORCE_ABI(...) code for the new function "
projects/clr/hipamd/src/hip_graph.cpp
+46 -8
Datei anzeigen
@@ -34,6 +34,18 @@ amd::Monitor g_streamSetLock{"StreamCaptureset"};
std::unordered_set<hip::Stream*> g_allCapturingStreams;
hipError_t ihipGraphDebugDotPrint(hipGraph_t graph, const char* path, unsigned int flags);
inline hipError_t ihipGraphUpload(hipGraphExec_t graphExec, hipStream_t stream) {
if (graphExec == nullptr) {
return hipErrorInvalidValue;
}
if (!hip::isValid(stream)) {
return hipErrorContextIsDestroyed;
}
return hipSuccess;
}
inline hipError_t ihipGraphAddNode(hip::GraphNode* graphNode, hip::Graph* graph,
hip::GraphNode* const* pDependencies, size_t numDependencies,
bool capture = true) {
@@ -1277,6 +1289,37 @@ hipError_t hipGraphInstantiateWithFlags(hipGraphExec_t* pGraphExec, hipGraph_t g
HIP_RETURN(status);
}
hipError_t hipGraphInstantiateWithParams(hipGraphExec_t* pGraphExec, hipGraph_t graph,
hipGraphInstantiateParams* instantiateParams) {
HIP_INIT_API(hipGraphInstantiateWithParams, pGraphExec, graph, instantiateParams);
if (pGraphExec == nullptr || graph == nullptr || instantiateParams == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
unsigned long long flags = instantiateParams->flags;
if (flags != 0 && flags != hipGraphInstantiateFlagAutoFreeOnLaunch &&
flags != hipGraphInstantiateFlagUpload && flags != hipGraphInstantiateFlagDeviceLaunch &&
flags != hipGraphInstantiateFlagUseNodePriority) {
HIP_RETURN(hipErrorInvalidValue);
}
hip::GraphExec* ge;
hipError_t status = ihipGraphInstantiate(&ge, reinterpret_cast<hip::Graph*>(graph), flags);
*pGraphExec = reinterpret_cast<hipGraphExec_t>(ge);
if(status != hipSuccess){
HIP_RETURN(status);
}
if (flags == hipGraphInstantiateFlagUpload) {
hipError_t status = ihipGraphUpload(*pGraphExec, instantiateParams->uploadStream);
HIP_RETURN(status);
}
HIP_RETURN(hipSuccess);
}
hipError_t hipGraphExecDestroy(hipGraphExec_t pGraphExec) {
HIP_INIT_API(hipGraphExecDestroy, pGraphExec);
if (pGraphExec == nullptr) {
@@ -2635,15 +2678,10 @@ hipError_t hipGraphNodeGetEnabled(hipGraphExec_t hGraphExec, hipGraphNode_t hNod
hipError_t hipGraphUpload(hipGraphExec_t graphExec, hipStream_t stream) {
HIP_INIT_API(hipGraphUpload, graphExec, stream);
if (graphExec == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
if (!hip::isValid(stream)) {
return hipErrorContextIsDestroyed;
}
// TODO: stream is known before launch, do preperatory work with graph optimizations. pre-allocate
// memory for memAlloc nodes if any when support is added with mempool feature
HIP_RETURN(hipSuccess);
hipError_t status = ihipGraphUpload(graphExec, stream);
HIP_RETURN(status);
}
hipError_t hipGraphAddNode(hipGraphNode_t *pGraphNode, hipGraph_t graph,
@@ -2789,7 +2827,7 @@ hipError_t hipGraphAddExternalSemaphoresSignalNode(hipGraphNode_t* pGraphNode, h
const hipGraphNode_t* pDependencies, size_t numDependencies,
const hipExternalSemaphoreSignalNodeParams* nodeParams) {
HIP_INIT_API(hipGraphAddExternalSemaphoresSignalNode, pGraphNode, graph, pDependencies,
numDependencies, nodeParams);
numDependencies, nodeParams);
hip::GraphNode* node = new hip::hipGraphExternalSemSignalNode(nodeParams);
hipError_t status = ihipGraphAddNode(node, reinterpret_cast<hip::Graph*>(graph),
reinterpret_cast<hip::GraphNode* const*>(pDependencies), numDependencies);
projects/clr/hipamd/src/hip_hcc.map.in
+7
Datei anzeigen
@@ -547,3 +547,10 @@ global:
local:
*;
} hip_5.6;
hip_6.1 {
global:
hipGraphInstantiateWithParams;
local:
*;
} hip_6.0;
projects/clr/hipamd/src/hip_table_interface.cpp
+5
Datei anzeigen
@@ -661,6 +661,11 @@ hipError_t hipGraphInstantiateWithFlags(hipGraphExec_t* pGraphExec, hipGraph_t g
unsigned long long flags) {
return hip::GetHipDispatchTable()->hipGraphInstantiateWithFlags_fn(pGraphExec, graph, flags);
}
hipError_t hipGraphInstantiateWithParams(hipGraphExec_t* pGraphExec, hipGraph_t graph,
hipGraphInstantiateParams* instantiateParams) {
return hip::GetHipDispatchTable()->hipGraphInstantiateWithParams_fn(pGraphExec, graph,
instantiateParams);
}
hipError_t hipGraphKernelNodeCopyAttributes(hipGraphNode_t hSrc, hipGraphNode_t hDst) {
return hip::GetHipDispatchTable()->hipGraphKernelNodeCopyAttributes_fn(hSrc, hDst);
}