From 576a6f6fbe4dcd47d000fc1279213e7f85415b32 Mon Sep 17 00:00:00 2001 From: Sebastian Luzynski Date: Fri, 3 Oct 2025 08:49:01 +0000 Subject: [PATCH] SWDEV-546322 Implement hipMipmappedArrayGetMemoryRequirements Signed-off-by: Sebastian Luzynski --- .../include/hip/amd_detail/hip_api_trace.hpp | 11 +++- .../include/hip/amd_detail/hip_prof_str.h | 29 +++++++++- projects/clr/hipamd/src/amdhip.def | 1 + projects/clr/hipamd/src/hip_api_trace.cpp | 11 +++- projects/clr/hipamd/src/hip_conversions.hpp | 7 ++- projects/clr/hipamd/src/hip_hcc.map.in | 1 + projects/clr/hipamd/src/hip_memory.cpp | 33 +++++++++++ .../clr/hipamd/src/hip_table_interface.cpp | 7 +++ .../hipMipmappedArrayGetMemoryRequirements.cc | 58 +++++++++++++++++++ projects/hip/include/hip/hip_runtime_api.h | 26 +++++++++ .../nvidia_detail/nvidia_hip_runtime_api.h | 12 +++- .../rocprofiler-sdk/cxx/enum_string.hpp | 6 ++ .../include/rocprofiler-sdk/hip/api_args.h | 8 +++ .../rocprofiler-sdk/hip/runtime_api_id.h | 3 + .../source/lib/rocprofiler-sdk/hip/abi.cpp | 6 ++ .../rocprofiler-sdk/hip/details/format.hpp | 9 +++ .../lib/rocprofiler-sdk/hip/hip.def.cpp | 3 + 17 files changed, 221 insertions(+), 10 deletions(-) create mode 100644 projects/hip-tests/catch/unit/texture/hipMipmappedArrayGetMemoryRequirements.cc 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