diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp b/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp index b0d5b83658..53f0dc748a 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp @@ -63,7 +63,7 @@ #define HIP_API_TABLE_STEP_VERSION 0 #define HIP_COMPILER_API_TABLE_STEP_VERSION 0 #define HIP_TOOLS_API_TABLE_STEP_VERSION 0 -#define HIP_RUNTIME_API_TABLE_STEP_VERSION 8 +#define HIP_RUNTIME_API_TABLE_STEP_VERSION 9 // HIP API interface // HIP compiler dispatch functions @@ -1025,6 +1025,8 @@ typedef hipError_t (*t_hipGraphBatchMemOpNodeSetParams)(hipGraphNode_t hNode, hipBatchMemOpNodeParams* nodeParams); typedef hipError_t (*t_hipGraphExecBatchMemOpNodeSetParams)( hipGraphExec_t hGraphExec, hipGraphNode_t hNode, const hipBatchMemOpNodeParams* nodeParams); +typedef hipError_t (*t_hipEventRecordWithFlags)(hipEvent_t event, hipStream_t stream, unsigned int flags); + // HIP Compiler dispatch table struct HipCompilerDispatchTable { // HIP_COMPILER_API_TABLE_STEP_VERSION == 0 @@ -1549,8 +1551,11 @@ struct HipDispatchTable { t_hipGraphBatchMemOpNodeSetParams hipGraphBatchMemOpNodeSetParams_fn; t_hipGraphExecBatchMemOpNodeSetParams hipGraphExecBatchMemOpNodeSetParams_fn; + // HIP_RUNTIME_API_TABLE_STEP_VERSION == 9 + t_hipEventRecordWithFlags hipEventRecordWithFlags_fn; + // DO NOT EDIT ABOVE! - // HIP_RUNTIME_API_TABLE_STEP_VERSION == 7 + // HIP_RUNTIME_API_TABLE_STEP_VERSION == 10 // ******************************************************************************************* // // diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h index f25c782dd9..9f083b7482 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h @@ -424,13 +424,14 @@ enum hip_api_id_t { HIP_API_ID_hipMemcpyDtoA = 404, HIP_API_ID_hipMemcpyHtoAAsync = 405, HIP_API_ID_hipSetValidDevices = 406, - HIP_API_ID_hipExtHostAlloc = 407, + HIP_API_ID_RESERVED_407 = 407, HIP_API_ID_hipStreamBatchMemOp = 408, HIP_API_ID_hipGraphAddBatchMemOpNode = 409, HIP_API_ID_hipGraphBatchMemOpNodeGetParams = 410, HIP_API_ID_hipGraphBatchMemOpNodeSetParams = 411, HIP_API_ID_hipGraphExecBatchMemOpNodeSetParams = 412, - HIP_API_ID_LAST = 412, + HIP_API_ID_hipEventRecordWithFlags = 413, + HIP_API_ID_LAST = 413, HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice), HIP_API_ID_hipGetDeviceProperties = HIP_API_ID_CONCAT(HIP_API_ID_,hipGetDeviceProperties), @@ -549,7 +550,9 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipEventElapsedTime: return "hipEventElapsedTime"; case HIP_API_ID_hipEventQuery: return "hipEventQuery"; case HIP_API_ID_hipEventRecord: return "hipEventRecord"; + case HIP_API_ID_hipEventRecordWithFlags: return "hipEventRecordWithFlags"; 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"; @@ -583,6 +586,7 @@ static inline const char* hip_api_name(const uint32_t id) { 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_hipGraphAddBatchMemOpNode: return "hipGraphAddBatchMemOpNode"; case HIP_API_ID_hipGraphAddChildGraphNode: return "hipGraphAddChildGraphNode"; case HIP_API_ID_hipGraphAddDependencies: return "hipGraphAddDependencies"; case HIP_API_ID_hipGraphAddEmptyNode: return "hipGraphAddEmptyNode"; @@ -600,6 +604,8 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipGraphAddMemcpyNodeToSymbol: return "hipGraphAddMemcpyNodeToSymbol"; case HIP_API_ID_hipGraphAddMemsetNode: return "hipGraphAddMemsetNode"; case HIP_API_ID_hipGraphAddNode: return "hipGraphAddNode"; + case HIP_API_ID_hipGraphBatchMemOpNodeGetParams: return "hipGraphBatchMemOpNodeGetParams"; + case HIP_API_ID_hipGraphBatchMemOpNodeSetParams: return "hipGraphBatchMemOpNodeSetParams"; case HIP_API_ID_hipGraphChildGraphNodeGetGraph: return "hipGraphChildGraphNodeGetGraph"; case HIP_API_ID_hipGraphClone: return "hipGraphClone"; case HIP_API_ID_hipGraphCreate: return "hipGraphCreate"; @@ -610,6 +616,7 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipGraphEventRecordNodeSetEvent: return "hipGraphEventRecordNodeSetEvent"; case HIP_API_ID_hipGraphEventWaitNodeGetEvent: return "hipGraphEventWaitNodeGetEvent"; case HIP_API_ID_hipGraphEventWaitNodeSetEvent: return "hipGraphEventWaitNodeSetEvent"; + case HIP_API_ID_hipGraphExecBatchMemOpNodeSetParams: return "hipGraphExecBatchMemOpNodeSetParams"; case HIP_API_ID_hipGraphExecChildGraphNodeSetParams: return "hipGraphExecChildGraphNodeSetParams"; case HIP_API_ID_hipGraphExecDestroy: return "hipGraphExecDestroy"; case HIP_API_ID_hipGraphExecEventRecordNodeSetEvent: return "hipGraphExecEventRecordNodeSetEvent"; @@ -817,6 +824,7 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipSignalExternalSemaphoresAsync: return "hipSignalExternalSemaphoresAsync"; case HIP_API_ID_hipStreamAddCallback: return "hipStreamAddCallback"; case HIP_API_ID_hipStreamAttachMemAsync: return "hipStreamAttachMemAsync"; + case HIP_API_ID_hipStreamBatchMemOp: return "hipStreamBatchMemOp"; case HIP_API_ID_hipStreamBeginCapture: return "hipStreamBeginCapture"; case HIP_API_ID_hipStreamBeginCaptureToGraph: return "hipStreamBeginCaptureToGraph"; case HIP_API_ID_hipStreamCreate: return "hipStreamCreate"; @@ -862,12 +870,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"; - case HIP_API_ID_hipStreamBatchMemOp: return "hipStreamBatchMemOp"; - case HIP_API_ID_hipGraphAddBatchMemOpNode: return "hipGraphAddBatchMemOpNode"; - case HIP_API_ID_hipGraphBatchMemOpNodeGetParams: return "hipGraphBatchMemOpNodeGetParams"; - case HIP_API_ID_hipGraphBatchMemOpNodeSetParams: return "hipGraphBatchMemOpNodeSetParams"; - case HIP_API_ID_hipGraphExecBatchMemOpNodeSetParams: return "hipGraphExecBatchMemOpNodeSetParams"; }; return "unknown"; }; @@ -956,6 +958,7 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipEventElapsedTime", name) == 0) return HIP_API_ID_hipEventElapsedTime; if (strcmp("hipEventQuery", name) == 0) return HIP_API_ID_hipEventQuery; if (strcmp("hipEventRecord", name) == 0) return HIP_API_ID_hipEventRecord; + if (strcmp("hipEventRecordWithFlags", name) == 0) return HIP_API_ID_hipEventRecordWithFlags; 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; @@ -991,6 +994,7 @@ static inline uint32_t hipApiIdByName(const char* name) { 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("hipGraphAddBatchMemOpNode", name) == 0) return HIP_API_ID_hipGraphAddBatchMemOpNode; if (strcmp("hipGraphAddChildGraphNode", name) == 0) return HIP_API_ID_hipGraphAddChildGraphNode; if (strcmp("hipGraphAddDependencies", name) == 0) return HIP_API_ID_hipGraphAddDependencies; if (strcmp("hipGraphAddEmptyNode", name) == 0) return HIP_API_ID_hipGraphAddEmptyNode; @@ -1008,6 +1012,8 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipGraphAddMemcpyNodeToSymbol", name) == 0) return HIP_API_ID_hipGraphAddMemcpyNodeToSymbol; if (strcmp("hipGraphAddMemsetNode", name) == 0) return HIP_API_ID_hipGraphAddMemsetNode; if (strcmp("hipGraphAddNode", name) == 0) return HIP_API_ID_hipGraphAddNode; + if (strcmp("hipGraphBatchMemOpNodeGetParams", name) == 0) return HIP_API_ID_hipGraphBatchMemOpNodeGetParams; + if (strcmp("hipGraphBatchMemOpNodeSetParams", name) == 0) return HIP_API_ID_hipGraphBatchMemOpNodeSetParams; if (strcmp("hipGraphChildGraphNodeGetGraph", name) == 0) return HIP_API_ID_hipGraphChildGraphNodeGetGraph; if (strcmp("hipGraphClone", name) == 0) return HIP_API_ID_hipGraphClone; if (strcmp("hipGraphCreate", name) == 0) return HIP_API_ID_hipGraphCreate; @@ -1018,6 +1024,7 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipGraphEventRecordNodeSetEvent", name) == 0) return HIP_API_ID_hipGraphEventRecordNodeSetEvent; if (strcmp("hipGraphEventWaitNodeGetEvent", name) == 0) return HIP_API_ID_hipGraphEventWaitNodeGetEvent; if (strcmp("hipGraphEventWaitNodeSetEvent", name) == 0) return HIP_API_ID_hipGraphEventWaitNodeSetEvent; + if (strcmp("hipGraphExecBatchMemOpNodeSetParams", name) == 0) return HIP_API_ID_hipGraphExecBatchMemOpNodeSetParams; if (strcmp("hipGraphExecChildGraphNodeSetParams", name) == 0) return HIP_API_ID_hipGraphExecChildGraphNodeSetParams; if (strcmp("hipGraphExecDestroy", name) == 0) return HIP_API_ID_hipGraphExecDestroy; if (strcmp("hipGraphExecEventRecordNodeSetEvent", name) == 0) return HIP_API_ID_hipGraphExecEventRecordNodeSetEvent; @@ -1225,6 +1232,7 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipSignalExternalSemaphoresAsync", name) == 0) return HIP_API_ID_hipSignalExternalSemaphoresAsync; if (strcmp("hipStreamAddCallback", name) == 0) return HIP_API_ID_hipStreamAddCallback; if (strcmp("hipStreamAttachMemAsync", name) == 0) return HIP_API_ID_hipStreamAttachMemAsync; + if (strcmp("hipStreamBatchMemOp", name) == 0) return HIP_API_ID_hipStreamBatchMemOp; 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; @@ -1270,11 +1278,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("hipStreamBatchMemOp", name) == 0) return HIP_API_ID_hipStreamBatchMemOp; - if (strcmp("hipGraphAddBatchMemOpNode", name) == 0) return HIP_API_ID_hipGraphAddBatchMemOpNode; - if (strcmp("hipGraphBatchMemOpNodeGetParams", name) == 0) return HIP_API_ID_hipGraphBatchMemOpNodeGetParams; - if (strcmp("hipGraphBatchMemOpNodeSetParams", name) == 0) return HIP_API_ID_hipGraphBatchMemOpNodeSetParams; - if (strcmp("hipGraphExecBatchMemOpNodeSetParams", name) == 0) return HIP_API_ID_hipGraphExecBatchMemOpNodeSetParams; return HIP_API_ID_NONE; } @@ -1672,6 +1675,11 @@ typedef struct hip_api_data_s { hipEvent_t event; hipStream_t stream; } hipEventRecord; + struct { + hipEvent_t event; + hipStream_t stream; + unsigned int flags; + } hipEventRecordWithFlags; struct { hipEvent_t event; } hipEventSynchronize; @@ -1858,6 +1866,16 @@ typedef struct hip_api_data_s { size_t size__val; const void* symbol; } hipGetSymbolSize; + struct { + hipGraphNode_t* phGraphNode; + hipGraphNode_t phGraphNode__val; + hipGraph_t hGraph; + const hipGraphNode_t* dependencies; + hipGraphNode_t dependencies__val; + size_t numDependencies; + const hipBatchMemOpNodeParams* nodeParams; + hipBatchMemOpNodeParams nodeParams__val; + } hipGraphAddBatchMemOpNode; struct { hipGraphNode_t* pGraphNode; hipGraphNode_t pGraphNode__val; @@ -2028,6 +2046,16 @@ typedef struct hip_api_data_s { hipGraphNodeParams* nodeParams; hipGraphNodeParams nodeParams__val; } hipGraphAddNode; + struct { + hipGraphNode_t hNode; + hipBatchMemOpNodeParams* nodeParams_out; + hipBatchMemOpNodeParams nodeParams_out__val; + } hipGraphBatchMemOpNodeGetParams; + struct { + hipGraphNode_t hNode; + hipBatchMemOpNodeParams* nodeParams; + hipBatchMemOpNodeParams nodeParams__val; + } hipGraphBatchMemOpNodeSetParams; struct { hipGraphNode_t node; hipGraph_t* pGraph; @@ -2073,6 +2101,12 @@ typedef struct hip_api_data_s { hipGraphNode_t node; hipEvent_t event; } hipGraphEventWaitNodeSetEvent; + struct { + hipGraphExec_t hGraphExec; + hipGraphNode_t hNode; + const hipBatchMemOpNodeParams* nodeParams; + hipBatchMemOpNodeParams nodeParams__val; + } hipGraphExecBatchMemOpNodeSetParams; struct { hipGraphExec_t hGraphExec; hipGraphNode_t node; @@ -3361,6 +3395,13 @@ typedef struct hip_api_data_s { size_t length; unsigned int flags; } hipStreamAttachMemAsync; + struct { + hipStream_t stream; + unsigned int count; + hipStreamBatchMemOpParams* paramArray; + hipStreamBatchMemOpParams paramArray__val; + unsigned int flags; + } hipStreamBatchMemOp; struct { hipStream_t stream; hipStreamCaptureMode mode; @@ -3630,39 +3671,6 @@ typedef struct hip_api_data_s { unsigned int numExtSems; hipStream_t stream; } hipWaitExternalSemaphoresAsync; - struct { - hipStream_t stream; - unsigned int count; - hipStreamBatchMemOpParams* paramArray; - hipStreamBatchMemOpParams paramArray__val; - unsigned int flags; - } hipStreamBatchMemOp; - struct { - hipGraphNode_t* phGraphNode; - hipGraphNode_t phGraphNode__val; - hipGraph_t hGraph; - const hipGraphNode_t* dependencies; - hipGraphNode_t dependencies__val; - size_t numDependencies; - const hipBatchMemOpNodeParams* nodeParams; - hipBatchMemOpNodeParams nodeParams__val; - } hipGraphAddBatchMemOpNode; - struct { - hipGraphNode_t hNode; - hipBatchMemOpNodeParams* nodeParams_out; - hipBatchMemOpNodeParams nodeParams_out__val; - } hipGraphBatchMemOpNodeGetParams; - struct { - hipGraphNode_t hNode; - hipBatchMemOpNodeParams* nodeParams; - hipBatchMemOpNodeParams nodeParams__val; - } hipGraphBatchMemOpNodeSetParams; - struct { - hipGraphExec_t hGraphExec; - hipGraphNode_t hNode; - const hipBatchMemOpNodeParams* nodeParams; - hipBatchMemOpNodeParams nodeParams__val; - } hipGraphExecBatchMemOpNodeSetParams; } args; uint64_t *phase_data; } hip_api_data_t; @@ -4074,6 +4082,12 @@ typedef struct hip_api_data_s { cb_data.args.hipEventRecord.event = (hipEvent_t)event; \ cb_data.args.hipEventRecord.stream = (hipStream_t)stream; \ }; +// hipEventRecordWithFlags[('hipEvent_t', 'event'), ('hipStream_t', 'stream'), ('unsigned int', 'flags')] +#define INIT_hipEventRecordWithFlags_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipEventRecordWithFlags.event = (hipEvent_t)event; \ + cb_data.args.hipEventRecordWithFlags.stream = (hipStream_t)stream; \ + cb_data.args.hipEventRecordWithFlags.flags = (unsigned int)flags; \ +}; // hipEventSynchronize[('hipEvent_t', 'event')] #define INIT_hipEventSynchronize_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipEventSynchronize.event = (hipEvent_t)event; \ @@ -4270,6 +4284,14 @@ typedef struct hip_api_data_s { cb_data.args.hipGetSymbolSize.size = (size_t*)sizePtr; \ cb_data.args.hipGetSymbolSize.symbol = (const void*)symbol; \ }; +// hipGraphAddBatchMemOpNode[('hipGraphNode_t*', 'phGraphNode'), ('hipGraph_t', 'hGraph'), ('const hipGraphNode_t*', 'dependencies'), ('size_t', 'numDependencies'), ('const hipBatchMemOpNodeParams*', 'nodeParams')] +#define INIT_hipGraphAddBatchMemOpNode_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipGraphAddBatchMemOpNode.phGraphNode = (hipGraphNode_t*)phGraphNode; \ + cb_data.args.hipGraphAddBatchMemOpNode.hGraph = (hipGraph_t)hGraph; \ + cb_data.args.hipGraphAddBatchMemOpNode.dependencies = (const hipGraphNode_t*)dependencies; \ + cb_data.args.hipGraphAddBatchMemOpNode.numDependencies = (size_t)numDependencies; \ + cb_data.args.hipGraphAddBatchMemOpNode.nodeParams = (const hipBatchMemOpNodeParams*)nodeParams; \ +}; // hipGraphAddChildGraphNode[('hipGraphNode_t*', 'pGraphNode'), ('hipGraph_t', 'graph'), ('const hipGraphNode_t*', 'pDependencies'), ('size_t', 'numDependencies'), ('hipGraph_t', 'childGraph')] #define INIT_hipGraphAddChildGraphNode_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipGraphAddChildGraphNode.pGraphNode = (hipGraphNode_t*)pGraphNode; \ @@ -4415,6 +4437,16 @@ typedef struct hip_api_data_s { cb_data.args.hipGraphAddNode.numDependencies = (size_t)numDependencies; \ cb_data.args.hipGraphAddNode.nodeParams = (hipGraphNodeParams*)nodeParams; \ }; +// hipGraphBatchMemOpNodeGetParams[('hipGraphNode_t', 'hNode'), ('hipBatchMemOpNodeParams*', 'nodeParams_out')] +#define INIT_hipGraphBatchMemOpNodeGetParams_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipGraphBatchMemOpNodeGetParams.hNode = (hipGraphNode_t)hNode; \ + cb_data.args.hipGraphBatchMemOpNodeGetParams.nodeParams_out = (hipBatchMemOpNodeParams*)nodeParams_out; \ +}; +// hipGraphBatchMemOpNodeSetParams[('hipGraphNode_t', 'hNode'), ('hipBatchMemOpNodeParams*', 'nodeParams')] +#define INIT_hipGraphBatchMemOpNodeSetParams_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipGraphBatchMemOpNodeSetParams.hNode = (hipGraphNode_t)hNode; \ + cb_data.args.hipGraphBatchMemOpNodeSetParams.nodeParams = (hipBatchMemOpNodeParams*)nodeParams; \ +}; // hipGraphChildGraphNodeGetGraph[('hipGraphNode_t', 'node'), ('hipGraph_t*', 'pGraph')] #define INIT_hipGraphChildGraphNodeGetGraph_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipGraphChildGraphNodeGetGraph.node = (hipGraphNode_t)node; \ @@ -4464,6 +4496,12 @@ typedef struct hip_api_data_s { cb_data.args.hipGraphEventWaitNodeSetEvent.node = (hipGraphNode_t)node; \ cb_data.args.hipGraphEventWaitNodeSetEvent.event = (hipEvent_t)event; \ }; +// hipGraphExecBatchMemOpNodeSetParams[('hipGraphExec_t', 'hGraphExec'), ('hipGraphNode_t', 'hNode'), ('const hipBatchMemOpNodeParams*', 'nodeParams')] +#define INIT_hipGraphExecBatchMemOpNodeSetParams_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipGraphExecBatchMemOpNodeSetParams.hGraphExec = (hipGraphExec_t)hGraphExec; \ + cb_data.args.hipGraphExecBatchMemOpNodeSetParams.hNode = (hipGraphNode_t)hNode; \ + cb_data.args.hipGraphExecBatchMemOpNodeSetParams.nodeParams = (const hipBatchMemOpNodeParams*)nodeParams; \ +}; // hipGraphExecChildGraphNodeSetParams[('hipGraphExec_t', 'hGraphExec'), ('hipGraphNode_t', 'node'), ('hipGraph_t', 'childGraph')] #define INIT_hipGraphExecChildGraphNodeSetParams_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipGraphExecChildGraphNodeSetParams.hGraphExec = (hipGraphExec_t)hGraphExec; \ @@ -5802,6 +5840,13 @@ typedef struct hip_api_data_s { cb_data.args.hipStreamAttachMemAsync.length = (size_t)length; \ cb_data.args.hipStreamAttachMemAsync.flags = (unsigned int)flags; \ }; +// hipStreamBatchMemOp[('hipStream_t', 'stream'), ('unsigned int', 'count'), ('hipStreamBatchMemOpParams*', 'paramArray'), ('unsigned int', 'flags')] +#define INIT_hipStreamBatchMemOp_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipStreamBatchMemOp.stream = (hipStream_t)stream; \ + cb_data.args.hipStreamBatchMemOp.count = (unsigned int)count; \ + cb_data.args.hipStreamBatchMemOp.paramArray = (hipStreamBatchMemOpParams*)paramArray; \ + cb_data.args.hipStreamBatchMemOp.flags = (unsigned int)flags; \ +}; // hipStreamBeginCapture[('hipStream_t', 'stream'), ('hipStreamCaptureMode', 'mode')] #define INIT_hipStreamBeginCapture_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipStreamBeginCapture.stream = (hipStream_t)stream; \ @@ -5926,15 +5971,6 @@ typedef struct hip_api_data_s { cb_data.args.hipStreamWriteValue64.value = (uint64_t)value; \ cb_data.args.hipStreamWriteValue64.flags = (unsigned int)flags; \ }; - -// hipStreamBatchMemOp[('hipStream_t', 'stream'), ('unsigned int', 'count'), -// ('hipStreamBatchMemOpParams*', 'paramArray'), ('unsigned int', 'flags')] -#define INIT_hipStreamBatchMemOp_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipStreamBatchMemOp.stream = (hipStream_t)stream; \ - cb_data.args.hipStreamBatchMemOp.count = (unsigned int)count; \ - cb_data.args.hipStreamBatchMemOp.paramArray= (hipStreamBatchMemOpParams*)paramArray; \ - cb_data.args.hipStreamBatchMemOp.flags = (unsigned int)flags; \ -}; // hipTexRefGetAddress[('hipDeviceptr_t*', 'dev_ptr'), ('const textureReference*', 'texRef')] #define INIT_hipTexRefGetAddress_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipTexRefGetAddress.dev_ptr = (hipDeviceptr_t*)dptr; \ @@ -6069,36 +6105,6 @@ typedef struct hip_api_data_s { cb_data.args.hipWaitExternalSemaphoresAsync.numExtSems = (unsigned int)numExtSems; \ cb_data.args.hipWaitExternalSemaphoresAsync.stream = (hipStream_t)stream; \ }; - -// hipGraphAddBatchMemOpNode[('hipGraphNode_t*', 'phGraphNode'), ('hipGraph_t', 'hGraph'), -// ('hipGraphNode_t*', 'dependencies'), ('size_t', 'numDependencies'), -// ('hipBatchMemOpNodeParams*'), 'nodeParams')] -#define INIT_hipGraphAddBatchMemOpNode_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipGraphAddBatchMemOpNode.phGraphNode = (hipGraphNode_t*)phGraphNode; \ - cb_data.args.hipGraphAddBatchMemOpNode.hGraph = (hipGraph_t)hGraph; \ - cb_data.args.hipGraphAddBatchMemOpNode.dependencies= (hipGraphNode_t*)dependencies; \ - cb_data.args.hipGraphAddBatchMemOpNode.numDependencies = (size_t)numDependencies; \ - cb_data.args.hipGraphAddBatchMemOpNode.nodeParams = (hipBatchMemOpNodeParams*)nodeParams; \ -}; -// hipGraphBatchMemOpNodeGetParams[('hipGraphNode_t', hNode), -// ('hipBatchMemOpNodeParams*', 'nodeParams_out')] -#define INIT_hipGraphBatchMemOpNodeGetParams_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipGraphBatchMemOpNodeGetParams.hNode = (hipGraphNode_t)hNode; \ - cb_data.args.hipGraphBatchMemOpNodeGetParams.nodeParams_out = (hipBatchMemOpNodeParams*)nodeParams_out; \ -}; -// hipGraphBatchMemOpNodeSetParams[('hipGraphNode_t', hNode), -// ('hipBatchMemOpNodeParams*', 'nodeParams')] -#define INIT_hipGraphBatchMemOpNodeSetParams_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipGraphBatchMemOpNodeSetParams.hNode = (hipGraphNode_t)hNode; \ - cb_data.args.hipGraphBatchMemOpNodeSetParams.nodeParams = (hipBatchMemOpNodeParams*)nodeParams; \ -}; -// hipGraphExecBatchMemOpNodeSetParams[('hipGraphExec_t'. hGraphExec), -// ('hipGraphNode_t'. hNode), ('hipBatchMemOpNodeParams*', 'nodeParams')] -#define INIT_hipGraphExecBatchMemOpNodeSetParams_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipGraphExecBatchMemOpNodeSetParams.hGraphExec = (hipGraphExec_t)hGraphExec; \ - cb_data.args.hipGraphExecBatchMemOpNodeSetParams.hNode = (hipGraphNode_t)hNode; \ - cb_data.args.hipGraphExecBatchMemOpNodeSetParams.nodeParams= (hipBatchMemOpNodeParams*)nodeParams; \ -}; #define INIT_CB_ARGS_DATA(cb_id, cb_data) INIT_##cb_id##_CB_ARGS_DATA(cb_data) // Macros for non-public API primitives @@ -6472,6 +6478,9 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipEventRecord[('hipEvent_t', 'event'), ('hipStream_t', 'stream')] case HIP_API_ID_hipEventRecord: break; +// hipEventRecordWithFlags[('hipEvent_t', 'event'), ('hipStream_t', 'stream'), ('unsigned int', 'flags')] + case HIP_API_ID_hipEventRecordWithFlags: + break; // hipEventSynchronize[('hipEvent_t', 'event')] case HIP_API_ID_hipEventSynchronize: break; @@ -6608,6 +6617,12 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { case HIP_API_ID_hipGetSymbolSize: if (data->args.hipGetSymbolSize.size) data->args.hipGetSymbolSize.size__val = *(data->args.hipGetSymbolSize.size); break; +// hipGraphAddBatchMemOpNode[('hipGraphNode_t*', 'phGraphNode'), ('hipGraph_t', 'hGraph'), ('const hipGraphNode_t*', 'dependencies'), ('size_t', 'numDependencies'), ('const hipBatchMemOpNodeParams*', 'nodeParams')] + case HIP_API_ID_hipGraphAddBatchMemOpNode: + if (data->args.hipGraphAddBatchMemOpNode.phGraphNode) data->args.hipGraphAddBatchMemOpNode.phGraphNode__val = *(data->args.hipGraphAddBatchMemOpNode.phGraphNode); + if (data->args.hipGraphAddBatchMemOpNode.dependencies) data->args.hipGraphAddBatchMemOpNode.dependencies__val = *(data->args.hipGraphAddBatchMemOpNode.dependencies); + if (data->args.hipGraphAddBatchMemOpNode.nodeParams) data->args.hipGraphAddBatchMemOpNode.nodeParams__val = *(data->args.hipGraphAddBatchMemOpNode.nodeParams); + break; // hipGraphAddChildGraphNode[('hipGraphNode_t*', 'pGraphNode'), ('hipGraph_t', 'graph'), ('const hipGraphNode_t*', 'pDependencies'), ('size_t', 'numDependencies'), ('hipGraph_t', 'childGraph')] case HIP_API_ID_hipGraphAddChildGraphNode: if (data->args.hipGraphAddChildGraphNode.pGraphNode) data->args.hipGraphAddChildGraphNode.pGraphNode__val = *(data->args.hipGraphAddChildGraphNode.pGraphNode); @@ -6701,6 +6716,14 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { if (data->args.hipGraphAddNode.pDependencies) data->args.hipGraphAddNode.pDependencies__val = *(data->args.hipGraphAddNode.pDependencies); if (data->args.hipGraphAddNode.nodeParams) data->args.hipGraphAddNode.nodeParams__val = *(data->args.hipGraphAddNode.nodeParams); break; +// hipGraphBatchMemOpNodeGetParams[('hipGraphNode_t', 'hNode'), ('hipBatchMemOpNodeParams*', 'nodeParams_out')] + case HIP_API_ID_hipGraphBatchMemOpNodeGetParams: + if (data->args.hipGraphBatchMemOpNodeGetParams.nodeParams_out) data->args.hipGraphBatchMemOpNodeGetParams.nodeParams_out__val = *(data->args.hipGraphBatchMemOpNodeGetParams.nodeParams_out); + break; +// hipGraphBatchMemOpNodeSetParams[('hipGraphNode_t', 'hNode'), ('hipBatchMemOpNodeParams*', 'nodeParams')] + case HIP_API_ID_hipGraphBatchMemOpNodeSetParams: + if (data->args.hipGraphBatchMemOpNodeSetParams.nodeParams) data->args.hipGraphBatchMemOpNodeSetParams.nodeParams__val = *(data->args.hipGraphBatchMemOpNodeSetParams.nodeParams); + break; // hipGraphChildGraphNodeGetGraph[('hipGraphNode_t', 'node'), ('hipGraph_t*', 'pGraph')] case HIP_API_ID_hipGraphChildGraphNodeGetGraph: if (data->args.hipGraphChildGraphNodeGetGraph.pGraph) data->args.hipGraphChildGraphNodeGetGraph.pGraph__val = *(data->args.hipGraphChildGraphNodeGetGraph.pGraph); @@ -6737,6 +6760,10 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipGraphEventWaitNodeSetEvent[('hipGraphNode_t', 'node'), ('hipEvent_t', 'event')] case HIP_API_ID_hipGraphEventWaitNodeSetEvent: break; +// hipGraphExecBatchMemOpNodeSetParams[('hipGraphExec_t', 'hGraphExec'), ('hipGraphNode_t', 'hNode'), ('const hipBatchMemOpNodeParams*', 'nodeParams')] + case HIP_API_ID_hipGraphExecBatchMemOpNodeSetParams: + if (data->args.hipGraphExecBatchMemOpNodeSetParams.nodeParams) data->args.hipGraphExecBatchMemOpNodeSetParams.nodeParams__val = *(data->args.hipGraphExecBatchMemOpNodeSetParams.nodeParams); + break; // hipGraphExecChildGraphNodeSetParams[('hipGraphExec_t', 'hGraphExec'), ('hipGraphNode_t', 'node'), ('hipGraph_t', 'childGraph')] case HIP_API_ID_hipGraphExecChildGraphNodeSetParams: break; @@ -7515,6 +7542,10 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipStreamAttachMemAsync[('hipStream_t', 'stream'), ('void*', 'dev_ptr'), ('size_t', 'length'), ('unsigned int', 'flags')] case HIP_API_ID_hipStreamAttachMemAsync: break; +// hipStreamBatchMemOp[('hipStream_t', 'stream'), ('unsigned int', 'count'), ('hipStreamBatchMemOpParams*', 'paramArray'), ('unsigned int', 'flags')] + case HIP_API_ID_hipStreamBatchMemOp: + if (data->args.hipStreamBatchMemOp.paramArray) data->args.hipStreamBatchMemOp.paramArray__val = *(data->args.hipStreamBatchMemOp.paramArray); + break; // hipStreamBeginCapture[('hipStream_t', 'stream'), ('hipStreamCaptureMode', 'mode')] case HIP_API_ID_hipStreamBeginCapture: break; @@ -7596,34 +7627,6 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipStreamWriteValue64[('hipStream_t', 'stream'), ('void*', 'ptr'), ('uint64_t', 'value'), ('unsigned int', 'flags')] case HIP_API_ID_hipStreamWriteValue64: break; -// hipStreamBatchMemOp[('hipStream_t', 'stream'), ('unsigned int', 'count'), -// ('hipStreamBatchMemOpParams*', 'paramArray'), ('unsigned int', 'flags')] - case HIP_API_ID_hipStreamBatchMemOp: - if (data->args.hipStreamBatchMemOp.paramArray) data->args.hipStreamBatchMemOp.paramArray__val = *(data->args.hipStreamBatchMemOp.paramArray); - break; -// hipGraphAddBatchMemOpNode[('hipGraphNode_t*', 'phGraphNode'), ('hipGraph_t', 'hGraph'), -// ('hipGraphNode_t*', 'dependencies'), ('size_t', 'numDependencies'), -// ('hipBatchMemOpNodeParams*'), 'nodeParams')] - case HIP_API_ID_hipGraphAddBatchMemOpNode: - if (data->args.hipGraphAddBatchMemOpNode.phGraphNode) data->args.hipGraphAddBatchMemOpNode.phGraphNode__val = *(data->args.hipGraphAddBatchMemOpNode.phGraphNode); - if (data->args.hipGraphAddBatchMemOpNode.dependencies) data->args.hipGraphAddBatchMemOpNode.dependencies__val = *(data->args.hipGraphAddBatchMemOpNode.dependencies); - if (data->args.hipGraphAddBatchMemOpNode.nodeParams) data->args.hipGraphAddBatchMemOpNode.nodeParams__val = *(data->args.hipGraphAddBatchMemOpNode.nodeParams); - break; -// hipGraphBatchMemOpNodeGetParams[('hipGraphNode_t', hNode), -// ('hipBatchMemOpNodeParams*', 'nodeParams_out')] - case HIP_API_ID_hipGraphBatchMemOpNodeGetParams: - if (data->args.hipGraphBatchMemOpNodeGetParams.nodeParams_out) data->args.hipGraphBatchMemOpNodeGetParams.nodeParams_out__val = *(data->args.hipGraphBatchMemOpNodeGetParams.nodeParams_out); - break; -// hipGraphBatchMemOpNodeSetParams[('hipGraphNode_t', hNode), -// ('hipBatchMemOpNodeParams*', 'nodeParams')] - case HIP_API_ID_hipGraphBatchMemOpNodeSetParams: - if (data->args.hipGraphBatchMemOpNodeSetParams.nodeParams) data->args.hipGraphBatchMemOpNodeSetParams.nodeParams__val = *(data->args.hipGraphBatchMemOpNodeSetParams.nodeParams); - break; -// hipGraphExecBatchMemOpNodeSetParams[('hipGraphExec_t'. hGraphExec), -// ('hipGraphNode_t'. hNode), ('hipBatchMemOpNodeParams*', 'nodeParams')] - case HIP_API_ID_hipGraphExecBatchMemOpNodeSetParams: - if (data->args.hipGraphExecBatchMemOpNodeSetParams.nodeParams) data->args.hipGraphExecBatchMemOpNodeSetParams.nodeParams__val = *(data->args.hipGraphExecBatchMemOpNodeSetParams.nodeParams); - break; // hipTexRefGetAddress[('hipDeviceptr_t*', 'dev_ptr'), ('const textureReference*', 'texRef')] case HIP_API_ID_hipTexRefGetAddress: if (data->args.hipTexRefGetAddress.dev_ptr) data->args.hipTexRefGetAddress.dev_ptr__val = *(data->args.hipTexRefGetAddress.dev_ptr); @@ -8301,6 +8304,13 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipEventRecord.stream); oss << ")"; break; + case HIP_API_ID_hipEventRecordWithFlags: + oss << "hipEventRecordWithFlags("; + oss << "event="; roctracer::hip_support::detail::operator<<(oss, data->args.hipEventRecordWithFlags.event); + oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipEventRecordWithFlags.stream); + oss << ", flags="; roctracer::hip_support::detail::operator<<(oss, data->args.hipEventRecordWithFlags.flags); + oss << ")"; + break; case HIP_API_ID_hipEventSynchronize: oss << "hipEventSynchronize("; oss << "event="; roctracer::hip_support::detail::operator<<(oss, data->args.hipEventSynchronize.event); @@ -8563,6 +8573,18 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << ", symbol="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetSymbolSize.symbol); oss << ")"; break; + case HIP_API_ID_hipGraphAddBatchMemOpNode: + oss << "hipGraphAddBatchMemOpNode("; + if (data->args.hipGraphAddBatchMemOpNode.phGraphNode == NULL) oss << "phGraphNode=NULL"; + else { oss << "phGraphNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddBatchMemOpNode.phGraphNode__val); } + oss << ", hGraph="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddBatchMemOpNode.hGraph); + if (data->args.hipGraphAddBatchMemOpNode.dependencies == NULL) oss << ", dependencies=NULL"; + else { oss << ", dependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddBatchMemOpNode.dependencies__val); } + oss << ", numDependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddBatchMemOpNode.numDependencies); + if (data->args.hipGraphAddBatchMemOpNode.nodeParams == NULL) oss << ", nodeParams=NULL"; + else { oss << ", nodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddBatchMemOpNode.nodeParams__val); } + oss << ")"; + break; case HIP_API_ID_hipGraphAddChildGraphNode: oss << "hipGraphAddChildGraphNode("; if (data->args.hipGraphAddChildGraphNode.pGraphNode == NULL) oss << "pGraphNode=NULL"; @@ -8767,6 +8789,20 @@ 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.hipGraphAddNode.nodeParams__val); } oss << ")"; break; + case HIP_API_ID_hipGraphBatchMemOpNodeGetParams: + oss << "hipGraphBatchMemOpNodeGetParams("; + oss << "hNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphBatchMemOpNodeGetParams.hNode); + if (data->args.hipGraphBatchMemOpNodeGetParams.nodeParams_out == NULL) oss << ", nodeParams_out=NULL"; + else { oss << ", nodeParams_out="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphBatchMemOpNodeGetParams.nodeParams_out__val); } + oss << ")"; + break; + case HIP_API_ID_hipGraphBatchMemOpNodeSetParams: + oss << "hipGraphBatchMemOpNodeSetParams("; + oss << "hNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphBatchMemOpNodeSetParams.hNode); + if (data->args.hipGraphBatchMemOpNodeSetParams.nodeParams == NULL) oss << ", nodeParams=NULL"; + else { oss << ", nodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphBatchMemOpNodeSetParams.nodeParams__val); } + oss << ")"; + break; case HIP_API_ID_hipGraphChildGraphNodeGetGraph: oss << "hipGraphChildGraphNodeGetGraph("; oss << "node="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphChildGraphNodeGetGraph.node); @@ -8832,6 +8868,14 @@ 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.hipGraphEventWaitNodeSetEvent.event); oss << ")"; break; + case HIP_API_ID_hipGraphExecBatchMemOpNodeSetParams: + oss << "hipGraphExecBatchMemOpNodeSetParams("; + oss << "hGraphExec="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecBatchMemOpNodeSetParams.hGraphExec); + oss << ", hNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecBatchMemOpNodeSetParams.hNode); + if (data->args.hipGraphExecBatchMemOpNodeSetParams.nodeParams == NULL) oss << ", nodeParams=NULL"; + else { oss << ", nodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecBatchMemOpNodeSetParams.nodeParams__val); } + oss << ")"; + break; case HIP_API_ID_hipGraphExecChildGraphNodeSetParams: oss << "hipGraphExecChildGraphNodeSetParams("; oss << "hGraphExec="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecChildGraphNodeSetParams.hGraphExec); @@ -10540,6 +10584,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.hipStreamAttachMemAsync.flags); oss << ")"; break; + case HIP_API_ID_hipStreamBatchMemOp: + oss << "hipStreamBatchMemOp("; + oss << "stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBatchMemOp.stream); + oss << ", count="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBatchMemOp.count); + if (data->args.hipStreamBatchMemOp.paramArray == NULL) oss << ", paramArray=NULL"; + else { oss << ", paramArray="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBatchMemOp.paramArray__val); } + oss << ", flags="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBatchMemOp.flags); + oss << ")"; + break; case HIP_API_ID_hipStreamBeginCapture: oss << "hipStreamBeginCapture("; oss << "stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBeginCapture.stream); @@ -10703,15 +10756,6 @@ 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.hipStreamWriteValue64.flags); oss << ")"; break; - case HIP_API_ID_hipStreamBatchMemOp: - oss << "hipStreamBatchMemOp("; - oss << "stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBatchMemOp.stream); - oss << ", count="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBatchMemOp.count); - if (data->args.hipStreamBatchMemOp.paramArray == NULL) oss << ", paramArray=NULL"; - else { oss << ", paramArray="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBatchMemOp.paramArray__val); } - oss << ", flags="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBatchMemOp.flags); - oss << ")"; - break; case HIP_API_ID_hipTexRefGetAddress: oss << "hipTexRefGetAddress("; if (data->args.hipTexRefGetAddress.dev_ptr == NULL) oss << "dev_ptr=NULL"; @@ -10908,34 +10952,6 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipWaitExternalSemaphoresAsync.stream); oss << ")"; break; - case HIP_API_ID_hipGraphAddBatchMemOpNode: - oss << "hipGraphAddBatchMemOpNode("; - oss << "phGraphNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddBatchMemOpNode.phGraphNode); - oss << ", hGraph="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddBatchMemOpNode.hGraph); - oss << ", dependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddBatchMemOpNode.dependencies); - oss << ", numDependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddBatchMemOpNode.numDependencies); - oss << ", nodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddBatchMemOpNode.nodeParams); - oss << ")"; - break; - case HIP_API_ID_hipGraphBatchMemOpNodeGetParams: - oss << "hipGraphBatchMemOpNodeGetParams("; - oss << "hNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphBatchMemOpNodeGetParams.hNode); - oss << ", nodeParams_out="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphBatchMemOpNodeGetParams.nodeParams_out); - oss << ")"; - break; - case HIP_API_ID_hipGraphBatchMemOpNodeSetParams: - oss << "hipGraphBatchMemOpNodeSetParams("; - oss << "hNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphBatchMemOpNodeSetParams.hNode); - oss << ", nodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphBatchMemOpNodeSetParams.nodeParams); - oss << ")"; - break; - case HIP_API_ID_hipGraphExecBatchMemOpNodeSetParams: - oss << "hipGraphExecBatchMemOpNodeSetParams("; - oss << "hGraphExec="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecBatchMemOpNodeSetParams.hGraphExec); - oss << ", hNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecBatchMemOpNodeSetParams.hNode); - oss << ", nodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecBatchMemOpNodeSetParams.nodeParams); - oss << ")"; - break; default: oss << "unknown"; }; return strdup(oss.str().c_str()); diff --git a/projects/clr/hipamd/src/amdhip.def b/projects/clr/hipamd/src/amdhip.def index 5327ac5628..28d8d7be9f 100644 --- a/projects/clr/hipamd/src/amdhip.def +++ b/projects/clr/hipamd/src/amdhip.def @@ -484,3 +484,4 @@ hipGraphAddBatchMemOpNode hipGraphBatchMemOpNodeGetParams hipGraphBatchMemOpNodeSetParams hipGraphExecBatchMemOpNodeSetParams +hipEventRecordWithFlags diff --git a/projects/clr/hipamd/src/hip_api_trace.cpp b/projects/clr/hipamd/src/hip_api_trace.cpp index 506b23d5f7..73b352a891 100644 --- a/projects/clr/hipamd/src/hip_api_trace.cpp +++ b/projects/clr/hipamd/src/hip_api_trace.cpp @@ -813,6 +813,7 @@ hipError_t hipGraphBatchMemOpNodeSetParams(hipGraphNode_t hNode, hipBatchMemOpNodeParams* nodeParams); hipError_t hipGraphExecBatchMemOpNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t hNode, const hipBatchMemOpNodeParams* nodeParams); +hipError_t hipEventRecordWithFlags(hipEvent_t event, hipStream_t stream, unsigned flags); } // namespace hip namespace hip { @@ -1319,6 +1320,7 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) { ptrDispatchTable->hipGraphBatchMemOpNodeSetParams_fn = hip::hipGraphBatchMemOpNodeSetParams; ptrDispatchTable->hipGraphExecBatchMemOpNodeSetParams_fn = hip::hipGraphExecBatchMemOpNodeSetParams; + ptrDispatchTable->hipEventRecordWithFlags_fn = hip::hipEventRecordWithFlags; } #if HIP_ROCPROFILER_REGISTER > 0 @@ -1948,16 +1950,17 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipGraphAddBatchMemOpNode_fn, 464); HIP_ENFORCE_ABI(HipDispatchTable, hipGraphBatchMemOpNodeGetParams_fn, 465); HIP_ENFORCE_ABI(HipDispatchTable, hipGraphBatchMemOpNodeSetParams_fn, 466); HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExecBatchMemOpNodeSetParams_fn, 467); - +// HIP_RUNTIME_API_TABLE_STEP_VERSION == 9 +HIP_ENFORCE_ABI(HipDispatchTable, hipEventRecordWithFlags_fn, 468) // 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.: // // HIP_ENFORCE_ABI(, , 8) // // HIP_ENFORCE_ABI_VERSIONING(
, 9) <- 8 + 1 = 9 -HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 468) +HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 469) -static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 8, +static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 9, "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/projects/clr/hipamd/src/hip_event.cpp b/projects/clr/hipamd/src/hip_event.cpp index c04f470e6b..619af62467 100644 --- a/projects/clr/hipamd/src/hip_event.cpp +++ b/projects/clr/hipamd/src/hip_event.cpp @@ -388,7 +388,10 @@ hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) { HIP_RETURN(eStart->elapsedTime(*eStop, *ms), "Elapsed Time = ", *ms); } -hipError_t hipEventRecord_common(hipEvent_t event, hipStream_t stream) { +hipError_t hipEventRecord_common(hipEvent_t event, hipStream_t stream, unsigned int flags) { + if (!(flags == hipEventRecordDefault || flags == hipEventRecordExternal)){ + return hipErrorInvalidValue; + } hipError_t status = hipSuccess; if (event == nullptr) { return hipErrorInvalidHandle; @@ -410,6 +413,17 @@ hipError_t hipEventRecord_common(hipEvent_t event, hipStream_t stream) { if (!lastCapturedNodes.empty()) { e->SetNodesPrevToRecorded(lastCapturedNodes); } + if (flags == hipEventRecordExternal) { + hip::GraphNode* pGraphNode; + status = hipGraphAddEventRecordNode((hipGraphNode_t*) pGraphNode, (hipGraph_t) s->GetCaptureGraph(), + (hipGraphNode_t*) s->GetLastCapturedNodes().data(), + s->GetLastCapturedNodes().size(), (hipEvent_t) e); + if (status != hipSuccess) { + ClPrint(amd::LOG_ERROR, amd::LOG_API, "hipEventRecord add external event node failed"); + return status; + } + s->SetLastCapturedNode(pGraphNode); + } } else { if (g_devices[e->deviceId()]->devices()[0] != &hip_stream->device()) { return hipErrorInvalidHandle; @@ -421,13 +435,18 @@ hipError_t hipEventRecord_common(hipEvent_t event, hipStream_t stream) { hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { HIP_INIT_API(hipEventRecord, event, stream); - HIP_RETURN(hipEventRecord_common(event, stream)); + HIP_RETURN(hipEventRecord_common(event, stream, hipEventRecordDefault)); } hipError_t hipEventRecord_spt(hipEvent_t event, hipStream_t stream) { HIP_INIT_API(hipEventRecord, event, stream); PER_THREAD_DEFAULT_STREAM(stream); - HIP_RETURN(hipEventRecord_common(event, stream)); + HIP_RETURN(hipEventRecord_common(event, stream, hipEventRecordDefault)); +} + +hipError_t hipEventRecordWithFlags(hipEvent_t event, hipStream_t stream, unsigned int flags) { + HIP_INIT_API(hipEventRecordWithFlags, event, stream, flags); + HIP_RETURN(hipEventRecord_common(event, stream, flags)); } hipError_t hipEventSynchronize(hipEvent_t event) { diff --git a/projects/clr/hipamd/src/hip_hcc.map.in b/projects/clr/hipamd/src/hip_hcc.map.in index 785845f0c3..1ba5a6f89a 100644 --- a/projects/clr/hipamd/src/hip_hcc.map.in +++ b/projects/clr/hipamd/src/hip_hcc.map.in @@ -587,6 +587,7 @@ global: hipGraphBatchMemOpNodeGetParams; hipGraphBatchMemOpNodeSetParams; hipGraphExecBatchMemOpNodeSetParams; + hipEventRecordWithFlags; local: *; } hip_6.2; diff --git a/projects/clr/hipamd/src/hip_table_interface.cpp b/projects/clr/hipamd/src/hip_table_interface.cpp index 77092e1fa1..cc34a40804 100644 --- a/projects/clr/hipamd/src/hip_table_interface.cpp +++ b/projects/clr/hipamd/src/hip_table_interface.cpp @@ -1840,4 +1840,7 @@ hipError_t hipGraphExecBatchMemOpNodeSetParams(hipGraphExec_t hGraphExec, const hipBatchMemOpNodeParams* nodeParams) { return hip::GetHipDispatchTable()->hipGraphExecBatchMemOpNodeSetParams_fn(hGraphExec, hNode, nodeParams); +} +hipError_t hipEventRecordWithFlags(hipEvent_t event, hipStream_t stream, unsigned int flags) { + return hip::GetHipDispatchTable()->hipEventRecordWithFlags_fn(event, stream, flags); } \ No newline at end of file