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 4685890336..96031d3b48 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 @@ -1,5 +1,5 @@ /* - Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + Copyright (c) 2023 - 2024 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -61,7 +61,7 @@ // - Reset any of the *_STEP_VERSION defines to zero if the corresponding *_MAJOR_VERSION increases #define HIP_API_TABLE_STEP_VERSION 0 #define HIP_COMPILER_API_TABLE_STEP_VERSION 0 -#define HIP_RUNTIME_API_TABLE_STEP_VERSION 2 +#define HIP_RUNTIME_API_TABLE_STEP_VERSION 3 // HIP API interface typedef hipError_t (*t___hipPopCallConfiguration)(dim3* gridDim, dim3* blockDim, size_t* sharedMem, @@ -952,6 +952,8 @@ typedef hipError_t (*t_hipStreamBeginCaptureToGraph)(hipStream_t stream, hipGrap const hipGraphEdgeData* dependencyData, size_t numDependencies, hipStreamCaptureMode mode); +typedef hipError_t (*t_hipGetFuncBySymbol)(hipFunction_t* functionPtr, const void* symbolPtr); + // HIP Compiler dispatch table struct HipCompilerDispatchTable { size_t size; @@ -1417,4 +1419,5 @@ struct HipDispatchTable { t_hipTexRefGetArray hipTexRefGetArray_fn; t_hipGetProcAddress hipGetProcAddress_fn; t_hipStreamBeginCaptureToGraph hipStreamBeginCaptureToGraph_fn; + t_hipGetFuncBySymbol hipGetFuncBySymbol_fn; }; 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 55b2679050..fcebe3d9f6 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 @@ -409,7 +409,8 @@ enum hip_api_id_t { HIP_API_ID_hipTexRefGetArray = 389, HIP_API_ID_hipTexRefGetBorderColor = 390, HIP_API_ID_hipStreamBeginCaptureToGraph = 391, - HIP_API_ID_LAST = 391, + HIP_API_ID_hipGetFuncBySymbol = 392, + HIP_API_ID_LAST = 392, HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice), HIP_API_ID_hipGetDeviceProperties = HIP_API_ID_CONCAT(HIP_API_ID_,hipGetDeviceProperties), @@ -563,6 +564,7 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipGetDevicePropertiesR0000: return "hipGetDevicePropertiesR0000"; case HIP_API_ID_hipGetDevicePropertiesR0600: return "hipGetDevicePropertiesR0600"; case HIP_API_ID_hipGetErrorString: return "hipGetErrorString"; + case HIP_API_ID_hipGetFuncBySymbol: return "hipGetFuncBySymbol"; case HIP_API_ID_hipGetLastError: return "hipGetLastError"; case HIP_API_ID_hipGetMipmappedArrayLevel: return "hipGetMipmappedArrayLevel"; case HIP_API_ID_hipGetProcAddress: return "hipGetProcAddress"; @@ -957,6 +959,7 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipGetDevicePropertiesR0000", name) == 0) return HIP_API_ID_hipGetDevicePropertiesR0000; if (strcmp("hipGetDevicePropertiesR0600", name) == 0) return HIP_API_ID_hipGetDevicePropertiesR0600; 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; if (strcmp("hipGetMipmappedArrayLevel", name) == 0) return HIP_API_ID_hipGetMipmappedArrayLevel; if (strcmp("hipGetProcAddress", name) == 0) return HIP_API_ID_hipGetProcAddress; @@ -1786,6 +1789,11 @@ typedef struct hip_api_data_s { hipDeviceProp_tR0600 prop__val; int deviceId; } hipGetDevicePropertiesR0600; + struct { + hipFunction_t* functionPtr; + hipFunction_t functionPtr__val; + const void* symbolPtr; + } hipGetFuncBySymbol; struct { hipArray_t* levelArray; hipArray_t levelArray__val; @@ -4093,6 +4101,11 @@ typedef struct hip_api_data_s { // hipGetErrorString[] #define INIT_hipGetErrorString_CB_ARGS_DATA(cb_data) { \ }; +// hipGetFuncBySymbol[('hipFunction_t*', 'functionPtr'), ('const void*', 'symbolPtr')] +#define INIT_hipGetFuncBySymbol_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipGetFuncBySymbol.functionPtr = (hipFunction_t*)functionPtr; \ + cb_data.args.hipGetFuncBySymbol.symbolPtr = (const void*)symbolPtr; \ +}; // hipGetLastError[] #define INIT_hipGetLastError_CB_ARGS_DATA(cb_data) { \ }; @@ -6344,6 +6357,10 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipGetErrorString[] case HIP_API_ID_hipGetErrorString: break; +// hipGetFuncBySymbol[('hipFunction_t*', 'functionPtr'), ('const void*', 'symbolPtr')] + case HIP_API_ID_hipGetFuncBySymbol: + if (data->args.hipGetFuncBySymbol.functionPtr) data->args.hipGetFuncBySymbol.functionPtr__val = *(data->args.hipGetFuncBySymbol.functionPtr); + break; // hipGetLastError[] case HIP_API_ID_hipGetLastError: break; @@ -8225,6 +8242,13 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << "hipGetErrorString("; oss << ")"; break; + case HIP_API_ID_hipGetFuncBySymbol: + oss << "hipGetFuncBySymbol("; + if (data->args.hipGetFuncBySymbol.functionPtr == NULL) oss << "functionPtr=NULL"; + else { oss << "functionPtr="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetFuncBySymbol.functionPtr__val); } + oss << ", symbolPtr="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetFuncBySymbol.symbolPtr); + oss << ")"; + break; case HIP_API_ID_hipGetLastError: oss << "hipGetLastError("; oss << ")"; diff --git a/projects/clr/hipamd/src/amdhip.def b/projects/clr/hipamd/src/amdhip.def index d9764ab4af..cfc3ae6b5d 100644 --- a/projects/clr/hipamd/src/amdhip.def +++ b/projects/clr/hipamd/src/amdhip.def @@ -462,3 +462,4 @@ hipGraphExecExternalSemaphoresWaitNodeSetParams hipGraphAddNode hipGraphInstantiateWithParams hipStreamBeginCaptureToGraph +hipGetFuncBySymbol diff --git a/projects/clr/hipamd/src/hip_api_trace.cpp b/projects/clr/hipamd/src/hip_api_trace.cpp index 7a6d69bc22..1479ae1507 100644 --- a/projects/clr/hipamd/src/hip_api_trace.cpp +++ b/projects/clr/hipamd/src/hip_api_trace.cpp @@ -1,5 +1,5 @@ /* - Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + Copyright (c) 2023 - 2024 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -767,6 +767,7 @@ hipError_t hipStreamBeginCaptureToGraph(hipStream_t stream, hipGraph_t graph, const hipGraphNode_t* dependencies, const hipGraphEdgeData* dependencyData, size_t numDependencies, hipStreamCaptureMode mode); +hipError_t hipGetFuncBySymbol(hipFunction_t* functionPtr, const void* symbolPtr); } // namespace hip namespace hip { @@ -1242,6 +1243,7 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) { ptrDispatchTable->hipTexRefGetArray_fn = hip::hipTexRefGetArray; ptrDispatchTable->hipGetProcAddress_fn = hip::hipGetProcAddress; ptrDispatchTable->hipStreamBeginCaptureToGraph_fn = hip::hipStreamBeginCaptureToGraph; + ptrDispatchTable->hipGetFuncBySymbol_fn = hip::hipGetFuncBySymbol; } #if HIP_ROCPROFILER_REGISTER > 0 @@ -1803,6 +1805,7 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipTexRefGetBorderColor_fn, 440) HIP_ENFORCE_ABI(HipDispatchTable, hipTexRefGetArray_fn, 441) HIP_ENFORCE_ABI(HipDispatchTable, hipGetProcAddress_fn, 442) HIP_ENFORCE_ABI(HipDispatchTable, hipStreamBeginCaptureToGraph_fn, 443); +HIP_ENFORCE_ABI(HipDispatchTable, hipGetFuncBySymbol_fn, 444); // if HIP_ENFORCE_ABI entries are added for each new function pointer in the table, the number below @@ -1811,9 +1814,9 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipStreamBeginCaptureToGraph_fn, 443); // HIP_ENFORCE_ABI(, , 8) // // HIP_ENFORCE_ABI_VERSIONING(
, 9) <- 8 + 1 = 9 -HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 444) +HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 445) -static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 2, +static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 3, "If you get this error, add new HIP_ENFORCE_ABI(...) code for the new function " "pointers and then update this check so it is true"); #endif diff --git a/projects/clr/hipamd/src/hip_hcc.map.in b/projects/clr/hipamd/src/hip_hcc.map.in index 0ce6ffd2ab..9bfa41d49a 100644 --- a/projects/clr/hipamd/src/hip_hcc.map.in +++ b/projects/clr/hipamd/src/hip_hcc.map.in @@ -556,3 +556,10 @@ global: local: *; } hip_6.0; + +hip_6.2 { +global: + hipGetFuncBySymbol; +local: + *; +} hip_6.1; diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index 71ae362449..dfb082dca7 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/src/hip_module.cpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2015 - 2024 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -656,6 +656,18 @@ hipError_t hipModuleLaunchCooperativeKernelMultiDevice(hipFunctionLaunchParams* } +hipError_t hipGetFuncBySymbol(hipFunction_t* functionPtr, const void* symbolPtr) { + HIP_INIT_API(hipGetFuncBySymbol, functionPtr, symbolPtr); + + hipError_t hip_error = PlatformState::instance().getStatFunc(functionPtr, + symbolPtr, ihipGetDevice()); + + if ((hip_error != hipSuccess) || (functionPtr == nullptr)) { + HIP_RETURN(hipErrorInvalidDeviceFunction); + } + HIP_RETURN(hipSuccess); +} + hipError_t hipLaunchKernel_common(const void* hostFunction, dim3 gridDim, dim3 blockDim, void** args, size_t sharedMemBytes, hipStream_t stream) { diff --git a/projects/clr/hipamd/src/hip_table_interface.cpp b/projects/clr/hipamd/src/hip_table_interface.cpp index 5369bca662..7173691e44 100644 --- a/projects/clr/hipamd/src/hip_table_interface.cpp +++ b/projects/clr/hipamd/src/hip_table_interface.cpp @@ -1,5 +1,5 @@ /* - Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + Copyright (c) 2023 - 2024 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -1740,3 +1740,6 @@ hipError_t hipStreamBeginCaptureToGraph(hipStream_t stream, hipGraph_t graph, return hip::GetHipDispatchTable()->hipStreamBeginCaptureToGraph_fn( stream, graph, dependencies, dependencyData, numDependencies, mode); } +hipError_t hipGetFuncBySymbol(hipFunction_t* functionPtr, const void* symbolPtr) { + return hip::GetHipDispatchTable()->hipGetFuncBySymbol_fn(functionPtr, symbolPtr); +}