diff --git a/hipamd/include/hip/amd_detail/hip_prof_str.h b/hipamd/include/hip/amd_detail/hip_prof_str.h index c67a8886e0..128dc3747e 100644 --- a/hipamd/include/hip/amd_detail/hip_prof_str.h +++ b/hipamd/include/hip/amd_detail/hip_prof_str.h @@ -7,7 +7,12 @@ #define _HIP_PROF_STR_H #define HIP_PROF_VER 1 -#include +#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) // HIP API callbacks ID enumeration enum hip_api_id_t { @@ -18,7 +23,7 @@ enum hip_api_id_t { HIP_API_ID_hipArray3DCreate = 3, HIP_API_ID_hipArrayCreate = 4, HIP_API_ID_hipArrayDestroy = 5, - HIP_API_ID_hipChooseDevice = 6, + HIP_API_ID_hipChooseDeviceR0000 = 6, HIP_API_ID_hipConfigureCall = 7, HIP_API_ID_hipCtxCreate = 8, HIP_API_ID_hipCtxDestroy = 9, @@ -93,7 +98,7 @@ enum hip_api_id_t { HIP_API_ID_hipGetDevice = 78, HIP_API_ID_hipGetDeviceCount = 79, HIP_API_ID_hipGetDeviceFlags = 80, - HIP_API_ID_hipGetDevicePropertiesR0600 = 81, + HIP_API_ID_hipGetDevicePropertiesR0000 = 81, HIP_API_ID_RESERVED_82 = 82, HIP_API_ID_hipGetErrorString = 83, HIP_API_ID_hipGetLastError = 84, @@ -377,8 +382,24 @@ enum hip_api_id_t { HIP_API_ID_hipArrayGetInfo = 362, HIP_API_ID_hipStreamGetDevice = 363, HIP_API_ID_hipExternalMemoryGetMappedMipmappedArray = 364, - HIP_API_ID_hipDrvGraphAddMemcpyNode = 365, - HIP_API_ID_LAST = 365, + 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_hipBindTexture = HIP_API_ID_NONE, HIP_API_ID_hipBindTexture2D = HIP_API_ID_NONE, @@ -393,14 +414,11 @@ 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, @@ -419,6 +437,9 @@ 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) { @@ -430,7 +451,8 @@ 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_hipChooseDevice: return "hipChooseDevice"; + case HIP_API_ID_hipChooseDeviceR0000: return "hipChooseDeviceR0000"; + case HIP_API_ID_hipChooseDeviceR0600: return "hipChooseDeviceR0600"; case HIP_API_ID_hipConfigureCall: return "hipConfigureCall"; case HIP_API_ID_hipCreateSurfaceObject: return "hipCreateSurfaceObject"; case HIP_API_ID_hipCtxCreate: return "hipCtxCreate"; @@ -485,6 +507,10 @@ 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"; @@ -504,6 +530,7 @@ 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"; @@ -519,6 +546,7 @@ 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"; @@ -530,12 +558,13 @@ 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"; @@ -554,6 +583,8 @@ 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"; @@ -562,6 +593,10 @@ 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"; @@ -782,7 +817,6 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipUserObjectRelease: return "hipUserObjectRelease"; case HIP_API_ID_hipUserObjectRetain: return "hipUserObjectRetain"; case HIP_API_ID_hipWaitExternalSemaphoresAsync: return "hipWaitExternalSemaphoresAsync"; - case HIP_API_ID_hipExternalMemoryGetMappedMipmappedArray: return "hipExternalMemoryGetMappedMipmappedArray"; }; return "unknown"; }; @@ -798,7 +832,8 @@ 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("hipChooseDevice", name) == 0) return HIP_API_ID_hipChooseDevice; + if (strcmp("hipChooseDeviceR0000", name) == 0) return HIP_API_ID_hipChooseDeviceR0000; + if (strcmp("hipChooseDeviceR0600", name) == 0) return HIP_API_ID_hipChooseDeviceR0600; 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; @@ -853,6 +888,10 @@ 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; @@ -872,6 +911,7 @@ 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; @@ -887,6 +927,7 @@ 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; @@ -898,6 +939,8 @@ 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; @@ -921,6 +964,8 @@ 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; @@ -929,6 +974,10 @@ 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; @@ -1149,8 +1198,6 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipUserObjectRelease", name) == 0) return HIP_API_ID_hipUserObjectRelease; if (strcmp("hipUserObjectRetain", name) == 0) return HIP_API_ID_hipUserObjectRetain; if (strcmp("hipWaitExternalSemaphoresAsync", name) == 0) return HIP_API_ID_hipWaitExternalSemaphoresAsync; - if (strcmp("hipExternalMemoryGetMappedMipmappedArray", name) == 0) return HIP_API_ID_hipExternalMemoryGetMappedMipmappedArray; - if (strcmp("hipDrvGraphAddMemcpyNode", name) == 0) return HIP_API_ID_hipDrvGraphAddMemcpyNode; return HIP_API_ID_NONE; } @@ -1209,12 +1256,18 @@ 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; - } hipChooseDevice; + } hipChooseDeviceR0600; struct { dim3 gridDim; dim3 blockDim; @@ -1441,6 +1494,38 @@ 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; @@ -1560,6 +1645,13 @@ 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; @@ -1626,9 +1718,14 @@ typedef struct hip_api_data_s { unsigned int flags__val; } hipGetDeviceFlags; struct { - hipDeviceProp_tR0600* props; - hipDeviceProp_tR0600 props__val; - hipDevice_t device; + hipDeviceProp_tR0000* prop; + hipDeviceProp_tR0000 prop__val; + int device; + } hipGetDevicePropertiesR0000; + struct { + hipDeviceProp_tR0600* prop; + hipDeviceProp_tR0600 prop__val; + int deviceId; } hipGetDevicePropertiesR0600; struct { hipArray_t* levelArray; @@ -1689,6 +1786,26 @@ 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; @@ -1849,6 +1966,18 @@ 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; @@ -1907,6 +2036,26 @@ 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; @@ -3262,24 +3411,6 @@ 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; @@ -3330,10 +3461,15 @@ typedef struct hip_api_data_s { cb_data.args.hipArrayGetInfo.flags = (unsigned int*)flags; \ cb_data.args.hipArrayGetInfo.array = (hipArray_t)array; \ }; -// 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; \ +// 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; \ }; // hipConfigureCall[('dim3', 'gridDim'), ('dim3', 'blockDim'), ('size_t', 'sharedMem'), ('hipStream_t', 'stream')] #define INIT_hipConfigureCall_CB_ARGS_DATA(cb_data) { \ @@ -3588,6 +3724,18 @@ 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; \ @@ -3708,10 +3856,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; \ @@ -3784,10 +3932,15 @@ typedef struct hip_api_data_s { #define INIT_hipGetDeviceFlags_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipGetDeviceFlags.flags = (unsigned int*)flags; \ }; -// hipGetDeviceProperties[('hipDeviceProp_tR0600*', 'props'), ('hipDevice_t', 'device')] +// 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')] #define INIT_hipGetDevicePropertiesR0600_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipGetDevicePropertiesR0600.props = (hipDeviceProp_tR0600*)props; \ - cb_data.args.hipGetDevicePropertiesR0600.device = (hipDevice_t)device; \ + cb_data.args.hipGetDevicePropertiesR0600.prop = (hipDeviceProp_tR0600*)prop; \ + cb_data.args.hipGetDevicePropertiesR0600.deviceId = (int)device; \ }; // hipGetErrorString[] #define INIT_hipGetErrorString_CB_ARGS_DATA(cb_data) { \ @@ -3849,6 +4002,12 @@ 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; \ @@ -4003,6 +4162,12 @@ 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; \ @@ -4063,6 +4228,18 @@ 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; \ @@ -5453,15 +5630,6 @@ 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 @@ -5491,8 +5659,6 @@ 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() @@ -5501,12 +5667,8 @@ 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() @@ -5583,10 +5745,15 @@ 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; -// 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); +// 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); break; // hipConfigureCall[('dim3', 'gridDim'), ('dim3', 'blockDim'), ('size_t', 'sharedMem'), ('hipStream_t', 'stream')] case HIP_API_ID_hipConfigureCall: @@ -5783,6 +5950,26 @@ 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); @@ -5862,6 +6049,7 @@ 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')] @@ -5917,9 +6105,13 @@ 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; -// hipGetDevicePropertiesR0600[('hipDeviceProp_tR0600*', 'props'), ('hipDevice_t', 'device')] +// 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')] case HIP_API_ID_hipGetDevicePropertiesR0600: - if (data->args.hipGetDevicePropertiesR0600.props) data->args.hipGetDevicePropertiesR0600.props__val = *(data->args.hipGetDevicePropertiesR0600.props); + if (data->args.hipGetDevicePropertiesR0600.prop) data->args.hipGetDevicePropertiesR0600.prop__val = *(data->args.hipGetDevicePropertiesR0600.prop); break; // hipGetErrorString[] case HIP_API_ID_hipGetErrorString: @@ -5964,6 +6156,18 @@ 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); @@ -6062,6 +6266,14 @@ 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); @@ -6092,6 +6304,22 @@ 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); @@ -6941,12 +7169,6 @@ 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; }; } @@ -6980,7 +7202,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, (void*)data->args.hipArray3DCreate.array__val); } + else { oss << "array="; roctracer::hip_support::detail::operator<<(oss, 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 << ")"; @@ -6989,13 +7211,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, (void*)data->args.hipArrayCreate.pHandle__val); } + else { oss << "pHandle="; roctracer::hip_support::detail::operator<<(oss, 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 << ")"; @@ -7020,16 +7242,23 @@ 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_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); } + 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); } oss << ")"; break; case HIP_API_ID_hipConfigureCall: @@ -7372,6 +7601,46 @@ 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"; @@ -7531,12 +7800,13 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da break; case HIP_API_ID_hipExternalMemoryGetMappedMipmappedArray: oss << "hipExternalMemoryGetMappedMipmappedArray("; - oss << "mipmap="; roctracer::hip_support::detail::operator<<(oss, data->args.hipExternalMemoryGetMappedMipmappedArray.mipmap); + if (data->args.hipExternalMemoryGetMappedMipmappedArray.mipmap == NULL) oss << "mipmap=NULL"; + else { oss << "mipmap="; roctracer::hip_support::detail::operator<<(oss, data->args.hipExternalMemoryGetMappedMipmappedArray.mipmap__val); } 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); @@ -7632,11 +7902,18 @@ 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.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); + 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); oss << ")"; break; case HIP_API_ID_hipGetErrorString: @@ -7722,6 +7999,30 @@ 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"; @@ -7928,6 +8229,22 @@ 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); @@ -8002,6 +8319,34 @@ 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); @@ -8503,7 +8848,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, (void*)data->args.hipMallocArray.array__val); } + else { oss << "array="; roctracer::hip_support::detail::operator<<(oss, 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); @@ -9803,18 +10148,6 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipWaitExternalSemaphoresAsync.stream); oss << ")"; break; - case HIP_API_ID_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/hipamd/src/CMakeLists.txt b/hipamd/src/CMakeLists.txt index 2275fa342d..116c4d459a 100644 --- a/hipamd/src/CMakeLists.txt +++ b/hipamd/src/CMakeLists.txt @@ -214,12 +214,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" @@ -233,19 +233,29 @@ if(USE_PROF_API) ") endif() - 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} + 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} 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} ${PROF_API_SRC} ${PROF_API_STR_IN} ${PROF_API_STR} - DEPENDS ${PROF_API_STR_IN} ${PROF_API_NEWHDR} ${PROF_API_GEN} + 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} COMMENT "Generating profiling primitives: ${PROF_API_STR}") add_custom_target(gen-prof-api-str-header ALL DEPENDS ${PROF_API_STR} - SOURCES ${PROF_API_NEWHDR}) + SOURCES ${PROF_API_NEWHDR}.i) set_target_properties(amdhip64 PROPERTIES PUBLIC_HEADER ${PROF_API_STR}) diff --git a/hipamd/src/amdhip.def b/hipamd/src/amdhip.def index c0e5a41b35..490e44c85e 100644 --- a/hipamd/src/amdhip.def +++ b/hipamd/src/amdhip.def @@ -1,5 +1,7 @@ EXPORTS hipChooseDevice +hipChooseDeviceR0000 +hipChooseDeviceR0600 hipCtxCreate hipCtxDestroy hipCtxDisablePeerAccess @@ -63,6 +65,7 @@ hipFuncSetSharedMemConfig hipGetDevice hipGetDeviceCount hipGetDeviceProperties +hipGetDevicePropertiesR0000 hipGetDevicePropertiesR0600 hipGetErrorName hipGetErrorString diff --git a/hipamd/src/hip_device.cpp b/hipamd/src/hip_device.cpp index e0e34d9f33..b617d2cf0b 100644 --- a/hipamd/src/hip_device.cpp +++ b/hipamd/src/hip_device.cpp @@ -19,6 +19,7 @@ THE SOFTWARE. */ #include +#include #include "hip_internal.hpp" #include "hip_mempool_impl.hpp" @@ -273,7 +274,7 @@ hipError_t hipDeviceGetUuid(hipUUID* uuid, hipDevice_t device) { HIP_RETURN(hipSuccess); } -hipError_t ihipGetDeviceProperties(hipDeviceProp_tR0600* props, hipDevice_t device) { +hipError_t ihipGetDeviceProperties(hipDeviceProp_tR0600* props, int device) { if (props == nullptr) { return hipErrorInvalidValue; } @@ -457,107 +458,27 @@ hipError_t ihipGetDeviceProperties(hipDeviceProp_tR0600* props, hipDevice_t devi return hipSuccess; } -hipError_t hipGetDevicePropertiesR0600(hipDeviceProp_tR0600* props, hipDevice_t device) { - HIP_INIT_API(hipGetDevicePropertiesR0600, props, device); +hipError_t hipGetDevicePropertiesR0600(hipDeviceProp_tR0600* prop, int device) { + HIP_INIT_API(hipGetDevicePropertiesR0600, prop, device); - HIP_RETURN(ihipGetDeviceProperties(props, device)); + HIP_RETURN(ihipGetDeviceProperties(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; +hipError_t hipGetDevicePropertiesR0000(hipDeviceProp_tR0000* prop, int device) { + HIP_INIT_API(hipGetDevicePropertiesR0000, prop, device); -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 (prop == nullptr) { + HIP_RETURN(hipErrorInvalidValue); } if (unsigned(device) >= g_devices.size()) { - return hipErrorInvalidDevice; + HIP_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_t deviceProps = {0}; + hipDeviceProp_tR0000 deviceProps = {0}; const auto& info = deviceHandle->info(); const auto& isa = deviceHandle->isa(); @@ -644,7 +565,13 @@ extern "C" hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, hipDevice_t deviceProps.pageableMemoryAccess = info.hmmCpuMemoryAccessible_; deviceProps.pageableMemoryAccessUsesHostPageTables = info.hostUnifiedMemory_; - *props = deviceProps; - return hipSuccess; + *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); } } // namespace hip + diff --git a/hipamd/src/hip_device_deprecated.cpp b/hipamd/src/hip_device_deprecated.cpp deleted file mode 100644 index 8a4f319b44..0000000000 --- a/hipamd/src/hip_device_deprecated.cpp +++ /dev/null @@ -1,83 +0,0 @@ -// 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/hipamd/src/hip_device_runtime.cpp b/hipamd/src/hip_device_runtime.cpp index 7ece571b4a..026dafa2c3 100644 --- a/hipamd/src/hip_device_runtime.cpp +++ b/hipamd/src/hip_device_runtime.cpp @@ -22,25 +22,35 @@ #include "hip_internal.hpp" +#undef hipChooseDevice +#undef hipDeviceProp_t + namespace hip { -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) { - HIP_RETURN(hipErrorInvalidValue); + return hipErrorInvalidValue; } *device = 0; cl_uint maxMatchedCount = 0; int count = 0; - HIP_RETURN_ONFAIL(ihipDeviceGetCount(&count)); + IHIP_RETURN_ONFAIL(ihipDeviceGetCount(&count)); for (cl_int i = 0; i < count; ++i) { - hipDeviceProp_t currentProp = {0}; + DeviceProp currentProp = {0}; cl_uint validPropCount = 0; cl_uint matchedCount = 0; - hipError_t err = ihipGetDeviceProperties(¤tProp, i); + hipError_t err = hipSuccess; + + if constexpr (std::is_same_v){ + err = ihipGetDeviceProperties(¤tProp, i); + } + else { + err = hipGetDevicePropertiesR0000(¤tProp, i); + } + if (properties->major != 0) { validPropCount++; if (currentProp.major >= properties->major) { @@ -132,9 +142,25 @@ hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* 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); @@ -150,7 +176,7 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) } // FIXME: should we cache the props, or just select from deviceHandle->info_? - hipDeviceProp_t prop = {0}; + hipDeviceProp_tR0600 prop = {0}; HIP_RETURN_ONFAIL(ihipGetDeviceProperties(&prop, device)); constexpr auto int32_max = static_cast(std::numeric_limits::max()); @@ -442,7 +468,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_t prop; + hipDeviceProp_tR0600 prop; HIP_RETURN_ONFAIL(ihipDeviceGet(&dev, i)); HIP_RETURN_ONFAIL(ihipGetDeviceProperties(&prop, dev)); @@ -482,7 +508,7 @@ hipError_t hipDeviceGetLimit(size_t* pValue, hipLimit_t limit) { switch (limit) { case hipLimitMallocHeapSize: - hipDeviceProp_t prop; + hipDeviceProp_tR0600 prop; HIP_RETURN_ONFAIL(ihipGetDeviceProperties(&prop, ihipGetDevice())); *pValue = prop.totalGlobalMem; break; @@ -511,7 +537,7 @@ hipError_t hipDeviceGetPCIBusId(char* pciBusId, int len, int device) { HIP_RETURN(hipErrorInvalidValue); } - hipDeviceProp_t prop; + hipDeviceProp_tR0600 prop; HIP_RETURN_ONFAIL(ihipGetDeviceProperties(&prop, device)); snprintf(pciBusId, len, "%04x:%02x:%02x.0", prop.pciDomainID, prop.pciBusID, prop.pciDeviceID); diff --git a/hipamd/src/hip_find_defs.py b/hipamd/src/hip_find_defs.py deleted file mode 100755 index d98d5bbdf3..0000000000 --- a/hipamd/src/hip_find_defs.py +++ /dev/null @@ -1,83 +0,0 @@ -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/hipamd/src/hip_gl.cpp b/hipamd/src/hip_gl.cpp index 07b25dfa8b..3cdab56bd3 100644 --- a/hipamd/src/hip_gl.cpp +++ b/hipamd/src/hip_gl.cpp @@ -20,6 +20,7 @@ #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/hipamd/src/hip_hcc.map.in b/hipamd/src/hip_hcc.map.in index 1bec29f25c..ce160ccaff 100644 --- a/hipamd/src/hip_hcc.map.in +++ b/hipamd/src/hip_hcc.map.in @@ -1,6 +1,7 @@ hip_4.2 { global: hipChooseDevice; + hipChooseDeviceR0000; hipCtxCreate; hipCtxDestroy; hipCtxDisablePeerAccess; @@ -62,6 +63,7 @@ global: hipGetDevice; hipGetDeviceCount; hipGetDeviceProperties; + hipGetDevicePropertiesR0000; hipGetErrorName; hipGetErrorString; hipGetLastError; @@ -530,7 +532,8 @@ local: hip_6.0 { global: + hipChooseDeviceR0600; hipGetDevicePropertiesR0600; local: *; -} hip_5.6; \ No newline at end of file +} hip_5.6; diff --git a/hipamd/src/hip_prof_gen.py b/hipamd/src/hip_prof_gen.py index 9085272a5e..7949d4c2fc 100755 --- a/hipamd/src/hip_prof_gen.py +++ b/hipamd/src/hip_prof_gen.py @@ -398,6 +398,10 @@ 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: @@ -407,6 +411,9 @@ 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') @@ -415,6 +422,7 @@ 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') @@ -422,18 +430,30 @@ 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') @@ -663,7 +683,6 @@ api_map = { 'hipGetErrorString': '', 'hipMallocHost': '', 'hipModuleLoadDataEx': '', - 'hipGetDeviceProperties': '', 'hipConfigureCall': '', 'hipHccModuleLaunchKernel': '', 'hipExtModuleLaunchKernel': '', @@ -693,6 +712,8 @@ 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']))