SWDEV-546308 - Implement hipKernelGetParamInfo API (#1783)

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`

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
 
   // ******************************************************************************************* //
   //

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);

projects/clr/hipamd/src/amdhip.def

@@ -522,3 +522,4 @@ hipKernelGetLibrary
 hipKernelGetName
 hipOccupancyAvailableDynamicSMemPerBlock
 hipGetProcAddress_spt
+hipKernelGetParamInfo

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(<table>, <functor>, 8)
 //
 //  HIP_ENFORCE_ABI_VERSIONING(<table>, 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

projects/clr/hipamd/src/hip_hcc.map.in

@@ -644,6 +644,7 @@ global:
     hipKernelGetName;
     hipOccupancyAvailableDynamicSMemPerBlock;
     hipGetProcAddress_spt;
+    hipKernelGetParamInfo;
 local:
     *;
 } hip_7.1;

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<hipFunction_t>(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

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);
 }

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, &paramsize));
+      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, &paramsize));
+      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, &paramsize));
+      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<float> 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, &paramsize));
+      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));
-}
+}
+
+TEST_CASE("Unit_hipKernelGetParamInfo_Negative") {
+  size_t offset, paramsize;
+
+  SECTION("Kernel as nullptr") {
+    HIP_CHECK_ERROR(hipKernelGetParamInfo(nullptr, 0, &offset, &paramsize), 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, &paramsize), hipErrorInvalidValue);
+  }
+
+}

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

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) {

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)

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

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;

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

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