[hipBLASLt] Deprecate rocm-smi and use amdsmi (#8527)

## Motivation

rocm-smi will be removed from TheRock, and users/customers will move to
using AMD SMI library and API’s

<!-- Explain the purpose of this PR and the goals it aims to achieve.
-->

## Technical Details

- Change all rocm-smi to amdsmi

<!-- Explain the changes along with any relevant GitHub links. -->

## Test Plan
- I also add the unit tests:

`python3 -m pytest -v \
  Tensile/Tests/unit/test_detectAvailableGpus.py \
  Tensile/Tests/unit/test_get_gpu_max_frequency_smi.py \
  Tensile/Tests/unit/test_specs_amdsmi.py`

<!-- Explain any relevant testing done to verify this PR. -->

## Test Result

<img width="2131" height="732" alt="image"
src="https://github.com/user-attachments/assets/26b1e0a2-4c5a-4c66-a5af-d256ba94b4b1"
/>

<!-- Briefly summarize test outcomes. -->

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

---------

Co-authored-by: tomchengchitang <tom.tang@amd.com>
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>

Author: 3 people

projects/hipblaslt/clients/scripts/performance/specs.py

@@ -2,6 +2,7 @@
 # SPDX-License-Identifier: MIT
 
 from pathlib import Path
+import json
 import re
 import socket
 import subprocess
@@ -20,6 +21,83 @@ def search(pattern, string):
     return None
 
 
+def _run_amdsmi_json(cmd: list):
+    """
+    Runs an ``amd-smi`` command that emits JSON and returns the parsed object.
+
+    Returns None if amd-smi is unavailable or the output cannot be parsed.
+    """
+    try:
+        completed = subprocess.run(cmd, stdout=subprocess.PIPE, stderr=subprocess.PIPE)
+        if getattr(completed, "returncode", 0) != 0:
+            return None
+        return json.loads(completed.stdout.decode("utf-8", errors="replace"))
+    except (FileNotFoundError, json.JSONDecodeError, ValueError):
+        return None
+
+
+def get_amdsmi_specs(devicenum: int = 0) -> dict:
+    """
+    Collects per-device GPU specs using ``amd-smi`` structured JSON output.
+
+    Returns a dict with the following fields (any field that cannot be
+    resolved is returned as None):
+
+      vbios_version, gpuid, vram (bytes), performance_level,
+      memory_clk, system_clk
+    """
+    g = str(devicenum)
+    static = _run_amdsmi_json(
+        ["amd-smi", "static", "-g", g, "--asic", "--vbios", "--json"]
+    )
+    metric = _run_amdsmi_json(
+        [
+            "amd-smi", "metric", "-g", g,
+            "--mem-usage", "--clock", "--perf-level", "--json",
+        ]
+    )
+
+    result = {
+        "vbios_version": None,
+        "gpuid": None,
+        "vram": None,
+        "performance_level": None,
+        "memory_clk": None,
+        "system_clk": None,
+    }
+
+    if static and static.get("gpu_data"):
+        d = static["gpu_data"][0]
+        result["gpuid"] = d.get("asic", {}).get("device_id")
+        # rocm-smi --showvbios reported the board part number as VBIOS version.
+        result["vbios_version"] = d.get("ifwi", {}).get("part_number")
+
+    if metric and metric.get("gpu_data"):
+        d = metric["gpu_data"][0]
+
+        total = d.get("mem_usage", {}).get("total_vram", {}).get("value")
+        if total is not None:
+            # amd-smi reports total VRAM in MB; convert to bytes to match the
+            # downstream "/1024**3" GiB formatting used for rocm-smi output.
+            result["vram"] = int(total) * 1024 * 1024
+
+        perf = d.get("perf_level")
+        if isinstance(perf, str):
+            # Normalize e.g. "AMDSMI_DEV_PERF_LEVEL_AUTO" -> "auto" to match the
+            # lowercase value rocm-smi used to print.
+            result["performance_level"] = perf.split("_")[-1].lower()
+
+        clock = d.get("clock", {})
+        sclk = clock.get("gfx_0", {}).get("clk", {}).get("value")
+        if sclk is not None:
+            result["system_clk"] = f"{sclk}Mhz"
+        mclk = clock.get("mem_0", {}).get("clk", {}).get("value")
+        if mclk is not None:
+            result["memory_clk"] = f"{mclk}Mhz"
+
+    return result
+
+
 def _subprocess_helper(cmd: list) -> tuple:
     """
     This is a helper method which runs a command line argument
@@ -239,51 +317,26 @@ def get_machine_specs(filename: str, devicenum: int = 0):
             rocm_info = rocm_path.read_text()
         else:
             rocm_info = None
-        try:
-            rocm_smi = run(
-                [
-                    "rocm-smi",
-                    "--showvbios",
-                    "--showid",
-                    "--showproductname",
-                    "--showperflevel",
-                    "--showclocks",
-                    "--showmeminfo",
-                    "vram",
-                    "lshw",
-                ]
-            )
-        except FileNotFoundError as e:
-            rocm_smi = None
+        amdsmi_specs = get_amdsmi_specs(devicenum)
 
         try:
             rocminfo = run(["rocminfo"])
         except FileNotFoundError as e:
             rocminfo = None
 
-        device = rf"^GPU\[{devicenum}\]\s*: "
         hostname = socket.gethostname()
 
         if rocm_info is None:
             rocm_version = None
         else:
             rocm_version = rocm_info.strip()
 
-        if rocm_smi is None:
-            vbios_version = (
-                gpuid
-            ) = deviceinfo = vram = performance_level = memory_clk = system_clk = None
-        else:
-            if device is None:
-                device = ""
-            vbios_version = search(device + r"VBIOS version: (.*?)$", rocm_smi)
-            gpuid = search(device + r"Device ID: (.*?)$", rocm_smi)
-            if not gpuid:
-                gpuid = search(device + r"GPU ID: (.*?)$", rocm_smi)
-            vram = search(device + r".... Total Memory .B.: (\d+)$", rocm_smi)
-            performance_level = search(device + r"Performance Level: (.*?)$", rocm_smi)
-            memory_clk = search(device + r"mclk.*\((.*?)\)$", rocm_smi)
-            system_clk = search(device + r"sclk.*\((.*?)\)$", rocm_smi)
+        vbios_version = amdsmi_specs["vbios_version"]
+        gpuid = amdsmi_specs["gpuid"]
+        vram = amdsmi_specs["vram"]
+        performance_level = amdsmi_specs["performance_level"]
+        memory_clk = amdsmi_specs["memory_clk"]
+        system_clk = amdsmi_specs["system_clk"]
 
         architecture_name, internal_product_name = get_device_info()
 

projects/hipblaslt/tensilelite/Tensile/ClientWriter.py

@@ -331,10 +331,13 @@ def writeRunScript(path, forBenchmark, enableTileSelection, cxxCompiler: str, cC
       runScriptFile.write(os.path.join(globalParameters["CMakeBuildType"], \
           "client.exe") )
     else:
-      if globalParameters["PinClocks"] and globalParameters["ROCmSMIPath"]:
-        runScriptFile.write("%s -d 0 --setfan 255 --setsclk 7\n" % globalParameters["ROCmSMIPath"])
+      if globalParameters["PinClocks"] and globalParameters["AMDSMIPath"]:
+        # amd-smi set/reset require elevated privileges. Pin to max
+        # performance and run the fan at full speed for the benchmark.
+        runScriptFile.write("sudo %s set -g 0 --fan 255\n" % globalParameters["AMDSMIPath"])
+        runScriptFile.write("sudo %s set -g 0 --perf-level HIGH\n" % globalParameters["AMDSMIPath"])
         runScriptFile.write("sleep 1\n")
-        runScriptFile.write("%s -d 0 -a\n" % globalParameters["ROCmSMIPath"])
+        runScriptFile.write("%s metric -g 0 --clock\n" % globalParameters["AMDSMIPath"])
 
       runScriptFile.write("set +e\n")
 
@@ -368,9 +371,10 @@ def writeRunScript(path, forBenchmark, enableTileSelection, cxxCompiler: str, cC
 """)
 
     if os.name != "nt":
-      if globalParameters["PinClocks"] and globalParameters["ROCmSMIPath"]:
-        runScriptFile.write("%s -d 0 --resetclocks\n" % globalParameters["ROCmSMIPath"])
-        runScriptFile.write("%s -d 0 --setfan 50\n" % globalParameters["ROCmSMIPath"])
+      if globalParameters["PinClocks"] and globalParameters["AMDSMIPath"]:
+        # Reset clocks/overdrive to default and return fans to automatic
+        # (driver) control once the benchmark is done.
+        runScriptFile.write("sudo %s reset -g 0 --clocks --fans\n" % globalParameters["AMDSMIPath"])
   else:
     mxScaleFormatFlag = " --mx-scale-format {}".format(globalParameters["MXScaleFormat"]) if globalParameters["MXScaleFormat"] else ""
     for configFile in configPaths:

projects/hipblaslt/tensilelite/Tensile/Common/GlobalParameters.py

@@ -58,7 +58,7 @@
 # timing between GSU / non-GSU kernels
 globalParameters["PinClocks"] = False  # T=pin gpu clocks and fan, F=don't
 globalParameters["HardwareMonitor"] = (
-    True  # False: disable benchmarking client monitoring clocks using rocm-smi.
+    True  # False: disable benchmarking client monitoring clocks using amd-smi.
 )
 globalParameters["MinFlopsPerSync"] = (
     1  # Minimum number of flops per sync to increase stability for small problems
@@ -272,7 +272,7 @@
 )
 
 # internal, i.e., gets set during startup
-globalParameters["ROCmSMIPath"] = None  # /opt/rocm/bin/rocm-smi
+globalParameters["AMDSMIPath"] = None  # /usr/bin/amd-smi
 globalParameters["HipClangVersion"] = "0.0.0"
 
 # default runtime is selected based on operating system, user can override
@@ -718,7 +718,7 @@ def capRow(isaInfoMap, cap, capType):
 # e.g. RocProfCounter: 42 to pass silently.
 globalParameterTypeOverrides = {
     "ClientExecutionLockPath": {type(None), str},   # path or unset
-    "ROCmSMIPath":             {type(None), str},   # path, populated at startup
+    "AMDSMIPath":              {type(None), str},   # path, populated at startup
     "CmakeCxxCompiler":        {type(None), str},   # path, populated at startup
     "RocProfCounter":          {type(None), str},   # counter spec or None
 }
@@ -888,13 +888,19 @@ def assignGlobalParameters(config, isaInfoMap: Dict[IsaVersion, IsaInfo]):
 
     globalParameters["ROCmBinPath"] = os.path.join(globalParameters["ROCmPath"], "bin")
     try:
-        globalParameters["ROCmSMIPath"] = locateExe(globalParameters["ROCmBinPath"], "rocm-smi")
+        globalParameters["AMDSMIPath"] = locateExe(globalParameters["ROCmBinPath"], "amd-smi")
     except OSError:
-        if os.name == "nt":
-            # rocm-smi is not presently supported on Windows so do not require it.
-            pass
-        else:
-            raise
+        # amd-smi is only needed at runtime to pin clocks/fans during benchmarking
+        # and tuning; it is not required to build libraries or validate logic.
+        # It is also not presently supported on Windows. Treat a missing amd-smi as
+        # non-fatal: leave AMDSMIPath unset (None) so that clock pinning is skipped,
+        # rather than aborting the build in environments that do not ship amd-smi.
+        globalParameters["AMDSMIPath"] = None
+        if os.name != "nt":
+            printWarning(
+                "Could not locate amd-smi; GPU clock/fan pinning will be disabled. "
+                "Install the amdsmi package to enable it."
+            )
 
     if "AsanBuild" in config:
         globalParameters["AsanBuild"] = config["AsanBuild"]
@@ -951,13 +957,18 @@ def setupRestoreClocks():
     import atexit
 
     def restoreClocks():
-        # Clocks will only be pinned if rocm-smi is available, therefore
+        # Clocks will only be pinned if amd-smi is available, therefore
         # we only need to restore if found.
         if globalParameters["PinClocks"]:
-            rsmi = globalParameters["ROCmSMIPath"]
-            if rsmi is not None:
-                subprocess.call([rsmi, "-d", "0", "--resetclocks"])
-                subprocess.call([rsmi, "-d", "0", "--setfan", "50"])
+            asmi = globalParameters["AMDSMIPath"]
+            if asmi is not None:
+                # amd-smi set/reset require elevated privileges.
+                # Reset clocks/overdrive to default and return fans to
+                # automatic (driver) control.
+                cmd = [asmi, "reset", "-g", "0", "--clocks", "--fans"]
+                if hasattr(os, "geteuid") and os.geteuid() != 0:
+                    cmd = ["sudo", "-n"] + cmd
+                subprocess.call(cmd, stdout=subprocess.DEVNULL, stderr=subprocess.DEVNULL)
 
     atexit.register(restoreClocks)
 

projects/hipblaslt/tensilelite/Tensile/Common/Utilities.py

@@ -341,7 +341,7 @@ def isRhel8() -> bool:
         content = f.read()
     match = re.search(pattern, content, re.DOTALL)
     if match:
-        printWarning("Rhel8 environments may not support all tools for system queries such as rocm-smi.")
+        printWarning("Rhel8 environments may not support all tools for system queries such as amd-smi.")
         return True
     return False
 

projects/hipblaslt/tensilelite/Tensile/ParallelExecution.py

@@ -86,7 +86,7 @@
     results.csv (all results, same format as single-GPU)
 
 Process Flow:
-1. Detect available GPUs via rocm-smi or hipInfo
+1. Detect available GPUs via amd-smi or hipInfo
 2. Count problems in the config file
 3. Create per-GPU config files with:
    - Assigned device-idx (GPU index)
@@ -107,6 +107,7 @@
 - runClientParallel(): Main orchestration function
 """
 
+import json
 import os
 import re
 import shutil
@@ -121,19 +122,19 @@
 
 
 def detectAvailableGpus():
-    """Detect the number of available GPUs using rocm-smi."""
+    """Detect the number of available GPUs using amd-smi."""
     try:
         result = subprocess.run(
-            ["rocm-smi", "--showid"],
+            ["amd-smi", "list", "--json"],
             capture_output=True,
             text=True,
             timeout=10
         )
         if result.returncode == 0:
-            # Count unique GPU indices (GPU[N] appears multiple times per device)
-            gpu_indices = set(re.findall(r'GPU\[(\d+)\]', result.stdout))
-            if gpu_indices:
-                return len(gpu_indices)
+            # amd-smi list --json returns one object per GPU
+            gpus = json.loads(result.stdout)
+            if gpus:
+                return len(gpus)
     except Exception:
         pass
 

projects/hipblaslt/tensilelite/Tensile/Tensile.py

@@ -31,6 +31,7 @@
 import sys
 import argparse
 import glob
+import json
 
 from datetime import datetime
 from pathlib import Path
@@ -274,37 +275,25 @@ def get_gpu_max_frequency_smi(device_id):
     Get the maximum frequency of the specified GPU device
     '''
     try:
-        # Run rocm-smi command and capture output
-        result = subprocess.run(['rocm-smi', '-s'], capture_output=True, text=True)
+        # Run amd-smi command and capture the GFX clock info as JSON
+        result = subprocess.run(
+            ['amd-smi', 'metric', '-g', str(device_id), '--clock', '--json'],
+            capture_output=True, text=True)
 
         if result.returncode != 0:
-           print(f"Error running rocm-smi: {result.stderr}")
+           print(f"Error running amd-smi: {result.stderr}")
            return None
 
-        # Parse the output
-        lines = result.stdout.split('\n')
-        sclk_section = False
-        frequencies = []
-
-        # Look for the sclk section of the specified device
-        for line in lines:
-            line = line.split(" ")
-            if 'sclk' in line and f"GPU{device_id}" in line:
-                sclk_section = True
-                continue
+        data = json.loads(result.stdout)
+        clocks = data['gpu_data'][0]['clock']
 
-           # Parse frequencies in the sclk section
-            if sclk_section:
-                for part in line:
-                    if part.endswith("Mhz"):
-                        try:
-                            frequency = part.replace("Mhz", "")
-                            frequencies.append(int(frequency))
-                        except ValueError:
-                            print(f"Error parsing frequency: {part}")
-                        break
-                if "socclk" in line:
-                    break
+        # Collect the max GFX (sclk) clock across all gfx engines/partitions
+        frequencies = []
+        for name, info in clocks.items():
+            if name.startswith('gfx'):
+                max_clk = info.get('max_clk', {}).get('value')
+                if isinstance(max_clk, int):
+                    frequencies.append(max_clk)
 
         # Return the maximum frequency found
         return max(frequencies) if frequencies else None
@@ -625,7 +614,7 @@ def Tensile(userArgs):
         max_frequency = get_gpu_max_frequency(device_id)
 
         if not max_frequency or max_frequency <= 0:
-            max_frequency = get_gpu_max_frequency_smi(device_id) # Using rocm-smi just in case
+            max_frequency = get_gpu_max_frequency_smi(device_id) # Using amd-smi just in case
 
         if not max_frequency or max_frequency <= 0:
             print(f"Could not detect valid GPU frequency for device {device_id}")

projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/CommonUtilities/test_mut_Utilities_isRhel8_char.py

@@ -96,5 +96,5 @@ def test_isRhel8_emits_exact_warning_text_on_match(monkeypatch):
 
     assert result is True
     assert calls == [
-        ("Rhel8 environments may not support all tools for system queries such as rocm-smi.",)
+        ("Rhel8 environments may not support all tools for system queries such as amd-smi.",)
     ]