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 0c138c6c80..e477ad8b67 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,7 +85,6 @@ THE SOFTWARE. // Driver Entry Point API #define hipGetDriverEntryPoint __HIP_API_SPT(hipGetDriverEntryPoint) -#define hipGetProcAddress __HIP_API_SPT(hipGetProcAddress) #endif #ifdef __cplusplus @@ -198,8 +197,6 @@ 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 9941eb8a79..80209d7a9c 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,8 +1111,6 @@ 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 { @@ -1698,7 +1696,6 @@ 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 6965222310..7407d2d7bd 100644 --- a/projects/clr/hipamd/src/amdhip.def +++ b/projects/clr/hipamd/src/amdhip.def @@ -521,4 +521,3 @@ 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 9b495f573b..a4d73f364b 100644 --- a/projects/clr/hipamd/src/hip_api_trace.cpp +++ b/projects/clr/hipamd/src/hip_api_trace.cpp @@ -790,8 +790,6 @@ 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, @@ -1379,7 +1377,6 @@ 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; @@ -2107,14 +2104,13 @@ 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, 507) +HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 506) 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 2869822c14..07708033fc 100644 --- a/projects/clr/hipamd/src/hip_device.cpp +++ b/projects/clr/hipamd/src/hip_device.cpp @@ -766,83 +766,55 @@ hipError_t hipGetDevicePropertiesR0000(hipDeviceProp_tR0000* prop, int device) { HIP_RETURN(hipSuccess); } -hipError_t hipGetProcAddress_common(const char* symbol, void** pfn, int hipVersion, uint64_t flags, +hipError_t hipGetProcAddress(const char* symbol, void** pfn, int hipVersion, uint64_t flags, hipDriverProcAddressQueryResult* symbolStatus) { - if (symbol == nullptr || std::string_view{symbol}.empty() || pfn == nullptr) { - return hipErrorInvalidValue; - } + HIP_INIT_API(hipGetProcAddress, symbol, pfn, hipVersion, flags, symbolStatus); + std::string symbolString = symbol; - 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; + if (symbol == nullptr || symbolString == "" || pfn == nullptr) { + HIP_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; } - return hipSuccess; + HIP_RETURN(hipSuccess); } else if (symbolString == "hipAmdFileWrite") { *pfn = reinterpret_cast(&hipAmdFileWrite); if (symbolStatus != nullptr) { *symbolStatus = HIP_GET_PROC_ADDRESS_SUCCESS; } - return hipSuccess; + HIP_RETURN(hipSuccess); } void* handle = hip::PlatformState::instance().getDynamicLibraryHandle(); if (handle == nullptr) { - return hipErrorInvalidValue; - } - - if (checkSpt) { - symbolString += "_spt"; + HIP_RETURN(hipErrorInvalidValue); } *pfn = amd::Os::getSymbol(handle, symbolString.c_str()); - 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; + if (!(*pfn)) { + if (symbolStatus != nullptr) { + *symbolStatus = HIP_GET_PROC_ADDRESS_SYMBOL_NOT_FOUND; } + HIP_RETURN(hipErrorInvalidValue); } if (symbolStatus != nullptr) { *symbolStatus = HIP_GET_PROC_ADDRESS_SUCCESS; } - - 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)); + HIP_RETURN(hipSuccess); } } // namespace hip diff --git a/projects/clr/hipamd/src/hip_hcc.map.in b/projects/clr/hipamd/src/hip_hcc.map.in index fc915143e8..f26fc61be1 100644 --- a/projects/clr/hipamd/src/hip_hcc.map.in +++ b/projects/clr/hipamd/src/hip_hcc.map.in @@ -643,7 +643,6 @@ 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 3e378fb4d2..2165924089 100644 --- a/projects/clr/hipamd/src/hip_table_interface.cpp +++ b/projects/clr/hipamd/src/hip_table_interface.cpp @@ -1811,12 +1811,6 @@ 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 1728b22b1b..c1226c389c 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -604,19 +604,6 @@ 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,