SWDEV-554372 - Add 3 HIP_GET_PROC_ADDRESS_xxx flags (#1057)

projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h

@@ -85,6 +85,7 @@ THE SOFTWARE.
 
 // Driver Entry Point API
 #define hipGetDriverEntryPoint __HIP_API_SPT(hipGetDriverEntryPoint)
+#define hipGetProcAddress __HIP_API_SPT(hipGetProcAddress)
 #endif
 
 #ifdef __cplusplus
@@ -197,6 +198,8 @@ hipError_t hipStreamGetCaptureInfo_v2_spt(hipStream_t stream,
 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);
+hipError_t hipGetProcAddress_spt(const char* symbol, void** pfn, int hipVersion, uint64_t flags,
+                                 hipDriverProcAddressQueryResult* symbolStatus);
 
 
 #ifdef __cplusplus

projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp

@@ -1111,6 +1111,8 @@ typedef hipError_t (*t_hipLibraryEnumerateKernels)(hipKernel_t* kernels, unsigne
                                                    hipLibrary_t library);
 typedef hipError_t (*t_hipKernelGetLibrary)(hipLibrary_t* library, hipKernel_t kernel);
 typedef hipError_t (*t_hipKernelGetName)(const char** name, hipKernel_t kernel);
+typedef hipError_t (*t_hipGetProcAddress_spt)(const char* symbol, void** pfn, int hipVersion, uint64_t flags,
+                                              hipDriverProcAddressQueryResult* symbolStatus);
 
 // HIP Compiler dispatch table
 struct HipCompilerDispatchTable {
@@ -1696,6 +1698,7 @@ struct HipDispatchTable {
   
   // HIP_RUNTIME_API_TABLE_STEP_VERSION == 18
   t_hipOccupancyAvailableDynamicSMemPerBlock hipOccupancyAvailableDynamicSMemPerBlock_fn;
+  t_hipGetProcAddress_spt hipGetProcAddress_spt_fn;
 
   // DO NOT EDIT ABOVE!
   // HIP_RUNTIME_API_TABLE_STEP_VERSION == 19

projects/clr/hipamd/src/amdhip.def

@@ -521,3 +521,4 @@ hipLibraryEnumerateKernels
 hipKernelGetLibrary
 hipKernelGetName
 hipOccupancyAvailableDynamicSMemPerBlock
+hipGetProcAddress_spt

projects/clr/hipamd/src/hip_api_trace.cpp

@@ -790,6 +790,8 @@ hipError_t hipTexRefGetBorderColor(float* pBorderColor, const textureReference*
 hipError_t hipTexRefGetArray(hipArray_t* pArray, const textureReference* texRef);
 hipError_t hipGetProcAddress(const char* symbol, void** pfn, int hipVersion, uint64_t flags,
                              hipDriverProcAddressQueryResult* symbolStatus = NULL);
+hipError_t hipGetProcAddress_spt(const char* symbol, void** pfn, int hipVersion, uint64_t flags,
+                                 hipDriverProcAddressQueryResult* symbolStatus = NULL);
 hipError_t hipStreamBeginCaptureToGraph(hipStream_t stream, hipGraph_t graph,
                                         const hipGraphNode_t* dependencies,
                                         const hipGraphEdgeData* dependencyData,
@@ -1377,6 +1379,7 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) {
   ptrDispatchTable->hipTexRefGetBorderColor_fn = hip::hipTexRefGetBorderColor;
   ptrDispatchTable->hipTexRefGetArray_fn = hip::hipTexRefGetArray;
   ptrDispatchTable->hipGetProcAddress_fn = hip::hipGetProcAddress;
+  ptrDispatchTable->hipGetProcAddress_spt_fn = hip::hipGetProcAddress_spt;
   ptrDispatchTable->hipStreamBeginCaptureToGraph_fn = hip::hipStreamBeginCaptureToGraph;
   ptrDispatchTable->hipGetFuncBySymbol_fn = hip::hipGetFuncBySymbol;
   ptrDispatchTable->hipSetValidDevices_fn = hip::hipSetValidDevices;
@@ -2104,13 +2107,14 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipKernelGetLibrary_fn, 503);
 HIP_ENFORCE_ABI(HipDispatchTable, hipKernelGetName_fn, 504);
 // HIP_RUNTIME_API_TABLE_STEP_VERSION == 18
 HIP_ENFORCE_ABI(HipDispatchTable, hipOccupancyAvailableDynamicSMemPerBlock_fn, 505);
+HIP_ENFORCE_ABI(HipDispatchTable, hipGetProcAddress_spt_fn, 506);
 // 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, 506)
+HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 507)
 
 static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 18,
               "If you get this error, add new HIP_ENFORCE_ABI(...) code for the new function "

projects/clr/hipamd/src/hip_device.cpp

@@ -766,55 +766,83 @@ hipError_t hipGetDevicePropertiesR0000(hipDeviceProp_tR0000* prop, int device) {
   HIP_RETURN(hipSuccess);
 }
 
-hipError_t hipGetProcAddress(const char* symbol, void** pfn, int hipVersion, uint64_t flags,
+hipError_t hipGetProcAddress_common(const char* symbol, void** pfn, int hipVersion, uint64_t flags,
                              hipDriverProcAddressQueryResult* symbolStatus) {
-  HIP_INIT_API(hipGetProcAddress, symbol, pfn, hipVersion, flags, symbolStatus);
-
+  if (symbol == nullptr || std::string_view{symbol}.empty() || pfn == nullptr) {
+    return hipErrorInvalidValue;
+  }
   std::string symbolString = symbol;
 
-  if (symbol == nullptr || symbolString == "" || pfn == nullptr) {
-    HIP_RETURN(hipErrorInvalidValue);
+  if (flags != HIP_GET_PROC_ADDRESS_DEFAULT && flags != HIP_GET_PROC_ADDRESS_LEGACY_STREAM
+      && flags != HIP_GET_PROC_ADDRESS_PER_THREAD_DEFAULT_STREAM) {
+    return hipErrorInvalidValue;
   }
 
+  bool checkSpt = (flags == HIP_GET_PROC_ADDRESS_PER_THREAD_DEFAULT_STREAM);
   if (symbolString == "hipGetDeviceProperties") {
     if (hipVersion >= 600) {
       symbolString = "hipGetDevicePropertiesR0600";
     }
+    checkSpt = false;
   } else if (symbolString == "hipChooseDevice") {
     if (hipVersion >= 600) {
       symbolString = "hipChooseDeviceR0600";
     }
+    checkSpt = false;
   } else if (symbolString == "hipAmdFileRead") {
     *pfn = reinterpret_cast<void*>(&hipAmdFileRead);
     if (symbolStatus != nullptr) {
       *symbolStatus = HIP_GET_PROC_ADDRESS_SUCCESS;
     }
-    HIP_RETURN(hipSuccess);
+    return hipSuccess;
   } else if (symbolString == "hipAmdFileWrite") {
     *pfn = reinterpret_cast<void*>(&hipAmdFileWrite);
     if (symbolStatus != nullptr) {
       *symbolStatus = HIP_GET_PROC_ADDRESS_SUCCESS;
     }
-    HIP_RETURN(hipSuccess);
+    return hipSuccess;
   }
 
   void* handle = hip::PlatformState::instance().getDynamicLibraryHandle();
   if (handle == nullptr) {
-    HIP_RETURN(hipErrorInvalidValue);
+    return hipErrorInvalidValue;
+  }
+
+  if (checkSpt) {
+      symbolString += "_spt";
   }
 
   *pfn = amd::Os::getSymbol(handle, symbolString.c_str());
-  if (!(*pfn)) {
-    if (symbolStatus != nullptr) {
-      *symbolStatus = HIP_GET_PROC_ADDRESS_SYMBOL_NOT_FOUND;
+  if (*pfn == nullptr) {
+    if (checkSpt) {
+      *pfn = amd::Os::getSymbol(handle, symbol);
+    }
+    if (*pfn == nullptr) {
+      if (symbolStatus != nullptr) {
+        *symbolStatus = HIP_GET_PROC_ADDRESS_SYMBOL_NOT_FOUND;
+      }
+      return hipErrorInvalidValue;
     }
-    HIP_RETURN(hipErrorInvalidValue);
   }
 
   if (symbolStatus != nullptr) {
     *symbolStatus = HIP_GET_PROC_ADDRESS_SUCCESS;
   }
-  HIP_RETURN(hipSuccess);
+
+  return hipSuccess;
+}
+
+hipError_t hipGetProcAddress(const char* symbol, void** pfn, int hipVersion, uint64_t flags,
+                             hipDriverProcAddressQueryResult* symbolStatus) {
+  HIP_INIT_API(hipGetProcAddress, symbol, pfn, hipVersion, flags, symbolStatus);
+  HIP_RETURN(hipGetProcAddress_common(symbol, pfn, hipVersion, flags, symbolStatus));
+}
+
+hipError_t hipGetProcAddress_spt(const char* symbol, void** pfn, int hipVersion, uint64_t flags,
+                             hipDriverProcAddressQueryResult* symbolStatus) {
+  HIP_INIT_API(hipGetProcAddress, symbol, pfn, hipVersion, flags, symbolStatus);
+  flags = (flags == HIP_GET_PROC_ADDRESS_DEFAULT) ? HIP_GET_PROC_ADDRESS_PER_THREAD_DEFAULT_STREAM : flags;
+  HIP_RETURN(hipGetProcAddress_common(symbol, pfn, hipVersion, flags, symbolStatus));
 }
 
 }  // namespace hip

projects/clr/hipamd/src/hip_hcc.map.in

@@ -643,6 +643,7 @@ global:
     hipKernelGetLibrary;
     hipKernelGetName;
     hipOccupancyAvailableDynamicSMemPerBlock;
+    hipGetProcAddress_spt;
 local:
     *;
 } hip_7.1;

projects/clr/hipamd/src/hip_table_interface.cpp

@@ -1811,6 +1811,12 @@ extern "C" hipError_t hipGetProcAddress(const char* symbol, void** pfn, int hipV
   return hip::GetHipDispatchTable()->hipGetProcAddress_fn(symbol, pfn, hipVersion, flags,
                                                           symbolStatus);
 }
+extern "C" hipError_t hipGetProcAddress_spt(const char* symbol, void** pfn, int  hipVersion,
+                                            uint64_t flags,
+                                            hipDriverProcAddressQueryResult* symbolStatus) {
+  return hip::GetHipDispatchTable()->hipGetProcAddress_spt_fn(symbol, pfn, hipVersion, flags,
+                                                              symbolStatus);
+}
 hipError_t hipStreamBeginCaptureToGraph(hipStream_t stream, hipGraph_t graph,
                                         const hipGraphNode_t* dependencies,
                                         const hipGraphEdgeData* dependencyData,

projects/hip/include/hip/hip_runtime_api.h

@@ -604,6 +604,19 @@ typedef enum hipDeviceAttribute_t {
   // Extended attributes for vendors
 } hipDeviceAttribute_t;
 
+// Flags that can be used with hipGetProcAddress.
+/** Default flag. Equivalent to HIP_GET_PROC_ADDRESS_PER_THREAD_DEFAULT_STREAM if compiled with
+ *  -fgpu-default-stream=per-thread flag or HIP_API_PER_THREAD_DEFAULT_STREAM macro is
+ * defined.*/
+#define HIP_GET_PROC_ADDRESS_DEFAULT 0x0
+
+/** Search for all symbols except the corresponding per-thread versions.*/
+#define HIP_GET_PROC_ADDRESS_LEGACY_STREAM 0x1
+
+/** Search for all symbols including the per-thread versions. If a per-thread version cannot be
+ * found, returns the legacy version.*/
+#define HIP_GET_PROC_ADDRESS_PER_THREAD_DEFAULT_STREAM 0x2
+
 typedef enum hipDriverProcAddressQueryResult {
   HIP_GET_PROC_ADDRESS_SUCCESS = 0,
   HIP_GET_PROC_ADDRESS_SYMBOL_NOT_FOUND = 1,