AI/rocm-systems
2
0
Förgrening 0
Kod Ärenden Pull-förfrågningar Actions Packages Projekt Släpp Wiki Aktiviteter

SWDEV-477219 - implement hipEventRecordWithFlags

Bläddra i källkod 
Change-Id: Icf07e85fc8c15f921f6e7c9fbd31dd3856dc988b


[ROCm/clr commit: 7a4a22d454]
This commit is contained in:
Jimbo Xie
2025-02-06 00:12:48 -05:00
committad av Jiabao Xie
förälder 16f9dbff6c
incheckning 8a42a52d0f
7 ändrade filer med 206 tillägg och 158 borttagningar
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
Visa fil
@@ -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
// ******************************************************************************************* //
//
projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h
+166 -150
Visa fil
@@ -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());
projects/clr/hipamd/src/amdhip.def
+1
Visa fil
@@ -484,3 +484,4 @@ hipGraphAddBatchMemOpNode
hipGraphBatchMemOpNodeGetParams
hipGraphBatchMemOpNodeSetParams
hipGraphExecBatchMemOpNodeSetParams
hipEventRecordWithFlags
projects/clr/hipamd/src/hip_api_trace.cpp
+6 -3
Visa fil
@@ -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(<table>, <functor>, 8)
//
// HIP_ENFORCE_ABI_VERSIONING(<table>, 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
projects/clr/hipamd/src/hip_event.cpp
+22 -3
Visa fil
@@ -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) {
projects/clr/hipamd/src/hip_hcc.map.in
+1
Visa fil
@@ -587,6 +587,7 @@ global:
hipGraphBatchMemOpNodeGetParams;
hipGraphBatchMemOpNodeSetParams;
hipGraphExecBatchMemOpNodeSetParams;
hipEventRecordWithFlags;
local:
*;
} hip_6.2;
projects/clr/hipamd/src/hip_table_interface.cpp
+3
Visa fil
@@ -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);
}