From 96b7f0404eacccffa9ef35a50852fc413d486066 Mon Sep 17 00:00:00 2001 From: Evgeny Date: Thu, 8 Nov 2018 08:36:50 -0600 Subject: [PATCH] renaming HIP_INIT_CB_API to HIP_INIT_API [ROCm/clr commit: e362688adf0697b4a8cb74e8d949c6d18fdc0136] --- .../include/hip/hcc_detail/hip_prof_api.h | 8 +- .../include/hip/hcc_detail/hip_prof_str.h | 65 +++++++++++ projects/clr/hipamd/src/hip_context.cpp | 46 ++++---- projects/clr/hipamd/src/hip_device.cpp | 40 +++---- projects/clr/hipamd/src/hip_error.cpp | 8 +- projects/clr/hipamd/src/hip_event.cpp | 14 +-- projects/clr/hipamd/src/hip_hcc.cpp | 8 +- projects/clr/hipamd/src/hip_hcc_internal.h | 12 +-- projects/clr/hipamd/src/hip_memory.cpp | 102 +++++++++--------- projects/clr/hipamd/src/hip_module.cpp | 18 ++-- projects/clr/hipamd/src/hip_peer.cpp | 20 ++-- projects/clr/hipamd/src/hip_stream.cpp | 16 +-- projects/clr/hipamd/src/hip_surface.cpp | 4 +- projects/clr/hipamd/src/hip_texture.cpp | 40 +++---- 14 files changed, 230 insertions(+), 171 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_prof_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_prof_api.h index 98154873f3..8589bc0ee6 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_prof_api.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_prof_api.h @@ -127,9 +127,11 @@ typedef activity_sync_callback_t hip_act_callback_t; // HIP API callbacks spawner object macro #define HIP_CB_SPAWNER_OBJECT(CB_ID) \ - hip_api_data_t api_data{}; \ - INIT_CB_ARGS_DATA(CB_ID, api_data); \ - api_callbacks_spawner_t __api_tracer(HIP_API_ID_##CB_ID, api_data); + if (HIP_API_ID_##CB_ID < HIP_API_ID_NUMBER) { \ + hip_api_data_t api_data{}; \ + INIT_CB_ARGS_DATA(CB_ID, api_data); \ + api_callbacks_spawner_t __api_tracer(HIP_API_ID_##CB_ID, api_data); \ + } typedef api_callbacks_table_templ #include +// Dummy API callbacks definition +#define INIT_hipHccModuleLaunchKernel_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipHccGetAccelerator_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipHccGetAcceleratorView_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipDeviceCanAccessPeer2_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipMemcpyPeer2_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipMemcpyPeerAsync2_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipCreateTextureObject_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipDestroyTextureObject_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipGetTextureObjectResourceDesc_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipGetTextureObjectResourceViewDesc_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipGetTextureObjectTextureDesc_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipBindTexture_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipBindTexture2D_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipBindTextureToArray_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipBindTextureToMipmappedArray_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipUnbindTexture_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipGetChannelDesc_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipGetTextureAlignmentOffset_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipGetTextureReference_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipTexRefSetFormat_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipTexRefSetFlags_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipTexRefSetFilterMode_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipTexRefSetAddressMode_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipTexRefSetArray_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipTexRefSetAddress_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipTexRefSetAddress2D_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipMemcpyHtoH_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipGetErrorName_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipGetErrorString_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipCreateSurfaceObject_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipDestroySurfaceObject_CB_ARGS_DATA(cb_data) {}; + // HIP API callbacks ID enumaration enum hip_api_id_t { HIP_API_ID_hipHostFree = 0, @@ -146,6 +179,38 @@ enum hip_api_id_t { HIP_API_ID_hipGetDeviceCount = 137, HIP_API_ID_NUMBER = 138, HIP_API_ID_ANY = 139, + + HIP_API_ID_hipHccModuleLaunchKernel = HIP_API_ID_NUMBER, + HIP_API_ID_hipHccGetAccelerator = HIP_API_ID_NUMBER, + HIP_API_ID_hipHccGetAcceleratorView = HIP_API_ID_NUMBER, + HIP_API_ID_hipDeviceCanAccessPeer2 = HIP_API_ID_NUMBER, + HIP_API_ID_hipMemcpyPeer2 = HIP_API_ID_NUMBER, + HIP_API_ID_hipMemcpyPeerAsync2 = HIP_API_ID_NUMBER, + HIP_API_ID_hipCreateTextureObject = HIP_API_ID_NUMBER, + HIP_API_ID_hipDestroyTextureObject = HIP_API_ID_NUMBER, + HIP_API_ID_hipGetTextureObjectResourceDesc = HIP_API_ID_NUMBER, + HIP_API_ID_hipGetTextureObjectResourceViewDesc = HIP_API_ID_NUMBER, + HIP_API_ID_hipGetTextureObjectTextureDesc = HIP_API_ID_NUMBER, + HIP_API_ID_hipBindTexture = HIP_API_ID_NUMBER, + HIP_API_ID_hipBindTexture2D = HIP_API_ID_NUMBER, + HIP_API_ID_hipBindTextureToArray = HIP_API_ID_NUMBER, + HIP_API_ID_hipBindTextureToMipmappedArray = HIP_API_ID_NUMBER, + HIP_API_ID_hipUnbindTexture = HIP_API_ID_NUMBER, + HIP_API_ID_hipGetChannelDesc = HIP_API_ID_NUMBER, + HIP_API_ID_hipGetTextureAlignmentOffset = HIP_API_ID_NUMBER, + HIP_API_ID_hipGetTextureReference = HIP_API_ID_NUMBER, + HIP_API_ID_hipTexRefSetFormat = HIP_API_ID_NUMBER, + HIP_API_ID_hipTexRefSetFlags = HIP_API_ID_NUMBER, + HIP_API_ID_hipTexRefSetFilterMode = HIP_API_ID_NUMBER, + HIP_API_ID_hipTexRefSetAddressMode = HIP_API_ID_NUMBER, + HIP_API_ID_hipTexRefSetArray = HIP_API_ID_NUMBER, + HIP_API_ID_hipTexRefSetAddress = HIP_API_ID_NUMBER, + HIP_API_ID_hipTexRefSetAddress2D = HIP_API_ID_NUMBER, + HIP_API_ID_hipMemcpyHtoH = HIP_API_ID_NUMBER, + HIP_API_ID_hipGetErrorName = HIP_API_ID_NUMBER, + HIP_API_ID_hipGetErrorString = HIP_API_ID_NUMBER, + HIP_API_ID_hipCreateSurfaceObject = HIP_API_ID_NUMBER, + HIP_API_ID_hipDestroySurfaceObject = HIP_API_ID_NUMBER, }; // Return HIP API string diff --git a/projects/clr/hipamd/src/hip_context.cpp b/projects/clr/hipamd/src/hip_context.cpp index f03d3bc7b5..e85c406cfc 100644 --- a/projects/clr/hipamd/src/hip_context.cpp +++ b/projects/clr/hipamd/src/hip_context.cpp @@ -40,7 +40,7 @@ void ihipCtxStackUpdate() { } hipError_t hipInit(unsigned int flags) { - HIP_INIT_CB_API(hipInit, flags); + HIP_INIT_API(hipInit, flags); hipError_t e = hipSuccess; @@ -53,7 +53,7 @@ hipError_t hipInit(unsigned int flags) { } hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int flags, hipDevice_t device) { - HIP_INIT_CB_API(hipCtxCreate, ctx, flags, device); // FIXME - review if we want to init + HIP_INIT_API(hipCtxCreate, ctx, flags, device); // FIXME - review if we want to init hipError_t e = hipSuccess; auto deviceHandle = ihipGetDevice(device); { @@ -71,7 +71,7 @@ hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int flags, hipDevice_t device) { } hipError_t hipDeviceGet(hipDevice_t* device, int deviceId) { - HIP_INIT_CB_API(hipDeviceGet, device, deviceId); // FIXME - review if we want to init + HIP_INIT_API(hipDeviceGet, device, deviceId); // FIXME - review if we want to init auto deviceHandle = ihipGetDevice(deviceId); @@ -86,7 +86,7 @@ hipError_t hipDeviceGet(hipDevice_t* device, int deviceId) { }; hipError_t hipDriverGetVersion(int* driverVersion) { - HIP_INIT_CB_API(hipDriverGetVersion, driverVersion); + HIP_INIT_API(hipDriverGetVersion, driverVersion); hipError_t e = hipSuccess; if (driverVersion) { *driverVersion = 4; @@ -98,7 +98,7 @@ hipError_t hipDriverGetVersion(int* driverVersion) { } hipError_t hipRuntimeGetVersion(int* runtimeVersion) { - HIP_INIT_CB_API(hipRuntimeGetVersion, runtimeVersion); + HIP_INIT_API(hipRuntimeGetVersion, runtimeVersion); hipError_t e = hipSuccess; if (runtimeVersion) { *runtimeVersion = HIP_VERSION_PATCH; @@ -110,7 +110,7 @@ hipError_t hipRuntimeGetVersion(int* runtimeVersion) { } hipError_t hipCtxDestroy(hipCtx_t ctx) { - HIP_INIT_CB_API(hipCtxDestroy, ctx); + HIP_INIT_API(hipCtxDestroy, ctx); hipError_t e = hipSuccess; ihipCtx_t* currentCtx = ihipGetTlsDefaultCtx(); ihipCtx_t* primaryCtx = ((ihipDevice_t*)ctx->getDevice())->_primaryCtx; @@ -134,7 +134,7 @@ hipError_t hipCtxDestroy(hipCtx_t ctx) { } hipError_t hipCtxPopCurrent(hipCtx_t* ctx) { - HIP_INIT_CB_API(hipCtxPopCurrent, ctx); + HIP_INIT_API(hipCtxPopCurrent, ctx); hipError_t e = hipSuccess; ihipCtx_t* currentCtx = ihipGetTlsDefaultCtx(); auto deviceHandle = currentCtx->getDevice(); @@ -155,7 +155,7 @@ hipError_t hipCtxPopCurrent(hipCtx_t* ctx) { } hipError_t hipCtxPushCurrent(hipCtx_t ctx) { - HIP_INIT_CB_API(hipCtxPushCurrent, ctx); + HIP_INIT_API(hipCtxPushCurrent, ctx); hipError_t e = hipSuccess; if (ctx != NULL) { // TODO- is this check needed? ihipSetTlsDefaultCtx(ctx); @@ -168,7 +168,7 @@ hipError_t hipCtxPushCurrent(hipCtx_t ctx) { } hipError_t hipCtxGetCurrent(hipCtx_t* ctx) { - HIP_INIT_CB_API(hipCtxGetCurrent, ctx); + HIP_INIT_API(hipCtxGetCurrent, ctx); hipError_t e = hipSuccess; if ((tls_getPrimaryCtx) || tls_ctxStack.empty()) { *ctx = ihipGetTlsDefaultCtx(); @@ -179,7 +179,7 @@ hipError_t hipCtxGetCurrent(hipCtx_t* ctx) { } hipError_t hipCtxSetCurrent(hipCtx_t ctx) { - HIP_INIT_CB_API(hipCtxSetCurrent, ctx); + HIP_INIT_API(hipCtxSetCurrent, ctx); hipError_t e = hipSuccess; if (ctx == NULL) { tls_ctxStack.pop(); @@ -192,7 +192,7 @@ hipError_t hipCtxSetCurrent(hipCtx_t ctx) { } hipError_t hipCtxGetDevice(hipDevice_t* device) { - HIP_INIT_CB_API(hipCtxGetDevice, device); + HIP_INIT_API(hipCtxGetDevice, device); hipError_t e = hipSuccess; ihipCtx_t* ctx = ihipGetTlsDefaultCtx(); @@ -208,7 +208,7 @@ hipError_t hipCtxGetDevice(hipDevice_t* device) { } hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) { - HIP_INIT_CB_API(hipCtxGetApiVersion, apiVersion); + HIP_INIT_API(hipCtxGetApiVersion, apiVersion); if (apiVersion) { *apiVersion = 4; @@ -218,7 +218,7 @@ hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) { } hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) { - HIP_INIT_CB_API(hipCtxGetCacheConfig, cacheConfig); + HIP_INIT_API(hipCtxGetCacheConfig, cacheConfig); *cacheConfig = hipFuncCachePreferNone; @@ -226,7 +226,7 @@ hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) { } hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) { - HIP_INIT_CB_API(hipCtxSetCacheConfig, cacheConfig); + HIP_INIT_API(hipCtxSetCacheConfig, cacheConfig); // Nop, AMD does not support variable cache configs. @@ -234,7 +234,7 @@ hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) { } hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) { - HIP_INIT_CB_API(hipCtxSetSharedMemConfig, config); + HIP_INIT_API(hipCtxSetSharedMemConfig, config); // Nop, AMD does not support variable shared mem configs. @@ -242,7 +242,7 @@ hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) { } hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig) { - HIP_INIT_CB_API(hipCtxGetSharedMemConfig, pConfig); + HIP_INIT_API(hipCtxGetSharedMemConfig, pConfig); *pConfig = hipSharedMemBankSizeFourByte; @@ -250,12 +250,12 @@ hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig) { } hipError_t hipCtxSynchronize(void) { - HIP_INIT_CB_API(hipCtxSynchronize, 1); + HIP_INIT_API(hipCtxSynchronize, 1); return ihipLogStatus(ihipSynchronize()); // TODP Shall check validity of ctx? } hipError_t hipCtxGetFlags(unsigned int* flags) { - HIP_INIT_CB_API(hipCtxGetFlags, flags); + HIP_INIT_API(hipCtxGetFlags, flags); hipError_t e = hipSuccess; ihipCtx_t* tempCtx; tempCtx = ihipGetTlsDefaultCtx(); @@ -264,7 +264,7 @@ hipError_t hipCtxGetFlags(unsigned int* flags) { } hipError_t hipDevicePrimaryCtxGetState(hipDevice_t dev, unsigned int* flags, int* active) { - HIP_INIT_CB_API(hipDevicePrimaryCtxGetState, dev, flags, active); + HIP_INIT_API(hipDevicePrimaryCtxGetState, dev, flags, active); hipError_t e = hipSuccess; auto deviceHandle = ihipGetDevice(dev); @@ -286,7 +286,7 @@ hipError_t hipDevicePrimaryCtxGetState(hipDevice_t dev, unsigned int* flags, int } hipError_t hipDevicePrimaryCtxRelease(hipDevice_t dev) { - HIP_INIT_CB_API(hipDevicePrimaryCtxRelease, dev); + HIP_INIT_API(hipDevicePrimaryCtxRelease, dev); hipError_t e = hipSuccess; auto deviceHandle = ihipGetDevice(dev); @@ -297,7 +297,7 @@ hipError_t hipDevicePrimaryCtxRelease(hipDevice_t dev) { } hipError_t hipDevicePrimaryCtxRetain(hipCtx_t* pctx, hipDevice_t dev) { - HIP_INIT_CB_API(hipDevicePrimaryCtxRetain, pctx, dev); + HIP_INIT_API(hipDevicePrimaryCtxRetain, pctx, dev); hipError_t e = hipSuccess; auto deviceHandle = ihipGetDevice(dev); @@ -309,7 +309,7 @@ hipError_t hipDevicePrimaryCtxRetain(hipCtx_t* pctx, hipDevice_t dev) { } hipError_t hipDevicePrimaryCtxReset(hipDevice_t dev) { - HIP_INIT_CB_API(hipDevicePrimaryCtxReset, dev); + HIP_INIT_API(hipDevicePrimaryCtxReset, dev); hipError_t e = hipSuccess; auto deviceHandle = ihipGetDevice(dev); @@ -322,7 +322,7 @@ hipError_t hipDevicePrimaryCtxReset(hipDevice_t dev) { } hipError_t hipDevicePrimaryCtxSetFlags(hipDevice_t dev, unsigned int flags) { - HIP_INIT_CB_API(hipDevicePrimaryCtxSetFlags, dev, flags); + HIP_INIT_API(hipDevicePrimaryCtxSetFlags, dev, flags); hipError_t e = hipSuccess; auto deviceHandle = ihipGetDevice(dev); diff --git a/projects/clr/hipamd/src/hip_device.cpp b/projects/clr/hipamd/src/hip_device.cpp index 4d45f441e8..246ce8cf6f 100644 --- a/projects/clr/hipamd/src/hip_device.cpp +++ b/projects/clr/hipamd/src/hip_device.cpp @@ -30,7 +30,7 @@ THE SOFTWARE. //------------------------------------------------------------------------------------------------- // TODO - does this initialize HIP runtime? hipError_t hipGetDevice(int* deviceId) { - HIP_INIT_CB_API(hipGetDevice, deviceId); + HIP_INIT_API(hipGetDevice, deviceId); hipError_t e = hipSuccess; @@ -69,12 +69,12 @@ hipError_t ihipGetDeviceCount(int* count) { } hipError_t hipGetDeviceCount(int* count) { - HIP_INIT_CB_API(hipGetDeviceCount, count); + HIP_INIT_API(hipGetDeviceCount, count); return ihipLogStatus(ihipGetDeviceCount(count)); } hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig) { - HIP_INIT_CB_API(hipDeviceSetCacheConfig, cacheConfig); + HIP_INIT_API(hipDeviceSetCacheConfig, cacheConfig); // Nop, AMD does not support variable cache configs. @@ -82,7 +82,7 @@ hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig) { } hipError_t hipDeviceGetCacheConfig(hipFuncCache_t* cacheConfig) { - HIP_INIT_CB_API(hipDeviceGetCacheConfig, cacheConfig); + HIP_INIT_API(hipDeviceGetCacheConfig, cacheConfig); if (cacheConfig == nullptr) { return ihipLogStatus(hipErrorInvalidValue); @@ -94,7 +94,7 @@ hipError_t hipDeviceGetCacheConfig(hipFuncCache_t* cacheConfig) { } hipError_t hipDeviceGetLimit(size_t* pValue, hipLimit_t limit) { - HIP_INIT_CB_API(hipDeviceGetLimit, pValue, limit); + HIP_INIT_API(hipDeviceGetLimit, pValue, limit); if (pValue == nullptr) { return ihipLogStatus(hipErrorInvalidValue); } @@ -107,7 +107,7 @@ hipError_t hipDeviceGetLimit(size_t* pValue, hipLimit_t limit) { } hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t cacheConfig) { - HIP_INIT_CB_API(hipFuncSetCacheConfig, cacheConfig); + HIP_INIT_API(hipFuncSetCacheConfig, cacheConfig); // Nop, AMD does not support variable cache configs. @@ -115,7 +115,7 @@ hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t cacheConfig) { } hipError_t hipDeviceSetSharedMemConfig(hipSharedMemConfig config) { - HIP_INIT_CB_API(hipDeviceSetSharedMemConfig, config); + HIP_INIT_API(hipDeviceSetSharedMemConfig, config); // Nop, AMD does not support variable shared mem configs. @@ -123,7 +123,7 @@ hipError_t hipDeviceSetSharedMemConfig(hipSharedMemConfig config) { } hipError_t hipDeviceGetSharedMemConfig(hipSharedMemConfig* pConfig) { - HIP_INIT_CB_API(hipDeviceGetSharedMemConfig, pConfig); + HIP_INIT_API(hipDeviceGetSharedMemConfig, pConfig); *pConfig = hipSharedMemBankSizeFourByte; @@ -131,7 +131,7 @@ hipError_t hipDeviceGetSharedMemConfig(hipSharedMemConfig* pConfig) { } hipError_t hipSetDevice(int deviceId) { - HIP_INIT_CB_API(hipSetDevice, deviceId); + HIP_INIT_API(hipSetDevice, deviceId); if ((deviceId < 0) || (deviceId >= g_deviceCnt)) { return ihipLogStatus(hipErrorInvalidDevice); } else { @@ -142,12 +142,12 @@ hipError_t hipSetDevice(int deviceId) { } hipError_t hipDeviceSynchronize(void) { - HIP_INIT_SPECIAL_CB_API(hipDeviceSynchronize, TRACE_SYNC); + HIP_INIT_SPECIAL_API(hipDeviceSynchronize, TRACE_SYNC); return ihipLogStatus(ihipSynchronize()); } hipError_t hipDeviceReset(void) { - HIP_INIT_CB_API(hipDeviceReset, ); + HIP_INIT_API(hipDeviceReset, ); auto* ctx = ihipGetTlsDefaultCtx(); @@ -287,7 +287,7 @@ hipError_t ihipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device } hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) { - HIP_INIT_CB_API(hipDeviceGetAttribute, pi, attr, device); + HIP_INIT_API(hipDeviceGetAttribute, pi, attr, device); if ((device < 0) || (device >= g_deviceCnt)) { return ihipLogStatus(hipErrorInvalidDevice); } @@ -314,7 +314,7 @@ hipError_t ihipGetDeviceProperties(hipDeviceProp_t* props, int device) { } hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, int device) { - HIP_INIT_CB_API(hipGetDeviceProperties, props, device); + HIP_INIT_API(hipGetDeviceProperties, props, device); if ((device < 0) || (device >= g_deviceCnt)) { return ihipLogStatus(hipErrorInvalidDevice); } @@ -322,7 +322,7 @@ hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, int device) { } hipError_t hipSetDeviceFlags(unsigned int flags) { - HIP_INIT_CB_API(hipSetDeviceFlags, flags); + HIP_INIT_API(hipSetDeviceFlags, flags); hipError_t e = hipSuccess; @@ -367,7 +367,7 @@ hipError_t hipSetDeviceFlags(unsigned int flags) { }; hipError_t hipDeviceComputeCapability(int* major, int* minor, hipDevice_t device) { - HIP_INIT_CB_API(hipDeviceComputeCapability, major, minor, device); + HIP_INIT_API(hipDeviceComputeCapability, major, minor, device); hipError_t e = hipSuccess; if ((device < 0) || (device >= g_deviceCnt)) { e = hipErrorInvalidDevice; @@ -380,7 +380,7 @@ hipError_t hipDeviceComputeCapability(int* major, int* minor, hipDevice_t device hipError_t hipDeviceGetName(char* name, int len, hipDevice_t device) { // Cast to void* here to avoid printing garbage in debug modes. - HIP_INIT_CB_API(hipDeviceGetName, (void*)name, len, device); + HIP_INIT_API(hipDeviceGetName, (void*)name, len, device); hipError_t e = hipSuccess; if ((device < 0) || (device >= g_deviceCnt)) { e = hipErrorInvalidDevice; @@ -394,7 +394,7 @@ hipError_t hipDeviceGetName(char* name, int len, hipDevice_t device) { hipError_t hipDeviceGetPCIBusId(char* pciBusId, int len, int device) { // Cast to void* here to avoid printing garbage in debug modes. - HIP_INIT_CB_API(hipDeviceGetPCIBusId, (void*)pciBusId, len, device); + HIP_INIT_API(hipDeviceGetPCIBusId, (void*)pciBusId, len, device); hipError_t e = hipErrorInvalidValue; if ((device < 0) || (device >= g_deviceCnt)) { e = hipErrorInvalidDevice; @@ -413,7 +413,7 @@ hipError_t hipDeviceGetPCIBusId(char* pciBusId, int len, int device) { } hipError_t hipDeviceTotalMem(size_t* bytes, hipDevice_t device) { - HIP_INIT_CB_API(hipDeviceTotalMem, bytes, device); + HIP_INIT_API(hipDeviceTotalMem, bytes, device); hipError_t e = hipSuccess; if ((device < 0) || (device >= g_deviceCnt)) { e = hipErrorInvalidDevice; @@ -425,7 +425,7 @@ hipError_t hipDeviceTotalMem(size_t* bytes, hipDevice_t device) { } hipError_t hipDeviceGetByPCIBusId(int* device, const char* pciBusId) { - HIP_INIT_CB_API(hipDeviceGetByPCIBusId, device, pciBusId); + HIP_INIT_API(hipDeviceGetByPCIBusId, device, pciBusId); hipDeviceProp_t tempProp; int deviceCount = 0; hipError_t e = hipErrorInvalidValue; @@ -451,7 +451,7 @@ hipError_t hipDeviceGetByPCIBusId(int* device, const char* pciBusId) { } hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* prop) { - HIP_INIT_CB_API(hipChooseDevice, device, prop); + HIP_INIT_API(hipChooseDevice, device, prop); hipDeviceProp_t tempProp; hipError_t e = hipSuccess; if ((device == NULL) || (prop == NULL)) { diff --git a/projects/clr/hipamd/src/hip_error.cpp b/projects/clr/hipamd/src/hip_error.cpp index ec1e2fbb02..6f1184d92f 100644 --- a/projects/clr/hipamd/src/hip_error.cpp +++ b/projects/clr/hipamd/src/hip_error.cpp @@ -30,7 +30,7 @@ THE SOFTWARE. //--- hipError_t hipGetLastError() { - HIP_INIT_CB_API(hipGetLastError); + HIP_INIT_API(hipGetLastError); // Return last error, but then reset the state: hipError_t e = ihipLogStatus(tls_lastHipError); @@ -39,20 +39,20 @@ hipError_t hipGetLastError() { } hipError_t hipPeekAtLastError() { - HIP_INIT_CB_API(hipPeekAtLastError); + HIP_INIT_API(hipPeekAtLastError); // peek at last error, but don't reset it. return ihipLogStatus(tls_lastHipError); } const char* hipGetErrorName(hipError_t hip_error) { - HIP_INIT_API(hip_error); + HIP_INIT_API(hipGetErrorName, hip_error); return ihipErrorString(hip_error); } const char* hipGetErrorString(hipError_t hip_error) { - HIP_INIT_API(hip_error); + HIP_INIT_API(hipGetErrorString, hip_error); // TODO - return a message explaining the error. // TODO - This should be set up to return the same string reported in the the doxygen comments, diff --git a/projects/clr/hipamd/src/hip_event.cpp b/projects/clr/hipamd/src/hip_event.cpp index 206ad16a79..ea014ab292 100644 --- a/projects/clr/hipamd/src/hip_event.cpp +++ b/projects/clr/hipamd/src/hip_event.cpp @@ -95,20 +95,20 @@ hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags) { } hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) { - HIP_INIT_CB_API(hipEventCreateWithFlags, event, flags); + HIP_INIT_API(hipEventCreateWithFlags, event, flags); return ihipLogStatus(ihipEventCreate(event, flags)); } hipError_t hipEventCreate(hipEvent_t* event) { - HIP_INIT_CB_API(hipEventCreate, event); + HIP_INIT_API(hipEventCreate, event); return ihipLogStatus(ihipEventCreate(event, 0)); } hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { - HIP_INIT_SPECIAL_CB_API(hipEventRecord, TRACE_SYNC, event, stream); + HIP_INIT_SPECIAL_API(hipEventRecord, TRACE_SYNC, event, stream); auto ecd = event->locked_copyCrit(); @@ -153,7 +153,7 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { hipError_t hipEventDestroy(hipEvent_t event) { - HIP_INIT_CB_API(hipEventDestroy, event); + HIP_INIT_API(hipEventDestroy, event); if (event) { delete event; @@ -165,7 +165,7 @@ hipError_t hipEventDestroy(hipEvent_t event) { } hipError_t hipEventSynchronize(hipEvent_t event) { - HIP_INIT_SPECIAL_CB_API(hipEventSynchronize, TRACE_SYNC, event); + HIP_INIT_SPECIAL_API(hipEventSynchronize, TRACE_SYNC, event); if (!(event->_flags & hipEventReleaseToSystem)) { tprintf(DB_WARN, @@ -198,7 +198,7 @@ hipError_t hipEventSynchronize(hipEvent_t event) { } hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) { - HIP_INIT_CB_API(hipEventElapsedTime, ms, start, stop); + HIP_INIT_API(hipEventElapsedTime, ms, start, stop); hipError_t status = hipSuccess; @@ -255,7 +255,7 @@ hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) { } hipError_t hipEventQuery(hipEvent_t event) { - HIP_INIT_SPECIAL_CB_API(hipEventQuery, TRACE_QUERY, event); + HIP_INIT_SPECIAL_API(hipEventQuery, TRACE_QUERY, event); if (!(event->_flags & hipEventReleaseToSystem)) { tprintf(DB_WARN, diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index ab25a04795..f99d171efa 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -2288,7 +2288,7 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes //------------------------------------------------------------------------------------------------- // Profiler, really these should live elsewhere: hipError_t hipProfilerStart() { - HIP_INIT_CB_API(hipProfilerStart); + HIP_INIT_API(hipProfilerStart); #if COMPILE_HIP_ATP_MARKER amdtResumeProfiling(AMDT_ALL_PROFILING); #endif @@ -2298,7 +2298,7 @@ hipError_t hipProfilerStart() { hipError_t hipProfilerStop() { - HIP_INIT_CB_API(hipProfilerStop); + HIP_INIT_API(hipProfilerStop); #if COMPILE_HIP_ATP_MARKER amdtStopProfiling(AMDT_ALL_PROFILING); #endif @@ -2313,7 +2313,7 @@ hipError_t hipProfilerStop() { //--- hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc) { - HIP_INIT_API(deviceId, acc); + HIP_INIT_API(hipHccGetAccelerator, deviceId, acc); const ihipDevice_t* device = ihipGetDevice(deviceId); hipError_t err; @@ -2329,7 +2329,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(hipHccGetAcceleratorView, stream, av); if (stream == hipStreamNull) { ihipCtx_t* device = ihipGetTlsDefaultCtx(); diff --git a/projects/clr/hipamd/src/hip_hcc_internal.h b/projects/clr/hipamd/src/hip_hcc_internal.h index b9667776d3..bc407c3df5 100644 --- a/projects/clr/hipamd/src/hip_hcc_internal.h +++ b/projects/clr/hipamd/src/hip_hcc_internal.h @@ -297,11 +297,7 @@ extern uint64_t recordApiTrace(std::string* fullStr, const std::string& apiStr); // This macro should be called at the beginning of every HIP API. // It initializes the hip runtime (exactly once), and // generates a trace string that can be output to stderr or to ATP file. -#define HIP_INIT_API(...) \ - HIP_INIT() \ - API_TRACE(0, __VA_ARGS__); - -#define HIP_INIT_CB_API(cid, ...) \ +#define HIP_INIT_API(cid, ...) \ HIP_INIT() \ API_TRACE(0, __VA_ARGS__); \ HIP_CB_SPAWNER_OBJECT(cid); @@ -310,11 +306,7 @@ extern uint64_t recordApiTrace(std::string* fullStr, const std::string& apiStr); // Like above, but will trace with a specified "special" bit. // Replace HIP_INIT_API with this call inside HIP APIs that launch work on the GPU: // kernel launches, copy commands, memory sets, etc. -#define HIP_INIT_SPECIAL_API(tbit, ...) \ - HIP_INIT() \ - API_TRACE((HIP_TRACE_API & (1 << tbit)), __VA_ARGS__); - -#define HIP_INIT_SPECIAL_CB_API(cid, tbit, ...) \ +#define HIP_INIT_SPECIAL_API(cid, tbit, ...) \ HIP_INIT() \ API_TRACE((HIP_TRACE_API & (1 << tbit)), __VA_ARGS__); \ HIP_CB_SPAWNER_OBJECT(cid); diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index b31192f5c5..b389b4f93a 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -159,7 +159,7 @@ void* allocAndSharePtr(const char* msg, size_t sizeBytes, ihipCtx_t* ctx, bool s // TODO - add more info here when available. // hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void* ptr) { - HIP_INIT_CB_API(hipPointerGetAttributes, attributes, ptr); + HIP_INIT_API(hipPointerGetAttributes, attributes, ptr); hipError_t e = hipSuccess; if ((attributes == nullptr) || (ptr == nullptr)) { @@ -206,7 +206,7 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void hipError_t hipHostGetDevicePointer(void** devicePointer, void* hostPointer, unsigned flags) { - HIP_INIT_CB_API(hipHostGetDevicePointer, devicePointer, hostPointer, flags); + HIP_INIT_API(hipHostGetDevicePointer, devicePointer, hostPointer, flags); hipError_t e = hipSuccess; @@ -237,7 +237,7 @@ hipError_t hipHostGetDevicePointer(void** devicePointer, void* hostPointer, unsi hipError_t hipMalloc(void** ptr, size_t sizeBytes) { - HIP_INIT_SPECIAL_CB_API(hipMalloc, (TRACE_MEM), ptr, sizeBytes); + HIP_INIT_SPECIAL_API(hipMalloc, (TRACE_MEM), ptr, sizeBytes); HIP_SET_DEVICE(); hipError_t hip_status = hipSuccess; @@ -266,7 +266,7 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) { hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { - HIP_INIT_SPECIAL_CB_API(hipHostMalloc, (TRACE_MEM), ptr, sizeBytes, flags); + HIP_INIT_SPECIAL_API(hipHostMalloc, (TRACE_MEM), ptr, sizeBytes, flags); HIP_SET_DEVICE(); hipError_t hip_status = hipSuccess; @@ -398,7 +398,7 @@ hipError_t ihipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t heigh // width in bytes hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height) { - HIP_INIT_SPECIAL_CB_API(hipMallocPitch, (TRACE_MEM), ptr, pitch, width, height); + HIP_INIT_SPECIAL_API(hipMallocPitch, (TRACE_MEM), ptr, pitch, width, height); HIP_SET_DEVICE(); hipError_t hip_status = hipSuccess; @@ -409,7 +409,7 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height } hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) { - HIP_INIT_CB_API(hipMalloc3D, pitchedDevPtr, &extent); + HIP_INIT_API(hipMalloc3D, pitchedDevPtr, &extent); HIP_SET_DEVICE(); hipError_t hip_status = hipSuccess; @@ -444,7 +444,7 @@ extern void getChannelOrderAndType(const hipChannelFormatDesc& desc, hsa_ext_image_channel_type_t* channelType); hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) { - HIP_INIT_SPECIAL_CB_API(hipArrayCreate, (TRACE_MEM), array, pAllocateArray); + HIP_INIT_SPECIAL_API(hipArrayCreate, (TRACE_MEM), array, pAllocateArray); HIP_SET_DEVICE(); hipError_t hip_status = hipSuccess; if (pAllocateArray->width > 0) { @@ -554,7 +554,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_SPECIAL_CB_API(hipMallocArray, (TRACE_MEM), array, desc, width, height, flags); + HIP_INIT_SPECIAL_API(hipMallocArray, (TRACE_MEM), array, desc, width, height, flags); HIP_SET_DEVICE(); hipError_t hip_status = hipSuccess; if (width > 0) { @@ -635,7 +635,7 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, si } hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) { - HIP_INIT_SPECIAL_CB_API(hipArray3DCreate, (TRACE_MEM), array, pAllocateArray); + HIP_INIT_SPECIAL_API(hipArray3DCreate, (TRACE_MEM), array, pAllocateArray); hipError_t hip_status = hipSuccess; auto ctx = ihipGetTlsDefaultCtx(); @@ -762,7 +762,7 @@ hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc* - HIP_INIT_CB_API(hipMalloc3DArray, array, desc, &extent, flags); + HIP_INIT_API(hipMalloc3DArray, array, desc, &extent, flags); HIP_SET_DEVICE(); hipError_t hip_status = hipSuccess; @@ -846,7 +846,7 @@ hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc* } hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) { - HIP_INIT_CB_API(hipHostGetFlags, flagsPtr, hostPtr); + HIP_INIT_API(hipHostGetFlags, flagsPtr, hostPtr); hipError_t hip_status = hipSuccess; @@ -874,7 +874,7 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) { // TODO - need to fix several issues here related to P2P access, host memory fallback. hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) { - HIP_INIT_CB_API(hipHostRegister, hostPtr, sizeBytes, flags); + HIP_INIT_API(hipHostRegister, hostPtr, sizeBytes, flags); hipError_t hip_status = hipSuccess; @@ -931,7 +931,7 @@ hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) } hipError_t hipHostUnregister(void* hostPtr) { - HIP_INIT_CB_API(hipHostUnregister, hostPtr); + HIP_INIT_API(hipHostUnregister, hostPtr); auto ctx = ihipGetTlsDefaultCtx(); hipError_t hip_status = hipSuccess; if (hostPtr == NULL) { @@ -966,7 +966,7 @@ inline hipDeviceptr_t agent_address_for_symbol(const char* symbolName) { hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t count, size_t offset, hipMemcpyKind kind) { - HIP_INIT_SPECIAL_CB_API(hipMemcpyToSymbol, (TRACE_MCMD), symbolName, src, count, offset, kind); + HIP_INIT_SPECIAL_API(hipMemcpyToSymbol, (TRACE_MCMD), symbolName, src, count, offset, kind); if (symbolName == nullptr) { return ihipLogStatus(hipErrorInvalidSymbol); @@ -999,7 +999,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_SPECIAL_CB_API(hipMemcpyFromSymbol, (TRACE_MCMD), symbolName, dst, count, offset, kind); + HIP_INIT_SPECIAL_API(hipMemcpyFromSymbol, (TRACE_MCMD), symbolName, dst, count, offset, kind); if (symbolName == nullptr) { return ihipLogStatus(hipErrorInvalidSymbol); @@ -1031,7 +1031,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_SPECIAL_CB_API(hipMemcpyToSymbolAsync, (TRACE_MCMD), symbolName, src, count, offset, kind, stream); + HIP_INIT_SPECIAL_API(hipMemcpyToSymbolAsync, (TRACE_MCMD), symbolName, src, count, offset, kind, stream); if (symbolName == nullptr) { return ihipLogStatus(hipErrorInvalidSymbol); @@ -1066,7 +1066,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_SPECIAL_CB_API(hipMemcpyFromSymbolAsync, (TRACE_MCMD), symbolName, dst, count, offset, kind, stream); + HIP_INIT_SPECIAL_API(hipMemcpyFromSymbolAsync, (TRACE_MCMD), symbolName, dst, count, offset, kind, stream); if (symbolName == nullptr) { return ihipLogStatus(hipErrorInvalidSymbol); @@ -1101,7 +1101,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t co //--- hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) { - HIP_INIT_SPECIAL_CB_API(hipMemcpy, (TRACE_MCMD), dst, src, sizeBytes, kind); + HIP_INIT_SPECIAL_API(hipMemcpy, (TRACE_MCMD), dst, src, sizeBytes, kind); hipError_t e = hipSuccess; @@ -1128,7 +1128,7 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) { - HIP_INIT_SPECIAL_CB_API(hipMemcpyHtoD, (TRACE_MCMD), dst, src, sizeBytes); + HIP_INIT_SPECIAL_API(hipMemcpyHtoD, (TRACE_MCMD), dst, src, sizeBytes); hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); @@ -1147,7 +1147,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_SPECIAL_CB_API(hipMemcpyDtoH, (TRACE_MCMD), dst, src, sizeBytes); + HIP_INIT_SPECIAL_API(hipMemcpyDtoH, (TRACE_MCMD), dst, src, sizeBytes); hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); @@ -1166,7 +1166,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_SPECIAL_CB_API(hipMemcpyDtoD, (TRACE_MCMD), dst, src, sizeBytes); + HIP_INIT_SPECIAL_API(hipMemcpyDtoD, (TRACE_MCMD), dst, src, sizeBytes); hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); @@ -1184,7 +1184,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_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes); + HIP_INIT_SPECIAL_API(hipMemcpyHtoH, (TRACE_MCMD), dst, src, sizeBytes); hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); @@ -1203,13 +1203,13 @@ 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_SPECIAL_CB_API(hipMemcpyAsync, (TRACE_MCMD), dst, src, sizeBytes, kind, stream); + HIP_INIT_SPECIAL_API(hipMemcpyAsync, (TRACE_MCMD), dst, src, sizeBytes, kind, stream); return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, kind, stream)); } hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, hipStream_t stream) { - HIP_INIT_SPECIAL_CB_API(hipMemcpyHtoDAsync, (TRACE_MCMD), dst, src, sizeBytes, stream); + HIP_INIT_SPECIAL_API(hipMemcpyHtoDAsync, (TRACE_MCMD), dst, src, sizeBytes, stream); return ihipLogStatus( hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyHostToDevice, stream)); @@ -1217,14 +1217,14 @@ hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, h hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream) { - HIP_INIT_SPECIAL_CB_API(hipMemcpyDtoDAsync, (TRACE_MCMD), dst, src, sizeBytes, stream); + HIP_INIT_SPECIAL_API(hipMemcpyDtoDAsync, (TRACE_MCMD), dst, src, sizeBytes, stream); return ihipLogStatus( hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDeviceToDevice, stream)); } hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream) { - HIP_INIT_SPECIAL_CB_API(hipMemcpyDtoHAsync, (TRACE_MCMD), dst, src, sizeBytes, stream); + HIP_INIT_SPECIAL_API(hipMemcpyDtoHAsync, (TRACE_MCMD), dst, src, sizeBytes, stream); return ihipLogStatus( hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDeviceToHost, stream)); @@ -1232,7 +1232,7 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, h 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_SPECIAL_CB_API(hipMemcpy2DToArray, (TRACE_MCMD), dst, wOffset, hOffset, src, spitch, width, height, kind); + HIP_INIT_SPECIAL_API(hipMemcpy2DToArray, (TRACE_MCMD), dst, wOffset, hOffset, src, spitch, width, height, kind); hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); @@ -1284,7 +1284,7 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t count, hipMemcpyKind kind) { - HIP_INIT_SPECIAL_CB_API(hipMemcpyToArray, (TRACE_MCMD), dst, wOffset, hOffset, src, count, kind); + HIP_INIT_SPECIAL_API(hipMemcpyToArray, (TRACE_MCMD), dst, wOffset, hOffset, src, count, kind); hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); @@ -1303,7 +1303,7 @@ hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffset, size_t hOffset, size_t count, hipMemcpyKind kind) { - HIP_INIT_SPECIAL_CB_API(hipMemcpyFromArray, (TRACE_MCMD), dst, srcArray, wOffset, hOffset, count, kind); + HIP_INIT_SPECIAL_API(hipMemcpyFromArray, (TRACE_MCMD), dst, srcArray, wOffset, hOffset, count, kind); hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); @@ -1321,7 +1321,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_SPECIAL_CB_API(hipMemcpyHtoA, (TRACE_MCMD), dstArray, dstOffset, srcHost, count); + HIP_INIT_SPECIAL_API(hipMemcpyHtoA, (TRACE_MCMD), dstArray, dstOffset, srcHost, count); hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); @@ -1339,7 +1339,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_SPECIAL_CB_API(hipMemcpyAtoH, (TRACE_MCMD), dst, srcArray, srcOffset, count); + HIP_INIT_SPECIAL_API(hipMemcpyAtoH, (TRACE_MCMD), dst, srcArray, srcOffset, count); hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); @@ -1358,7 +1358,7 @@ hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t } hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p) { - HIP_INIT_SPECIAL_CB_API(hipMemcpy3D, (TRACE_MCMD), p); + HIP_INIT_SPECIAL_API(hipMemcpy3D, (TRACE_MCMD), p); hipError_t e = hipSuccess; if (p) { size_t byteSize; @@ -1626,7 +1626,7 @@ hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { - HIP_INIT_SPECIAL_CB_API(hipMemcpy2D, (TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind); + HIP_INIT_SPECIAL_API(hipMemcpy2D, (TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind); hipError_t e = hipSuccess; e = ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind); return ihipLogStatus(e); @@ -1634,7 +1634,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_SPECIAL_CB_API(hipMemcpy2DAsync, (TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind, stream); + HIP_INIT_SPECIAL_API(hipMemcpy2DAsync, (TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind, stream); if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return ihipLogStatus(hipErrorInvalidValue); hipError_t e = hipSuccess; int isLocked = 0; @@ -1673,7 +1673,7 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp } hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) { - HIP_INIT_SPECIAL_CB_API(hipMemcpyParam2D, (TRACE_MCMD), pCopy); + HIP_INIT_SPECIAL_API(hipMemcpyParam2D, (TRACE_MCMD), pCopy); hipError_t e = hipSuccess; if (pCopy == nullptr) { e = hipErrorInvalidValue; @@ -1685,7 +1685,7 @@ hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) { // TODO-sync: function is async unless target is pinned host memory - then these are fully sync. hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) { - HIP_INIT_SPECIAL_CB_API(hipMemsetAsync, (TRACE_MCMD), dst, value, sizeBytes, stream); + HIP_INIT_SPECIAL_API(hipMemsetAsync, (TRACE_MCMD), dst, value, sizeBytes, stream); hipError_t e = hipSuccess; @@ -1697,7 +1697,7 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st }; hipError_t hipMemset(void* dst, int value, size_t sizeBytes) { - HIP_INIT_SPECIAL_CB_API(hipMemset, (TRACE_MCMD), dst, value, sizeBytes); + HIP_INIT_SPECIAL_API(hipMemset, (TRACE_MCMD), dst, value, sizeBytes); hipError_t e = hipSuccess; @@ -1713,7 +1713,7 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes) { } hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height) { - HIP_INIT_SPECIAL_CB_API(hipMemset2D, (TRACE_MCMD), dst, pitch, value, width, height); + HIP_INIT_SPECIAL_API(hipMemset2D, (TRACE_MCMD), dst, pitch, value, width, height); hipError_t e = hipSuccess; @@ -1732,7 +1732,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_SPECIAL_CB_API(hipMemset2DAsync, (TRACE_MCMD), dst, pitch, value, width, height, stream); + HIP_INIT_SPECIAL_API(hipMemset2DAsync, (TRACE_MCMD), dst, pitch, value, width, height, stream); hipError_t e = hipSuccess; @@ -1749,7 +1749,7 @@ hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, si }; hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes) { - HIP_INIT_SPECIAL_CB_API(hipMemsetD8, (TRACE_MCMD), dst, value, sizeBytes); + HIP_INIT_SPECIAL_API(hipMemsetD8, (TRACE_MCMD), dst, value, sizeBytes); hipError_t e = hipSuccess; @@ -1766,7 +1766,7 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ) { - HIP_INIT_SPECIAL_CB_API(hipMemset3D, (TRACE_MCMD), &pitchedDevPtr, value, &extent); + HIP_INIT_SPECIAL_API(hipMemset3D, (TRACE_MCMD), &pitchedDevPtr, value, &extent); hipError_t e = hipSuccess; hipStream_t stream = hipStreamNull; @@ -1785,7 +1785,7 @@ hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ,hipStream_t stream ) { - HIP_INIT_SPECIAL_CB_API(hipMemset3DAsync, (TRACE_MCMD), &pitchedDevPtr, value, &extent); + HIP_INIT_SPECIAL_API(hipMemset3DAsync, (TRACE_MCMD), &pitchedDevPtr, value, &extent); hipError_t e = hipSuccess; // TODO - call an ihip memset so HIP_TRACE is correct. @@ -1801,7 +1801,7 @@ hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent e } hipError_t hipMemGetInfo(size_t* free, size_t* total) { - HIP_INIT_CB_API(hipMemGetInfo, free, total); + HIP_INIT_API(hipMemGetInfo, free, total); hipError_t e = hipSuccess; @@ -1835,7 +1835,7 @@ hipError_t hipMemGetInfo(size_t* free, size_t* total) { } hipError_t hipMemPtrGetInfo(void* ptr, size_t* size) { - HIP_INIT_CB_API(hipMemPtrGetInfo, ptr, size); + HIP_INIT_API(hipMemPtrGetInfo, ptr, size); hipError_t e = hipSuccess; @@ -1860,7 +1860,7 @@ hipError_t hipMemPtrGetInfo(void* ptr, size_t* size) { hipError_t hipFree(void* ptr) { - HIP_INIT_SPECIAL_CB_API(hipFree, (TRACE_MEM), ptr); + HIP_INIT_SPECIAL_API(hipFree, (TRACE_MEM), ptr); hipError_t hipStatus = hipErrorInvalidDevicePointer; @@ -1892,7 +1892,7 @@ hipError_t hipFree(void* ptr) { hipError_t hipHostFree(void* ptr) { - HIP_INIT_SPECIAL_CB_API(hipHostFree, (TRACE_MEM), ptr); + HIP_INIT_SPECIAL_API(hipHostFree, (TRACE_MEM), ptr); // Synchronize to ensure all work has finished. ihipGetTlsDefaultCtx()->locked_waitAllStreams(); // ignores non-blocking streams, this waits @@ -1927,7 +1927,7 @@ hipError_t hipHostFree(void* ptr) { hipError_t hipFreeHost(void* ptr) { return hipHostFree(ptr); } hipError_t hipFreeArray(hipArray* array) { - HIP_INIT_SPECIAL_CB_API(hipFreeArray, (TRACE_MEM), array); + HIP_INIT_SPECIAL_API(hipFreeArray, (TRACE_MEM), array); hipError_t hipStatus = hipErrorInvalidDevicePointer; @@ -1955,7 +1955,7 @@ hipError_t hipFreeArray(hipArray* array) { } hipError_t hipMemGetAddressRange(hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr) { - HIP_INIT_CB_API(hipMemGetAddressRange, pbase, psize, dptr); + HIP_INIT_API(hipMemGetAddressRange, pbase, psize, dptr); hipError_t hipStatus = hipSuccess; hc::accelerator acc; #if (__hcc_workweek__ >= 17332) @@ -1976,7 +1976,7 @@ hipError_t hipMemGetAddressRange(hipDeviceptr_t* pbase, size_t* psize, hipDevice // TODO: IPC implementaiton: hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr) { - HIP_INIT_CB_API(hipIpcGetMemHandle, handle, devPtr); + HIP_INIT_API(hipIpcGetMemHandle, handle, devPtr); hipError_t hipStatus = hipSuccess; // Get the size of allocated pointer size_t psize = 0u; @@ -2012,7 +2012,7 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr) { } hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags) { - HIP_INIT_CB_API(hipIpcOpenMemHandle, devPtr, &handle, flags); + HIP_INIT_API(hipIpcOpenMemHandle, devPtr, &handle, flags); hipError_t hipStatus = hipSuccess; if (devPtr == NULL) { hipStatus = hipErrorInvalidValue; @@ -2042,7 +2042,7 @@ hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned } hipError_t hipIpcCloseMemHandle(void* devPtr) { - HIP_INIT_CB_API(hipIpcCloseMemHandle, devPtr); + HIP_INIT_API(hipIpcCloseMemHandle, devPtr); hipError_t hipStatus = hipSuccess; if (devPtr == NULL) { hipStatus = hipErrorInvalidValue; diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index 539fbed9af..db89f5f3f8 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/src/hip_module.cpp @@ -97,7 +97,7 @@ std::string& FunctionSymbol(hipFunction_t f) { return f->_name; }; } hipError_t hipModuleUnload(hipModule_t hmod) { - HIP_INIT_CB_API(hipModuleUnload, hmod); + HIP_INIT_API(hipModuleUnload, hmod); // TODO - improve this synchronization so it is thread-safe. // Currently we want for all inflight activity to complete, but don't prevent another @@ -231,7 +231,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, uint32_t gridDimX, uint32_t gr uint32_t gridDimZ, uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ, uint32_t sharedMemBytes, hipStream_t hStream, void** kernelParams, void** extra) { - HIP_INIT_CB_API(hipModuleLaunchKernel, f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, + HIP_INIT_API(hipModuleLaunchKernel, f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra); return ihipLogStatus(ihipModuleLaunchKernel( f, blockDimX * gridDimX, blockDimY * gridDimY, gridDimZ * blockDimZ, blockDimX, blockDimY, @@ -245,7 +245,7 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, uint32_t localWorkSizeZ, size_t sharedMemBytes, hipStream_t hStream, void** kernelParams, void** extra, hipEvent_t startEvent, hipEvent_t stopEvent) { - HIP_INIT_API(f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, + HIP_INIT_API(hipHccModuleLaunchKernel, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra); return ihipLogStatus(ihipModuleLaunchKernel( f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY, @@ -461,13 +461,13 @@ hipError_t ihipModuleGetFunction(hipFunction_t* func, hipModule_t hmod, const ch } hipError_t hipModuleGetFunction(hipFunction_t* hfunc, hipModule_t hmod, const char* name) { - HIP_INIT_CB_API(hipModuleGetFunction, hfunc, hmod, name); + HIP_INIT_API(hipModuleGetFunction, hfunc, hmod, name); return ihipLogStatus(ihipModuleGetFunction(hfunc, hmod, name)); } hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod, const char* name) { - HIP_INIT_CB_API(hipModuleGetGlobal, dptr, bytes, hmod, name); + HIP_INIT_API(hipModuleGetGlobal, dptr, bytes, hmod, name); if (!dptr || !bytes) return ihipLogStatus(hipErrorInvalidValue); @@ -559,12 +559,12 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) { } hipError_t hipModuleLoadData(hipModule_t* module, const void* image) { - HIP_INIT_CB_API(hipModuleLoadData, module, image); + HIP_INIT_API(hipModuleLoadData, module, image); return ihipLogStatus(ihipModuleLoadData(module,image)); } hipError_t hipModuleLoad(hipModule_t* module, const char* fname) { - HIP_INIT_CB_API(hipModuleLoad, module, fname); + HIP_INIT_API(hipModuleLoad, module, fname); if (!fname) return ihipLogStatus(hipErrorInvalidValue); @@ -579,12 +579,12 @@ hipError_t hipModuleLoad(hipModule_t* module, const char* fname) { hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* image, unsigned int numOptions, hipJitOption* options, void** optionValues) { - HIP_INIT_CB_API(hipModuleLoadDataEx, module, image, numOptions, options, optionValues); + HIP_INIT_API(hipModuleLoadDataEx, module, image, numOptions, options, optionValues); return ihipLogStatus(ihipModuleLoadData(module, image)); } hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name) { - HIP_INIT_CB_API(hipModuleGetTexRef, texRef, hmod, name); + HIP_INIT_API(hipModuleGetTexRef, texRef, hmod, name); hipError_t ret = hipErrorNotFound; if (!texRef) return ihipLogStatus(hipErrorInvalidValue); diff --git a/projects/clr/hipamd/src/hip_peer.cpp b/projects/clr/hipamd/src/hip_peer.cpp index 97c3fd4052..cffb895c57 100644 --- a/projects/clr/hipamd/src/hip_peer.cpp +++ b/projects/clr/hipamd/src/hip_peer.cpp @@ -73,7 +73,7 @@ hipError_t ihipDeviceCanAccessPeer(int* canAccessPeer, hipCtx_t thisCtx, hipCtx_ */ //--- hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t peerCtx) { - HIP_INIT_API(canAccessPeer, thisCtx, peerCtx); + HIP_INIT_API(hipDeviceCanAccessPeer2, canAccessPeer, thisCtx, peerCtx); return ihipLogStatus(ihipDeviceCanAccessPeer(canAccessPeer, thisCtx, peerCtx)); } @@ -150,7 +150,7 @@ hipError_t ihipEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) { //--- 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(hipMemcpyPeer2, dst, dstCtx, src, srcCtx, sizeBytes); // TODO - move to ihip memory copy implementaion. // HCC has a unified memory architecture so device specifiers are not required. @@ -161,7 +161,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(hipMemcpyPeerAsync2, dst, dstDevice, src, srcDevice, sizeBytes, stream); // TODO - move to ihip memory copy implementaion. // HCC has a unified memory architecture so device specifiers are not required. @@ -175,21 +175,21 @@ hipError_t hipMemcpyPeerAsync(void* dst, hipCtx_t dstDevice, const void* src, hi //============================================================================= hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, int deviceId, int peerDeviceId) { - HIP_INIT_CB_API(hipDeviceCanAccessPeer, canAccessPeer, deviceId, peerDeviceId); + HIP_INIT_API(hipDeviceCanAccessPeer, canAccessPeer, deviceId, peerDeviceId); return ihipLogStatus(ihipDeviceCanAccessPeer(canAccessPeer, ihipGetPrimaryCtx(deviceId), ihipGetPrimaryCtx(peerDeviceId))); } hipError_t hipDeviceDisablePeerAccess(int peerDeviceId) { - HIP_INIT_CB_API(hipDeviceDisablePeerAccess, peerDeviceId); + HIP_INIT_API(hipDeviceDisablePeerAccess, peerDeviceId); return ihipLogStatus(ihipDisablePeerAccess(ihipGetPrimaryCtx(peerDeviceId))); } hipError_t hipDeviceEnablePeerAccess(int peerDeviceId, unsigned int flags) { - HIP_INIT_CB_API(hipDeviceEnablePeerAccess, peerDeviceId, flags); + HIP_INIT_API(hipDeviceEnablePeerAccess, peerDeviceId, flags); return ihipLogStatus(ihipEnablePeerAccess(ihipGetPrimaryCtx(peerDeviceId), flags)); } @@ -197,7 +197,7 @@ hipError_t hipDeviceEnablePeerAccess(int peerDeviceId, unsigned int flags) { hipError_t hipMemcpyPeer(void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes) { - HIP_INIT_CB_API(hipMemcpyPeer, dst, dstDevice, src, srcDevice, sizeBytes); + HIP_INIT_API(hipMemcpyPeer, dst, dstDevice, src, srcDevice, sizeBytes); return ihipLogStatus(hipMemcpyPeer(dst, ihipGetPrimaryCtx(dstDevice), src, ihipGetPrimaryCtx(srcDevice), sizeBytes)); } @@ -205,18 +205,18 @@ hipError_t hipMemcpyPeer(void* dst, int dstDevice, const void* src, int srcDevic hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes, hipStream_t stream) { - HIP_INIT_CB_API(hipMemcpyPeerAsync, dst, dstDevice, src, srcDevice, sizeBytes, stream); + HIP_INIT_API(hipMemcpyPeerAsync, dst, dstDevice, src, srcDevice, sizeBytes, stream); return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDefault, stream)); } hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) { - HIP_INIT_CB_API(hipCtxEnablePeerAccess, peerCtx, flags); + HIP_INIT_API(hipCtxEnablePeerAccess, peerCtx, flags); return ihipLogStatus(ihipEnablePeerAccess(peerCtx, flags)); } hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx) { - HIP_INIT_CB_API(hipCtxDisablePeerAccess, peerCtx); + HIP_INIT_API(hipCtxDisablePeerAccess, peerCtx); return ihipLogStatus(ihipDisablePeerAccess(peerCtx)); } diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index c710aba6c8..86c323a8b6 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -90,14 +90,14 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priorit //--- hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags) { - HIP_INIT_CB_API(hipStreamCreateWithFlags, stream, flags); + HIP_INIT_API(hipStreamCreateWithFlags, stream, flags); return ihipLogStatus(ihipStreamCreate(stream, flags, priority_normal)); } //--- hipError_t hipStreamCreate(hipStream_t* stream) { - HIP_INIT_CB_API(hipStreamCreate, stream); + HIP_INIT_API(hipStreamCreate, stream); return ihipLogStatus(ihipStreamCreate(stream, hipStreamDefault, priority_normal)); } @@ -121,7 +121,7 @@ hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPrio } hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) { - HIP_INIT_SPECIAL_CB_API(hipStreamWaitEvent, TRACE_SYNC, stream, event, flags); + HIP_INIT_SPECIAL_API(hipStreamWaitEvent, TRACE_SYNC, stream, event, flags); hipError_t e = hipSuccess; @@ -152,7 +152,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int //--- hipError_t hipStreamQuery(hipStream_t stream) { - HIP_INIT_SPECIAL_CB_API(hipStreamQuery, TRACE_QUERY, stream); + HIP_INIT_SPECIAL_API(hipStreamQuery, TRACE_QUERY, stream); // Use default stream if 0 specified: if (stream == hipStreamNull) { @@ -175,7 +175,7 @@ hipError_t hipStreamQuery(hipStream_t stream) { //--- hipError_t hipStreamSynchronize(hipStream_t stream) { - HIP_INIT_SPECIAL_CB_API(hipStreamSynchronize, TRACE_SYNC, stream); + HIP_INIT_SPECIAL_API(hipStreamSynchronize, TRACE_SYNC, stream); return ihipLogStatus(ihipStreamSynchronize(stream)); } @@ -186,7 +186,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream) { * @return #hipSuccess, #hipErrorInvalidResourceHandle */ hipError_t hipStreamDestroy(hipStream_t stream) { - HIP_INIT_CB_API(hipStreamDestroy, stream); + HIP_INIT_API(hipStreamDestroy, stream); hipError_t e = hipSuccess; @@ -214,7 +214,7 @@ hipError_t hipStreamDestroy(hipStream_t stream) { //--- hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags) { - HIP_INIT_CB_API(hipStreamGetFlags, stream, flags); + HIP_INIT_API(hipStreamGetFlags, stream, flags); if (flags == NULL) { return ihipLogStatus(hipErrorInvalidValue); @@ -250,7 +250,7 @@ hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) { //--- hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData, unsigned int flags) { - HIP_INIT_CB_API(hipStreamAddCallback, stream, callback, userData, flags); + HIP_INIT_API(hipStreamAddCallback, stream, callback, userData, flags); hipError_t e = hipSuccess; // Create a thread in detached mode to handle callback diff --git a/projects/clr/hipamd/src/hip_surface.cpp b/projects/clr/hipamd/src/hip_surface.cpp index 6ac0dde3e3..9acd827f73 100644 --- a/projects/clr/hipamd/src/hip_surface.cpp +++ b/projects/clr/hipamd/src/hip_surface.cpp @@ -41,7 +41,7 @@ void saveSurfaceInfo(const hipSurface* pSurface, const hipResourceDesc* pResDesc // Surface Object APIs hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject, const hipResourceDesc* pResDesc) { - HIP_INIT_API(pSurfObject, pResDesc); + HIP_INIT_API(hipCreateSurfaceObject, pSurfObject, pResDesc); hipError_t hip_status = hipSuccess; auto ctx = ihipGetTlsDefaultCtx(); @@ -71,7 +71,7 @@ hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject, } hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject) { - HIP_INIT_API(surfaceObject); + HIP_INIT_API(hipDestroySurfaceObject, surfaceObject); hipError_t hip_status = hipSuccess; diff --git a/projects/clr/hipamd/src/hip_texture.cpp b/projects/clr/hipamd/src/hip_texture.cpp index 521e0e24de..316fba20cd 100644 --- a/projects/clr/hipamd/src/hip_texture.cpp +++ b/projects/clr/hipamd/src/hip_texture.cpp @@ -202,7 +202,7 @@ bool getHipTextureObject(hipTextureObject_t* pTexObject, hsa_ext_image_t& image, 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(hipCreateTextureObject, pTexObject, pResDesc, pTexDesc, pResViewDesc); hipError_t hip_status = hipSuccess; auto ctx = ihipGetTlsDefaultCtx(); @@ -311,7 +311,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou } hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject) { - HIP_INIT_API(textureObject); + HIP_INIT_API(hipDestroyTextureObject, textureObject); hipError_t hip_status = hipSuccess; @@ -335,7 +335,7 @@ hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject) { hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc, hipTextureObject_t textureObject) { - HIP_INIT_API(pResDesc, textureObject); + HIP_INIT_API(hipGetTextureObjectResourceDesc, pResDesc, textureObject); hipError_t hip_status = hipSuccess; auto ctx = ihipGetTlsDefaultCtx(); @@ -350,7 +350,7 @@ hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc, hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc, hipTextureObject_t textureObject) { - HIP_INIT_API(pResViewDesc, textureObject); + HIP_INIT_API(hipGetTextureObjectResourceViewDesc, pResViewDesc, textureObject); hipError_t hip_status = hipSuccess; auto ctx = ihipGetTlsDefaultCtx(); @@ -366,7 +366,7 @@ hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc, hipTextureObject_t textureObject) { - HIP_INIT_API(pTexDesc, textureObject); + HIP_INIT_API(hipGetTextureObjectTextureDesc, pTexDesc, textureObject); hipError_t hip_status = hipSuccess; @@ -444,7 +444,7 @@ hipError_t ihipBindTextureImpl(int dim, enum hipTextureReadMode readMode, size_t 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(hipBindTexture, offset, tex, devPtr, desc, size); hipError_t hip_status = hipSuccess; // TODO: hipReadModeElementType is default. hip_status = ihipBindTextureImpl(hipTextureType1D, hipReadModeElementType, offset, devPtr, desc, @@ -517,7 +517,7 @@ hipError_t ihipBindTexture2DImpl(int dim, enum hipTextureReadMode readMode, size 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(hipBindTexture2D, offset, tex, devPtr, desc, width, height, pitch); hipError_t hip_status = hipSuccess; hip_status = ihipBindTexture2DImpl(hipTextureType2D, hipReadModeElementType, offset, devPtr, desc, width, height, tex); @@ -613,7 +613,7 @@ hipError_t ihipBindTextureToArrayImpl(int dim, enum hipTextureReadMode readMode, hipError_t hipBindTextureToArray(textureReference* tex, hipArray_const_t array, const hipChannelFormatDesc* desc) { - HIP_INIT_API(tex, array, desc); + HIP_INIT_API(hipBindTextureToArray, tex, array, desc); hipError_t hip_status = hipSuccess; // TODO: hipReadModeElementType is default. hip_status = @@ -624,7 +624,7 @@ hipError_t hipBindTextureToArray(textureReference* tex, hipArray_const_t array, hipError_t hipBindTextureToMipmappedArray(textureReference* tex, hipMipmappedArray_const_t mipmappedArray, const hipChannelFormatDesc* desc) { - HIP_INIT_API(tex, mipmappedArray, desc); + HIP_INIT_API(hipBindTextureToMipmappedArray, tex, mipmappedArray, desc); hipError_t hip_status = hipSuccess; return ihipLogStatus(hip_status); } @@ -652,14 +652,14 @@ hipError_t ihipUnbindTextureImpl(const hipTextureObject_t& textureObject) { } hipError_t hipUnbindTexture(const textureReference* tex) { - HIP_INIT_API(tex); + HIP_INIT_API(hipUnbindTexture, tex); hipError_t hip_status = hipSuccess; hip_status = ihipUnbindTextureImpl(tex->textureObject); return ihipLogStatus(hip_status); } hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_const_t array) { - HIP_INIT_API(desc, array); + HIP_INIT_API(hipGetChannelDesc, desc, array); hipError_t hip_status = hipSuccess; auto ctx = ihipGetTlsDefaultCtx(); @@ -670,7 +670,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(hipGetTextureAlignmentOffset, offset, tex); hipError_t hip_status = hipSuccess; @@ -683,7 +683,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(hipGetTextureReference, tex, symbol); hipError_t hip_status = hipSuccess; @@ -694,7 +694,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(hipTexRefSetFormat, tex, fmt, NumPackedComponents); hipError_t hip_status = hipSuccess; tex->format = fmt; tex->numChannels = NumPackedComponents; @@ -702,28 +702,28 @@ 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(hipTexRefSetFlags, tex, flags); hipError_t hip_status = hipSuccess; tex->normalized = flags; return ihipLogStatus(hip_status); } hipError_t hipTexRefSetFilterMode(textureReference* tex, hipTextureFilterMode fm) { - HIP_INIT_API(tex, fm); + HIP_INIT_API(hipTexRefSetFilterMode, tex, fm); hipError_t hip_status = hipSuccess; tex->filterMode = fm; return ihipLogStatus(hip_status); } hipError_t hipTexRefSetAddressMode(textureReference* tex, int dim, hipTextureAddressMode am) { - HIP_INIT_API(tex, dim, am); + HIP_INIT_API(hipTexRefSetAddressMode, tex, dim, am); hipError_t hip_status = hipSuccess; tex->addressMode[dim] = am; return ihipLogStatus(hip_status); } hipError_t hipTexRefSetArray(textureReference* tex, hipArray_const_t array, unsigned int flags) { - HIP_INIT_API(tex, array, flags); + HIP_INIT_API(hipTexRefSetArray, tex, array, flags); hipError_t hip_status = hipSuccess; hip_status = ihipBindTextureToArrayImpl(array->textureType, hipReadModeElementType, array, @@ -734,7 +734,7 @@ hipError_t hipTexRefSetArray(textureReference* tex, hipArray_const_t array, unsi hipError_t hipTexRefSetAddress(size_t* offset, textureReference* tex, hipDeviceptr_t devPtr, size_t size) { - HIP_INIT_API(offset, tex, devPtr, size); + HIP_INIT_API(hipTexRefSetAddress, offset, tex, devPtr, size); hipError_t hip_status = hipSuccess; // TODO: hipReadModeElementType is default. hip_status = ihipBindTextureImpl(hipTextureType1D, hipReadModeElementType, offset, devPtr, NULL, @@ -744,7 +744,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(hipTexRefSetAddress2D, tex, desc, devPtr, pitch); size_t offset; hipError_t hip_status = hipSuccess; // TODO: hipReadModeElementType is default.