AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

SWDEV-444021 - Implement hipGetFuncBySymbol

Przeglądaj źródła 
Change-Id: I7ef13d02c5b5c6ce2386ccb92b5602d005b35988


[ROCm/clr commit: 684fd60c8f]
This commit is contained in:
Satyanvesh Dittakavi
2024-02-15 09:42:48 +00:00
rodzic 2d8d4d5821
commit 15f01694f4
7 zmienionych plików z 61 dodań i 8 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp
+5 -2
Wyświetl plik
@@ -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;
};
projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h
+25 -1
Wyświetl plik
@@ -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 << ")";
projects/clr/hipamd/src/amdhip.def
+1
Wyświetl plik
@@ -462,3 +462,4 @@ hipGraphExecExternalSemaphoresWaitNodeSetParams
hipGraphAddNode
hipGraphInstantiateWithParams
hipStreamBeginCaptureToGraph
hipGetFuncBySymbol
projects/clr/hipamd/src/hip_api_trace.cpp
+6 -3
Wyświetl plik
@@ -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(<table>, <functor>, 8)
//
// HIP_ENFORCE_ABI_VERSIONING(<table>, 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
projects/clr/hipamd/src/hip_hcc.map.in
+7
Wyświetl plik
@@ -556,3 +556,10 @@ global:
local:
*;
} hip_6.0;
hip_6.2 {
global:
hipGetFuncBySymbol;
local:
*;
} hip_6.1;
projects/clr/hipamd/src/hip_module.cpp
+13 -1
Wyświetl plik
@@ -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) {
projects/clr/hipamd/src/hip_table_interface.cpp
+4 -1
Wyświetl plik
@@ -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);
}