diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h index 355c3c90d3..68973cc659 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h @@ -364,7 +364,13 @@ enum hip_api_id_t { HIP_API_ID_hipGraphNodeGetEnabled = 351, HIP_API_ID_hipGraphNodeSetEnabled = 352, HIP_API_ID_hipPointerSetAttribute = 353, - HIP_API_ID_LAST = 353, + HIP_API_ID_hipGraphAddMemAllocNode = 354, + HIP_API_ID_hipGraphAddMemFreeNode = 355, + HIP_API_ID_hipGraphMemAllocNodeGetParams = 356, + HIP_API_ID_hipGraphMemFreeNodeGetParams = 357, + HIP_API_ID_hipModuleLaunchCooperativeKernel = 358, + HIP_API_ID_hipModuleLaunchCooperativeKernelMultiDevice = 359, + HIP_API_ID_LAST = 359, HIP_API_ID_hipArray3DGetDescriptor = HIP_API_ID_NONE, HIP_API_ID_hipArrayGetDescriptor = HIP_API_ID_NONE, @@ -518,6 +524,8 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipGraphAddEventWaitNode: return "hipGraphAddEventWaitNode"; 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_hipGraphAddMemcpyNode1D: return "hipGraphAddMemcpyNode1D"; case HIP_API_ID_hipGraphAddMemcpyNodeFromSymbol: return "hipGraphAddMemcpyNodeFromSymbol"; @@ -558,6 +566,8 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipGraphKernelNodeSetAttribute: return "hipGraphKernelNodeSetAttribute"; case HIP_API_ID_hipGraphKernelNodeSetParams: return "hipGraphKernelNodeSetParams"; case HIP_API_ID_hipGraphLaunch: return "hipGraphLaunch"; + case HIP_API_ID_hipGraphMemAllocNodeGetParams: return "hipGraphMemAllocNodeGetParams"; + case HIP_API_ID_hipGraphMemFreeNodeGetParams: return "hipGraphMemFreeNodeGetParams"; case HIP_API_ID_hipGraphMemcpyNodeGetParams: return "hipGraphMemcpyNodeGetParams"; case HIP_API_ID_hipGraphMemcpyNodeSetParams: return "hipGraphMemcpyNodeSetParams"; case HIP_API_ID_hipGraphMemcpyNodeSetParams1D: return "hipGraphMemcpyNodeSetParams1D"; @@ -694,6 +704,8 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipModuleGetFunction: return "hipModuleGetFunction"; case HIP_API_ID_hipModuleGetGlobal: return "hipModuleGetGlobal"; case HIP_API_ID_hipModuleGetTexRef: return "hipModuleGetTexRef"; + case HIP_API_ID_hipModuleLaunchCooperativeKernel: return "hipModuleLaunchCooperativeKernel"; + case HIP_API_ID_hipModuleLaunchCooperativeKernelMultiDevice: return "hipModuleLaunchCooperativeKernelMultiDevice"; case HIP_API_ID_hipModuleLaunchKernel: return "hipModuleLaunchKernel"; case HIP_API_ID_hipModuleLoad: return "hipModuleLoad"; case HIP_API_ID_hipModuleLoadData: return "hipModuleLoadData"; @@ -874,6 +886,8 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipGraphAddEventWaitNode", name) == 0) return HIP_API_ID_hipGraphAddEventWaitNode; 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; + if (strcmp("hipGraphAddMemFreeNode", name) == 0) return HIP_API_ID_hipGraphAddMemFreeNode; if (strcmp("hipGraphAddMemcpyNode", name) == 0) return HIP_API_ID_hipGraphAddMemcpyNode; if (strcmp("hipGraphAddMemcpyNode1D", name) == 0) return HIP_API_ID_hipGraphAddMemcpyNode1D; if (strcmp("hipGraphAddMemcpyNodeFromSymbol", name) == 0) return HIP_API_ID_hipGraphAddMemcpyNodeFromSymbol; @@ -914,6 +928,8 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipGraphKernelNodeSetAttribute", name) == 0) return HIP_API_ID_hipGraphKernelNodeSetAttribute; if (strcmp("hipGraphKernelNodeSetParams", name) == 0) return HIP_API_ID_hipGraphKernelNodeSetParams; if (strcmp("hipGraphLaunch", name) == 0) return HIP_API_ID_hipGraphLaunch; + if (strcmp("hipGraphMemAllocNodeGetParams", name) == 0) return HIP_API_ID_hipGraphMemAllocNodeGetParams; + if (strcmp("hipGraphMemFreeNodeGetParams", name) == 0) return HIP_API_ID_hipGraphMemFreeNodeGetParams; if (strcmp("hipGraphMemcpyNodeGetParams", name) == 0) return HIP_API_ID_hipGraphMemcpyNodeGetParams; if (strcmp("hipGraphMemcpyNodeSetParams", name) == 0) return HIP_API_ID_hipGraphMemcpyNodeSetParams; if (strcmp("hipGraphMemcpyNodeSetParams1D", name) == 0) return HIP_API_ID_hipGraphMemcpyNodeSetParams1D; @@ -1050,6 +1066,8 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipModuleGetFunction", name) == 0) return HIP_API_ID_hipModuleGetFunction; if (strcmp("hipModuleGetGlobal", name) == 0) return HIP_API_ID_hipModuleGetGlobal; if (strcmp("hipModuleGetTexRef", name) == 0) return HIP_API_ID_hipModuleGetTexRef; + if (strcmp("hipModuleLaunchCooperativeKernel", name) == 0) return HIP_API_ID_hipModuleLaunchCooperativeKernel; + if (strcmp("hipModuleLaunchCooperativeKernelMultiDevice", name) == 0) return HIP_API_ID_hipModuleLaunchCooperativeKernelMultiDevice; if (strcmp("hipModuleLaunchKernel", name) == 0) return HIP_API_ID_hipModuleLaunchKernel; if (strcmp("hipModuleLoad", name) == 0) return HIP_API_ID_hipModuleLoad; if (strcmp("hipModuleLoadData", name) == 0) return HIP_API_ID_hipModuleLoadData; @@ -1657,6 +1675,25 @@ typedef struct hip_api_data_s { const hipKernelNodeParams* pNodeParams; hipKernelNodeParams pNodeParams__val; } hipGraphAddKernelNode; + struct { + hipGraphNode_t* pGraphNode; + hipGraphNode_t pGraphNode__val; + hipGraph_t graph; + const hipGraphNode_t* pDependencies; + hipGraphNode_t pDependencies__val; + size_t numDependencies; + hipMemAllocNodeParams* pNodeParams; + hipMemAllocNodeParams pNodeParams__val; + } hipGraphAddMemAllocNode; + struct { + hipGraphNode_t* pGraphNode; + hipGraphNode_t pGraphNode__val; + hipGraph_t graph; + const hipGraphNode_t* pDependencies; + hipGraphNode_t pDependencies__val; + size_t numDependencies; + void* dev_ptr; + } hipGraphAddMemFreeNode; struct { hipGraphNode_t* pGraphNode; hipGraphNode_t pGraphNode__val; @@ -1915,6 +1952,15 @@ typedef struct hip_api_data_s { hipGraphExec_t graphExec; hipStream_t stream; } hipGraphLaunch; + struct { + hipGraphNode_t node; + hipMemAllocNodeParams* pNodeParams; + hipMemAllocNodeParams pNodeParams__val; + } hipGraphMemAllocNodeGetParams; + struct { + hipGraphNode_t node; + void* dev_ptr; + } hipGraphMemFreeNodeGetParams; struct { hipGraphNode_t node; hipMemcpy3DParms* pNodeParams; @@ -2774,6 +2820,25 @@ typedef struct hip_api_data_s { const char* name; char name__val; } hipModuleGetTexRef; + struct { + hipFunction_t f; + unsigned int gridDimX; + unsigned int gridDimY; + unsigned int gridDimZ; + unsigned int blockDimX; + unsigned int blockDimY; + unsigned int blockDimZ; + unsigned int sharedMemBytes; + hipStream_t stream; + void** kernelParams; + void* kernelParams__val; + } hipModuleLaunchCooperativeKernel; + struct { + hipFunctionLaunchParams* launchParamsList; + hipFunctionLaunchParams launchParamsList__val; + unsigned int numDevices; + unsigned int flags; + } hipModuleLaunchCooperativeKernelMultiDevice; struct { hipFunction_t f; unsigned int gridDimX; @@ -3725,6 +3790,12 @@ typedef struct hip_api_data_s { cb_data.args.hipGraphAddKernelNode.numDependencies = (size_t)numDependencies; \ cb_data.args.hipGraphAddKernelNode.pNodeParams = (const hipKernelNodeParams*)pNodeParams; \ }; +// hipGraphAddMemAllocNode[('hipGraphNode_t*', 'pGraphNode'), ('hipGraph_t', 'graph'), ('const hipGraphNode_t*', 'pDependencies'), ('size_t', 'numDependencies'), ('hipMemAllocNodeParams*', 'pNodeParams')] +#define INIT_hipGraphAddMemAllocNode_CB_ARGS_DATA(cb_data) { \ +}; +// hipGraphAddMemFreeNode[('hipGraphNode_t*', 'pGraphNode'), ('hipGraph_t', 'graph'), ('const hipGraphNode_t*', 'pDependencies'), ('size_t', 'numDependencies'), ('void*', 'dev_ptr')] +#define INIT_hipGraphAddMemFreeNode_CB_ARGS_DATA(cb_data) { \ +}; // hipGraphAddMemcpyNode[('hipGraphNode_t*', 'pGraphNode'), ('hipGraph_t', 'graph'), ('const hipGraphNode_t*', 'pDependencies'), ('size_t', 'numDependencies'), ('const hipMemcpy3DParms*', 'pCopyParams')] #define INIT_hipGraphAddMemcpyNode_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipGraphAddMemcpyNode.pGraphNode = (hipGraphNode_t*)pGraphNode; \ @@ -3982,6 +4053,12 @@ typedef struct hip_api_data_s { cb_data.args.hipGraphLaunch.graphExec = (hipGraphExec_t)graphExec; \ cb_data.args.hipGraphLaunch.stream = (hipStream_t)stream; \ }; +// hipGraphMemAllocNodeGetParams[('hipGraphNode_t', 'node'), ('hipMemAllocNodeParams*', 'pNodeParams')] +#define INIT_hipGraphMemAllocNodeGetParams_CB_ARGS_DATA(cb_data) { \ +}; +// hipGraphMemFreeNodeGetParams[('hipGraphNode_t', 'node'), ('void*', 'dev_ptr')] +#define INIT_hipGraphMemFreeNodeGetParams_CB_ARGS_DATA(cb_data) { \ +}; // hipGraphMemcpyNodeGetParams[('hipGraphNode_t', 'node'), ('hipMemcpy3DParms*', 'pNodeParams')] #define INIT_hipGraphMemcpyNodeGetParams_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipGraphMemcpyNodeGetParams.node = (hipGraphNode_t)node; \ @@ -4875,6 +4952,25 @@ typedef struct hip_api_data_s { cb_data.args.hipModuleGetTexRef.hmod = (hipModule_t)hmod; \ cb_data.args.hipModuleGetTexRef.name = (name) ? strdup(name) : NULL; \ }; +// hipModuleLaunchCooperativeKernel[('hipFunction_t', 'f'), ('unsigned int', 'gridDimX'), ('unsigned int', 'gridDimY'), ('unsigned int', 'gridDimZ'), ('unsigned int', 'blockDimX'), ('unsigned int', 'blockDimY'), ('unsigned int', 'blockDimZ'), ('unsigned int', 'sharedMemBytes'), ('hipStream_t', 'stream'), ('void**', 'kernelParams')] +#define INIT_hipModuleLaunchCooperativeKernel_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipModuleLaunchCooperativeKernel.f = (hipFunction_t)f; \ + cb_data.args.hipModuleLaunchCooperativeKernel.gridDimX = (unsigned int)gridDimX; \ + cb_data.args.hipModuleLaunchCooperativeKernel.gridDimY = (unsigned int)gridDimY; \ + cb_data.args.hipModuleLaunchCooperativeKernel.gridDimZ = (unsigned int)gridDimZ; \ + cb_data.args.hipModuleLaunchCooperativeKernel.blockDimX = (unsigned int)blockDimX; \ + cb_data.args.hipModuleLaunchCooperativeKernel.blockDimY = (unsigned int)blockDimY; \ + cb_data.args.hipModuleLaunchCooperativeKernel.blockDimZ = (unsigned int)blockDimZ; \ + cb_data.args.hipModuleLaunchCooperativeKernel.sharedMemBytes = (unsigned int)sharedMemBytes; \ + cb_data.args.hipModuleLaunchCooperativeKernel.stream = (hipStream_t)stream; \ + cb_data.args.hipModuleLaunchCooperativeKernel.kernelParams = (void**)kernelParams; \ +}; +// hipModuleLaunchCooperativeKernelMultiDevice[('hipFunctionLaunchParams*', 'launchParamsList'), ('unsigned int', 'numDevices'), ('unsigned int', 'flags')] +#define INIT_hipModuleLaunchCooperativeKernelMultiDevice_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipModuleLaunchCooperativeKernelMultiDevice.launchParamsList = (hipFunctionLaunchParams*)launchParamsList; \ + cb_data.args.hipModuleLaunchCooperativeKernelMultiDevice.numDevices = (unsigned int)numDevices; \ + cb_data.args.hipModuleLaunchCooperativeKernelMultiDevice.flags = (unsigned int)flags; \ +}; // hipModuleLaunchKernel[('hipFunction_t', 'f'), ('unsigned int', 'gridDimX'), ('unsigned int', 'gridDimY'), ('unsigned int', 'gridDimZ'), ('unsigned int', 'blockDimX'), ('unsigned int', 'blockDimY'), ('unsigned int', 'blockDimZ'), ('unsigned int', 'sharedMemBytes'), ('hipStream_t', 'stream'), ('void**', 'kernelParams'), ('void**', 'extra')] #define INIT_hipModuleLaunchKernel_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipModuleLaunchKernel.f = (hipFunction_t)f; \ @@ -5767,6 +5863,17 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { if (data->args.hipGraphAddKernelNode.pDependencies) data->args.hipGraphAddKernelNode.pDependencies__val = *(data->args.hipGraphAddKernelNode.pDependencies); if (data->args.hipGraphAddKernelNode.pNodeParams) data->args.hipGraphAddKernelNode.pNodeParams__val = *(data->args.hipGraphAddKernelNode.pNodeParams); break; +// hipGraphAddMemAllocNode[('hipGraphNode_t*', 'pGraphNode'), ('hipGraph_t', 'graph'), ('const hipGraphNode_t*', 'pDependencies'), ('size_t', 'numDependencies'), ('hipMemAllocNodeParams*', 'pNodeParams')] + case HIP_API_ID_hipGraphAddMemAllocNode: + if (data->args.hipGraphAddMemAllocNode.pGraphNode) data->args.hipGraphAddMemAllocNode.pGraphNode__val = *(data->args.hipGraphAddMemAllocNode.pGraphNode); + if (data->args.hipGraphAddMemAllocNode.pDependencies) data->args.hipGraphAddMemAllocNode.pDependencies__val = *(data->args.hipGraphAddMemAllocNode.pDependencies); + if (data->args.hipGraphAddMemAllocNode.pNodeParams) data->args.hipGraphAddMemAllocNode.pNodeParams__val = *(data->args.hipGraphAddMemAllocNode.pNodeParams); + break; +// hipGraphAddMemFreeNode[('hipGraphNode_t*', 'pGraphNode'), ('hipGraph_t', 'graph'), ('const hipGraphNode_t*', 'pDependencies'), ('size_t', 'numDependencies'), ('void*', 'dev_ptr')] + case HIP_API_ID_hipGraphAddMemFreeNode: + if (data->args.hipGraphAddMemFreeNode.pGraphNode) data->args.hipGraphAddMemFreeNode.pGraphNode__val = *(data->args.hipGraphAddMemFreeNode.pGraphNode); + if (data->args.hipGraphAddMemFreeNode.pDependencies) data->args.hipGraphAddMemFreeNode.pDependencies__val = *(data->args.hipGraphAddMemFreeNode.pDependencies); + break; // hipGraphAddMemcpyNode[('hipGraphNode_t*', 'pGraphNode'), ('hipGraph_t', 'graph'), ('const hipGraphNode_t*', 'pDependencies'), ('size_t', 'numDependencies'), ('const hipMemcpy3DParms*', 'pCopyParams')] case HIP_API_ID_hipGraphAddMemcpyNode: if (data->args.hipGraphAddMemcpyNode.pGraphNode) data->args.hipGraphAddMemcpyNode.pGraphNode__val = *(data->args.hipGraphAddMemcpyNode.pGraphNode); @@ -5928,6 +6035,13 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipGraphLaunch[('hipGraphExec_t', 'graphExec'), ('hipStream_t', 'stream')] case HIP_API_ID_hipGraphLaunch: break; +// hipGraphMemAllocNodeGetParams[('hipGraphNode_t', 'node'), ('hipMemAllocNodeParams*', 'pNodeParams')] + case HIP_API_ID_hipGraphMemAllocNodeGetParams: + if (data->args.hipGraphMemAllocNodeGetParams.pNodeParams) data->args.hipGraphMemAllocNodeGetParams.pNodeParams__val = *(data->args.hipGraphMemAllocNodeGetParams.pNodeParams); + break; +// hipGraphMemFreeNodeGetParams[('hipGraphNode_t', 'node'), ('void*', 'dev_ptr')] + case HIP_API_ID_hipGraphMemFreeNodeGetParams: + break; // hipGraphMemcpyNodeGetParams[('hipGraphNode_t', 'node'), ('hipMemcpy3DParms*', 'pNodeParams')] case HIP_API_ID_hipGraphMemcpyNodeGetParams: if (data->args.hipGraphMemcpyNodeGetParams.pNodeParams) data->args.hipGraphMemcpyNodeGetParams.pNodeParams__val = *(data->args.hipGraphMemcpyNodeGetParams.pNodeParams); @@ -6438,6 +6552,14 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { if (data->args.hipModuleGetTexRef.texRef) data->args.hipModuleGetTexRef.texRef__val = *(data->args.hipModuleGetTexRef.texRef); if (data->args.hipModuleGetTexRef.name) data->args.hipModuleGetTexRef.name__val = *(data->args.hipModuleGetTexRef.name); break; +// hipModuleLaunchCooperativeKernel[('hipFunction_t', 'f'), ('unsigned int', 'gridDimX'), ('unsigned int', 'gridDimY'), ('unsigned int', 'gridDimZ'), ('unsigned int', 'blockDimX'), ('unsigned int', 'blockDimY'), ('unsigned int', 'blockDimZ'), ('unsigned int', 'sharedMemBytes'), ('hipStream_t', 'stream'), ('void**', 'kernelParams')] + case HIP_API_ID_hipModuleLaunchCooperativeKernel: + if (data->args.hipModuleLaunchCooperativeKernel.kernelParams) data->args.hipModuleLaunchCooperativeKernel.kernelParams__val = *(data->args.hipModuleLaunchCooperativeKernel.kernelParams); + break; +// hipModuleLaunchCooperativeKernelMultiDevice[('hipFunctionLaunchParams*', 'launchParamsList'), ('unsigned int', 'numDevices'), ('unsigned int', 'flags')] + case HIP_API_ID_hipModuleLaunchCooperativeKernelMultiDevice: + if (data->args.hipModuleLaunchCooperativeKernelMultiDevice.launchParamsList) data->args.hipModuleLaunchCooperativeKernelMultiDevice.launchParamsList__val = *(data->args.hipModuleLaunchCooperativeKernelMultiDevice.launchParamsList); + break; // hipModuleLaunchKernel[('hipFunction_t', 'f'), ('unsigned int', 'gridDimX'), ('unsigned int', 'gridDimY'), ('unsigned int', 'gridDimZ'), ('unsigned int', 'blockDimX'), ('unsigned int', 'blockDimY'), ('unsigned int', 'blockDimZ'), ('unsigned int', 'sharedMemBytes'), ('hipStream_t', 'stream'), ('void**', 'kernelParams'), ('void**', 'extra')] case HIP_API_ID_hipModuleLaunchKernel: if (data->args.hipModuleLaunchKernel.kernelParams) data->args.hipModuleLaunchKernel.kernelParams__val = *(data->args.hipModuleLaunchKernel.kernelParams); @@ -7474,6 +7596,29 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da else { oss << ", pNodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddKernelNode.pNodeParams__val); } oss << ")"; break; + case HIP_API_ID_hipGraphAddMemAllocNode: + oss << "hipGraphAddMemAllocNode("; + if (data->args.hipGraphAddMemAllocNode.pGraphNode == NULL) oss << "pGraphNode=NULL"; + else { oss << "pGraphNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddMemAllocNode.pGraphNode__val); } + oss << ", graph="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddMemAllocNode.graph); + if (data->args.hipGraphAddMemAllocNode.pDependencies == NULL) oss << ", pDependencies=NULL"; + else { oss << ", pDependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddMemAllocNode.pDependencies__val); } + oss << ", numDependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddMemAllocNode.numDependencies); + if (data->args.hipGraphAddMemAllocNode.pNodeParams == NULL) oss << ", pNodeParams=NULL"; + else { oss << ", pNodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddMemAllocNode.pNodeParams__val); } + oss << ")"; + break; + case HIP_API_ID_hipGraphAddMemFreeNode: + oss << "hipGraphAddMemFreeNode("; + if (data->args.hipGraphAddMemFreeNode.pGraphNode == NULL) oss << "pGraphNode=NULL"; + else { oss << "pGraphNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddMemFreeNode.pGraphNode__val); } + oss << ", graph="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddMemFreeNode.graph); + if (data->args.hipGraphAddMemFreeNode.pDependencies == NULL) oss << ", pDependencies=NULL"; + else { oss << ", pDependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddMemFreeNode.pDependencies__val); } + oss << ", numDependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddMemFreeNode.numDependencies); + oss << ", dev_ptr="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphAddMemFreeNode.dev_ptr); + oss << ")"; + break; case HIP_API_ID_hipGraphAddMemcpyNode: oss << "hipGraphAddMemcpyNode("; if (data->args.hipGraphAddMemcpyNode.pGraphNode == NULL) oss << "pGraphNode=NULL"; @@ -7812,6 +7957,19 @@ 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.hipGraphLaunch.stream); oss << ")"; break; + case HIP_API_ID_hipGraphMemAllocNodeGetParams: + oss << "hipGraphMemAllocNodeGetParams("; + oss << "node="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphMemAllocNodeGetParams.node); + if (data->args.hipGraphMemAllocNodeGetParams.pNodeParams == NULL) oss << ", pNodeParams=NULL"; + else { oss << ", pNodeParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphMemAllocNodeGetParams.pNodeParams__val); } + oss << ")"; + break; + case HIP_API_ID_hipGraphMemFreeNodeGetParams: + oss << "hipGraphMemFreeNodeGetParams("; + oss << "node="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphMemFreeNodeGetParams.node); + oss << ", dev_ptr="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphMemFreeNodeGetParams.dev_ptr); + oss << ")"; + break; case HIP_API_ID_hipGraphMemcpyNodeGetParams: oss << "hipGraphMemcpyNodeGetParams("; oss << "node="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGraphMemcpyNodeGetParams.node); @@ -8943,6 +9101,29 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da else { oss << ", name="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleGetTexRef.name__val); } oss << ")"; break; + case HIP_API_ID_hipModuleLaunchCooperativeKernel: + oss << "hipModuleLaunchCooperativeKernel("; + oss << "f="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLaunchCooperativeKernel.f); + oss << ", gridDimX="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLaunchCooperativeKernel.gridDimX); + oss << ", gridDimY="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLaunchCooperativeKernel.gridDimY); + oss << ", gridDimZ="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLaunchCooperativeKernel.gridDimZ); + oss << ", blockDimX="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLaunchCooperativeKernel.blockDimX); + oss << ", blockDimY="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLaunchCooperativeKernel.blockDimY); + oss << ", blockDimZ="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLaunchCooperativeKernel.blockDimZ); + oss << ", sharedMemBytes="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLaunchCooperativeKernel.sharedMemBytes); + oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLaunchCooperativeKernel.stream); + if (data->args.hipModuleLaunchCooperativeKernel.kernelParams == NULL) oss << ", kernelParams=NULL"; + else { oss << ", kernelParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLaunchCooperativeKernel.kernelParams__val); } + oss << ")"; + break; + case HIP_API_ID_hipModuleLaunchCooperativeKernelMultiDevice: + oss << "hipModuleLaunchCooperativeKernelMultiDevice("; + if (data->args.hipModuleLaunchCooperativeKernelMultiDevice.launchParamsList == NULL) oss << "launchParamsList=NULL"; + else { oss << "launchParamsList="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLaunchCooperativeKernelMultiDevice.launchParamsList__val); } + oss << ", numDevices="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLaunchCooperativeKernelMultiDevice.numDevices); + oss << ", flags="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLaunchCooperativeKernelMultiDevice.flags); + oss << ")"; + break; case HIP_API_ID_hipModuleLaunchKernel: oss << "hipModuleLaunchKernel("; oss << "f="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLaunchKernel.f); @@ -9084,9 +9265,9 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da break; case HIP_API_ID_hipPointerSetAttribute: oss << "hipPointerSetAttribute("; - oss << "value=" << data->args.hipPointerSetAttribute.value; - oss << ", attribute=" << data->args.hipPointerSetAttribute.attribute; - oss << ", ptr=" << data->args.hipPointerSetAttribute.ptr; + oss << "value="; roctracer::hip_support::detail::operator<<(oss, data->args.hipPointerSetAttribute.value); + oss << ", attribute="; roctracer::hip_support::detail::operator<<(oss, data->args.hipPointerSetAttribute.attribute); + oss << ", ptr="; roctracer::hip_support::detail::operator<<(oss, data->args.hipPointerSetAttribute.ptr); oss << ")"; break; case HIP_API_ID_hipProfilerStart: diff --git a/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h index 712d793c51..f561ed3876 100644 --- a/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -437,6 +437,7 @@ typedef struct cudaArray* hipArray_t; typedef struct cudaArray* hipArray_const_t; typedef struct cudaFuncAttributes hipFuncAttributes; typedef struct cudaLaunchParams hipLaunchParams; +typedef CUDA_LAUNCH_PARAMS hipFunctionLaunchParams; #define hipFunction_attribute CUfunction_attribute #define hipPointer_attribute CUpointer_attribute #define hip_Memcpy2D CUDA_MEMCPY2D @@ -2807,11 +2808,29 @@ inline static hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim, cudaLaunchCooperativeKernel(f, gridDim, blockDim, kernelParams, sharedMemBytes, stream)); } +inline static hipError_t hipModuleLaunchCooperativeKernel(hipFunction_t f, unsigned int gridDimX, + unsigned int gridDimY, unsigned int gridDimZ, + unsigned int blockDimX, unsigned int blockDimY, + unsigned int blockDimZ, unsigned int sharedMemBytes, + hipStream_t stream, void** kernelParams) { + return hipCUResultTohipError(cuLaunchCooperativeKernel(f, gridDimX, gridDimY, gridDimZ, + blockDimX, blockDimY, blockDimZ, + sharedMemBytes, stream,kernelParams)); +} + inline static hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices, unsigned int flags) { return hipCUDAErrorTohipError(cudaLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags)); } +inline static hipError_t hipModuleLaunchCooperativeKernelMultiDevice( + hipFunctionLaunchParams* launchParamsList, + unsigned int numDevices, + unsigned int flags) { + return hipCUResultTohipError(cuLaunchCooperativeKernelMultiDevice(launchParamsList, + numDevices, flags)); +} + inline static hipError_t hipImportExternalSemaphore(hipExternalSemaphore_t* extSem_out, const hipExternalSemaphoreHandleDesc* semHandleDesc) { return hipCUDAErrorTohipError(cudaImportExternalSemaphore(extSem_out,(const struct cudaExternalSemaphoreHandleDesc*)semHandleDesc)); diff --git a/projects/clr/hipamd/src/amdhip.def b/projects/clr/hipamd/src/amdhip.def index d5ad2076bb..38489b9de6 100644 --- a/projects/clr/hipamd/src/amdhip.def +++ b/projects/clr/hipamd/src/amdhip.def @@ -158,6 +158,8 @@ hipModuleGetGlobal hipModuleGetTexRef hipModuleLaunchKernel hipModuleLaunchKernelExt +hipModuleLaunchCooperativeKernel +hipModuleLaunchCooperativeKernelMultiDevice hipLaunchCooperativeKernel hipLaunchCooperativeKernelMultiDevice hipHccModuleLaunchKernel diff --git a/projects/clr/hipamd/src/hip_hcc.def.in b/projects/clr/hipamd/src/hip_hcc.def.in index cad7b46eb3..e93f213cce 100644 --- a/projects/clr/hipamd/src/hip_hcc.def.in +++ b/projects/clr/hipamd/src/hip_hcc.def.in @@ -158,6 +158,8 @@ hipModuleGetGlobal hipModuleGetTexRef hipModuleLaunchKernel hipModuleLaunchKernelExt +hipModuleLaunchCooperativeKernel +hipModuleLaunchCooperativeKernelMultiDevice hipLaunchCooperativeKernel hipLaunchCooperativeKernelMultiDevice hipHccModuleLaunchKernel diff --git a/projects/clr/hipamd/src/hip_hcc.map.in b/projects/clr/hipamd/src/hip_hcc.map.in index 8f9f76f312..1f6c40ea78 100644 --- a/projects/clr/hipamd/src/hip_hcc.map.in +++ b/projects/clr/hipamd/src/hip_hcc.map.in @@ -505,3 +505,11 @@ global: local: *; } hip_5.2; + +hip_5.5 { +global: + hipModuleLaunchCooperativeKernel; + hipModuleLaunchCooperativeKernelMultiDevice; +local: + *; +} hip_5.3; diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index bb400ad1b5..e14b33462d 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/src/hip_module.cpp @@ -473,6 +473,158 @@ hipError_t hipModuleLaunchKernelExt(hipFunction_t f, uint32_t globalWorkSizeX, extra, startEvent, stopEvent)); } +hipError_t hipModuleLaunchCooperativeKernel(hipFunction_t f, unsigned int gridDimX, + unsigned int gridDimY, unsigned int gridDimZ, + unsigned int blockDimX, unsigned int blockDimY, + unsigned int blockDimZ, unsigned int sharedMemBytes, + hipStream_t stream, void** kernelParams) { + HIP_INIT_API(hipModuleLaunchCooperativeKernel, f, gridDimX, gridDimY, gridDimZ, blockDimX, + blockDimY, blockDimZ, sharedMemBytes, stream, kernelParams); + + if (!hip::isValid(stream)) { + HIP_RETURN(hipErrorInvalidValue); + } + + size_t globalWorkSizeX = static_cast(gridDimX) * blockDimX; + size_t globalWorkSizeY = static_cast(gridDimY) * blockDimY; + size_t globalWorkSizeZ = static_cast(gridDimZ) * blockDimZ; + if (globalWorkSizeX > std::numeric_limits::max() || + globalWorkSizeY > std::numeric_limits::max() || + globalWorkSizeZ > std::numeric_limits::max()) { + HIP_RETURN(hipErrorInvalidConfiguration); + } + HIP_RETURN(ihipModuleLaunchKernel(f, static_cast(globalWorkSizeX), + static_cast(globalWorkSizeY), + static_cast(globalWorkSizeZ), blockDimX, blockDimY, + blockDimZ, sharedMemBytes, stream, kernelParams, nullptr, nullptr, + nullptr, 0, amd::NDRangeKernelCommand::CooperativeGroups)); +} + +hipError_t ihipModuleLaunchCooperativeKernelMultiDevice(hipFunctionLaunchParams* launchParamsList, + unsigned int numDevices, + unsigned int flags, + uint32_t extFlags) { + int numActiveGPUs = 0; + hipError_t result = hipSuccess; + result = ihipDeviceGetCount(&numActiveGPUs); + + if (numDevices > numActiveGPUs) { + return hipErrorInvalidValue; + } + + uint64_t allGridSize = 0; + std::vector mgpu_list(numDevices); + + for (int i = 0; i < numDevices; ++i) { + uint32_t blockDims = 0; + const hipFunctionLaunchParams& launch = launchParamsList[i]; + blockDims = launch.blockDimX * launch.blockDimY * launch.blockDimZ; + allGridSize += launch.gridDimX * launch.gridDimY * launch.gridDimZ * blockDims; + + // Make sure block dimensions are valid + if (0 == blockDims) { + return hipErrorInvalidConfiguration; + } + if (launch.hStream != nullptr) { + // Validate devices to make sure it dosn't have duplicates + amd::HostQueue* queue = reinterpret_cast(launch.hStream)->asHostQueue(); + auto device = &queue->vdev()->device(); + for (int j = 0; j < numDevices; ++j) { + if (mgpu_list[j] == device) { + return hipErrorInvalidDevice; + } + } + mgpu_list[i] = device; + } else { + return hipErrorInvalidResourceHandle; + } + } + uint64_t prevGridSize = 0; + uint32_t firstDevice = 0; + + // Sync the execution streams on all devices + if ((flags & hipCooperativeLaunchMultiDeviceNoPreSync) == 0) { + for (int i = 0; i < numDevices; ++i) { + amd::HostQueue* queue = + reinterpret_cast(launchParamsList[i].hStream)->asHostQueue(); + queue->finish(); + } + } + + for (int i = 0; i < numDevices; ++i) { + const hipFunctionLaunchParams& launch = launchParamsList[i]; + amd::HostQueue* queue = reinterpret_cast(launch.hStream)->asHostQueue(); + + if (i == 0) { + // The order of devices in the launch may not match the order in the global array + for (size_t dev = 0; dev < g_devices.size(); ++dev) { + // Find the matching device + if (&queue->vdev()->device() == g_devices[dev]->devices()[0]) { + // Save ROCclr index of the first device in the launch + firstDevice = queue->vdev()->device().index(); + break; + } + } + } + + size_t globalWorkSizeX = static_cast(launch.gridDimX) * launch.blockDimX; + size_t globalWorkSizeY = static_cast(launch.gridDimY) * launch.blockDimY; + size_t globalWorkSizeZ = static_cast(launch.gridDimZ) * launch.blockDimZ; + if (globalWorkSizeX > std::numeric_limits::max() || + globalWorkSizeY > std::numeric_limits::max() || + globalWorkSizeZ > std::numeric_limits::max()) { + return hipErrorInvalidConfiguration; + } + result = ihipModuleLaunchKernel( + launch.function, static_cast(globalWorkSizeX), + static_cast(globalWorkSizeY), + static_cast(globalWorkSizeZ), launch.blockDimX, launch.blockDimY, + launch.blockDimZ, launch.sharedMemBytes, launch.hStream, launch.kernelParams, + nullptr, nullptr, nullptr, flags, extFlags, + i, numDevices, prevGridSize, allGridSize, firstDevice); + if (result != hipSuccess) { + break; + } + prevGridSize += globalWorkSizeX * globalWorkSizeY * globalWorkSizeZ; + } + + // Sync the execution streams on all devices + if ((flags & hipCooperativeLaunchMultiDeviceNoPostSync) == 0) { + for (int i = 0; i < numDevices; ++i) { + amd::HostQueue* queue = + reinterpret_cast(launchParamsList[i].hStream)->asHostQueue(); + queue->finish(); + } + } + + return result; +} + +hipError_t hipModuleLaunchCooperativeKernelMultiDevice(hipFunctionLaunchParams* launchParamsList, + unsigned int numDevices, + unsigned int flags) { + HIP_INIT_API(hipModuleLaunchCooperativeKernelMultiDevice, launchParamsList, numDevices, flags); + + if (launchParamsList == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + // Validate all streams passed by user + for (int i = 0; i < numDevices; ++i) { + if (!hip::isValid(launchParamsList[i].hStream)) { + HIP_RETURN(hipErrorInvalidValue); + } + } + + HIP_RETURN(ihipModuleLaunchCooperativeKernelMultiDevice( + launchParamsList, + numDevices, + flags, + (amd::NDRangeKernelCommand::CooperativeGroups | + amd::NDRangeKernelCommand::CooperativeMultiDeviceGroups))); + +} + extern "C" hipError_t hipLaunchKernel_common(const void* hostFunction, dim3 gridDim, dim3 blockDim, void** args, size_t sharedMemBytes, hipStream_t stream) { @@ -557,61 +709,19 @@ hipError_t hipLaunchCooperativeKernel_spt(const void* f, dim3 gridDim, dim3 bloc hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices, unsigned int flags, uint32_t extFlags) { - int numActiveGPUs = 0; - hipError_t result = hipSuccess; - result = ihipDeviceGetCount(&numActiveGPUs); - - if ((numDevices > numActiveGPUs) || (launchParamsList == nullptr)) { + if (launchParamsList == nullptr) { return hipErrorInvalidValue; } - // Validate all streams passed by user + + std::vector functionLaunchParamsList(numDevices); + // Convert hipLaunchParams to hipFunctionLaunchParams for (int i = 0; i < numDevices; ++i) { - if (!hip::isValid(launchParamsList[i].stream)) { + hipLaunchParams& launch = launchParamsList[i]; + // Validate stream passed by user + if (!hip::isValid(launch.stream)) { return hipErrorInvalidValue; } - } - uint64_t allGridSize = 0; - std::vector mgpu_list(numDevices); - - for (int i = 0; i < numDevices; ++i) { - uint32_t blockDims = 0; - const hipLaunchParams& launch = launchParamsList[i]; - blockDims = launch.blockDim.x * launch.blockDim.y * launch.blockDim.z; - allGridSize += launch.gridDim.x * launch.gridDim.y * launch.gridDim.z * blockDims; - - // Make sure block dimensions are valid - if (0 == blockDims) { - return hipErrorInvalidConfiguration; - } - if (launch.stream != nullptr) { - // Validate devices to make sure it dosn't have duplicates - amd::HostQueue* queue = reinterpret_cast(launch.stream)->asHostQueue(); - auto device = &queue->vdev()->device(); - for (int j = 0; j < numDevices; ++j) { - if (mgpu_list[j] == device) { - return hipErrorInvalidDevice; - } - } - mgpu_list[i] = device; - } else { - return hipErrorInvalidResourceHandle; - } - } - uint64_t prevGridSize = 0; - uint32_t firstDevice = 0; - - // Sync the execution streams on all devices - if ((flags & hipCooperativeLaunchMultiDeviceNoPreSync) == 0) { - for (int i = 0; i < numDevices; ++i) { - amd::HostQueue* queue = - reinterpret_cast(launchParamsList[i].stream)->asHostQueue(); - queue->finish(); - } - } - - for (int i = 0; i < numDevices; ++i) { - const hipLaunchParams& launch = launchParamsList[i]; amd::HostQueue* queue = reinterpret_cast(launch.stream)->asHostQueue(); hipFunction_t func = nullptr; // The order of devices in the launch may not match the order in the global array @@ -619,46 +729,29 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL // Find the matching device and request the kernel function if (&queue->vdev()->device() == g_devices[dev]->devices()[0]) { IHIP_RETURN_ONFAIL(PlatformState::instance().getStatFunc(&func, launch.func, dev)); - // Save ROCclr index of the first device in the launch - if (i == 0) { - firstDevice = queue->vdev()->device().index(); - } break; } } if (func == nullptr) { - result = hipErrorInvalidDeviceFunction; - HIP_RETURN(result); + return hipErrorInvalidDeviceFunction; } - size_t globalWorkSizeX = static_cast(launch.gridDim.x) * launch.blockDim.x; - size_t globalWorkSizeY = static_cast(launch.gridDim.y) * launch.blockDim.y; - size_t globalWorkSizeZ = static_cast(launch.gridDim.z) * launch.blockDim.z; - if (globalWorkSizeX > std::numeric_limits::max() || - globalWorkSizeY > std::numeric_limits::max() || - globalWorkSizeZ > std::numeric_limits::max()) { - HIP_RETURN(hipErrorInvalidConfiguration); - } - result = ihipModuleLaunchKernel( - func, static_cast(globalWorkSizeX), static_cast(globalWorkSizeY), - static_cast(globalWorkSizeZ), launch.blockDim.x, launch.blockDim.y, - launch.blockDim.z, launch.sharedMem, launch.stream, launch.args, nullptr, nullptr, nullptr, - flags, extFlags, i, numDevices, prevGridSize, allGridSize, firstDevice); - if (result != hipSuccess) { - break; - } - prevGridSize += globalWorkSizeX * globalWorkSizeY * globalWorkSizeZ; + + functionLaunchParamsList[i].function = func; + functionLaunchParamsList[i].gridDimX = launch.gridDim.x; + functionLaunchParamsList[i].gridDimY = launch.gridDim.y; + functionLaunchParamsList[i].gridDimZ = launch.gridDim.z; + functionLaunchParamsList[i].blockDimX = launch.blockDim.x; + functionLaunchParamsList[i].blockDimY = launch.blockDim.y; + functionLaunchParamsList[i].blockDimZ = launch.blockDim.z; + functionLaunchParamsList[i].sharedMemBytes = launch.sharedMem; + functionLaunchParamsList[i].hStream = launch.stream; + functionLaunchParamsList[i].kernelParams = launch.args; } - // Sync the execution streams on all devices - if ((flags & hipCooperativeLaunchMultiDeviceNoPostSync) == 0) { - for (int i = 0; i < numDevices; ++i) { - amd::HostQueue* queue = - reinterpret_cast(launchParamsList[i].stream)->asHostQueue(); - queue->finish(); - } - } - - return result; + return ihipModuleLaunchCooperativeKernelMultiDevice(functionLaunchParamsList.data(), + functionLaunchParamsList.size(), + flags, + extFlags); } hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices,