diff --git a/hipamd/include/hip/hcc_detail/hip_prof_str.h b/hipamd/include/hip/hcc_detail/hip_prof_str.h index 512e8a3188..9fe681c18d 100644 --- a/hipamd/include/hip/hcc_detail/hip_prof_str.h +++ b/hipamd/include/hip/hcc_detail/hip_prof_str.h @@ -5,6 +5,7 @@ #include // Dummy API callbacks definition +#define INIT_NONE_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) {}; @@ -35,6 +36,9 @@ #define INIT_hipGetErrorString_CB_ARGS_DATA(cb_data) {}; #define INIT_hipCreateSurfaceObject_CB_ARGS_DATA(cb_data) {}; #define INIT_hipDestroySurfaceObject_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipStreamCreateWithPriority_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipDeviceGetStreamPriorityRange_CB_ARGS_DATA(cb_data) {}; +#define INIT_hipStreamGetPriority_CB_ARGS_DATA(cb_data) {}; // HIP API callbacks ID enumaration enum hip_api_id_t { @@ -180,6 +184,7 @@ enum hip_api_id_t { HIP_API_ID_NUMBER = 139, HIP_API_ID_ANY = 140, + HIP_API_ID_NONE = 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, @@ -210,6 +215,9 @@ enum hip_api_id_t { HIP_API_ID_hipGetErrorString = HIP_API_ID_NUMBER, HIP_API_ID_hipCreateSurfaceObject = HIP_API_ID_NUMBER, HIP_API_ID_hipDestroySurfaceObject = HIP_API_ID_NUMBER, + HIP_API_ID_hipStreamCreateWithPriority = HIP_API_ID_NUMBER, + HIP_API_ID_hipDeviceGetStreamPriorityRange = HIP_API_ID_NUMBER, + HIP_API_ID_hipStreamGetPriority = HIP_API_ID_NUMBER, }; // Return HIP API string @@ -339,7 +347,6 @@ static const char* hip_api_name(const uint32_t& id) { case HIP_API_ID_hipCtxEnablePeerAccess: return "hipCtxEnablePeerAccess"; case HIP_API_ID_hipMemcpyDtoHAsync: return "hipMemcpyDtoHAsync"; case HIP_API_ID_hipModuleLaunchKernel: return "hipModuleLaunchKernel"; - case HIP_API_ID_hipHccModuleLaunchKernel: return "hipHccModuleLaunchKernel"; case HIP_API_ID_hipModuleGetTexRef: return "hipModuleGetTexRef"; case HIP_API_ID_hipRemoveActivityCallback: return "hipRemoveActivityCallback"; case HIP_API_ID_hipDeviceGetLimit: return "hipDeviceGetLimit"; @@ -2498,29 +2505,4 @@ const char* hipApiString(hip_api_id_t id, const hip_api_data_t* data) { }; #endif -// HIP API activity record type -// Base record type -struct hip_act_record_t { - uint32_t domain; // activity domain id - uint32_t op_id; // operation id, dispatch/copy/barrier - uint32_t activity_kind; // activity kind - uint64_t correlation_id; // activity correlation ID - uint64_t begin_ns; // host begin timestamp, nano-seconds - uint64_t end_ns; // host end timestamp, nano-seconds -}; -// Async record type -struct hip_async_record_t : hip_act_record_t { - int device_id; - uint64_t stream_id; -}; -// Dispatch record type -struct hip_dispatch_record_t : hip_async_record_t {}; -// Barrier record type -struct hip_barrier_record_t : hip_async_record_t {}; -// Memcpy record type -struct hip_copy_record_t : hip_async_record_t { - size_t bytes; -}; -// Generic async operation record -typedef hip_copy_record_t hip_ops_record_t; #endif // _HIP_CBSTR diff --git a/hipamd/src/hip_clang.cpp b/hipamd/src/hip_clang.cpp index 44080884e7..fe08bbe45c 100644 --- a/hipamd/src/hip_clang.cpp +++ b/hipamd/src/hip_clang.cpp @@ -159,7 +159,7 @@ extern "C" void __hipRegisterFunction( dim3* gridDim, int* wSize) { - HIP_INIT_API(modules, hostFunction, deviceFunction, deviceName); + HIP_INIT_API(NONE, modules, hostFunction, deviceFunction, deviceName); std::vector functions{g_deviceCnt}; assert(modules && modules->size() >= g_deviceCnt); @@ -214,7 +214,7 @@ hipError_t hipSetupArgument( size_t size, size_t offset) { - HIP_INIT_API(arg, size, offset); + HIP_INIT_API(hipSetupArgument, arg, size, offset); auto ctx = ihipGetTlsDefaultCtx(); LockedAccessor_CtxCrit_t crit(ctx->criticalData()); auto& arguments = crit->_execStack.top()._arguments; @@ -229,7 +229,7 @@ hipError_t hipSetupArgument( hipError_t hipLaunchByPtr(const void *hostFunction) { - HIP_INIT_API(hostFunction); + HIP_INIT_API(hipLaunchByPtr, hostFunction); ihipExec_t exec; { auto ctx = ihipGetTlsDefaultCtx(); diff --git a/hipamd/src/hip_stream.cpp b/hipamd/src/hip_stream.cpp index 86c323a8b6..5331b41576 100644 --- a/hipamd/src/hip_stream.cpp +++ b/hipamd/src/hip_stream.cpp @@ -104,7 +104,7 @@ hipError_t hipStreamCreate(hipStream_t* stream) { //--- hipError_t hipStreamCreateWithPriority(hipStream_t* stream, unsigned int flags, int priority) { - HIP_INIT_API(stream, flags, priority); + HIP_INIT_API(hipStreamCreateWithPriority, stream, flags, priority); // clamp priority to range [priority_high:priority_low] priority = (priority < priority_high ? priority_high : (priority > priority_low ? priority_low : priority)); @@ -113,7 +113,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 != NULL) *leastPriority = priority_low; if (greatestPriority != NULL) *greatestPriority = priority_high; @@ -229,7 +229,7 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags) { //-- hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) { - HIP_INIT_API(stream, priority); + HIP_INIT_API(hipStreamGetPriority, stream, priority); if (priority == NULL) { return ihipLogStatus(hipErrorInvalidValue);