Skip to content

Commit 684fd60

Browse files
satyanveshdsatyanveshd
committed
SWDEV-444021 - Implement hipGetFuncBySymbol
Change-Id: I7ef13d02c5b5c6ce2386ccb92b5602d005b35988
1 parent 43e2bb6 commit 684fd60

File tree

7 files changed

+61
-8
lines changed

7 files changed

+61
-8
lines changed

‎hipamd/include/hip/amd_detail/hip_api_trace.hpp‎

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
2+
Copyright (c) 2023 - 2024 Advanced Micro Devices, Inc. All rights reserved.
33
44
Permission is hereby granted, free of charge, to any person obtaining a copy
55
of this software and associated documentation files (the "Software"), to deal
@@ -61,7 +61,7 @@
6161
// - Reset any of the *_STEP_VERSION defines to zero if the corresponding *_MAJOR_VERSION increases
6262
#define HIP_API_TABLE_STEP_VERSION 0
6363
#define HIP_COMPILER_API_TABLE_STEP_VERSION 0
64-
#define HIP_RUNTIME_API_TABLE_STEP_VERSION 2
64+
#define HIP_RUNTIME_API_TABLE_STEP_VERSION 3
6565

6666
// HIP API interface
6767
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
952952
const hipGraphEdgeData* dependencyData,
953953
size_t numDependencies,
954954
hipStreamCaptureMode mode);
955+
typedef hipError_t (*t_hipGetFuncBySymbol)(hipFunction_t* functionPtr, const void* symbolPtr);
956+
955957
// HIP Compiler dispatch table
956958
struct HipCompilerDispatchTable {
957959
size_t size;
@@ -1417,4 +1419,5 @@ struct HipDispatchTable {
14171419
t_hipTexRefGetArray hipTexRefGetArray_fn;
14181420
t_hipGetProcAddress hipGetProcAddress_fn;
14191421
t_hipStreamBeginCaptureToGraph hipStreamBeginCaptureToGraph_fn;
1422+
t_hipGetFuncBySymbol hipGetFuncBySymbol_fn;
14201423
};

‎hipamd/include/hip/amd_detail/hip_prof_str.h‎

Lines changed: 25 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -409,7 +409,8 @@ enum hip_api_id_t {
409409
HIP_API_ID_hipTexRefGetArray = 389,
410410
HIP_API_ID_hipTexRefGetBorderColor = 390,
411411
HIP_API_ID_hipStreamBeginCaptureToGraph = 391,
412-
HIP_API_ID_LAST = 391,
412+
HIP_API_ID_hipGetFuncBySymbol = 392,
413+
HIP_API_ID_LAST = 392,
413414

414415
HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice),
415416
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) {
563564
case HIP_API_ID_hipGetDevicePropertiesR0000: return "hipGetDevicePropertiesR0000";
564565
case HIP_API_ID_hipGetDevicePropertiesR0600: return "hipGetDevicePropertiesR0600";
565566
case HIP_API_ID_hipGetErrorString: return "hipGetErrorString";
567+
case HIP_API_ID_hipGetFuncBySymbol: return "hipGetFuncBySymbol";
566568
case HIP_API_ID_hipGetLastError: return "hipGetLastError";
567569
case HIP_API_ID_hipGetMipmappedArrayLevel: return "hipGetMipmappedArrayLevel";
568570
case HIP_API_ID_hipGetProcAddress: return "hipGetProcAddress";
@@ -957,6 +959,7 @@ static inline uint32_t hipApiIdByName(const char* name) {
957959
if (strcmp("hipGetDevicePropertiesR0000", name) == 0) return HIP_API_ID_hipGetDevicePropertiesR0000;
958960
if (strcmp("hipGetDevicePropertiesR0600", name) == 0) return HIP_API_ID_hipGetDevicePropertiesR0600;
959961
if (strcmp("hipGetErrorString", name) == 0) return HIP_API_ID_hipGetErrorString;
962+
if (strcmp("hipGetFuncBySymbol", name) == 0) return HIP_API_ID_hipGetFuncBySymbol;
960963
if (strcmp("hipGetLastError", name) == 0) return HIP_API_ID_hipGetLastError;
961964
if (strcmp("hipGetMipmappedArrayLevel", name) == 0) return HIP_API_ID_hipGetMipmappedArrayLevel;
962965
if (strcmp("hipGetProcAddress", name) == 0) return HIP_API_ID_hipGetProcAddress;
@@ -1786,6 +1789,11 @@ typedef struct hip_api_data_s {
17861789
hipDeviceProp_tR0600 prop__val;
17871790
int deviceId;
17881791
} hipGetDevicePropertiesR0600;
1792+
struct {
1793+
hipFunction_t* functionPtr;
1794+
hipFunction_t functionPtr__val;
1795+
const void* symbolPtr;
1796+
} hipGetFuncBySymbol;
17891797
struct {
17901798
hipArray_t* levelArray;
17911799
hipArray_t levelArray__val;
@@ -4093,6 +4101,11 @@ typedef struct hip_api_data_s {
40934101
// hipGetErrorString[]
40944102
#define INIT_hipGetErrorString_CB_ARGS_DATA(cb_data) { \
40954103
};
4104+
// hipGetFuncBySymbol[('hipFunction_t*', 'functionPtr'), ('const void*', 'symbolPtr')]
4105+
#define INIT_hipGetFuncBySymbol_CB_ARGS_DATA(cb_data) { \
4106+
cb_data.args.hipGetFuncBySymbol.functionPtr = (hipFunction_t*)functionPtr; \
4107+
cb_data.args.hipGetFuncBySymbol.symbolPtr = (const void*)symbolPtr; \
4108+
};
40964109
// hipGetLastError[]
40974110
#define INIT_hipGetLastError_CB_ARGS_DATA(cb_data) { \
40984111
};
@@ -6344,6 +6357,10 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
63446357
// hipGetErrorString[]
63456358
case HIP_API_ID_hipGetErrorString:
63466359
break;
6360+
// hipGetFuncBySymbol[('hipFunction_t*', 'functionPtr'), ('const void*', 'symbolPtr')]
6361+
case HIP_API_ID_hipGetFuncBySymbol:
6362+
if (data->args.hipGetFuncBySymbol.functionPtr) data->args.hipGetFuncBySymbol.functionPtr__val = *(data->args.hipGetFuncBySymbol.functionPtr);
6363+
break;
63476364
// hipGetLastError[]
63486365
case HIP_API_ID_hipGetLastError:
63496366
break;
@@ -8225,6 +8242,13 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
82258242
oss << "hipGetErrorString(";
82268243
oss << ")";
82278244
break;
8245+
case HIP_API_ID_hipGetFuncBySymbol:
8246+
oss << "hipGetFuncBySymbol(";
8247+
if (data->args.hipGetFuncBySymbol.functionPtr == NULL) oss << "functionPtr=NULL";
8248+
else { oss << "functionPtr="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetFuncBySymbol.functionPtr__val); }
8249+
oss << ", symbolPtr="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetFuncBySymbol.symbolPtr);
8250+
oss << ")";
8251+
break;
82288252
case HIP_API_ID_hipGetLastError:
82298253
oss << "hipGetLastError(";
82308254
oss << ")";

‎hipamd/src/amdhip.def‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -462,3 +462,4 @@ hipGraphExecExternalSemaphoresWaitNodeSetParams
462462
hipGraphAddNode
463463
hipGraphInstantiateWithParams
464464
hipStreamBeginCaptureToGraph
465+
hipGetFuncBySymbol

‎hipamd/src/hip_api_trace.cpp‎

Lines changed: 6 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
2+
Copyright (c) 2023 - 2024 Advanced Micro Devices, Inc. All rights reserved.
33
44
Permission is hereby granted, free of charge, to any person obtaining a copy
55
of this software and associated documentation files (the "Software"), to deal
@@ -767,6 +767,7 @@ hipError_t hipStreamBeginCaptureToGraph(hipStream_t stream, hipGraph_t graph,
767767
const hipGraphNode_t* dependencies,
768768
const hipGraphEdgeData* dependencyData,
769769
size_t numDependencies, hipStreamCaptureMode mode);
770+
hipError_t hipGetFuncBySymbol(hipFunction_t* functionPtr, const void* symbolPtr);
770771
} // namespace hip
771772

772773
namespace hip {
@@ -1242,6 +1243,7 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) {
12421243
ptrDispatchTable->hipTexRefGetArray_fn = hip::hipTexRefGetArray;
12431244
ptrDispatchTable->hipGetProcAddress_fn = hip::hipGetProcAddress;
12441245
ptrDispatchTable->hipStreamBeginCaptureToGraph_fn = hip::hipStreamBeginCaptureToGraph;
1246+
ptrDispatchTable->hipGetFuncBySymbol_fn = hip::hipGetFuncBySymbol;
12451247
}
12461248

12471249
#if HIP_ROCPROFILER_REGISTER > 0
@@ -1803,6 +1805,7 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipTexRefGetBorderColor_fn, 440)
18031805
HIP_ENFORCE_ABI(HipDispatchTable, hipTexRefGetArray_fn, 441)
18041806
HIP_ENFORCE_ABI(HipDispatchTable, hipGetProcAddress_fn, 442)
18051807
HIP_ENFORCE_ABI(HipDispatchTable, hipStreamBeginCaptureToGraph_fn, 443);
1808+
HIP_ENFORCE_ABI(HipDispatchTable, hipGetFuncBySymbol_fn, 444);
18061809

18071810

18081811
// 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);
18111814
// HIP_ENFORCE_ABI(<table>, <functor>, 8)
18121815
//
18131816
// HIP_ENFORCE_ABI_VERSIONING(<table>, 9) <- 8 + 1 = 9
1814-
HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 444)
1817+
HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 445)
18151818

1816-
static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 2,
1819+
static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 3,
18171820
"If you get this error, add new HIP_ENFORCE_ABI(...) code for the new function "
18181821
"pointers and then update this check so it is true");
18191822
#endif

‎hipamd/src/hip_hcc.map.in‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -556,3 +556,10 @@ global:
556556
local:
557557
*;
558558
} hip_6.0;
559+
560+
hip_6.2 {
561+
global:
562+
hipGetFuncBySymbol;
563+
local:
564+
*;
565+
} hip_6.1;

‎hipamd/src/hip_module.cpp‎

Lines changed: 13 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
/* Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc.
1+
/* Copyright (c) 2015 - 2024 Advanced Micro Devices, Inc.
22
33
Permission is hereby granted, free of charge, to any person obtaining a copy
44
of this software and associated documentation files (the "Software"), to deal
@@ -656,6 +656,18 @@ hipError_t hipModuleLaunchCooperativeKernelMultiDevice(hipFunctionLaunchParams*
656656

657657
}
658658

659+
hipError_t hipGetFuncBySymbol(hipFunction_t* functionPtr, const void* symbolPtr) {
660+
HIP_INIT_API(hipGetFuncBySymbol, functionPtr, symbolPtr);
661+
662+
hipError_t hip_error = PlatformState::instance().getStatFunc(functionPtr,
663+
symbolPtr, ihipGetDevice());
664+
665+
if ((hip_error != hipSuccess) || (functionPtr == nullptr)) {
666+
HIP_RETURN(hipErrorInvalidDeviceFunction);
667+
}
668+
HIP_RETURN(hipSuccess);
669+
}
670+
659671
hipError_t hipLaunchKernel_common(const void* hostFunction, dim3 gridDim, dim3 blockDim,
660672
void** args, size_t sharedMemBytes,
661673
hipStream_t stream) {

‎hipamd/src/hip_table_interface.cpp‎

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
2+
Copyright (c) 2023 - 2024 Advanced Micro Devices, Inc. All rights reserved.
33
44
Permission is hereby granted, free of charge, to any person obtaining a copy
55
of this software and associated documentation files (the "Software"), to deal
@@ -1740,3 +1740,6 @@ hipError_t hipStreamBeginCaptureToGraph(hipStream_t stream, hipGraph_t graph,
17401740
return hip::GetHipDispatchTable()->hipStreamBeginCaptureToGraph_fn(
17411741
stream, graph, dependencies, dependencyData, numDependencies, mode);
17421742
}
1743+
hipError_t hipGetFuncBySymbol(hipFunction_t* functionPtr, const void* symbolPtr) {
1744+
return hip::GetHipDispatchTable()->hipGetFuncBySymbol_fn(functionPtr, symbolPtr);
1745+
}

0 commit comments

Comments
 (0)