diff --git a/projects/clr/CHANGELOG.md b/projects/clr/CHANGELOG.md index 344180dd44..a2b70532b0 100644 --- a/projects/clr/CHANGELOG.md +++ b/projects/clr/CHANGELOG.md @@ -41,6 +41,7 @@ Full documentation for HIP is available at [rocm.docs.amd.com](https://rocm.docs - `hipDrvLaunchKernelEx` dispatches the device kernel represented by a HIP function object. - `hipMemGetHandleForAddressRange` gets a handle for the address range requested. - `num_threads` Total number of threads in the group. The legacy API size is alias. + - `hipGetDriverEntryPoint ` gets function pointer of a HIP API. * New support for Open Compute Project (OCP) floating-point `FP4`/`FP6`/`FP8` as the following. For details, see [Low precision floating point document](https://rocm.docs.amd.com/projects/HIP/en/latest/reference/low_fp_types.html). - Data types for `FP4`/`FP6`/`FP8`. - HIP APIs for `FP4`/`FP6`/`FP8`, which are compatible with corresponding CUDA APIs. diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h index 77b3bee5a4..ea4691d327 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h @@ -82,6 +82,9 @@ THE SOFTWARE. #define hipStreamIsCapturing __HIP_API_SPT(hipStreamIsCapturing) #define hipStreamGetCaptureInfo __HIP_API_SPT(hipStreamGetCaptureInfo) #define hipStreamGetCaptureInfo_v2 __HIP_API_SPT(hipStreamGetCaptureInfo_v2) + + // Driver Entry Point API + #define hipGetDriverEntryPoint __HIP_API_SPT(hipGetDriverEntryPoint) #endif #ifdef __cplusplus @@ -192,6 +195,8 @@ hipError_t hipStreamGetCaptureInfo_v2_spt(hipStream_t stream, hipStreamCaptureSt const hipGraphNode_t** dependencies_out, size_t* numDependencies_out); hipError_t hipLaunchHostFunc_spt(hipStream_t stream, hipHostFn_t fn, void* userData); +hipError_t hipGetDriverEntryPoint_spt(const char* symbol, void** funcPtr, unsigned long long flags, + hipDriverEntryPointQueryResult* status); #ifdef __cplusplus diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp b/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp index ca3eac3355..83eb82a96a 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp @@ -1073,6 +1073,13 @@ typedef hipError_t (*t_hipMemcpy3DBatchAsync) (size_t numOps, struct hipMemcpy3D typedef hipError_t (*t_hipMemcpy3DPeer) (hipMemcpy3DPeerParms *p); typedef hipError_t (*t_hipMemcpy3DPeerAsync) (hipMemcpy3DPeerParms *p, hipStream_t stream); +typedef hipError_t (*t_hipGetDriverEntryPoint)(const char* symbol, void** funcPtr, + unsigned long long flags, + hipDriverEntryPointQueryResult* status); +typedef hipError_t (*t_hipGetDriverEntryPoint_spt)(const char* symbol, void** funcPtr, + unsigned long long flags, + hipDriverEntryPointQueryResult* status); + // HIP Compiler dispatch table struct HipCompilerDispatchTable { // HIP_COMPILER_API_TABLE_STEP_VERSION == 0 @@ -1629,6 +1636,8 @@ struct HipDispatchTable { t_hipMemcpy3DBatchAsync hipMemcpy3DBatchAsync_fn; t_hipMemcpy3DPeer hipMemcpy3DPeer_fn; t_hipMemcpy3DPeerAsync hipMemcpy3DPeerAsync_fn; + t_hipGetDriverEntryPoint hipGetDriverEntryPoint_fn; + t_hipGetDriverEntryPoint_spt hipGetDriverEntryPoint_spt_fn; // HIP_RUNTIME_API_TABLE_STEP_VERSION = 14 // removed HIP_MEMSET_NODE_PARAMS replaced by hipMemsetParams diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h index 102848d2e5..503c3616da 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h @@ -452,7 +452,8 @@ enum hip_api_id_t { HIP_API_ID_hipMemcpy3DPeer = 432, HIP_API_ID_hipMemcpy3DPeerAsync = 433, HIP_API_ID_hipMemcpyBatchAsync = 434, - HIP_API_ID_LAST = 434, + HIP_API_ID_hipGetDriverEntryPoint = 435, + HIP_API_ID_LAST = 435, HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice), HIP_API_ID_hipGetDeviceProperties = HIP_API_ID_CONCAT(HIP_API_ID_,hipGetDeviceProperties), @@ -465,6 +466,7 @@ enum hip_api_id_t { HIP_API_ID_hipDestroyTextureObject = HIP_API_ID_NONE, HIP_API_ID_hipDeviceGetCount = HIP_API_ID_NONE, HIP_API_ID_hipDeviceGetTexture1DLinearMaxWidth = HIP_API_ID_NONE, + HIP_API_ID_hipGetDriverEntryPoint_spt = HIP_API_ID_NONE, HIP_API_ID_hipGetTextureAlignmentOffset = HIP_API_ID_NONE, HIP_API_ID_hipGetTextureObjectResourceDesc = HIP_API_ID_NONE, HIP_API_ID_hipGetTextureObjectResourceViewDesc = HIP_API_ID_NONE, @@ -602,6 +604,7 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipGetDeviceFlags: return "hipGetDeviceFlags"; case HIP_API_ID_hipGetDevicePropertiesR0000: return "hipGetDevicePropertiesR0000"; case HIP_API_ID_hipGetDevicePropertiesR0600: return "hipGetDevicePropertiesR0600"; + case HIP_API_ID_hipGetDriverEntryPoint: return "hipGetDriverEntryPoint"; case HIP_API_ID_hipGetErrorString: return "hipGetErrorString"; case HIP_API_ID_hipGetFuncBySymbol: return "hipGetFuncBySymbol"; case HIP_API_ID_hipGetLastError: return "hipGetLastError"; @@ -1031,6 +1034,7 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipGetDeviceFlags", name) == 0) return HIP_API_ID_hipGetDeviceFlags; if (strcmp("hipGetDevicePropertiesR0000", name) == 0) return HIP_API_ID_hipGetDevicePropertiesR0000; if (strcmp("hipGetDevicePropertiesR0600", name) == 0) return HIP_API_ID_hipGetDevicePropertiesR0600; + if (strcmp("hipGetDriverEntryPoint", name) == 0) return HIP_API_ID_hipGetDriverEntryPoint; if (strcmp("hipGetErrorString", name) == 0) return HIP_API_ID_hipGetErrorString; if (strcmp("hipGetFuncBySymbol", name) == 0) return HIP_API_ID_hipGetFuncBySymbol; if (strcmp("hipGetLastError", name) == 0) return HIP_API_ID_hipGetLastError; @@ -1908,6 +1912,15 @@ typedef struct hip_api_data_s { hipDeviceProp_tR0600 prop__val; int deviceId; } hipGetDevicePropertiesR0600; + struct { + const char* symbol; + char symbol__val; + void** funcPtr; + void* funcPtr__val; + unsigned long long flags; + hipDriverEntryPointQueryResult* driverStatus; + hipDriverEntryPointQueryResult driverStatus__val; + } hipGetDriverEntryPoint; struct { hipFunction_t* functionPtr; hipFunction_t functionPtr__val; @@ -4481,6 +4494,13 @@ typedef struct hip_api_data_s { cb_data.args.hipGetDevicePropertiesR0600.prop = (hipDeviceProp_tR0600*)prop; \ cb_data.args.hipGetDevicePropertiesR0600.deviceId = (int)device; \ }; +// hipGetDriverEntryPoint[('const char*', 'symbol'), ('void**', 'funcPtr'), ('unsigned long long', 'flags'), ('hipDriverEntryPointQueryResult*', 'driverStatus')] +#define INIT_hipGetDriverEntryPoint_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipGetDriverEntryPoint.symbol = (symbol) ? strdup(symbol) : NULL; \ + cb_data.args.hipGetDriverEntryPoint.funcPtr = (void**)funcPtr; \ + cb_data.args.hipGetDriverEntryPoint.flags = (unsigned long long)flags; \ + cb_data.args.hipGetDriverEntryPoint.driverStatus = (hipDriverEntryPointQueryResult*)status; \ +}; // hipGetErrorString[] #define INIT_hipGetErrorString_CB_ARGS_DATA(cb_data) { \ }; @@ -6496,6 +6516,8 @@ typedef struct hip_api_data_s { #define INIT_hipDeviceGetCount_CB_ARGS_DATA(cb_data) {}; // hipDeviceGetTexture1DLinearMaxWidth() #define INIT_hipDeviceGetTexture1DLinearMaxWidth_CB_ARGS_DATA(cb_data) {}; +// hipGetDriverEntryPoint_spt() +#define INIT_hipGetDriverEntryPoint_spt_CB_ARGS_DATA(cb_data) {}; // hipGetTextureAlignmentOffset() #define INIT_hipGetTextureAlignmentOffset_CB_ARGS_DATA(cb_data) {}; // hipGetTextureObjectResourceDesc() @@ -6969,6 +6991,12 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { case HIP_API_ID_hipGetDevicePropertiesR0600: if (data->args.hipGetDevicePropertiesR0600.prop) data->args.hipGetDevicePropertiesR0600.prop__val = *(data->args.hipGetDevicePropertiesR0600.prop); break; +// hipGetDriverEntryPoint[('const char*', 'symbol'), ('void**', 'funcPtr'), ('unsigned long long', 'flags'), ('hipDriverEntryPointQueryResult*', 'driverStatus')] + case HIP_API_ID_hipGetDriverEntryPoint: + if (data->args.hipGetDriverEntryPoint.symbol) data->args.hipGetDriverEntryPoint.symbol__val = *(data->args.hipGetDriverEntryPoint.symbol); + if (data->args.hipGetDriverEntryPoint.funcPtr) data->args.hipGetDriverEntryPoint.funcPtr__val = *(data->args.hipGetDriverEntryPoint.funcPtr); + if (data->args.hipGetDriverEntryPoint.driverStatus) data->args.hipGetDriverEntryPoint.driverStatus__val = *(data->args.hipGetDriverEntryPoint.driverStatus); + break; // hipGetErrorString[] case HIP_API_ID_hipGetErrorString: break; @@ -9002,6 +9030,17 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << ", deviceId="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetDevicePropertiesR0600.deviceId); oss << ")"; break; + case HIP_API_ID_hipGetDriverEntryPoint: + oss << "hipGetDriverEntryPoint("; + if (data->args.hipGetDriverEntryPoint.symbol == NULL) oss << "symbol=NULL"; + else { oss << "symbol="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetDriverEntryPoint.symbol__val); } + if (data->args.hipGetDriverEntryPoint.funcPtr == NULL) oss << ", funcPtr=NULL"; + else { oss << ", funcPtr="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetDriverEntryPoint.funcPtr__val); } + oss << ", flags="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetDriverEntryPoint.flags); + if (data->args.hipGetDriverEntryPoint.driverStatus == NULL) oss << ", driverStatus=NULL"; + else { oss << ", driverStatus="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetDriverEntryPoint.driverStatus__val); } + oss << ")"; + break; case HIP_API_ID_hipGetErrorString: oss << "hipGetErrorString("; oss << ")"; diff --git a/projects/clr/hipamd/src/amdhip.def b/projects/clr/hipamd/src/amdhip.def index fde675a3da..a98899f526 100644 --- a/projects/clr/hipamd/src/amdhip.def +++ b/projects/clr/hipamd/src/amdhip.def @@ -505,4 +505,6 @@ hipModuleLoadFatBinary hipMemcpyBatchAsync hipMemcpy3DBatchAsync hipMemcpy3DPeer -hipMemcpy3DPeerAsync \ No newline at end of file +hipMemcpy3DPeerAsync +hipGetDriverEntryPoint +hipGetDriverEntryPoint_spt diff --git a/projects/clr/hipamd/src/hip_api_trace.cpp b/projects/clr/hipamd/src/hip_api_trace.cpp index 08dd982f9a..5ed8e3206d 100644 --- a/projects/clr/hipamd/src/hip_api_trace.cpp +++ b/projects/clr/hipamd/src/hip_api_trace.cpp @@ -186,6 +186,10 @@ hipError_t hipGetDeviceCount(int* count); hipError_t hipGetDeviceFlags(unsigned int* flags); hipError_t hipGetDevicePropertiesR0600(hipDeviceProp_tR0600* prop, int deviceId); hipError_t hipGetDevicePropertiesR0000(hipDeviceProp_tR0000* prop, int device); +hipError_t hipGetDriverEntryPoint(const char* symbol, void** funcPtr, unsigned long long flags, + hipDriverEntryPointQueryResult* status); +hipError_t hipGetDriverEntryPoint_spt(const char* symbol, void** funcPtr, unsigned long long flags, + hipDriverEntryPointQueryResult* status); const char* hipGetErrorName(hipError_t hip_error); const char* hipGetErrorString(hipError_t hipError); hipError_t hipGetLastError(void); @@ -1340,6 +1344,8 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) { ptrDispatchTable->hipGetStreamDeviceId_fn = hip::hipGetStreamDeviceId; ptrDispatchTable->hipDrvGraphAddMemsetNode_fn = hip::hipDrvGraphAddMemsetNode; ptrDispatchTable->hipGetDevicePropertiesR0000_fn = hip::hipGetDevicePropertiesR0000; + ptrDispatchTable->hipGetDriverEntryPoint_fn = hip::hipGetDriverEntryPoint; + ptrDispatchTable->hipGetDriverEntryPoint_spt_fn = hip::hipGetDriverEntryPoint_spt; ptrDispatchTable->hipExtGetLastError_fn = hip::hipExtGetLastError; ptrDispatchTable->hipTexRefGetBorderColor_fn = hip::hipTexRefGetBorderColor; ptrDispatchTable->hipTexRefGetArray_fn = hip::hipTexRefGetArray; @@ -2042,13 +2048,15 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipMemcpyBatchAsync_fn, 487); HIP_ENFORCE_ABI(HipDispatchTable, hipMemcpy3DBatchAsync_fn, 488); HIP_ENFORCE_ABI(HipDispatchTable, hipMemcpy3DPeer_fn, 489); HIP_ENFORCE_ABI(HipDispatchTable, hipMemcpy3DPeerAsync_fn, 490); +HIP_ENFORCE_ABI(HipDispatchTable, hipGetDriverEntryPoint_fn, 491); +HIP_ENFORCE_ABI(HipDispatchTable, hipGetDriverEntryPoint_spt_fn, 492); // if HIP_ENFORCE_ABI entries are added for each new function pointer in the table, the number below // will be +1 of the number in the last HIP_ENFORCE_ABI line. E.g.: // // HIP_ENFORCE_ABI(, , 8) // // HIP_ENFORCE_ABI_VERSIONING(
, 9) <- 8 + 1 = 9 -HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 491) +HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 493) static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 14, "If you get this error, add new HIP_ENFORCE_ABI(...) code for the new function " diff --git a/projects/clr/hipamd/src/hip_device_runtime.cpp b/projects/clr/hipamd/src/hip_device_runtime.cpp index 19e96101fa..21d890252c 100644 --- a/projects/clr/hipamd/src/hip_device_runtime.cpp +++ b/projects/clr/hipamd/src/hip_device_runtime.cpp @@ -21,6 +21,7 @@ #include #include "hip_internal.hpp" +#include "hip_platform.hpp" #undef hipChooseDevice #undef hipDeviceProp_t @@ -716,6 +717,60 @@ hipError_t hipGetDeviceFlags(unsigned int* flags) { HIP_RETURN(hipSuccess); } +hipError_t hipGetDriverEntryPoint_common(const char* symbol, void** funcPtr, unsigned long long flags, + hipDriverEntryPointQueryResult* status) { + std::string symbolString = symbol; + if (symbol == nullptr || symbolString == "" || funcPtr == nullptr) { + return hipErrorInvalidValue; + } + + if (flags != hipEnableDefault && flags != hipEnableLegacyStream + && flags != hipEnablePerThreadDefaultStream) { + return hipErrorInvalidValue; + } + + void* handle = hip::PlatformState::instance().getDynamicLibraryHandle(); + if (handle == nullptr) { + return hipErrorInvalidValue; + } + + if (flags == hipEnablePerThreadDefaultStream) { + symbolString += "_spt"; + } + + *funcPtr = amd::Os::getSymbol(handle, symbolString.c_str()); + if (funcPtr == nullptr) { + if (flags == hipEnablePerThreadDefaultStream) { + *funcPtr = amd::Os::getSymbol(handle, symbol); + } + if (funcPtr == nullptr) { + if (status != nullptr) { + *status = hipDriverEntryPointSymbolNotFound; + } + return hipErrorInvalidValue; + } + } + + if (status != nullptr) { + *status = hipDriverEntryPointSuccess; + } + + return hipSuccess; +} + +hipError_t hipGetDriverEntryPoint(const char* symbol, void** funcPtr, unsigned long long flags, + hipDriverEntryPointQueryResult* status) { + HIP_INIT_API(hipGetDriverEntryPoint, symbol, funcPtr, flags, status); + HIP_RETURN(hipGetDriverEntryPoint_common(symbol, funcPtr, flags, status)); +} + +hipError_t hipGetDriverEntryPoint_spt(const char* symbol, void** funcPtr, unsigned long long flags, + hipDriverEntryPointQueryResult* status) { + HIP_INIT_API(hipGetDriverEntryPoint, symbol, funcPtr, flags, status); + flags = (flags == hipEnableDefault) ? hipEnablePerThreadDefaultStream : flags; + HIP_RETURN(hipGetDriverEntryPoint_common(symbol, funcPtr, flags, status)); +} + hipError_t hipSetDevice(int device) { HIP_INIT_API_NO_RETURN(hipSetDevice, device); diff --git a/projects/clr/hipamd/src/hip_hcc.map.in b/projects/clr/hipamd/src/hip_hcc.map.in index 06ec368c6c..09d1d9e84a 100644 --- a/projects/clr/hipamd/src/hip_hcc.map.in +++ b/projects/clr/hipamd/src/hip_hcc.map.in @@ -622,6 +622,8 @@ global: hipMemcpy3DBatchAsync; hipMemcpy3DPeer; hipMemcpy3DPeerAsync; + hipGetDriverEntryPoint; + hipGetDriverEntryPoint_spt; local: *; } hip_6.5; \ No newline at end of file diff --git a/projects/clr/hipamd/src/hip_table_interface.cpp b/projects/clr/hipamd/src/hip_table_interface.cpp index f3616a18ee..2106ee2a45 100644 --- a/projects/clr/hipamd/src/hip_table_interface.cpp +++ b/projects/clr/hipamd/src/hip_table_interface.cpp @@ -419,6 +419,16 @@ extern "C" hipError_t hipGetDevicePropertiesR0600(hipDeviceProp_tR0600* prop, in extern "C" hipError_t hipGetDevicePropertiesR0000(hipDeviceProp_tR0000* prop, int device) { return hip::GetHipDispatchTable()->hipGetDevicePropertiesR0000_fn(prop, device); } +hipError_t hipGetDriverEntryPoint(const char* symbol, void** funcPtr, unsigned long long flags, + hipDriverEntryPointQueryResult* status) { + return hip::GetHipDispatchTable()->hipGetDriverEntryPoint_fn(symbol, funcPtr, flags, + status); +} +hipError_t hipGetDriverEntryPoint_spt(const char* symbol, void** funcPtr, unsigned long long flags, + hipDriverEntryPointQueryResult* status) { + return hip::GetHipDispatchTable()->hipGetDriverEntryPoint_spt_fn(symbol, funcPtr, flags, + status); +} const char* hipGetErrorName(hipError_t hip_error) { return hip::GetHipDispatchTable()->hipGetErrorName_fn(hip_error); }