SWDEV-546285 - add hipGetDriverEntryPoint (#855)

[ROCm/clr commit: 789e2029ca]
2025-08-15 20:08:21 -07:00
commit ad9eb56dd4
@@ -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.
@@ -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
@@ -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
@@ -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 << ")";
@@ -505,4 +505,6 @@ hipModuleLoadFatBinary
 hipMemcpyBatchAsync
 hipMemcpy3DBatchAsync
 hipMemcpy3DPeer
-hipMemcpy3DPeerAsync
+hipMemcpy3DPeerAsync
+hipGetDriverEntryPoint
+hipGetDriverEntryPoint_spt
@@ -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(<table>, <functor>, 8)
 //
 //  HIP_ENFORCE_ABI_VERSIONING(<table>, 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 "
@@ -21,6 +21,7 @@
 #include <hip/hip_runtime.h>

 #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);

@@ -622,6 +622,8 @@ global:
    hipMemcpy3DBatchAsync;
    hipMemcpy3DPeer;
    hipMemcpy3DPeerAsync;
+    hipGetDriverEntryPoint;
+    hipGetDriverEntryPoint_spt;
 local:
    *;
 } hip_6.5;
@@ -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);
 }