SWDEV-546322 Implement hipMipmappedArrayGetMemoryRequirements

sluzynsk-amd · sluzynsk-amd · commit 76228e11ac87 · 2025-10-24T11:20:01.000Z
Signed-off-by: Sebastian Luzynski &lt;Sebastian.Luzynski@amd.com&gt;
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
@@ -1103,8 +1103,9 @@ typedef hipError_t (*t_hipLibraryLoadFromFile)(hipLibrary_t* library, const char
 typedef hipError_t (*t_hipLibraryUnload)(hipLibrary_t library);
 typedef hipError_t (*t_hipLibraryGetKernel)(hipKernel_t* pKernel, hipLibrary_t library,
                                             const char* name);
-typedef hipError_t (*t_hipLibraryGetKernelCount)(unsigned int *count,
-                                                 hipLibrary_t library);
+typedef hipError_t (*t_hipLibraryGetKernelCount)(unsigned int* count, hipLibrary_t library);
+typedef hipError_t (*t_hipMipmappedArrayGetMemoryRequirements)(
+    hipArrayMemoryRequirements* memoryRequirements, hipMipmappedArray_t mipmap, hipDevice_t device);
 
 // HIP Compiler dispatch table
 struct HipCompilerDispatchTable {
@@ -1682,6 +1683,7 @@ struct HipDispatchTable {
 
   // HIP_RUNTIME_API_TABLE_STEP_VERSION = 16
   t_hipStreamCopyAttributes hipStreamCopyAttributes_fn;
+  t_hipMipmappedArrayGetMemoryRequirements hipMipmappedArrayGetMemoryRequirements_fn;
 
   // DO NOT EDIT ABOVE!
   // HIP_RUNTIME_API_TABLE_STEP_VERSION == 17
@@ -1721,4 +1723,4 @@ struct HipToolsDispatchTable {
   // 4) GENERATE COMMENT FOR NEXT STEP VERSION
   // 5) ADD "DO NOT EDIT ABOVE!" COMMENT
   // ******************************************************************************************* //
-};
+};
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
@@ -463,7 +463,8 @@ enum hip_api_id_t {
   HIP_API_ID_hipLibraryGetKernelCount = 443,
   HIP_API_ID_hipMemGetHandleForAddressRange = 444,
   HIP_API_ID_hipStreamCopyAttributes = 445,
-  HIP_API_ID_LAST = 445,
+  HIP_API_ID_hipMipmappedArrayGetMemoryRequirements = 446,
+  HIP_API_ID_LAST = 446,
 
   HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice),
   HIP_API_ID_hipGetDeviceProperties = HIP_API_ID_CONCAT(HIP_API_ID_,hipGetDeviceProperties),
@@ -850,6 +851,7 @@ static inline const char* hip_api_name(const uint32_t id) {
     case HIP_API_ID_hipMipmappedArrayCreate: return "hipMipmappedArrayCreate";
     case HIP_API_ID_hipMipmappedArrayDestroy: return "hipMipmappedArrayDestroy";
     case HIP_API_ID_hipMipmappedArrayGetLevel: return "hipMipmappedArrayGetLevel";
+    case HIP_API_ID_hipMipmappedArrayGetMemoryRequirements: return "hipMipmappedArrayGetMemoryRequirements";
     case HIP_API_ID_hipModuleGetFunction: return "hipModuleGetFunction";
     case HIP_API_ID_hipModuleGetFunctionCount: return "hipModuleGetFunctionCount";
     case HIP_API_ID_hipModuleGetGlobal: return "hipModuleGetGlobal";
@@ -1289,6 +1291,7 @@ static inline uint32_t hipApiIdByName(const char* name) {
   if (strcmp("hipMipmappedArrayCreate", name) == 0) return HIP_API_ID_hipMipmappedArrayCreate;
   if (strcmp("hipMipmappedArrayDestroy", name) == 0) return HIP_API_ID_hipMipmappedArrayDestroy;
   if (strcmp("hipMipmappedArrayGetLevel", name) == 0) return HIP_API_ID_hipMipmappedArrayGetLevel;
+  if (strcmp("hipMipmappedArrayGetMemoryRequirements", name) == 0) return HIP_API_ID_hipMipmappedArrayGetMemoryRequirements;
   if (strcmp("hipModuleGetFunction", name) == 0) return HIP_API_ID_hipModuleGetFunction;
   if (strcmp("hipModuleGetFunctionCount", name) == 0) return HIP_API_ID_hipModuleGetFunctionCount;
   if (strcmp("hipModuleGetGlobal", name) == 0) return HIP_API_ID_hipModuleGetGlobal;
@@ -3509,6 +3512,12 @@ typedef struct hip_api_data_s {
       hipMipmappedArray_t hMipMappedArray;
       unsigned int level;
     } hipMipmappedArrayGetLevel;
+    struct {
+      hipArrayMemoryRequirements* memoryRequirements;
+      hipArrayMemoryRequirements memoryRequirements__val;
+      hipMipmappedArray_t mipmap;
+      hipDevice_t device;
+    } hipMipmappedArrayGetMemoryRequirements;
     struct {
       hipFunction_t* function;
       hipFunction_t function__val;
@@ -6174,6 +6183,12 @@ typedef struct hip_api_data_s {
   cb_data.args.hipMipmappedArrayGetLevel.hMipMappedArray = (hipMipmappedArray_t)mipmapped_array_ptr; \
   cb_data.args.hipMipmappedArrayGetLevel.level = (unsigned int)mip_level; \
 };
+// hipMipmappedArrayGetMemoryRequirements[('hipArrayMemoryRequirements*', 'memoryRequirements'), ('hipMipmappedArray_t', 'mipmap'), ('hipDevice_t', 'device')]
+#define INIT_hipMipmappedArrayGetMemoryRequirements_CB_ARGS_DATA(cb_data) { \
+  cb_data.args.hipMipmappedArrayGetMemoryRequirements.memoryRequirements = (hipArrayMemoryRequirements*)memoryRequirements; \
+  cb_data.args.hipMipmappedArrayGetMemoryRequirements.mipmap = (hipMipmappedArray_t)mipmap; \
+  cb_data.args.hipMipmappedArrayGetMemoryRequirements.device = (hipDevice_t)device; \
+};
 // hipModuleGetFunction[('hipFunction_t*', 'function'), ('hipModule_t', 'module'), ('const char*', 'kname')]
 #define INIT_hipModuleGetFunction_CB_ARGS_DATA(cb_data) { \
   cb_data.args.hipModuleGetFunction.function = (hipFunction_t*)hfunc; \
@@ -8094,6 +8109,10 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
     case HIP_API_ID_hipMipmappedArrayGetLevel:
       if (data->args.hipMipmappedArrayGetLevel.pLevelArray) data->args.hipMipmappedArrayGetLevel.pLevelArray__val = *(data->args.hipMipmappedArrayGetLevel.pLevelArray);
       break;
+// hipMipmappedArrayGetMemoryRequirements[('hipArrayMemoryRequirements*', 'memoryRequirements'), ('hipMipmappedArray_t', 'mipmap'), ('hipDevice_t', 'device')]
+    case HIP_API_ID_hipMipmappedArrayGetMemoryRequirements:
+      if (data->args.hipMipmappedArrayGetMemoryRequirements.memoryRequirements) data->args.hipMipmappedArrayGetMemoryRequirements.memoryRequirements__val = *(data->args.hipMipmappedArrayGetMemoryRequirements.memoryRequirements);
+      break;
 // hipModuleGetFunction[('hipFunction_t*', 'function'), ('hipModule_t', 'module'), ('const char*', 'kname')]
     case HIP_API_ID_hipModuleGetFunction:
       if (data->args.hipModuleGetFunction.function) data->args.hipModuleGetFunction.function__val = *(data->args.hipModuleGetFunction.function);
@@ -11284,6 +11303,14 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
       oss << ", level="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMipmappedArrayGetLevel.level);
       oss << ")";
     break;
+    case HIP_API_ID_hipMipmappedArrayGetMemoryRequirements:
+      oss << "hipMipmappedArrayGetMemoryRequirements(";
+      if (data->args.hipMipmappedArrayGetMemoryRequirements.memoryRequirements == NULL) oss << "memoryRequirements=NULL";
+      else { oss << "memoryRequirements="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMipmappedArrayGetMemoryRequirements.memoryRequirements__val); }
+      oss << ", mipmap="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMipmappedArrayGetMemoryRequirements.mipmap);
+      oss << ", device="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMipmappedArrayGetMemoryRequirements.device);
+      oss << ")";
+    break;
     case HIP_API_ID_hipModuleGetFunction:
       oss << "hipModuleGetFunction(";
       if (data->args.hipModuleGetFunction.function == NULL) oss << "function=NULL";
diff --git a/projects/clr/hipamd/src/amdhip.def b/projects/clr/hipamd/src/amdhip.def
@@ -517,3 +517,4 @@ hipLibraryUnload
 hipLibraryGetKernel
 hipLibraryGetKernelCount
 hipStreamCopyAttributes
+hipMipmappedArrayGetMemoryRequirements
diff --git a/projects/clr/hipamd/src/hip_api_trace.cpp b/projects/clr/hipamd/src/hip_api_trace.cpp
@@ -875,6 +875,8 @@ hipError_t hipLibraryLoadFromFile(hipLibrary_t* library, const char* fileName,
 hipError_t hipLibraryUnload(hipLibrary_t library);
 hipError_t hipLibraryGetKernel(hipKernel_t* pKernel, hipLibrary_t library, const char* name);
 hipError_t hipLibraryGetKernelCount(unsigned int* count, hipLibrary_t library);
+hipError_t hipMipmappedArrayGetMemoryRequirements(hipArrayMemoryRequirements* memoryRequirements,
+                                                  hipMipmappedArray_t mipmap, hipDevice_t device);
 }  // namespace hip
 
 namespace hip {
@@ -1219,6 +1221,8 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) {
   ptrDispatchTable->hipMemsetD8Async_fn = hip::hipMemsetD8Async;
   ptrDispatchTable->hipMipmappedArrayCreate_fn = hip::hipMipmappedArrayCreate;
   ptrDispatchTable->hipMipmappedArrayDestroy_fn = hip::hipMipmappedArrayDestroy;
+  ptrDispatchTable->hipMipmappedArrayGetMemoryRequirements_fn =
+     hip::hipMipmappedArrayGetMemoryRequirements;
   ptrDispatchTable->hipMipmappedArrayGetLevel_fn = hip::hipMipmappedArrayGetLevel;
   ptrDispatchTable->hipModuleGetFunction_fn = hip::hipModuleGetFunction;
   ptrDispatchTable->hipModuleGetFunctionCount_fn = hip::hipModuleGetFunctionCount;
@@ -2088,13 +2092,14 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipLibraryGetKernel_fn, 499);
 HIP_ENFORCE_ABI(HipDispatchTable, hipLibraryGetKernelCount_fn, 500);
 // HIP_RUNTIME_API_TABLE_STEP_VERSION == 16
 HIP_ENFORCE_ABI(HipDispatchTable, hipStreamCopyAttributes_fn, 501);
+HIP_ENFORCE_ABI(HipDispatchTable, hipMipmappedArrayGetMemoryRequirements_fn, 502);
 // 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, 502)
+HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 503)
 
 static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 16,
               "If you get this error, add new HIP_ENFORCE_ABI(...) code for the new function "
diff --git a/projects/clr/hipamd/src/hip_conversions.hpp b/projects/clr/hipamd/src/hip_conversions.hpp
@@ -994,4 +994,9 @@ inline hipMemcpy3DParms getMemcpy3DParms(const hipMemcpy3DPeerParms& desc) {
   params.kind = hipMemcpyDeviceToDevice;
   return params;
 }
-};  // namespace hip
+
+inline hipArrayMemoryRequirements getArrayMemoryRequirements(
+    const hipArrayMemoryRequirements& desc) {
+  return hipArrayMemoryRequirements { desc.alignment, desc.size };
+}
+};  // namespace hip
diff --git a/projects/clr/hipamd/src/hip_hcc.map.in b/projects/clr/hipamd/src/hip_hcc.map.in
@@ -639,6 +639,7 @@ global:
     hipLibraryGetKernel;
     hipLibraryGetKernelCount;
     hipStreamCopyAttributes;
+    hipMipmappedArrayGetMemoryRequirements;
 local:
     *;
-} hip_7.1;
+} hip_7.1;
diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp
@@ -4264,6 +4264,32 @@ hipError_t ihipMipmappedArrayGetLevel(hipArray_t* level_array_pptr,
   return hipSuccess;
 }
 
+hipError_t ihipMipmappedArrayGetMemoryRequirements(hipArrayMemoryRequirements* memoryRequirements,
+                                                  hipMipmappedArray_t mipmap,
+                                                  hipDevice_t device) {
+  if (memoryRequirements == nullptr) {
+    return hipErrorInvalidValue;
+  }
+  if (mipmap == nullptr) {
+    return hipErrorInvalidHandle;
+  }
+
+  cl_mem cl_mem_obj = reinterpret_cast<cl_mem>(mipmap->data);
+  if (is_valid(cl_mem_obj) == false) {
+    return hipErrorInvalidValue;
+  }
+
+  amd::Image* image = as_amd(cl_mem_obj)->asImage();
+  if (image == nullptr) {
+    return hipErrorInvalidValue;
+  }
+
+  memoryRequirements->alignment = image->getAlignment();
+  memoryRequirements->size = image->getSize();
+
+  return hipSuccess;
+}
+
 hipError_t hipMipmappedArrayCreate(hipMipmappedArray_t* mipmapped_array_pptr,
                                    HIP_ARRAY3D_DESCRIPTOR* mipmapped_array_desc_ptr,
                                    unsigned int num_mipmap_levels) {
@@ -4288,6 +4314,13 @@ hipError_t hipMipmappedArrayGetLevel(hipArray_t* level_array_pptr,
   HIP_RETURN(ihipMipmappedArrayGetLevel(level_array_pptr, mipmapped_array_ptr, mip_level));
 }
 
+hipError_t hipMipmappedArrayGetMemoryRequirements(hipArrayMemoryRequirements* memoryRequirements,
+                                                  hipMipmappedArray_t mipmap,
+                                                  hipDevice_t device) {
+  HIP_INIT_API(hipMipmappedArrayGetMemoryRequirements, memoryRequirements, mipmap, device);
+  HIP_RETURN(ihipMipmappedArrayGetMemoryRequirements(memoryRequirements, mipmap, device));
+}
+
 hipError_t hipMallocMipmappedArray(hipMipmappedArray_t* mipmappedArray,
                                    const hipChannelFormatDesc* desc, hipExtent extent,
                                    unsigned int numLevels, unsigned int flags) {
diff --git a/projects/clr/hipamd/src/hip_table_interface.cpp b/projects/clr/hipamd/src/hip_table_interface.cpp
@@ -2040,4 +2040,9 @@ hipError_t hipLibraryGetKernel(hipKernel_t* pKernel, hipLibrary_t library, const
 hipError_t hipLibraryGetKernelCount(unsigned int *count, hipLibrary_t library) {
   return hip::GetHipDispatchTable()->hipLibraryGetKernelCount_fn(count,
                                                                  library);
-}
+}
+hipError_t hipMipmappedArrayGetMemoryRequirements(hipArrayMemoryRequirements* memoryRequirements,
+                                                  hipMipmappedArray_t mipmap, hipDevice_t device) {
+  return hip::GetHipDispatchTable()->hipMipmappedArrayGetMemoryRequirements_fn(memoryRequirements,
+                                                                               mipmap, device);
+}
diff --git a/projects/hip-tests/catch/unit/texture/CMakeLists.txt b/projects/hip-tests/catch/unit/texture/CMakeLists.txt
@@ -147,6 +147,7 @@ set(TEST_SRC
     tex2DLod.cc
     tex2DLayeredLod.cc
     tex3DLod.cc
+    hipMipmappedArrayGetMemoryRequirements.cc
 )
 endif()
 
diff --git a/projects/hip-tests/catch/unit/texture/hipMipmappedArrayGetMemoryRequirements.cc b/projects/hip-tests/catch/unit/texture/hipMipmappedArrayGetMemoryRequirements.cc
@@ -0,0 +1,58 @@
+/*
+Copyright (c) 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
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_array_common.hh>
+#include <hip_test_common.hh>
+
+TEST_CASE("Unit_hipMipmappedArrayGetMemoryRequirements_Negative_Parameters") {
+  CHECK_IMAGE_SUPPORT;
+
+#ifdef __linux__
+  HipTest::HIP_SKIP_TEST("Mipmap APIs are not supported on Linux");
+  return;
+#endif  //__linux__
+
+  const int device_id = 0;
+  hipArrayMemoryRequirements memoryRequirements{};
+  hipmipmappedArray array;
+  HIP_ARRAY3D_DESCRIPTOR desc = {};
+  using vec_info = vector_info<float>;
+  desc.Format = vec_info::format;
+  desc.NumChannels = vec_info::size;
+  desc.Width = 4;
+  desc.Height = 4;
+  desc.Depth = 6;
+  desc.Flags = 0;
+
+  unsigned int levels = 1 + std::log2(desc.Depth);
+
+  HIP_CHECK(hipFree(0));
+  HIP_CHECK(hipMipmappedArrayCreate(&array, &desc, levels));
+
+  SECTION("memoryRequirements is nullptr") {
+    HIP_CHECK_ERROR(hipMipmappedArrayGetMemoryRequirements(nullptr, array, device_id), hipErrorInvalidValue);
+  }
+
+  SECTION("mipmap is nullptr") {
+    HIP_CHECK_ERROR(hipMipmappedArrayGetMemoryRequirements(&memoryRequirements, nullptr, device_id), hipErrorInvalidHandle);
+  }
+}
diff --git a/projects/hip/include/hip/driver_types.h b/projects/hip/include/hip/driver_types.h
@@ -131,6 +131,13 @@ typedef struct hip_Memcpy2D {
   size_t Height;                ///< Height of the 2D memory copy
 } hip_Memcpy2D;
 #endif  // !defined(__HIPCC_RTC__)
+/**
+ * Struct representing array memory requirements.
+ */
+typedef struct hipArrayMemoryRequirements {
+  size_t alignment;
+  size_t size;
+} hipArrayMemoryRequirements;
 /**
  * HIP mipmapped array
  */
diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h
@@ -5813,6 +5813,24 @@ hipError_t hipMemcpy3DPeer(hipMemcpy3DPeerParms* p);
  * @returns #hipSuccess, #hipErrorInvalidValue, hipErrorInvalidDevice
  */
 hipError_t hipMemcpy3DPeerAsync(hipMemcpy3DPeerParms* p, hipStream_t stream __dparm(0));
+
+/**
+ * @brief Returns the memory requirements of a HIP mipmapped array.
+ *
+ * @param[out] memoryRequirements Pointer to hipArrayMemoryRequirements
+ * @param[in] mipmap HIP mipmapped array to get the memory requirements of
+ * @param[in] device Device to get the memory requirements for
+ *
+ * @returns #hipSuccess, #hipErrorInvalidValue
+ *
+ * Returns the memory requirements of a HIP mipmapped array in memoryRequirements.
+ *
+ * The returned value in hipArrayMemoryRequirements::size represents the total size of the HIP
+ mipmapped array. The returned value in hipArrayMemoryRequirements::alignment represents the
+ alignment necessary for mapping the HIP mipmapped array.
+ */
+hipError_t hipMipmappedArrayGetMemoryRequirements(hipArrayMemoryRequirements* memoryRequirements,
+                                                  hipMipmappedArray_t mipmap, hipDevice_t device);
 // doxygen end Memory
 /**
  * @}
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
@@ -2070,8 +2070,8 @@ inline static hipError_t hipMemPrefetchAsync_v2(const void* dev_ptr, size_t coun
 inline static hipError_t hipMemAdvise_v2(const void* dev_ptr, size_t count, hipMemoryAdvise advice,
                                          hipMemLocation location) {
 #if CUDA_VERSION >= 13000
-  return hipCUDAErrorTohipError(cudaMemAdvise(dev_ptr, count,
-      hipMemoryAdviseTocudaMemoryAdvise(advice), location));
+  return hipCUDAErrorTohipError(
+      cudaMemAdvise(dev_ptr, count, hipMemoryAdviseTocudaMemoryAdvise(advice), location));
 #else
   return hipCUDAErrorTohipError(
       cudaMemAdvise_v2(dev_ptr, count, hipMemoryAdviseTocudaMemoryAdvise(advice), location));
@@ -2422,6 +2422,12 @@ inline static hipError_t hipMemcpy3DPeer(hipMemcpy3DPeerParms* p) {
 inline static hipError_t hipMemcpy3DPeerAsync(hipMemcpy3DPeerParms* p, hipStream_t stream) {
   return hipCUDAErrorTohipError(cudaMemcpy3DPeerAsync(p, stream));
 }
+inline static hipError_t hipMipmappedArrayGetMemoryRequirements(
+    hipArrayMemoryRequirements* memoryRequirements, hipMipmappedArray_t mipmap,
+    hipDevice_t device) {
+  return hipCUDAErrorTohipError(
+      cudaMipmappedArrayGetMemoryRequirements(memoryRequirements, mipmap, device));
+}
 
 __HIP_DEPRECATED inline static hipError_t hipMemcpyToArray(hipArray_t dst, size_t wOffset,
                                                            size_t hOffset, const void* src,
@@ -3303,7 +3309,7 @@ inline static hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* fla
   return hipCUDAErrorTohipError(cudaStreamGetFlags(stream, flags));
 }
 
-inline static hipError_t hipStreamGetId(hipStream_t stream, unsigned long long *streamId) {
+inline static hipError_t hipStreamGetId(hipStream_t stream, unsigned long long* streamId) {
   return hipCUDAErrorTohipError(cudaStreamGetId(stream, streamId));
 }
 
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
@@ -989,6 +989,7 @@ ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipLibraryGetKernelCount)
 #endif
 #if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 16
 ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamCopyAttributes)
+ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipMipmappedArrayGetMemoryRequirements)
 #endif
 #if HIP_RUNTIME_API_TABLE_STEP_VERSION == 0
 static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 442);
@@ -1023,7 +1024,7 @@ static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 496);
 #elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 15
 static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 501);
 #elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 16
-static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 502);
+static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 503);
 #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
@@ -3330,6 +3330,14 @@ typedef union rocprofiler_hip_api_args_t
         hipStream_t dst;
         hipStream_t src;
     } hipStreamCopyAttributes;
+    struct
+    {
+        size_t* dynamicSmemSize;
+        const void* f;
+        int numBlocks;
+        int blockSize;
+    } hipMipmappedArrayGetMemoryRequirements;
+    
 #endif
 } rocprofiler_hip_api_args_t;
 
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
diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/abi.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/abi.cpp
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

Original file line number	Diff line number	Diff line change
`@@ -147,6 +147,7 @@ set(TEST_SRC`
`147`	`147`	`tex2DLod.cc`
`148`	`148`	`tex2DLayeredLod.cc`
`149`	`149`	`tex3DLod.cc`
	`150`	`+ hipMipmappedArrayGetMemoryRequirements.cc`
`150`	`151`	`)`
`151`	`152`	`endif()`
`152`	`153`