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 c5984f0aa72e..6c68a24a9d2d 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 22
+#define HIP_RUNTIME_API_TABLE_STEP_VERSION 23
// HIP API interface
// HIP compiler dispatch functions
@@ -1123,6 +1123,8 @@ typedef hipError_t (*t_hipMemSetMemPool)(hipMemLocation* location, hipMemAllocat
hipMemPool_t pool);
typedef hipError_t (*t_hipMemGetMemPool)(hipMemPool_t* pool, hipMemLocation* location,
hipMemAllocationType type);
+typedef hipError_t (*t_hipMipmappedArrayGetMemoryRequirements)(
+ hipArrayMemoryRequirements* memoryRequirements, hipMipmappedArray_t mipmap, hipDevice_t device);
// HIP Compiler dispatch table
struct HipCompilerDispatchTable {
// HIP_COMPILER_API_TABLE_STEP_VERSION == 0
@@ -1723,8 +1725,11 @@ struct HipDispatchTable {
t_hipMemSetMemPool hipMemSetMemPool_fn;
t_hipMemGetMemPool hipMemGetMemPool_fn;
- // DO NOT EDIT ABOVE!
// HIP_RUNTIME_API_TABLE_STEP_VERSION == 23
+ t_hipMipmappedArrayGetMemoryRequirements hipMipmappedArrayGetMemoryRequirements_fn;
+
+ // DO NOT EDIT ABOVE!
+ // HIP_RUNTIME_API_TABLE_STEP_VERSION == 24
// ******************************************************************************************* //
//
@@ -1761,4 +1766,4 @@ struct HipToolsDispatchTable {
// 4) GENERATE COMMENT FOR NEXT STEP VERSION
// 5) ADD "DO NOT EDIT ABOVE!" COMMENT
// ******************************************************************************************* //
-};
\ No newline at end of file
+};
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 f7f7083c8225..4537e23f8a39 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
@@ -473,7 +473,8 @@ enum hip_api_id_t {
HIP_API_ID_hipExtSetLoggingParams = 453,
HIP_API_ID_hipMemSetMemPool = 454,
HIP_API_ID_hipMemGetMemPool = 455,
- HIP_API_ID_LAST = 455,
+ HIP_API_ID_hipMipmappedArrayGetMemoryRequirements = 456,
+ HIP_API_ID_LAST = 456,
HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice),
HIP_API_ID_hipGetDeviceProperties = HIP_API_ID_CONCAT(HIP_API_ID_,hipGetDeviceProperties),
@@ -869,6 +870,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";
@@ -1318,6 +1320,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;
@@ -3581,6 +3584,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;
@@ -6300,6 +6309,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; \
@@ -8262,6 +8277,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);
@@ -11520,6 +11539,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
index d67b8ad745b5..5d5f7132205a 100644
--- a/projects/clr/hipamd/src/amdhip.def
+++ b/projects/clr/hipamd/src/amdhip.def
@@ -528,3 +528,4 @@ hipExtEnableLogging
hipExtSetLoggingParams
hipMemSetMemPool
hipMemGetMemPool
+hipMipmappedArrayGetMemoryRequirements
diff --git a/projects/clr/hipamd/src/hip_api_trace.cpp b/projects/clr/hipamd/src/hip_api_trace.cpp
index 355db7a511bb..5bc8695edd5c 100644
--- a/projects/clr/hipamd/src/hip_api_trace.cpp
+++ b/projects/clr/hipamd/src/hip_api_trace.cpp
@@ -891,6 +891,8 @@ hipError_t hipExtSetLoggingParams(size_t log_level, size_t log_size, size_t log_
hipError_t hipMemSetMemPool(hipMemLocation* location, hipMemAllocationType type, hipMemPool_t pool);
hipError_t hipMemGetMemPool(hipMemPool_t* pool, hipMemLocation* location,
hipMemAllocationType type);
+hipError_t hipMipmappedArrayGetMemoryRequirements(hipArrayMemoryRequirements* memoryRequirements,
+ hipMipmappedArray_t mipmap, hipDevice_t device);
} // namespace hip
namespace hip {
@@ -1235,6 +1237,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;
@@ -2132,15 +2136,18 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipExtSetLoggingParams_fn, 510);
// HIP_RUNTIME_API_TABLE_STEP_VERSION == 22
HIP_ENFORCE_ABI(HipDispatchTable, hipMemSetMemPool_fn, 511);
HIP_ENFORCE_ABI(HipDispatchTable, hipMemGetMemPool_fn, 512);
+// HIP_RUNTIME_API_TABLE_STEP_VERSION == 23
+HIP_ENFORCE_ABI(HipDispatchTable, hipMipmappedArrayGetMemoryRequirements_fn, 513);
+
// 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, 513)
+HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 514)
-static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 22,
+static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 23,
"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_conversions.hpp b/projects/clr/hipamd/src/hip_conversions.hpp
index 7a9f07344d96..7db85fdf5d02 100644
--- 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
\ No newline at end of file
+
+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
index e3c8147c2466..4ed786bd2f94 100644
--- a/projects/clr/hipamd/src/hip_hcc.map.in
+++ b/projects/clr/hipamd/src/hip_hcc.map.in
@@ -650,6 +650,7 @@ global:
hipExtSetLoggingParams;
hipMemSetMemPool;
hipMemGetMemPool;
+ hipMipmappedArrayGetMemoryRequirements;
local:
*;
} hip_7.1;
diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp
index 911d3256c6a6..a083140c91b1 100644
--- a/projects/clr/hipamd/src/hip_memory.cpp
+++ b/projects/clr/hipamd/src/hip_memory.cpp
@@ -4502,6 +4502,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(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) {
@@ -4526,6 +4552,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
index 3f9fcdba9896..b021f11a4a1d 100644
--- a/projects/clr/hipamd/src/hip_table_interface.cpp
+++ b/projects/clr/hipamd/src/hip_table_interface.cpp
@@ -3180,4 +3180,11 @@ hipError_t hipMemGetMemPool(hipMemPool_t* pool, hipMemLocation* location,
TRY;
return hip::GetHipDispatchTable()->hipMemGetMemPool_fn(pool, location, type);
CATCH;
+}
+hipError_t hipMipmappedArrayGetMemoryRequirements(hipArrayMemoryRequirements* memoryRequirements,
+ hipMipmappedArray_t mipmap, hipDevice_t device) {
+ TRY;
+ return hip::GetHipDispatchTable()->hipMipmappedArrayGetMemoryRequirements_fn(memoryRequirements,
+ mipmap, device);
+ CATCH;
}
\ No newline at end of file
diff --git a/projects/hip-tests/catch/unit/texture/hipMipmappedArrayGetMemoryRequirements.cc b/projects/hip-tests/catch/unit/texture/hipMipmappedArrayGetMemoryRequirements.cc
new file mode 100644
index 000000000000..f7365f5aa959
--- /dev/null
+++ 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
+#include
+
+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;
+ 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/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h
index 5148481b3d92..14660ec45877 100644
--- a/projects/hip/include/hip/hip_runtime_api.h
+++ b/projects/hip/include/hip/hip_runtime_api.h
@@ -2017,6 +2017,14 @@ typedef struct HIP_LAUNCH_CONFIG_st {
unsigned int numAttrs; ///< Number of attributes
} HIP_LAUNCH_CONFIG;
+/**
+ * Struct representing array memory requirements.
+ */
+typedef struct hipArrayMemoryRequirements {
+ size_t alignment;
+ size_t size;
+} hipArrayMemoryRequirements;
+
/**
* Requested handle type for address range.
*/
@@ -5844,6 +5852,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
index 3526b20000d2..7869eeefc997 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
@@ -2100,8 +2100,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));
@@ -2477,6 +2477,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,
@@ -3385,7 +3391,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
index 4be88eecd11c..d4ecd39ba7aa 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
@@ -1060,6 +1060,10 @@ ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipExtSetLoggingParams)
ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipMemSetMemPool)
ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipMemGetMemPool)
#endif
+#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 23
+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);
#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 1
@@ -1106,6 +1110,8 @@ static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 508);
static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 511);
#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 22
static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 513);
+#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 23
+static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 514);
#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 92a3f165513e..e4b8e5498d88 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
@@ -3411,6 +3411,14 @@ typedef union rocprofiler_hip_api_args_t
hipMemAllocationType type;
} hipMemGetMemPool;
#endif
+#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 23
+ struct
+ {
+ hipArrayMemoryRequirements* memoryRequirements;
+ hipMipmappedArray_t mipmap;
+ hipDevice_t device;
+ } hipMipmappedArrayGetMemoryRequirements;
+#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 68b9f56bcc22..f7b07dd9454b 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
@@ -584,6 +584,9 @@ typedef enum rocprofiler_hip_runtime_api_id_t // NOLINT(performance-enum-size)
#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 22
ROCPROFILER_HIP_RUNTIME_API_ID_hipMemSetMemPool,
ROCPROFILER_HIP_RUNTIME_API_ID_hipMemGetMemPool,
+#endif
+#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 23
+ ROCPROFILER_HIP_RUNTIME_API_ID_hipMipmappedArrayGetMemoryRequirements,
#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 89998681d7ce..e943a43d5ab1 100644
--- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/abi.cpp
+++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/abi.cpp
@@ -633,6 +633,10 @@ ROCP_SDK_ENFORCE_ABI(::HipDispatchTable, hipMemSetMemPool_fn, 511);
ROCP_SDK_ENFORCE_ABI(::HipDispatchTable, hipMemGetMemPool_fn, 512);
#endif
+#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 23
+ROCP_SDK_ENFORCE_ABI(::HipDispatchTable, hipMipmappedArrayGetMemoryRequirements_fn, 513);
+#endif
+
#if HIP_RUNTIME_API_TABLE_STEP_VERSION == 0
ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 442)
#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 1
@@ -679,6 +683,8 @@ ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 508)
ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 511)
#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 22
ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 513)
+#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 23
+ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 514)
#else
INTERNAL_CI_ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 0)
#endif
diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/details/format.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/details/format.hpp
index 98efeb6b2441..d32958be2092 100644
--- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/details/format.hpp
+++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/details/format.hpp
@@ -630,6 +630,15 @@ struct formatter : rocprofiler::hip::details::ba
}
};
#endif
+
+#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 23
+ROCP_SDK_HIP_FORMATTER(hipArrayMemoryRequirements,
+ "{}alignment={}, size={}{}",
+ '{',
+ v.alignment,
+ v.size,
+ '}')
+#endif
} // namespace fmt
#undef ROCP_SDK_HIP_FORMATTER
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 82b119e53f5c..c987584ce130 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
@@ -660,6 +660,9 @@ HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNT
HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNTIME_API_ID_hipMemGetMemPool, hipMemGetMemPool, hipMemGetMemPool_fn, pool, location, type);
#endif
+#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 23
+HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNTIME_API_ID_hipMipmappedArrayGetMemoryRequirements, hipMipmappedArrayGetMemoryRequirements, hipMipmappedArrayGetMemoryRequirements_fn, memoryRequirements, mipmap, device);
+#endif
// clang-format on
#else