AI/rocm-systems
2
0
Fork 0
Code Issues Pull-Requests Actions Pakete Projekte Releases Wiki Aktivität

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

Quellcode durchsuchen
Dieser Commit ist enthalten in:
Todd tiantuo Li
2025-11-10 23:29:40 -08:00
committet von GitHub
Ursprung 683a63d9ec
Commit cf536a8c1a
14 geänderte Dateien mit 111 neuen und 17 gelöschten Zeilen
Patch-Datei herunterladen Vergleichs-Datei herunterladen Alle Dateien ausklappen Alle Dateien einklappen
projects/clr/CHANGELOG.md
+4
Datei anzeigen
@@ -17,6 +17,10 @@ Full documentation for HIP is available at [rocm.docs.amd.com](https://rocm.docs
- `hipLibraryGetKernelCount` gets kernel count in library
- `hipStreamCopyAttributes` copies attributes from source stream to destination stream
- `hipOccupancyAvailableDynamicSMemPerBlock` Returns dynamic shared memory available per block when launching numBlocks blocks on CU.
* Support for the following flags in `hipGetProcAddress`, enabling searching for the per-thread version symbols.
- `HIP_GET_PROC_ADDRESS_DEFAULT`
- `HIP_GET_PROC_ADDRESS_LEGACY_STREAM`
- `HIP_GET_PROC_ADDRESS_PER_THREAD_DEFAULT_STREAM`
## HIP 7.1 for ROCm 7.1
projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h
+3
Datei anzeigen
@@ -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
+7 -2
Datei anzeigen
@@ -63,7 +63,7 @@
#define HIP_API_TABLE_STEP_VERSION 0
#define HIP_COMPILER_API_TABLE_STEP_VERSION 0
#define HIP_TOOLS_API_TABLE_STEP_VERSION 0
#define HIP_RUNTIME_API_TABLE_STEP_VERSION 18
#define HIP_RUNTIME_API_TABLE_STEP_VERSION 19
// HIP API interface
// HIP compiler dispatch functions
@@ -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 {
@@ -1697,8 +1699,11 @@ struct HipDispatchTable {
// HIP_RUNTIME_API_TABLE_STEP_VERSION == 18
t_hipOccupancyAvailableDynamicSMemPerBlock hipOccupancyAvailableDynamicSMemPerBlock_fn;
// DO NOT EDIT ABOVE!
// HIP_RUNTIME_API_TABLE_STEP_VERSION == 19
t_hipGetProcAddress_spt hipGetProcAddress_spt_fn;
// DO NOT EDIT ABOVE!
// HIP_RUNTIME_API_TABLE_STEP_VERSION == 20
// ******************************************************************************************* //
//
projects/clr/hipamd/src/amdhip.def
+1
Datei anzeigen
@@ -521,3 +521,4 @@ hipLibraryEnumerateKernels
hipKernelGetLibrary
hipKernelGetName
hipOccupancyAvailableDynamicSMemPerBlock
hipGetProcAddress_spt
projects/clr/hipamd/src/hip_api_trace.cpp
+7 -2
Datei anzeigen
@@ -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,15 +2107,17 @@ 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_RUNTIME_API_TABLE_STEP_VERSION == 19
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,
static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 19,
"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
projects/clr/hipamd/src/hip_device.cpp
+41 -13
Datei anzeigen
@@ -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
+1
Datei anzeigen
@@ -643,6 +643,7 @@ global:
hipKernelGetLibrary;
hipKernelGetName;
hipOccupancyAvailableDynamicSMemPerBlock;
hipGetProcAddress_spt;
local:
*;
} hip_7.1;
projects/clr/hipamd/src/hip_table_interface.cpp
+6
Datei anzeigen
@@ -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
+13
Datei anzeigen
@@ -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,
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/enum_string.hpp
+5
Datei anzeigen
@@ -998,6 +998,9 @@ ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipKernelGetName)
#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 18
ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipOccupancyAvailableDynamicSMemPerBlock)
#endif
#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 19
ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipGetProcAddress_spt)
#endif
#if HIP_RUNTIME_API_TABLE_STEP_VERSION == 0
static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 442);
#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 1
@@ -1036,6 +1039,8 @@ static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 502);
static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 505);
#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 18
static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 506);
#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 19
static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 507);
#else
# if !defined(ROCPROFILER_UNSAFE_NO_VERSION_CHECK) && \
(defined(ROCPROFILER_CI) && ROCPROFILER_CI > 0)
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h
+10
Datei anzeigen
@@ -3358,6 +3358,16 @@ typedef union rocprofiler_hip_api_args_t
int blockSize;
} hipOccupancyAvailableDynamicSMemPerBlock;
#endif
#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 19
struct
{
const char* symbol;
void** pfn;
int hipVersion;
uint64_t flags;
hipDriverProcAddressQueryResult* symbolStatus;
} hipGetProcAddress_spt;
#endif
} rocprofiler_hip_api_args_t;
ROCPROFILER_EXTERN_C_FINI
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/runtime_api_id.h
+3
Datei anzeigen
@@ -569,6 +569,9 @@ typedef enum rocprofiler_hip_runtime_api_id_t // NOLINT(performance-enum-size)
#endif
#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 18
ROCPROFILER_HIP_RUNTIME_API_ID_hipOccupancyAvailableDynamicSMemPerBlock,
#endif
#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 19
ROCPROFILER_HIP_RUNTIME_API_ID_hipGetProcAddress_spt,
#endif
ROCPROFILER_HIP_RUNTIME_API_ID_LAST,
} rocprofiler_hip_runtime_api_id_t;
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/abi.cpp
+6
Datei anzeigen
@@ -615,6 +615,10 @@ ROCP_SDK_ENFORCE_ABI(::HipDispatchTable, hipKernelGetName_fn, 504);
ROCP_SDK_ENFORCE_ABI(::HipDispatchTable, hipOccupancyAvailableDynamicSMemPerBlock_fn, 505);
#endif
#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 19
ROCP_SDK_ENFORCE_ABI(::HipDispatchTable, hipGetProcAddress_spt_fn, 506);
#endif
#if HIP_RUNTIME_API_TABLE_STEP_VERSION == 0
ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 442)
#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 1
@@ -653,6 +657,8 @@ ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 502)
ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 505)
#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 18
ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 506)
#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 19
ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 507)
#else
INTERNAL_CI_ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 0)
#endif
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp
+4
Datei anzeigen
@@ -641,6 +641,10 @@ HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNT
#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 18
HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNTIME_API_ID_hipOccupancyAvailableDynamicSMemPerBlock, hipOccupancyAvailableDynamicSMemPerBlock, hipOccupancyAvailableDynamicSMemPerBlock_fn, dynamicSmemSize, f, numBlocks, blockSize);
#endif
#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 19
HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNTIME_API_ID_hipGetProcAddress_spt, hipGetProcAddress_spt, hipGetProcAddress_spt_fn, symbol, pfn, hipVersion, flags, symbolStatus);
#endif
// clang-format on
#else