From f13155426d340e1a6e1f11bca1d0dcabc0ad2c53 Mon Sep 17 00:00:00 2001 From: Mike Qin Date: Thu, 20 Aug 2026 13:57:32 -0400 Subject: [PATCH 1/6] add changelog --- projects/clr/CHANGELOG.md | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/projects/clr/CHANGELOG.md b/projects/clr/CHANGELOG.md index ff83537c71a8..d4691959fce1 100644 --- a/projects/clr/CHANGELOG.md +++ b/projects/clr/CHANGELOG.md @@ -54,6 +54,15 @@ Copy operations at or below the 256-row threshold are unchanged. - Streamlined batch grouping: * Removed the `AgentGroup/src_agent` mapping for D2D broadcasts. * Processed `H2D` and `D2H` LINEAR operations directly, bypassing the broadcast map. + +## HIP 8.0 for ROCm 8.0 + +### Added +* An optional `HIP_FORCE_API_VERSION` macro can be used to select an older version of the HIP APIs. For example, `HIP_FORCE_API_VERSION=600` selects the HIP 6.0 APIs. This macro must be defined before including `hip_runtime_api.h`. + +### Deprecated + +* `hipMemAdvise` is now an alias for `hipMemAdvise_v2`. To use the previous version of `hipMemAdvise`, define `HIP_FORCE_API_VERSION` as a non-zero value less than 800. ## HIP 7.14 for ROCm 7.14 From 4d0c075c1db4304ff0faa3c1ecca73a146ed17f8 Mon Sep 17 00:00:00 2001 From: Mike Qin Date: Tue, 25 Aug 2026 11:46:11 -0400 Subject: [PATCH 2/6] use wrapper function to version the APIs Do not pollute sources with #define macros. This commit no longer relies on macro pollution in existing clr code base. Also migrate the header generation script into Python3. Traversal order of Python dictionary has changed so the order is slightly different than before. Introduce HIP_ABI_IMPL for clr, hrr and hip-tests. --- projects/clr/CMakeLists.txt | 2 +- .../include/hip/amd_detail/hip_api_trace.hpp | 2 +- .../include/hip/amd_detail/hip_prof_str.h | 66 ++++----- projects/clr/hipamd/src/CMakeLists.txt | 16 +- projects/clr/hipamd/src/hip_api_trace.cpp | 4 +- projects/clr/hipamd/src/hip_comgr_helper.cpp | 2 +- .../clr/hipamd/src/hip_deprecated_abi.cpp | 29 ++++ projects/clr/hipamd/src/hip_device.cpp | 12 +- .../clr/hipamd/src/hip_device_runtime.cpp | 7 - projects/clr/hipamd/src/hip_platform.cpp | 12 +- projects/clr/hipamd/src/hip_prof_gen.py | 51 ++++--- .../clr/hipamd/src/hip_table_interface.cpp | 5 +- .../hipamd/src/hrr/hip_capture_generated.cpp | 2 +- .../hipamd/src/hrr/hip_capture_metadata.cpp | 1 + .../profiler/hip_clr_dispatch_wrappers.cpp | 4 +- .../hip-tests/catch/cmake/hip-tests.cmake | 5 + projects/hip/docs/doxygen/Doxyfile | 1 + projects/hip/include/hip/hip_common.h | 6 + projects/hip/include/hip/hip_runtime_api.h | 137 +++++++++++++++--- projects/hrr/playback/CMakeLists.txt | 2 + 20 files changed, 252 insertions(+), 114 deletions(-) create mode 100644 projects/clr/hipamd/src/hip_deprecated_abi.cpp diff --git a/projects/clr/CMakeLists.txt b/projects/clr/CMakeLists.txt index 14c7f6e79d51..5cd6eb65cffb 100644 --- a/projects/clr/CMakeLists.txt +++ b/projects/clr/CMakeLists.txt @@ -2,7 +2,7 @@ # # SPDX-License-Identifier: MIT -cmake_minimum_required(VERSION 3.16.8) +cmake_minimum_required(VERSION 3.27) project(clr) ########## 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 12343a427e14..9796873a4319 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 @@ -99,7 +99,7 @@ typedef hipError_t (*t_hipBindTextureToArray)(const textureReference* tex, hipAr typedef hipError_t (*t_hipBindTextureToMipmappedArray)(const textureReference* tex, hipMipmappedArray_const_t mipmappedArray, const hipChannelFormatDesc* desc); -typedef hipError_t (*t_hipChooseDevice)(int* device, const hipDeviceProp_t* prop); +typedef hipError_t (*t_hipChooseDevice)(int* device, const hipDeviceProp_tR0600* prop); typedef hipError_t (*t_hipChooseDeviceR0000)(int* device, const hipDeviceProp_tR0000* properties); typedef hipError_t (*t_hipConfigureCall)(dim3 gridDim, dim3 blockDim, size_t sharedMem, hipStream_t stream); 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 5c3c85109722..d6d54777ed8b 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 @@ -16,9 +16,6 @@ #include #include "amd_hip_gl_interop.h" -#define HIP_API_ID_CONCAT_HELPER(a,b) a##b -#define HIP_API_ID_CONCAT(a,b) HIP_API_ID_CONCAT_HELPER(a,b) - // HIP API callbacks ID enumeration enum hip_api_id_t { HIP_API_ID_NONE = 0, @@ -511,8 +508,6 @@ enum hip_api_id_t { HIP_API_ID_hipInitDevice = 486, HIP_API_ID_LAST = 486, - HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice), - HIP_API_ID_hipGetDeviceProperties = HIP_API_ID_CONCAT(HIP_API_ID_,hipGetDeviceProperties), HIP_API_ID_hipBindTexture = HIP_API_ID_NONE, HIP_API_ID_hipBindTexture2D = HIP_API_ID_NONE, @@ -541,9 +536,6 @@ enum hip_api_id_t { HIP_API_ID_hipUnbindTexture = HIP_API_ID_NONE, }; -#undef HIP_API_ID_CONCAT_HELPER -#undef HIP_API_ID_CONCAT - // Return the HIP API string for a given callback ID static inline const char* hip_api_name(const uint32_t id) { switch(id) { @@ -594,6 +586,7 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipDeviceGetExecutionCtx: return "hipDeviceGetExecutionCtx"; case HIP_API_ID_hipDeviceGetGraphMemAttribute: return "hipDeviceGetGraphMemAttribute"; case HIP_API_ID_hipDeviceGetLimit: return "hipDeviceGetLimit"; + case HIP_API_ID_hipDeviceGetLuid: return "hipDeviceGetLuid"; case HIP_API_ID_hipDeviceGetMemPool: return "hipDeviceGetMemPool"; case HIP_API_ID_hipDeviceGetName: return "hipDeviceGetName"; case HIP_API_ID_hipDeviceGetP2PAttribute: return "hipDeviceGetP2PAttribute"; @@ -601,7 +594,6 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipDeviceGetSharedMemConfig: return "hipDeviceGetSharedMemConfig"; case HIP_API_ID_hipDeviceGetStreamPriorityRange: return "hipDeviceGetStreamPriorityRange"; case HIP_API_ID_hipDeviceGetUuid: return "hipDeviceGetUuid"; - case HIP_API_ID_hipDeviceGetLuid: return "hipDeviceGetLuid"; case HIP_API_ID_hipDeviceGraphMemTrim: return "hipDeviceGraphMemTrim"; case HIP_API_ID_hipDevicePrimaryCtxGetState: return "hipDevicePrimaryCtxGetState"; case HIP_API_ID_hipDevicePrimaryCtxRelease: return "hipDevicePrimaryCtxRelease"; @@ -1074,6 +1066,7 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipDeviceGetExecutionCtx", name) == 0) return HIP_API_ID_hipDeviceGetExecutionCtx; if (strcmp("hipDeviceGetGraphMemAttribute", name) == 0) return HIP_API_ID_hipDeviceGetGraphMemAttribute; if (strcmp("hipDeviceGetLimit", name) == 0) return HIP_API_ID_hipDeviceGetLimit; + if (strcmp("hipDeviceGetLuid", name) == 0) return HIP_API_ID_hipDeviceGetLuid; if (strcmp("hipDeviceGetMemPool", name) == 0) return HIP_API_ID_hipDeviceGetMemPool; if (strcmp("hipDeviceGetName", name) == 0) return HIP_API_ID_hipDeviceGetName; if (strcmp("hipDeviceGetP2PAttribute", name) == 0) return HIP_API_ID_hipDeviceGetP2PAttribute; @@ -1081,7 +1074,6 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipDeviceGetSharedMemConfig", name) == 0) return HIP_API_ID_hipDeviceGetSharedMemConfig; if (strcmp("hipDeviceGetStreamPriorityRange", name) == 0) return HIP_API_ID_hipDeviceGetStreamPriorityRange; if (strcmp("hipDeviceGetUuid", name) == 0) return HIP_API_ID_hipDeviceGetUuid; - if (strcmp("hipDeviceGetLuid", name) == 0) return HIP_API_ID_hipDeviceGetLuid; if (strcmp("hipDeviceGraphMemTrim", name) == 0) return HIP_API_ID_hipDeviceGraphMemTrim; if (strcmp("hipDevicePrimaryCtxGetState", name) == 0) return HIP_API_ID_hipDevicePrimaryCtxGetState; if (strcmp("hipDevicePrimaryCtxRelease", name) == 0) return HIP_API_ID_hipDevicePrimaryCtxRelease; @@ -1746,6 +1738,13 @@ typedef struct hip_api_data_s { size_t pValue__val; enum hipLimit_t limit; } hipDeviceGetLimit; + struct { + char* luid; + char luid__val; + unsigned int* deviceNodeMask; + unsigned int deviceNodeMask__val; + hipDevice_t device; + } hipDeviceGetLuid; struct { hipMemPool_t* mem_pool; hipMemPool_t mem_pool__val; @@ -1785,13 +1784,6 @@ typedef struct hip_api_data_s { hipUUID uuid__val; hipDevice_t device; } hipDeviceGetUuid; - struct { - char* luid; - char luid__val; - unsigned int* deviceNodeMask; - unsigned int deviceNodeMask__val; - hipDevice_t device; - } hipDeviceGetLuid; struct { int device; } hipDeviceGraphMemTrim; @@ -4657,6 +4649,12 @@ typedef struct hip_api_data_s { cb_data.args.hipDeviceGetLimit.pValue = (size_t*)pValue; \ cb_data.args.hipDeviceGetLimit.limit = (hipLimit_t)limit; \ }; +// hipDeviceGetLuid[('char*', 'luid'), ('unsigned int*', 'deviceNodeMask'), ('hipDevice_t', 'device')] +#define INIT_hipDeviceGetLuid_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipDeviceGetLuid.luid = (char*)luid; \ + cb_data.args.hipDeviceGetLuid.deviceNodeMask = (unsigned int*)deviceNodeMask; \ + cb_data.args.hipDeviceGetLuid.device = (hipDevice_t)device; \ +}; // hipDeviceGetMemPool[('hipMemPool_t*', 'mem_pool'), ('int', 'device')] #define INIT_hipDeviceGetMemPool_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipDeviceGetMemPool.mem_pool = (hipMemPool_t*)mem_pool; \ @@ -4695,12 +4693,6 @@ typedef struct hip_api_data_s { cb_data.args.hipDeviceGetUuid.uuid = (hipUUID*)uuid; \ cb_data.args.hipDeviceGetUuid.device = (hipDevice_t)device; \ }; -// hipDeviceGetLuid[('char*', 'luid'), ('unsigned int*', 'deviceNodeMask'), ('hipDevice_t', 'device')] -#define INIT_hipDeviceGetLuid_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDeviceGetLuid.luid = (char*)luid; \ - cb_data.args.hipDeviceGetLuid.deviceNodeMask = (unsigned int*)deviceNodeMask; \ - cb_data.args.hipDeviceGetLuid.device = (hipDevice_t)device; \ -}; // hipDeviceGraphMemTrim[('int', 'device')] #define INIT_hipDeviceGraphMemTrim_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipDeviceGraphMemTrim.device = (int)device; \ @@ -7579,6 +7571,11 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { case HIP_API_ID_hipDeviceGetLimit: if (data->args.hipDeviceGetLimit.pValue) data->args.hipDeviceGetLimit.pValue__val = *(data->args.hipDeviceGetLimit.pValue); break; +// hipDeviceGetLuid[('char*', 'luid'), ('unsigned int*', 'deviceNodeMask'), ('hipDevice_t', 'device')] + case HIP_API_ID_hipDeviceGetLuid: + data->args.hipDeviceGetLuid.luid = (data->args.hipDeviceGetLuid.luid) ? strdup(data->args.hipDeviceGetLuid.luid) : NULL; + if (data->args.hipDeviceGetLuid.deviceNodeMask) data->args.hipDeviceGetLuid.deviceNodeMask__val = *(data->args.hipDeviceGetLuid.deviceNodeMask); + break; // hipDeviceGetMemPool[('hipMemPool_t*', 'mem_pool'), ('int', 'device')] case HIP_API_ID_hipDeviceGetMemPool: if (data->args.hipDeviceGetMemPool.mem_pool) data->args.hipDeviceGetMemPool.mem_pool__val = *(data->args.hipDeviceGetMemPool.mem_pool); @@ -7608,11 +7605,6 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { case HIP_API_ID_hipDeviceGetUuid: if (data->args.hipDeviceGetUuid.uuid) data->args.hipDeviceGetUuid.uuid__val = *(data->args.hipDeviceGetUuid.uuid); break; -// hipDeviceGetLuid[('char*', 'luid'), ('unsigned int*', 'deviceNodeMask'), ('hipDevice_t', 'device')] - case HIP_API_ID_hipDeviceGetLuid: - if (data->args.hipDeviceGetLuid.luid) data->args.hipDeviceGetLuid.luid__val = *(data->args.hipDeviceGetLuid.luid); - if (data->args.hipDeviceGetLuid.deviceNodeMask) data->args.hipDeviceGetLuid.deviceNodeMask__val = *(data->args.hipDeviceGetLuid.deviceNodeMask); - break; // hipDeviceGraphMemTrim[('int', 'device')] case HIP_API_ID_hipDeviceGraphMemTrim: break; @@ -9629,6 +9621,15 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << ", limit="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDeviceGetLimit.limit); oss << ")"; break; + case HIP_API_ID_hipDeviceGetLuid: + oss << "hipDeviceGetLuid("; + if (data->args.hipDeviceGetLuid.luid == NULL) oss << "luid=NULL"; + else { oss << "luid="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDeviceGetLuid.luid__val); } + if (data->args.hipDeviceGetLuid.deviceNodeMask == NULL) oss << ", deviceNodeMask=NULL"; + else { oss << ", deviceNodeMask="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDeviceGetLuid.deviceNodeMask__val); } + oss << ", device="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDeviceGetLuid.device); + oss << ")"; + break; case HIP_API_ID_hipDeviceGetMemPool: oss << "hipDeviceGetMemPool("; if (data->args.hipDeviceGetMemPool.mem_pool == NULL) oss << "mem_pool=NULL"; @@ -9682,15 +9683,6 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << ", device="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDeviceGetUuid.device); oss << ")"; break; - case HIP_API_ID_hipDeviceGetLuid: - oss << "hipDeviceGetLuid("; - if (data->args.hipDeviceGetLuid.luid == NULL) oss << "luid=NULL"; - else { oss << "luid="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDeviceGetLuid.luid__val); } - if (data->args.hipDeviceGetLuid.deviceNodeMask == NULL) oss << ", deviceNodeMask=NULL"; - else { oss << ", deviceNodeMask="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDeviceGetLuid.deviceNodeMask__val); } - oss << ", device="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDeviceGetLuid.device); - oss << ")"; - break; case HIP_API_ID_hipDeviceGraphMemTrim: oss << "hipDeviceGraphMemTrim("; oss << "device="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDeviceGraphMemTrim.device); diff --git a/projects/clr/hipamd/src/CMakeLists.txt b/projects/clr/hipamd/src/CMakeLists.txt index b83ab78b45e8..236edcadea84 100644 --- a/projects/clr/hipamd/src/CMakeLists.txt +++ b/projects/clr/hipamd/src/CMakeLists.txt @@ -125,6 +125,7 @@ target_sources(amdhip64 PRIVATE hip_api_trace.cpp hip_table_interface.cpp hip_table_interface_c.cpp + hip_deprecated_abi.cpp hip_comgr_helper.cpp hip_library.cpp profiler/hip_clr_profiler.cpp @@ -173,6 +174,8 @@ target_include_directories(amdhip64 ${PROJECT_BINARY_DIR}/include) target_compile_definitions(amdhip64 PRIVATE __HIP_PLATFORM_AMD__) + +target_compile_definitions(amdhip64 PRIVATE HIP_ABI_IMPL) target_link_libraries(amdhip64 PRIVATE ${OPENGL_LIBRARIES}) target_link_libraries(amdhip64 PRIVATE ${CMAKE_DL_LIBS}) # Add link to comgr, hsa-runtime and other required libraries in target files @@ -264,16 +267,21 @@ if(USE_PROF_API) ") endif() + + if(MSVC) + set(pp_out /P /Fi:${PROF_API_NEWHDR}.i) + else() + set(pp_out -E -o ${PROF_API_NEWHDR}.i) # note: -E before input + endif() + add_custom_command(OUTPUT ${PROF_API_NEWHDR}.i COMMAND ${CMAKE_COMMAND} -E cat ${PROF_API_HDR} ${PROF_GL_HDR} > ${PROF_API_NEWHDR} COMMAND ${CMAKE_C_COMPILER} "-D$,;-D>" - "-I$,;-I>" + "$,PREPEND,-I>" # Directory name may contain spaces. "-DHIP_INCLUDE_HIP_HIP_RUNTIME_PT_API_H=1" - ${c_flags} $ - ${CPP_EXTRA_C_FLAGS} - -E ${PROF_API_NEWHDR} -o ${PROF_API_NEWHDR}.i + ${pp_out} ${PROF_API_NEWHDR} COMMAND_EXPAND_LISTS VERBATIM IMPLICIT_DEPENDS C ${PROF_API_HDR} ${PROF_GL_HDR} ${PROF_API_DEPRECATED} DEPENDS ${PROF_API_HDR} ${PROF_GL_HDR} ${PROF_API_DEPRECATED} diff --git a/projects/clr/hipamd/src/hip_api_trace.cpp b/projects/clr/hipamd/src/hip_api_trace.cpp index cc48f9f08ae1..90397a537496 100644 --- a/projects/clr/hipamd/src/hip_api_trace.cpp +++ b/projects/clr/hipamd/src/hip_api_trace.cpp @@ -64,7 +64,7 @@ hipError_t hipBindTextureToArray(const textureReference* tex, hipArray_const_t a hipError_t hipBindTextureToMipmappedArray(const textureReference* tex, hipMipmappedArray_const_t mipmappedArray, const hipChannelFormatDesc* desc); -hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* prop); +hipError_t hipChooseDeviceR0600(int* device, const hipDeviceProp_tR0600* prop); hipError_t hipChooseDeviceR0000(int* device, const hipDeviceProp_tR0000* properties); hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, hipStream_t stream); hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResourceDesc* pResDesc, @@ -975,7 +975,7 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) { ptrDispatchTable->hipBindTexture2D_fn = hip::hipBindTexture2D; ptrDispatchTable->hipBindTextureToArray_fn = hip::hipBindTextureToArray; ptrDispatchTable->hipBindTextureToMipmappedArray_fn = hip::hipBindTextureToMipmappedArray; - ptrDispatchTable->hipChooseDevice_fn = hip::hipChooseDevice; + ptrDispatchTable->hipChooseDevice_fn = hip::hipChooseDeviceR0600; ptrDispatchTable->hipChooseDeviceR0000_fn = hip::hipChooseDeviceR0000; ptrDispatchTable->hipConfigureCall_fn = hip::hipConfigureCall; ptrDispatchTable->hipCreateSurfaceObject_fn = hip::hipCreateSurfaceObject; diff --git a/projects/clr/hipamd/src/hip_comgr_helper.cpp b/projects/clr/hipamd/src/hip_comgr_helper.cpp index 5054eddc17e4..aaf43a3d4c9f 100644 --- a/projects/clr/hipamd/src/hip_comgr_helper.cpp +++ b/projects/clr/hipamd/src/hip_comgr_helper.cpp @@ -1109,7 +1109,7 @@ bool LinkProgram::LinkComplete(void** bin_out, size_t* size_out) { if (!findIsa()) { return false; } - + hip::comgr_helper::ComgrDataSetUniqueHandle link_output; if (link_output.Create() != AMD_COMGR_STATUS_SUCCESS) { return false; diff --git a/projects/clr/hipamd/src/hip_deprecated_abi.cpp b/projects/clr/hipamd/src/hip_deprecated_abi.cpp new file mode 100644 index 000000000000..ab70a53899d3 --- /dev/null +++ b/projects/clr/hipamd/src/hip_deprecated_abi.cpp @@ -0,0 +1,29 @@ +/* + * Copyright (c) Advanced Micro Devices, Inc., or its affiliates. + * + * SPDX-License-Identifier: MIT + */ + +// These symbols are no longer provided as APIs, but they are supported as part +// of the ABI. If the application insists on using them, they need to call the +// R0000 version. + +// We name this file .cpp, so that we can use CXXFLAGS. +extern "C" { + +typedef int hipDevice_t; +typedef struct hipDeviceProp_tR0000 hipDeviceProp_tR0000; +typedef int hipError_t; + +hipError_t hipGetDevicePropertiesR0000(hipDeviceProp_tR0000* props, hipDevice_t device); +hipError_t hipChooseDeviceR0000(int* device, const hipDeviceProp_tR0000* properties); + +hipError_t hipGetDeviceProperties(hipDeviceProp_tR0000* props, hipDevice_t device) { + return hipGetDevicePropertiesR0000(props, device); +} + +hipError_t hipChooseDevice(int* device, const hipDeviceProp_tR0000* properties) { + return hipChooseDeviceR0000(device, properties); +} + +} diff --git a/projects/clr/hipamd/src/hip_device.cpp b/projects/clr/hipamd/src/hip_device.cpp index 6bceb17afb45..ea82c7be48e0 100644 --- a/projects/clr/hipamd/src/hip_device.cpp +++ b/projects/clr/hipamd/src/hip_device.cpp @@ -13,9 +13,6 @@ #include "hip_mempool_impl.hpp" #include "hip_platform.hpp" -#undef hipGetDeviceProperties -#undef hipDeviceProp_t - namespace hip { // ================================================================================================ @@ -627,7 +624,7 @@ hipError_t hipDeviceGetLuid(char* luid, unsigned int* deviceNodeMask, hipDevice_ } // ================================================================================================ -hipError_t ihipGetDeviceProperties(hipDeviceProp_tR0600* props, int device) { +hipError_t ihipGetDeviceProperties(hipDeviceProp_t* props, int device) { if (props == nullptr) { return hipErrorInvalidValue; } @@ -831,6 +828,8 @@ hipError_t ihipGetDeviceProperties(hipDeviceProp_tR0600* props, int device) { hipError_t hipGetDevicePropertiesR0600(hipDeviceProp_tR0600* prop, int device) { HIP_INIT_API(hipGetDevicePropertiesR0600, prop, device); + // For now, we can forward to the internal/unstable API. In the future, if + // hipDeviceProp_t is different, this will fail to compile. HIP_RETURN(ihipGetDeviceProperties(prop, device)); } @@ -1026,8 +1025,3 @@ hipError_t hipGetProcAddress_spt(const char* symbol, void** pfn, int hipVersion, } } // namespace hip - -// ================================================================================================ -extern "C" hipError_t hipGetDeviceProperties(hipDeviceProp_tR0000* props, hipDevice_t device) { - return hip::hipGetDevicePropertiesR0000(props, device); -} diff --git a/projects/clr/hipamd/src/hip_device_runtime.cpp b/projects/clr/hipamd/src/hip_device_runtime.cpp index 5cb9c55235af..bdcf7e0dad0d 100644 --- a/projects/clr/hipamd/src/hip_device_runtime.cpp +++ b/projects/clr/hipamd/src/hip_device_runtime.cpp @@ -9,9 +9,6 @@ #include "hip_internal.hpp" #include "hip_platform.hpp" -#undef hipChooseDevice -#undef hipDeviceProp_t - namespace hip { hipError_t hipGetDevicePropertiesR0000(hipDeviceProp_tR0000* prop, int device); @@ -919,7 +916,3 @@ hipError_t hipSetValidDevices(int* device_arr, int len) { HIP_RETURN(hipSuccess); } } // namespace hip - -extern "C" hipError_t hipChooseDevice(int* device, const hipDeviceProp_tR0000* properties) { - return hip::hipChooseDeviceR0000(device, properties); -} diff --git a/projects/clr/hipamd/src/hip_platform.cpp b/projects/clr/hipamd/src/hip_platform.cpp index 89fe79784687..1c182f350625 100644 --- a/projects/clr/hipamd/src/hip_platform.cpp +++ b/projects/clr/hipamd/src/hip_platform.cpp @@ -27,7 +27,7 @@ hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor( const auto* wrkGrpInfo = kernel->getDeviceKernel(device)->workGroupInfo(); const int maxWorkGroupSize = static_cast(device.info().maxWorkGroupSize_); - + if (!bCalcPotentialBlkSz) { if (inputBlockSize <= 0) { return hipErrorInvalidValue; @@ -41,7 +41,7 @@ hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor( } else if (inputBlockSize > maxWorkGroupSize || inputBlockSize <= 0) { inputBlockSize = maxWorkGroupSize; } - + // Find wave occupancy per CU => simd_per_cu * GPR usage // Limited by SPI 32 per CU, hence 8 per SIMD const size_t MaxWavesPerSimd = (device.isa().versionMajor() <= 9) ? 8 : 16; @@ -214,7 +214,7 @@ void __hipRegisterFunction(void** modules, const void* hostFunction, char* devic const char* deviceName, unsigned int threadLimit, uint3* tid, uint3* bid, dim3* blockDim, dim3* gridDim, int* wSize) { auto* fat_binary_modules = reinterpret_cast(modules); - + static const bool enable_deferred_loading = []() { const char* var = getenv("HIP_ENABLE_DEFERRED_LOADING"); return var ? atoi(var) != 0 : true; @@ -231,7 +231,7 @@ void __hipRegisterFunction(void** modules, const void* hostFunction, char* devic if (!enable_deferred_loading) { HIP_INIT_VOID(); - + for (size_t dev_idx = 0; dev_idx < g_devices.size(); ++dev_idx) { hipFunction_t hfunc = nullptr; hipError_t hip_error = platform.StatCO().GetFunc(&hfunc, hostFunction, dev_idx); @@ -479,7 +479,7 @@ hipError_t ihipCreateGlobalVarObj(const char* name, hipModule_t hmod, amd::Memor LogPrintfError("Cannot get Device Function for module: 0x%x", hmod); HIP_RETURN(hipErrorInvalidDeviceFunction); } - + // Find the global Symbols if (!dev_program->createGlobalVarObj(amd_mem_obj, dptr, bytes, name)) { LogPrintfError("Cannot create Global Var obj for symbol: %s", name); @@ -524,7 +524,7 @@ hipError_t hipOccupancyAvailableDynamicSMemPerBlock(size_t* dynamicSmemSize, con const int staticSharedMemoryUsage = wrkGrpInfo->usedLDSSize_; const int maxDynamicSharedSizeBytes = wrkGrpInfo->maxDynamicSharedSizeBytes_; const int maxNumBlocks = prop.maxThreadsPerMultiProcessor / blockSize; - const int maxSharedMemoryPerMultiProcessor = prop.maxSharedMemoryPerMultiProcessor - + const int maxSharedMemoryPerMultiProcessor = prop.maxSharedMemoryPerMultiProcessor - staticSharedMemoryUsage * std::min(numBlocks, maxNumBlocks); const int maxDynamicSmemSize = std::min(maxSharedMemoryPerMultiProcessor / maxNumBlocks, maxDynamicSharedSizeBytes); diff --git a/projects/clr/hipamd/src/hip_prof_gen.py b/projects/clr/hipamd/src/hip_prof_gen.py index 1110e76384e9..2f577de68c83 100755 --- a/projects/clr/hipamd/src/hip_prof_gen.py +++ b/projects/clr/hipamd/src/hip_prof_gen.py @@ -55,8 +55,8 @@ def filtr_api_name(name): return name def filtr_api_decl(record): - record = re.sub("\s__dparm\([^\)]*\)", r'', record); - record = re.sub("\(void\*\)", r'', record); + record = re.sub(r"\s__dparm\([^\)]*\)", r'', record); + record = re.sub(r"\(void\*\)", r'', record); return record # Normalizing API arguments @@ -83,7 +83,7 @@ def list_api_args(args_str): for arg_pair in args_str.split(','): if arg_pair == 'void': continue arg_pair = re.sub(r'\s*=\s*\S+$','', arg_pair); - m = re.match("^(.*)\s(\S+)$", arg_pair); + m = re.match(r"^(.*)\s(\S+)$", arg_pair); if m: arg_type = norm_api_types(m.group(1)) arg_name = m.group(2) @@ -127,9 +127,9 @@ def parse_api(inp_file_p, out): global line_num inp_file = inp_file_p - beg_pattern = re.compile("^(hipError_t|const char\s*\*)\s+([^\(]+)\("); - api_pattern = re.compile("^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)"); - end_pattern = re.compile("Texture"); + beg_pattern = re.compile(r"^(hipError_t|const char\s*\*)\s+([^\(]+)\("); + api_pattern = re.compile(r"^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)"); + end_pattern = re.compile(r"Texture"); hidden_pattern = re.compile(r'__attribute__\(\(visibility\("hidden"\)\)\)') nms_open_pattern = re.compile(r'namespace hip_impl {') nms_close_pattern = re.compile(r'}') @@ -161,7 +161,7 @@ def parse_api(inp_file_p, out): found = 1 if found != 0: - record = re.sub("\s__dparm\([^\)]*\)", '', record); + record = re.sub(r"\s__dparm\([^\)]*\)", '', record); m = api_pattern.match(record) if m: found = 0 @@ -200,13 +200,13 @@ def parse_content(inp_file_p, api_map, out): inp_file = inp_file_p # API method begin pattern - beg_pattern = re.compile("^(hipError_t|const char\s*\*)\s+[^\(]+\("); + beg_pattern = re.compile(r"^(hipError_t|const char\s*\*)\s+[^\(]+\("); # API declaration pattern - decl_pattern = re.compile("^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)\s*;"); + decl_pattern = re.compile(r"^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)\s*;"); # API definition pattern - api_pattern = re.compile("^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)\s*{"); + api_pattern = re.compile(r"^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)\s*{"); # API init macro pattern - init_pattern = re.compile("(^\s*HIP_INIT_API[^\s]*\s*)\((([^,]+)(,.*|)|)(\);|,)\s*$"); + init_pattern = re.compile(r"(^\s*HIP_INIT_API[^\s]*\s*)\((([^,]+)(,.*|)|)(\);|,)\s*$"); # Open input file inp = open(inp_file, 'r') @@ -401,8 +401,8 @@ def generate_prof_header(f, api_map, callback_ids, opts_map): priv_lst.append(name) message("Private: " + name) - f.write('\n#define HIP_API_ID_CONCAT_HELPER(a,b) a##b\n'); - f.write('#define HIP_API_ID_CONCAT(a,b) HIP_API_ID_CONCAT_HELPER(a,b)\n'); + # f.write('\n#define HIP_API_ID_CONCAT_HELPER(a,b) a##b\n'); + # f.write('#define HIP_API_ID_CONCAT(a,b) HIP_API_ID_CONCAT_HELPER(a,b)\n'); # Generating the callbacks ID enumaration f.write('\n// HIP API callbacks ID enumeration\n') @@ -412,7 +412,11 @@ def generate_prof_header(f, api_map, callback_ids, opts_map): cb_id_map = {} last_cb_id = 0 - versioned_functions = set() + + # In the past, we rely on macro pollutions and `HIP_API_ID_CONCAT()` + # to version the symbols, we no longer need that. + # + # versioned_functions = set() for name, cb_id in callback_ids: if not name in api_map: f.write(' HIP_API_ID_RESERVED_' + str(cb_id) + ' = ' + str(cb_id) + ',\n') @@ -420,29 +424,30 @@ def generate_prof_header(f, api_map, callback_ids, opts_map): f.write(' HIP_API_ID_' + name + ' = ' + str(cb_id) + ',\n') cb_id_map[name] = cb_id if cb_id > last_cb_id: last_cb_id = cb_id - m = re.match(r'(.*)R[0-9][0-9][0-9][0-9]$', name) - if m: versioned_functions.add(m.group(1)) + # m = re.match(r'(.*)R[0-9][0-9][0-9][0-9]$', name) + #if m: versioned_functions.add(m.group(1)) for name in sorted(api_map.keys()): if not name in cb_id_map: last_cb_id += 1 f.write(' HIP_API_ID_' + name + ' = ' + str(last_cb_id) + ',\n') - m = re.match(r'(.*)R[0-9][0-9][0-9][0-9]$', name) - if m: versioned_functions.add(m.group(1)) + # m = re.match(r'(.*)R[0-9][0-9][0-9][0-9]$', name) + # if m: versioned_functions.add(m.group(1)) f.write(' HIP_API_ID_LAST = ' + str(last_cb_id) + ',\n') f.write('\n') - for name in sorted(versioned_functions): - f.write(' HIP_API_ID_' + name + ' = ' + 'HIP_API_ID_CONCAT(HIP_API_ID_,' + name + '),\n') + # for name in sorted(versioned_functions): + # f.write(' HIP_API_ID_' + name + ' = ' + 'HIP_API_ID_CONCAT(HIP_API_ID_,' + name + '),\n') + f.write('\n') for name in sorted(priv_lst): f.write(' HIP_API_ID_' + name + ' = HIP_API_ID_NONE,\n') f.write('};\n') - f.write('\n#undef HIP_API_ID_CONCAT_HELPER\n'); - f.write('#undef HIP_API_ID_CONCAT\n'); + # f.write('\n#undef HIP_API_ID_CONCAT_HELPER\n'); + # f.write('#undef HIP_API_ID_CONCAT\n'); # Generating the method to return API name by ID f.write('\n// Return the HIP API string for a given callback ID\n') @@ -635,7 +640,7 @@ def generate_prof_header(f, api_map, callback_ids, opts_map): " ./src ./include/hip/amd_detail/hip_prof_str.h ./include/hip/amd_detail/hip_prof_str.h.new"); # API header file given as an argument -src_pat = "\.cpp$" +src_pat = r"\.cpp$" api_hfile = sys.argv[1] if not os.path.isfile(api_hfile): fatal("input file '" + api_hfile + "' not found") diff --git a/projects/clr/hipamd/src/hip_table_interface.cpp b/projects/clr/hipamd/src/hip_table_interface.cpp index 67c35c8f5ac5..cfd668ec5716 100644 --- a/projects/clr/hipamd/src/hip_table_interface.cpp +++ b/projects/clr/hipamd/src/hip_table_interface.cpp @@ -4,7 +4,10 @@ * SPDX-License-Identifier: MIT */ +// This is the entry point of all APIs we need to support. This file needs to +// guarantee ABI stability so we need to see all versions. #include +#include // For hipDeviceProp_tR0000 #include "hip_internal.hpp" #include "utils/flags.hpp" #include "utils/debug.hpp" @@ -179,7 +182,7 @@ hipError_t hipBindTextureToMipmappedArray(const textureReference* tex, return hip::GetHipDispatchTable()->hipBindTextureToMipmappedArray_fn(tex, mipmappedArray, desc); CATCH; } -extern "C" hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* prop) { +extern "C" hipError_t hipChooseDeviceR0600(int* device, const hipDeviceProp_tR0600* prop) { TRY; return hip::GetHipDispatchTable()->hipChooseDevice_fn(device, prop); CATCH; diff --git a/projects/clr/hipamd/src/hrr/hip_capture_generated.cpp b/projects/clr/hipamd/src/hrr/hip_capture_generated.cpp index 28a2b099ac4e..b4f2363d0400 100644 --- a/projects/clr/hipamd/src/hrr/hip_capture_generated.cpp +++ b/projects/clr/hipamd/src/hrr/hip_capture_generated.cpp @@ -289,7 +289,7 @@ static hipError_t capture_hipBindTextureToMipmappedArray(const textureReference* } // Generated shim -static hipError_t capture_hipChooseDevice(int* device, const hipDeviceProp_t* prop) { +static hipError_t capture_hipChooseDevice(int* device, const hipDeviceProp_tR0600* prop) { hipError_t r = g_real_table.hipChooseDevice_fn(device, prop); if (r == hipSuccess) { hrr_args_hipChooseDevice a{}; diff --git a/projects/clr/hipamd/src/hrr/hip_capture_metadata.cpp b/projects/clr/hipamd/src/hrr/hip_capture_metadata.cpp index bb3c6e8a2673..b802df68d816 100644 --- a/projects/clr/hipamd/src/hrr/hip_capture_metadata.cpp +++ b/projects/clr/hipamd/src/hrr/hip_capture_metadata.cpp @@ -18,6 +18,7 @@ namespace hip { class Device; extern std::vector g_devices; +// ihip is internal/unstable API extern hipError_t ihipGetDeviceProperties(hipDeviceProp_t* props, hipDevice_t device); } // namespace hip diff --git a/projects/clr/hipamd/src/profiler/hip_clr_dispatch_wrappers.cpp b/projects/clr/hipamd/src/profiler/hip_clr_dispatch_wrappers.cpp index 1aa356e39c65..c0cf04bacdcd 100644 --- a/projects/clr/hipamd/src/profiler/hip_clr_dispatch_wrappers.cpp +++ b/projects/clr/hipamd/src/profiler/hip_clr_dispatch_wrappers.cpp @@ -183,7 +183,7 @@ static hipError_t hipBindTextureToMipmappedArrayLayer(const textureReference* te } // api_id = 11 -static hipError_t hipChooseDeviceLayer(int* device, const hipDeviceProp_t* prop) { +static hipError_t hipChooseDeviceLayerR0600(int* device, const hipDeviceProp_tR0600* prop) { auto* _rec = HipGetActiveRecordExt(11u); auto _r = g_next.hipChooseDevice_fn(device, prop); _rec->end_ns = NowNs(); @@ -5592,7 +5592,7 @@ void HipProfilerBuildWrapperTableExt(HipDispatchTable* tbl) { g_wrapper_tbl.hipBindTexture2D_fn = hipBindTexture2DLayer; g_wrapper_tbl.hipBindTextureToArray_fn = hipBindTextureToArrayLayer; g_wrapper_tbl.hipBindTextureToMipmappedArray_fn = hipBindTextureToMipmappedArrayLayer; - g_wrapper_tbl.hipChooseDevice_fn = hipChooseDeviceLayer; + g_wrapper_tbl.hipChooseDevice_fn = hipChooseDeviceLayerR0600; g_wrapper_tbl.hipChooseDeviceR0000_fn = hipChooseDeviceR0000Layer; g_wrapper_tbl.hipConfigureCall_fn = hipConfigureCallLayer; g_wrapper_tbl.hipCreateSurfaceObject_fn = hipCreateSurfaceObjectLayer; diff --git a/projects/hip-tests/catch/cmake/hip-tests.cmake b/projects/hip-tests/catch/cmake/hip-tests.cmake index b74bf35005f4..4b0fb46f179b 100644 --- a/projects/hip-tests/catch/cmake/hip-tests.cmake +++ b/projects/hip-tests/catch/cmake/hip-tests.cmake @@ -108,6 +108,11 @@ function(hip_gen_exe_target) # Add dependency on build_tests to build it on this custom target add_dependencies(${_TEST_TARGET_NAME} ${_EXE_NAME}) + # Test deprecated functions too. We would like to stay on the + # latest API, so define HIP_ABI_IMPL to access deprecated + # APIs. The cost of this is that we lose access to wrapper + # functions. + target_compile_definitions(${_EXE_NAME} PRIVATE HIP_ABI_IMPL) if (DEFINED _COMPILE_OPTIONS) target_compile_options(${_EXE_NAME} PUBLIC ${_COMPILE_OPTIONS}) diff --git a/projects/hip/docs/doxygen/Doxyfile b/projects/hip/docs/doxygen/Doxyfile index d287645c62b2..449217fb8deb 100644 --- a/projects/hip/docs/doxygen/Doxyfile +++ b/projects/hip/docs/doxygen/Doxyfile @@ -2198,6 +2198,7 @@ INCLUDE_FILE_PATTERNS = # This tag requires that the tag ENABLE_PREPROCESSING is set to YES. PREDEFINED = "__HIP_PLATFORM_AMD__" \ + "HIP_FORCE_API_VERSION=99999" \ "DOXYGEN_SHOULD_INCLUDE_THIS=1" \ "DOXYGEN_SHOULD_SKIP_THIS=1" \ "__dparm(x)=" \ diff --git a/projects/hip/include/hip/hip_common.h b/projects/hip/include/hip/hip_common.h index 4989ce8b1906..56fec6994161 100644 --- a/projects/hip/include/hip/hip_common.h +++ b/projects/hip/include/hip/hip_common.h @@ -81,4 +81,10 @@ #pragma clang diagnostic pop #endif +#define HIP_API_VERSION (HIP_VERSION_MAJOR * 100 + HIP_VERSION_MINOR) + +#ifndef HIP_FORCE_API_VERSION +# define HIP_FORCE_API_VERSION HIP_API_VERSION +#endif + #endif diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index fc98f112c932..b88a8c1ca252 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -30,6 +30,10 @@ #include #include +#if HIP_FORCE_API_VERSION > HIP_API_VERSION +# error "HIP_FORCE_API_VERSION is newer than the installed HIP headers" +#endif + enum { HIP_SUCCESS = 0, HIP_ERROR_INVALID_VALUE, @@ -84,9 +88,23 @@ typedef struct hipUUID_t { //--- // Common headers for both NVCC and HIP-Clang paths: -#define hipGetDeviceProperties hipGetDevicePropertiesR0600 -#define hipDeviceProp_t hipDeviceProp_tR0600 -#define hipChooseDevice hipChooseDeviceR0600 +/* + Versioning struct. Because each public API must use the versioned struct to + ensure ABI compatibility, we must typedef the actual struct name here. +*/ + +#if HIP_FORCE_API_VERSION < 600 + +#include +// Legacy version, HIP_ABI_IMPL is on the latest version, so it must not see this. +typedef hipDeviceProp_tR0000 hipDeviceProp_t; +// Dummy struct for hip_prof_str.h +typedef struct hipDeviceProp_tR0600 {} hipDeviceProp_tR0600; + +#endif + +#if HIP_FORCE_API_VERSION >= 600 +// Latest version. /** * hipDeviceProp @@ -237,6 +255,10 @@ typedef struct hipDeviceProp_t { int asicRevision; ///< Revision of the GPU in this device } hipDeviceProp_t; +// Latest version, HIP_ABI_IMPL should see this. +typedef hipDeviceProp_t hipDeviceProp_tR0600; +#endif + /** * hipMemoryType (for pointer attributes) * @@ -396,7 +418,7 @@ typedef enum __HIP_NODISCARD hipError_t { hipErrorInvalidResourceType = 914, ///< Resource type is not valid for the operation. hipErrorInvalidResourceConfiguration = 915, ///< Resource configuration is not valid for ///< the operation. - hipErrorStreamDetached = 916, ///< The stream is detached. + hipErrorStreamDetached = 916, ///< The stream is detached. hipErrorUnknown = 999, ///< Unknown error. // HSA Runtime Error Codes start here. hipErrorRuntimeMemory = 1052, ///< HSA runtime memory call returned error. Typically not seen @@ -657,11 +679,16 @@ enum hipGPUDirectRDMAWritesOrdering { #include #include #include -#if defined(_MSC_VER) + +#if defined(HIP_ABI_IMPL) +// Nothing is deprecated if we are implementing the HIP ABI. +#define HIP_DEPRECATED(x) +#elif defined(_MSC_VER) #define HIP_DEPRECATED(msg) __declspec(deprecated(msg)) #else // !defined(_MSC_VER) #define HIP_DEPRECATED(msg) __attribute__((deprecated(msg))) #endif // !defined(_MSC_VER) + #define HIP_DEPRECATED_MSG \ "This API is marked as deprecated and might not be supported in future releases. For more " \ "details please refer " \ @@ -2199,6 +2226,7 @@ typedef enum hipMemRangeFlags { hipMemRangeFlagsMax = 0x7fffffff } hipMemRangeFlags; + // Doxygen end group GlobalDefs /** * @} @@ -2527,6 +2555,14 @@ hipError_t hipDeviceSetMemPool(int device, hipMemPool_t mem_pool); * change and might have outstanding issues. */ hipError_t hipDeviceGetMemPool(hipMemPool_t* mem_pool, int device); + +// hipDeviceProp_t and hipGetDeviceProperties/hipChooseDevice are versioned. +// See HIP_FORCE_API_VERSION in hip_common.h. + +#if HIP_FORCE_API_VERSION >= 600 + +hipError_t hipGetDevicePropertiesR0600(hipDeviceProp_tR0600* prop, int deviceId); + /** * @brief Returns device properties. * @@ -2540,7 +2576,13 @@ hipError_t hipDeviceGetMemPool(hipMemPool_t* mem_pool, int device); * * Populates hipGetDeviceProperties with information for the specified device. */ -hipError_t hipGetDeviceProperties(hipDeviceProp_t* prop, int deviceId); +static inline hipError_t hipGetDeviceProperties(hipDeviceProp_t* prop, int deviceId) { + return hipGetDevicePropertiesR0600(prop, deviceId); +} + +#endif + + /** * @brief Gets the maximum width for 1D linear textures on the specified device * @@ -2680,6 +2722,24 @@ hipError_t hipDeviceSetSharedMemConfig(hipSharedMemConfig config); * */ hipError_t hipSetDeviceFlags(unsigned flags); + +#if HIP_FORCE_API_VERSION >= 600 +hipError_t hipChooseDeviceR0600(int* device, const hipDeviceProp_tR0600* prop); + +/** + * @brief Device which matches hipDeviceProp_t is returned + * + * @param [out] device Pointer of the device + * @param [in] prop Pointer of the properties + * + * @returns #hipSuccess, #hipErrorInvalidValue + */ +static inline hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* prop) { + return hipChooseDeviceR0600(device, prop); +} + +#endif + /** * @brief Initialize the specified device to be used for GPU executions. * @@ -2703,16 +2763,9 @@ hipError_t hipSetDeviceFlags(unsigned flags); * * @see hipSetDevice, hipSetDeviceFlags */ + hipError_t hipInitDevice(int device, unsigned int deviceFlags, unsigned int flags); -/** - * @brief Device which matches hipDeviceProp_t is returned - * - * @param [out] device Pointer of the device - * @param [in] prop Pointer of the properties - * - * @returns #hipSuccess, #hipErrorInvalidValue - */ -hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* prop); + /** * @brief Returns the link type and hop count between two devices * @@ -4285,6 +4338,37 @@ hipError_t hipDrvMemDiscardAndPrefetchBatchAsync(hipDeviceptr_t* dptrs, size_t* size_t* prefetchLocIdxs, size_t numPrefetchLocs, unsigned long long flags, hipStream_t stream); + +// Deprecated since 800. Introduced in 700. +#if (HIP_FORCE_API_VERSION >= 700 && HIP_FORCE_API_VERSION < 800) || defined(HIP_ABI_IMPL) + +/** + * @brief Advise about the usage of a given memory range to HIP. + * + * @param [in] dev_ptr pointer to memory to set the advice for + * @param [in] count size in bytes of the memory range, it should be CPU page size alligned. + * @param [in] advice advice to be applied for the specified memory range + * @param [in] location location to apply the advice for + * + * @returns #hipSuccess, #hipErrorInvalidValue + * + * This HIP API advises about the usage to be applied on unified memory allocation in the + * range starting from the pointer address devPtr, with the size of count bytes. + * The memory range must refer to managed memory allocated via the API hipMallocManaged, and the + * range will be handled with proper round down and round up respectively in the driver to + * be aligned to CPU page size, the same way as corresponding CUDA API behaves in CUDA version 8.0 + * and afterwards. + * + * @note This API is implemented on Linux and is under development on Microsoft Windows. + */ +hipError_t hipMemAdvise_v2(const void* dev_ptr, size_t count, hipMemoryAdvise advice, + hipMemLocation location); + +#endif + +// Deprecated since 800. +#if HIP_FORCE_API_VERSION < 800 || defined(HIP_ABI_IMPL) + /** * @brief Advise about the usage of a given memory range to HIP. * @@ -4305,6 +4389,15 @@ hipError_t hipDrvMemDiscardAndPrefetchBatchAsync(hipDeviceptr_t* dptrs, size_t* * @note This API is implemented on Linux and is under development on Microsoft Windows. */ hipError_t hipMemAdvise(const void* dev_ptr, size_t count, hipMemoryAdvise advice, int device); + +#endif + +// Latest +#if HIP_FORCE_API_VERSION >= 800 && !defined(HIP_ABI_IMPL) + +hipError_t hipMemAdvise_v2(const void* dev_ptr, size_t count, hipMemoryAdvise advice, + hipMemLocation location); + /** * @brief Advise about the usage of a given memory range to HIP. * @@ -4323,9 +4416,15 @@ hipError_t hipMemAdvise(const void* dev_ptr, size_t count, hipMemoryAdvise advic * and afterwards. * * @note This API is implemented on Linux and is under development on Microsoft Windows. + * @note Since 7.16, this API is an alias to the old hipMemAdvise_v2 API. */ -hipError_t hipMemAdvise_v2(const void* dev_ptr, size_t count, hipMemoryAdvise advice, - hipMemLocation location); +static inline hipError_t hipMemAdvise(const void* dev_ptr, size_t count, hipMemoryAdvise advice, + hipMemLocation location) { + return hipMemAdvise_v2(dev_ptr, count, advice, location); +} + +#endif + /** * @brief Query an attribute of a given memory range in HIP. * @@ -10028,7 +10127,7 @@ hipError_t hipMemAddressReserve(void** ptr, size_t size, size_t alignment, void* * @param [in] prop - properties of the allocation. * @param [in] flags - currently unused, must be zero. * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotSupported - * + * * This API creates a memory allocation on the target device specified through the prop structure. * The prop allocation type must be specified as either #hipMemAllocationTypePinned or * #hipMemAllocationTypeUncached. @@ -10327,6 +10426,7 @@ hipError_t hipExtDisableLogging(); * @see hipExtEnableLogging, hipExtDisableLogging */ hipError_t hipExtSetLoggingParams(size_t log_level, size_t log_size, size_t log_mask); + /** * @} */ @@ -11079,7 +11179,6 @@ static inline hipError_t hipMallocManaged(T** devPtr, size_t size, return hipMallocManaged((void**)devPtr, size, flags); } - #endif #endif // doxygen end HIP API diff --git a/projects/hrr/playback/CMakeLists.txt b/projects/hrr/playback/CMakeLists.txt index 12e47de2879a..72b4cbe83f66 100644 --- a/projects/hrr/playback/CMakeLists.txt +++ b/projects/hrr/playback/CMakeLists.txt @@ -31,6 +31,7 @@ target_include_directories(hrr_reader PUBLIC $ $ ) +target_compile_definitions(hrr_reader PRIVATE HIP_ABI_IMPL) target_compile_definitions(hrr_reader PRIVATE HRR_API_ARGS_IMPLEMENTATION $<$:NOMINMAX WIN32_LEAN_AND_MEAN> @@ -54,6 +55,7 @@ if(HRR_BUILD_PLAYBACK) target_include_directories(hrr-playback PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} ) + target_compile_definitions(hrr-playback PRIVATE HIP_ABI_IMPL) target_compile_definitions(hrr-playback PRIVATE __HIP_PLATFORM_AMD__ "HRR_BUILD_TYPE=\"$>,$,unknown>\"" From 6ce26abcff4edfd37b1faf977f619bd64c9cb27c Mon Sep 17 00:00:00 2001 From: Mike Qin Date: Tue, 1 Sep 2026 22:38:08 -0400 Subject: [PATCH 3/6] move hipError_t typedef up. --- projects/hip/include/hip/hip_runtime_api.h | 280 ++++++++++----------- 1 file changed, 140 insertions(+), 140 deletions(-) diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index b88a8c1ca252..914efb73fa12 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -88,6 +88,146 @@ typedef struct hipUUID_t { //--- // Common headers for both NVCC and HIP-Clang paths: +// Ignoring error-code return values from hip APIs is discouraged. On C++17, +// we can make that yield a warning +#if __cplusplus >= 201703L +#define __HIP_NODISCARD [[nodiscard]] +#else +#define __HIP_NODISCARD +#endif + +/** + * HIP error type + * + */ +// Developer note - when updating these, update the hipErrorName and hipErrorString functions in +// NVCC and HIP-Clang paths Also update the hipCUDAErrorTohipError function in NVCC path. + +typedef enum __HIP_NODISCARD hipError_t { + hipSuccess = 0, ///< Successful completion. + hipErrorInvalidValue = 1, ///< One or more of the parameters passed to the API call is NULL + ///< or not in an acceptable range. + hipErrorOutOfMemory = 2, ///< out of memory range. + // Deprecated + hipErrorMemoryAllocation = 2, ///< Memory allocation error. + hipErrorNotInitialized = 3, ///< Invalid not initialized + // Deprecated + hipErrorInitializationError = 3, + hipErrorDeinitialized = 4, ///< Deinitialized + hipErrorProfilerDisabled = 5, + hipErrorProfilerNotInitialized = 6, + hipErrorProfilerAlreadyStarted = 7, + hipErrorProfilerAlreadyStopped = 8, + hipErrorInvalidConfiguration = 9, ///< Invalide configuration + hipErrorInvalidPitchValue = 12, ///< Invalid pitch value + hipErrorInvalidSymbol = 13, ///< Invalid symbol + hipErrorInvalidDevicePointer = 17, ///< Invalid Device Pointer + hipErrorInvalidMemcpyDirection = 21, ///< Invalid memory copy direction + hipErrorInsufficientDriver = 35, + hipErrorMissingConfiguration = 52, + hipErrorPriorLaunchFailure = 53, + hipErrorInvalidDeviceFunction = 98, ///< Invalid device function + hipErrorNoDevice = 100, ///< Call to hipGetDeviceCount returned 0 devices + hipErrorInvalidDevice = 101, ///< DeviceID must be in range from 0 to compute-devices. + hipErrorInvalidImage = 200, ///< Invalid image + hipErrorInvalidContext = 201, ///< Produced when input context is invalid. + hipErrorContextAlreadyCurrent = 202, + hipErrorMapFailed = 205, + // Deprecated + hipErrorMapBufferObjectFailed = 205, ///< Produced when the IPC memory attach failed from ROCr. + hipErrorUnmapFailed = 206, + hipErrorArrayIsMapped = 207, + hipErrorAlreadyMapped = 208, + hipErrorNoBinaryForGpu = 209, + hipErrorAlreadyAcquired = 210, + hipErrorNotMapped = 211, + hipErrorNotMappedAsArray = 212, + hipErrorNotMappedAsPointer = 213, + hipErrorECCNotCorrectable = 214, + hipErrorUnsupportedLimit = 215, ///< Unsupported limit + hipErrorContextAlreadyInUse = 216, ///< The context is already in use + hipErrorPeerAccessUnsupported = 217, + hipErrorInvalidKernelFile = 218, ///< In CUDA DRV, it is CUDA_ERROR_INVALID_PTX + hipErrorInvalidGraphicsContext = 219, + hipErrorInvalidSource = 300, ///< Invalid source. + hipErrorFileNotFound = 301, ///< the file is not found. + hipErrorSharedObjectSymbolNotFound = 302, + hipErrorSharedObjectInitFailed = 303, ///< Failed to initialize shared object. + hipErrorOperatingSystem = 304, ///< Not the correct operating system + hipErrorInvalidHandle = 400, ///< Invalide handle + // Deprecated + hipErrorInvalidResourceHandle = 400, ///< Resource handle (hipEvent_t or hipStream_t) invalid. + hipErrorIllegalState = 401, ///< Resource required is not in a valid state to perform operation. + hipErrorNotFound = 500, ///< Not found + hipErrorNotReady = 600, ///< Indicates that asynchronous operations enqueued earlier are not + ///< ready. This is not actually an error, but is used to distinguish + ///< from hipSuccess (which indicates completion). APIs that return + ///< this error include hipEventQuery and hipStreamQuery. + hipErrorIllegalAddress = 700, + hipErrorLaunchOutOfResources = 701, ///< Out of resources error. + hipErrorLaunchTimeOut = 702, ///< Timeout for the launch. + hipErrorPeerAccessAlreadyEnabled = 704, ///< Peer access was already enabled from the current + ///< device. + hipErrorPeerAccessNotEnabled = 705, ///< Peer access was never enabled from the current device. + hipErrorSetOnActiveProcess = 708, ///< The process is active. + hipErrorContextIsDestroyed = 709, ///< The context is already destroyed + hipErrorAssert = 710, ///< Produced when the kernel calls assert. + hipErrorHostMemoryAlreadyRegistered = 712, ///< Produced when trying to lock a page-locked + ///< memory. + hipErrorHostMemoryNotRegistered = 713, ///< Produced when trying to unlock a non-page-locked + ///< memory. + hipErrorLaunchFailure = 719, ///< An exception occurred on the device while executing a kernel. + hipErrorCooperativeLaunchTooLarge = 720, ///< This error indicates that the number of blocks + ///< launched per grid for a kernel that was launched + ///< via cooperative launch APIs exceeds the maximum + ///< number of allowed blocks for the current device. + hipErrorNotSupported = 801, ///< Produced when the hip API is not supported/implemented + hipErrorStreamCaptureUnsupported = 900, ///< The operation is not permitted when the stream + ///< is capturing. + hipErrorStreamCaptureInvalidated = 901, ///< The current capture sequence on the stream + ///< has been invalidated due to a previous error. + hipErrorStreamCaptureMerge = 902, ///< The operation would have resulted in a merge of + ///< two independent capture sequences. + hipErrorStreamCaptureUnmatched = 903, ///< The capture was not initiated in this stream. + hipErrorStreamCaptureUnjoined = 904, ///< The capture sequence contains a fork that was not + ///< joined to the primary stream. + hipErrorStreamCaptureIsolation = 905, ///< A dependency would have been created which crosses + ///< the capture sequence boundary. Only implicit + ///< in-stream ordering dependencies are allowed + ///< to cross the boundary + hipErrorStreamCaptureImplicit = 906, ///< The operation would have resulted in a disallowed + ///< implicit dependency on a current capture sequence + ///< from hipStreamLegacy. + hipErrorCapturedEvent = 907, ///< The operation is not permitted on an event which was last + ///< recorded in a capturing stream. + hipErrorStreamCaptureWrongThread = 908, ///< A stream capture sequence not initiated with + ///< the hipStreamCaptureModeRelaxed argument to + ///< hipStreamBeginCapture was passed to + ///< hipStreamEndCapture in a different thread. + hipErrorGraphExecUpdateFailure = 910, ///< This error indicates that the graph update + ///< not performed because it included changes which + ///< violated constraintsspecific to instantiated graph + ///< update. + hipErrorInvalidChannelDescriptor = 911, ///< Invalid channel descriptor. + hipErrorInvalidTexture = 912, ///< Invalid texture. + hipErrorInvalidResourceType = 914, ///< Resource type is not valid for the operation. + hipErrorInvalidResourceConfiguration = 915, ///< Resource configuration is not valid for + ///< the operation. + hipErrorStreamDetached = 916, ///< The stream is detached. + hipErrorUnknown = 999, ///< Unknown error. + // HSA Runtime Error Codes start here. + hipErrorRuntimeMemory = 1052, ///< HSA runtime memory call returned error. Typically not seen + ///< in production systems. + hipErrorRuntimeOther = 1053, ///< HSA runtime call other than memory returned error. Typically + ///< not seen in production systems. + hipErrorInvalidClusterSize = 1054, ///< The specified cluster size is invalid, for instance + ///< when passing launch configurations to occupancy + ///< calculations + hipErrorTbd ///< Marker that more error codes are needed. +} hipError_t; + +#undef __HIP_NODISCARD + /* Versioning struct. Because each public API must use the versioned struct to ensure ABI compatibility, we must typedef the actual struct name here. @@ -293,146 +433,6 @@ typedef struct hipPointerAttribute_t { /* peers? */ } hipPointerAttribute_t; -// Ignoring error-code return values from hip APIs is discouraged. On C++17, -// we can make that yield a warning -#if __cplusplus >= 201703L -#define __HIP_NODISCARD [[nodiscard]] -#else -#define __HIP_NODISCARD -#endif - -/** - * HIP error type - * - */ -// Developer note - when updating these, update the hipErrorName and hipErrorString functions in -// NVCC and HIP-Clang paths Also update the hipCUDAErrorTohipError function in NVCC path. - -typedef enum __HIP_NODISCARD hipError_t { - hipSuccess = 0, ///< Successful completion. - hipErrorInvalidValue = 1, ///< One or more of the parameters passed to the API call is NULL - ///< or not in an acceptable range. - hipErrorOutOfMemory = 2, ///< out of memory range. - // Deprecated - hipErrorMemoryAllocation = 2, ///< Memory allocation error. - hipErrorNotInitialized = 3, ///< Invalid not initialized - // Deprecated - hipErrorInitializationError = 3, - hipErrorDeinitialized = 4, ///< Deinitialized - hipErrorProfilerDisabled = 5, - hipErrorProfilerNotInitialized = 6, - hipErrorProfilerAlreadyStarted = 7, - hipErrorProfilerAlreadyStopped = 8, - hipErrorInvalidConfiguration = 9, ///< Invalide configuration - hipErrorInvalidPitchValue = 12, ///< Invalid pitch value - hipErrorInvalidSymbol = 13, ///< Invalid symbol - hipErrorInvalidDevicePointer = 17, ///< Invalid Device Pointer - hipErrorInvalidMemcpyDirection = 21, ///< Invalid memory copy direction - hipErrorInsufficientDriver = 35, - hipErrorMissingConfiguration = 52, - hipErrorPriorLaunchFailure = 53, - hipErrorInvalidDeviceFunction = 98, ///< Invalid device function - hipErrorNoDevice = 100, ///< Call to hipGetDeviceCount returned 0 devices - hipErrorInvalidDevice = 101, ///< DeviceID must be in range from 0 to compute-devices. - hipErrorInvalidImage = 200, ///< Invalid image - hipErrorInvalidContext = 201, ///< Produced when input context is invalid. - hipErrorContextAlreadyCurrent = 202, - hipErrorMapFailed = 205, - // Deprecated - hipErrorMapBufferObjectFailed = 205, ///< Produced when the IPC memory attach failed from ROCr. - hipErrorUnmapFailed = 206, - hipErrorArrayIsMapped = 207, - hipErrorAlreadyMapped = 208, - hipErrorNoBinaryForGpu = 209, - hipErrorAlreadyAcquired = 210, - hipErrorNotMapped = 211, - hipErrorNotMappedAsArray = 212, - hipErrorNotMappedAsPointer = 213, - hipErrorECCNotCorrectable = 214, - hipErrorUnsupportedLimit = 215, ///< Unsupported limit - hipErrorContextAlreadyInUse = 216, ///< The context is already in use - hipErrorPeerAccessUnsupported = 217, - hipErrorInvalidKernelFile = 218, ///< In CUDA DRV, it is CUDA_ERROR_INVALID_PTX - hipErrorInvalidGraphicsContext = 219, - hipErrorInvalidSource = 300, ///< Invalid source. - hipErrorFileNotFound = 301, ///< the file is not found. - hipErrorSharedObjectSymbolNotFound = 302, - hipErrorSharedObjectInitFailed = 303, ///< Failed to initialize shared object. - hipErrorOperatingSystem = 304, ///< Not the correct operating system - hipErrorInvalidHandle = 400, ///< Invalide handle - // Deprecated - hipErrorInvalidResourceHandle = 400, ///< Resource handle (hipEvent_t or hipStream_t) invalid. - hipErrorIllegalState = 401, ///< Resource required is not in a valid state to perform operation. - hipErrorNotFound = 500, ///< Not found - hipErrorNotReady = 600, ///< Indicates that asynchronous operations enqueued earlier are not - ///< ready. This is not actually an error, but is used to distinguish - ///< from hipSuccess (which indicates completion). APIs that return - ///< this error include hipEventQuery and hipStreamQuery. - hipErrorIllegalAddress = 700, - hipErrorLaunchOutOfResources = 701, ///< Out of resources error. - hipErrorLaunchTimeOut = 702, ///< Timeout for the launch. - hipErrorPeerAccessAlreadyEnabled = 704, ///< Peer access was already enabled from the current - ///< device. - hipErrorPeerAccessNotEnabled = 705, ///< Peer access was never enabled from the current device. - hipErrorSetOnActiveProcess = 708, ///< The process is active. - hipErrorContextIsDestroyed = 709, ///< The context is already destroyed - hipErrorAssert = 710, ///< Produced when the kernel calls assert. - hipErrorHostMemoryAlreadyRegistered = 712, ///< Produced when trying to lock a page-locked - ///< memory. - hipErrorHostMemoryNotRegistered = 713, ///< Produced when trying to unlock a non-page-locked - ///< memory. - hipErrorLaunchFailure = 719, ///< An exception occurred on the device while executing a kernel. - hipErrorCooperativeLaunchTooLarge = 720, ///< This error indicates that the number of blocks - ///< launched per grid for a kernel that was launched - ///< via cooperative launch APIs exceeds the maximum - ///< number of allowed blocks for the current device. - hipErrorNotSupported = 801, ///< Produced when the hip API is not supported/implemented - hipErrorStreamCaptureUnsupported = 900, ///< The operation is not permitted when the stream - ///< is capturing. - hipErrorStreamCaptureInvalidated = 901, ///< The current capture sequence on the stream - ///< has been invalidated due to a previous error. - hipErrorStreamCaptureMerge = 902, ///< The operation would have resulted in a merge of - ///< two independent capture sequences. - hipErrorStreamCaptureUnmatched = 903, ///< The capture was not initiated in this stream. - hipErrorStreamCaptureUnjoined = 904, ///< The capture sequence contains a fork that was not - ///< joined to the primary stream. - hipErrorStreamCaptureIsolation = 905, ///< A dependency would have been created which crosses - ///< the capture sequence boundary. Only implicit - ///< in-stream ordering dependencies are allowed - ///< to cross the boundary - hipErrorStreamCaptureImplicit = 906, ///< The operation would have resulted in a disallowed - ///< implicit dependency on a current capture sequence - ///< from hipStreamLegacy. - hipErrorCapturedEvent = 907, ///< The operation is not permitted on an event which was last - ///< recorded in a capturing stream. - hipErrorStreamCaptureWrongThread = 908, ///< A stream capture sequence not initiated with - ///< the hipStreamCaptureModeRelaxed argument to - ///< hipStreamBeginCapture was passed to - ///< hipStreamEndCapture in a different thread. - hipErrorGraphExecUpdateFailure = 910, ///< This error indicates that the graph update - ///< not performed because it included changes which - ///< violated constraintsspecific to instantiated graph - ///< update. - hipErrorInvalidChannelDescriptor = 911, ///< Invalid channel descriptor. - hipErrorInvalidTexture = 912, ///< Invalid texture. - hipErrorInvalidResourceType = 914, ///< Resource type is not valid for the operation. - hipErrorInvalidResourceConfiguration = 915, ///< Resource configuration is not valid for - ///< the operation. - hipErrorStreamDetached = 916, ///< The stream is detached. - hipErrorUnknown = 999, ///< Unknown error. - // HSA Runtime Error Codes start here. - hipErrorRuntimeMemory = 1052, ///< HSA runtime memory call returned error. Typically not seen - ///< in production systems. - hipErrorRuntimeOther = 1053, ///< HSA runtime call other than memory returned error. Typically - ///< not seen in production systems. - hipErrorInvalidClusterSize = 1054, ///< The specified cluster size is invalid, for instance - ///< when passing launch configurations to occupancy - ///< calculations - hipErrorTbd ///< Marker that more error codes are needed. -} hipError_t; - -#undef __HIP_NODISCARD - /** * hipDeviceAttribute_t * hipDeviceAttributeUnused number: 5 From 53ebcd1ad6191240972124217ba670f25ef62181 Mon Sep 17 00:00:00 2001 From: Pengda Xie Date: Tue, 25 Aug 2026 06:04:43 +0000 Subject: [PATCH 4/6] fix(clr): resolve type-alias macros in hip_prof_gen signature check hipStreamSet/GetAttribute (and graph kernel-node attribute APIs) are defined in source via the hipStreamAttrID/hipStreamAttrValue alias macros, while the etalon is parsed from the preprocessed header where cpp has expanded them to hipLaunchAttributeID/hipLaunchAttributeValue. The string type-compare saw a mismatch, treated the APIs as non-public overloads, and dropped their profiling callbacks from hip_prof_str.h. On mismatch, lazily build a map of object-like #define type aliases and re-compare after resolving transitively. The alias #defines live in the source API headers that the build cat's into new_header.h before preprocessing; because they are cat'd (not #include'd) their bodies carry no per-source cpp line-marker, so scan the raw cat'd header directly (the .i input minus its suffix), plus any hip headers named by line-markers. Only suppresses spelling differences; real type drift still warns. --- .../include/hip/amd_detail/hip_prof_str.h | 12 ++ projects/clr/hipamd/src/hip_prof_gen.py | 107 +++++++++++++++--- 2 files changed, 105 insertions(+), 14 deletions(-) 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 d6d54777ed8b..ebc0d87fd6b2 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 @@ -5542,6 +5542,9 @@ typedef struct hip_api_data_s { }; // hipGraphKernelNodeGetAttribute[('hipGraphNode_t', 'hNode'), ('hipLaunchAttributeID', 'attr'), ('hipLaunchAttributeValue*', 'value')] #define INIT_hipGraphKernelNodeGetAttribute_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipGraphKernelNodeGetAttribute.hNode = (hipGraphNode_t)hNode; \ + cb_data.args.hipGraphKernelNodeGetAttribute.attr = (hipLaunchAttributeID)attr; \ + cb_data.args.hipGraphKernelNodeGetAttribute.value = (hipLaunchAttributeValue*)value; \ }; // hipGraphKernelNodeGetParams[('hipGraphNode_t', 'node'), ('hipKernelNodeParams*', 'pNodeParams')] #define INIT_hipGraphKernelNodeGetParams_CB_ARGS_DATA(cb_data) { \ @@ -5550,6 +5553,9 @@ typedef struct hip_api_data_s { }; // hipGraphKernelNodeSetAttribute[('hipGraphNode_t', 'hNode'), ('hipLaunchAttributeID', 'attr'), ('const hipLaunchAttributeValue*', 'value')] #define INIT_hipGraphKernelNodeSetAttribute_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipGraphKernelNodeSetAttribute.hNode = (hipGraphNode_t)hNode; \ + cb_data.args.hipGraphKernelNodeSetAttribute.attr = (hipLaunchAttributeID)attr; \ + cb_data.args.hipGraphKernelNodeSetAttribute.value = (const hipLaunchAttributeValue*)value; \ }; // hipGraphKernelNodeSetParams[('hipGraphNode_t', 'node'), ('const hipKernelNodeParams*', 'pNodeParams')] #define INIT_hipGraphKernelNodeSetParams_CB_ARGS_DATA(cb_data) { \ @@ -7086,6 +7092,9 @@ typedef struct hip_api_data_s { }; // hipStreamGetAttribute[('hipStream_t', 'stream'), ('hipLaunchAttributeID', 'attr'), ('hipLaunchAttributeValue*', 'value_out')] #define INIT_hipStreamGetAttribute_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipStreamGetAttribute.stream = (hipStream_t)stream; \ + cb_data.args.hipStreamGetAttribute.attr = (hipLaunchAttributeID)attr; \ + cb_data.args.hipStreamGetAttribute.value_out = (hipLaunchAttributeValue*)value_out; \ }; // hipStreamGetCaptureInfo[('hipStream_t', 'stream'), ('hipStreamCaptureStatus*', 'pCaptureStatus'), ('unsigned long long*', 'pId')] #define INIT_hipStreamGetCaptureInfo_CB_ARGS_DATA(cb_data) { \ @@ -7139,6 +7148,9 @@ typedef struct hip_api_data_s { }; // hipStreamSetAttribute[('hipStream_t', 'stream'), ('hipLaunchAttributeID', 'attr'), ('const hipLaunchAttributeValue*', 'value')] #define INIT_hipStreamSetAttribute_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipStreamSetAttribute.stream = (hipStream_t)stream; \ + cb_data.args.hipStreamSetAttribute.attr = (hipLaunchAttributeID)attr; \ + cb_data.args.hipStreamSetAttribute.value = (const hipLaunchAttributeValue*)value; \ }; // hipStreamSynchronize[('hipStream_t', 'stream')] #define INIT_hipStreamSynchronize_CB_ARGS_DATA(cb_data) { \ diff --git a/projects/clr/hipamd/src/hip_prof_gen.py b/projects/clr/hipamd/src/hip_prof_gen.py index 2f577de68c83..d1b45a0f7ee5 100755 --- a/projects/clr/hipamd/src/hip_prof_gen.py +++ b/projects/clr/hipamd/src/hip_prof_gen.py @@ -55,8 +55,8 @@ def filtr_api_name(name): return name def filtr_api_decl(record): - record = re.sub(r"\s__dparm\([^\)]*\)", r'', record); - record = re.sub(r"\(void\*\)", r'', record); + record = re.sub(r'\s__dparm\([^\)]*\)', r'', record); + record = re.sub(r'\(void\*\)', r'', record); return record # Normalizing API arguments @@ -75,6 +75,79 @@ def norm_api_types(type_str): type_str = re.sub(r'^unsigned$', r'unsigned int', type_str) return type_str +# None until built; {} once built (even if empty). +type_aliases = None + +# Building the alias map from the headers behind the preprocessed input. +def build_type_aliases(preproc_file): + hdrs = [] + seen = set() + define_pat = re.compile(r'^\s*#\s*define\s+([A-Za-z_]\w*)\s+([A-Za-z_]\w*)\s*$') + marker_pat = re.compile(r'^#\s+\d+\s+"([^"]+)"') + # The raw cat'd input header (the .i file's source) is the authoritative place + # the cat'd source #defines land; scan it first, independent of line-markers. + if preproc_file.endswith('.i'): + cat_hdr = preproc_file[:-len('.i')] + if os.path.isfile(cat_hdr): + seen.add(cat_hdr) + hdrs.append(cat_hdr) + try: + with open(preproc_file, 'r') as f: + for line in f: + m = marker_pat.match(line) + if m: + path = m.group(1) + # Restrict to hip headers; skip system headers and duplicates. + if path not in seen and 'hip' in os.path.basename(path).lower() and os.path.isfile(path): + seen.add(path) + hdrs.append(path) + except IOError: + return {} + + parent = {} + def find(name): + root = name + while parent.get(root, root) != root: + root = parent[root] + while parent.get(name, name) != root: + parent[name], name = root, parent[name] + return root + def union(alias, target): + ra, rt = find(alias), find(target) + if ra != rt: + parent[ra] = rt + + defined = set() + for path in hdrs: + try: + with open(path, 'r') as f: + for line in f: + m = define_pat.match(line) + if m: + alias, target = m.group(1), m.group(2) + parent.setdefault(alias, alias) + parent.setdefault(target, target) + union(alias, target) + defined.add(alias) + except IOError: + continue + + aliases = {} + for alias in defined: + canonical = find(alias) + if canonical != alias: + aliases[alias] = canonical + return aliases + +# Rewriting alias type names in a normalized type string to their canonical form. +def resolve_type_aliases(type_str): + global type_aliases + if type_aliases is None: + type_aliases = build_type_aliases(api_hfile) + for alias, canonical in type_aliases.items(): + type_str = re.sub(r'\b' + re.escape(alias) + r'\b', lambda m: canonical, type_str) + return type_str + # Creating a list of arguments [(type, name), ...] def list_api_args(args_str): args_str = filtr_api_args(args_str) @@ -83,7 +156,7 @@ def list_api_args(args_str): for arg_pair in args_str.split(','): if arg_pair == 'void': continue arg_pair = re.sub(r'\s*=\s*\S+$','', arg_pair); - m = re.match(r"^(.*)\s(\S+)$", arg_pair); + m = re.match(r'^(.*)\s(\S+)$', arg_pair); if m: arg_type = norm_api_types(m.group(1)) arg_name = m.group(2) @@ -127,9 +200,9 @@ def parse_api(inp_file_p, out): global line_num inp_file = inp_file_p - beg_pattern = re.compile(r"^(hipError_t|const char\s*\*)\s+([^\(]+)\("); - api_pattern = re.compile(r"^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)"); - end_pattern = re.compile(r"Texture"); + beg_pattern = re.compile(r'^(hipError_t|const char\s*\*)\s+([^\(]+)\('); + api_pattern = re.compile(r'^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)'); + end_pattern = re.compile(r'Texture'); hidden_pattern = re.compile(r'__attribute__\(\(visibility\("hidden"\)\)\)') nms_open_pattern = re.compile(r'namespace hip_impl {') nms_close_pattern = re.compile(r'}') @@ -161,7 +234,7 @@ def parse_api(inp_file_p, out): found = 1 if found != 0: - record = re.sub(r"\s__dparm\([^\)]*\)", '', record); + record = re.sub(r'\s__dparm\([^\)]*\)', '', record); m = api_pattern.match(record) if m: found = 0 @@ -200,13 +273,13 @@ def parse_content(inp_file_p, api_map, out): inp_file = inp_file_p # API method begin pattern - beg_pattern = re.compile(r"^(hipError_t|const char\s*\*)\s+[^\(]+\("); + beg_pattern = re.compile(r'^(hipError_t|const char\s*\*)\s+[^\(]+\('); # API declaration pattern - decl_pattern = re.compile(r"^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)\s*;"); + decl_pattern = re.compile(r'^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)\s*;'); # API definition pattern - api_pattern = re.compile(r"^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)\s*{"); + api_pattern = re.compile(r'^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)\s*{'); # API init macro pattern - init_pattern = re.compile(r"(^\s*HIP_INIT_API[^\s]*\s*)\((([^,]+)(,.*|)|)(\);|,)\s*$"); + init_pattern = re.compile(r'(^\s*HIP_INIT_API[^\s]*\s*)\((([^,]+)(,.*|)|)(\);|,)\s*$'); # Open input file inp = open(inp_file, 'r') @@ -274,7 +347,13 @@ def parse_content(inp_file_p, api_map, out): api_types = filtr_api_types(api_args) # Normalizing etalon arguments eta_types = filtr_api_types(eta_args) - if (api_types == eta_types) or ((types_check_mode == 0) and (not api_name in out)): + types_match = (api_types == eta_types) + # A pure spelling difference (source uses a type-alias macro that the + # preprocessed etalon already expanded) is a false mismatch. Resolve + # aliases and re-compare before treating it as an overload. + if not types_match: + types_match = (resolve_type_aliases(api_types) == resolve_type_aliases(eta_types)) + if types_match or ((types_check_mode == 0) and (not api_name in out)): # API is already found and not is mismatched if (api_name in out): fatal("API redefined \"" + api_name + "\", record \"" + record + "\"") @@ -357,7 +436,7 @@ def parse_src(api_map, src_path, src_patt, out): message(file) content = '' filename = os.path.basename(file) - if re.match("hip_table_interface.cpp", filename): + if re.match(r'hip_table_interface\.cpp$', filename): message("SKIP FILE:" + filename) else: content = parse_content(file, api_map, out); @@ -640,7 +719,7 @@ def generate_prof_header(f, api_map, callback_ids, opts_map): " ./src ./include/hip/amd_detail/hip_prof_str.h ./include/hip/amd_detail/hip_prof_str.h.new"); # API header file given as an argument -src_pat = r"\.cpp$" +src_pat = r'\.cpp$' api_hfile = sys.argv[1] if not os.path.isfile(api_hfile): fatal("input file '" + api_hfile + "' not found") From e1ed95cb48b9ef5b4f369774e73ae71e27674c4b Mon Sep 17 00:00:00 2001 From: Pengda Xie Date: Tue, 8 Sep 2026 09:50:44 +0000 Subject: [PATCH 5/6] remove redundant code --- projects/clr/hipamd/src/hip_prof_gen.py | 13 ++++--------- 1 file changed, 4 insertions(+), 9 deletions(-) diff --git a/projects/clr/hipamd/src/hip_prof_gen.py b/projects/clr/hipamd/src/hip_prof_gen.py index d1b45a0f7ee5..11b80d19b9b6 100755 --- a/projects/clr/hipamd/src/hip_prof_gen.py +++ b/projects/clr/hipamd/src/hip_prof_gen.py @@ -125,19 +125,14 @@ def union(alias, target): m = define_pat.match(line) if m: alias, target = m.group(1), m.group(2) - parent.setdefault(alias, alias) - parent.setdefault(target, target) union(alias, target) defined.add(alias) except IOError: continue - aliases = {} - for alias in defined: - canonical = find(alias) - if canonical != alias: - aliases[alias] = canonical - return aliases + for sym in defined: + find(sym) + return parent # Rewriting alias type names in a normalized type string to their canonical form. def resolve_type_aliases(type_str): @@ -145,7 +140,7 @@ def resolve_type_aliases(type_str): if type_aliases is None: type_aliases = build_type_aliases(api_hfile) for alias, canonical in type_aliases.items(): - type_str = re.sub(r'\b' + re.escape(alias) + r'\b', lambda m: canonical, type_str) + type_str = re.sub(r'\b' + re.escape(alias) + r'\b', canonical, type_str) return type_str # Creating a list of arguments [(type, name), ...] From ac45bd5613ccfd3ba51072e5fa6c3b54eb08d2c8 Mon Sep 17 00:00:00 2001 From: Mike Qin Date: Wed, 9 Sep 2026 18:36:09 -0400 Subject: [PATCH 6/6] add hip-tests that ensures R0000 struct is used with 500 API --- .../catch/config/configs/unit/device.yaml | 3 ++ .../catch/unit/device/CMakeLists.txt | 3 +- .../device/hipGetDeviceProperties_R0000.cc | 47 +++++++++++++++++++ 3 files changed, 52 insertions(+), 1 deletion(-) create mode 100644 projects/hip-tests/catch/unit/device/hipGetDeviceProperties_R0000.cc diff --git a/projects/hip-tests/catch/config/configs/unit/device.yaml b/projects/hip-tests/catch/config/configs/unit/device.yaml index 2bc10accf6a8..762c3a1ea757 100644 --- a/projects/hip-tests/catch/config/configs/unit/device.yaml +++ b/projects/hip-tests/catch/config/configs/unit/device.yaml @@ -105,6 +105,9 @@ device: Unit_hipGetDeviceProperties_NegTst: <<: *level_2 tags: [query, P0] + Unit_hipGetDeviceProperties_R0000: + <<: *level_2 + tags: [query, P0] Unit_hipRuntimeGetVersion_Positive: <<: *level_2 tags: [init, smoke, P0] diff --git a/projects/hip-tests/catch/unit/device/CMakeLists.txt b/projects/hip-tests/catch/unit/device/CMakeLists.txt index 4e972c7e77a3..d3ebfc11feb1 100644 --- a/projects/hip-tests/catch/unit/device/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/device/CMakeLists.txt @@ -15,6 +15,7 @@ set(TEST_SRC hipGetDeviceAttribute.cc hipGetDeviceCount.cc hipGetDeviceProperties.cc + hipGetDeviceProperties_R0000.cc hipRuntimeGetVersion.cc hipGetSetDeviceFlags.cc hipSetGetDevice.cc @@ -128,4 +129,4 @@ if(HIP_PLATFORM MATCHES "amd" AND BUILD_SHARED_LIBS AND UNIX) target_link_libraries(hipGetProcAddressIpcEventImport ${GPGPU_LINKER_LIBRARIES}) add_dependencies(DeviceTest hipGetProcAddressIpcEventImport) set_property(GLOBAL APPEND PROPERTY G_INSTALL_EXE_TARGETS hipGetProcAddressIpcEventImport) -endif() \ No newline at end of file +endif() diff --git a/projects/hip-tests/catch/unit/device/hipGetDeviceProperties_R0000.cc b/projects/hip-tests/catch/unit/device/hipGetDeviceProperties_R0000.cc new file mode 100644 index 000000000000..1eacc1a423b2 --- /dev/null +++ b/projects/hip-tests/catch/unit/device/hipGetDeviceProperties_R0000.cc @@ -0,0 +1,47 @@ +/* + * Copyright (c) Advanced Micro Devices, Inc., or its affiliates. + * + * SPDX-License-Identifier: MIT + */ + +// This test doesn't test the ABI, it tests if the API is alias to the R0000 when HIP_FORCE_API_VERSION is set. +// HIP Nvidia doesn't support this mechanism, so we can't test this. + +#if HT_AMD + +#undef HIP_ABI_IMPL +#define HIP_FORCE_API_VERSION 500 +#include +#include + +#ifdef ENABLE_YAML_TAGS +#include "hip_test_config.hh" + +#define SECOND_ARG(a, b, ...) b +#define GET_TAGS(...) SECOND_ARG(__VA_ARGS__) +#define HIP_TEST_CASE(name) TEST_CASE(#name, GET_TAGS(name)) +#define HIP_TEMPLATE_TEST_CASE(name, ...) TEMPLATE_TEST_CASE(#name, GET_TAGS(name), __VA_ARGS__) + +#else +#define GET_TAGS(...) +#define HIP_TEST_CASE(name) TEST_CASE(#name, "") +#define HIP_TEMPLATE_TEST_CASE(name, ...) TEMPLATE_TEST_CASE(#name, "", __VA_ARGS__) +#endif + +/** + * Test Description + * ------------------------ + * - Test if R0000 is used for hipDeviceProp_t + * Test source + * ------------------------ + * - unit/device/hipGetDeviceProperties_R0000.cc + * Test requirements + * ------------------------ + * - Platform specific (AMD) + * - HIP_VERSION <= 6.0 + */ +HIP_TEST_CASE(Unit_hipGetDeviceProperties_R0000) { + REQUIRE(sizeof(hipDeviceProp_t) == sizeof(hipDeviceProp_tR0000)); +} + +#endif