From 1ae22b2688d9cb0cf4b84d84e5dccbc5447128b4 Mon Sep 17 00:00:00 2001 From: foreman Date: Mon, 7 Oct 2019 11:55:30 -0400 Subject: [PATCH] P4 to Git Change 2009236 by eshcherb@evgeny-hip on 2019/10/07 11:49:55 SWDEV-197287 - HIP tracing layer instrumentation hip_prof_str.h - generated header; as a next step a genrator will be integrated in HIP makefile (similar to HIP-HCC) prof_protocol.h - temporarily included in HIP sources. Provided by roctracer-proto package. HIP-HCC CMake has parameter to set a path to the header. Affected files ... ... //depot/stg/opencl/drivers/opencl/api/hip/build/Makefile.hip#23 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_context.cpp#23 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_device.cpp#24 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_device_runtime.cpp#19 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_error.cpp#3 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_event.cpp#17 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_intercept.cpp#1 add ... //depot/stg/opencl/drivers/opencl/api/hip/hip_internal.hpp#44 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#74 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#43 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_peer.cpp#7 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#44 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_prof_api.h#1 add ... //depot/stg/opencl/drivers/opencl/api/hip/hip_profile.cpp#3 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_stream.cpp#25 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_surface.cpp#5 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_texture.cpp#21 edit ... //depot/stg/opencl/drivers/opencl/api/hip/prof_protocol.h#1 add ... //depot/stg/opencl/drivers/opencl/hip_prof_gen.py#1 add ... //depot/stg/opencl/drivers/opencl/make/hip.git/include/hip/hcc_detail/hip_prof_str.h#2 delete [ROCm/hip commit: 6e7f1dea66bf87325b7c7b94a332e85b99527cee] --- projects/hip/api/hip/hip_context.cpp | 44 ++-- projects/hip/api/hip/hip_device.cpp | 18 +- projects/hip/api/hip/hip_device_runtime.cpp | 38 +-- projects/hip/api/hip/hip_error.cpp | 4 +- projects/hip/api/hip/hip_event.cpp | 16 +- projects/hip/api/hip/hip_intercept.cpp | 58 +++++ projects/hip/api/hip/hip_internal.hpp | 6 +- projects/hip/api/hip/hip_memory.cpp | 106 ++++---- projects/hip/api/hip/hip_module.cpp | 35 +-- projects/hip/api/hip/hip_peer.cpp | 20 +- projects/hip/api/hip/hip_platform.cpp | 14 +- projects/hip/api/hip/hip_prof_api.h | 255 ++++++++++++++++++++ projects/hip/api/hip/hip_profile.cpp | 4 +- projects/hip/api/hip/hip_stream.cpp | 20 +- projects/hip/api/hip/hip_surface.cpp | 4 +- projects/hip/api/hip/hip_texture.cpp | 50 ++-- projects/hip/api/hip/prof_protocol.h | 90 +++++++ 17 files changed, 596 insertions(+), 186 deletions(-) create mode 100644 projects/hip/api/hip/hip_intercept.cpp create mode 100644 projects/hip/api/hip/hip_prof_api.h create mode 100644 projects/hip/api/hip/prof_protocol.h diff --git a/projects/hip/api/hip/hip_context.cpp b/projects/hip/api/hip/hip_context.cpp index 759676e83d..ed741c4883 100644 --- a/projects/hip/api/hip/hip_context.cpp +++ b/projects/hip/api/hip/hip_context.cpp @@ -112,13 +112,13 @@ amd::HostQueue* getNullStream() { using namespace hip; hipError_t hipInit(unsigned int flags) { - HIP_INIT_API(flags); + HIP_INIT_API(hipInit, flags); HIP_RETURN(hipSuccess); } hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device) { - HIP_INIT_API(ctx, flags, device); + HIP_INIT_API(hipCtxCreate, ctx, flags, device); if (static_cast(device) >= g_devices.size()) { HIP_RETURN(hipErrorInvalidValue); @@ -134,7 +134,7 @@ hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device) } hipError_t hipCtxSetCurrent(hipCtx_t ctx) { - HIP_INIT_API(ctx); + HIP_INIT_API(hipCtxSetCurrent, ctx); if (ctx == nullptr) { if(!g_ctxtStack.empty()) { @@ -152,7 +152,7 @@ hipError_t hipCtxSetCurrent(hipCtx_t ctx) { } hipError_t hipCtxGetCurrent(hipCtx_t* ctx) { - HIP_INIT_API(ctx); + HIP_INIT_API(hipCtxGetCurrent, ctx); *ctx = reinterpret_cast(hip::getCurrentContext()); @@ -160,7 +160,7 @@ hipError_t hipCtxGetCurrent(hipCtx_t* ctx) { } hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig) { - HIP_INIT_API(pConfig); + HIP_INIT_API(hipCtxGetSharedMemConfig, pConfig); *pConfig = hipSharedMemBankSizeFourByte; @@ -168,7 +168,7 @@ hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig) { } hipError_t hipRuntimeGetVersion(int *runtimeVersion) { - HIP_INIT_API(runtimeVersion); + HIP_INIT_API(hipRuntimeGetVersion, runtimeVersion); if (!runtimeVersion) { HIP_RETURN(hipErrorInvalidValue); @@ -180,7 +180,7 @@ hipError_t hipRuntimeGetVersion(int *runtimeVersion) { } hipError_t hipCtxDestroy(hipCtx_t ctx) { - HIP_INIT_API(ctx); + HIP_INIT_API(hipCtxDestroy, ctx); amd::Context* amdContext = reinterpret_cast(as_amd(ctx)); if (amdContext == nullptr) { @@ -207,7 +207,7 @@ hipError_t hipCtxDestroy(hipCtx_t ctx) { } hipError_t hipCtxPopCurrent(hipCtx_t* ctx) { - HIP_INIT_API(ctx); + HIP_INIT_API(hipCtxPopCurrent, ctx); amd::Context* amdContext = reinterpret_cast(as_amd(ctx)); if (amdContext == nullptr) { @@ -225,7 +225,7 @@ hipError_t hipCtxPopCurrent(hipCtx_t* ctx) { } hipError_t hipCtxPushCurrent(hipCtx_t ctx) { - HIP_INIT_API(ctx); + HIP_INIT_API(hipCtxPushCurrent, ctx); amd::Context* amdContext = reinterpret_cast(as_amd(ctx)); if (amdContext == nullptr) { @@ -239,7 +239,7 @@ hipError_t hipCtxPushCurrent(hipCtx_t ctx) { } hipError_t hipDriverGetVersion(int* driverVersion) { - HIP_INIT_API(driverVersion); + HIP_INIT_API(hipDriverGetVersion, driverVersion); auto* deviceHandle = g_devices[0]->devices()[0]; const auto& info = deviceHandle->info(); @@ -255,7 +255,7 @@ hipError_t hipDriverGetVersion(int* driverVersion) { } hipError_t hipCtxGetDevice(hipDevice_t* device) { - HIP_INIT_API(device); + HIP_INIT_API(hipCtxGetDevice, device); if (device != nullptr) { for (unsigned int i = 0; i < g_devices.size(); i++) { @@ -272,7 +272,7 @@ hipError_t hipCtxGetDevice(hipDevice_t* device) { } hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) { - HIP_INIT_API(apiVersion); + HIP_INIT_API(hipCtxGetApiVersion, apiVersion); assert(0 && "Unimplemented"); @@ -280,7 +280,7 @@ hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) { } hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) { - HIP_INIT_API(cacheConfig); + HIP_INIT_API(hipCtxGetCacheConfig, cacheConfig); assert(0 && "Unimplemented"); @@ -288,7 +288,7 @@ hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) { } hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) { - HIP_INIT_API(cacheConfig); + HIP_INIT_API(hipCtxSetCacheConfig, cacheConfig); assert(0 && "Unimplemented"); @@ -296,7 +296,7 @@ hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) { } hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) { - HIP_INIT_API(config); + HIP_INIT_API(hipCtxSetSharedMemConfig, config); assert(0 && "Unimplemented"); @@ -304,7 +304,7 @@ hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) { } hipError_t hipCtxSynchronize(void) { - HIP_INIT_API(1); + HIP_INIT_API(hipCtxSynchronize, 1); assert(0 && "Unimplemented"); @@ -312,7 +312,7 @@ hipError_t hipCtxSynchronize(void) { } hipError_t hipCtxGetFlags(unsigned int* flags) { - HIP_INIT_API(flags); + HIP_INIT_API(hipCtxGetFlags, flags); assert(0 && "Unimplemented"); @@ -320,7 +320,7 @@ hipError_t hipCtxGetFlags(unsigned int* flags) { } hipError_t hipDevicePrimaryCtxGetState(hipDevice_t dev, unsigned int* flags, int* active) { - HIP_INIT_API(dev, flags, active); + HIP_INIT_API(hipDevicePrimaryCtxGetState, dev, flags, active); if (static_cast(dev) >= g_devices.size()) { HIP_RETURN(hipErrorInvalidDevice); @@ -338,7 +338,7 @@ hipError_t hipDevicePrimaryCtxGetState(hipDevice_t dev, unsigned int* flags, int } hipError_t hipDevicePrimaryCtxRelease(hipDevice_t dev) { - HIP_INIT_API(dev); + HIP_INIT_API(hipDevicePrimaryCtxRelease, dev); if (static_cast(dev) >= g_devices.size()) { HIP_RETURN(hipErrorInvalidDevice); @@ -348,7 +348,7 @@ hipError_t hipDevicePrimaryCtxRelease(hipDevice_t dev) { } hipError_t hipDevicePrimaryCtxRetain(hipCtx_t* pctx, hipDevice_t dev) { - HIP_INIT_API(pctx, dev); + HIP_INIT_API(hipDevicePrimaryCtxRetain, pctx, dev); if (static_cast(dev) >= g_devices.size()) { HIP_RETURN(hipErrorInvalidDevice); @@ -363,13 +363,13 @@ hipError_t hipDevicePrimaryCtxRetain(hipCtx_t* pctx, hipDevice_t dev) { } hipError_t hipDevicePrimaryCtxReset(hipDevice_t dev) { - HIP_INIT_API(dev); + HIP_INIT_API(hipDevicePrimaryCtxReset, dev); HIP_RETURN(hipSuccess); } hipError_t hipDevicePrimaryCtxSetFlags(hipDevice_t dev, unsigned int flags) { - HIP_INIT_API(dev, flags); + HIP_INIT_API(hipDevicePrimaryCtxSetFlags, dev, flags); if (static_cast(dev) >= g_devices.size()) { HIP_RETURN(hipErrorInvalidDevice); diff --git a/projects/hip/api/hip/hip_device.cpp b/projects/hip/api/hip/hip_device.cpp index 4dfb447404..39c02aa0c7 100644 --- a/projects/hip/api/hip/hip_device.cpp +++ b/projects/hip/api/hip/hip_device.cpp @@ -25,7 +25,7 @@ THE SOFTWARE. #include "hip_internal.hpp" hipError_t hipDeviceGet(hipDevice_t *device, int deviceId) { - HIP_INIT_API(device, deviceId); + HIP_INIT_API(hipDeviceGet, device, deviceId); if (device != nullptr) { *device = deviceId; @@ -38,7 +38,7 @@ hipError_t hipDeviceGet(hipDevice_t *device, int deviceId) { hipError_t hipFuncSetCacheConfig (const void* func, hipFuncCache_t cacheConfig) { - HIP_INIT_API(cacheConfig); + HIP_INIT_API(hipFuncSetCacheConfig, cacheConfig); // No way to set cache config yet. @@ -47,7 +47,7 @@ hipError_t hipFuncSetCacheConfig (const void* func, hipFuncCache_t cacheConfig) hipError_t hipDeviceTotalMem (size_t *bytes, hipDevice_t device) { - HIP_INIT_API(bytes, device); + HIP_INIT_API(hipDeviceTotalMem, bytes, device); if (device < 0 || static_cast(device) >= g_devices.size()) { HIP_RETURN(hipErrorInvalidDevice); @@ -67,7 +67,7 @@ hipError_t hipDeviceTotalMem (size_t *bytes, hipDevice_t device) { hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device) { - HIP_INIT_API(major, minor, device); + HIP_INIT_API(hipDeviceComputeCapability, major, minor, device); if (device < 0 || static_cast(device) >= g_devices.size()) { HIP_RETURN(hipErrorInvalidDevice); @@ -86,7 +86,7 @@ hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device } hipError_t hipDeviceGetCount(int* count) { - HIP_INIT_API(count); + HIP_INIT_API(NONE, count); HIP_RETURN(ihipDeviceGetCount(count)); } @@ -108,7 +108,7 @@ hipError_t ihipDeviceGetCount(int* count) { hipError_t hipDeviceGetName(char *name, int len, hipDevice_t device) { - HIP_INIT_API((void*)name, len, device); + HIP_INIT_API(hipDeviceGetName, (void*)name, len, device); if (device < 0 || static_cast(device) >= g_devices.size()) { HIP_RETURN(hipErrorInvalidDevice); @@ -134,7 +134,7 @@ hipError_t hipDeviceGetName(char *name, int len, hipDevice_t device) { } hipError_t hipGetDeviceProperties ( hipDeviceProp_t* props, hipDevice_t device ) { - HIP_INIT_API(props, device); + HIP_INIT_API(hipGetDeviceProperties, props, device); if (props == nullptr) { HIP_RETURN(hipErrorInvalidValue); @@ -218,7 +218,7 @@ hipError_t hipGetDeviceProperties ( hipDeviceProp_t* props, hipDevice_t device ) } hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc) { - HIP_INIT_API(deviceId, acc); + HIP_INIT_API(NONE, deviceId, acc); assert(0 && "Unimplemented"); @@ -226,7 +226,7 @@ hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc) { } hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** av) { - HIP_INIT_API(stream, av); + HIP_INIT_API(NONE, stream, av); assert(0 && "Unimplemented"); diff --git a/projects/hip/api/hip/hip_device_runtime.cpp b/projects/hip/api/hip/hip_device_runtime.cpp index 6aa9eae29e..b16be0980d 100644 --- a/projects/hip/api/hip/hip_device_runtime.cpp +++ b/projects/hip/api/hip/hip_device_runtime.cpp @@ -26,7 +26,7 @@ THE SOFTWARE. hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* properties) { - HIP_INIT_API(device, properties); + HIP_INIT_API(hipChooseDevice, device, properties); if (device == nullptr || properties == nullptr) { HIP_RETURN(hipErrorInvalidValue); @@ -140,7 +140,7 @@ hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* properties) { hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) { - HIP_INIT_API(pi, attr, device); + HIP_INIT_API(hipDeviceGetAttribute, pi, attr, device); if (pi == nullptr) { HIP_RETURN(hipErrorInvalidValue); @@ -289,7 +289,7 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) hipError_t hipDeviceGetByPCIBusId(int* device, const char*pciBusIdstr) { - HIP_INIT_API(device, pciBusIdstr); + HIP_INIT_API(hipDeviceGetByPCIBusId, device, pciBusIdstr); if (device == nullptr || pciBusIdstr == nullptr) { HIP_RETURN(hipErrorInvalidValue); @@ -319,7 +319,7 @@ hipError_t hipDeviceGetByPCIBusId(int* device, const char*pciBusIdstr) { } hipError_t hipDeviceGetCacheConfig ( hipFuncCache_t * cacheConfig ) { - HIP_INIT_API(cacheConfig); + HIP_INIT_API(hipDeviceGetCacheConfig, cacheConfig); if(cacheConfig == nullptr) { HIP_RETURN(hipErrorInvalidValue); @@ -332,7 +332,7 @@ hipError_t hipDeviceGetCacheConfig ( hipFuncCache_t * cacheConfig ) { hipError_t hipDeviceGetLimit ( size_t* pValue, hipLimit_t limit ) { - HIP_INIT_API(pValue, limit); + HIP_INIT_API(hipDeviceGetLimit, pValue, limit); if(pValue == nullptr) { HIP_RETURN(hipErrorInvalidValue); @@ -357,7 +357,7 @@ hipError_t hipDeviceGetP2PAttribute ( int* value, hipDeviceP2PAttr attr, int sr hipError_t hipDeviceGetPCIBusId ( char* pciBusId, int len, int device ) { - HIP_INIT_API((void*)pciBusId, len, device); + HIP_INIT_API(hipDeviceGetPCIBusId, (void*)pciBusId, len, device); int count; ihipDeviceGetCount(&count); @@ -381,7 +381,7 @@ hipError_t hipDeviceGetPCIBusId ( char* pciBusId, int len, int device ) { } hipError_t hipDeviceGetSharedMemConfig ( hipSharedMemConfig * pConfig ) { - HIP_INIT_API(pConfig); + HIP_INIT_API(hipDeviceGetSharedMemConfig, pConfig); *pConfig = hipSharedMemBankSizeFourByte; @@ -389,7 +389,7 @@ hipError_t hipDeviceGetSharedMemConfig ( hipSharedMemConfig * pConfig ) { } hipError_t hipDeviceReset ( void ) { - HIP_INIT_API(); + HIP_INIT_API(hipDeviceReset); /* FIXME */ @@ -397,7 +397,7 @@ hipError_t hipDeviceReset ( void ) { } hipError_t hipDeviceSetCacheConfig ( hipFuncCache_t cacheConfig ) { - HIP_INIT_API(cacheConfig); + HIP_INIT_API(hipDeviceSetCacheConfig, cacheConfig); // No way to set cache config yet. @@ -409,7 +409,7 @@ hipError_t hipDeviceSetLimit ( hipLimit_t limit, size_t value ) { } hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config ) { - HIP_INIT_API(config); + HIP_INIT_API(hipDeviceSetSharedMemConfig, config); // No way to set cache config yet. @@ -417,7 +417,7 @@ hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config ) { } hipError_t hipDeviceSynchronize ( void ) { - HIP_INIT_API(); + HIP_INIT_API(hipDeviceSynchronize); hip::syncStreams(); @@ -442,7 +442,7 @@ int ihipGetDevice() { } hipError_t hipGetDevice ( int* deviceId ) { - HIP_INIT_API(deviceId); + HIP_INIT_API(hipGetDevice, deviceId); if (deviceId != nullptr) { int dev = ihipGetDevice(); @@ -455,7 +455,7 @@ hipError_t hipGetDevice ( int* deviceId ) { } hipError_t hipGetDeviceCount ( int* count ) { - HIP_INIT_API(count); + HIP_INIT_API(hipGetDeviceCount, count); HIP_RETURN(ihipDeviceGetCount(count)); } @@ -465,7 +465,7 @@ hipError_t hipGetDeviceFlags ( unsigned int* flags ) { } hipError_t hipIpcGetEventHandle ( hipIpcEventHandle_t* handle, hipEvent_t event ) { - HIP_INIT_API(handle, event); + HIP_INIT_API(NONE, handle, event); assert(0 && "Unimplemented"); @@ -473,7 +473,7 @@ hipError_t hipIpcGetEventHandle ( hipIpcEventHandle_t* handle, hipEvent_t event } hipError_t hipIpcOpenEventHandle ( hipEvent_t* event, hipIpcEventHandle_t handle ) { - HIP_INIT_API(event, handle); + HIP_INIT_API(NONE, event, handle); assert(0 && "Unimplemented"); @@ -481,7 +481,7 @@ hipError_t hipIpcOpenEventHandle ( hipEvent_t* event, hipIpcEventHandle_t handle } hipError_t hipSetDevice ( int device ) { - HIP_INIT_API(device); + HIP_INIT_API(hipSetDevice, device); if (static_cast(device) < g_devices.size()) { hip::setCurrentContext(device); @@ -492,7 +492,7 @@ hipError_t hipSetDevice ( int device ) { } hipError_t hipSetDeviceFlags ( unsigned int flags ) { - HIP_INIT_API(flags); + HIP_INIT_API(hipSetDeviceFlags, flags); /* FIXME */ /* Not all of Ctx may be implemented */ @@ -508,7 +508,7 @@ hipError_t hipSetDeviceFlags ( unsigned int flags ) { } hipError_t hipSetValidDevices ( int* device_arr, int len ) { - HIP_INIT_API(device_arr, len); + HIP_INIT_API(NONE, device_arr, len); assert(0 && "Unimplemented"); @@ -516,7 +516,7 @@ hipError_t hipSetValidDevices ( int* device_arr, int len ) { } hipError_t hipExtGetLinkTypeAndHopCount(int device1, int device2, uint32_t* linktype, uint32_t* hopcount) { - HIP_INIT_API(device1, device2, linktype, hopcount); + HIP_INIT_API(hipExtGetLinkTypeAndHopCount, device1, device2, linktype, hopcount); const int numDevices = static_cast(g_devices.size()); diff --git a/projects/hip/api/hip/hip_error.cpp b/projects/hip/api/hip/hip_error.cpp index 5f76e560c0..77b1b8148a 100644 --- a/projects/hip/api/hip/hip_error.cpp +++ b/projects/hip/api/hip/hip_error.cpp @@ -26,7 +26,7 @@ THE SOFTWARE. hipError_t hipGetLastError() { - HIP_INIT_API(); + HIP_INIT_API(hipGetLastError); hipError_t err = hip::g_lastError; hip::g_lastError = hipSuccess; return err; @@ -34,7 +34,7 @@ hipError_t hipGetLastError() hipError_t hipPeekAtLastError() { - HIP_INIT_API(); + HIP_INIT_API(hipPeekAtLastError); hipError_t err = hip::g_lastError; HIP_RETURN(err); } diff --git a/projects/hip/api/hip/hip_event.cpp b/projects/hip/api/hip/hip_event.cpp index 81ae93e279..be2c698f46 100644 --- a/projects/hip/api/hip/hip_event.cpp +++ b/projects/hip/api/hip/hip_event.cpp @@ -182,19 +182,19 @@ hipError_t ihipEventQuery(hipEvent_t event) { } hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) { - HIP_INIT_API(event, flags); + HIP_INIT_API(hipEventCreateWithFlags, event, flags); HIP_RETURN(ihipEventCreateWithFlags(event, flags)); -} +} hipError_t hipEventCreate(hipEvent_t* event) { - HIP_INIT_API(event); + HIP_INIT_API(hipEventCreate, event); HIP_RETURN(ihipEventCreateWithFlags(event, 0)); } hipError_t hipEventDestroy(hipEvent_t event) { - HIP_INIT_API(event); + HIP_INIT_API(hipEventDestroy, event); if (event == nullptr) { HIP_RETURN(hipErrorInvalidResourceHandle); @@ -206,7 +206,7 @@ hipError_t hipEventDestroy(hipEvent_t event) { } hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) { - HIP_INIT_API(ms, start, stop); + HIP_INIT_API(hipEventElapsedTime, ms, start, stop); if (start == nullptr || stop == nullptr) { HIP_RETURN(hipErrorInvalidResourceHandle); @@ -223,7 +223,7 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) { } hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { - HIP_INIT_API(event, stream); + HIP_INIT_API(hipEventRecord, event, stream); if (event == nullptr) { HIP_RETURN(hipErrorInvalidResourceHandle); @@ -246,7 +246,7 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { } hipError_t hipEventSynchronize(hipEvent_t event) { - HIP_INIT_API(event); + HIP_INIT_API(hipEventSynchronize, event); if (event == nullptr) { HIP_RETURN(hipErrorInvalidResourceHandle); @@ -258,7 +258,7 @@ hipError_t hipEventSynchronize(hipEvent_t event) { } hipError_t hipEventQuery(hipEvent_t event) { - HIP_INIT_API(event); + HIP_INIT_API(hipEventQuery, event); HIP_RETURN(ihipEventQuery(event)); } diff --git a/projects/hip/api/hip/hip_intercept.cpp b/projects/hip/api/hip/hip_intercept.cpp new file mode 100644 index 0000000000..b6aed14701 --- /dev/null +++ b/projects/hip/api/hip/hip_intercept.cpp @@ -0,0 +1,58 @@ +/* +Copyright (c) 2019 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hip/hip_runtime.h" +#include "hip_prof_api.h" + +// HIP API callback/activity + +api_callbacks_table_t callbacks_table; + +extern std::string& FunctionName(const hipFunction_t f); +const char* hipKernelNameRef(const hipFunction_t f) { return FunctionName(f).c_str(); } + +hipError_t hipRegisterApiCallback(uint32_t id, void* fun, void* arg) { + return callbacks_table.set_callback(id, reinterpret_cast(fun), arg) ? + hipSuccess : hipErrorInvalidValue; +} + +hipError_t hipRemoveApiCallback(uint32_t id) { + return callbacks_table.set_callback(id, NULL, NULL) ? hipSuccess : hipErrorInvalidValue; +} + +hipError_t hipRegisterActivityCallback(uint32_t id, void* fun, void* arg) { + return callbacks_table.set_activity(id, reinterpret_cast(fun), arg) ? + hipSuccess : hipErrorInvalidValue; +} + +hipError_t hipRemoveActivityCallback(uint32_t id) { + return callbacks_table.set_activity(id, NULL, NULL) ? hipSuccess : hipErrorInvalidValue; +} + +hipError_t hipEnableTracing(bool enabled) { + callbacks_table.set_enabled(enabled); + return hipSuccess; +} + +const char* hipApiName(uint32_t id) { + return hip_api_name(id); +} diff --git a/projects/hip/api/hip/hip_internal.hpp b/projects/hip/api/hip/hip_internal.hpp index 4d34bcbaf5..ba4a300739 100644 --- a/projects/hip/api/hip/hip_internal.hpp +++ b/projects/hip/api/hip/hip_internal.hpp @@ -24,6 +24,7 @@ THE SOFTWARE. #define HIP_SRC_HIP_INTERNAL_H #include "cl_common.hpp" +#include "hip_prof_api.h" #include "trace_helper.h" #include "utils/debug.hpp" #include @@ -51,13 +52,14 @@ typedef struct ihipIpcMemHandle_st { } // This macro should be called at the beginning of every HIP API. -#define HIP_INIT_API(...) \ +#define HIP_INIT_API(cid, ...) \ LogPrintfInfo("[%zx] %s ( %s )", std::this_thread::get_id(), __func__, ToString( __VA_ARGS__ ).c_str()); \ amd::Thread* thread = amd::Thread::current(); \ if (!CL_CHECK_THREAD(thread)) { \ HIP_RETURN(hipErrorOutOfMemory); \ } \ - HIP_INIT(); + HIP_INIT(); \ + HIP_CB_SPAWNER_OBJECT(cid); #define HIP_RETURN(ret) \ hip::g_lastError = ret; \ diff --git a/projects/hip/api/hip/hip_memory.cpp b/projects/hip/api/hip/hip_memory.cpp index fad9d2aefe..5abfb7e365 100644 --- a/projects/hip/api/hip/hip_memory.cpp +++ b/projects/hip/api/hip/hip_memory.cpp @@ -200,7 +200,7 @@ hipError_t ihipMemset(void* dst, int value, size_t valueSize, size_t sizeBytes, } hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flags) { - HIP_INIT_API(ptr, sizeBytes, flags); + HIP_INIT_API(hipExtMallocWithFlags, ptr, sizeBytes, flags); if (flags != hipDeviceMallocDefault && flags != hipDeviceMallocFinegrained) { @@ -211,13 +211,13 @@ hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flag } hipError_t hipMalloc(void** ptr, size_t sizeBytes) { - HIP_INIT_API(ptr, sizeBytes); + HIP_INIT_API(hipMalloc, ptr, sizeBytes); HIP_RETURN(ihipMalloc(ptr, sizeBytes, 0)); } hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { - HIP_INIT_API(ptr, sizeBytes, flags); + HIP_INIT_API(hipHostMalloc, ptr, sizeBytes, flags); if (ptr == nullptr) { HIP_RETURN(hipErrorInvalidValue); @@ -241,7 +241,7 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { } hipError_t hipFree(void* ptr) { - HIP_INIT_API(ptr); + HIP_INIT_API(hipFree, ptr); if (ptr == nullptr) { HIP_RETURN(hipSuccess); @@ -256,7 +256,7 @@ hipError_t hipFree(void* ptr) { } hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) { - HIP_INIT_API(dst, src, sizeBytes, kind); + HIP_INIT_API(hipMemcpy, dst, src, sizeBytes, kind); hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); @@ -264,31 +264,31 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind } hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) { - HIP_INIT_API(dst, value, sizeBytes, stream); + HIP_INIT_API(hipMemsetAsync, dst, value, sizeBytes, stream); HIP_RETURN(ihipMemset(dst, value, sizeof(char), sizeBytes, stream, true)); } hipError_t hipMemsetD32Async(hipDeviceptr_t dst, int value, size_t count, hipStream_t stream) { - HIP_INIT_API(dst, value, count, stream); + HIP_INIT_API(hipMemsetD32Async, dst, value, count, stream); HIP_RETURN(ihipMemset(dst, value, sizeof(int), count * sizeof(int), stream, true)); } hipError_t hipMemset(void* dst, int value, size_t sizeBytes) { - HIP_INIT_API(dst, value, sizeBytes); + HIP_INIT_API(hipMemset, dst, value, sizeBytes); HIP_RETURN(ihipMemset(dst, value, sizeof(char), sizeBytes, nullptr)); } hipError_t hipMemsetD32(hipDeviceptr_t dst, int value, size_t count) { - HIP_INIT_API(dst, value, count); + HIP_INIT_API(hipMemsetD32, dst, value, count); HIP_RETURN(ihipMemset(dst, value, sizeof(int), count * sizeof(int), nullptr)); } hipError_t hipMemPtrGetInfo(void *ptr, size_t *size) { - HIP_INIT_API(ptr, size); + HIP_INIT_API(hipMemPtrGetInfo, ptr, size); size_t offset = 0; amd::Memory* svmMem = getMemoryObject(ptr, offset); @@ -303,7 +303,7 @@ hipError_t hipMemPtrGetInfo(void *ptr, size_t *size) { } hipError_t hipHostFree(void* ptr) { - HIP_INIT_API(ptr); + HIP_INIT_API(hipHostFree, ptr); if (amd::SvmBuffer::malloced(ptr)) { amd::SvmBuffer::free(*hip::getCurrentContext(), ptr); @@ -313,7 +313,7 @@ hipError_t hipHostFree(void* ptr) { } hipError_t hipFreeArray(hipArray* array) { - HIP_INIT_API(array); + HIP_INIT_API(hipFreeArray, array); if (amd::SvmBuffer::malloced(array->data)) { amd::SvmBuffer::free(*hip::getCurrentContext(), array->data); @@ -323,7 +323,7 @@ hipError_t hipFreeArray(hipArray* array) { } hipError_t hipMemGetAddressRange(hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr) { - HIP_INIT_API(pbase, psize, dptr); + HIP_INIT_API(hipMemGetAddressRange, pbase, psize, dptr); // Since we are using SVM buffer DevicePtr and HostPtr is the same void* ptr = dptr; @@ -341,7 +341,7 @@ hipError_t hipMemGetAddressRange(hipDeviceptr_t* pbase, size_t* psize, hipDevice } hipError_t hipMemGetInfo(size_t* free, size_t* total) { - HIP_INIT_API(free, total); + HIP_INIT_API(hipMemGetInfo, free, total); size_t freeMemory[2]; amd::Device* device = hip::getCurrentContext()->devices()[0]; @@ -394,14 +394,14 @@ hipError_t ihipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t heigh hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height) { - HIP_INIT_API(ptr, pitch, width, height); + HIP_INIT_API(hipMallocPitch, ptr, pitch, width, height); const cl_image_format image_format = { CL_R, CL_UNSIGNED_INT8 }; HIP_RETURN(ihipMallocPitch(ptr, pitch, width, height, 1, CL_MEM_OBJECT_IMAGE2D, &image_format)); } hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) { - HIP_INIT_API(pitchedDevPtr, &extent); + HIP_INIT_API(hipMalloc3D, pitchedDevPtr, &extent); size_t pitch = 0; @@ -424,7 +424,7 @@ hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) { } hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent) { - HIP_INIT_API(pitchedDevPtr, value, &extent); + HIP_INIT_API(hipMemset3D, pitchedDevPtr, value, &extent); void *dst = pitchedDevPtr.ptr; size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth; @@ -433,7 +433,7 @@ hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent) } hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) { - HIP_INIT_API(array, pAllocateArray); + HIP_INIT_API(hipArrayCreate, array, pAllocateArray); if (pAllocateArray->Width == 0) { HIP_RETURN(hipErrorInvalidValue); @@ -461,7 +461,7 @@ hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocat hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags) { - HIP_INIT_API(array, desc, width, height, flags); + HIP_INIT_API(hipMallocArray, array, desc, width, height, flags); if (width == 0) { HIP_RETURN(hipErrorInvalidValue); @@ -503,7 +503,7 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, } hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAllocateArray) { - HIP_INIT_API(array, pAllocateArray); + HIP_INIT_API(hipArray3DCreate, array, pAllocateArray); *array = (hipArray*)malloc(sizeof(hipArray)); array[0]->type = pAllocateArray->Flags; @@ -531,7 +531,7 @@ hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAll hipError_t hipMalloc3DArray(hipArray_t* array, const struct hipChannelFormatDesc* desc, struct hipExtent extent, unsigned int flags) { - HIP_INIT_API(array, desc, &extent, flags); + HIP_INIT_API(hipMalloc3DArray, array, desc, &extent, flags); *array = (hipArray*)malloc(sizeof(hipArray)); array[0]->type = flags; @@ -569,7 +569,7 @@ hipError_t hipMalloc3DArray(hipArray_t* array, const struct hipChannelFormatDesc } hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) { - HIP_INIT_API(flagsPtr, hostPtr); + HIP_INIT_API(hipHostGetFlags, flagsPtr, hostPtr); if (flagsPtr == nullptr || hostPtr == nullptr) { @@ -589,7 +589,7 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) { } hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) { - HIP_INIT_API(hostPtr, sizeBytes, flags); + HIP_INIT_API(hipHostRegister, hostPtr, sizeBytes, flags); if(hostPtr != nullptr) { amd::Memory* mem = new (*hip::host_context) amd::Buffer(*hip::host_context, CL_MEM_USE_HOST_PTR | CL_MEM_SVM_ATOMICS, sizeBytes); @@ -616,7 +616,7 @@ hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) } hipError_t hipHostUnregister(void* hostPtr) { - HIP_INIT_API(hostPtr); + HIP_INIT_API(hipHostUnregister, hostPtr); if (amd::SvmBuffer::malloced(hostPtr)) { hip::syncStreams(); @@ -651,7 +651,7 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) { hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t count, size_t offset, hipMemcpyKind kind) { - HIP_INIT_API(symbolName, src, count, offset, kind); + HIP_INIT_API(hipMemcpyToSymbol, symbolName, src, count, offset, kind); size_t sym_size = 0; hipDeviceptr_t device_ptr = nullptr; @@ -675,7 +675,7 @@ hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t cou hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, size_t count, size_t offset, hipMemcpyKind kind) { - HIP_INIT_API(symbolName, dst, count, offset, kind); + HIP_INIT_API(hipMemcpyFromSymbol, symbolName, dst, count, offset, kind); size_t sym_size = 0; hipDeviceptr_t device_ptr = nullptr; @@ -699,7 +699,7 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, size_t count, hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, size_t count, size_t offset, hipMemcpyKind kind, hipStream_t stream) { - HIP_INIT_API(symbolName, src, count, offset, kind, stream); + HIP_INIT_API(hipMemcpyToSymbolAsync, symbolName, src, count, offset, kind, stream); size_t sym_size = 0; hipDeviceptr_t device_ptr = nullptr; @@ -723,7 +723,7 @@ hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, size_ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t count, size_t offset, hipMemcpyKind kind, hipStream_t stream) { - HIP_INIT_API(symbolName, dst, count, offset, kind, stream); + HIP_INIT_API(hipMemcpyFromSymbolAsync, symbolName, dst, count, offset, kind, stream); size_t sym_size = 0; hipDeviceptr_t device_ptr = nullptr; @@ -746,7 +746,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t co } hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) { - HIP_INIT_API(dst, src, sizeBytes); + HIP_INIT_API(hipMemcpyHtoD, dst, src, sizeBytes); hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); @@ -755,7 +755,7 @@ hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) { } hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) { - HIP_INIT_API(dst, src, sizeBytes); + HIP_INIT_API(hipMemcpyDtoH, dst, src, sizeBytes); hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); @@ -764,7 +764,7 @@ hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) { } hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes) { - HIP_INIT_API(dst, src, sizeBytes); + HIP_INIT_API(hipMemcpyDtoD, dst, src, sizeBytes); hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); @@ -773,7 +773,7 @@ hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeByte } hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) { - HIP_INIT_API(dst, src, sizeBytes); + HIP_INIT_API(NONE, dst, src, sizeBytes); hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); @@ -783,7 +783,7 @@ hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) { hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream) { - HIP_INIT_API(dst, src, sizeBytes, kind, stream); + HIP_INIT_API(hipMemcpyAsync, dst, src, sizeBytes, kind, stream); amd::HostQueue* queue = hip::getQueue(stream); @@ -793,7 +793,7 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, hipStream_t stream) { - HIP_INIT_API(dst, src, sizeBytes, stream); + HIP_INIT_API(hipMemcpyHtoDAsync, dst, src, sizeBytes, stream); amd::HostQueue* queue = hip::getQueue(stream); @@ -803,7 +803,7 @@ hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream) { - HIP_INIT_API(dst, src, sizeBytes, stream); + HIP_INIT_API(hipMemcpyDtoDAsync, dst, src, sizeBytes, stream); amd::HostQueue* queue = hip::getQueue(stream); @@ -813,7 +813,7 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream) { - HIP_INIT_API(dst, src, sizeBytes, stream); + HIP_INIT_API(hipMemcpyDtoHAsync, dst, src, sizeBytes, stream); amd::HostQueue* queue = hip::getQueue(stream); @@ -888,7 +888,7 @@ hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch } hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) { - HIP_INIT_API(pCopy); + HIP_INIT_API(hipMemcpyParam2D, pCopy); hipError_t e = hipSuccess; if (pCopy == nullptr) { e = hipErrorInvalidValue; @@ -903,7 +903,7 @@ hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) { hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { - HIP_INIT_API(dst, dpitch, src, spitch, width, height, kind); + HIP_INIT_API(hipMemcpy2D, dst, dpitch, src, spitch, width, height, kind); hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); @@ -914,7 +914,7 @@ hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) { - HIP_INIT_API(dst, dpitch, src, spitch, width, height, kind, stream); + HIP_INIT_API(hipMemcpy2DAsync, dst, dpitch, src, spitch, width, height, kind, stream); amd::HostQueue* queue = hip::getQueue(stream); @@ -923,7 +923,7 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { - HIP_INIT_API(dst, wOffset, hOffset, src, spitch, width, height, kind); + HIP_INIT_API(hipMemcpy2DToArray, dst, wOffset, hOffset, src, spitch, width, height, kind); if (dst->data == nullptr) { HIP_RETURN(hipErrorUnknown); @@ -1018,7 +1018,7 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con hipError_t hipMemcpyToArray(hipArray* dstArray, size_t wOffset, size_t hOffset, const void* src, size_t count, hipMemcpyKind kind) { - HIP_INIT_API(dstArray, wOffset, hOffset, src, count, kind); + HIP_INIT_API(hipMemcpyToArray, dstArray, wOffset, hOffset, src, count, kind); hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); @@ -1058,7 +1058,7 @@ hipError_t hipMemcpyToArray(hipArray* dstArray, size_t wOffset, size_t hOffset, hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffset, size_t hOffset, size_t count, hipMemcpyKind kind) { - HIP_INIT_API(dst, srcArray, wOffset, hOffset, count, kind); + HIP_INIT_API(hipMemcpyFromArray, dst, srcArray, wOffset, hOffset, count, kind); hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); @@ -1097,7 +1097,7 @@ hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffs } hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHost, size_t count) { - HIP_INIT_API(dstArray, dstOffset, srcHost, count); + HIP_INIT_API(hipMemcpyHtoA, dstArray, dstOffset, srcHost, count); hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); @@ -1134,7 +1134,7 @@ hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHo } hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t count) { - HIP_INIT_API(dst, srcArray, srcOffset, count); + HIP_INIT_API(hipMemcpyAtoH, dst, srcArray, srcOffset, count); hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); @@ -1171,7 +1171,7 @@ hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t } hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p) { - HIP_INIT_API(p); + HIP_INIT_API(hipMemcpy3D, p); hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); @@ -1325,7 +1325,7 @@ hipError_t ihipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t } hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height) { - HIP_INIT_API(dst, pitch, value, width, height); + HIP_INIT_API(hipMemset2D, dst, pitch, value, width, height); hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); @@ -1334,7 +1334,7 @@ hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, size_t height, hipStream_t stream) { - HIP_INIT_API(dst, pitch, value, width, height, stream); + HIP_INIT_API(hipMemset2DAsync, dst, pitch, value, width, height, stream); amd::HostQueue* queue = hip::getQueue(stream); @@ -1342,7 +1342,7 @@ hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, } hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes) { - HIP_INIT_API(dst, value, sizeBytes); + HIP_INIT_API(hipMemsetD8, dst, value, sizeBytes); if (dst == nullptr) { HIP_RETURN(hipErrorInvalidValue); @@ -1377,7 +1377,7 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes } hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* dev_ptr) { - HIP_INIT_API(handle, dev_ptr); + HIP_INIT_API(hipIpcGetMemHandle, handle, dev_ptr); size_t offset = 0; amd::Memory* amd_mem_obj = nullptr; @@ -1408,7 +1408,7 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* dev_ptr) { } hipError_t hipIpcOpenMemHandle(void** dev_ptr, hipIpcMemHandle_t handle, unsigned int flags) { - HIP_INIT_API(dev_ptr, &handle, flags); + HIP_INIT_API(hipIpcOpenMemHandle, dev_ptr, &handle, flags); amd::Memory* amd_mem_obj = nullptr; amd::Device* device = nullptr; @@ -1434,7 +1434,7 @@ hipError_t hipIpcOpenMemHandle(void** dev_ptr, hipIpcMemHandle_t handle, unsigne } hipError_t hipIpcCloseMemHandle(void* dev_ptr) { - HIP_INIT_API(dev_ptr); + HIP_INIT_API(hipIpcCloseMemHandle, dev_ptr); size_t offset = 0; amd::Device* device = nullptr; @@ -1479,7 +1479,7 @@ hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannel } hipError_t hipHostGetDevicePointer(void** devicePointer, void* hostPointer, unsigned flags) { - HIP_INIT_API(devicePointer, hostPointer, flags); + HIP_INIT_API(hipHostGetDevicePointer, devicePointer, hostPointer, flags); size_t offset = 0; @@ -1493,7 +1493,7 @@ hipError_t hipHostGetDevicePointer(void** devicePointer, void* hostPointer, unsi } hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void* ptr) { - HIP_INIT_API(attributes, ptr); + HIP_INIT_API(hipPointerGetAttributes, attributes, ptr); size_t offset = 0; amd::Memory* memObj = getMemoryObject(ptr, offset); diff --git a/projects/hip/api/hip/hip_module.cpp b/projects/hip/api/hip/hip_module.cpp index 5f0fe540a9..ad637d3967 100644 --- a/projects/hip/api/hip/hip_module.cpp +++ b/projects/hip/api/hip/hip_module.cpp @@ -30,6 +30,11 @@ THE SOFTWARE. hipError_t ihipModuleLoadData(hipModule_t *module, const void *image); +const std::string& FunctionName(const hipFunction_t f) +{ + return hip::Function::asFunction(f)->function_->name(); +} + static uint64_t ElfSize(const void *emi) { const Elf64_Ehdr *ehdr = (const Elf64_Ehdr*)emi; @@ -53,7 +58,7 @@ static uint64_t ElfSize(const void *emi) hipError_t hipModuleLoad(hipModule_t* module, const char* fname) { - HIP_INIT_API(module, fname); + HIP_INIT_API(hipModuleLoad, module, fname); if (!fname) { HIP_RETURN(hipErrorInvalidValue); @@ -81,7 +86,7 @@ bool ihipModuleUnregisterGlobal(hipModule_t hmod) { hipError_t hipModuleUnload(hipModule_t hmod) { - HIP_INIT_API(hmod); + HIP_INIT_API(hipModuleUnload, hmod); if (hmod == nullptr) { HIP_RETURN(hipErrorUnknown); @@ -100,7 +105,7 @@ hipError_t hipModuleUnload(hipModule_t hmod) hipError_t hipModuleLoadData(hipModule_t *module, const void *image) { - HIP_INIT_API(module, image); + HIP_INIT_API(hipModuleLoadData, module, image); HIP_RETURN(ihipModuleLoadData(module, image)); } @@ -200,7 +205,7 @@ hipError_t ihipModuleLoadData(hipModule_t *module, const void *image) hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod, const char *name) { - HIP_INIT_API(hfunc, hmod, name); + HIP_INIT_API(hipModuleGetFunction, hfunc, hmod, name); amd::Program* program = as_amd(reinterpret_cast(hmod)); @@ -222,7 +227,7 @@ hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod, const ch hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod, const char* name) { - HIP_INIT_API(dptr, bytes, hmod, name); + HIP_INIT_API(hipModuleGetGlobal, dptr, bytes, hmod, name); /* Get address and size for the global symbol */ if (!PlatformState::instance().getGlobalVar(name, ihipGetDevice(), dptr, @@ -235,7 +240,7 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t h hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func) { - HIP_INIT_API(attr, func); + HIP_INIT_API(hipFuncGetAttributes, attr, func); if (!PlatformState::instance().getFuncAttr(func, attr)) { HIP_RETURN(hipErrorUnknown); @@ -253,7 +258,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags = 0, uint32_t params = 0) { - HIP_INIT_API(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, + HIP_INIT_API(NONE, f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, flags, params); hip::Function* function = hip::Function::asFunction(f); @@ -347,7 +352,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, uint32_t sharedMemBytes, hipStream_t hStream, void **kernelParams, void **extra) { - HIP_INIT_API(f, gridDimX, gridDimY, gridDimZ, + HIP_INIT_API(hipModuleLaunchKernel, f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra); @@ -364,7 +369,7 @@ hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, hipStream_t hStream, void** kernelParams, void** extra, hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags) { - HIP_INIT_API(f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, + HIP_INIT_API(NONE, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, flags); @@ -383,7 +388,7 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t gridDimX, hipEvent_t startEvent, hipEvent_t stopEvent) { - HIP_INIT_API(f, gridDimX, gridDimY, gridDimZ, + HIP_INIT_API(NONE, f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent); @@ -400,7 +405,7 @@ hipError_t hipModuleLaunchKernelExt(hipFunction_t f, uint32_t gridDimX, hipEvent_t startEvent, hipEvent_t stopEvent) { - HIP_INIT_API(f, gridDimX, gridDimY, gridDimZ, + HIP_INIT_API(NONE, f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent); @@ -413,7 +418,7 @@ hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim, dim3 blockDim, void **kernelParams, uint32_t sharedMemBytes, hipStream_t hStream) { - HIP_INIT_API(f, gridDim, blockDim, + HIP_INIT_API(hipLaunchCooperativeKernel, f, gridDim, blockDim, sharedMemBytes, hStream); int deviceId = ihipGetDevice(); @@ -464,7 +469,7 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices, unsigned int flags) { - HIP_INIT_API(launchParamsList, numDevices, flags); + HIP_INIT_API(hipLaunchCooperativeKernelMultiDevice, launchParamsList, numDevices, flags); return ihipLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags, (amd::NDRangeKernelCommand::CooperativeGroups | @@ -473,13 +478,13 @@ hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsLi hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices, unsigned int flags) { - HIP_INIT_API(launchParamsList, numDevices, flags); + HIP_INIT_API(hipExtLaunchMultiKernelMultiDevice, launchParamsList, numDevices, flags); return ihipLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags, 0); } hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name) { - HIP_INIT_API(texRef, hmod, name); + HIP_INIT_API(hipModuleGetTexRef, texRef, hmod, name); /* input args check */ if ((texRef == nullptr) || (name == nullptr)) { diff --git a/projects/hip/api/hip/hip_peer.cpp b/projects/hip/api/hip/hip_peer.cpp index cdfec74f0e..95cf206e9c 100644 --- a/projects/hip/api/hip/hip_peer.cpp +++ b/projects/hip/api/hip/hip_peer.cpp @@ -25,7 +25,7 @@ THE SOFTWARE. #include "hip_internal.hpp" hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t peerCtx) { - HIP_INIT_API(canAccessPeer, thisCtx, peerCtx); + HIP_INIT_API(NONE, canAccessPeer, thisCtx, peerCtx); assert(0 && "Unimplemented"); @@ -34,7 +34,7 @@ hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t hipError_t hipMemcpyPeer(void* dst, hipCtx_t dstCtx, const void* src, hipCtx_t srcCtx, size_t sizeBytes) { - HIP_INIT_API(dst, dstCtx, src, srcCtx, sizeBytes); + HIP_INIT_API(NONE, dst, dstCtx, src, srcCtx, sizeBytes); assert(0 && "Unimplemented"); @@ -43,7 +43,7 @@ hipError_t hipMemcpyPeer(void* dst, hipCtx_t dstCtx, const void* src, hipCtx_t s hipError_t hipMemcpyPeerAsync(void* dst, hipCtx_t dstDevice, const void* src, hipCtx_t srcDevice, size_t sizeBytes, hipStream_t stream) { - HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes, stream); + HIP_INIT_API(NONE, dst, dstDevice, src, srcDevice, sizeBytes, stream); assert(0 && "Unimplemented"); @@ -51,7 +51,7 @@ hipError_t hipMemcpyPeerAsync(void* dst, hipCtx_t dstDevice, const void* src, hi } hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, int deviceId, int peerDeviceId) { - HIP_INIT_API(canAccessPeer, deviceId, peerDeviceId); + HIP_INIT_API(hipDeviceCanAccessPeer, canAccessPeer, deviceId, peerDeviceId); amd::Device* device = nullptr; amd::Device* peer_device = nullptr; @@ -83,39 +83,39 @@ hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, int deviceId, int peerDevi } hipError_t hipDeviceDisablePeerAccess(int peerDeviceId) { - HIP_INIT_API(peerDeviceId); + HIP_INIT_API(hipDeviceDisablePeerAccess, peerDeviceId); HIP_RETURN(hipSuccess); } hipError_t hipDeviceEnablePeerAccess(int peerDeviceId, unsigned int flags) { - HIP_INIT_API(peerDeviceId, flags); + HIP_INIT_API(hipDeviceEnablePeerAccess, peerDeviceId, flags); HIP_RETURN(hipSuccess); } hipError_t hipMemcpyPeer(void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes) { - HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes); + HIP_INIT_API(hipMemcpyPeer, dst, dstDevice, src, srcDevice, sizeBytes); HIP_RETURN(hipMemcpy(dst, src, sizeBytes, hipMemcpyDeviceToDevice)); } hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes, hipStream_t stream) { - HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes, stream); + HIP_INIT_API(hipMemcpyPeerAsync, dst, dstDevice, src, srcDevice, sizeBytes, stream); HIP_RETURN(hipMemcpyAsync(dst, src, sizeBytes, hipMemcpyDeviceToDevice, stream)); } hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) { - HIP_INIT_API(peerCtx, flags); + HIP_INIT_API(hipCtxEnablePeerAccess, peerCtx, flags); HIP_RETURN(hipSuccess); } hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx) { - HIP_INIT_API(peerCtx); + HIP_INIT_API(hipCtxDisablePeerAccess, peerCtx); HIP_RETURN(hipSuccess); } diff --git a/projects/hip/api/hip/hip_platform.cpp b/projects/hip/api/hip/hip_platform.cpp index 940bed4d5d..b6ea5070e2 100644 --- a/projects/hip/api/hip/hip_platform.cpp +++ b/projects/hip/api/hip/hip_platform.cpp @@ -425,7 +425,7 @@ extern "C" hipError_t hipConfigureCall( size_t sharedMem, hipStream_t stream) { - HIP_INIT_API(gridDim, blockDim, sharedMem, stream); + HIP_INIT_API(NONE, gridDim, blockDim, sharedMem, stream); PlatformState::instance().configureCall(gridDim, blockDim, sharedMem, stream); @@ -438,7 +438,7 @@ extern "C" hipError_t __hipPushCallConfiguration( size_t sharedMem, hipStream_t stream) { - HIP_INIT_API(gridDim, blockDim, sharedMem, stream); + HIP_INIT_API(NONE, gridDim, blockDim, sharedMem, stream); PlatformState::instance().configureCall(gridDim, blockDim, sharedMem, stream); @@ -449,7 +449,7 @@ extern "C" hipError_t __hipPopCallConfiguration(dim3 *gridDim, dim3 *blockDim, size_t *sharedMem, hipStream_t *stream) { - HIP_INIT_API(gridDim, blockDim, sharedMem, stream); + HIP_INIT_API(NONE, gridDim, blockDim, sharedMem, stream); ihipExec_t exec; PlatformState::instance().popExec(exec); @@ -466,7 +466,7 @@ extern "C" hipError_t hipSetupArgument( size_t size, size_t offset) { - HIP_INIT_API(arg, size, offset); + HIP_INIT_API(NONE, arg, size, offset); PlatformState::instance().setupArgument(arg, size, offset); @@ -475,7 +475,7 @@ extern "C" hipError_t hipSetupArgument( extern "C" hipError_t hipLaunchByPtr(const void *hostFunction) { - HIP_INIT_API(hostFunction); + HIP_INIT_API(NONE, hostFunction); int deviceId = ihipGetDevice(); hipFunction_t func = PlatformState::instance().getFunc(hostFunction, deviceId); @@ -506,7 +506,7 @@ extern "C" hipError_t hipLaunchKernel(const void *hostFunction, size_t sharedMemBytes, hipStream_t stream) { - HIP_INIT_API(hostFunction, gridDim, blockDim, args, sharedMemBytes, + HIP_INIT_API(NONE, hostFunction, gridDim, blockDim, args, sharedMemBytes, stream); int deviceId = ihipGetDevice(); @@ -567,7 +567,7 @@ hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor(int* numBlocks, int blockSize, size_t dynamicSMemSize) { - HIP_INIT_API(f, blockSize, dynamicSMemSize); + HIP_INIT_API(NONE, f, blockSize, dynamicSMemSize); int deviceId = ihipGetDevice(); hipFunction_t func = PlatformState::instance().getFunc(f, deviceId); if (func == nullptr) { diff --git a/projects/hip/api/hip/hip_prof_api.h b/projects/hip/api/hip/hip_prof_api.h new file mode 100644 index 0000000000..32679f403a --- /dev/null +++ b/projects/hip/api/hip/hip_prof_api.h @@ -0,0 +1,255 @@ +/* +Copyright (c) 2019 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef HIP_SRC_HIP_PROF_API_H +#define HIP_SRC_HIP_PROF_API_H + +#include +#include +#include + +#include "hip/hcc_detail/hip_prof_str.h" + +template +class api_callbacks_table_templ { + public: + typedef std::recursive_mutex mutex_t; + + typedef Record record_t; + typedef Fun fun_t; + typedef Act act_t; + + // HIP API callbacks table + struct hip_cb_table_entry_t { + volatile std::atomic sync; + volatile std::atomic sem; + act_t act; + void* a_arg; + fun_t fun; + void* arg; + }; + + struct hip_cb_table_t { + hip_cb_table_entry_t arr[HIP_API_ID_NUMBER]; + }; + + api_callbacks_table_templ() { + memset(&callbacks_table_, 0, sizeof(callbacks_table_)); + } + + bool set_activity(uint32_t id, act_t fun, void* arg) { + std::lock_guard lock(mutex_); + bool ret = true; + if (id == HIP_API_ID_ANY) { + for (unsigned i = 0; i < HIP_API_ID_NUMBER; ++i) set_activity(i, fun, arg); + } else if (id < HIP_API_ID_NUMBER) { + cb_sync(id); + callbacks_table_.arr[id].act = fun; + callbacks_table_.arr[id].a_arg = arg; + enabled_ = true; + cb_release(id); + } else { + ret = false; + } + + return ret; + } + + bool set_callback(uint32_t id, fun_t fun, void* arg) { + std::lock_guard lock(mutex_); + bool ret = true; + if (id == HIP_API_ID_ANY) { + for (unsigned i = 0; i < HIP_API_ID_NUMBER; ++i) set_callback(i, fun, arg); + } else if (id < HIP_API_ID_NUMBER) { + cb_sync(id); + callbacks_table_.arr[id].fun = fun; + callbacks_table_.arr[id].arg = arg; + enabled_ = true; + cb_release(id); + } else { + ret = false; + } + + return ret; + } + + void set_enabled(const bool& enabled) { + enabled_ = enabled; + } + + inline hip_cb_table_entry_t& entry(const uint32_t& id) { + return callbacks_table_.arr[id]; + } + + inline void sem_sync(const uint32_t& id) { + sem_increment(id); + if (entry(id).sync.load() == true) sync_wait(id); + } + + inline void sem_release(const uint32_t& id) { + sem_decrement(id); + } + + inline bool is_enabled() const { + return enabled_; + } + + private: + inline void cb_sync(const uint32_t& id) { + entry(id).sync.store(true); + while (entry(id).sem.load() != 0) {} + } + + inline void cb_release(const uint32_t& id) { + entry(id).sync.store(false); + } + + inline void sem_increment(const uint32_t& id) { + const uint32_t prev = entry(id).sem.fetch_add(1); + if (prev == UINT32_MAX) { + std::cerr << "sem overflow id = " << id << std::endl << std::flush; + abort(); + } + } + + inline void sem_decrement(const uint32_t& id) { + const uint32_t prev = entry(id).sem.fetch_sub(1); + if (prev == 0) { + std::cerr << "sem corrupted id = " << id << std::endl << std::flush; + abort(); + } + } + + void sync_wait(const uint32_t& id) { + sem_decrement(id); + while (entry(id).sync.load() == true) {} + sem_increment(id); + } + + mutex_t mutex_; + hip_cb_table_t callbacks_table_; + bool enabled_; +}; + + +#if USE_PROF_API +#include "prof_protocol.h" + +static const uint32_t HIP_DOMAIN_ID = ACTIVITY_DOMAIN_HIP_API; +typedef activity_record_t hip_api_record_t; +typedef activity_rtapi_callback_t hip_api_callback_t; +typedef activity_sync_callback_t hip_act_callback_t; + +// HIP API callbacks spawner object macro +#define HIP_CB_SPAWNER_OBJECT(CB_ID) \ + api_callbacks_spawner_t __api_tracer(HIP_API_ID_##CB_ID); \ + { \ + hip_api_data_t* api_data = __api_tracer.get_api_data_ptr(); \ + if (api_data != NULL) { \ + hip_api_data_t& api_data_ref = *api_data; \ + INIT_CB_ARGS_DATA(CB_ID, api_data_ref); \ + __api_tracer.call(); \ + } \ + } + +typedef api_callbacks_table_templ api_callbacks_table_t; +extern api_callbacks_table_t callbacks_table; + +template +class api_callbacks_spawner_t { + public: + api_callbacks_spawner_t(const hip_api_id_t& cid) : + api_data_(NULL) + { + if (!is_enabled()) return; + + if (cid_ >= HIP_API_ID_NUMBER) { + fprintf(stderr, "HIP %s bad id %d\n", __FUNCTION__, cid_); + abort(); + } + callbacks_table.sem_sync(cid_); + + hip_act_callback_t act = entry(cid_).act; + if (act != NULL) api_data_ = (hip_api_data_t*) act(cid_, NULL, NULL, NULL); + } + + void call() { + hip_api_callback_t fun = entry(cid_).fun; + void* arg = entry(cid_).arg; + if (fun != NULL) fun(HIP_DOMAIN_ID, cid_, api_data_, arg); + } + + ~api_callbacks_spawner_t() { + if (api_data_ == NULL) return; + + hip_act_callback_t act = entry(cid_).act; + void* a_arg = entry(cid_).a_arg; + hip_api_callback_t fun = entry(cid_).fun; + void* arg = entry(cid_).arg; + if (act != NULL) act(cid_, NULL, api_data_, a_arg); + if (fun != NULL) fun(HIP_DOMAIN_ID, cid_, api_data_, arg); + + callbacks_table.sem_release(cid_); + } + + hip_api_data_t* get_api_data_ptr() { + return api_data_; + } + + bool is_enabled() const { + return callbacks_table.is_enabled(); + } + + private: + inline api_callbacks_table_t::hip_cb_table_entry_t& entry(const uint32_t& id) { + return callbacks_table.entry(id); + } + + hip_api_data_t* api_data_; +}; + +template <> +class api_callbacks_spawner_t { + public: + api_callbacks_spawner_t(const hip_api_id_t& cid) {} + void call() {} + hip_api_data_t* get_api_data_ptr() { return NULL; } + bool is_enabled() const { return false; } +}; + +#else + +#define HIP_CB_SPAWNER_OBJECT(x) do {} while(0) + +class api_callbacks_table_t { + public: + typedef void* act_t; + typedef void* fun_t; + bool set_activity(uint32_t id, act_t fun, void* arg) { return false; } + bool set_callback(uint32_t id, fun_t fun, void* arg) { return false; } +}; + +#endif + +#endif // HIP_SRC_HIP_PROF_API_H diff --git a/projects/hip/api/hip/hip_profile.cpp b/projects/hip/api/hip/hip_profile.cpp index 8b3d42cbbc..b8c8f28e94 100644 --- a/projects/hip/api/hip/hip_profile.cpp +++ b/projects/hip/api/hip/hip_profile.cpp @@ -25,7 +25,7 @@ THE SOFTWARE. #include "hip_internal.hpp" hipError_t hipProfilerStart() { - HIP_INIT_API(); + HIP_INIT_API(hipProfilerStart); assert(0 && "Unimplemented"); @@ -34,7 +34,7 @@ hipError_t hipProfilerStart() { hipError_t hipProfilerStop() { - HIP_INIT_API(); + HIP_INIT_API(hipProfilerStop); assert(0 && "Unimplemented"); diff --git a/projects/hip/api/hip/hip_stream.cpp b/projects/hip/api/hip/hip_stream.cpp index ab26b4fc2e..14439ebc40 100644 --- a/projects/hip/api/hip/hip_stream.cpp +++ b/projects/hip/api/hip/hip_stream.cpp @@ -120,19 +120,19 @@ static hipError_t ihipStreamCreate(hipStream_t *stream, unsigned int flags, amd: } hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) { - HIP_INIT_API(stream, flags); + HIP_INIT_API(hipStreamCreateWithFlags, stream, flags); HIP_RETURN(ihipStreamCreate(stream, flags, amd::CommandQueue::Priority::Normal)); } hipError_t hipStreamCreate(hipStream_t *stream) { - HIP_INIT_API(stream); + HIP_INIT_API(hipStreamCreate, stream); HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, amd::CommandQueue::Priority::Normal)); } hipError_t hipStreamCreateWithPriority(hipStream_t* stream, unsigned int flags, int priority) { - HIP_INIT_API(stream, flags, priority); + HIP_INIT_API(hipStreamCreateWithPriority, stream, flags, priority); if (priority > static_cast(amd::CommandQueue::Priority::High)) { priority = static_cast(amd::CommandQueue::Priority::High); @@ -144,7 +144,7 @@ hipError_t hipStreamCreateWithPriority(hipStream_t* stream, unsigned int flags, } hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPriority) { - HIP_INIT_API(leastPriority, greatestPriority); + HIP_INIT_API(hipDeviceGetStreamPriorityRange, leastPriority, greatestPriority); if (leastPriority != nullptr) { *leastPriority = static_cast(amd::CommandQueue::Priority::Normal); @@ -157,7 +157,7 @@ hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPrio } hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags) { - HIP_INIT_API(stream, flags); + HIP_INIT_API(hipStreamGetFlags, stream, flags); hip::Stream* hStream = reinterpret_cast(stream); @@ -171,7 +171,7 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags) { } hipError_t hipStreamSynchronize(hipStream_t stream) { - HIP_INIT_API(stream); + HIP_INIT_API(hipStreamSynchronize, stream); amd::HostQueue* hostQueue = hip::getQueue(stream); hostQueue->finish(); @@ -180,7 +180,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream) { } hipError_t hipStreamDestroy(hipStream_t stream) { - HIP_INIT_API(stream); + HIP_INIT_API(hipStreamDestroy, stream); if (stream == nullptr) { HIP_RETURN(hipErrorInvalidResourceHandle); @@ -199,7 +199,7 @@ hipError_t hipStreamDestroy(hipStream_t stream) { } hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) { - HIP_INIT_API(stream, event, flags); + HIP_INIT_API(hipStreamWaitEvent, stream, event, flags); amd::HostQueue* queue; @@ -219,7 +219,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int } hipError_t hipStreamQuery(hipStream_t stream) { - HIP_INIT_API(stream); + HIP_INIT_API(hipStreamQuery, stream); amd::HostQueue* hostQueue; if (stream == nullptr) { @@ -242,7 +242,7 @@ hipError_t hipStreamQuery(hipStream_t stream) { hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData, unsigned int flags) { - HIP_INIT_API(stream, callback, userData, flags); + HIP_INIT_API(hipStreamAddCallback, stream, callback, userData, flags); amd::HostQueue* hostQueue = reinterpret_cast (stream)->asHostQueue(); diff --git a/projects/hip/api/hip/hip_surface.cpp b/projects/hip/api/hip/hip_surface.cpp index 2323209cae..42741add7a 100644 --- a/projects/hip/api/hip/hip_surface.cpp +++ b/projects/hip/api/hip/hip_surface.cpp @@ -47,7 +47,7 @@ using namespace hip; hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject, const hipResourceDesc* pResDesc) { - HIP_INIT_API(pSurfObject, pResDesc); + HIP_INIT_API(NONE, pSurfObject, pResDesc); hipSurface* pSurface = new hipSurface(pResDesc); assert(pSurface != nullptr); @@ -82,7 +82,7 @@ hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject, hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject) { - HIP_INIT_API(surfaceObject); + HIP_INIT_API(NONE, surfaceObject); amd::ScopedLock lock(surfaceLock); hipSurface* pSurface = surfaceHash[surfaceObject]; diff --git a/projects/hip/api/hip/hip_texture.cpp b/projects/hip/api/hip/hip_texture.cpp index 5809a75200..28328932c3 100644 --- a/projects/hip/api/hip/hip_texture.cpp +++ b/projects/hip/api/hip/hip_texture.cpp @@ -215,7 +215,7 @@ hip::TextureObject* ihipCreateTextureObject(const hipResourceDesc& resDesc, amd: hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResourceDesc* pResDesc, const hipTextureDesc* pTexDesc, const hipResourceViewDesc* pResViewDesc) { - HIP_INIT_API(pTexObject, pResDesc, pTexDesc, pResViewDesc); + HIP_INIT_API(NONE, pTexObject, pResDesc, pTexDesc, pResViewDesc); amd::Device* device = hip::getCurrentContext()->devices()[0]; @@ -307,7 +307,7 @@ void ihipDestroyTextureObject(hip::TextureObject* texture) { } hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject) { - HIP_INIT_API(textureObject); + HIP_INIT_API(NONE, textureObject); hip::TextureObject* texture = reinterpret_cast(textureObject); @@ -318,7 +318,7 @@ hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject) { hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc, hipTextureObject_t textureObject) { - HIP_INIT_API(pResDesc, textureObject); + HIP_INIT_API(NONE, pResDesc, textureObject); hip::TextureObject* texture = reinterpret_cast(textureObject); @@ -331,7 +331,7 @@ hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc, hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc, hipTextureObject_t textureObject) { - HIP_INIT_API(pResViewDesc, textureObject); + HIP_INIT_API(NONE, pResViewDesc, textureObject); assert(0 && "Unimplemented"); @@ -340,7 +340,7 @@ hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc, hipTextureObject_t textureObject) { - HIP_INIT_API(pTexDesc, textureObject); + HIP_INIT_API(NONE, pTexDesc, textureObject); assert(0 && "Unimplemented"); @@ -416,7 +416,7 @@ hipError_t ihipBindTexture(cl_mem_object_type type, hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* devPtr, const hipChannelFormatDesc* desc, size_t size) { - HIP_INIT_API(offset, tex, devPtr, desc, size); + HIP_INIT_API(NONE, offset, tex, devPtr, desc, size); if (desc == nullptr) { HIP_RETURN(hipErrorInvalidValue); @@ -432,14 +432,14 @@ hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* dev hipError_t hipBindTexture2D(size_t* offset, textureReference* tex, const void* devPtr, const hipChannelFormatDesc* desc, size_t width, size_t height, size_t pitch) { - HIP_INIT_API(offset, tex, devPtr, desc, width, height, pitch); + HIP_INIT_API(NONE, offset, tex, devPtr, desc, width, height, pitch); HIP_RETURN(ihipBindTexture(CL_MEM_OBJECT_IMAGE2D, offset, tex, devPtr, desc, width, height, pitch)); } hipError_t hipBindTextureToArray(textureReference* tex, hipArray_const_t array, const hipChannelFormatDesc* desc) { - HIP_INIT_API(tex, array, desc); + HIP_INIT_API(NONE, tex, array, desc); assert(0 && "Unimplemented"); @@ -449,7 +449,7 @@ hipError_t hipBindTextureToArray(textureReference* tex, hipArray_const_t array, hipError_t ihipBindTextureImpl(TlsData* tls, int dim, enum hipTextureReadMode readMode, size_t* offset, const void* devPtr, const struct hipChannelFormatDesc* desc, size_t size, textureReference* tex) { - HIP_INIT_API(dim, readMode, offset, devPtr, size, tex); + HIP_INIT_API(NONE, dim, readMode, offset, devPtr, size, tex); assert(1 == dim); @@ -460,7 +460,7 @@ hipError_t ihipBindTextureToArrayImpl(TlsData* tls, int dim, enum hipTextureRead hipArray_const_t array, const struct hipChannelFormatDesc& desc, textureReference* tex) { - HIP_INIT_API(dim, readMode, &desc, array, tex); + HIP_INIT_API(NONE, dim, readMode, &desc, array, tex); cl_mem_object_type clType; size_t offset = 0; @@ -483,7 +483,7 @@ hipError_t ihipBindTextureToArrayImpl(TlsData* tls, int dim, enum hipTextureRead hipError_t hipBindTextureToMipmappedArray(textureReference* tex, hipMipmappedArray_const_t mipmappedArray, const hipChannelFormatDesc* desc) { - HIP_INIT_API(tex, mipmappedArray, desc); + HIP_INIT_API(NONE, tex, mipmappedArray, desc); assert(0 && "Unimplemented"); @@ -491,7 +491,7 @@ hipError_t hipBindTextureToMipmappedArray(textureReference* tex, } hipError_t hipUnbindTexture(const textureReference* tex) { - HIP_INIT_API(tex); + HIP_INIT_API(NONE, tex); ihipDestroyTextureObject(reinterpret_cast(tex->textureObject)); @@ -499,7 +499,7 @@ hipError_t hipUnbindTexture(const textureReference* tex) { } hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_const_t array) { - HIP_INIT_API(desc, array); + HIP_INIT_API(NONE, desc, array); if (desc != nullptr) { *desc = array->desc; @@ -509,7 +509,7 @@ hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_const_t array) } hipError_t hipGetTextureAlignmentOffset(size_t* offset, const textureReference* tex) { - HIP_INIT_API(offset, tex); + HIP_INIT_API(NONE, offset, tex); if ((offset == nullptr) || (tex == nullptr)) { HIP_RETURN(hipErrorInvalidValue); @@ -521,7 +521,7 @@ hipError_t hipGetTextureAlignmentOffset(size_t* offset, const textureReference* } hipError_t hipGetTextureReference(const textureReference** tex, const void* symbol) { - HIP_INIT_API(tex, symbol); + HIP_INIT_API(NONE, tex, symbol); assert(0 && "Unimplemented"); @@ -529,7 +529,7 @@ hipError_t hipGetTextureReference(const textureReference** tex, const void* symb } hipError_t hipTexRefSetFormat(textureReference* tex, hipArray_Format fmt, int NumPackedComponents) { - HIP_INIT_API(tex, fmt, NumPackedComponents); + HIP_INIT_API(NONE, tex, fmt, NumPackedComponents); if (tex == nullptr) { HIP_RETURN(hipErrorInvalidImage); @@ -542,7 +542,7 @@ hipError_t hipTexRefSetFormat(textureReference* tex, hipArray_Format fmt, int Nu } hipError_t hipTexRefSetFlags(textureReference* tex, unsigned int flags) { - HIP_INIT_API(tex, flags); + HIP_INIT_API(NONE, tex, flags); if (tex == nullptr) { HIP_RETURN(hipErrorInvalidImage); @@ -554,7 +554,7 @@ hipError_t hipTexRefSetFlags(textureReference* tex, unsigned int flags) { } hipError_t hipTexRefSetFilterMode(textureReference* tex, hipTextureFilterMode fm) { - HIP_INIT_API(tex, fm); + HIP_INIT_API(NONE, tex, fm); if (tex == nullptr) { HIP_RETURN(hipErrorInvalidImage); @@ -566,7 +566,7 @@ hipError_t hipTexRefSetFilterMode(textureReference* tex, hipTextureFilterMode fm } hipError_t hipTexRefGetAddressMode(hipTextureAddressMode* am, textureReference tex, int dim) { - HIP_INIT_API(am, &tex, dim); + HIP_INIT_API(NONE, am, &tex, dim); if ((am == nullptr) || (dim >= 3)) { HIP_RETURN(hipErrorInvalidValue); @@ -578,7 +578,7 @@ hipError_t hipTexRefGetAddressMode(hipTextureAddressMode* am, textureReference t } hipError_t hipTexRefSetAddressMode(textureReference* tex, int dim, hipTextureAddressMode am) { - HIP_INIT_API(tex, dim, am); + HIP_INIT_API(NONE, tex, dim, am); if (tex == nullptr) { HIP_RETURN(hipErrorInvalidImage); @@ -590,7 +590,7 @@ hipError_t hipTexRefSetAddressMode(textureReference* tex, int dim, hipTextureAdd } hipError_t hipTexRefGetArray(hipArray_t* array, textureReference tex) { - HIP_INIT_API(array, &tex); + HIP_INIT_API(NONE, array, &tex); hip::TextureObject* texture = nullptr; @@ -613,7 +613,7 @@ hipError_t hipTexRefGetArray(hipArray_t* array, textureReference tex) { } hipError_t hipTexRefSetArray(textureReference* tex, hipArray_const_t array, unsigned int flags) { - HIP_INIT_API(tex, array, flags); + HIP_INIT_API(NONE, tex, array, flags); size_t offset = 0; @@ -626,7 +626,7 @@ hipError_t hipTexRefSetArray(textureReference* tex, hipArray_const_t array, unsi } hipError_t hipTexRefGetAddress(hipDeviceptr_t* dev_ptr, textureReference tex) { - HIP_INIT_API(dev_ptr, &tex); + HIP_INIT_API(NONE, dev_ptr, &tex); hip::TextureObject* texture = nullptr; device::Memory* dev_mem = nullptr; @@ -648,7 +648,7 @@ hipError_t hipTexRefGetAddress(hipDeviceptr_t* dev_ptr, textureReference tex) { hipError_t hipTexRefSetAddress(size_t* offset, textureReference* tex, hipDeviceptr_t devPtr, size_t size) { - HIP_INIT_API(offset, tex, devPtr, size); + HIP_INIT_API(NONE, offset, tex, devPtr, size); if (tex == nullptr) { HIP_RETURN(hipErrorInvalidImage); @@ -664,7 +664,7 @@ hipError_t hipTexRefSetAddress(size_t* offset, textureReference* tex, hipDevicep hipError_t hipTexRefSetAddress2D(textureReference* tex, const HIP_ARRAY_DESCRIPTOR* desc, hipDeviceptr_t devPtr, size_t pitch) { - HIP_INIT_API(tex, desc, devPtr, pitch); + HIP_INIT_API(NONE, tex, desc, devPtr, pitch); if (desc == nullptr) { HIP_RETURN(hipErrorInvalidValue); diff --git a/projects/hip/api/hip/prof_protocol.h b/projects/hip/api/hip/prof_protocol.h new file mode 100644 index 0000000000..f058968ca3 --- /dev/null +++ b/projects/hip/api/hip/prof_protocol.h @@ -0,0 +1,90 @@ +/* +Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef INC_EXT_PROF_PROTOCOL_H_ +#define INC_EXT_PROF_PROTOCOL_H_ + +#include + +// Traced API domains +typedef enum { + ACTIVITY_DOMAIN_HSA_API = 0, // HSA API domain + ACTIVITY_DOMAIN_HSA_OPS = 1, // HSA async activity domain + ACTIVITY_DOMAIN_HCC_OPS = 2, // HCC async activity domain + ACTIVITY_DOMAIN_HIP_API = 3, // HIP API domain + ACTIVITY_DOMAIN_EXT_API = 4, // External ID domain + ACTIVITY_DOMAIN_ROCTX = 5, // ROCTX domain + ACTIVITY_DOMAIN_NUMBER +} activity_domain_t; + +// Extension API opcodes +typedef enum { + ACTIVITY_EXT_OP_MARK = 0, + ACTIVITY_EXT_OP_EXTERN_ID = 1 +} activity_ext_op_t; + +// API calback type +typedef void (*activity_rtapi_callback_t)(uint32_t domain, uint32_t cid, const void* data, void* arg); +typedef uint32_t activity_kind_t; +typedef uint32_t activity_op_t; + +// API callback phase +typedef enum { + ACTIVITY_API_PHASE_ENTER = 0, + ACTIVITY_API_PHASE_EXIT = 1 +} activity_api_phase_t; + +// Trace record types +// Correlation id +typedef uint64_t activity_correlation_id_t; + +// Activity record type +struct activity_record_t { + uint32_t domain; // activity domain id + activity_kind_t kind; // activity kind + activity_op_t op; // activity op + activity_correlation_id_t correlation_id; // activity ID + uint64_t begin_ns; // host begin timestamp + uint64_t end_ns; // host end timestamp + union { + struct { + int device_id; // device id + uint64_t queue_id; // queue id + }; + struct { + uint32_t process_id; // device id + uint32_t thread_id; // thread id + }; + struct { + activity_correlation_id_t external_id; // external correlatino id + }; + }; + size_t bytes; // data size bytes +}; + +// Activity sync calback type +typedef void* (*activity_sync_callback_t)(uint32_t cid, activity_record_t* record, const void* data, void* arg); +// Activity async calback type +typedef void (*activity_id_callback_t)(activity_correlation_id_t id); +typedef void (*activity_async_callback_t)(uint32_t op, void* record, void* arg); + +#endif // INC_EXT_PROF_PROTOCOL_H_