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 e477ad8b67..0c138c6c80 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 @@ -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 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 80209d7a9c..9941eb8a79 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 @@ -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 diff --git a/projects/clr/hipamd/src/amdhip.def b/projects/clr/hipamd/src/amdhip.def index 7407d2d7bd..6965222310 100644 --- a/projects/clr/hipamd/src/amdhip.def +++ b/projects/clr/hipamd/src/amdhip.def @@ -521,3 +521,4 @@ hipLibraryEnumerateKernels hipKernelGetLibrary hipKernelGetName hipOccupancyAvailableDynamicSMemPerBlock +hipGetProcAddress_spt diff --git a/projects/clr/hipamd/src/hip_api_trace.cpp b/projects/clr/hipamd/src/hip_api_trace.cpp index a4d73f364b..9b495f573b 100644 --- a/projects/clr/hipamd/src/hip_api_trace.cpp +++ b/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(, , 8) // // HIP_ENFORCE_ABI_VERSIONING(
, 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 " diff --git a/projects/clr/hipamd/src/hip_device.cpp b/projects/clr/hipamd/src/hip_device.cpp index 07708033fc..2869822c14 100644 --- a/projects/clr/hipamd/src/hip_device.cpp +++ b/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(&hipAmdFileRead); if (symbolStatus != nullptr) { *symbolStatus = HIP_GET_PROC_ADDRESS_SUCCESS; } - HIP_RETURN(hipSuccess); + return hipSuccess; } else if (symbolString == "hipAmdFileWrite") { *pfn = reinterpret_cast(&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 diff --git a/projects/clr/hipamd/src/hip_hcc.map.in b/projects/clr/hipamd/src/hip_hcc.map.in index f26fc61be1..fc915143e8 100644 --- a/projects/clr/hipamd/src/hip_hcc.map.in +++ b/projects/clr/hipamd/src/hip_hcc.map.in @@ -643,6 +643,7 @@ global: hipKernelGetLibrary; hipKernelGetName; hipOccupancyAvailableDynamicSMemPerBlock; + hipGetProcAddress_spt; local: *; } hip_7.1; diff --git a/projects/clr/hipamd/src/hip_table_interface.cpp b/projects/clr/hipamd/src/hip_table_interface.cpp index 2165924089..3e378fb4d2 100644 --- a/projects/clr/hipamd/src/hip_table_interface.cpp +++ b/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, diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index c1226c389c..1728b22b1b 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/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,