Merge ~tchavadar/ubuntu/+source/rocm-llvm:ubuntu/devel into ubuntu/+source/rocm-llvm:ubuntu/devel

[Talha Can Havadar]

Merge request to debian is also created here: https://salsa.debian.org/rocm-team/rocm-llvm/-/merge_requests/3

This upgrades the baseline for libamd-comgr and rocm-device-libs to rocm-7.1.0 version of https://github.com/ROCm/llvm-project

We need 7.1.0 because we are targeting to have at least ROCm 7.1.0 available in ubuntu resolute with whole stack.

[Talha Can Havadar (tchavadar)]

Just cleaned up the noise coming from gbp import-orig

[Andreas Hasenack (ahasenack)]

This commit is changing two files, and not just the one claimed by the commit mesage:

commit aa6b5d7f9a950b908b63c4c8cefdc9638e09fc07 (tchavadar/ubuntu/devel)
Author: Talha Can Havadar <email address hidden>
Date: Thu Dec 18 12:06:44 2025 +0100

    d/changelog: update changelog for resolute

It's also changing the maintainer in d/control ("update-maintainer"). You could amend the commit message, or separate the d/control change into another commit.

[Andreas Hasenack (ahasenack)]

Do you have a PPA build with this branch, and autopkgtests triggered? I wanted to check build logs, and also run autopkgtests with that build.

[Talha Can Havadar (tchavadar)]

Hello Andreas, thank you very much for the review, yes I built this package on ppa: https://launchpad.net/~tchavadar/+archive/ubuntu/rocm-with-llvm-21

Didnt triggered the autopkgtests though, gonna create a new ppa with proposed enabled and trigger the autopkgtests there now

[Talha Can Havadar (tchavadar)]

Uploaded the package to this ppa: https://launchpad.net/~tchavadar/+archive/ubuntu/lp-rocm-llvm-7.1.0-upgrade/+packages (enabled proposed)

Triggered the autopkgtest for the package, right now they are in queue

[Talha Can Havadar (tchavadar)]

I am not sure whats going on with autopkgtest.ubuntu.com but these simple tests are still running suprisingly

[Andreas Hasenack (ahasenack)]

Probably related to the LP outage. Can you trigger them again, or does that give you back an error along the lines of "this test is already running"?

[Talha Can Havadar (tchavadar)]

yes exactly:
```
You submitted an invalid request:

Test already running:
```

[Talha Can Havadar (tchavadar)]

this is from my local run: https://pastebin.ubuntu.com/p/RTvGYndFrz/

only amd64 though

Adding summary here (for when pastebin expires):
```
autopkgtest [13:46:34]: @@@@@@@@@@@@@@@@@@@@ summary
hipcc.sh PASS (superficial)
clang.sh PASS (superficial)
hipconfig.sh PASS (superficial)
gcc.sh PASS (superficial)
g++.sh PASS (superficial)
```

[Andreas Hasenack (ahasenack)]

Please hop into the #debcrafters MM channel with the UUIDs of the stale running tests, and ask them to kill them, so you can trigger new runs

[Andreas Hasenack (ahasenack)]

- ubuntu version: ok
- new orig tarball (dfsg): created correctly, matches upstream, OK
- upstream changes: no release notes, but diff is small, inspected manually: OK
- autopkgtests: thanks for the local run, that will be enough here. If there are failures in other arches, we will deal with them after the upload
- build logs: OK, particularly no new symbols in the library this produces

Since architectures were removed, this might require an archive admin intervention before it migrates. Please keep an eye out for it in https://ubuntu-archive-team.ubuntu.com/proposed-migration/update_excuses.html#rocm-llvm (the link may take a few hours to become live).

review: Approve

[Andreas Hasenack (ahasenack)]

Sponsored for resolute:

Uploading rocm-llvm_7.1.0+dfsg-0ubuntu1.dsc
Uploading rocm-llvm_7.1.0+dfsg.orig.tar.xz
Uploading rocm-llvm_7.1.0+dfsg-0ubuntu1.debian.tar.xz
Uploading rocm-llvm_7.1.0+dfsg-0ubuntu1_source.buildinfo
Uploading rocm-llvm_7.1.0+dfsg-0ubuntu1_source.changes

[Talha Can Havadar (tchavadar)]

Thank you very much Andreas I will be following update_excuses

[Andreas Hasenack (ahasenack)]

rocm-llvm is now stuck because of this[1]:

trying: rocm-llvm
skipped: rocm-llvm (34, 0, 499)
    got: 74+682: a-8:a-682:a-4:a-28:i-7:p-14:r-4:s-9
    * ppc64el: hip-utils, libamdhip64-dev, libhipsparse-dev, librccl-dev, librocrand-dev, libstarpu-contrib-dev, libstarpu-dev, libstdgpu-hip-dev

That comes from the update_output.txt report. It's a bit cryptic, but it means that if the rocm-llvm new binaries are installed on a resolute system, certain packages will be removed, and it lists those packages in the case of ppc64el (but there are removals in other architectures too, it's the number after the "minus" sign in the "got" line).

To see what is going on, one way is to bring up a resolute container, update it, install current rocm-llvm binaries and those packages listed, and then update to rocm-llvm binaries from proposed, and see why we are getting removals.

Since this upload dropped architectures, it's very likely that reverse-dependencies will have to be updated and uploaded too, also dropping those architectures.

1. https://ubuntu-archive-team.ubuntu.com/proposed-migration/update_output.txt#:~:text=skipped%3A%20rocm%2Dllvm%20(34%2C%200%2C%20499)

[Andreas Hasenack (ahasenack)]

I'm thinking this is mostly about the dropped architectures. I did a quick check on amd64 and arm64, and the only thing I noticed is that libstarpu-contrib-dev and libstarpu-dev conflict with each other, but that isn't new.

I would start with the dropped architectures, and once that is resolved, see what remains in terms of migration problems.

[Talha Can Havadar (tchavadar)]

Installed these packages (list from rocm meta package that we are planning to release) in a resolute container and enabled resolute-proposed
```
Depends: ${misc:Depends},
         amd-dbgapi-dev,
         hip-utils,
         hipcc,
         hipify-perl,
         libamd-smi-dev [amd64],
         libamdhip64-dev,
         libhipblas-common-dev,
         libhipblas-dev,
         libhipcub-dev,
         libhipfft-dev,
         libhiprand-dev,
         libhipsolver-dev,
         libhipsparse-dev,
         libmiopen-dev,
         liboam-dev,
         librccl-dev,
         librocalution-dev,
         librocblas-dev,
         librocfft-dev,
         librocm-smi-dev,
         librocprim-dev,
         librocrand-dev,
         librocsolver-dev,
         librocsparse-dev,
         librocthrust-dev,
         libroctx-dev,
         rocminfo,
```

Then I installed the binaries we provided in rocm-llvm, it didnt try to delete anything on amd64
```
root@pretty-antelope:~# apt install libamd-comgr-dev=7.1.0+dfsg-0ubuntu1 hipcc=7.1.0+dfsg-0ubuntu1 rocm-device-libs-21=7.1.0+dfsg-0ubuntu1 libamd-comgr3=7.1.0+dfsg-0ubuntu1
Upgrading:
  hipcc libamd-comgr-dev libamd-comgr3 rocm-device-libs-21

Summary:
  Upgrading: 4, Installing: 0, Removing: 0, Not Upgrading: 56
```

Should I investigate each package being removed in other architectures and provide update for them to not build for that arch?

[Andreas Hasenack (ahasenack)]

> Should I investigate each package being removed in other architectures and provide update for them to not
> build for that arch?

I think so, that would be a good starting point and clear up the report, leaving only the non-obvious problems.

[Andreas Hasenack (ahasenack)]

> I think so, that would be a good starting point and clear up the report, leaving only the non-obvious
> problems.

On second note, please don't do that. I just conferred with other archive admins, and the best way out is to remove the binaries of these other architectures in the reverse build dependencies.

[Talha Can Havadar (tchavadar)]

I was just planning to find some tool or write my own to create MPs for all ahahaha, phew :S thank you!

[Andreas Hasenack (ahasenack)]

Talha, here is something we missed: the architectures line in rocm-device-libs-21:

Package: rocm-device-libs-21
Architecture: any
Section: libs

This is what caused the package to build on ppc64el, and many of those errors.

Could you please prepare a new PR/upload with that line also updated to match amd64/arm64 like the other binary packages, and test a build on a ppa with ALL architectures enabled, to confirm it will only build on amd64/amd64v3/arm64?

[Talha Can Havadar (tchavadar)]

yep of course, I'm on it, gonna create a new merge proposal

diff --git a/amd/comgr/CMakeLists.txt b/amd/comgr/CMakeLists.txt
index cfa170f..2eeb8c2 100644
--- a/amd/comgr/CMakeLists.txt
+++ b/amd/comgr/CMakeLists.txt
@@ -38,7 +38,6 @@ endif(EXISTS "${CMAKE_SOURCE_DIR}/../../.git")
 
 include(GNUInstallDirs)
 include(CMakePackageConfigHelpers)
-option(FILE_REORG_BACKWARD_COMPATIBILITY "Enable File Reorg with backward compatibility" OFF)
 # Optionally, build Compiler Support with ccache.
 set(ROCM_CCACHE_BUILD OFF CACHE BOOL "Set to ON for a ccache enabled build")
 if (ROCM_CCACHE_BUILD)
@@ -327,26 +326,6 @@ install(FILES
   COMPONENT amd-comgr
   DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/${amd_comgr_NAME})
 
-#File reorg Backward compatibility function
-if(FILE_REORG_BACKWARD_COMPATIBILITY)
-# To enable/disable #error in wrapper header files
- if(NOT DEFINED ROCM_HEADER_WRAPPER_WERROR)
-      if(DEFINED ENV{ROCM_HEADER_WRAPPER_WERROR})
-          set(ROCM_HEADER_WRAPPER_WERROR "$ENV{ROCM_HEADER_WRAPPER_WERROR}"
-                  CACHE STRING "Header wrapper warnings as errors.")
-      else()
-          set(ROCM_HEADER_WRAPPER_WERROR "OFF" CACHE STRING "Header wrapper warnings as errors.")
-      endif()
-  endif()
-  if(ROCM_HEADER_WRAPPER_WERROR)
-      set(deprecated_error 1)
-  else()
-      set(deprecated_error 0)
-  endif()
-
-  include(comgr-backward-compat.cmake)
-endif()
-
 if(ENABLE_ASAN_PACKAGING)
   install(FILES
     "LICENSE.txt"
diff --git a/amd/comgr/comgr-backward-compat.cmake b/amd/comgr/comgr-backward-compat.cmake
deleted file mode 100644
index c4f5992..0000000
--- a/amd/comgr/comgr-backward-compat.cmake
+++ /dev/null
@@ -1,96 +0,0 @@
-# Copyright (c) 2022 Advanced Micro Devices, Inc. All Rights Reserved.
-# Permission is hereby granted, free of charge, to any person obtaining a copy
-# of this software and associated documentation files (the "Software"), to deal
-# in the Software without restriction, including without limitation the rights
-# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
-# copies of the Software, and to permit persons to whom the Software is
-# furnished to do so, subject to the following conditions:
-#
-# The above copyright notice and this permission notice shall be included in
-# all copies or substantial portions of the Software.
-#
-# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
-# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
-# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
-# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
-# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
-# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
-# THE SOFTWARE.
-
-cmake_minimum_required(VERSION 3.16.8)
-
-set(COMGR_BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR})
-set(COMGR_WRAPPER_DIR ${COMGR_BUILD_DIR}/wrapper_dir)
-set(COMGR_WRAPPER_INC_DIR ${COMGR_WRAPPER_DIR}/include)
-
-#Function to generate header template file
-function(create_header_template)
-    file(WRITE ${COMGR_WRAPPER_DIR}/header.hpp.in "/*
-    Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
-
-    Permission is hereby granted, free of charge, to any person obtaining a copy
-    of this software and associated documentation files (the \"Software\"), to deal
-    in the Software without restriction, including without limitation the rights
-    to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
-    copies of the Software, and to permit persons to whom the Software is
-    furnished to do so, subject to the following conditions:
-
-    The above copyright notice and this permission notice shall be included in
-    all copies or substantial portions of the Software.
-
-    THE SOFTWARE IS PROVIDED \"AS IS\", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
-    IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
-    FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
-    AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
-    LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
-    OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
-   THE SOFTWARE.
-   */
-
-#ifndef @include_guard@
-#define @include_guard@
-
-#ifndef ROCM_HEADER_WRAPPER_WERROR
-#define ROCM_HEADER_WRAPPER_WERROR @deprecated_error@
-#endif
-#if ROCM_HEADER_WRAPPER_WERROR  /* ROCM_HEADER_WRAPPER_WERROR 1 */
-#error \"This file is deprecated. Use file from include path /opt/rocm-ver/include/ and prefix with amd_comgr\"
-#else   /* ROCM_HEADER_WRAPPER_WERROR 0 */
-#if defined(__GNUC__)
-#warning \"This file is deprecated. Use file from include path /opt/rocm-ver/include/ and prefix with amd_comgr\"
-#else
-#pragma message(\"This file is deprecated. Use file from include path /opt/rocm-ver/include/ and prefix with amd_comgr\")
-#endif
-#endif /* ROCM_HEADER_WRAPPER_WERROR */
-
-@include_statements@
-
-#endif")
-endfunction()
-
-#use header template file and generate wrapper header files
-function(generate_wrapper_header)
-  file(MAKE_DIRECTORY ${COMGR_WRAPPER_INC_DIR})
-  #find all header files(*.h) from include
-  file(GLOB include_files ${COMGR_BUILD_DIR}/include/*.h)
-  #Generate wrapper header files for each files in the list
-  foreach(header_file ${include_files})
-    # set include guard
-    get_filename_component(INC_GAURD_NAME ${header_file} NAME_WE)
-    string(TOUPPER ${INC_GAURD_NAME} INC_GAURD_NAME)
-    set(include_guard "${include_guard}COMGR_WRAPPER_INCLUDE_${INC_GAURD_NAME}_H")
-    #set #include statement
-    get_filename_component(file_name ${header_file} NAME)
-    set(include_statements "${include_statements}#include \"${amd_comgr_NAME}/${file_name}\"\n")
-    configure_file(${COMGR_WRAPPER_DIR}/header.hpp.in ${COMGR_WRAPPER_INC_DIR}/${file_name})
-    unset(include_statements)
-    unset(include_guard)
-  endforeach()
-
-endfunction()
-
-#Creater a template for header file
-create_header_template()
-#Use template header file and generater wrapper header files
-generate_wrapper_header()
-install(DIRECTORY ${COMGR_WRAPPER_INC_DIR} COMPONENT amd-comgr DESTINATION .)
diff --git a/amd/comgr/src/comgr-isa-metadata.def b/amd/comgr/src/comgr-isa-metadata.def
index a46ee01..c926978 100644
--- a/amd/comgr/src/comgr-isa-metadata.def
+++ b/amd/comgr/src/comgr-isa-metadata.def
@@ -68,10 +68,10 @@ HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx904",           false,  true, EF_AMDGPU_MAC
 HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx906",            true,  true, EF_AMDGPU_MACH_AMDGCN_GFX906,          true,  true, 65536,  32,  4,   40, 1024,   16, 800, 102,    4, 256, 256)
 HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx908",            true,  true, EF_AMDGPU_MACH_AMDGCN_GFX908,          true,  true, 65536,  32,  4,   40, 1024,   16, 800, 102,    4, 256, 256)
 HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx909",           false,  true, EF_AMDGPU_MACH_AMDGCN_GFX909,          true,  true, 65536,  32,  4,   40, 1024,   16, 800, 102,    4, 256, 256)
-HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx90a",            true,  true, EF_AMDGPU_MACH_AMDGCN_GFX90A,          true,  true, 65536,  32,  4,   40, 1024,   16, 800, 102,    4, 256, 256)
+HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx90a",            true,  true, EF_AMDGPU_MACH_AMDGCN_GFX90A,          true,  true, 65536,  32,  4,   40, 1024,   16, 800, 102,    8, 512, 512)
 HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx90c",           false,  true, EF_AMDGPU_MACH_AMDGCN_GFX90C,          true,  true, 65536,  32,  4,   40, 1024,   16, 800, 102,    4, 256, 256)
-HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx942",            true,  true, EF_AMDGPU_MACH_AMDGCN_GFX942,          true, false, 65536,  32,  4,   40, 1024,   16, 800, 102,    4, 256, 256)
-HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx950",            true,  true, EF_AMDGPU_MACH_AMDGCN_GFX950,          true, false, 65536,  32,  4,   40, 1024,   16, 800, 102,    4, 256, 256)
+HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx942",            true,  true, EF_AMDGPU_MACH_AMDGCN_GFX942,          true, false, 65536,  32,  4,   40, 1024,   16, 800, 102,    8, 512, 512)
+HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx950",            true,  true, EF_AMDGPU_MACH_AMDGCN_GFX950,          true, false, 65536,  32,  4,   40, 1024,   16, 800, 102,    8, 512, 512)
 HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx1010",          false,  true, EF_AMDGPU_MACH_AMDGCN_GFX1010,         true,  true, 65536,  32,  4,   40, 1024,  106, 800, 106,    8, 256, 256)
 HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx1011",          false,  true, EF_AMDGPU_MACH_AMDGCN_GFX1011,         true,  true, 65536,  32,  4,   40, 1024,  106, 800, 106,    8, 256, 256)
 HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx1012",          false,  true, EF_AMDGPU_MACH_AMDGCN_GFX1012,         true,  true, 65536,  32,  4,   40, 1024,  106, 800, 106,    8, 256, 256)
@@ -95,7 +95,7 @@ HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx1200",          false, false, EF_AMDGPU_MAC
 HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx1201",          false, false, EF_AMDGPU_MACH_AMDGCN_GFX1201,         true,  true, 65536,  32,  4,   40, 1024,  106, 800, 106,   24, 1536, 256)
 
 HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx9-generic",     false,  true, EF_AMDGPU_MACH_AMDGCN_GFX9_GENERIC,    true,  true, 65536,  32,  4,   40, 1024,   16, 800, 102,    4, 256, 256)
-HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx9-4-generic",   true,   true, EF_AMDGPU_MACH_AMDGCN_GFX9_4_GENERIC,  true, false, 65536,  32,  4,   40, 1024,   16, 800, 102,    4, 256, 256)
+HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx9-4-generic",   true,   true, EF_AMDGPU_MACH_AMDGCN_GFX9_4_GENERIC,  true, false, 65536,  32,  4,   40, 1024,   16, 800, 102,    8, 512, 512)
 HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx10-1-generic",  false,  true, EF_AMDGPU_MACH_AMDGCN_GFX10_1_GENERIC, true,  true, 65536,  32,  4,   40, 1024,  106, 800, 106,    8, 256, 256)
 HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx10-3-generic",  false, false, EF_AMDGPU_MACH_AMDGCN_GFX10_3_GENERIC, true,  true, 65536,  32,  4,   40, 1024,  106, 800, 106,    8, 256, 256)
 HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx11-generic",    false, false, EF_AMDGPU_MACH_AMDGCN_GFX11_GENERIC,   true,  true, 65536,  32,  4,   40, 1024,  106, 800, 106,   16, 1024, 256)
diff --git a/amd/comgr/src/comgr.cpp b/amd/comgr/src/comgr.cpp
index 28c2c4e..053c030 100644
--- a/amd/comgr/src/comgr.cpp
+++ b/amd/comgr/src/comgr.cpp
@@ -344,8 +344,17 @@ amd_comgr_status_t COMGR::parseTargetIdentifier(StringRef IdentStr,
   Ident.Processor = Ident.Features[0];
   Ident.Features.erase(Ident.Features.begin());
 
-  size_t IsaIndex;
 
+  // TODO: Add a LIT test for this
+  if (IdentStr == "amdgcn-amd-amdhsa--amdgcnspirv" ||
+      IdentStr == "amdgcn-amd-amdhsa-unknown-amdgcnspirv") {
+    // Features not supported for SPIR-V
+    if (!Ident.Features.empty())
+      return AMD_COMGR_STATUS_ERROR_INVALID_ARGUMENT;
+    return AMD_COMGR_STATUS_SUCCESS;
+  }
+
+  size_t IsaIndex;
   amd_comgr_status_t Status = metadata::getIsaIndex(IdentStr, IsaIndex);
   if (Status != AMD_COMGR_STATUS_SUCCESS) {
     return Status;
@@ -1017,6 +1026,10 @@ amd_comgr_status_t AMD_COMGR_API
     return AMD_COMGR_STATUS_SUCCESS;
   }
 
+  if (StringRef(IsaName) == "amdgcn-amd-amdhsa--amdgcnspirv") {
+    return ActionP->setIsaName(IsaName);
+  }
+
   if (!metadata::isValidIsaName(IsaName)) {
     return AMD_COMGR_STATUS_ERROR_INVALID_ARGUMENT;
   }
@@ -2129,6 +2142,8 @@ amd_comgr_populate_name_expression_map(amd_comgr_data_t Data, size_t *Count) {
     if (!RelaRangeOrError) {
       llvm::logAllUnhandledErrors(RelaRangeOrError.takeError(), llvm::errs(),
                                   "RelaRange creation error: ");
+      for (auto *Ptr : NameExpDataVec)
+        delete Ptr;
       return AMD_COMGR_STATUS_ERROR;
     }
     auto RelaRange = std::move(RelaRangeOrError.get());
@@ -2149,6 +2164,8 @@ amd_comgr_populate_name_expression_map(amd_comgr_data_t Data, size_t *Count) {
     if (!RodataOrError) {
       llvm::logAllUnhandledErrors(RodataOrError.takeError(), llvm::errs(),
                                   "Rodata creation error: ");
+      for (auto *Ptr : NameExpDataVec)
+        delete Ptr;
       return AMD_COMGR_STATUS_ERROR;
     }
     auto Rodata = std::move(RodataOrError.get());
@@ -2179,6 +2196,8 @@ amd_comgr_populate_name_expression_map(amd_comgr_data_t Data, size_t *Count) {
       }
     }
 
+    for (auto *Ptr : NameExpDataVec)
+      delete Ptr;
   } // end AMD_COMGR_DATA_KIND_EXECUTABLE conditional
 
   *Count = DataP->NameExpressionMap.size();
diff --git a/amd/device-libs/opencl/src/devenq/schedule_rocm.cl b/amd/device-libs/opencl/src/devenq/schedule_rocm.cl
index 731da88..209eebe 100644
--- a/amd/device-libs/opencl/src/devenq/schedule_rocm.cl
+++ b/amd/device-libs/opencl/src/devenq/schedule_rocm.cl
@@ -61,16 +61,23 @@ min_command(uint slot_num, __global AmdAqlWrap* wraps)
     return minCommand;
 }
 
+static inline bool
+check_pcie_support(__global SchedulerParam* param) {
+  #define kInvalidWriteIndex (ulong)(-1)
+  return (param->write_index == kInvalidWriteIndex) ? true : false;
+}
+
 static inline void
 EnqueueDispatch(__global hsa_kernel_dispatch_packet_t* aqlPkt, __global SchedulerParam* param)
 {
     __global hsa_queue_t* child_queue = param->child_queue;
 
-
-    // ulong index = __ockl_hsa_queue_add_write_index(child_queue, 1, __ockl_memory_order_relaxed);
-    // The original code seen above relies on PCIe 3 atomics, which might not be supported on some systems, so use a device side global
-    // for workaround.
-    ulong index = atomic_fetch_add_explicit((__global atomic_ulong*)&param->write_index, (ulong)1, memory_order_relaxed, memory_scope_device);
+    ulong index;
+    if (check_pcie_support(param)) {
+      index = __ockl_hsa_queue_add_write_index(child_queue, 1, __ockl_memory_order_relaxed);
+    } else {
+      index = atomic_fetch_add_explicit((__global atomic_ulong*)&param->write_index, (ulong)1, memory_order_relaxed, memory_scope_device);
+    }
 
     const ulong queueMask = child_queue->size - 1;
     __global hsa_kernel_dispatch_packet_t* dispatch_packet = &(((__global hsa_kernel_dispatch_packet_t*)(child_queue->base_address))[index & queueMask]);
@@ -82,17 +89,20 @@ EnqueueScheduler(__global SchedulerParam* param)
 {
     __global hsa_queue_t* child_queue = param->child_queue;
 
-    // ulong index = __ockl_hsa_queue_add_write_index(child_queue, 1, __ockl_memory_order_relaxed);
-    // The original code seen above relies on PCIe 3 atomics, which might not be supported on some systems, so use a device side global
-    // for workaround.
-    ulong index = atomic_fetch_add_explicit((__global atomic_ulong*)&param->write_index, (ulong)1, memory_order_relaxed, memory_scope_device);
+    ulong index;
+    if (check_pcie_support(param)) {
+      index = __ockl_hsa_queue_add_write_index(child_queue, 1, __ockl_memory_order_relaxed);
+    } else {
+      index = atomic_fetch_add_explicit((__global atomic_ulong*)&param->write_index, (ulong)1, memory_order_relaxed, memory_scope_device);
+    }
 
     const ulong queueMask = child_queue->size - 1;
     __global hsa_kernel_dispatch_packet_t* dispatch_packet = &(((__global hsa_kernel_dispatch_packet_t*)(child_queue->base_address))[index & queueMask]);
     *dispatch_packet = param->scheduler_aql;
 
-     // This is part of the PCIe 3 atomics workaround, to write the final write_index value back to the child_queue
-    __ockl_hsa_queue_store_write_index(child_queue, index + 1, __ockl_memory_order_relaxed);
+    if (!check_pcie_support(param)) {
+      __ockl_hsa_queue_store_write_index(child_queue, index + 1, __ockl_memory_order_relaxed);
+    }
 
     __ockl_hsa_signal_store(child_queue->doorbell_signal, index, __ockl_memory_order_release);
 }
diff --git a/amd/hipcc/bin/hipvars.pm b/amd/hipcc/bin/hipvars.pm
new file mode 100644
index 0000000..93fba75
--- /dev/null
+++ b/amd/hipcc/bin/hipvars.pm
@@ -0,0 +1,171 @@
+#!/usr/bin/env perl
+# Copyright (c) 2020 - 2021 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in
+# all copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+# THE SOFTWARE.
+
+package hipvars;
+use warnings;
+use Getopt::Long;
+use Cwd;
+use File::Basename;
+
+$HIP_BASE_VERSION_MAJOR = "6";
+$HIP_BASE_VERSION_MINOR = "2";
+$HIP_BASE_VERSION_PATCH = "0";
+
+#---
+# Function to parse config file
+sub parse_config_file {
+    my ($file, $config) = @_;
+    if (open (CONFIG, "$file")) {
+        while (<CONFIG>) {
+            my $config_line=$_;
+            chop ($config_line);
+            $config_line =~ s/^\s*//;
+            $config_line =~ s/\s*$//;
+            if (($config_line !~ /^#/) && ($config_line ne "")) {
+                my ($name, $value) = split (/=/, $config_line);
+                $$config{$name} = $value;
+            }
+        }
+        close(CONFIG);
+    }
+}
+
+#---
+# Function to check if executable can be run
+sub can_run {
+    my ($exe) = @_;
+    `$exe --version 2>&1`;
+    if ($? == 0) {
+        return 1;
+    } else {
+        return 0;
+    }
+}
+
+$isWindows =  ($^O eq 'MSWin32' or $^O eq 'msys');
+$doubleQuote = "\"";
+
+#
+# TODO: Fix rpath LDFLAGS settings
+#
+# Since this hipcc script gets installed at two uneven hierarchical levels,
+# linked by symlink, the absolute path of this script should be used to
+# derive HIP_PATH, as dirname $0 could be /opt/rocm/bin or /opt/rocm/hip/bin
+# depending on how it gets invoked.
+# ROCM_PATH which points to <rocm_install_dir> is determined based on whether
+# we find bin/rocm_agent_enumerator in the parent of HIP_PATH or not. If it is found,
+# ROCM_PATH is defined relative to HIP_PATH else it is hardcoded to /opt/rocm.
+#
+$HIP_PATH=$ENV{'HIP_PATH'} // dirname(Cwd::abs_path("$0/../")); # use parent directory of hipcc
+if ($isWindows and defined $ENV{'HIP_PATH'}) {
+  $HIP_PATH =~ s/^"(.*)"$/$1/;
+  $HIP_PATH =~ s/\\/\//g;
+}
+if (-e "$HIP_PATH/bin/rocm_agent_enumerator") {
+    $ROCM_PATH=$ENV{'ROCM_PATH'} // "$HIP_PATH"; # use HIP_PATH
+}elsif (-e "$HIP_PATH/../bin/rocm_agent_enumerator") { # case for backward compatibility
+    $ROCM_PATH=$ENV{'ROCM_PATH'} // dirname("$HIP_PATH"); # use parent directory of HIP_PATH
+} else {
+    $ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm";
+}
+$CUDA_PATH=$ENV{'CUDA_PATH'} // '/usr/local/cuda';
+if ($isWindows and defined $ENV{'CUDA_PATH'}) {
+  $CUDA_PATH =~ s/^"(.*)"$/$1/;
+  $CUDA_PATH =~ s/\\/\//g;
+}
+
+# Windows/Distro's have a different structure, all binaries are with hipcc
+if ($isWindows or -e "$HIP_PATH/bin/clang") {
+    $HIP_CLANG_PATH=$ENV{'HIP_CLANG_PATH'} // "$HIP_PATH/bin";
+} else {
+    $HIP_CLANG_PATH=$ENV{'HIP_CLANG_PATH'} // "$ROCM_PATH/lib/llvm/bin";
+}
+# HIP_ROCCLR_HOME is used by Windows builds
+$HIP_ROCCLR_HOME=$ENV{'HIP_ROCCLR_HOME'};
+
+if (defined $HIP_ROCCLR_HOME) {
+    $HIP_INFO_PATH= "$HIP_ROCCLR_HOME/lib/.hipInfo";
+} else {
+    $HIP_INFO_PATH= "$HIP_PATH/lib/.hipInfo"; # use actual file
+}
+#---
+#HIP_PLATFORM controls whether to use nvidia or amd platform:
+$HIP_PLATFORM=$ENV{'HIP_PLATFORM'};
+# Read .hipInfo
+my %hipInfo = ();
+parse_config_file("$HIP_INFO_PATH", \%hipInfo);
+# Prioritize Env first, otherwise use the hipInfo config file
+$HIP_COMPILER = $ENV{'HIP_COMPILER'} // $hipInfo{'HIP_COMPILER'} // "clang";
+$HIP_RUNTIME = $ENV{'HIP_RUNTIME'} // $hipInfo{'HIP_RUNTIME'} // "rocclr";
+
+# If using ROCclr runtime, need to find HIP_ROCCLR_HOME
+if (defined $HIP_RUNTIME and $HIP_RUNTIME eq "rocclr" and !defined $HIP_ROCCLR_HOME) {
+    my $hipvars_dir = dirname(Cwd::abs_path($0));
+    if (-e "$hipvars_dir/../lib/bitcode") {
+        $HIP_ROCCLR_HOME = Cwd::abs_path($hipvars_dir . "/.."); #FILE_REORG Backward compatibility
+    } elsif (-e "$hipvars_dir/lib/bitcode") {
+        $HIP_ROCCLR_HOME = Cwd::abs_path($hipvars_dir);
+    } else {
+        $HIP_ROCCLR_HOME = $HIP_PATH; # use HIP_PATH
+    }
+}
+
+if (not defined $HIP_PLATFORM) {
+    if (can_run($doubleQuote . "$HIP_CLANG_PATH/clang++" . $doubleQuote) or can_run("amdclang++")) {
+        $HIP_PLATFORM = "amd";
+    } elsif (can_run($doubleQuote . "$CUDA_PATH/bin/nvcc" . $doubleQuote) or can_run("nvcc")) {
+        $HIP_PLATFORM = "nvidia";
+        $HIP_COMPILER = "nvcc";
+        $HIP_RUNTIME = "cuda";
+    } else {
+        # Default to amd for now
+        $HIP_PLATFORM = "amd";
+    }
+} elsif ($HIP_PLATFORM eq "hcc") {
+    $HIP_PLATFORM = "amd";
+    warn("Warning: HIP_PLATFORM=hcc is deprecated. Please use HIP_PLATFORM=amd. \n")
+} elsif ($HIP_PLATFORM eq "nvcc") {
+    $HIP_PLATFORM = "nvidia";
+    $HIP_COMPILER = "nvcc";
+    $HIP_RUNTIME = "cuda";
+    warn("Warning: HIP_PLATFORM=nvcc is deprecated. Please use HIP_PLATFORM=nvidia. \n")
+}
+
+if ($HIP_COMPILER eq "clang") {
+    # Windows does not have clang at linux default path
+    if (defined $HIP_ROCCLR_HOME and (-e "$HIP_ROCCLR_HOME/bin/clang" or -e "$HIP_ROCCLR_HOME/bin/clang.exe")) {
+        $HIP_CLANG_PATH = "$HIP_ROCCLR_HOME/bin";
+    }
+}
+
+#---
+# Read .hipVersion
+my %hipVersion = ();
+if ($isWindows) {
+    parse_config_file("$hipvars::HIP_PATH/bin/.hipVersion", \%hipVersion);
+} else {
+    parse_config_file("$hipvars::HIP_PATH/share/hip/version", \%hipVersion);
+}
+$HIP_VERSION_MAJOR = $hipVersion{'HIP_VERSION_MAJOR'} // $HIP_BASE_VERSION_MAJOR;
+$HIP_VERSION_MINOR = $hipVersion{'HIP_VERSION_MINOR'} // $HIP_BASE_VERSION_MINOR;
+$HIP_VERSION_PATCH = $hipVersion{'HIP_VERSION_PATCH'} // $HIP_BASE_VERSION_PATCH;
+$HIP_VERSION_GITHASH = $hipVersion{'HIP_VERSION_GITHASH'} // 0;
+$HIP_VERSION="$HIP_VERSION_MAJOR.$HIP_VERSION_MINOR.$HIP_VERSION_PATCH-$HIP_VERSION_GITHASH";
diff --git a/debian/changelog b/debian/changelog
index c8062a6..1de8861 100644
--- a/debian/changelog
+++ b/debian/changelog
@@ -1,3 +1,10 @@
+rocm-llvm (7.1.0+dfsg-0ubuntu1) resolute; urgency=medium
+
+  * New upstream version 7.1.0
+  * d/control: remove unsupported architectures riscv64 and ppc64el
+
+ -- Talha Can Havadar <[e-mail]>  Thu, 15 Jan 2026 15:09:35 +0100
+
 rocm-llvm (7.0.2+dfsg-2) unstable; urgency=medium
 
   * Change HIP_CLANG_PATH from /usr/bin to /usr/lib/llvm-21/bin because
diff --git a/debian/control b/debian/control
index 802a0c6..df5edcc 100644
--- a/debian/control
+++ b/debian/control
@@ -1,7 +1,8 @@
 Source: rocm-llvm
 Section: devel
 Priority: optional
-Maintainer: Debian ROCm Team <[e-mail]>
+Maintainer: Ubuntu Developers <[e-mail]>
+XSBC-Original-Maintainer: Debian ROCm Team <[e-mail]>
 Uploaders: Xuanteng Huang <[e-mail]>,
            Christian Kastner <[e-mail]>,
            Cordell Bloor <[e-mail]>
@@ -28,7 +29,7 @@ Rules-Requires-Root: no
 
 Package: hipcc
 Section: devel
-Architecture: amd64 arm64 ppc64el riscv64
+Architecture: amd64 arm64
 Depends: ${misc:Depends},
          ${shlibs:Depends},
          clang-21,
@@ -77,7 +78,7 @@ Description: AMD specific device-side language runtime libraries
  Heterogeneous Compute built-in library.
 
 Package: libamd-comgr3
-Architecture: amd64 arm64 ppc64el riscv64
+Architecture: amd64 arm64
 Section: libs
 Depends: ${misc:Depends},
          ${shlibs:Depends},
@@ -88,7 +89,7 @@ Description: ROCm code object manager
  header file /usr/include/amd_comgr/amd_comgr.h .
 
 Package: libamd-comgr-dev
-Architecture: amd64 arm64 ppc64el riscv64
+Architecture: amd64 arm64
 Section: libdevel
 Depends: libamd-comgr3 (= ${binary:Version}),
          ${misc:Depends},