AI/rocm-systems
2
0
Derivar 0
Código Questões Pedidos de integração Operações Pacotes Planeamentos Lançamentos Wiki Trabalho

SWDEV-427855 - hipamd change for profiler and TF fix

Ver a proveniência 
This reverts commit 57cb840058.

Change-Id: Id69e47a1afd336ae1edb9c8e173be27e7b9dcc8d
Este cometimento está contido em:
Rahul Garg
2023-11-29 15:50:10 -05:00
ascendente 3c6505c2d5
cometimento afc28b091e
10 ficheiros modificados com 539 adições e 381 eliminações
Descarregar ficheiro patch Descarregar ficheiro diff Expandir todos os ficheiros Colapsar todos os ficheiros
hipamd/include/hip/amd_detail/hip_prof_str.h
+435 -102
Ver ficheiro
@@ -7,7 +7,12 @@
#define _HIP_PROF_STR_H
#define HIP_PROF_VER 1
#include <hip/amd_detail/amd_hip_gl_interop.h>
#include <hip/hip_runtime_api.h>
#include <hip/hip_deprecated.h>
#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());
hipamd/src/CMakeLists.txt
+18 -8
Ver ficheiro
@@ -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$<JOIN:$<TARGET_PROPERTY:amdhip64,COMPILE_DEFINITIONS>,;-D>"
"-I$<JOIN:$<TARGET_PROPERTY:amdhip64,INCLUDE_DIRECTORIES>,;-I>"
"-DHIP_INCLUDE_HIP_HIP_RUNTIME_PT_API_H=1"
${c_flags}
$<TARGET_PROPERTY:amdhip64,COMPILE_OPTIONS>
${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})
hipamd/src/amdhip.def
+3
Ver ficheiro
@@ -1,5 +1,7 @@
EXPORTS
hipChooseDevice
hipChooseDeviceR0000
hipChooseDeviceR0600
hipCtxCreate
hipCtxDestroy
hipCtxDisablePeerAccess
@@ -63,6 +65,7 @@ hipFuncSetSharedMemConfig
hipGetDevice
hipGetDeviceCount
hipGetDeviceProperties
hipGetDevicePropertiesR0000
hipGetDevicePropertiesR0600
hipGetErrorName
hipGetErrorString
hipamd/src/hip_device.cpp
+19 -92
Ver ficheiro
@@ -19,6 +19,7 @@
THE SOFTWARE. */
#include <hip/hip_runtime.h>
#include <hip/hip_deprecated.h>
#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<uint64_t>(std::numeric_limits<int32_t>::max());
constexpr auto uint16_max = static_cast<uint64_t>(std::numeric_limits<uint16_t>::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
hipamd/src/hip_device_deprecated.cpp
-83
Ver ficheiro
@@ -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);
hipamd/src/hip_device_runtime.cpp
+37 -11
Ver ficheiro
@@ -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 <typename DeviceProp>
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(&currentProp, i);
hipError_t err = hipSuccess;
if constexpr (std::is_same_v<DeviceProp, hipDeviceProp_tR0600>){
err = ihipGetDeviceProperties(&currentProp, i);
}
else {
err = hipGetDevicePropertiesR0000(&currentProp, 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<uint64_t>(std::numeric_limits<int32_t>::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);
hipamd/src/hip_find_defs.py
-83
Ver ficheiro
@@ -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()
hipamd/src/hip_gl.cpp
+1
Ver ficheiro
@@ -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"
hipamd/src/hip_hcc.map.in
+4 -1
Ver ficheiro
@@ -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;
} hip_5.6;
hipamd/src/hip_prof_gen.py
+22 -1
Ver ficheiro
@@ -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 <hip/hip_runtime_api.h>\n')
f.write('#include <hip/hip_deprecated.h>\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']))