From 07dd4c85e7e7309ad4cc4778e12f2404a0b66dab Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi <53337087+satyanveshd@users.noreply.github.com> Date: Wed, 12 Nov 2025 14:09:26 +0530 Subject: [PATCH] SWDEV-546308 - Implement hipKernelGetParamInfo API (#1783) --- projects/clr/CHANGELOG.md | 1 + .../include/hip/amd_detail/hip_api_trace.hpp | 9 +++- .../include/hip/amd_detail/hip_prof_str.h | 35 +++++++++++++- projects/clr/hipamd/src/amdhip.def | 1 + projects/clr/hipamd/src/hip_api_trace.cpp | 9 +++- projects/clr/hipamd/src/hip_hcc.map.in | 1 + projects/clr/hipamd/src/hip_library.cpp | 26 ++++++++++ .../clr/hipamd/src/hip_table_interface.cpp | 5 ++ .../catch/unit/library/loadlib_co.cc | 48 ++++++++++++++++++- projects/hip/include/hip/hip_runtime_api.h | 13 +++++ .../nvidia_detail/nvidia_hip_runtime_api.h | 7 ++- .../rocprofiler-sdk/cxx/enum_string.hpp | 5 ++ .../include/rocprofiler-sdk/hip/api_args.h | 9 ++++ .../rocprofiler-sdk/hip/runtime_api_id.h | 3 ++ .../source/lib/rocprofiler-sdk/hip/abi.cpp | 6 +++ .../lib/rocprofiler-sdk/hip/hip.def.cpp | 5 ++ 16 files changed, 176 insertions(+), 7 deletions(-) diff --git a/projects/clr/CHANGELOG.md b/projects/clr/CHANGELOG.md index 8ce55ac982..755f5ef4fd 100644 --- a/projects/clr/CHANGELOG.md +++ b/projects/clr/CHANGELOG.md @@ -17,6 +17,7 @@ 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. + - `hipKernelGetParamInfo` returns the offset and size of a kernel parameter * 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` 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 e5729afa5a..f1864eaca4 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 @@ -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 19 +#define HIP_RUNTIME_API_TABLE_STEP_VERSION 20 // HIP API interface // HIP compiler dispatch functions @@ -1114,6 +1114,8 @@ 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); +typedef hipError_t (*t_hipKernelGetParamInfo)(hipKernel_t kernel, size_t paramIndex, + size_t* paramOffset, size_t* paramSize); // HIP Compiler dispatch table struct HipCompilerDispatchTable { // HIP_COMPILER_API_TABLE_STEP_VERSION == 0 @@ -1702,8 +1704,11 @@ struct HipDispatchTable { // HIP_RUNTIME_API_TABLE_STEP_VERSION == 19 t_hipGetProcAddress_spt hipGetProcAddress_spt_fn; - // DO NOT EDIT ABOVE! // HIP_RUNTIME_API_TABLE_STEP_VERSION == 20 + t_hipKernelGetParamInfo hipKernelGetParamInfo_fn; + + // DO NOT EDIT ABOVE! + // HIP_RUNTIME_API_TABLE_STEP_VERSION == 21 // ******************************************************************************************* // // 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 8d2c6c02ba..7aef3e7f52 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 @@ -467,7 +467,8 @@ enum hip_api_id_t { HIP_API_ID_hipLibraryEnumerateKernels = 447, HIP_API_ID_hipKernelGetName = 448, HIP_API_ID_hipOccupancyAvailableDynamicSMemPerBlock = 449, - HIP_API_ID_LAST = 449, + HIP_API_ID_hipKernelGetParamInfo = 450, + HIP_API_ID_LAST = 450, HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice), HIP_API_ID_hipGetDeviceProperties = HIP_API_ID_CONCAT(HIP_API_ID_,hipGetDeviceProperties), @@ -733,6 +734,7 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipIpcOpenMemHandle: return "hipIpcOpenMemHandle"; case HIP_API_ID_hipKernelGetLibrary: return "hipKernelGetLibrary"; case HIP_API_ID_hipKernelGetName: return "hipKernelGetName"; + case HIP_API_ID_hipKernelGetParamInfo: return "hipKernelGetParamInfo"; case HIP_API_ID_hipLaunchByPtr: return "hipLaunchByPtr"; case HIP_API_ID_hipLaunchCooperativeKernel: return "hipLaunchCooperativeKernel"; case HIP_API_ID_hipLaunchCooperativeKernelMultiDevice: return "hipLaunchCooperativeKernelMultiDevice"; @@ -1176,6 +1178,7 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipIpcOpenMemHandle", name) == 0) return HIP_API_ID_hipIpcOpenMemHandle; if (strcmp("hipKernelGetLibrary", name) == 0) return HIP_API_ID_hipKernelGetLibrary; if (strcmp("hipKernelGetName", name) == 0) return HIP_API_ID_hipKernelGetName; + if (strcmp("hipKernelGetParamInfo", name) == 0) return HIP_API_ID_hipKernelGetParamInfo; if (strcmp("hipLaunchByPtr", name) == 0) return HIP_API_ID_hipLaunchByPtr; if (strcmp("hipLaunchCooperativeKernel", name) == 0) return HIP_API_ID_hipLaunchCooperativeKernel; if (strcmp("hipLaunchCooperativeKernelMultiDevice", name) == 0) return HIP_API_ID_hipLaunchCooperativeKernelMultiDevice; @@ -2694,6 +2697,14 @@ typedef struct hip_api_data_s { const char* name__val; hipKernel_t kernel; } hipKernelGetName; + struct { + hipKernel_t kernel; + size_t paramIndex; + size_t* paramOffset; + size_t paramOffset__val; + size_t* paramSize; + size_t paramSize__val; + } hipKernelGetParamInfo; struct { const void* hostFunction; } hipLaunchByPtr; @@ -5352,6 +5363,13 @@ typedef struct hip_api_data_s { cb_data.args.hipKernelGetName.name = (const char**)name; \ cb_data.args.hipKernelGetName.kernel = (hipKernel_t)kernel; \ }; +// hipKernelGetParamInfo[('hipKernel_t', 'kernel'), ('size_t', 'paramIndex'), ('size_t*', 'paramOffset'), ('size_t*', 'paramSize')] +#define INIT_hipKernelGetParamInfo_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipKernelGetParamInfo.kernel = (hipKernel_t)kernel; \ + cb_data.args.hipKernelGetParamInfo.paramIndex = (size_t)paramIndex; \ + cb_data.args.hipKernelGetParamInfo.paramOffset = (size_t*)paramOffset; \ + cb_data.args.hipKernelGetParamInfo.paramSize = (size_t*)paramSize; \ +}; // hipLaunchByPtr[('const void*', 'hostFunction')] #define INIT_hipLaunchByPtr_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipLaunchByPtr.hostFunction = (const void*)hostFunction; \ @@ -7698,6 +7716,11 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { case HIP_API_ID_hipKernelGetName: if (data->args.hipKernelGetName.name) data->args.hipKernelGetName.name__val = *(data->args.hipKernelGetName.name); break; +// hipKernelGetParamInfo[('hipKernel_t', 'kernel'), ('size_t', 'paramIndex'), ('size_t*', 'paramOffset'), ('size_t*', 'paramSize')] + case HIP_API_ID_hipKernelGetParamInfo: + if (data->args.hipKernelGetParamInfo.paramOffset) data->args.hipKernelGetParamInfo.paramOffset__val = *(data->args.hipKernelGetParamInfo.paramOffset); + if (data->args.hipKernelGetParamInfo.paramSize) data->args.hipKernelGetParamInfo.paramSize__val = *(data->args.hipKernelGetParamInfo.paramSize); + break; // hipLaunchByPtr[('const void*', 'hostFunction')] case HIP_API_ID_hipLaunchByPtr: break; @@ -10289,6 +10312,16 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << ", kernel="; roctracer::hip_support::detail::operator<<(oss, data->args.hipKernelGetName.kernel); oss << ")"; break; + case HIP_API_ID_hipKernelGetParamInfo: + oss << "hipKernelGetParamInfo("; + oss << "kernel="; roctracer::hip_support::detail::operator<<(oss, data->args.hipKernelGetParamInfo.kernel); + oss << ", paramIndex="; roctracer::hip_support::detail::operator<<(oss, data->args.hipKernelGetParamInfo.paramIndex); + if (data->args.hipKernelGetParamInfo.paramOffset == NULL) oss << ", paramOffset=NULL"; + else { oss << ", paramOffset="; roctracer::hip_support::detail::operator<<(oss, data->args.hipKernelGetParamInfo.paramOffset__val); } + if (data->args.hipKernelGetParamInfo.paramSize == NULL) oss << ", paramSize=NULL"; + else { oss << ", paramSize="; roctracer::hip_support::detail::operator<<(oss, data->args.hipKernelGetParamInfo.paramSize__val); } + oss << ")"; + break; case HIP_API_ID_hipLaunchByPtr: oss << "hipLaunchByPtr("; oss << "hostFunction="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLaunchByPtr.hostFunction); diff --git a/projects/clr/hipamd/src/amdhip.def b/projects/clr/hipamd/src/amdhip.def index 6965222310..dff046d1ac 100644 --- a/projects/clr/hipamd/src/amdhip.def +++ b/projects/clr/hipamd/src/amdhip.def @@ -522,3 +522,4 @@ hipKernelGetLibrary hipKernelGetName hipOccupancyAvailableDynamicSMemPerBlock hipGetProcAddress_spt +hipKernelGetParamInfo diff --git a/projects/clr/hipamd/src/hip_api_trace.cpp b/projects/clr/hipamd/src/hip_api_trace.cpp index 2232f0c71e..6399924ec2 100644 --- a/projects/clr/hipamd/src/hip_api_trace.cpp +++ b/projects/clr/hipamd/src/hip_api_trace.cpp @@ -883,6 +883,8 @@ hipError_t hipKernelGetLibrary(hipLibrary_t* library, hipKernel_t kernel); hipError_t hipKernelGetName(const char** name, hipKernel_t kernel); hipError_t hipOccupancyAvailableDynamicSMemPerBlock(size_t* dynamicSmemSize, const void* f, int numBlocks, int blockSize); +hipError_t hipKernelGetParamInfo(hipKernel_t kernel, size_t paramIndex, size_t* paramOffset, + size_t* paramSize); } // namespace hip namespace hip { @@ -1429,6 +1431,7 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) { ptrDispatchTable->hipKernelGetLibrary_fn = hip::hipKernelGetLibrary; ptrDispatchTable->hipKernelGetName_fn = hip::hipKernelGetName; ptrDispatchTable->hipOccupancyAvailableDynamicSMemPerBlock_fn = hip::hipOccupancyAvailableDynamicSMemPerBlock; + ptrDispatchTable->hipKernelGetParamInfo_fn = hip::hipKernelGetParamInfo; } #if HIP_ROCPROFILER_REGISTER > 0 @@ -2109,15 +2112,17 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipKernelGetName_fn, 504); HIP_ENFORCE_ABI(HipDispatchTable, hipOccupancyAvailableDynamicSMemPerBlock_fn, 505); // HIP_RUNTIME_API_TABLE_STEP_VERSION == 19 HIP_ENFORCE_ABI(HipDispatchTable, hipGetProcAddress_spt_fn, 506); +// HIP_RUNTIME_API_TABLE_STEP_VERSION == 20 +HIP_ENFORCE_ABI(HipDispatchTable, hipKernelGetParamInfo_fn, 507); // 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, 508) -static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 19, +static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 20, "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 fc915143e8..37437589e8 100644 --- a/projects/clr/hipamd/src/hip_hcc.map.in +++ b/projects/clr/hipamd/src/hip_hcc.map.in @@ -644,6 +644,7 @@ global: hipKernelGetName; hipOccupancyAvailableDynamicSMemPerBlock; hipGetProcAddress_spt; + hipKernelGetParamInfo; local: *; } hip_7.1; diff --git a/projects/clr/hipamd/src/hip_library.cpp b/projects/clr/hipamd/src/hip_library.cpp index 900a909b07..bfccc4e4f1 100644 --- a/projects/clr/hipamd/src/hip_library.cpp +++ b/projects/clr/hipamd/src/hip_library.cpp @@ -270,4 +270,30 @@ hipError_t hipKernelGetName(const char** name, hipKernel_t kernel) { HIP_RETURN(ret); } +hipError_t hipKernelGetParamInfo(hipKernel_t kernel, size_t paramIndex, size_t* paramOffset, + size_t* paramSize ) { + HIP_INIT_API(hipKernelGetParamInfo, kernel, paramIndex, paramOffset, paramSize); + if (kernel == nullptr || paramOffset == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + const auto* const d_function = hip::DeviceFunc::asFunction(reinterpret_cast(kernel)); + if (d_function == nullptr) { + HIP_RETURN(hipErrorInvalidHandle); + } + const auto* const d_kernel = d_function->kernel(); + if (d_kernel == nullptr) { + HIP_RETURN(hipErrorInvalidDeviceFunction); + } + const amd::KernelSignature& signature = d_kernel->signature(); + if (paramIndex >= signature.numParameters()) { + HIP_RETURN(hipErrorInvalidValue); + } + const amd::KernelParameterDescriptor& desc = signature.at(paramIndex); + *paramOffset = desc.offset_; + if (paramSize != nullptr) { + *paramSize = desc.size_; + } + HIP_RETURN(hipSuccess); +} + } // namespace hip diff --git a/projects/clr/hipamd/src/hip_table_interface.cpp b/projects/clr/hipamd/src/hip_table_interface.cpp index 3e378fb4d2..443ba81ba5 100644 --- a/projects/clr/hipamd/src/hip_table_interface.cpp +++ b/projects/clr/hipamd/src/hip_table_interface.cpp @@ -2062,4 +2062,9 @@ hipError_t hipKernelGetLibrary(hipLibrary_t* library, hipKernel_t kernel) { } hipError_t hipKernelGetName(const char** name, hipKernel_t kernel) { return hip::GetHipDispatchTable()->hipKernelGetName_fn(name, kernel); +} +hipError_t hipKernelGetParamInfo(hipKernel_t kernel, size_t paramIndex, size_t* paramOffset, + size_t* paramSize) { + return hip::GetHipDispatchTable()->hipKernelGetParamInfo_fn(kernel, paramIndex, paramOffset, + paramSize); } \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/library/loadlib_co.cc b/projects/hip-tests/catch/unit/library/loadlib_co.cc index f59e4577ee..268958d7c8 100644 --- a/projects/hip-tests/catch/unit/library/loadlib_co.cc +++ b/projects/hip-tests/catch/unit/library/loadlib_co.cc @@ -76,6 +76,13 @@ TEST_CASE("Unit_hip_library_load_co") { HIP_CHECK(hipLibraryGetKernelCount(&count, library)); REQUIRE(count == 3); + size_t offset, paramsize; + for (size_t k = 0; k < count; ++k) { + HIP_CHECK(hipKernelGetParamInfo(function, k, &offset, ¶msize)); + REQUIRE(offset == k * sizeof(float*)); + REQUIRE(paramsize == sizeof(float*)); + } + void* args[] = {&d_out, &d_in1, &d_in2}; HIP_CHECK(hipLaunchKernel(function, 1, size, args, 0, stream)); @@ -104,6 +111,13 @@ TEST_CASE("Unit_hip_library_load_co") { HIP_CHECK(hipLibraryGetKernelCount(&count, library)); REQUIRE(count == 3); + size_t offset, paramsize; + for (size_t k = 0; k < count; ++k) { + HIP_CHECK(hipKernelGetParamInfo(function, k, &offset, ¶msize)); + REQUIRE(offset == k * sizeof(float*)); + REQUIRE(paramsize == sizeof(float*)); + } + void* args[] = {&d_out, &d_in1, &d_in2}; HIP_CHECK(hipLaunchKernel(function, 1, size, args, 0, stream)); @@ -132,6 +146,13 @@ TEST_CASE("Unit_hip_library_load_co") { HIP_CHECK(hipLibraryGetKernelCount(&count, library)); REQUIRE(count == 3); + size_t offset, paramsize; + for (size_t k = 0; k < count; ++k) { + HIP_CHECK(hipKernelGetParamInfo(function, k, &offset, ¶msize)); + REQUIRE(offset == k * sizeof(float*)); + REQUIRE(paramsize == sizeof(float*)); + } + void* args[] = {&d_out, &d_in1, &d_in2}; HIP_CHECK(hipLaunchKernel(function, 1, size, args, 0, stream)); @@ -169,9 +190,13 @@ TEST_CASE("Unit_hip_library_load_co") { std::vector out(size, 0); + size_t offset, paramsize; for (int k = 0; k < num_kernels; k++) { const char* kName = nullptr; HIP_CHECK(hipKernelGetName(&kName, functions[k])); + HIP_CHECK(hipKernelGetParamInfo(functions[k], k, &offset, ¶msize)); + REQUIRE(paramsize == sizeof(float*)); + REQUIRE(offset == k * sizeof(float*)); HIP_CHECK(hipLaunchKernel(functions[k], 1, size, args, 0, stream)); HIP_CHECK(hipStreamSynchronize(stream)); HIP_CHECK(hipMemcpy(out.data(), d_out, sizeof(float) * size, hipMemcpyDeviceToHost)); @@ -186,4 +211,25 @@ TEST_CASE("Unit_hip_library_load_co") { HIP_CHECK(hipFree(d_in1)); HIP_CHECK(hipFree(d_in2)); HIP_CHECK(hipFree(d_out)); -} \ No newline at end of file +} + +TEST_CASE("Unit_hipKernelGetParamInfo_Negative") { + size_t offset, paramsize; + + SECTION("Kernel as nullptr") { + HIP_CHECK_ERROR(hipKernelGetParamInfo(nullptr, 0, &offset, ¶msize), hipErrorInvalidValue); + } + + std::string lib_co = "library_code_load.code"; + hipLibrary_t library; + hipKernel_t function; + + HIP_CHECK( + hipLibraryLoadFromFile(&library, lib_co.data(), nullptr, nullptr, 0, nullptr, nullptr, 0)); + HIP_CHECK(hipLibraryGetKernel(&function, library, "mul_kernel")); + + SECTION("Param offset as nullptr") { + HIP_CHECK_ERROR(hipKernelGetParamInfo(function, 0, nullptr, ¶msize), hipErrorInvalidValue); + } + +} diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index 1728b22b1b..b2b4dc3e93 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -6466,6 +6466,19 @@ hipError_t hipKernelGetLibrary(hipLibrary_t* library, hipKernel_t kernel); */ hipError_t hipKernelGetName(const char** name, hipKernel_t kernel); +/** + * @brief Returns the offset and size of a kernel parameter + * + * @param [in] kernel Kernel handle to retrieve parameter info + * @param [in] paramIndex Index of the parameter + * @param [out] paramOffset returns the offset of the parameter + * @param [out] paramSize Optionally returns the size of the parameter + * + * @return #hipSuccess, #hipErrorInvalidValue +*/ +hipError_t hipKernelGetParamInfo(hipKernel_t kernel, size_t paramIndex, size_t* paramOffset, + size_t* paramSize); + /** * @brief Find out attributes for a given function. * @ingroup Execution diff --git a/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h b/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h index 1fc60e8680..8f6c295aab 100644 --- a/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015 - 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - 2025 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 @@ -3764,6 +3764,11 @@ inline static hipError_t hipKernelGetName(const char** name, hipKernel_t kernel) return hipCUResultTohipError(cuKernelGetName(name, kernel)); } +inline static hipError_t hipKernelGetParamInfo(hipKernel_t kernel, size_t paramIndex, size_t* paramOffset, + size_t* paramSize) { + return hipCUResultTohipError(cuKernelGetParamInfo(kernel, paramIndex, paramOffset, paramSize)); +} + inline static hipError_t hipLaunchKernel(const void* function_address, dim3 numBlocks, dim3 dimBlocks, void** args, size_t sharedMemBytes, hipStream_t stream) { diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/enum_string.hpp b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/enum_string.hpp index d00fb3a593..514d56e632 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/enum_string.hpp +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/enum_string.hpp @@ -1001,6 +1001,9 @@ ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipOccupancyAvailableDynam #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 >= 20 +ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipKernelGetParamInfo) +#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 @@ -1041,6 +1044,8 @@ static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 505); 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); +#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 20 +static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 508); #else # if !defined(ROCPROFILER_UNSAFE_NO_VERSION_CHECK) && \ (defined(ROCPROFILER_CI) && ROCPROFILER_CI > 0) diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h index 51c8c8344f..17634ce0c8 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h @@ -3368,6 +3368,15 @@ typedef union rocprofiler_hip_api_args_t hipDriverProcAddressQueryResult* symbolStatus; } hipGetProcAddress_spt; #endif +#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 20 + struct + { + hipKernel_t kernel; + size_t paramIndex; + size_t* paramOffset; + size_t* paramSize; + } hipKernelGetParamInfo; +#endif } rocprofiler_hip_api_args_t; ROCPROFILER_EXTERN_C_FINI diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/runtime_api_id.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/runtime_api_id.h index f1e9be8e9a..4c2475ddaf 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/runtime_api_id.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/runtime_api_id.h @@ -572,6 +572,9 @@ typedef enum rocprofiler_hip_runtime_api_id_t // NOLINT(performance-enum-size) #endif #if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 19 ROCPROFILER_HIP_RUNTIME_API_ID_hipGetProcAddress_spt, +#endif +#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 20 + ROCPROFILER_HIP_RUNTIME_API_ID_hipKernelGetParamInfo, #endif ROCPROFILER_HIP_RUNTIME_API_ID_LAST, } rocprofiler_hip_runtime_api_id_t; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/abi.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/abi.cpp index a906876f60..5fa4222ad1 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/abi.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/abi.cpp @@ -619,6 +619,10 @@ ROCP_SDK_ENFORCE_ABI(::HipDispatchTable, hipOccupancyAvailableDynamicSMemPerBloc ROCP_SDK_ENFORCE_ABI(::HipDispatchTable, hipGetProcAddress_spt_fn, 506); #endif +#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 20 +ROCP_SDK_ENFORCE_ABI(::HipDispatchTable, hipKernelGetParamInfo_fn, 507); +#endif + #if HIP_RUNTIME_API_TABLE_STEP_VERSION == 0 ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 442) #elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 1 @@ -659,6 +663,8 @@ ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 505) ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 506) #elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 19 ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 507) +#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 20 +ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 508) #else INTERNAL_CI_ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 0) #endif diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp index b012103fb0..44cf89d6c1 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp @@ -645,6 +645,11 @@ HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNT #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 + +#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 20 +HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNTIME_API_ID_hipKernelGetParamInfo, hipKernelGetParamInfo, hipKernelGetParamInfo_fn, kernel, paramIndex, paramOffset, paramSize); +#endif + // clang-format on #else