From 45dccd6eab1b87c2668aa562fe0491cd65df105c Mon Sep 17 00:00:00 2001 From: Anusha GodavarthySurya Date: Mon, 22 Apr 2024 14:53:13 +0000 Subject: [PATCH] SWDEV-453739 - Added definition for undeclared methods Change-Id: I3f10e0a3fb4bbaf9d873a6a988847f05ac43a312 --- .../include/hip/amd_detail/hip_api_trace.hpp | 25 +- hipamd/include/hip/amd_detail/hip_prof_str.h | 253 ++++++++++++++++-- hipamd/src/amdhip.def | 7 + hipamd/src/hip_api_trace.cpp | 55 +++- hipamd/src/hip_hcc.map.in | 7 + hipamd/src/hip_table_interface.cpp | 33 +++ 6 files changed, 344 insertions(+), 36 deletions(-) diff --git a/hipamd/include/hip/amd_detail/hip_api_trace.hpp b/hipamd/include/hip/amd_detail/hip_api_trace.hpp index 957ea75627..08889fa5f3 100644 --- a/hipamd/include/hip/amd_detail/hip_api_trace.hpp +++ b/hipamd/include/hip/amd_detail/hip_api_trace.hpp @@ -953,7 +953,6 @@ typedef hipError_t (*t_hipStreamBeginCaptureToGraph)(hipStream_t stream, hipGrap size_t numDependencies, hipStreamCaptureMode mode); typedef hipError_t (*t_hipGetFuncBySymbol)(hipFunction_t* functionPtr, const void* symbolPtr); - typedef hipError_t (*t_hipDrvGraphAddMemFreeNode)(hipGraphNode_t* phGraphNode, hipGraph_t hGraph, const hipGraphNode_t* dependencies, size_t numDependencies, hipDeviceptr_t dptr); @@ -965,6 +964,23 @@ typedef hipError_t (*t_hipDrvGraphExecMemcpyNodeSetParams)(hipGraphExec_t hGraph typedef hipError_t (*t_hipDrvGraphExecMemsetNodeSetParams)(hipGraphExec_t hGraphExec, hipGraphNode_t hNode, const HIP_MEMSET_NODE_PARAMS* memsetParams, hipCtx_t ctx); +typedef hipError_t (*t_hipSetValidDevices)(int* device_arr, int len); +typedef hipError_t (*t_hipMemcpyAtoD)(hipDeviceptr_t dstDevice, hipArray_t srcArray, + size_t srcOffset, size_t ByteCount); +typedef hipError_t (*t_hipMemcpyDtoA)(hipArray_t dstArray, size_t dstOffset, + hipDeviceptr_t srcDevice, size_t ByteCount); +typedef hipError_t (*t_hipMemcpyAtoA)(hipArray_t dstArray, size_t dstOffset, hipArray_t srcArray, + size_t srcOffset, size_t ByteCount); +typedef hipError_t (*t_hipMemcpyAtoHAsync)(void* dstHost, hipArray_t srcArray, size_t srcOffset, + size_t ByteCount, hipStream_t stream); +typedef hipError_t (*t_hipMemcpyHtoAAsync)(hipArray_t dstArray, size_t dstOffset, + const void* srcHost, size_t ByteCount, + hipStream_t stream); +typedef hipError_t (*t_hipMemcpy2DArrayToArray)(hipArray_t dst, size_t wOffsetDst, + size_t hOffsetDst, hipArray_const_t src, + size_t wOffsetSrc, size_t hOffsetSrc, size_t width, + size_t height, hipMemcpyKind kind); + // HIP Compiler dispatch table struct HipCompilerDispatchTable { @@ -1435,4 +1451,11 @@ struct HipDispatchTable { t_hipDrvGraphAddMemFreeNode hipDrvGraphAddMemFreeNode_fn; t_hipDrvGraphExecMemcpyNodeSetParams hipDrvGraphExecMemcpyNodeSetParams_fn; t_hipDrvGraphExecMemsetNodeSetParams hipDrvGraphExecMemsetNodeSetParams_fn; + t_hipSetValidDevices hipSetValidDevices_fn; + t_hipMemcpyAtoD hipMemcpyAtoD_fn; + t_hipMemcpyDtoA hipMemcpyDtoA_fn; + t_hipMemcpyAtoA hipMemcpyAtoA_fn; + t_hipMemcpyAtoHAsync hipMemcpyAtoHAsync_fn; + t_hipMemcpyHtoAAsync hipMemcpyHtoAAsync_fn; + t_hipMemcpy2DArrayToArray hipMemcpy2DArrayToArray_fn; }; diff --git a/hipamd/include/hip/amd_detail/hip_prof_str.h b/hipamd/include/hip/amd_detail/hip_prof_str.h index fcebe3d9f6..3cf8291215 100644 --- a/hipamd/include/hip/amd_detail/hip_prof_str.h +++ b/hipamd/include/hip/amd_detail/hip_prof_str.h @@ -410,7 +410,21 @@ enum hip_api_id_t { HIP_API_ID_hipTexRefGetBorderColor = 390, HIP_API_ID_hipStreamBeginCaptureToGraph = 391, HIP_API_ID_hipGetFuncBySymbol = 392, - HIP_API_ID_LAST = 392, + HIP_API_ID_RESERVED_393 = 393, + HIP_API_ID_RESERVED_394 = 394, + HIP_API_ID_RESERVED_395 = 395, + HIP_API_ID_RESERVED_396 = 396, + HIP_API_ID_RESERVED_397 = 397, + HIP_API_ID_RESERVED_398 = 398, + HIP_API_ID_RESERVED_399 = 399, + HIP_API_ID_hipMemcpy2DArrayToArray = 400, + HIP_API_ID_hipMemcpyAtoA = 401, + HIP_API_ID_hipMemcpyAtoD = 402, + HIP_API_ID_hipMemcpyAtoHAsync = 403, + HIP_API_ID_hipMemcpyDtoA = 404, + HIP_API_ID_hipMemcpyHtoAAsync = 405, + HIP_API_ID_hipSetValidDevices = 406, + HIP_API_ID_LAST = 406, HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice), HIP_API_ID_hipGetDeviceProperties = HIP_API_ID_CONCAT(HIP_API_ID_,hipGetDeviceProperties), @@ -427,13 +441,6 @@ enum hip_api_id_t { HIP_API_ID_hipGetTextureObjectResourceViewDesc = HIP_API_ID_NONE, HIP_API_ID_hipGetTextureObjectTextureDesc = HIP_API_ID_NONE, HIP_API_ID_hipGetTextureReference = HIP_API_ID_NONE, - HIP_API_ID_hipMemcpy2DArrayToArray = HIP_API_ID_NONE, - HIP_API_ID_hipMemcpyAtoA = HIP_API_ID_NONE, - HIP_API_ID_hipMemcpyAtoD = HIP_API_ID_NONE, - HIP_API_ID_hipMemcpyAtoHAsync = HIP_API_ID_NONE, - HIP_API_ID_hipMemcpyDtoA = HIP_API_ID_NONE, - HIP_API_ID_hipMemcpyHtoAAsync = HIP_API_ID_NONE, - HIP_API_ID_hipSetValidDevices = HIP_API_ID_NONE, HIP_API_ID_hipTexObjectCreate = HIP_API_ID_NONE, HIP_API_ID_hipTexObjectDestroy = HIP_API_ID_NONE, HIP_API_ID_hipTexObjectGetResourceDesc = HIP_API_ID_NONE, @@ -725,6 +732,7 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipMemUnmap: return "hipMemUnmap"; case HIP_API_ID_hipMemcpy: return "hipMemcpy"; case HIP_API_ID_hipMemcpy2D: return "hipMemcpy2D"; + case HIP_API_ID_hipMemcpy2DArrayToArray: return "hipMemcpy2DArrayToArray"; case HIP_API_ID_hipMemcpy2DAsync: return "hipMemcpy2DAsync"; case HIP_API_ID_hipMemcpy2DFromArray: return "hipMemcpy2DFromArray"; case HIP_API_ID_hipMemcpy2DFromArrayAsync: return "hipMemcpy2DFromArrayAsync"; @@ -733,7 +741,11 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipMemcpy3D: return "hipMemcpy3D"; case HIP_API_ID_hipMemcpy3DAsync: return "hipMemcpy3DAsync"; case HIP_API_ID_hipMemcpyAsync: return "hipMemcpyAsync"; + case HIP_API_ID_hipMemcpyAtoA: return "hipMemcpyAtoA"; + case HIP_API_ID_hipMemcpyAtoD: return "hipMemcpyAtoD"; case HIP_API_ID_hipMemcpyAtoH: return "hipMemcpyAtoH"; + case HIP_API_ID_hipMemcpyAtoHAsync: return "hipMemcpyAtoHAsync"; + case HIP_API_ID_hipMemcpyDtoA: return "hipMemcpyDtoA"; case HIP_API_ID_hipMemcpyDtoD: return "hipMemcpyDtoD"; case HIP_API_ID_hipMemcpyDtoDAsync: return "hipMemcpyDtoDAsync"; case HIP_API_ID_hipMemcpyDtoH: return "hipMemcpyDtoH"; @@ -742,6 +754,7 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipMemcpyFromSymbol: return "hipMemcpyFromSymbol"; case HIP_API_ID_hipMemcpyFromSymbolAsync: return "hipMemcpyFromSymbolAsync"; case HIP_API_ID_hipMemcpyHtoA: return "hipMemcpyHtoA"; + case HIP_API_ID_hipMemcpyHtoAAsync: return "hipMemcpyHtoAAsync"; case HIP_API_ID_hipMemcpyHtoD: return "hipMemcpyHtoD"; case HIP_API_ID_hipMemcpyHtoDAsync: return "hipMemcpyHtoDAsync"; case HIP_API_ID_hipMemcpyParam2D: return "hipMemcpyParam2D"; @@ -793,6 +806,7 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipRuntimeGetVersion: return "hipRuntimeGetVersion"; case HIP_API_ID_hipSetDevice: return "hipSetDevice"; case HIP_API_ID_hipSetDeviceFlags: return "hipSetDeviceFlags"; + case HIP_API_ID_hipSetValidDevices: return "hipSetValidDevices"; case HIP_API_ID_hipSetupArgument: return "hipSetupArgument"; case HIP_API_ID_hipSignalExternalSemaphoresAsync: return "hipSignalExternalSemaphoresAsync"; case HIP_API_ID_hipStreamAddCallback: return "hipStreamAddCallback"; @@ -1120,6 +1134,7 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipMemUnmap", name) == 0) return HIP_API_ID_hipMemUnmap; if (strcmp("hipMemcpy", name) == 0) return HIP_API_ID_hipMemcpy; if (strcmp("hipMemcpy2D", name) == 0) return HIP_API_ID_hipMemcpy2D; + if (strcmp("hipMemcpy2DArrayToArray", name) == 0) return HIP_API_ID_hipMemcpy2DArrayToArray; if (strcmp("hipMemcpy2DAsync", name) == 0) return HIP_API_ID_hipMemcpy2DAsync; if (strcmp("hipMemcpy2DFromArray", name) == 0) return HIP_API_ID_hipMemcpy2DFromArray; if (strcmp("hipMemcpy2DFromArrayAsync", name) == 0) return HIP_API_ID_hipMemcpy2DFromArrayAsync; @@ -1128,7 +1143,11 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipMemcpy3D", name) == 0) return HIP_API_ID_hipMemcpy3D; if (strcmp("hipMemcpy3DAsync", name) == 0) return HIP_API_ID_hipMemcpy3DAsync; if (strcmp("hipMemcpyAsync", name) == 0) return HIP_API_ID_hipMemcpyAsync; + if (strcmp("hipMemcpyAtoA", name) == 0) return HIP_API_ID_hipMemcpyAtoA; + if (strcmp("hipMemcpyAtoD", name) == 0) return HIP_API_ID_hipMemcpyAtoD; if (strcmp("hipMemcpyAtoH", name) == 0) return HIP_API_ID_hipMemcpyAtoH; + if (strcmp("hipMemcpyAtoHAsync", name) == 0) return HIP_API_ID_hipMemcpyAtoHAsync; + if (strcmp("hipMemcpyDtoA", name) == 0) return HIP_API_ID_hipMemcpyDtoA; if (strcmp("hipMemcpyDtoD", name) == 0) return HIP_API_ID_hipMemcpyDtoD; if (strcmp("hipMemcpyDtoDAsync", name) == 0) return HIP_API_ID_hipMemcpyDtoDAsync; if (strcmp("hipMemcpyDtoH", name) == 0) return HIP_API_ID_hipMemcpyDtoH; @@ -1137,6 +1156,7 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipMemcpyFromSymbol", name) == 0) return HIP_API_ID_hipMemcpyFromSymbol; if (strcmp("hipMemcpyFromSymbolAsync", name) == 0) return HIP_API_ID_hipMemcpyFromSymbolAsync; if (strcmp("hipMemcpyHtoA", name) == 0) return HIP_API_ID_hipMemcpyHtoA; + if (strcmp("hipMemcpyHtoAAsync", name) == 0) return HIP_API_ID_hipMemcpyHtoAAsync; if (strcmp("hipMemcpyHtoD", name) == 0) return HIP_API_ID_hipMemcpyHtoD; if (strcmp("hipMemcpyHtoDAsync", name) == 0) return HIP_API_ID_hipMemcpyHtoDAsync; if (strcmp("hipMemcpyParam2D", name) == 0) return HIP_API_ID_hipMemcpyParam2D; @@ -1188,6 +1208,7 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipRuntimeGetVersion", name) == 0) return HIP_API_ID_hipRuntimeGetVersion; if (strcmp("hipSetDevice", name) == 0) return HIP_API_ID_hipSetDevice; if (strcmp("hipSetDeviceFlags", name) == 0) return HIP_API_ID_hipSetDeviceFlags; + if (strcmp("hipSetValidDevices", name) == 0) return HIP_API_ID_hipSetValidDevices; if (strcmp("hipSetupArgument", name) == 0) return HIP_API_ID_hipSetupArgument; if (strcmp("hipSignalExternalSemaphoresAsync", name) == 0) return HIP_API_ID_hipSignalExternalSemaphoresAsync; if (strcmp("hipStreamAddCallback", name) == 0) return HIP_API_ID_hipStreamAddCallback; @@ -2809,6 +2830,17 @@ typedef struct hip_api_data_s { size_t height; hipMemcpyKind kind; } hipMemcpy2D; + struct { + hipArray_t dst; + size_t wOffsetDst; + size_t hOffsetDst; + hipArray_const_t src; + size_t wOffsetSrc; + size_t hOffsetSrc; + size_t width; + size_t height; + hipMemcpyKind kind; + } hipMemcpy2DArrayToArray; struct { void* dst; size_t dpitch; @@ -2877,12 +2909,38 @@ typedef struct hip_api_data_s { hipMemcpyKind kind; hipStream_t stream; } hipMemcpyAsync; + struct { + hipArray_t dstArray; + size_t dstOffset; + hipArray_t srcArray; + size_t srcOffset; + size_t ByteCount; + } hipMemcpyAtoA; + struct { + hipDeviceptr_t dstDevice; + hipArray_t srcArray; + size_t srcOffset; + size_t ByteCount; + } hipMemcpyAtoD; struct { void* dst; hipArray_t srcArray; size_t srcOffset; size_t count; } hipMemcpyAtoH; + struct { + void* dstHost; + hipArray_t srcArray; + size_t srcOffset; + size_t ByteCount; + hipStream_t stream; + } hipMemcpyAtoHAsync; + struct { + hipArray_t dstArray; + size_t dstOffset; + hipDeviceptr_t srcDevice; + size_t ByteCount; + } hipMemcpyDtoA; struct { hipDeviceptr_t dst; hipDeviceptr_t src; @@ -2934,6 +2992,13 @@ typedef struct hip_api_data_s { const void* srcHost; size_t count; } hipMemcpyHtoA; + struct { + hipArray_t dstArray; + size_t dstOffset; + const void* srcHost; + size_t ByteCount; + hipStream_t stream; + } hipMemcpyHtoAAsync; struct { hipDeviceptr_t dst; void* src; @@ -3249,6 +3314,11 @@ typedef struct hip_api_data_s { struct { unsigned int flags; } hipSetDeviceFlags; + struct { + int* device_arr; + int device_arr__val; + int len; + } hipSetValidDevices; struct { const void* arg; size_t size; @@ -3858,6 +3928,11 @@ typedef struct hip_api_data_s { }; // hipDrvGraphAddMemFreeNode[('hipGraphNode_t*', 'phGraphNode'), ('hipGraph_t', 'hGraph'), ('const hipGraphNode_t*', 'dependencies'), ('size_t', 'numDependencies'), ('hipDeviceptr_t', 'dptr')] #define INIT_hipDrvGraphAddMemFreeNode_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipDrvGraphAddMemFreeNode.phGraphNode = (hipGraphNode_t*)phGraphNode; \ + cb_data.args.hipDrvGraphAddMemFreeNode.hGraph = (hipGraph_t)hGraph; \ + cb_data.args.hipDrvGraphAddMemFreeNode.dependencies = (const hipGraphNode_t*)dependencies; \ + cb_data.args.hipDrvGraphAddMemFreeNode.numDependencies = (size_t)numDependencies; \ + cb_data.args.hipDrvGraphAddMemFreeNode.dptr = (hipDeviceptr_t)dptr; \ }; // hipDrvGraphAddMemcpyNode[('hipGraphNode_t*', 'phGraphNode'), ('hipGraph_t', 'hGraph'), ('const hipGraphNode_t*', 'dependencies'), ('size_t', 'numDependencies'), ('const HIP_MEMCPY3D*', 'copyParams'), ('hipCtx_t', 'ctx')] #define INIT_hipDrvGraphAddMemcpyNode_CB_ARGS_DATA(cb_data) { \ @@ -3879,9 +3954,17 @@ typedef struct hip_api_data_s { }; // hipDrvGraphExecMemcpyNodeSetParams[('hipGraphExec_t', 'hGraphExec'), ('hipGraphNode_t', 'hNode'), ('const HIP_MEMCPY3D*', 'copyParams'), ('hipCtx_t', 'ctx')] #define INIT_hipDrvGraphExecMemcpyNodeSetParams_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipDrvGraphExecMemcpyNodeSetParams.hGraphExec = (hipGraphExec_t)hGraphExec; \ + cb_data.args.hipDrvGraphExecMemcpyNodeSetParams.hNode = (hipGraphNode_t)hNode; \ + cb_data.args.hipDrvGraphExecMemcpyNodeSetParams.copyParams = (const HIP_MEMCPY3D*)copyParams; \ + cb_data.args.hipDrvGraphExecMemcpyNodeSetParams.ctx = (hipCtx_t)ctx; \ }; // hipDrvGraphExecMemsetNodeSetParams[('hipGraphExec_t', 'hGraphExec'), ('hipGraphNode_t', 'hNode'), ('const HIP_MEMSET_NODE_PARAMS*', 'memsetParams'), ('hipCtx_t', 'ctx')] #define INIT_hipDrvGraphExecMemsetNodeSetParams_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipDrvGraphExecMemsetNodeSetParams.hGraphExec = (hipGraphExec_t)hGraphExec; \ + cb_data.args.hipDrvGraphExecMemsetNodeSetParams.hNode = (hipGraphNode_t)hNode; \ + cb_data.args.hipDrvGraphExecMemsetNodeSetParams.memsetParams = (const HIP_MEMSET_NODE_PARAMS*)memsetParams; \ + cb_data.args.hipDrvGraphExecMemsetNodeSetParams.ctx = (hipCtx_t)ctx; \ }; // hipDrvGraphMemcpyNodeGetParams[('hipGraphNode_t', 'hNode'), ('HIP_MEMCPY3D*', 'nodeParams')] #define INIT_hipDrvGraphMemcpyNodeGetParams_CB_ARGS_DATA(cb_data) { \ @@ -5098,6 +5181,18 @@ typedef struct hip_api_data_s { cb_data.args.hipMemcpy2D.height = (size_t)height; \ cb_data.args.hipMemcpy2D.kind = (hipMemcpyKind)kind; \ }; +// hipMemcpy2DArrayToArray[('hipArray_t', 'dst'), ('size_t', 'wOffsetDst'), ('size_t', 'hOffsetDst'), ('hipArray_const_t', 'src'), ('size_t', 'wOffsetSrc'), ('size_t', 'hOffsetSrc'), ('size_t', 'width'), ('size_t', 'height'), ('hipMemcpyKind', 'kind')] +#define INIT_hipMemcpy2DArrayToArray_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemcpy2DArrayToArray.dst = (hipArray_t)dst; \ + cb_data.args.hipMemcpy2DArrayToArray.wOffsetDst = (size_t)wOffsetDst; \ + cb_data.args.hipMemcpy2DArrayToArray.hOffsetDst = (size_t)hOffsetDst; \ + cb_data.args.hipMemcpy2DArrayToArray.src = (hipArray_const_t)src; \ + cb_data.args.hipMemcpy2DArrayToArray.wOffsetSrc = (size_t)wOffsetSrc; \ + cb_data.args.hipMemcpy2DArrayToArray.hOffsetSrc = (size_t)hOffsetSrc; \ + cb_data.args.hipMemcpy2DArrayToArray.width = (size_t)width; \ + cb_data.args.hipMemcpy2DArrayToArray.height = (size_t)height; \ + cb_data.args.hipMemcpy2DArrayToArray.kind = (hipMemcpyKind)kind; \ +}; // hipMemcpy2DAsync[('void*', 'dst'), ('size_t', 'dpitch'), ('const void*', 'src'), ('size_t', 'spitch'), ('size_t', 'width'), ('size_t', 'height'), ('hipMemcpyKind', 'kind'), ('hipStream_t', 'stream')] #define INIT_hipMemcpy2DAsync_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipMemcpy2DAsync.dst = (void*)dst; \ @@ -5172,6 +5267,21 @@ typedef struct hip_api_data_s { cb_data.args.hipMemcpyAsync.kind = (hipMemcpyKind)kind; \ cb_data.args.hipMemcpyAsync.stream = (hipStream_t)stream; \ }; +// hipMemcpyAtoA[('hipArray_t', 'dstArray'), ('size_t', 'dstOffset'), ('hipArray_t', 'srcArray'), ('size_t', 'srcOffset'), ('size_t', 'ByteCount')] +#define INIT_hipMemcpyAtoA_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemcpyAtoA.dstArray = (hipArray_t)dstArray; \ + cb_data.args.hipMemcpyAtoA.dstOffset = (size_t)dstOffset; \ + cb_data.args.hipMemcpyAtoA.srcArray = (hipArray_t)srcArray; \ + cb_data.args.hipMemcpyAtoA.srcOffset = (size_t)srcOffset; \ + cb_data.args.hipMemcpyAtoA.ByteCount = (size_t)ByteCount; \ +}; +// hipMemcpyAtoD[('hipDeviceptr_t', 'dstDevice'), ('hipArray_t', 'srcArray'), ('size_t', 'srcOffset'), ('size_t', 'ByteCount')] +#define INIT_hipMemcpyAtoD_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemcpyAtoD.dstDevice = (hipDeviceptr_t)dstDevice; \ + cb_data.args.hipMemcpyAtoD.srcArray = (hipArray_t)srcArray; \ + cb_data.args.hipMemcpyAtoD.srcOffset = (size_t)srcOffset; \ + cb_data.args.hipMemcpyAtoD.ByteCount = (size_t)ByteCount; \ +}; // hipMemcpyAtoH[('void*', 'dst'), ('hipArray_t', 'srcArray'), ('size_t', 'srcOffset'), ('size_t', 'count')] #define INIT_hipMemcpyAtoH_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipMemcpyAtoH.dst = (void*)dstHost; \ @@ -5179,6 +5289,21 @@ typedef struct hip_api_data_s { cb_data.args.hipMemcpyAtoH.srcOffset = (size_t)srcOffset; \ cb_data.args.hipMemcpyAtoH.count = (size_t)ByteCount; \ }; +// hipMemcpyAtoHAsync[('void*', 'dstHost'), ('hipArray_t', 'srcArray'), ('size_t', 'srcOffset'), ('size_t', 'ByteCount'), ('hipStream_t', 'stream')] +#define INIT_hipMemcpyAtoHAsync_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemcpyAtoHAsync.dstHost = (void*)dstHost; \ + cb_data.args.hipMemcpyAtoHAsync.srcArray = (hipArray_t)srcArray; \ + cb_data.args.hipMemcpyAtoHAsync.srcOffset = (size_t)srcOffset; \ + cb_data.args.hipMemcpyAtoHAsync.ByteCount = (size_t)ByteCount; \ + cb_data.args.hipMemcpyAtoHAsync.stream = (hipStream_t)stream; \ +}; +// hipMemcpyDtoA[('hipArray_t', 'dstArray'), ('size_t', 'dstOffset'), ('hipDeviceptr_t', 'srcDevice'), ('size_t', 'ByteCount')] +#define INIT_hipMemcpyDtoA_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemcpyDtoA.dstArray = (hipArray_t)dstArray; \ + cb_data.args.hipMemcpyDtoA.dstOffset = (size_t)dstOffset; \ + cb_data.args.hipMemcpyDtoA.srcDevice = (hipDeviceptr_t)srcDevice; \ + cb_data.args.hipMemcpyDtoA.ByteCount = (size_t)ByteCount; \ +}; // hipMemcpyDtoD[('hipDeviceptr_t', 'dst'), ('hipDeviceptr_t', 'src'), ('size_t', 'sizeBytes')] #define INIT_hipMemcpyDtoD_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipMemcpyDtoD.dst = (hipDeviceptr_t)dstDevice; \ @@ -5238,6 +5363,14 @@ typedef struct hip_api_data_s { cb_data.args.hipMemcpyHtoA.srcHost = (const void*)srcHost; \ cb_data.args.hipMemcpyHtoA.count = (size_t)ByteCount; \ }; +// hipMemcpyHtoAAsync[('hipArray_t', 'dstArray'), ('size_t', 'dstOffset'), ('const void*', 'srcHost'), ('size_t', 'ByteCount'), ('hipStream_t', 'stream')] +#define INIT_hipMemcpyHtoAAsync_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemcpyHtoAAsync.dstArray = (hipArray_t)dstArray; \ + cb_data.args.hipMemcpyHtoAAsync.dstOffset = (size_t)dstOffset; \ + cb_data.args.hipMemcpyHtoAAsync.srcHost = (const void*)srcHost; \ + cb_data.args.hipMemcpyHtoAAsync.ByteCount = (size_t)ByteCount; \ + cb_data.args.hipMemcpyHtoAAsync.stream = (hipStream_t)stream; \ +}; // hipMemcpyHtoD[('hipDeviceptr_t', 'dst'), ('void*', 'src'), ('size_t', 'sizeBytes')] #define INIT_hipMemcpyHtoD_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipMemcpyHtoD.dst = (hipDeviceptr_t)dstDevice; \ @@ -5576,6 +5709,11 @@ typedef struct hip_api_data_s { #define INIT_hipSetDeviceFlags_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipSetDeviceFlags.flags = (unsigned int)flags; \ }; +// hipSetValidDevices[('int*', 'device_arr'), ('int', 'len')] +#define INIT_hipSetValidDevices_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipSetValidDevices.device_arr = (int*)device_arr; \ + cb_data.args.hipSetValidDevices.len = (int)len; \ +}; // hipSetupArgument[('const void*', 'arg'), ('size_t', 'size'), ('size_t', 'offset')] #define INIT_hipSetupArgument_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipSetupArgument.arg = (const void*)arg; \ @@ -5888,20 +6026,6 @@ typedef struct hip_api_data_s { #define INIT_hipGetTextureObjectTextureDesc_CB_ARGS_DATA(cb_data) {}; // hipGetTextureReference() #define INIT_hipGetTextureReference_CB_ARGS_DATA(cb_data) {}; -// hipMemcpy2DArrayToArray() -#define INIT_hipMemcpy2DArrayToArray_CB_ARGS_DATA(cb_data) {}; -// hipMemcpyAtoA() -#define INIT_hipMemcpyAtoA_CB_ARGS_DATA(cb_data) {}; -// hipMemcpyAtoD() -#define INIT_hipMemcpyAtoD_CB_ARGS_DATA(cb_data) {}; -// hipMemcpyAtoHAsync() -#define INIT_hipMemcpyAtoHAsync_CB_ARGS_DATA(cb_data) {}; -// hipMemcpyDtoA() -#define INIT_hipMemcpyDtoA_CB_ARGS_DATA(cb_data) {}; -// hipMemcpyHtoAAsync() -#define INIT_hipMemcpyHtoAAsync_CB_ARGS_DATA(cb_data) {}; -// hipSetValidDevices() -#define INIT_hipSetValidDevices_CB_ARGS_DATA(cb_data) {}; // hipTexObjectCreate() #define INIT_hipTexObjectCreate_CB_ARGS_DATA(cb_data) {}; // hipTexObjectDestroy() @@ -7013,6 +7137,9 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipMemcpy2D[('void*', 'dst'), ('size_t', 'dpitch'), ('const void*', 'src'), ('size_t', 'spitch'), ('size_t', 'width'), ('size_t', 'height'), ('hipMemcpyKind', 'kind')] case HIP_API_ID_hipMemcpy2D: break; +// hipMemcpy2DArrayToArray[('hipArray_t', 'dst'), ('size_t', 'wOffsetDst'), ('size_t', 'hOffsetDst'), ('hipArray_const_t', 'src'), ('size_t', 'wOffsetSrc'), ('size_t', 'hOffsetSrc'), ('size_t', 'width'), ('size_t', 'height'), ('hipMemcpyKind', 'kind')] + case HIP_API_ID_hipMemcpy2DArrayToArray: + break; // hipMemcpy2DAsync[('void*', 'dst'), ('size_t', 'dpitch'), ('const void*', 'src'), ('size_t', 'spitch'), ('size_t', 'width'), ('size_t', 'height'), ('hipMemcpyKind', 'kind'), ('hipStream_t', 'stream')] case HIP_API_ID_hipMemcpy2DAsync: break; @@ -7039,9 +7166,21 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipMemcpyAsync[('void*', 'dst'), ('const void*', 'src'), ('size_t', 'sizeBytes'), ('hipMemcpyKind', 'kind'), ('hipStream_t', 'stream')] case HIP_API_ID_hipMemcpyAsync: break; +// hipMemcpyAtoA[('hipArray_t', 'dstArray'), ('size_t', 'dstOffset'), ('hipArray_t', 'srcArray'), ('size_t', 'srcOffset'), ('size_t', 'ByteCount')] + case HIP_API_ID_hipMemcpyAtoA: + break; +// hipMemcpyAtoD[('hipDeviceptr_t', 'dstDevice'), ('hipArray_t', 'srcArray'), ('size_t', 'srcOffset'), ('size_t', 'ByteCount')] + case HIP_API_ID_hipMemcpyAtoD: + break; // hipMemcpyAtoH[('void*', 'dst'), ('hipArray_t', 'srcArray'), ('size_t', 'srcOffset'), ('size_t', 'count')] case HIP_API_ID_hipMemcpyAtoH: break; +// hipMemcpyAtoHAsync[('void*', 'dstHost'), ('hipArray_t', 'srcArray'), ('size_t', 'srcOffset'), ('size_t', 'ByteCount'), ('hipStream_t', 'stream')] + case HIP_API_ID_hipMemcpyAtoHAsync: + break; +// hipMemcpyDtoA[('hipArray_t', 'dstArray'), ('size_t', 'dstOffset'), ('hipDeviceptr_t', 'srcDevice'), ('size_t', 'ByteCount')] + case HIP_API_ID_hipMemcpyDtoA: + break; // hipMemcpyDtoD[('hipDeviceptr_t', 'dst'), ('hipDeviceptr_t', 'src'), ('size_t', 'sizeBytes')] case HIP_API_ID_hipMemcpyDtoD: break; @@ -7066,6 +7205,9 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipMemcpyHtoA[('hipArray_t', 'dstArray'), ('size_t', 'dstOffset'), ('const void*', 'srcHost'), ('size_t', 'count')] case HIP_API_ID_hipMemcpyHtoA: break; +// hipMemcpyHtoAAsync[('hipArray_t', 'dstArray'), ('size_t', 'dstOffset'), ('const void*', 'srcHost'), ('size_t', 'ByteCount'), ('hipStream_t', 'stream')] + case HIP_API_ID_hipMemcpyHtoAAsync: + break; // hipMemcpyHtoD[('hipDeviceptr_t', 'dst'), ('void*', 'src'), ('size_t', 'sizeBytes')] case HIP_API_ID_hipMemcpyHtoD: break; @@ -7253,6 +7395,10 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipSetDeviceFlags[('unsigned int', 'flags')] case HIP_API_ID_hipSetDeviceFlags: break; +// hipSetValidDevices[('int*', 'device_arr'), ('int', 'len')] + case HIP_API_ID_hipSetValidDevices: + if (data->args.hipSetValidDevices.device_arr) data->args.hipSetValidDevices.device_arr__val = *(data->args.hipSetValidDevices.device_arr); + break; // hipSetupArgument[('const void*', 'arg'), ('size_t', 'size'), ('size_t', 'offset')] case HIP_API_ID_hipSetupArgument: break; @@ -9586,6 +9732,19 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << ", kind="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2D.kind); oss << ")"; break; + case HIP_API_ID_hipMemcpy2DArrayToArray: + oss << "hipMemcpy2DArrayToArray("; + oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DArrayToArray.dst); + oss << ", wOffsetDst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DArrayToArray.wOffsetDst); + oss << ", hOffsetDst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DArrayToArray.hOffsetDst); + oss << ", src="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DArrayToArray.src); + oss << ", wOffsetSrc="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DArrayToArray.wOffsetSrc); + oss << ", hOffsetSrc="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DArrayToArray.hOffsetSrc); + oss << ", width="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DArrayToArray.width); + oss << ", height="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DArrayToArray.height); + oss << ", kind="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DArrayToArray.kind); + oss << ")"; + break; case HIP_API_ID_hipMemcpy2DAsync: oss << "hipMemcpy2DAsync("; oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DAsync.dst); @@ -9670,6 +9829,23 @@ 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.hipMemcpyAsync.stream); oss << ")"; break; + case HIP_API_ID_hipMemcpyAtoA: + oss << "hipMemcpyAtoA("; + oss << "dstArray="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoA.dstArray); + oss << ", dstOffset="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoA.dstOffset); + oss << ", srcArray="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoA.srcArray); + oss << ", srcOffset="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoA.srcOffset); + oss << ", ByteCount="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoA.ByteCount); + oss << ")"; + break; + case HIP_API_ID_hipMemcpyAtoD: + oss << "hipMemcpyAtoD("; + oss << "dstDevice="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoD.dstDevice); + oss << ", srcArray="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoD.srcArray); + oss << ", srcOffset="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoD.srcOffset); + oss << ", ByteCount="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoD.ByteCount); + oss << ")"; + break; case HIP_API_ID_hipMemcpyAtoH: oss << "hipMemcpyAtoH("; oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoH.dst); @@ -9678,6 +9854,23 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << ", count="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoH.count); oss << ")"; break; + case HIP_API_ID_hipMemcpyAtoHAsync: + oss << "hipMemcpyAtoHAsync("; + oss << "dstHost="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoHAsync.dstHost); + oss << ", srcArray="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoHAsync.srcArray); + oss << ", srcOffset="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoHAsync.srcOffset); + oss << ", ByteCount="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoHAsync.ByteCount); + oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoHAsync.stream); + oss << ")"; + break; + case HIP_API_ID_hipMemcpyDtoA: + oss << "hipMemcpyDtoA("; + oss << "dstArray="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyDtoA.dstArray); + oss << ", dstOffset="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyDtoA.dstOffset); + oss << ", srcDevice="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyDtoA.srcDevice); + oss << ", ByteCount="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyDtoA.ByteCount); + oss << ")"; + break; case HIP_API_ID_hipMemcpyDtoD: oss << "hipMemcpyDtoD("; oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyDtoD.dst); @@ -9745,6 +9938,15 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << ", count="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyHtoA.count); oss << ")"; break; + case HIP_API_ID_hipMemcpyHtoAAsync: + oss << "hipMemcpyHtoAAsync("; + oss << "dstArray="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyHtoAAsync.dstArray); + oss << ", dstOffset="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyHtoAAsync.dstOffset); + oss << ", srcHost="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyHtoAAsync.srcHost); + oss << ", ByteCount="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyHtoAAsync.ByteCount); + oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyHtoAAsync.stream); + oss << ")"; + break; case HIP_API_ID_hipMemcpyHtoD: oss << "hipMemcpyHtoD("; oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyHtoD.dst); @@ -10168,6 +10370,13 @@ 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.hipSetDeviceFlags.flags); oss << ")"; break; + case HIP_API_ID_hipSetValidDevices: + oss << "hipSetValidDevices("; + if (data->args.hipSetValidDevices.device_arr == NULL) oss << "device_arr=NULL"; + else { oss << "device_arr="; roctracer::hip_support::detail::operator<<(oss, data->args.hipSetValidDevices.device_arr__val); } + oss << ", len="; roctracer::hip_support::detail::operator<<(oss, data->args.hipSetValidDevices.len); + oss << ")"; + break; case HIP_API_ID_hipSetupArgument: oss << "hipSetupArgument("; oss << "arg="; roctracer::hip_support::detail::operator<<(oss, data->args.hipSetupArgument.arg); diff --git a/hipamd/src/amdhip.def b/hipamd/src/amdhip.def index c1017af7af..5ad7942013 100644 --- a/hipamd/src/amdhip.def +++ b/hipamd/src/amdhip.def @@ -466,3 +466,10 @@ hipGetFuncBySymbol hipDrvGraphAddMemFreeNode hipDrvGraphExecMemcpyNodeSetParams hipDrvGraphExecMemsetNodeSetParams +hipSetValidDevices +hipMemcpyAtoD +hipMemcpyDtoA +hipMemcpyAtoA +hipMemcpyAtoHAsync +hipMemcpyHtoAAsync +hipMemcpy2DArrayToArray diff --git a/hipamd/src/hip_api_trace.cpp b/hipamd/src/hip_api_trace.cpp index 05598dd7f1..551660041a 100644 --- a/hipamd/src/hip_api_trace.cpp +++ b/hipamd/src/hip_api_trace.cpp @@ -775,6 +775,20 @@ hipError_t hipDrvGraphExecMemcpyNodeSetParams(hipGraphExec_t hGraphExec, hipGrap const HIP_MEMCPY3D* copyParams, hipCtx_t ctx); hipError_t hipDrvGraphExecMemsetNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t hNode, const HIP_MEMSET_NODE_PARAMS* memsetParams, hipCtx_t ctx); +hipError_t hipSetValidDevices(int* device_arr, int len); +hipError_t hipMemcpyAtoD(hipDeviceptr_t dstDevice, hipArray_t srcArray, size_t srcOffset, + size_t ByteCount); +hipError_t hipMemcpyDtoA(hipArray_t dstArray, size_t dstOffset, hipDeviceptr_t srcDevice, + size_t ByteCount); +hipError_t hipMemcpyAtoA(hipArray_t dstArray, size_t dstOffset, hipArray_t srcArray, + size_t srcOffset, size_t ByteCount); +hipError_t hipMemcpyAtoHAsync(void* dstHost, hipArray_t srcArray, size_t srcOffset, + size_t ByteCount, hipStream_t stream); +hipError_t hipMemcpyHtoAAsync(hipArray_t dstArray, size_t dstOffset, const void* srcHost, + size_t ByteCount, hipStream_t stream); +hipError_t hipMemcpy2DArrayToArray(hipArray_t dst, size_t wOffsetDst, size_t hOffsetDst, + hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, + size_t width, size_t height, hipMemcpyKind kind); } // namespace hip namespace hip { @@ -1254,6 +1268,13 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) { ptrDispatchTable->hipDrvGraphAddMemFreeNode_fn = hip::hipDrvGraphAddMemFreeNode; ptrDispatchTable->hipDrvGraphExecMemcpyNodeSetParams_fn = hip::hipDrvGraphExecMemcpyNodeSetParams; ptrDispatchTable->hipDrvGraphExecMemsetNodeSetParams_fn = hip::hipDrvGraphExecMemsetNodeSetParams; + ptrDispatchTable->hipSetValidDevices_fn = hip::hipSetValidDevices; + ptrDispatchTable->hipMemcpyAtoD_fn = hip::hipMemcpyAtoD; + ptrDispatchTable->hipMemcpyDtoA_fn = hip::hipMemcpyDtoA; + ptrDispatchTable->hipMemcpyAtoA_fn = hip::hipMemcpyAtoA; + ptrDispatchTable->hipMemcpyAtoHAsync_fn = hip::hipMemcpyAtoHAsync; + ptrDispatchTable->hipMemcpyHtoAAsync_fn = hip::hipMemcpyHtoAAsync; + ptrDispatchTable->hipMemcpy2DArrayToArray_fn = hip::hipMemcpy2DArrayToArray; } #if HIP_ROCPROFILER_REGISTER > 0 @@ -1800,25 +1821,33 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipStreamGetCaptureInfo_v2_spt_fn, 425) HIP_ENFORCE_ABI(HipDispatchTable, hipLaunchHostFunc_spt_fn, 426) HIP_ENFORCE_ABI(HipDispatchTable, hipGetStreamDeviceId_fn, 427) HIP_ENFORCE_ABI(HipDispatchTable, hipDrvGraphAddMemsetNode_fn, 428) -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphAddExternalSemaphoresWaitNode_fn, 429); -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphAddExternalSemaphoresSignalNode_fn, 430); -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExternalSemaphoresSignalNodeSetParams_fn, 431); -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExternalSemaphoresWaitNodeSetParams_fn, 432); -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExternalSemaphoresSignalNodeGetParams_fn, 433); -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExternalSemaphoresWaitNodeGetParams_fn, 434); -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExecExternalSemaphoresSignalNodeSetParams_fn, 435); -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExecExternalSemaphoresWaitNodeSetParams_fn, 436); -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphAddNode_fn, 437); -HIP_ENFORCE_ABI(HipDispatchTable, hipGraphInstantiateWithParams_fn, 438); +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphAddExternalSemaphoresWaitNode_fn, 429) +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphAddExternalSemaphoresSignalNode_fn, 430) +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExternalSemaphoresSignalNodeSetParams_fn, 431) +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExternalSemaphoresWaitNodeSetParams_fn, 432) +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExternalSemaphoresSignalNodeGetParams_fn, 433) +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExternalSemaphoresWaitNodeGetParams_fn, 434) +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExecExternalSemaphoresSignalNodeSetParams_fn, 435) +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExecExternalSemaphoresWaitNodeSetParams_fn, 436) +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphAddNode_fn, 437) +HIP_ENFORCE_ABI(HipDispatchTable, hipGraphInstantiateWithParams_fn, 438) HIP_ENFORCE_ABI(HipDispatchTable, hipExtGetLastError_fn, 439) HIP_ENFORCE_ABI(HipDispatchTable, hipTexRefGetBorderColor_fn, 440) HIP_ENFORCE_ABI(HipDispatchTable, hipTexRefGetArray_fn, 441) HIP_ENFORCE_ABI(HipDispatchTable, hipGetProcAddress_fn, 442) -HIP_ENFORCE_ABI(HipDispatchTable, hipStreamBeginCaptureToGraph_fn, 443); -HIP_ENFORCE_ABI(HipDispatchTable, hipGetFuncBySymbol_fn, 444); +HIP_ENFORCE_ABI(HipDispatchTable, hipStreamBeginCaptureToGraph_fn, 443) +HIP_ENFORCE_ABI(HipDispatchTable, hipGetFuncBySymbol_fn, 444) HIP_ENFORCE_ABI(HipDispatchTable, hipDrvGraphAddMemFreeNode_fn, 445) HIP_ENFORCE_ABI(HipDispatchTable, hipDrvGraphExecMemcpyNodeSetParams_fn, 446) HIP_ENFORCE_ABI(HipDispatchTable, hipDrvGraphExecMemsetNodeSetParams_fn, 447) +HIP_ENFORCE_ABI(HipDispatchTable, hipSetValidDevices_fn, 448) +HIP_ENFORCE_ABI(HipDispatchTable, hipMemcpyAtoD_fn, 449) +HIP_ENFORCE_ABI(HipDispatchTable, hipMemcpyDtoA_fn, 450) +HIP_ENFORCE_ABI(HipDispatchTable, hipMemcpyAtoA_fn, 451) +HIP_ENFORCE_ABI(HipDispatchTable, hipMemcpyAtoHAsync_fn, 452) +HIP_ENFORCE_ABI(HipDispatchTable, hipMemcpyHtoAAsync_fn, 453) +HIP_ENFORCE_ABI(HipDispatchTable, hipMemcpy2DArrayToArray_fn, 454) + // 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.: @@ -1826,7 +1855,7 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipDrvGraphExecMemsetNodeSetParams_fn, 447) // HIP_ENFORCE_ABI(, , 8) // // HIP_ENFORCE_ABI_VERSIONING(
, 9) <- 8 + 1 = 9 -HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 448) +HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 455) static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 3, "If you get this error, add new HIP_ENFORCE_ABI(...) code for the new function " diff --git a/hipamd/src/hip_hcc.map.in b/hipamd/src/hip_hcc.map.in index 2ec315cba0..4efc55b23f 100644 --- a/hipamd/src/hip_hcc.map.in +++ b/hipamd/src/hip_hcc.map.in @@ -563,6 +563,13 @@ global: hipDrvGraphExecMemcpyNodeSetParams; hipDrvGraphExecMemsetNodeSetParams; hipDrvGraphAddMemFreeNode; + hipSetValidDevices; + hipMemcpyAtoD; + hipMemcpyDtoA; + hipMemcpyAtoA; + hipMemcpyAtoHAsync; + hipMemcpyHtoAAsync; + hipMemcpy2DArrayToArray; local: *; } hip_6.1; diff --git a/hipamd/src/hip_table_interface.cpp b/hipamd/src/hip_table_interface.cpp index 0b92c8696c..201790af43 100644 --- a/hipamd/src/hip_table_interface.cpp +++ b/hipamd/src/hip_table_interface.cpp @@ -1760,3 +1760,36 @@ hipError_t hipDrvGraphExecMemcpyNodeSetParams(hipGraphExec_t hGraphExec, hipGrap return hip::GetHipDispatchTable()->hipDrvGraphExecMemcpyNodeSetParams_fn(hGraphExec, hNode, copyParams, ctx); } +hipError_t hipSetValidDevices(int* device_arr, int len) { + return hip::GetHipDispatchTable()->hipSetValidDevices_fn(device_arr, len); +} +hipError_t hipMemcpyAtoD(hipDeviceptr_t dstDevice, hipArray_t srcArray, size_t srcOffset, + size_t ByteCount) { + return hip::GetHipDispatchTable()->hipMemcpyAtoD_fn(dstDevice, srcArray, srcOffset, + ByteCount); +} +hipError_t hipMemcpyDtoA(hipArray_t dstArray, size_t dstOffset, hipDeviceptr_t srcDevice, + size_t ByteCount) { + return hip::GetHipDispatchTable()->hipMemcpyDtoA_fn(dstArray, dstOffset, srcDevice, ByteCount); +} +hipError_t hipMemcpyAtoA(hipArray_t dstArray, size_t dstOffset, hipArray_t srcArray, + size_t srcOffset, size_t ByteCount) { + return hip::GetHipDispatchTable()->hipMemcpyAtoA_fn(dstArray, dstOffset, srcArray, srcOffset, + ByteCount); +} +hipError_t hipMemcpyAtoHAsync(void* dstHost, hipArray_t srcArray, size_t srcOffset, + size_t ByteCount, hipStream_t stream) { + return hip::GetHipDispatchTable()->hipMemcpyAtoHAsync_fn(dstHost, srcArray, srcOffset, ByteCount, + stream); +} +hipError_t hipMemcpyHtoAAsync(hipArray_t dstArray, size_t dstOffset, const void* srcHost, + size_t ByteCount, hipStream_t stream) { + return hip::GetHipDispatchTable()->hipMemcpyHtoAAsync_fn(dstArray, dstOffset, srcHost, ByteCount, + stream); +} +hipError_t hipMemcpy2DArrayToArray(hipArray_t dst, size_t wOffsetDst, size_t hOffsetDst, + hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, + size_t width, size_t height, hipMemcpyKind kind) { + return hip::GetHipDispatchTable()->hipMemcpy2DArrayToArray_fn( + dst, wOffsetDst, hOffsetDst, src, wOffsetSrc, hOffsetSrc, width, height, kind); +}