From e1889b77b4091cab754d8a279762df1c7f976047 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 16 Nov 2023 00:20:54 -0500 Subject: [PATCH] SWDEV-427855 - Revert "SWDEV-427855 - hipamd change for profiler and TF fix" This reverts commit 7478e9072799acb35356aa720cf82c16569dc648. Reason for revert: Change-Id: I856b7ea30744f1b7bb099b6adbce2155201be539 [ROCm/clr commit: 57cb8400583baa3aa2d30b53372894b3683c6474] --- .../include/hip/amd_detail/hip_prof_str.h | 537 ++++-------------- projects/clr/hipamd/src/CMakeLists.txt | 26 +- projects/clr/hipamd/src/amdhip.def | 3 - projects/clr/hipamd/src/hip_device.cpp | 112 +++- .../clr/hipamd/src/hip_device_deprecated.cpp | 83 +++ .../clr/hipamd/src/hip_device_runtime.cpp | 46 +- projects/clr/hipamd/src/hip_find_defs.py | 83 +++ projects/clr/hipamd/src/hip_gl.cpp | 1 - projects/clr/hipamd/src/hip_hcc.map.in | 5 +- projects/clr/hipamd/src/hip_prof_gen.py | 23 +- 10 files changed, 381 insertions(+), 538 deletions(-) create mode 100644 projects/clr/hipamd/src/hip_device_deprecated.cpp create mode 100755 projects/clr/hipamd/src/hip_find_defs.py 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 128dc3747e..c67a8886e0 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 @@ -7,12 +7,7 @@ #define _HIP_PROF_STR_H #define HIP_PROF_VER 1 -#include -#include -#include "amd_hip_gl_interop.h" - -#define HIP_API_ID_CONCAT_HELPER(a,b) a##b -#define HIP_API_ID_CONCAT(a,b) HIP_API_ID_CONCAT_HELPER(a,b) +#include // HIP API callbacks ID enumeration enum hip_api_id_t { @@ -23,7 +18,7 @@ enum hip_api_id_t { HIP_API_ID_hipArray3DCreate = 3, HIP_API_ID_hipArrayCreate = 4, HIP_API_ID_hipArrayDestroy = 5, - HIP_API_ID_hipChooseDeviceR0000 = 6, + HIP_API_ID_hipChooseDevice = 6, HIP_API_ID_hipConfigureCall = 7, HIP_API_ID_hipCtxCreate = 8, HIP_API_ID_hipCtxDestroy = 9, @@ -98,7 +93,7 @@ enum hip_api_id_t { HIP_API_ID_hipGetDevice = 78, HIP_API_ID_hipGetDeviceCount = 79, HIP_API_ID_hipGetDeviceFlags = 80, - HIP_API_ID_hipGetDevicePropertiesR0000 = 81, + HIP_API_ID_hipGetDevicePropertiesR0600 = 81, HIP_API_ID_RESERVED_82 = 82, HIP_API_ID_hipGetErrorString = 83, HIP_API_ID_hipGetLastError = 84, @@ -382,24 +377,8 @@ enum hip_api_id_t { HIP_API_ID_hipArrayGetInfo = 362, HIP_API_ID_hipStreamGetDevice = 363, HIP_API_ID_hipExternalMemoryGetMappedMipmappedArray = 364, - HIP_API_ID_hipChooseDeviceR0600 = 365, - HIP_API_ID_hipDrvGraphAddMemcpyNode = 366, - HIP_API_ID_hipDrvGraphAddMemsetNode = 367, - HIP_API_ID_hipDrvGraphMemcpyNodeGetParams = 368, - HIP_API_ID_hipDrvGraphMemcpyNodeSetParams = 369, - HIP_API_ID_hipGetDevicePropertiesR0600 = 370, - HIP_API_ID_hipGraphAddExternalSemaphoresSignalNode = 371, - HIP_API_ID_hipGraphAddExternalSemaphoresWaitNode = 372, - HIP_API_ID_hipGraphExecExternalSemaphoresSignalNodeSetParams = 373, - HIP_API_ID_hipGraphExecExternalSemaphoresWaitNodeSetParams = 374, - HIP_API_ID_hipGraphExternalSemaphoresSignalNodeGetParams = 375, - HIP_API_ID_hipGraphExternalSemaphoresSignalNodeSetParams = 376, - HIP_API_ID_hipGraphExternalSemaphoresWaitNodeGetParams = 377, - HIP_API_ID_hipGraphExternalSemaphoresWaitNodeSetParams = 378, - HIP_API_ID_LAST = 378, - - HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice), - HIP_API_ID_hipGetDeviceProperties = HIP_API_ID_CONCAT(HIP_API_ID_,hipGetDeviceProperties), + HIP_API_ID_hipDrvGraphAddMemcpyNode = 365, + HIP_API_ID_LAST = 365, HIP_API_ID_hipBindTexture = HIP_API_ID_NONE, HIP_API_ID_hipBindTexture2D = HIP_API_ID_NONE, @@ -414,11 +393,14 @@ enum hip_api_id_t { 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_hipMemcpyArrayToArray = 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_hipMemcpyFromArrayAsync = HIP_API_ID_NONE, HIP_API_ID_hipMemcpyHtoAAsync = HIP_API_ID_NONE, + HIP_API_ID_hipMemcpyToArrayAsync = 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, @@ -437,9 +419,6 @@ enum hip_api_id_t { HIP_API_ID_hipUnbindTexture = HIP_API_ID_NONE, }; -#undef HIP_API_ID_CONCAT_HELPER -#undef HIP_API_ID_CONCAT - // Return the HIP API string for a given callback ID static inline const char* hip_api_name(const uint32_t id) { switch(id) { @@ -451,8 +430,7 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipArrayDestroy: return "hipArrayDestroy"; case HIP_API_ID_hipArrayGetDescriptor: return "hipArrayGetDescriptor"; case HIP_API_ID_hipArrayGetInfo: return "hipArrayGetInfo"; - case HIP_API_ID_hipChooseDeviceR0000: return "hipChooseDeviceR0000"; - case HIP_API_ID_hipChooseDeviceR0600: return "hipChooseDeviceR0600"; + case HIP_API_ID_hipChooseDevice: return "hipChooseDevice"; case HIP_API_ID_hipConfigureCall: return "hipConfigureCall"; case HIP_API_ID_hipCreateSurfaceObject: return "hipCreateSurfaceObject"; case HIP_API_ID_hipCtxCreate: return "hipCtxCreate"; @@ -507,10 +485,6 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipDeviceSynchronize: return "hipDeviceSynchronize"; case HIP_API_ID_hipDeviceTotalMem: return "hipDeviceTotalMem"; case HIP_API_ID_hipDriverGetVersion: return "hipDriverGetVersion"; - case HIP_API_ID_hipDrvGraphAddMemcpyNode: return "hipDrvGraphAddMemcpyNode"; - case HIP_API_ID_hipDrvGraphAddMemsetNode: return "hipDrvGraphAddMemsetNode"; - case HIP_API_ID_hipDrvGraphMemcpyNodeGetParams: return "hipDrvGraphMemcpyNodeGetParams"; - case HIP_API_ID_hipDrvGraphMemcpyNodeSetParams: return "hipDrvGraphMemcpyNodeSetParams"; case HIP_API_ID_hipDrvMemcpy2DUnaligned: return "hipDrvMemcpy2DUnaligned"; case HIP_API_ID_hipDrvMemcpy3D: return "hipDrvMemcpy3D"; case HIP_API_ID_hipDrvMemcpy3DAsync: return "hipDrvMemcpy3DAsync"; @@ -530,7 +504,6 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipExtStreamCreateWithCUMask: return "hipExtStreamCreateWithCUMask"; case HIP_API_ID_hipExtStreamGetCUMask: return "hipExtStreamGetCUMask"; case HIP_API_ID_hipExternalMemoryGetMappedBuffer: return "hipExternalMemoryGetMappedBuffer"; - case HIP_API_ID_hipExternalMemoryGetMappedMipmappedArray: return "hipExternalMemoryGetMappedMipmappedArray"; case HIP_API_ID_hipFree: return "hipFree"; case HIP_API_ID_hipFreeArray: return "hipFreeArray"; case HIP_API_ID_hipFreeAsync: return "hipFreeAsync"; @@ -546,7 +519,6 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipGetDevice: return "hipGetDevice"; case HIP_API_ID_hipGetDeviceCount: return "hipGetDeviceCount"; case HIP_API_ID_hipGetDeviceFlags: return "hipGetDeviceFlags"; - case HIP_API_ID_hipGetDevicePropertiesR0000: return "hipGetDevicePropertiesR0000"; case HIP_API_ID_hipGetDevicePropertiesR0600: return "hipGetDevicePropertiesR0600"; case HIP_API_ID_hipGetErrorString: return "hipGetErrorString"; case HIP_API_ID_hipGetLastError: return "hipGetLastError"; @@ -558,13 +530,12 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipGraphAddEmptyNode: return "hipGraphAddEmptyNode"; case HIP_API_ID_hipGraphAddEventRecordNode: return "hipGraphAddEventRecordNode"; case HIP_API_ID_hipGraphAddEventWaitNode: return "hipGraphAddEventWaitNode"; - case HIP_API_ID_hipGraphAddExternalSemaphoresSignalNode: return "hipGraphAddExternalSemaphoresSignalNode"; - case HIP_API_ID_hipGraphAddExternalSemaphoresWaitNode: return "hipGraphAddExternalSemaphoresWaitNode"; case HIP_API_ID_hipGraphAddHostNode: return "hipGraphAddHostNode"; case HIP_API_ID_hipGraphAddKernelNode: return "hipGraphAddKernelNode"; case HIP_API_ID_hipGraphAddMemAllocNode: return "hipGraphAddMemAllocNode"; case HIP_API_ID_hipGraphAddMemFreeNode: return "hipGraphAddMemFreeNode"; case HIP_API_ID_hipGraphAddMemcpyNode: return "hipGraphAddMemcpyNode"; + case HIP_API_ID_hipDrvGraphAddMemcpyNode: return "hipDrvGraphAddMemcpyNode"; case HIP_API_ID_hipGraphAddMemcpyNode1D: return "hipGraphAddMemcpyNode1D"; case HIP_API_ID_hipGraphAddMemcpyNodeFromSymbol: return "hipGraphAddMemcpyNodeFromSymbol"; case HIP_API_ID_hipGraphAddMemcpyNodeToSymbol: return "hipGraphAddMemcpyNodeToSymbol"; @@ -583,8 +554,6 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipGraphExecDestroy: return "hipGraphExecDestroy"; case HIP_API_ID_hipGraphExecEventRecordNodeSetEvent: return "hipGraphExecEventRecordNodeSetEvent"; case HIP_API_ID_hipGraphExecEventWaitNodeSetEvent: return "hipGraphExecEventWaitNodeSetEvent"; - case HIP_API_ID_hipGraphExecExternalSemaphoresSignalNodeSetParams: return "hipGraphExecExternalSemaphoresSignalNodeSetParams"; - case HIP_API_ID_hipGraphExecExternalSemaphoresWaitNodeSetParams: return "hipGraphExecExternalSemaphoresWaitNodeSetParams"; case HIP_API_ID_hipGraphExecHostNodeSetParams: return "hipGraphExecHostNodeSetParams"; case HIP_API_ID_hipGraphExecKernelNodeSetParams: return "hipGraphExecKernelNodeSetParams"; case HIP_API_ID_hipGraphExecMemcpyNodeSetParams: return "hipGraphExecMemcpyNodeSetParams"; @@ -593,10 +562,6 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipGraphExecMemcpyNodeSetParamsToSymbol: return "hipGraphExecMemcpyNodeSetParamsToSymbol"; case HIP_API_ID_hipGraphExecMemsetNodeSetParams: return "hipGraphExecMemsetNodeSetParams"; case HIP_API_ID_hipGraphExecUpdate: return "hipGraphExecUpdate"; - case HIP_API_ID_hipGraphExternalSemaphoresSignalNodeGetParams: return "hipGraphExternalSemaphoresSignalNodeGetParams"; - case HIP_API_ID_hipGraphExternalSemaphoresSignalNodeSetParams: return "hipGraphExternalSemaphoresSignalNodeSetParams"; - case HIP_API_ID_hipGraphExternalSemaphoresWaitNodeGetParams: return "hipGraphExternalSemaphoresWaitNodeGetParams"; - case HIP_API_ID_hipGraphExternalSemaphoresWaitNodeSetParams: return "hipGraphExternalSemaphoresWaitNodeSetParams"; case HIP_API_ID_hipGraphGetEdges: return "hipGraphGetEdges"; case HIP_API_ID_hipGraphGetNodes: return "hipGraphGetNodes"; case HIP_API_ID_hipGraphGetRootNodes: return "hipGraphGetRootNodes"; @@ -817,6 +782,7 @@ 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_hipExternalMemoryGetMappedMipmappedArray: return "hipExternalMemoryGetMappedMipmappedArray"; }; return "unknown"; }; @@ -832,8 +798,7 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipArrayDestroy", name) == 0) return HIP_API_ID_hipArrayDestroy; if (strcmp("hipArrayGetDescriptor", name) == 0) return HIP_API_ID_hipArrayGetDescriptor; if (strcmp("hipArrayGetInfo", name) == 0) return HIP_API_ID_hipArrayGetInfo; - if (strcmp("hipChooseDeviceR0000", name) == 0) return HIP_API_ID_hipChooseDeviceR0000; - if (strcmp("hipChooseDeviceR0600", name) == 0) return HIP_API_ID_hipChooseDeviceR0600; + if (strcmp("hipChooseDevice", name) == 0) return HIP_API_ID_hipChooseDevice; if (strcmp("hipConfigureCall", name) == 0) return HIP_API_ID_hipConfigureCall; if (strcmp("hipCreateSurfaceObject", name) == 0) return HIP_API_ID_hipCreateSurfaceObject; if (strcmp("hipCtxCreate", name) == 0) return HIP_API_ID_hipCtxCreate; @@ -888,10 +853,6 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipDeviceSynchronize", name) == 0) return HIP_API_ID_hipDeviceSynchronize; if (strcmp("hipDeviceTotalMem", name) == 0) return HIP_API_ID_hipDeviceTotalMem; if (strcmp("hipDriverGetVersion", name) == 0) return HIP_API_ID_hipDriverGetVersion; - if (strcmp("hipDrvGraphAddMemcpyNode", name) == 0) return HIP_API_ID_hipDrvGraphAddMemcpyNode; - if (strcmp("hipDrvGraphAddMemsetNode", name) == 0) return HIP_API_ID_hipDrvGraphAddMemsetNode; - if (strcmp("hipDrvGraphMemcpyNodeGetParams", name) == 0) return HIP_API_ID_hipDrvGraphMemcpyNodeGetParams; - if (strcmp("hipDrvGraphMemcpyNodeSetParams", name) == 0) return HIP_API_ID_hipDrvGraphMemcpyNodeSetParams; if (strcmp("hipDrvMemcpy2DUnaligned", name) == 0) return HIP_API_ID_hipDrvMemcpy2DUnaligned; if (strcmp("hipDrvMemcpy3D", name) == 0) return HIP_API_ID_hipDrvMemcpy3D; if (strcmp("hipDrvMemcpy3DAsync", name) == 0) return HIP_API_ID_hipDrvMemcpy3DAsync; @@ -911,7 +872,6 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipExtStreamCreateWithCUMask", name) == 0) return HIP_API_ID_hipExtStreamCreateWithCUMask; if (strcmp("hipExtStreamGetCUMask", name) == 0) return HIP_API_ID_hipExtStreamGetCUMask; if (strcmp("hipExternalMemoryGetMappedBuffer", name) == 0) return HIP_API_ID_hipExternalMemoryGetMappedBuffer; - if (strcmp("hipExternalMemoryGetMappedMipmappedArray", name) == 0) return HIP_API_ID_hipExternalMemoryGetMappedMipmappedArray; if (strcmp("hipFree", name) == 0) return HIP_API_ID_hipFree; if (strcmp("hipFreeArray", name) == 0) return HIP_API_ID_hipFreeArray; if (strcmp("hipFreeAsync", name) == 0) return HIP_API_ID_hipFreeAsync; @@ -927,7 +887,6 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipGetDevice", name) == 0) return HIP_API_ID_hipGetDevice; if (strcmp("hipGetDeviceCount", name) == 0) return HIP_API_ID_hipGetDeviceCount; if (strcmp("hipGetDeviceFlags", name) == 0) return HIP_API_ID_hipGetDeviceFlags; - if (strcmp("hipGetDevicePropertiesR0000", name) == 0) return HIP_API_ID_hipGetDevicePropertiesR0000; if (strcmp("hipGetDevicePropertiesR0600", name) == 0) return HIP_API_ID_hipGetDevicePropertiesR0600; if (strcmp("hipGetErrorString", name) == 0) return HIP_API_ID_hipGetErrorString; if (strcmp("hipGetLastError", name) == 0) return HIP_API_ID_hipGetLastError; @@ -939,8 +898,6 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipGraphAddEmptyNode", name) == 0) return HIP_API_ID_hipGraphAddEmptyNode; if (strcmp("hipGraphAddEventRecordNode", name) == 0) return HIP_API_ID_hipGraphAddEventRecordNode; if (strcmp("hipGraphAddEventWaitNode", name) == 0) return HIP_API_ID_hipGraphAddEventWaitNode; - if (strcmp("hipGraphAddExternalSemaphoresSignalNode", name) == 0) return HIP_API_ID_hipGraphAddExternalSemaphoresSignalNode; - if (strcmp("hipGraphAddExternalSemaphoresWaitNode", name) == 0) return HIP_API_ID_hipGraphAddExternalSemaphoresWaitNode; if (strcmp("hipGraphAddHostNode", name) == 0) return HIP_API_ID_hipGraphAddHostNode; if (strcmp("hipGraphAddKernelNode", name) == 0) return HIP_API_ID_hipGraphAddKernelNode; if (strcmp("hipGraphAddMemAllocNode", name) == 0) return HIP_API_ID_hipGraphAddMemAllocNode; @@ -964,8 +921,6 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipGraphExecDestroy", name) == 0) return HIP_API_ID_hipGraphExecDestroy; if (strcmp("hipGraphExecEventRecordNodeSetEvent", name) == 0) return HIP_API_ID_hipGraphExecEventRecordNodeSetEvent; if (strcmp("hipGraphExecEventWaitNodeSetEvent", name) == 0) return HIP_API_ID_hipGraphExecEventWaitNodeSetEvent; - if (strcmp("hipGraphExecExternalSemaphoresSignalNodeSetParams", name) == 0) return HIP_API_ID_hipGraphExecExternalSemaphoresSignalNodeSetParams; - if (strcmp("hipGraphExecExternalSemaphoresWaitNodeSetParams", name) == 0) return HIP_API_ID_hipGraphExecExternalSemaphoresWaitNodeSetParams; if (strcmp("hipGraphExecHostNodeSetParams", name) == 0) return HIP_API_ID_hipGraphExecHostNodeSetParams; if (strcmp("hipGraphExecKernelNodeSetParams", name) == 0) return HIP_API_ID_hipGraphExecKernelNodeSetParams; if (strcmp("hipGraphExecMemcpyNodeSetParams", name) == 0) return HIP_API_ID_hipGraphExecMemcpyNodeSetParams; @@ -974,10 +929,6 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipGraphExecMemcpyNodeSetParamsToSymbol", name) == 0) return HIP_API_ID_hipGraphExecMemcpyNodeSetParamsToSymbol; if (strcmp("hipGraphExecMemsetNodeSetParams", name) == 0) return HIP_API_ID_hipGraphExecMemsetNodeSetParams; if (strcmp("hipGraphExecUpdate", name) == 0) return HIP_API_ID_hipGraphExecUpdate; - if (strcmp("hipGraphExternalSemaphoresSignalNodeGetParams", name) == 0) return HIP_API_ID_hipGraphExternalSemaphoresSignalNodeGetParams; - if (strcmp("hipGraphExternalSemaphoresSignalNodeSetParams", name) == 0) return HIP_API_ID_hipGraphExternalSemaphoresSignalNodeSetParams; - if (strcmp("hipGraphExternalSemaphoresWaitNodeGetParams", name) == 0) return HIP_API_ID_hipGraphExternalSemaphoresWaitNodeGetParams; - if (strcmp("hipGraphExternalSemaphoresWaitNodeSetParams", name) == 0) return HIP_API_ID_hipGraphExternalSemaphoresWaitNodeSetParams; if (strcmp("hipGraphGetEdges", name) == 0) return HIP_API_ID_hipGraphGetEdges; if (strcmp("hipGraphGetNodes", name) == 0) return HIP_API_ID_hipGraphGetNodes; if (strcmp("hipGraphGetRootNodes", name) == 0) return HIP_API_ID_hipGraphGetRootNodes; @@ -1198,6 +1149,8 @@ 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("hipExternalMemoryGetMappedMipmappedArray", name) == 0) return HIP_API_ID_hipExternalMemoryGetMappedMipmappedArray; + if (strcmp("hipDrvGraphAddMemcpyNode", name) == 0) return HIP_API_ID_hipDrvGraphAddMemcpyNode; return HIP_API_ID_NONE; } @@ -1256,18 +1209,12 @@ typedef struct hip_api_data_s { unsigned int flags__val; hipArray_t array; } hipArrayGetInfo; - struct { - int* device; - int device__val; - const hipDeviceProp_tR0000* prop; - hipDeviceProp_tR0000 prop__val; - } hipChooseDeviceR0000; struct { int* device; int device__val; const hipDeviceProp_tR0600* prop; hipDeviceProp_tR0600 prop__val; - } hipChooseDeviceR0600; + } hipChooseDevice; struct { dim3 gridDim; dim3 blockDim; @@ -1494,38 +1441,6 @@ typedef struct hip_api_data_s { int* driverVersion; int driverVersion__val; } hipDriverGetVersion; - struct { - hipGraphNode_t* phGraphNode; - hipGraphNode_t phGraphNode__val; - hipGraph_t hGraph; - const hipGraphNode_t* dependencies; - hipGraphNode_t dependencies__val; - size_t numDependencies; - const HIP_MEMCPY3D* copyParams; - HIP_MEMCPY3D copyParams__val; - hipCtx_t ctx; - } hipDrvGraphAddMemcpyNode; - struct { - hipGraphNode_t* phGraphNode; - hipGraphNode_t phGraphNode__val; - hipGraph_t hGraph; - const hipGraphNode_t* dependencies; - hipGraphNode_t dependencies__val; - size_t numDependencies; - const HIP_MEMSET_NODE_PARAMS* memsetParams; - HIP_MEMSET_NODE_PARAMS memsetParams__val; - hipCtx_t ctx; - } hipDrvGraphAddMemsetNode; - struct { - hipGraphNode_t hNode; - HIP_MEMCPY3D* nodeParams; - HIP_MEMCPY3D nodeParams__val; - } hipDrvGraphMemcpyNodeGetParams; - struct { - hipGraphNode_t hNode; - const HIP_MEMCPY3D* nodeParams; - HIP_MEMCPY3D nodeParams__val; - } hipDrvGraphMemcpyNodeSetParams; struct { const hip_Memcpy2D* pCopy; hip_Memcpy2D pCopy__val; @@ -1645,13 +1560,6 @@ typedef struct hip_api_data_s { const hipExternalMemoryBufferDesc* bufferDesc; hipExternalMemoryBufferDesc bufferDesc__val; } hipExternalMemoryGetMappedBuffer; - struct { - hipMipmappedArray_t* mipmap; - hipMipmappedArray_t mipmap__val; - hipExternalMemory_t extMem; - const hipExternalMemoryMipmappedArrayDesc* mipmapDesc; - hipExternalMemoryMipmappedArrayDesc mipmapDesc__val; - } hipExternalMemoryGetMappedMipmappedArray; struct { void* ptr; } hipFree; @@ -1718,14 +1626,9 @@ typedef struct hip_api_data_s { unsigned int flags__val; } hipGetDeviceFlags; struct { - hipDeviceProp_tR0000* prop; - hipDeviceProp_tR0000 prop__val; - int device; - } hipGetDevicePropertiesR0000; - struct { - hipDeviceProp_tR0600* prop; - hipDeviceProp_tR0600 prop__val; - int deviceId; + hipDeviceProp_tR0600* props; + hipDeviceProp_tR0600 props__val; + hipDevice_t device; } hipGetDevicePropertiesR0600; struct { hipArray_t* levelArray; @@ -1786,26 +1689,6 @@ typedef struct hip_api_data_s { size_t numDependencies; hipEvent_t event; } hipGraphAddEventWaitNode; - struct { - hipGraphNode_t* pGraphNode; - hipGraphNode_t pGraphNode__val; - hipGraph_t graph; - const hipGraphNode_t* pDependencies; - hipGraphNode_t pDependencies__val; - size_t numDependencies; - const hipExternalSemaphoreSignalNodeParams* nodeParams; - hipExternalSemaphoreSignalNodeParams nodeParams__val; - } hipGraphAddExternalSemaphoresSignalNode; - struct { - hipGraphNode_t* pGraphNode; - hipGraphNode_t pGraphNode__val; - hipGraph_t graph; - const hipGraphNode_t* pDependencies; - hipGraphNode_t pDependencies__val; - size_t numDependencies; - const hipExternalSemaphoreWaitNodeParams* nodeParams; - hipExternalSemaphoreWaitNodeParams nodeParams__val; - } hipGraphAddExternalSemaphoresWaitNode; struct { hipGraphNode_t* pGraphNode; hipGraphNode_t pGraphNode__val; @@ -1966,18 +1849,6 @@ typedef struct hip_api_data_s { hipGraphNode_t hNode; hipEvent_t event; } hipGraphExecEventWaitNodeSetEvent; - struct { - hipGraphExec_t hGraphExec; - hipGraphNode_t hNode; - const hipExternalSemaphoreSignalNodeParams* nodeParams; - hipExternalSemaphoreSignalNodeParams nodeParams__val; - } hipGraphExecExternalSemaphoresSignalNodeSetParams; - struct { - hipGraphExec_t hGraphExec; - hipGraphNode_t hNode; - const hipExternalSemaphoreWaitNodeParams* nodeParams; - hipExternalSemaphoreWaitNodeParams nodeParams__val; - } hipGraphExecExternalSemaphoresWaitNodeSetParams; struct { hipGraphExec_t hGraphExec; hipGraphNode_t node; @@ -2036,26 +1907,6 @@ typedef struct hip_api_data_s { hipGraphExecUpdateResult* updateResult_out; hipGraphExecUpdateResult updateResult_out__val; } hipGraphExecUpdate; - struct { - hipGraphNode_t hNode; - hipExternalSemaphoreSignalNodeParams* params_out; - hipExternalSemaphoreSignalNodeParams params_out__val; - } hipGraphExternalSemaphoresSignalNodeGetParams; - struct { - hipGraphNode_t hNode; - const hipExternalSemaphoreSignalNodeParams* nodeParams; - hipExternalSemaphoreSignalNodeParams nodeParams__val; - } hipGraphExternalSemaphoresSignalNodeSetParams; - struct { - hipGraphNode_t hNode; - hipExternalSemaphoreWaitNodeParams* params_out; - hipExternalSemaphoreWaitNodeParams params_out__val; - } hipGraphExternalSemaphoresWaitNodeGetParams; - struct { - hipGraphNode_t hNode; - const hipExternalSemaphoreWaitNodeParams* nodeParams; - hipExternalSemaphoreWaitNodeParams nodeParams__val; - } hipGraphExternalSemaphoresWaitNodeSetParams; struct { hipGraph_t graph; hipGraphNode_t* from; @@ -3411,6 +3262,24 @@ typedef struct hip_api_data_s { unsigned int numExtSems; hipStream_t stream; } hipWaitExternalSemaphoresAsync; + struct { + hipMipmappedArray_t* mipmap; + hipExternalMemory_t extMem; + const hipExternalMemoryMipmappedArrayDesc* mipmapDesc; + hipExternalMemoryMipmappedArrayDesc mipmapDesc__val; + } hipExternalMemoryGetMappedMipmappedArray; + + struct { + hipGraphNode_t* phGraphNode; + hipGraphNode_t phGraphNode__val; + hipGraph_t hGraph; + const hipGraphNode_t* dependencies; + hipGraphNode_t dependencies__val; + size_t numDependencies; + const HIP_MEMCPY3D* copyParams; + HIP_MEMCPY3D copyParams__val; + hipCtx_t ctx; + } hipDrvGraphAddMemcpyNode; } args; uint64_t *phase_data; } hip_api_data_t; @@ -3461,15 +3330,10 @@ typedef struct hip_api_data_s { cb_data.args.hipArrayGetInfo.flags = (unsigned int*)flags; \ cb_data.args.hipArrayGetInfo.array = (hipArray_t)array; \ }; -// hipChooseDeviceR0000[('int*', 'device'), ('const hipDeviceProp_tR0000*', 'prop')] -#define INIT_hipChooseDeviceR0000_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipChooseDeviceR0000.device = (int*)device; \ - cb_data.args.hipChooseDeviceR0000.prop = (const hipDeviceProp_tR0000*)properties; \ -}; -// hipChooseDeviceR0600[('int*', 'device'), ('const hipDeviceProp_tR0600*', 'prop')] -#define INIT_hipChooseDeviceR0600_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipChooseDeviceR0600.device = (int*)device; \ - cb_data.args.hipChooseDeviceR0600.prop = (const hipDeviceProp_tR0600*)properties; \ +// hipChooseDevice[('int*', 'device'), ('const hipDeviceProp_tR0600*', 'prop')] +#define INIT_hipChooseDevice_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipChooseDevice.device = (int*)device; \ + cb_data.args.hipChooseDevice.prop = (const hipDeviceProp_tR0600*)properties; \ }; // hipConfigureCall[('dim3', 'gridDim'), ('dim3', 'blockDim'), ('size_t', 'sharedMem'), ('hipStream_t', 'stream')] #define INIT_hipConfigureCall_CB_ARGS_DATA(cb_data) { \ @@ -3724,18 +3588,6 @@ typedef struct hip_api_data_s { #define INIT_hipDriverGetVersion_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipDriverGetVersion.driverVersion = (int*)driverVersion; \ }; -// 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) { \ -}; -// hipDrvGraphAddMemsetNode[('hipGraphNode_t*', 'phGraphNode'), ('hipGraph_t', 'hGraph'), ('const hipGraphNode_t*', 'dependencies'), ('size_t', 'numDependencies'), ('const HIP_MEMSET_NODE_PARAMS*', 'memsetParams'), ('hipCtx_t', 'ctx')] -#define INIT_hipDrvGraphAddMemsetNode_CB_ARGS_DATA(cb_data) { \ -}; -// hipDrvGraphMemcpyNodeGetParams[('hipGraphNode_t', 'hNode'), ('HIP_MEMCPY3D*', 'nodeParams')] -#define INIT_hipDrvGraphMemcpyNodeGetParams_CB_ARGS_DATA(cb_data) { \ -}; -// hipDrvGraphMemcpyNodeSetParams[('hipGraphNode_t', 'hNode'), ('const HIP_MEMCPY3D*', 'nodeParams')] -#define INIT_hipDrvGraphMemcpyNodeSetParams_CB_ARGS_DATA(cb_data) { \ -}; // hipDrvMemcpy2DUnaligned[('const hip_Memcpy2D*', 'pCopy')] #define INIT_hipDrvMemcpy2DUnaligned_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipDrvMemcpy2DUnaligned.pCopy = (const hip_Memcpy2D*)pCopy; \ @@ -3856,10 +3708,10 @@ typedef struct hip_api_data_s { }; // hipExternalMemoryGetMappedMipmappedArray[('hipMipmappedArray_t*', 'mipmap'), ('hipExternalMemory_t', 'extMem'), ('const hipExternalMemoryMipmappedArrayDesc*', 'mipmapDesc')] #define INIT_hipExternalMemoryGetMappedMipmappedArray_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipExternalMemoryGetMappedMipmappedArray.mipmap = (hipMipmappedArray_t*)mipmap; \ - cb_data.args.hipExternalMemoryGetMappedMipmappedArray.extMem = (hipExternalMemory_t)extMem; \ - cb_data.args.hipExternalMemoryGetMappedMipmappedArray.mipmapDesc = (const hipExternalMemoryMipmappedArrayDesc*)mipmapDesc; \ -}; + cb_data.args.hipExternalMemoryGetMappedMipmappedArray.mipmap = (hipMipmappedArray_t*)mipmap; \ + cb_data.args.hipExternalMemoryGetMappedMipmappedArray.extMem = (hipExternalMemory_t)extMem; \ + cb_data.args.hipExternalMemoryGetMappedMipmappedArray.mipmapDesc = (const hipExternalMemoryMipmappedArrayDesc*)mipmapDesc; \ + }; // hipFree[('void*', 'ptr')] #define INIT_hipFree_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipFree.ptr = (void*)ptr; \ @@ -3932,15 +3784,10 @@ typedef struct hip_api_data_s { #define INIT_hipGetDeviceFlags_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipGetDeviceFlags.flags = (unsigned int*)flags; \ }; -// hipGetDevicePropertiesR0000[('hipDeviceProp_tR0000*', 'prop'), ('int', 'device')] -#define INIT_hipGetDevicePropertiesR0000_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipGetDevicePropertiesR0000.prop = (hipDeviceProp_tR0000*)prop; \ - cb_data.args.hipGetDevicePropertiesR0000.device = (int)device; \ -}; -// hipGetDevicePropertiesR0600[('hipDeviceProp_tR0600*', 'prop'), ('int', 'deviceId')] +// hipGetDeviceProperties[('hipDeviceProp_tR0600*', 'props'), ('hipDevice_t', 'device')] #define INIT_hipGetDevicePropertiesR0600_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipGetDevicePropertiesR0600.prop = (hipDeviceProp_tR0600*)prop; \ - cb_data.args.hipGetDevicePropertiesR0600.deviceId = (int)device; \ + cb_data.args.hipGetDevicePropertiesR0600.props = (hipDeviceProp_tR0600*)props; \ + cb_data.args.hipGetDevicePropertiesR0600.device = (hipDevice_t)device; \ }; // hipGetErrorString[] #define INIT_hipGetErrorString_CB_ARGS_DATA(cb_data) { \ @@ -4002,12 +3849,6 @@ typedef struct hip_api_data_s { cb_data.args.hipGraphAddEventWaitNode.numDependencies = (size_t)numDependencies; \ cb_data.args.hipGraphAddEventWaitNode.event = (hipEvent_t)event; \ }; -// hipGraphAddExternalSemaphoresSignalNode[('hipGraphNode_t*', 'pGraphNode'), ('hipGraph_t', 'graph'), ('const hipGraphNode_t*', 'pDependencies'), ('size_t', 'numDependencies'), ('const hipExternalSemaphoreSignalNodeParams*', 'nodeParams')] -#define INIT_hipGraphAddExternalSemaphoresSignalNode_CB_ARGS_DATA(cb_data) { \ -}; -// hipGraphAddExternalSemaphoresWaitNode[('hipGraphNode_t*', 'pGraphNode'), ('hipGraph_t', 'graph'), ('const hipGraphNode_t*', 'pDependencies'), ('size_t', 'numDependencies'), ('const hipExternalSemaphoreWaitNodeParams*', 'nodeParams')] -#define INIT_hipGraphAddExternalSemaphoresWaitNode_CB_ARGS_DATA(cb_data) { \ -}; // hipGraphAddHostNode[('hipGraphNode_t*', 'pGraphNode'), ('hipGraph_t', 'graph'), ('const hipGraphNode_t*', 'pDependencies'), ('size_t', 'numDependencies'), ('const hipHostNodeParams*', 'pNodeParams')] #define INIT_hipGraphAddHostNode_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipGraphAddHostNode.pGraphNode = (hipGraphNode_t*)pGraphNode; \ @@ -4162,12 +4003,6 @@ typedef struct hip_api_data_s { cb_data.args.hipGraphExecEventWaitNodeSetEvent.hNode = (hipGraphNode_t)hNode; \ cb_data.args.hipGraphExecEventWaitNodeSetEvent.event = (hipEvent_t)event; \ }; -// hipGraphExecExternalSemaphoresSignalNodeSetParams[('hipGraphExec_t', 'hGraphExec'), ('hipGraphNode_t', 'hNode'), ('const hipExternalSemaphoreSignalNodeParams*', 'nodeParams')] -#define INIT_hipGraphExecExternalSemaphoresSignalNodeSetParams_CB_ARGS_DATA(cb_data) { \ -}; -// hipGraphExecExternalSemaphoresWaitNodeSetParams[('hipGraphExec_t', 'hGraphExec'), ('hipGraphNode_t', 'hNode'), ('const hipExternalSemaphoreWaitNodeParams*', 'nodeParams')] -#define INIT_hipGraphExecExternalSemaphoresWaitNodeSetParams_CB_ARGS_DATA(cb_data) { \ -}; // hipGraphExecHostNodeSetParams[('hipGraphExec_t', 'hGraphExec'), ('hipGraphNode_t', 'node'), ('const hipHostNodeParams*', 'pNodeParams')] #define INIT_hipGraphExecHostNodeSetParams_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipGraphExecHostNodeSetParams.hGraphExec = (hipGraphExec_t)hGraphExec; \ @@ -4228,18 +4063,6 @@ typedef struct hip_api_data_s { cb_data.args.hipGraphExecUpdate.hErrorNode_out = (hipGraphNode_t*)hErrorNode_out; \ cb_data.args.hipGraphExecUpdate.updateResult_out = (hipGraphExecUpdateResult*)updateResult_out; \ }; -// hipGraphExternalSemaphoresSignalNodeGetParams[('hipGraphNode_t', 'hNode'), ('hipExternalSemaphoreSignalNodeParams*', 'params_out')] -#define INIT_hipGraphExternalSemaphoresSignalNodeGetParams_CB_ARGS_DATA(cb_data) { \ -}; -// hipGraphExternalSemaphoresSignalNodeSetParams[('hipGraphNode_t', 'hNode'), ('const hipExternalSemaphoreSignalNodeParams*', 'nodeParams')] -#define INIT_hipGraphExternalSemaphoresSignalNodeSetParams_CB_ARGS_DATA(cb_data) { \ -}; -// hipGraphExternalSemaphoresWaitNodeGetParams[('hipGraphNode_t', 'hNode'), ('hipExternalSemaphoreWaitNodeParams*', 'params_out')] -#define INIT_hipGraphExternalSemaphoresWaitNodeGetParams_CB_ARGS_DATA(cb_data) { \ -}; -// hipGraphExternalSemaphoresWaitNodeSetParams[('hipGraphNode_t', 'hNode'), ('const hipExternalSemaphoreWaitNodeParams*', 'nodeParams')] -#define INIT_hipGraphExternalSemaphoresWaitNodeSetParams_CB_ARGS_DATA(cb_data) { \ -}; // hipGraphGetEdges[('hipGraph_t', 'graph'), ('hipGraphNode_t*', 'from'), ('hipGraphNode_t*', 'to'), ('size_t*', 'numEdges')] #define INIT_hipGraphGetEdges_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipGraphGetEdges.graph = (hipGraph_t)graph; \ @@ -5630,6 +5453,15 @@ typedef struct hip_api_data_s { cb_data.args.hipWaitExternalSemaphoresAsync.numExtSems = (unsigned int)numExtSems; \ cb_data.args.hipWaitExternalSemaphoresAsync.stream = (hipStream_t)stream; \ }; +// hipDrvGraphAddMemcpyNode[('hipGraphNode_t*', 'phGraphNode'), ('hipGraph_t', 'hGraph'), ('const hipGraphNode_t*', 'dependencies'), ('size_t', 'numDependencies'), ('const hipMemcpy3DParms*', 'copyParams'), ('hipCtx_t', 'ctx')] +#define INIT_hipDrvGraphAddMemcpyNode_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipDrvGraphAddMemcpyNode.phGraphNode = (hipGraphNode_t*)phGraphNode; \ + cb_data.args.hipDrvGraphAddMemcpyNode.hGraph = (hipGraph_t)hGraph; \ + cb_data.args.hipDrvGraphAddMemcpyNode.dependencies = (const hipGraphNode_t*)dependencies; \ + cb_data.args.hipDrvGraphAddMemcpyNode.numDependencies = (size_t)numDependencies; \ + cb_data.args.hipDrvGraphAddMemcpyNode.copyParams = (const HIP_MEMCPY3D*)copyParams; \ + cb_data.args.hipDrvGraphAddMemcpyNode.ctx = (hipCtx_t)ctx; \ +}; #define INIT_CB_ARGS_DATA(cb_id, cb_data) INIT_##cb_id##_CB_ARGS_DATA(cb_data) // Macros for non-public API primitives @@ -5659,6 +5491,8 @@ typedef struct hip_api_data_s { #define INIT_hipGetTextureReference_CB_ARGS_DATA(cb_data) {}; // hipMemcpy2DArrayToArray() #define INIT_hipMemcpy2DArrayToArray_CB_ARGS_DATA(cb_data) {}; +// hipMemcpyArrayToArray() +#define INIT_hipMemcpyArrayToArray_CB_ARGS_DATA(cb_data) {}; // hipMemcpyAtoA() #define INIT_hipMemcpyAtoA_CB_ARGS_DATA(cb_data) {}; // hipMemcpyAtoD() @@ -5667,8 +5501,12 @@ typedef struct hip_api_data_s { #define INIT_hipMemcpyAtoHAsync_CB_ARGS_DATA(cb_data) {}; // hipMemcpyDtoA() #define INIT_hipMemcpyDtoA_CB_ARGS_DATA(cb_data) {}; +// hipMemcpyFromArrayAsync() +#define INIT_hipMemcpyFromArrayAsync_CB_ARGS_DATA(cb_data) {}; // hipMemcpyHtoAAsync() #define INIT_hipMemcpyHtoAAsync_CB_ARGS_DATA(cb_data) {}; +// hipMemcpyToArrayAsync() +#define INIT_hipMemcpyToArrayAsync_CB_ARGS_DATA(cb_data) {}; // hipSetValidDevices() #define INIT_hipSetValidDevices_CB_ARGS_DATA(cb_data) {}; // hipTexObjectCreate() @@ -5745,15 +5583,10 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { if (data->args.hipArrayGetInfo.extent) data->args.hipArrayGetInfo.extent__val = *(data->args.hipArrayGetInfo.extent); if (data->args.hipArrayGetInfo.flags) data->args.hipArrayGetInfo.flags__val = *(data->args.hipArrayGetInfo.flags); break; -// hipChooseDeviceR0000[('int*', 'device'), ('const hipDeviceProp_tR0000*', 'prop')] - case HIP_API_ID_hipChooseDeviceR0000: - if (data->args.hipChooseDeviceR0000.device) data->args.hipChooseDeviceR0000.device__val = *(data->args.hipChooseDeviceR0000.device); - if (data->args.hipChooseDeviceR0000.prop) data->args.hipChooseDeviceR0000.prop__val = *(data->args.hipChooseDeviceR0000.prop); - break; -// hipChooseDeviceR0600[('int*', 'device'), ('const hipDeviceProp_tR0600*', 'prop')] - case HIP_API_ID_hipChooseDeviceR0600: - if (data->args.hipChooseDeviceR0600.device) data->args.hipChooseDeviceR0600.device__val = *(data->args.hipChooseDeviceR0600.device); - if (data->args.hipChooseDeviceR0600.prop) data->args.hipChooseDeviceR0600.prop__val = *(data->args.hipChooseDeviceR0600.prop); +// hipChooseDevice[('int*', 'device'), ('const hipDeviceProp_tR0600*', 'prop')] + case HIP_API_ID_hipChooseDevice: + if (data->args.hipChooseDevice.device) data->args.hipChooseDevice.device__val = *(data->args.hipChooseDevice.device); + if (data->args.hipChooseDevice.prop) data->args.hipChooseDevice.prop__val = *(data->args.hipChooseDevice.prop); break; // hipConfigureCall[('dim3', 'gridDim'), ('dim3', 'blockDim'), ('size_t', 'sharedMem'), ('hipStream_t', 'stream')] case HIP_API_ID_hipConfigureCall: @@ -5950,26 +5783,6 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { case HIP_API_ID_hipDriverGetVersion: if (data->args.hipDriverGetVersion.driverVersion) data->args.hipDriverGetVersion.driverVersion__val = *(data->args.hipDriverGetVersion.driverVersion); break; -// hipDrvGraphAddMemcpyNode[('hipGraphNode_t*', 'phGraphNode'), ('hipGraph_t', 'hGraph'), ('const hipGraphNode_t*', 'dependencies'), ('size_t', 'numDependencies'), ('const HIP_MEMCPY3D*', 'copyParams'), ('hipCtx_t', 'ctx')] - case HIP_API_ID_hipDrvGraphAddMemcpyNode: - if (data->args.hipDrvGraphAddMemcpyNode.phGraphNode) data->args.hipDrvGraphAddMemcpyNode.phGraphNode__val = *(data->args.hipDrvGraphAddMemcpyNode.phGraphNode); - if (data->args.hipDrvGraphAddMemcpyNode.dependencies) data->args.hipDrvGraphAddMemcpyNode.dependencies__val = *(data->args.hipDrvGraphAddMemcpyNode.dependencies); - if (data->args.hipDrvGraphAddMemcpyNode.copyParams) data->args.hipDrvGraphAddMemcpyNode.copyParams__val = *(data->args.hipDrvGraphAddMemcpyNode.copyParams); - break; -// hipDrvGraphAddMemsetNode[('hipGraphNode_t*', 'phGraphNode'), ('hipGraph_t', 'hGraph'), ('const hipGraphNode_t*', 'dependencies'), ('size_t', 'numDependencies'), ('const HIP_MEMSET_NODE_PARAMS*', 'memsetParams'), ('hipCtx_t', 'ctx')] - case HIP_API_ID_hipDrvGraphAddMemsetNode: - if (data->args.hipDrvGraphAddMemsetNode.phGraphNode) data->args.hipDrvGraphAddMemsetNode.phGraphNode__val = *(data->args.hipDrvGraphAddMemsetNode.phGraphNode); - if (data->args.hipDrvGraphAddMemsetNode.dependencies) data->args.hipDrvGraphAddMemsetNode.dependencies__val = *(data->args.hipDrvGraphAddMemsetNode.dependencies); - if (data->args.hipDrvGraphAddMemsetNode.memsetParams) data->args.hipDrvGraphAddMemsetNode.memsetParams__val = *(data->args.hipDrvGraphAddMemsetNode.memsetParams); - break; -// hipDrvGraphMemcpyNodeGetParams[('hipGraphNode_t', 'hNode'), ('HIP_MEMCPY3D*', 'nodeParams')] - case HIP_API_ID_hipDrvGraphMemcpyNodeGetParams: - if (data->args.hipDrvGraphMemcpyNodeGetParams.nodeParams) data->args.hipDrvGraphMemcpyNodeGetParams.nodeParams__val = *(data->args.hipDrvGraphMemcpyNodeGetParams.nodeParams); - break; -// hipDrvGraphMemcpyNodeSetParams[('hipGraphNode_t', 'hNode'), ('const HIP_MEMCPY3D*', 'nodeParams')] - case HIP_API_ID_hipDrvGraphMemcpyNodeSetParams: - if (data->args.hipDrvGraphMemcpyNodeSetParams.nodeParams) data->args.hipDrvGraphMemcpyNodeSetParams.nodeParams__val = *(data->args.hipDrvGraphMemcpyNodeSetParams.nodeParams); - break; // hipDrvMemcpy2DUnaligned[('const hip_Memcpy2D*', 'pCopy')] case HIP_API_ID_hipDrvMemcpy2DUnaligned: if (data->args.hipDrvMemcpy2DUnaligned.pCopy) data->args.hipDrvMemcpy2DUnaligned.pCopy__val = *(data->args.hipDrvMemcpy2DUnaligned.pCopy); @@ -6049,7 +5862,6 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { break; // hipExternalMemoryGetMappedMipmappedArray[('hipMipmappedArray_t*', 'mipmap'), ('hipExternalMemory_t', 'extMem'), ('const hipExternalMemoryMipmappedArrayDesc*', 'mipmapDesc')] case HIP_API_ID_hipExternalMemoryGetMappedMipmappedArray: - if (data->args.hipExternalMemoryGetMappedMipmappedArray.mipmap) data->args.hipExternalMemoryGetMappedMipmappedArray.mipmap__val = *(data->args.hipExternalMemoryGetMappedMipmappedArray.mipmap); if (data->args.hipExternalMemoryGetMappedMipmappedArray.mipmapDesc) data->args.hipExternalMemoryGetMappedMipmappedArray.mipmapDesc__val = *(data->args.hipExternalMemoryGetMappedMipmappedArray.mipmapDesc); break; // hipFree[('void*', 'ptr')] @@ -6105,13 +5917,9 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { case HIP_API_ID_hipGetDeviceFlags: if (data->args.hipGetDeviceFlags.flags) data->args.hipGetDeviceFlags.flags__val = *(data->args.hipGetDeviceFlags.flags); break; -// hipGetDevicePropertiesR0000[('hipDeviceProp_tR0000*', 'prop'), ('int', 'device')] - case HIP_API_ID_hipGetDevicePropertiesR0000: - if (data->args.hipGetDevicePropertiesR0000.prop) data->args.hipGetDevicePropertiesR0000.prop__val = *(data->args.hipGetDevicePropertiesR0000.prop); - break; -// hipGetDevicePropertiesR0600[('hipDeviceProp_tR0600*', 'prop'), ('int', 'deviceId')] +// hipGetDevicePropertiesR0600[('hipDeviceProp_tR0600*', 'props'), ('hipDevice_t', 'device')] case HIP_API_ID_hipGetDevicePropertiesR0600: - if (data->args.hipGetDevicePropertiesR0600.prop) data->args.hipGetDevicePropertiesR0600.prop__val = *(data->args.hipGetDevicePropertiesR0600.prop); + if (data->args.hipGetDevicePropertiesR0600.props) data->args.hipGetDevicePropertiesR0600.props__val = *(data->args.hipGetDevicePropertiesR0600.props); break; // hipGetErrorString[] case HIP_API_ID_hipGetErrorString: @@ -6156,18 +5964,6 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { if (data->args.hipGraphAddEventWaitNode.pGraphNode) data->args.hipGraphAddEventWaitNode.pGraphNode__val = *(data->args.hipGraphAddEventWaitNode.pGraphNode); if (data->args.hipGraphAddEventWaitNode.pDependencies) data->args.hipGraphAddEventWaitNode.pDependencies__val = *(data->args.hipGraphAddEventWaitNode.pDependencies); break; -// hipGraphAddExternalSemaphoresSignalNode[('hipGraphNode_t*', 'pGraphNode'), ('hipGraph_t', 'graph'), ('const hipGraphNode_t*', 'pDependencies'), ('size_t', 'numDependencies'), ('const hipExternalSemaphoreSignalNodeParams*', 'nodeParams')] - case HIP_API_ID_hipGraphAddExternalSemaphoresSignalNode: - if (data->args.hipGraphAddExternalSemaphoresSignalNode.pGraphNode) data->args.hipGraphAddExternalSemaphoresSignalNode.pGraphNode__val = *(data->args.hipGraphAddExternalSemaphoresSignalNode.pGraphNode); - if (data->args.hipGraphAddExternalSemaphoresSignalNode.pDependencies) data->args.hipGraphAddExternalSemaphoresSignalNode.pDependencies__val = *(data->args.hipGraphAddExternalSemaphoresSignalNode.pDependencies); - if (data->args.hipGraphAddExternalSemaphoresSignalNode.nodeParams) data->args.hipGraphAddExternalSemaphoresSignalNode.nodeParams__val = *(data->args.hipGraphAddExternalSemaphoresSignalNode.nodeParams); - break; -// hipGraphAddExternalSemaphoresWaitNode[('hipGraphNode_t*', 'pGraphNode'), ('hipGraph_t', 'graph'), ('const hipGraphNode_t*', 'pDependencies'), ('size_t', 'numDependencies'), ('const hipExternalSemaphoreWaitNodeParams*', 'nodeParams')] - case HIP_API_ID_hipGraphAddExternalSemaphoresWaitNode: - if (data->args.hipGraphAddExternalSemaphoresWaitNode.pGraphNode) data->args.hipGraphAddExternalSemaphoresWaitNode.pGraphNode__val = *(data->args.hipGraphAddExternalSemaphoresWaitNode.pGraphNode); - if (data->args.hipGraphAddExternalSemaphoresWaitNode.pDependencies) data->args.hipGraphAddExternalSemaphoresWaitNode.pDependencies__val = *(data->args.hipGraphAddExternalSemaphoresWaitNode.pDependencies); - if (data->args.hipGraphAddExternalSemaphoresWaitNode.nodeParams) data->args.hipGraphAddExternalSemaphoresWaitNode.nodeParams__val = *(data->args.hipGraphAddExternalSemaphoresWaitNode.nodeParams); - break; // hipGraphAddHostNode[('hipGraphNode_t*', 'pGraphNode'), ('hipGraph_t', 'graph'), ('const hipGraphNode_t*', 'pDependencies'), ('size_t', 'numDependencies'), ('const hipHostNodeParams*', 'pNodeParams')] case HIP_API_ID_hipGraphAddHostNode: if (data->args.hipGraphAddHostNode.pGraphNode) data->args.hipGraphAddHostNode.pGraphNode__val = *(data->args.hipGraphAddHostNode.pGraphNode); @@ -6266,14 +6062,6 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipGraphExecEventWaitNodeSetEvent[('hipGraphExec_t', 'hGraphExec'), ('hipGraphNode_t', 'hNode'), ('hipEvent_t', 'event')] case HIP_API_ID_hipGraphExecEventWaitNodeSetEvent: break; -// hipGraphExecExternalSemaphoresSignalNodeSetParams[('hipGraphExec_t', 'hGraphExec'), ('hipGraphNode_t', 'hNode'), ('const hipExternalSemaphoreSignalNodeParams*', 'nodeParams')] - case HIP_API_ID_hipGraphExecExternalSemaphoresSignalNodeSetParams: - if (data->args.hipGraphExecExternalSemaphoresSignalNodeSetParams.nodeParams) data->args.hipGraphExecExternalSemaphoresSignalNodeSetParams.nodeParams__val = *(data->args.hipGraphExecExternalSemaphoresSignalNodeSetParams.nodeParams); - break; -// hipGraphExecExternalSemaphoresWaitNodeSetParams[('hipGraphExec_t', 'hGraphExec'), ('hipGraphNode_t', 'hNode'), ('const hipExternalSemaphoreWaitNodeParams*', 'nodeParams')] - case HIP_API_ID_hipGraphExecExternalSemaphoresWaitNodeSetParams: - if (data->args.hipGraphExecExternalSemaphoresWaitNodeSetParams.nodeParams) data->args.hipGraphExecExternalSemaphoresWaitNodeSetParams.nodeParams__val = *(data->args.hipGraphExecExternalSemaphoresWaitNodeSetParams.nodeParams); - break; // hipGraphExecHostNodeSetParams[('hipGraphExec_t', 'hGraphExec'), ('hipGraphNode_t', 'node'), ('const hipHostNodeParams*', 'pNodeParams')] case HIP_API_ID_hipGraphExecHostNodeSetParams: if (data->args.hipGraphExecHostNodeSetParams.pNodeParams) data->args.hipGraphExecHostNodeSetParams.pNodeParams__val = *(data->args.hipGraphExecHostNodeSetParams.pNodeParams); @@ -6304,22 +6092,6 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { if (data->args.hipGraphExecUpdate.hErrorNode_out) data->args.hipGraphExecUpdate.hErrorNode_out__val = *(data->args.hipGraphExecUpdate.hErrorNode_out); if (data->args.hipGraphExecUpdate.updateResult_out) data->args.hipGraphExecUpdate.updateResult_out__val = *(data->args.hipGraphExecUpdate.updateResult_out); break; -// hipGraphExternalSemaphoresSignalNodeGetParams[('hipGraphNode_t', 'hNode'), ('hipExternalSemaphoreSignalNodeParams*', 'params_out')] - case HIP_API_ID_hipGraphExternalSemaphoresSignalNodeGetParams: - if (data->args.hipGraphExternalSemaphoresSignalNodeGetParams.params_out) data->args.hipGraphExternalSemaphoresSignalNodeGetParams.params_out__val = *(data->args.hipGraphExternalSemaphoresSignalNodeGetParams.params_out); - break; -// hipGraphExternalSemaphoresSignalNodeSetParams[('hipGraphNode_t', 'hNode'), ('const hipExternalSemaphoreSignalNodeParams*', 'nodeParams')] - case HIP_API_ID_hipGraphExternalSemaphoresSignalNodeSetParams: - if (data->args.hipGraphExternalSemaphoresSignalNodeSetParams.nodeParams) data->args.hipGraphExternalSemaphoresSignalNodeSetParams.nodeParams__val = *(data->args.hipGraphExternalSemaphoresSignalNodeSetParams.nodeParams); - break; -// hipGraphExternalSemaphoresWaitNodeGetParams[('hipGraphNode_t', 'hNode'), ('hipExternalSemaphoreWaitNodeParams*', 'params_out')] - case HIP_API_ID_hipGraphExternalSemaphoresWaitNodeGetParams: - if (data->args.hipGraphExternalSemaphoresWaitNodeGetParams.params_out) data->args.hipGraphExternalSemaphoresWaitNodeGetParams.params_out__val = *(data->args.hipGraphExternalSemaphoresWaitNodeGetParams.params_out); - break; -// hipGraphExternalSemaphoresWaitNodeSetParams[('hipGraphNode_t', 'hNode'), ('const hipExternalSemaphoreWaitNodeParams*', 'nodeParams')] - case HIP_API_ID_hipGraphExternalSemaphoresWaitNodeSetParams: - if (data->args.hipGraphExternalSemaphoresWaitNodeSetParams.nodeParams) data->args.hipGraphExternalSemaphoresWaitNodeSetParams.nodeParams__val = *(data->args.hipGraphExternalSemaphoresWaitNodeSetParams.nodeParams); - break; // hipGraphGetEdges[('hipGraph_t', 'graph'), ('hipGraphNode_t*', 'from'), ('hipGraphNode_t*', 'to'), ('size_t*', 'numEdges')] case HIP_API_ID_hipGraphGetEdges: if (data->args.hipGraphGetEdges.from) data->args.hipGraphGetEdges.from__val = *(data->args.hipGraphGetEdges.from); @@ -7169,6 +6941,12 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { if (data->args.hipWaitExternalSemaphoresAsync.extSemArray) data->args.hipWaitExternalSemaphoresAsync.extSemArray__val = *(data->args.hipWaitExternalSemaphoresAsync.extSemArray); if (data->args.hipWaitExternalSemaphoresAsync.paramsArray) data->args.hipWaitExternalSemaphoresAsync.paramsArray__val = *(data->args.hipWaitExternalSemaphoresAsync.paramsArray); break; +// hipDrvGraphAddMemcpyNode[('hipGraphNode_t*', 'phGraphNode'), ('hipGraph_t', 'hGraph'), ('const hipGraphNode_t*', 'dependencies'), ('size_t', 'numDependencies'), ('const HIP_MEMCPY3D*', 'copyParams'), ('hipCtx_t', 'ctx')] + case HIP_API_ID_hipDrvGraphAddMemcpyNode: + if (data->args.hipDrvGraphAddMemcpyNode.phGraphNode) data->args.hipDrvGraphAddMemcpyNode.phGraphNode__val = *(data->args.hipDrvGraphAddMemcpyNode.phGraphNode); + if (data->args.hipDrvGraphAddMemcpyNode.dependencies) data->args.hipDrvGraphAddMemcpyNode.dependencies__val = *(data->args.hipDrvGraphAddMemcpyNode.dependencies); + if (data->args.hipDrvGraphAddMemcpyNode.copyParams) data->args.hipDrvGraphAddMemcpyNode.copyParams__val = *(data->args.hipDrvGraphAddMemcpyNode.copyParams); + break; default: break; }; } @@ -7202,7 +6980,7 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da case HIP_API_ID_hipArray3DCreate: oss << "hipArray3DCreate("; if (data->args.hipArray3DCreate.array == NULL) oss << "array=NULL"; - else { oss << "array="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArray3DCreate.array__val); } + else { oss << "array="; roctracer::hip_support::detail::operator<<(oss, (void*)data->args.hipArray3DCreate.array__val); } if (data->args.hipArray3DCreate.pAllocateArray == NULL) oss << ", pAllocateArray=NULL"; else { oss << ", pAllocateArray="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArray3DCreate.pAllocateArray__val); } oss << ")"; @@ -7211,13 +6989,13 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << "hipArray3DGetDescriptor("; if (data->args.hipArray3DGetDescriptor.pArrayDescriptor == NULL) oss << "pArrayDescriptor=NULL"; else { oss << "pArrayDescriptor="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArray3DGetDescriptor.pArrayDescriptor__val); } - oss << ", array="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArray3DGetDescriptor.array); + oss << "array="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArray3DGetDescriptor.array); oss << ")"; break; case HIP_API_ID_hipArrayCreate: oss << "hipArrayCreate("; if (data->args.hipArrayCreate.pHandle == NULL) oss << "pHandle=NULL"; - else { oss << "pHandle="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayCreate.pHandle__val); } + else { oss << "pHandle="; roctracer::hip_support::detail::operator<<(oss, (void*)data->args.hipArrayCreate.pHandle__val); } if (data->args.hipArrayCreate.pAllocateArray == NULL) oss << ", pAllocateArray=NULL"; else { oss << ", pAllocateArray="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayCreate.pAllocateArray__val); } oss << ")"; @@ -7242,23 +7020,16 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da else { oss << ", extent="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayGetInfo.extent__val); } if (data->args.hipArrayGetInfo.flags == NULL) oss << ", flags=NULL"; else { oss << ", flags="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayGetInfo.flags__val); } + if (data->args.hipArrayGetInfo.array == NULL) oss << ", array=NULL"; oss << ", array="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayGetInfo.array); oss << ")"; break; - case HIP_API_ID_hipChooseDeviceR0000: - oss << "hipChooseDeviceR0000("; - if (data->args.hipChooseDeviceR0000.device == NULL) oss << "device=NULL"; - else { oss << "device="; roctracer::hip_support::detail::operator<<(oss, data->args.hipChooseDeviceR0000.device__val); } - if (data->args.hipChooseDeviceR0000.prop == NULL) oss << ", prop=NULL"; - else { oss << ", prop="; roctracer::hip_support::detail::operator<<(oss, data->args.hipChooseDeviceR0000.prop__val); } - oss << ")"; - break; - case HIP_API_ID_hipChooseDeviceR0600: - oss << "hipChooseDeviceR0600("; - if (data->args.hipChooseDeviceR0600.device == NULL) oss << "device=NULL"; - else { oss << "device="; roctracer::hip_support::detail::operator<<(oss, data->args.hipChooseDeviceR0600.device__val); } - if (data->args.hipChooseDeviceR0600.prop == NULL) oss << ", prop=NULL"; - else { oss << ", prop="; roctracer::hip_support::detail::operator<<(oss, data->args.hipChooseDeviceR0600.prop__val); } + case HIP_API_ID_hipChooseDevice: + oss << "hipChooseDevice("; + if (data->args.hipChooseDevice.device == NULL) oss << "device=NULL"; + else { oss << "device="; roctracer::hip_support::detail::operator<<(oss, data->args.hipChooseDevice.device__val); } + if (data->args.hipChooseDevice.prop == NULL) oss << ", prop=NULL"; + else { oss << ", prop="; roctracer::hip_support::detail::operator<<(oss, data->args.hipChooseDevice.prop__val); } oss << ")"; break; case HIP_API_ID_hipConfigureCall: @@ -7601,46 +7372,6 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da else { oss << "driverVersion="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDriverGetVersion.driverVersion__val); } oss << ")"; break; - case HIP_API_ID_hipDrvGraphAddMemcpyNode: - oss << "hipDrvGraphAddMemcpyNode("; - if (data->args.hipDrvGraphAddMemcpyNode.phGraphNode == NULL) oss << "phGraphNode=NULL"; - else { oss << "phGraphNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.phGraphNode__val); } - oss << ", hGraph="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.hGraph); - if (data->args.hipDrvGraphAddMemcpyNode.dependencies == NULL) oss << ", dependencies=NULL"; - else { oss << ", dependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.dependencies__val); } - oss << ", numDependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.numDependencies); - if (data->args.hipDrvGraphAddMemcpyNode.copyParams == NULL) oss << ", copyParams=NULL"; - else { oss << ", copyParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.copyParams__val); } - oss << ", ctx="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.ctx); - oss << ")"; - break; - case HIP_API_ID_hipDrvGraphAddMemsetNode: - oss << "hipDrvGraphAddMemsetNode("; - if (data->args.hipDrvGraphAddMemsetNode.phGraphNode == NULL) oss << "phGraphNode=NULL"; - else { oss << "phGraphNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemsetNode.phGraphNode__val); } - oss << ", hGraph="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemsetNode.hGraph); - if (data->args.hipDrvGraphAddMemsetNode.dependencies == NULL) oss << ", dependencies=NULL"; - else { oss << ", dependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemsetNode.dependencies__val); } - oss << ", numDependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemsetNode.numDependencies); - if (data->args.hipDrvGraphAddMemsetNode.memsetParams == NULL) oss << ", memsetParams=NULL"; - else { oss << ", memsetParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemsetNode.memsetParams__val); } - oss << ", ctx="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemsetNode.ctx); - oss << ")"; - break; - case HIP_API_ID_hipDrvGraphMemcpyNodeGetParams: - oss << "hipDrvGraphMemcpyNodeGetParams("; - oss << "hNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphMemcpyNodeGetParams.hNode); - if (data->args.hipDrvGraphMemcpyNodeGetParams.nodeParams == NULL) oss << ", nodeParams=NULL"; - else { oss << ", nodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphMemcpyNodeGetParams.nodeParams__val); } - oss << ")"; - break; - case HIP_API_ID_hipDrvGraphMemcpyNodeSetParams: - oss << "hipDrvGraphMemcpyNodeSetParams("; - oss << "hNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphMemcpyNodeSetParams.hNode); - if (data->args.hipDrvGraphMemcpyNodeSetParams.nodeParams == NULL) oss << ", nodeParams=NULL"; - else { oss << ", nodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphMemcpyNodeSetParams.nodeParams__val); } - oss << ")"; - break; case HIP_API_ID_hipDrvMemcpy2DUnaligned: oss << "hipDrvMemcpy2DUnaligned("; if (data->args.hipDrvMemcpy2DUnaligned.pCopy == NULL) oss << "pCopy=NULL"; @@ -7800,13 +7531,12 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da break; case HIP_API_ID_hipExternalMemoryGetMappedMipmappedArray: oss << "hipExternalMemoryGetMappedMipmappedArray("; - if (data->args.hipExternalMemoryGetMappedMipmappedArray.mipmap == NULL) oss << "mipmap=NULL"; - else { oss << "mipmap="; roctracer::hip_support::detail::operator<<(oss, data->args.hipExternalMemoryGetMappedMipmappedArray.mipmap__val); } + oss << "mipmap="; roctracer::hip_support::detail::operator<<(oss, data->args.hipExternalMemoryGetMappedMipmappedArray.mipmap); oss << ", extMem="; roctracer::hip_support::detail::operator<<(oss, data->args.hipExternalMemoryGetMappedMipmappedArray.extMem); if (data->args.hipExternalMemoryGetMappedMipmappedArray.mipmapDesc == NULL) oss << ", mipmapDesc=NULL"; else { oss << ", mipmapDesc="; roctracer::hip_support::detail::operator<<(oss, data->args.hipExternalMemoryGetMappedMipmappedArray.mipmapDesc__val); } oss << ")"; - break; + break; case HIP_API_ID_hipFree: oss << "hipFree("; oss << "ptr="; roctracer::hip_support::detail::operator<<(oss, data->args.hipFree.ptr); @@ -7902,18 +7632,11 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da else { oss << "flags="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetDeviceFlags.flags__val); } oss << ")"; break; - case HIP_API_ID_hipGetDevicePropertiesR0000: - oss << "hipGetDevicePropertiesR0000("; - if (data->args.hipGetDevicePropertiesR0000.prop == NULL) oss << "prop=NULL"; - else { oss << "prop="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetDevicePropertiesR0000.prop__val); } - oss << ", device="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetDevicePropertiesR0000.device); - oss << ")"; - break; case HIP_API_ID_hipGetDevicePropertiesR0600: oss << "hipGetDevicePropertiesR0600("; - if (data->args.hipGetDevicePropertiesR0600.prop == NULL) oss << "prop=NULL"; - else { oss << "prop="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetDevicePropertiesR0600.prop__val); } - oss << ", deviceId="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetDevicePropertiesR0600.deviceId); + if (data->args.hipGetDevicePropertiesR0600.props == NULL) oss << "props=NULL"; + else { oss << "props="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetDevicePropertiesR0600.props__val); } + oss << ", device="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetDevicePropertiesR0600.device); oss << ")"; break; case HIP_API_ID_hipGetErrorString: @@ -7999,30 +7722,6 @@ 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.hipGraphAddEventWaitNode.event); oss << ")"; break; - case HIP_API_ID_hipGraphAddExternalSemaphoresSignalNode: - oss << "hipGraphAddExternalSemaphoresSignalNode("; - if (data->args.hipGraphAddExternalSemaphoresSignalNode.pGraphNode == NULL) oss << "pGraphNode=NULL"; - else { oss << "pGraphNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddExternalSemaphoresSignalNode.pGraphNode__val); } - oss << ", graph="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddExternalSemaphoresSignalNode.graph); - if (data->args.hipGraphAddExternalSemaphoresSignalNode.pDependencies == NULL) oss << ", pDependencies=NULL"; - else { oss << ", pDependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddExternalSemaphoresSignalNode.pDependencies__val); } - oss << ", numDependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddExternalSemaphoresSignalNode.numDependencies); - if (data->args.hipGraphAddExternalSemaphoresSignalNode.nodeParams == NULL) oss << ", nodeParams=NULL"; - else { oss << ", nodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddExternalSemaphoresSignalNode.nodeParams__val); } - oss << ")"; - break; - case HIP_API_ID_hipGraphAddExternalSemaphoresWaitNode: - oss << "hipGraphAddExternalSemaphoresWaitNode("; - if (data->args.hipGraphAddExternalSemaphoresWaitNode.pGraphNode == NULL) oss << "pGraphNode=NULL"; - else { oss << "pGraphNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddExternalSemaphoresWaitNode.pGraphNode__val); } - oss << ", graph="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddExternalSemaphoresWaitNode.graph); - if (data->args.hipGraphAddExternalSemaphoresWaitNode.pDependencies == NULL) oss << ", pDependencies=NULL"; - else { oss << ", pDependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddExternalSemaphoresWaitNode.pDependencies__val); } - oss << ", numDependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddExternalSemaphoresWaitNode.numDependencies); - if (data->args.hipGraphAddExternalSemaphoresWaitNode.nodeParams == NULL) oss << ", nodeParams=NULL"; - else { oss << ", nodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddExternalSemaphoresWaitNode.nodeParams__val); } - oss << ")"; - break; case HIP_API_ID_hipGraphAddHostNode: oss << "hipGraphAddHostNode("; if (data->args.hipGraphAddHostNode.pGraphNode == NULL) oss << "pGraphNode=NULL"; @@ -8229,22 +7928,6 @@ 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.hipGraphExecEventWaitNodeSetEvent.event); oss << ")"; break; - case HIP_API_ID_hipGraphExecExternalSemaphoresSignalNodeSetParams: - oss << "hipGraphExecExternalSemaphoresSignalNodeSetParams("; - oss << "hGraphExec="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecExternalSemaphoresSignalNodeSetParams.hGraphExec); - oss << ", hNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecExternalSemaphoresSignalNodeSetParams.hNode); - if (data->args.hipGraphExecExternalSemaphoresSignalNodeSetParams.nodeParams == NULL) oss << ", nodeParams=NULL"; - else { oss << ", nodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecExternalSemaphoresSignalNodeSetParams.nodeParams__val); } - oss << ")"; - break; - case HIP_API_ID_hipGraphExecExternalSemaphoresWaitNodeSetParams: - oss << "hipGraphExecExternalSemaphoresWaitNodeSetParams("; - oss << "hGraphExec="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecExternalSemaphoresWaitNodeSetParams.hGraphExec); - oss << ", hNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecExternalSemaphoresWaitNodeSetParams.hNode); - if (data->args.hipGraphExecExternalSemaphoresWaitNodeSetParams.nodeParams == NULL) oss << ", nodeParams=NULL"; - else { oss << ", nodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecExternalSemaphoresWaitNodeSetParams.nodeParams__val); } - oss << ")"; - break; case HIP_API_ID_hipGraphExecHostNodeSetParams: oss << "hipGraphExecHostNodeSetParams("; oss << "hGraphExec="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecHostNodeSetParams.hGraphExec); @@ -8319,34 +8002,6 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da else { oss << ", updateResult_out="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExecUpdate.updateResult_out__val); } oss << ")"; break; - case HIP_API_ID_hipGraphExternalSemaphoresSignalNodeGetParams: - oss << "hipGraphExternalSemaphoresSignalNodeGetParams("; - oss << "hNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExternalSemaphoresSignalNodeGetParams.hNode); - if (data->args.hipGraphExternalSemaphoresSignalNodeGetParams.params_out == NULL) oss << ", params_out=NULL"; - else { oss << ", params_out="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExternalSemaphoresSignalNodeGetParams.params_out__val); } - oss << ")"; - break; - case HIP_API_ID_hipGraphExternalSemaphoresSignalNodeSetParams: - oss << "hipGraphExternalSemaphoresSignalNodeSetParams("; - oss << "hNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExternalSemaphoresSignalNodeSetParams.hNode); - if (data->args.hipGraphExternalSemaphoresSignalNodeSetParams.nodeParams == NULL) oss << ", nodeParams=NULL"; - else { oss << ", nodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExternalSemaphoresSignalNodeSetParams.nodeParams__val); } - oss << ")"; - break; - case HIP_API_ID_hipGraphExternalSemaphoresWaitNodeGetParams: - oss << "hipGraphExternalSemaphoresWaitNodeGetParams("; - oss << "hNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExternalSemaphoresWaitNodeGetParams.hNode); - if (data->args.hipGraphExternalSemaphoresWaitNodeGetParams.params_out == NULL) oss << ", params_out=NULL"; - else { oss << ", params_out="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExternalSemaphoresWaitNodeGetParams.params_out__val); } - oss << ")"; - break; - case HIP_API_ID_hipGraphExternalSemaphoresWaitNodeSetParams: - oss << "hipGraphExternalSemaphoresWaitNodeSetParams("; - oss << "hNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExternalSemaphoresWaitNodeSetParams.hNode); - if (data->args.hipGraphExternalSemaphoresWaitNodeSetParams.nodeParams == NULL) oss << ", nodeParams=NULL"; - else { oss << ", nodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphExternalSemaphoresWaitNodeSetParams.nodeParams__val); } - oss << ")"; - break; case HIP_API_ID_hipGraphGetEdges: oss << "hipGraphGetEdges("; oss << "graph="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphGetEdges.graph); @@ -8848,7 +8503,7 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da case HIP_API_ID_hipMallocArray: oss << "hipMallocArray("; if (data->args.hipMallocArray.array == NULL) oss << "array=NULL"; - else { oss << "array="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMallocArray.array__val); } + else { oss << "array="; roctracer::hip_support::detail::operator<<(oss, (void*)data->args.hipMallocArray.array__val); } if (data->args.hipMallocArray.desc == NULL) oss << ", desc=NULL"; else { oss << ", desc="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMallocArray.desc__val); } oss << ", width="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMallocArray.width); @@ -10148,6 +9803,18 @@ 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_hipDrvGraphAddMemcpyNode: + oss << "hipDrvGraphAddMemcpyNode("; + if (data->args.hipDrvGraphAddMemcpyNode.phGraphNode == NULL) oss << "phGraphNode=NULL"; + else { oss << "phGraphNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.phGraphNode__val); } + oss << ", hGraph="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.hGraph); + if (data->args.hipDrvGraphAddMemcpyNode.dependencies == NULL) oss << ", dependencies=NULL"; + else { oss << ", dependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.dependencies__val); } + oss << ", numDependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.numDependencies); + if (data->args.hipDrvGraphAddMemcpyNode.copyParams == NULL) oss << ", copyParams=NULL"; + else { oss << ", copyParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.copyParams__val); } + oss << ")"; + break; default: oss << "unknown"; }; return strdup(oss.str().c_str()); diff --git a/projects/clr/hipamd/src/CMakeLists.txt b/projects/clr/hipamd/src/CMakeLists.txt index 5e96c42dae..6f727ba035 100644 --- a/projects/clr/hipamd/src/CMakeLists.txt +++ b/projects/clr/hipamd/src/CMakeLists.txt @@ -212,12 +212,12 @@ if(USE_PROF_API) set(PROF_API_STR "${PROJECT_BINARY_DIR}/include/hip/amd_detail/hip_prof_str.h") set(PROF_API_STR_IN "${CMAKE_SOURCE_DIR}/hipamd/include/hip/amd_detail/hip_prof_str.h") set(PROF_API_HDR "${HIP_COMMON_INCLUDE_DIR}/hip/hip_runtime_api.h") - set(PROF_GL_HDR "${CMAKE_SOURCE_DIR}/hipamd/include/hip/amd_detail/amd_hip_gl_interop.h") - set(PROF_API_DEPRECATED "${HIP_COMMON_INCLUDE_DIR}/hip/hip_deprecated.h") set(PROF_API_SRC "${CMAKE_CURRENT_SOURCE_DIR}") set(PROF_API_GEN "${CMAKE_CURRENT_SOURCE_DIR}/hip_prof_gen.py") set(PROF_API_LOG "${PROJECT_BINARY_DIR}/hip_prof_gen.log.txt") set(PROF_API_NEWHDR "${PROJECT_BINARY_DIR}/new_header.h") + set(PROF_API_NEWHDR_GEN "${CMAKE_CURRENT_SOURCE_DIR}/hip_find_defs.py") + set(PROF_API_DEPRECATED "${CMAKE_CURRENT_SOURCE_DIR}/hip_device_deprecated.cpp") find_package(Python3 COMPONENTS Interpreter REQUIRED) execute_process(COMMAND ${Python3_EXECUTABLE} -c "import CppHeaderParser" @@ -231,29 +231,19 @@ if(USE_PROF_API) ") endif() - add_custom_command(OUTPUT ${PROF_API_NEWHDR}.i - COMMAND ${CMAKE_COMMAND} -E cat ${PROF_API_HDR} ${PROF_GL_HDR} > ${PROF_API_NEWHDR} - COMMAND ${CMAKE_C_COMPILER} - "-D$,;-D>" - "-I$,;-I>" - "-DHIP_INCLUDE_HIP_HIP_RUNTIME_PT_API_H=1" - ${c_flags} - $ - ${CPP_EXTRA_C_FLAGS} - -E ${PROF_API_NEWHDR} -o ${PROF_API_NEWHDR}.i - COMMAND_EXPAND_LISTS VERBATIM - IMPLICIT_DEPENDS C ${PROF_API_HDR} ${PROF_GL_HDR} ${PROF_API_DEPRECATED} - DEPENDS ${PROF_API_HDR} ${PROF_GL_HDR} ${PROF_API_DEPRECATED} + add_custom_command(OUTPUT ${PROF_API_NEWHDR} + COMMAND ${Python3_EXECUTABLE} ${PROF_API_NEWHDR_GEN} --output ${PROF_API_NEWHDR} --deprecated ${PROF_API_DEPRECATED} ${PROF_API_HDR} + DEPENDS ${PROF_API_NEWHDR_GEN} ${PROF_API_HDR} COMMENT "Generating new header from hip_runtime_api.h") add_custom_command(OUTPUT ${PROF_API_STR} - COMMAND ${Python3_EXECUTABLE} ${PROF_API_GEN} -v -t --priv ${PROF_API_NEWHDR}.i ${PROF_API_SRC} ${PROF_API_STR_IN} ${PROF_API_STR} - DEPENDS ${PROF_API_STR_IN} ${PROF_API_NEWHDR}.i ${PROF_API_GEN} + COMMAND ${Python3_EXECUTABLE} ${PROF_API_GEN} -v -t --priv ${PROF_API_NEWHDR} ${PROF_API_SRC} ${PROF_API_STR_IN} ${PROF_API_STR} + DEPENDS ${PROF_API_STR_IN} ${PROF_API_NEWHDR} ${PROF_API_GEN} COMMENT "Generating profiling primitives: ${PROF_API_STR}") add_custom_target(gen-prof-api-str-header ALL DEPENDS ${PROF_API_STR} - SOURCES ${PROF_API_NEWHDR}.i) + SOURCES ${PROF_API_NEWHDR}) set_target_properties(amdhip64 PROPERTIES PUBLIC_HEADER ${PROF_API_STR}) diff --git a/projects/clr/hipamd/src/amdhip.def b/projects/clr/hipamd/src/amdhip.def index 97a23d5b31..a3bbd32766 100644 --- a/projects/clr/hipamd/src/amdhip.def +++ b/projects/clr/hipamd/src/amdhip.def @@ -1,7 +1,5 @@ EXPORTS hipChooseDevice -hipChooseDeviceR0000 -hipChooseDeviceR0600 hipCtxCreate hipCtxDestroy hipCtxDisablePeerAccess @@ -65,7 +63,6 @@ hipFuncSetSharedMemConfig hipGetDevice hipGetDeviceCount hipGetDeviceProperties -hipGetDevicePropertiesR0000 hipGetDevicePropertiesR0600 hipGetErrorName hipGetErrorString diff --git a/projects/clr/hipamd/src/hip_device.cpp b/projects/clr/hipamd/src/hip_device.cpp index 48c0e5e284..32cfdaf624 100644 --- a/projects/clr/hipamd/src/hip_device.cpp +++ b/projects/clr/hipamd/src/hip_device.cpp @@ -19,7 +19,6 @@ THE SOFTWARE. */ #include -#include #include "hip_internal.hpp" #include "hip_mempool_impl.hpp" @@ -276,7 +275,7 @@ hipError_t hipDeviceGetUuid(hipUUID* uuid, hipDevice_t device) { HIP_RETURN(hipSuccess); } -hipError_t ihipGetDeviceProperties(hipDeviceProp_tR0600* props, int device) { +hipError_t ihipGetDeviceProperties(hipDeviceProp_tR0600* props, hipDevice_t device) { if (props == nullptr) { return hipErrorInvalidValue; } @@ -460,27 +459,107 @@ hipError_t ihipGetDeviceProperties(hipDeviceProp_tR0600* props, int device) { return hipSuccess; } -hipError_t hipGetDevicePropertiesR0600(hipDeviceProp_tR0600* prop, int device) { - HIP_INIT_API(hipGetDevicePropertiesR0600, prop, device); +hipError_t hipGetDevicePropertiesR0600(hipDeviceProp_tR0600* props, hipDevice_t device) { + HIP_INIT_API(hipGetDevicePropertiesR0600, props, device); - HIP_RETURN(ihipGetDeviceProperties(prop, device)); + HIP_RETURN(ihipGetDeviceProperties(props, device)); } -hipError_t hipGetDevicePropertiesR0000(hipDeviceProp_tR0000* prop, int device) { - HIP_INIT_API(hipGetDevicePropertiesR0000, prop, device); +extern "C" typedef struct hipDeviceProp_t { + char name[256]; ///< Device name. + size_t totalGlobalMem; ///< Size of global memory region (in bytes). + size_t sharedMemPerBlock; ///< Size of shared memory region (in bytes). + int regsPerBlock; ///< Registers per block. + int warpSize; ///< Warp size. + int maxThreadsPerBlock; ///< Max work items per work group or workgroup max size. + int maxThreadsDim[3]; ///< Max number of threads in each dimension (XYZ) of a block. + int maxGridSize[3]; ///< Max grid dimensions (XYZ). + int clockRate; ///< Max clock frequency of the multiProcessors in khz. + int memoryClockRate; ///< Max global memory clock frequency in khz. + int memoryBusWidth; ///< Global memory bus width in bits. + size_t totalConstMem; ///< Size of shared memory region (in bytes). + int major; ///< Major compute capability. On HCC, this is an approximation and features may + ///< differ from CUDA CC. See the arch feature flags for portable ways to query + ///< feature caps. + int minor; ///< Minor compute capability. On HCC, this is an approximation and features may + ///< differ from CUDA CC. See the arch feature flags for portable ways to query + ///< feature caps. + int multiProcessorCount; ///< Number of multi-processors (compute units). + int l2CacheSize; ///< L2 cache size. + int maxThreadsPerMultiProcessor; ///< Maximum resident threads per multi-processor. + int computeMode; ///< Compute mode. + int clockInstructionRate; ///< Frequency in khz of the timer used by the device-side "clock*" + ///< instructions. New for HIP. + hipDeviceArch_t arch; ///< Architectural feature flags. New for HIP. + int concurrentKernels; ///< Device can possibly execute multiple kernels concurrently. + int pciDomainID; ///< PCI Domain ID + int pciBusID; ///< PCI Bus ID. + int pciDeviceID; ///< PCI Device ID. + size_t maxSharedMemoryPerMultiProcessor; ///< Maximum Shared Memory Per Multiprocessor. + int isMultiGpuBoard; ///< 1 if device is on a multi-GPU board, 0 if not. + int canMapHostMemory; ///< Check whether HIP can map host memory + int gcnArch; ///< DEPRECATED: use gcnArchName instead + char gcnArchName[256]; ///< AMD GCN Arch Name. + int integrated; ///< APU vs dGPU + int cooperativeLaunch; ///< HIP device supports cooperative launch + int cooperativeMultiDeviceLaunch; ///< HIP device supports cooperative launch on multiple + ///< devices + int maxTexture1DLinear; ///< Maximum size for 1D textures bound to linear memory + int maxTexture1D; ///< Maximum number of elements in 1D images + int maxTexture2D[2]; ///< Maximum dimensions (width, height) of 2D images, in image elements + int maxTexture3D[3]; ///< Maximum dimensions (width, height, depth) of 3D images, in image + ///< elements + unsigned int* hdpMemFlushCntl; ///< Addres of HDP_MEM_COHERENCY_FLUSH_CNTL register + unsigned int* hdpRegFlushCntl; ///< Addres of HDP_REG_COHERENCY_FLUSH_CNTL register + size_t memPitch; ///< Maximum pitch in bytes allowed by memory copies + size_t textureAlignment; ///< Alignment requirement for textures + size_t texturePitchAlignment; ///< Pitch alignment requirement for texture references bound to + ///< pitched memory + int kernelExecTimeoutEnabled; ///< Run time limit for kernels executed on the device + int ECCEnabled; ///< Device has ECC support enabled + int tccDriver; ///< 1:If device is Tesla device using TCC driver, else 0 + int cooperativeMultiDeviceUnmatchedFunc; ///< HIP device supports cooperative launch on + ///< multiple + /// devices with unmatched functions + int cooperativeMultiDeviceUnmatchedGridDim; ///< HIP device supports cooperative launch on + ///< multiple + /// devices with unmatched grid dimensions + int cooperativeMultiDeviceUnmatchedBlockDim; ///< HIP device supports cooperative launch on + ///< multiple + /// devices with unmatched block dimensions + int cooperativeMultiDeviceUnmatchedSharedMem; ///< HIP device supports cooperative launch on + ///< multiple + /// devices with unmatched shared memories + int isLargeBar; ///< 1: if it is a large PCI bar device, else 0 + int asicRevision; ///< Revision of the GPU in this device + int managedMemory; ///< Device supports allocating managed memory on this system + int directManagedMemAccessFromHost; ///< Host can directly access managed memory on the device + ///< without migration + int concurrentManagedAccess; ///< Device can coherently access managed memory concurrently with + ///< the CPU + int pageableMemoryAccess; ///< Device supports coherently accessing pageable memory + ///< without calling hipHostRegister on it + int pageableMemoryAccessUsesHostPageTables; ///< Device accesses pageable memory via the host's + ///< page tables +} hipDeviceProp_t; - if (prop == nullptr) { - HIP_RETURN(hipErrorInvalidValue); +extern "C" hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, hipDevice_t device) { + // Removing this API from tracing. + // This API is now in backwards compatibility mode and is not callable from newly compiled apps. + HIP_INIT_VOID(); + + if (props == nullptr) { + return hipErrorInvalidValue; } if (unsigned(device) >= g_devices.size()) { - HIP_RETURN(hipErrorInvalidDevice); + return hipErrorInvalidDevice; } auto* deviceHandle = g_devices[device]->devices()[0]; constexpr auto int32_max = static_cast(std::numeric_limits::max()); constexpr auto uint16_max = static_cast(std::numeric_limits::max()) + 1; - hipDeviceProp_tR0000 deviceProps = {0}; + hipDeviceProp_t deviceProps = {0}; const auto& info = deviceHandle->info(); const auto& isa = deviceHandle->isa(); @@ -567,11 +646,6 @@ hipError_t hipGetDevicePropertiesR0000(hipDeviceProp_tR0000* prop, int device) { deviceProps.pageableMemoryAccess = info.hmmCpuMemoryAccessible_; deviceProps.pageableMemoryAccessUsesHostPageTables = info.hostUnifiedMemory_; - *prop = deviceProps; - HIP_RETURN(hipSuccess); -} - -extern "C" hipError_t hipGetDeviceProperties(hipDeviceProp_tR0000* props, hipDevice_t device); -hipError_t hipGetDeviceProperties(hipDeviceProp_tR0000* props, hipDevice_t device) { - return hipGetDevicePropertiesR0000(props, device); -} + *props = deviceProps; + return hipSuccess; +} \ No newline at end of file diff --git a/projects/clr/hipamd/src/hip_device_deprecated.cpp b/projects/clr/hipamd/src/hip_device_deprecated.cpp new file mode 100644 index 0000000000..8a4f319b44 --- /dev/null +++ b/projects/clr/hipamd/src/hip_device_deprecated.cpp @@ -0,0 +1,83 @@ +// This file will add older hip functions used in the versioning system +// Find the deprecated functions and structs in hip_device.cpp + +// This struct is also kept in hip_device.cpp +extern "C" typedef struct hipDeviceProp_t { + char name[256]; ///< Device name. + size_t totalGlobalMem; ///< Size of global memory region (in bytes). + size_t sharedMemPerBlock; ///< Size of shared memory region (in bytes). + int regsPerBlock; ///< Registers per block. + int warpSize; ///< Warp size. + int maxThreadsPerBlock; ///< Max work items per work group or workgroup max size. + int maxThreadsDim[3]; ///< Max number of threads in each dimension (XYZ) of a block. + int maxGridSize[3]; ///< Max grid dimensions (XYZ). + int clockRate; ///< Max clock frequency of the multiProcessors in khz. + int memoryClockRate; ///< Max global memory clock frequency in khz. + int memoryBusWidth; ///< Global memory bus width in bits. + size_t totalConstMem; ///< Size of shared memory region (in bytes). + int major; ///< Major compute capability. On HCC, this is an approximation and features may + ///< differ from CUDA CC. See the arch feature flags for portable ways to query + ///< feature caps. + int minor; ///< Minor compute capability. On HCC, this is an approximation and features may + ///< differ from CUDA CC. See the arch feature flags for portable ways to query + ///< feature caps. + int multiProcessorCount; ///< Number of multi-processors (compute units). + int l2CacheSize; ///< L2 cache size. + int maxThreadsPerMultiProcessor; ///< Maximum resident threads per multi-processor. + int computeMode; ///< Compute mode. + int clockInstructionRate; ///< Frequency in khz of the timer used by the device-side "clock*" + ///< instructions. New for HIP. + hipDeviceArch_t arch; ///< Architectural feature flags. New for HIP. + int concurrentKernels; ///< Device can possibly execute multiple kernels concurrently. + int pciDomainID; ///< PCI Domain ID + int pciBusID; ///< PCI Bus ID. + int pciDeviceID; ///< PCI Device ID. + size_t maxSharedMemoryPerMultiProcessor; ///< Maximum Shared Memory Per Multiprocessor. + int isMultiGpuBoard; ///< 1 if device is on a multi-GPU board, 0 if not. + int canMapHostMemory; ///< Check whether HIP can map host memory + int gcnArch; ///< DEPRECATED: use gcnArchName instead + char gcnArchName[256]; ///< AMD GCN Arch Name. + int integrated; ///< APU vs dGPU + int cooperativeLaunch; ///< HIP device supports cooperative launch + int cooperativeMultiDeviceLaunch; ///< HIP device supports cooperative launch on multiple + ///< devices + int maxTexture1DLinear; ///< Maximum size for 1D textures bound to linear memory + int maxTexture1D; ///< Maximum number of elements in 1D images + int maxTexture2D[2]; ///< Maximum dimensions (width, height) of 2D images, in image elements + int maxTexture3D[3]; ///< Maximum dimensions (width, height, depth) of 3D images, in image + ///< elements + unsigned int* hdpMemFlushCntl; ///< Addres of HDP_MEM_COHERENCY_FLUSH_CNTL register + unsigned int* hdpRegFlushCntl; ///< Addres of HDP_REG_COHERENCY_FLUSH_CNTL register + size_t memPitch; ///< Maximum pitch in bytes allowed by memory copies + size_t textureAlignment; ///< Alignment requirement for textures + size_t texturePitchAlignment; ///< Pitch alignment requirement for texture references bound to + ///< pitched memory + int kernelExecTimeoutEnabled; ///< Run time limit for kernels executed on the device + int ECCEnabled; ///< Device has ECC support enabled + int tccDriver; ///< 1:If device is Tesla device using TCC driver, else 0 + int cooperativeMultiDeviceUnmatchedFunc; ///< HIP device supports cooperative launch on + ///< multiple + /// devices with unmatched functions + int cooperativeMultiDeviceUnmatchedGridDim; ///< HIP device supports cooperative launch on + ///< multiple + /// devices with unmatched grid dimensions + int cooperativeMultiDeviceUnmatchedBlockDim; ///< HIP device supports cooperative launch on + ///< multiple + /// devices with unmatched block dimensions + int cooperativeMultiDeviceUnmatchedSharedMem; ///< HIP device supports cooperative launch on + ///< multiple + /// devices with unmatched shared memories + int isLargeBar; ///< 1: if it is a large PCI bar device, else 0 + int asicRevision; ///< Revision of the GPU in this device + int managedMemory; ///< Device supports allocating managed memory on this system + int directManagedMemAccessFromHost; ///< Host can directly access managed memory on the device + ///< without migration + int concurrentManagedAccess; ///< Device can coherently access managed memory concurrently with + ///< the CPU + int pageableMemoryAccess; ///< Device supports coherently accessing pageable memory + ///< without calling hipHostRegister on it + int pageableMemoryAccessUsesHostPageTables; ///< Device accesses pageable memory via the host's + ///< page tables +} hipDeviceProp_t; + +extern "C" hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, hipDevice_t device); diff --git a/projects/clr/hipamd/src/hip_device_runtime.cpp b/projects/clr/hipamd/src/hip_device_runtime.cpp index 7e7cc53882..a728a8f946 100644 --- a/projects/clr/hipamd/src/hip_device_runtime.cpp +++ b/projects/clr/hipamd/src/hip_device_runtime.cpp @@ -22,33 +22,23 @@ #include "hip_internal.hpp" -#undef hipChooseDevice -#undef hipDeviceProp_t +hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* properties) { + HIP_INIT_API(hipChooseDevice, device, properties); -template -hipError_t ihipChooseDevice(int* device, const DeviceProp* properties) { if (device == nullptr || properties == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } *device = 0; cl_uint maxMatchedCount = 0; int count = 0; - IHIP_RETURN_ONFAIL(ihipDeviceGetCount(&count)); + HIP_RETURN_ONFAIL(ihipDeviceGetCount(&count)); for (cl_int i = 0; i < count; ++i) { - DeviceProp currentProp = {0}; + hipDeviceProp_t currentProp = {0}; cl_uint validPropCount = 0; cl_uint matchedCount = 0; - hipError_t err = hipSuccess; - - if constexpr (std::is_same_v){ - err = ihipGetDeviceProperties(¤tProp, i); - } - else { - err = hipGetDevicePropertiesR0000(¤tProp, i); - } - + hipError_t err = ihipGetDeviceProperties(¤tProp, i); if (properties->major != 0) { validPropCount++; if (currentProp.major >= properties->major) { @@ -140,25 +130,9 @@ hipError_t ihipChooseDevice(int* device, const DeviceProp* properties) { } } - return hipSuccess; -} - -hipError_t hipChooseDeviceR0600(int* device, const hipDeviceProp_tR0600* properties) { - HIP_INIT_API(hipChooseDeviceR0600, device, properties); - HIP_RETURN(ihipChooseDevice(device, properties)); -} - -hipError_t hipChooseDeviceR0000(int* device, const hipDeviceProp_tR0000* properties) { - HIP_INIT_API(hipChooseDeviceR0000, device, properties); - HIP_RETURN(ihipChooseDevice(device, properties)); HIP_RETURN(hipSuccess); } -extern "C" hipError_t hipChooseDevice(int* device, const hipDeviceProp_tR0000* properties); -hipError_t hipChooseDevice(int* device, const hipDeviceProp_tR0000* properties) { - return hipChooseDeviceR0000(device, properties); -} - hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) { HIP_INIT_API(hipDeviceGetAttribute, pi, attr, device); @@ -174,7 +148,7 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) } // FIXME: should we cache the props, or just select from deviceHandle->info_? - hipDeviceProp_tR0600 prop = {0}; + hipDeviceProp_t prop = {0}; HIP_RETURN_ONFAIL(ihipGetDeviceProperties(&prop, device)); constexpr auto int32_max = static_cast(std::numeric_limits::max()); @@ -466,7 +440,7 @@ hipError_t hipDeviceGetByPCIBusId(int* device, const char* pciBusIdstr) { HIP_RETURN_ONFAIL(ihipDeviceGetCount(&count)); for (cl_int i = 0; i < count; i++) { hipDevice_t dev; - hipDeviceProp_tR0600 prop; + hipDeviceProp_t prop; HIP_RETURN_ONFAIL(ihipDeviceGet(&dev, i)); HIP_RETURN_ONFAIL(ihipGetDeviceProperties(&prop, dev)); @@ -506,7 +480,7 @@ hipError_t hipDeviceGetLimit(size_t* pValue, hipLimit_t limit) { switch (limit) { case hipLimitMallocHeapSize: - hipDeviceProp_tR0600 prop; + hipDeviceProp_t prop; HIP_RETURN_ONFAIL(ihipGetDeviceProperties(&prop, ihipGetDevice())); *pValue = prop.totalGlobalMem; break; @@ -535,7 +509,7 @@ hipError_t hipDeviceGetPCIBusId(char* pciBusId, int len, int device) { HIP_RETURN(hipErrorInvalidValue); } - hipDeviceProp_tR0600 prop; + hipDeviceProp_t prop; HIP_RETURN_ONFAIL(ihipGetDeviceProperties(&prop, device)); snprintf(pciBusId, len, "%04x:%02x:%02x.0", prop.pciDomainID, prop.pciBusID, prop.pciDeviceID); diff --git a/projects/clr/hipamd/src/hip_find_defs.py b/projects/clr/hipamd/src/hip_find_defs.py new file mode 100755 index 0000000000..d98d5bbdf3 --- /dev/null +++ b/projects/clr/hipamd/src/hip_find_defs.py @@ -0,0 +1,83 @@ +import getopt, sys, os + +def write_new_header(): + arg_list = sys.argv[1:] # Files to read is dictated by arguments + optlist, files_to_read = getopt.getopt(arg_list, "od", ["output=", "deprecated="]) + + write_header = '' + deprecated_functions = '' + for arg, value in optlist: + if arg in ["-o", "--output"]: + write_header = value + elif arg in ["-p", "--deprecated"]: + deprecated_functions = value + + print(optlist) + if len(write_header) == 0: + print("hip_find_defs.py Command Line argument parsing incorrectly!") + return + + new_header_file = open(write_header, 'w') + + version_define_map = {} + + struct_mode = False + struct_string = '' + struct_name = '' + struct_depth = 0 + + for the_file in files_to_read: + header = open(the_file, 'r') + header_lines = header.readlines() + for line in header_lines: + #reading a struct + if struct_mode: + struct_string += line + if '{' in line: + struct_depth += 1 + elif '}' in line: + struct_depth -= 1 + + if struct_depth == 0: + struct_mode = False + if struct_name in version_define_map: + #new_header_file.write('\n') + #new_header_file.write(getOlderStruct(struct_name, version_define_map[struct_name], hip_device)) + new_header_file.write(struct_string.replace(struct_name, version_define_map[struct_name])) + else: + new_header_file.write(struct_string) + continue + + #finding defines used for versioning + if "#define" in line: + line_split = line.split() + if len(line_split) == 3 and line_split[1] in line_split[2] and line_split[2][-1].isnumeric(): + version_define_map[line_split[1]] = line_split[2] + continue + + #Looking for struct + if "typedef struct" in line and '{' in line: + struct_mode = True + struct_string = line + struct_depth = 1 + struct_name = line.replace('{', '').split()[-1] + continue + + #Looking for a typical function signature + if '(' in line and ')' in line and len(line.split('(')[0].split(' ')) == 2: + function_name = line.split('(')[0].split(' ')[1] + #If this function is one of the version functions, write the versioned function too + if function_name in version_define_map: + duplicate_line = line.replace(function_name, version_define_map[function_name]) + new_header_file.write(duplicate_line) + continue + new_header_file.write(line) + header.close() + + if os.path.exists(deprecated_functions): + deprecated_file = open(deprecated_functions, 'r') + new_header_file.write(deprecated_file.read()) + + new_header_file.close() + +write_new_header() diff --git a/projects/clr/hipamd/src/hip_gl.cpp b/projects/clr/hipamd/src/hip_gl.cpp index 22e961fbfa..5094085819 100644 --- a/projects/clr/hipamd/src/hip_gl.cpp +++ b/projects/clr/hipamd/src/hip_gl.cpp @@ -20,7 +20,6 @@ #include "top.hpp" #include "hip/hip_runtime.h" -#include "hip/hip_gl_interop.h" #include "hip_internal.hpp" #include "platform/interop_gl.hpp" #include "cl_common.hpp" diff --git a/projects/clr/hipamd/src/hip_hcc.map.in b/projects/clr/hipamd/src/hip_hcc.map.in index 58e2311add..3e92924d91 100644 --- a/projects/clr/hipamd/src/hip_hcc.map.in +++ b/projects/clr/hipamd/src/hip_hcc.map.in @@ -1,7 +1,6 @@ hip_4.2 { global: hipChooseDevice; - hipChooseDeviceR0000; hipCtxCreate; hipCtxDestroy; hipCtxDisablePeerAccess; @@ -63,7 +62,6 @@ global: hipGetDevice; hipGetDeviceCount; hipGetDeviceProperties; - hipGetDevicePropertiesR0000; hipGetErrorName; hipGetErrorString; hipGetLastError; @@ -532,8 +530,7 @@ local: hip_6.0 { global: - hipChooseDeviceR0600; hipGetDevicePropertiesR0600; local: *; -} hip_5.6; +} hip_5.6; \ No newline at end of file diff --git a/projects/clr/hipamd/src/hip_prof_gen.py b/projects/clr/hipamd/src/hip_prof_gen.py index 492ae46107..6dc5247609 100755 --- a/projects/clr/hipamd/src/hip_prof_gen.py +++ b/projects/clr/hipamd/src/hip_prof_gen.py @@ -393,10 +393,6 @@ def generate_prof_header(f, api_map, callback_ids, opts_map): f.write('#define _HIP_PROF_STR_H\n'); f.write('#define HIP_PROF_VER 1\n') - f.write('\n#include \n') - f.write('#include \n') - f.write('#include "amd_hip_gl_interop.h"\n') - # Check for non-public API for name in sorted(opts_map.keys()): if not name in api_map: @@ -406,9 +402,6 @@ def generate_prof_header(f, api_map, callback_ids, opts_map): priv_lst.append(name) message("Private: " + name) - f.write('\n#define HIP_API_ID_CONCAT_HELPER(a,b) a##b\n'); - f.write('#define HIP_API_ID_CONCAT(a,b) HIP_API_ID_CONCAT_HELPER(a,b)\n'); - # Generating the callbacks ID enumaration f.write('\n// HIP API callbacks ID enumeration\n') f.write('enum hip_api_id_t {\n') @@ -417,7 +410,6 @@ def generate_prof_header(f, api_map, callback_ids, opts_map): cb_id_map = {} last_cb_id = 0 - versioned_functions = set() for name, cb_id in callback_ids: if not name in api_map: f.write(' HIP_API_ID_RESERVED_' + str(cb_id) + ' = ' + str(cb_id) + ',\n') @@ -425,30 +417,18 @@ def generate_prof_header(f, api_map, callback_ids, opts_map): f.write(' HIP_API_ID_' + name + ' = ' + str(cb_id) + ',\n') cb_id_map[name] = cb_id if cb_id > last_cb_id: last_cb_id = cb_id - m = re.match(r'(.*)R[0-9][0-9][0-9][0-9]$', name) - if m: versioned_functions.add(m.group(1)) for name in sorted(api_map.keys()): if not name in cb_id_map: last_cb_id += 1 f.write(' HIP_API_ID_' + name + ' = ' + str(last_cb_id) + ',\n') - m = re.match(r'(.*)R[0-9][0-9][0-9][0-9]$', name) - if m: versioned_functions.add(m.group(1)) f.write(' HIP_API_ID_LAST = ' + str(last_cb_id) + ',\n') f.write('\n') - - for name in sorted(versioned_functions): - f.write(' HIP_API_ID_' + name + ' = ' + 'HIP_API_ID_CONCAT(HIP_API_ID_,' + name + '),\n') - f.write('\n') - for name in sorted(priv_lst): f.write(' HIP_API_ID_' + name + ' = HIP_API_ID_NONE,\n') f.write('};\n') - f.write('\n#undef HIP_API_ID_CONCAT_HELPER\n'); - f.write('#undef HIP_API_ID_CONCAT\n'); - # Generating the method to return API name by ID f.write('\n// Return the HIP API string for a given callback ID\n') f.write('static inline const char* hip_api_name(const uint32_t id) {\n') @@ -678,6 +658,7 @@ api_map = { 'hipGetErrorString': '', 'hipMallocHost': '', 'hipModuleLoadDataEx': '', + 'hipGetDeviceProperties': '', 'hipConfigureCall': '', 'hipHccModuleLaunchKernel': '', 'hipExtModuleLaunchKernel': '', @@ -707,8 +688,6 @@ for enum in cppHeader.enums: continue if value['name'] == 'HIP_API_ID_LAST': break - if type(value['value']) == str: - continue m = re.match(r'HIP_API_ID_(\S*)', value['name']) if m: api_callback_ids.append((m.group(1), value['value']))