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
 
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..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
@@ -16,9 +16,6 @@
 #include <hip/hip_deprecated.h>
 #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; \
@@ -5550,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) { \
@@ -5558,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) { \
@@ -7094,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) { \
@@ -7147,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) { \
@@ -7579,6 +7583,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 +7617,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 +9633,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 +9695,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$<JOIN:$<TARGET_PROPERTY:amdhip64,COMPILE_DEFINITIONS>,;-D>"
-        "-I$<JOIN:$<TARGET_PROPERTY:amdhip64,INCLUDE_DIRECTORIES>,;-I>"
+        "$<LIST:TRANSFORM,$<TARGET_PROPERTY:amdhip64,INCLUDE_DIRECTORIES>,PREPEND,-I>" # Directory name may contain spaces.
         "-DHIP_INCLUDE_HIP_HIP_RUNTIME_PT_API_H=1"
-        ${c_flags}
         $<TARGET_PROPERTY:amdhip64,COMPILE_OPTIONS>
-        ${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<int>(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<hip::FatBinaryInfo**>(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..11b80d19b9b6 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
@@ -75,6 +75,74 @@ 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)
+            union(alias, target)
+            defined.add(alias)
+    except IOError:
+      continue
+
+  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):
+  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', 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 +151,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 +195,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 +229,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 +268,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')
@@ -274,7 +342,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 +431,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);
@@ -401,8 +475,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 +486,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 +498,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 +714,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 <hip/amd_detail/hip_api_trace.hpp>
+#include <hip/hip_deprecated.h> // 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<hip::Device*> 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-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 <hip/hip_runtime_api.h>
+#include <catch2/catch_all.hpp>
+
+#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
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..914efb73fa12 100644
--- a/projects/hip/include/hip/hip_runtime_api.h
+++ b/projects/hip/include/hip/hip_runtime_api.h
@@ -30,6 +30,10 @@
 #include <hip/hip_common.h>
 #include <hip/linker_types.h>
 
+#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,163 @@ typedef struct hipUUID_t {
 //---
 // Common headers for both NVCC and HIP-Clang paths:
 
-#define hipGetDeviceProperties hipGetDevicePropertiesR0600
-#define hipDeviceProp_t hipDeviceProp_tR0600
-#define hipChooseDevice hipChooseDeviceR0600
+// 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.
+*/
+
+#if HIP_FORCE_API_VERSION < 600
+
+#include <hip/hip_deprecated.h>
+// 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 +395,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)
  *
@@ -271,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
@@ -657,11 +679,16 @@ enum hipGPUDirectRDMAWritesOrdering {
 #include <hip/driver_types.h>
 #include <hip/texture_types.h>
 #include <hip/surface_types.h>
-#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
   $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}>
   $<INSTALL_INTERFACE:include/hrr/playback>
 )
+target_compile_definitions(hrr_reader PRIVATE HIP_ABI_IMPL)
 target_compile_definitions(hrr_reader PRIVATE
   HRR_API_ARGS_IMPLEMENTATION
   $<$<BOOL:${WIN32}>: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=\"$<IF:$<BOOL:$<CONFIG>>,$<CONFIG>,unknown>\""