From f5c037a7350a8a778e412ced3e53fec8dccb60ce Mon Sep 17 00:00:00 2001 From: resonantsky Date: Thu, 2 Apr 2026 14:59:00 +0200 Subject: [PATCH 01/12] Added further rocblas support enhancements and performance-related best practice settings. --- scripts/rocm/rocm_mgr.py | 44 +++++++++++++- scripts/rocm/rocm_vars.py | 119 +++++++++++++++++++++++++++++++++++--- scripts/rocm_ext.py | 85 +++++++++++++++++++++------ 3 files changed, 221 insertions(+), 27 deletions(-) diff --git a/scripts/rocm/rocm_mgr.py b/scripts/rocm/rocm_mgr.py index b704fff03..9bbc0199f 100644 --- a/scripts/rocm/rocm_mgr.py +++ b/scripts/rocm/rocm_mgr.py @@ -1,4 +1,5 @@ import os +import re import sys from pathlib import Path from typing import Dict, Optional @@ -121,14 +122,22 @@ def _get_venv() -> str: return os.environ.get("VIRTUAL_ENV", "") or sys.prefix +def _get_root() -> str: + """App root — one level above the venv folder (e.g. E:\\Sd.Next).""" + return str(Path(_get_venv()).parent) + + def _expand_venv(value: str) -> str: - return value.replace("{VIRTUAL_ENV}", _get_venv()) + return value.replace("{VIRTUAL_ENV}", _get_venv()).replace("{ROOT}", _get_root()) def _collapse_venv(value: str) -> str: venv = _get_venv() + root = _get_root() if venv and value.startswith(venv): return "{VIRTUAL_ENV}" + value[len(venv):] + if root and value.startswith(root): + return "{ROOT}" + value[len(root):] return value @@ -365,6 +374,28 @@ def _user_db_summary(path: Path) -> dict: return out +def _extract_db_hash(db_path: Path) -> str: + """Derive the cache subfolder name from udb.txt filenames. + e.g. gfx1030_30.HIP.3_5_1_5454e9e2da.udb.txt → '3.5.1.5454e9e2da'""" + for f in db_path.glob("*.HIP.*.udb.txt"): + m = re.search(r'\.HIP\.([^.]+)\.udb\.txt$', f.name) + if m: + return m.group(1).replace("_", ".") + return "" + + +def _user_cache_summary(path: Path) -> dict: + """Return {filename: 'N KB'} for binary cache blobs in the resolved cache path.""" + out = {} + if not path.exists(): + return out + for f in sorted(path.iterdir()): + if f.is_file(): + kb = f.stat().st_size // 1024 + out[f.name] = f"{kb} KB" + return out + + def info() -> dict: config = load_config() db_path = Path(_expand_venv(config.get("MIOPEN_SYSTEM_DB_PATH", ""))) @@ -427,12 +458,23 @@ def info() -> dict: if ufiles: udb["files"] = ufiles + # --- User cache (~/.miopen/cache/) --- + cache_base = Path.home() / ".miopen" / "cache" + db_hash = _extract_db_hash(user_db_path) if user_db_path.exists() else "" + cache_path = cache_base / db_hash if db_hash else cache_base + ucache = {"path": str(cache_path), "exists": cache_path.exists()} + if cache_path.exists(): + cfiles = _user_cache_summary(cache_path) + if cfiles: + ucache["files"] = cfiles + return { "rocm": rocm_section, "torch": torch_section, "gpu": gpu_section, "system_db": sdb, "user_db": udb, + "user_cache": ucache, } diff --git a/scripts/rocm/rocm_vars.py b/scripts/rocm/rocm_vars.py index ac1e720c6..235aef0d0 100644 --- a/scripts/rocm/rocm_vars.py +++ b/scripts/rocm/rocm_vars.py @@ -3,6 +3,7 @@ from typing import Dict, Any, List, Tuple # --- General MIOpen/rocBLAS variables (dropdown/textbox/checkbox) --- GENERAL_VARS: Dict[str, Dict[str, Any]] = { + # ── GEMM backend selector + companion toggles ────────────────────────── "MIOPEN_GEMM_ENFORCE_BACKEND": { "default": "1", "desc": "Enforce GEMM backend", @@ -10,6 +11,29 @@ GENERAL_VARS: Dict[str, Dict[str, Any]] = { "options": [("1 - rocBLAS", "1"), ("5 - hipBLASLt", "5")], "restart_required": False, }, + "PYTORCH_ROCM_USE_ROCBLAS": { + "default": "1", + "desc": "PyTorch ROCm: prioritise rocBLAS for linear algebra", + "widget": "dropdown", + "options": [("0 - Off", "0"), ("1 - On", "1")], + "restart_required": True, + }, + "PYTORCH_HIPBLASLT_DISABLE": { + "default": "1", + "desc": "Disable PyTorch hipBLASLt dispatcher", + "widget": "dropdown", + "options": [("0 - Allow hipBLASLt", "0"), ("1 - Disable hipBLASLt", "1")], + "restart_required": True, + }, + "ROCBLAS_USE_HIPBLASLT": { + "default": "0", + "desc": "rocBLAS: use hipBLASLt backend (0 = Tensile)", + "widget": "dropdown", + "options": [("0 - Tensile (rocBLAS)", "0"), ("1 - hipBLASLt", "1")], + "restart_required": True, + }, + + # ── MIOpen behavioural settings ──────────────────────────────────────── "MIOPEN_FIND_MODE": { "default": "2", "desc": "MIOpen Find Mode", @@ -31,6 +55,15 @@ GENERAL_VARS: Dict[str, Dict[str, Any]] = { "options": [("0 - Off", "0"), ("1 - On", "1")], "restart_required": True, }, + "MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC": { + "default": "0", + "desc": "Deterministic convolution (reproducible results, may be slower)", + "widget": "dropdown", + "options": [("0 - Off", "0"), ("1 - On", "1")], + "restart_required": False, + }, + + # ── Paths / sizes ────────────────────────────────────────────────────── "MIOPEN_SYSTEM_DB_PATH": { "default": "{VIRTUAL_ENV}\\Lib\\site-packages\\_rocm_sdk_devel\\bin\\", "desc": "MIOpen system DB path", @@ -38,6 +71,75 @@ GENERAL_VARS: Dict[str, Dict[str, Any]] = { "options": None, "restart_required": True, }, + "MIOPEN_CONVOLUTION_MAX_WORKSPACE": { + "default": "1073741824", + "desc": "MIOpen convolution max workspace (bytes; 1 GB default)", + "widget": "textbox", + "options": None, + "restart_required": False, + }, + "ROCBLAS_TENSILE_LIBPATH": { + "default": "{VIRTUAL_ENV}\\Lib\\site-packages\\_rocm_sdk_devel\\bin\\rocblas\\library", + "desc": "rocBLAS Tensile library path", + "widget": "textbox", + "options": None, + "restart_required": True, + }, + "ROCBLAS_DEVICE_MEMORY_SIZE": { + "default": "", + "desc": "rocBLAS workspace size in bytes (empty = dynamic)", + "widget": "textbox", + "options": None, + "restart_required": False, + }, + "PYTORCH_TUNABLEOP_CACHE_DIR": { + "default": "{ROOT}\\models\\tunable", + "desc": "TunableOp: kernel profile cache directory", + "widget": "textbox", + "options": None, + "restart_required": False, + }, + + # ── rocBLAS settings ─────────────────────────────────────────────────── + "ROCBLAS_STREAM_ORDER_ALLOC": { + "default": "1", + "desc": "rocBLAS stream-ordered memory allocation", + "widget": "dropdown", + "options": [("0 - Standard", "0"), ("1 - Stream-ordered", "1")], + "restart_required": False, + }, + "ROCBLAS_DEFAULT_ATOMICS_MODE": { + "default": "1", + "desc": "rocBLAS default atomics mode (1 = allow non-deterministic for performance)", + "widget": "dropdown", + "options": [("0 - Off (deterministic)", "0"), ("1 - On (performance)", "1")], + "restart_required": False, + }, + "PYTORCH_TUNABLEOP_ROCBLAS_ENABLED": { + "default": "1", + "desc": "TunableOp: wrap and optimise rocBLAS GEMM calls", + "widget": "dropdown", + "options": [("0 - Off", "0"), ("1 - On", "1")], + "restart_required": False, + }, + "PYTORCH_TUNABLEOP_TUNING": { + "default": "0", + "desc": "TunableOp: tuning mode (1 = benchmark; 0 = use saved CSV)", + "widget": "dropdown", + "options": [("0 - Use saved CSV", "0"), ("1 - Benchmark new shapes", "1")], + "restart_required": False, + }, + + # ── hipBLASLt settings ───────────────────────────────────────────────── + "PYTORCH_TUNABLEOP_HIPBLASLT_ENABLED": { + "default": "0", + "desc": "TunableOp: benchmark hipBLASLt kernels", + "widget": "dropdown", + "options": [("0 - Off", "0"), ("1 - On", "1")], + "restart_required": False, + }, + + # ── Logging: MIOpen → rocBLAS → hipBLASLt ───────────────────────────── "MIOPEN_LOG_LEVEL": { "default": "0", "desc": "MIOpen log verbosity level", @@ -66,13 +168,6 @@ GENERAL_VARS: Dict[str, Dict[str, Any]] = { "options": [("0 - Off", "0"), ("1 - Error", "1"), ("2 - Trace", "2"), ("3 - Hints", "3"), ("4 - Info", "4"), ("5 - API Trace", "5")], "restart_required": False, }, - "MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC": { - "default": "0", - "desc": "Deterministic convolution (reproducible results, may be slower)", - "widget": "dropdown", - "options": [("0 - Off", "0"), ("1 - On", "1")], - "restart_required": False, - }, } # --- Solver toggles (inference/FWD only, RDNA2/3/4 compatible) --- @@ -251,3 +346,13 @@ SOLVER_GROUPS: List[Tuple[str, List[str]]] = [ "MIOPEN_DEBUG_CK_DEFAULT_KERNELS", ]), ] + +# Variables that are relevant only when hipBLASLt is the active GEMM backend. +# These are visually greyed-out in the UI when rocBLAS (MIOPEN_GEMM_ENFORCE_BACKEND="1") is selected. +HIPBLASLT_VARS: set = { + "PYTORCH_HIPBLASLT_DISABLE", + "ROCBLAS_USE_HIPBLASLT", + "PYTORCH_TUNABLEOP_HIPBLASLT_ENABLED", + "HIPBLASLT_LOG_LEVEL", +} + diff --git a/scripts/rocm_ext.py b/scripts/rocm_ext.py index 4d89f7e42..e8ddf24ed 100644 --- a/scripts/rocm_ext.py +++ b/scripts/rocm_ext.py @@ -19,7 +19,7 @@ class ROCmScript(scripts_manager.Script): if not shared.cmd_opts.use_rocm and not installer.torch_info.get('type') == 'rocm': # skip ui creation if not rocm return [] - from scripts.rocm import rocm_mgr, rocm_vars # pylint: disable=no-name-in-module + from scripts.rocm import rocm_mgr, rocm_vars, rocm_profiles # pylint: disable=no-name-in-module config = rocm_mgr.load_config() var_names = [] @@ -59,11 +59,25 @@ class ROCmScript(scripts_manager.Script): row("path", udb.get("path", "")) for fname, finfo in udb.get("files", {}).items(): row(fname, finfo) + section("User cache (~/.miopen/cache)") + ucache = d.get("user_cache", {}) + row("path", ucache.get("path", "")) + for fname, sz in ucache.get("files", {}).items(): + row(fname, sz) return f"{''.join(rows)}
" + def _build_style(unavailable, hipblaslt_disabled=False): + rules = [] + for v in (unavailable or []): + rules.append(f"#rocm_var_{v.lower()} label {{ text-decoration: line-through; opacity: 0.5; }}") + if hipblaslt_disabled: + for v in rocm_vars.HIPBLASLT_VARS: + rules.append(f"#rocm_var_{v.lower()} {{ opacity: 0.45; pointer-events: none; }}") + return f"" if rules else "" + with gr.Accordion('ROCm: Advanced Config', open=False, elem_id='rocm_config'): with gr.Row(): - gr.HTML("

Advanced configuration for ROCm users.


Set your database and solver selections based on GPU profile or individually.


Enable cuDNN in Backend Settings to activate MIOpen.

") + gr.HTML("

Advanced configuration for ROCm users.


For best performance ensure that cudnn and torch tunable ops are set to default in Backend Settings.

") with gr.Row(): btn_info = gr.Button("Refresh Info", variant="primary", elem_id="rocm_btn_info", size="sm") btn_apply = gr.Button("Apply", variant="primary", elem_id="rocm_btn_apply", size="sm") @@ -74,7 +88,10 @@ class ROCmScript(scripts_manager.Script): btn_rdna2 = gr.Button("RDNA2 (RX 6000)", elem_id="rocm_btn_rdna2") btn_rdna3 = gr.Button("RDNA3 (RX 7000)", elem_id="rocm_btn_rdna3") btn_rdna4 = gr.Button("RDNA4 (RX 9000)", elem_id="rocm_btn_rdna4") - style_out = gr.HTML("") + _init_gemm = config.get("MIOPEN_GEMM_ENFORCE_BACKEND", "1") + _init_arch = config.get(rocm_mgr._ARCH_KEY, "") + _init_unavailable = rocm_profiles.UNAVAILABLE.get(_init_arch, set()) if _init_arch else set() + style_out = gr.HTML(_build_style(_init_unavailable, _init_gemm == "1")) info_out = gr.HTML(value=_info_html, elem_id="rocm_info_table") # General vars (dropdowns, textboxes, checkboxes) @@ -106,13 +123,46 @@ class ROCmScript(scripts_manager.Script): for name, comp in zip(var_names, components): meta = rocm_vars.ROCM_ENV_VARS[name] - if meta["widget"] == "dropdown": + if meta["widget"] == "dropdown" and name != "MIOPEN_GEMM_ENFORCE_BACKEND": comp.change(fn=lambda v, n=name: _autosave_field(n, v), inputs=[comp], outputs=[], show_progress='hidden') + _GEMM_COMPANIONS = { + "PYTORCH_ROCM_USE_ROCBLAS": {"1": "1", "5": "0"}, + "PYTORCH_HIPBLASLT_DISABLE": {"1": "1", "5": "0"}, + "ROCBLAS_USE_HIPBLASLT": {"1": "0", "5": "1"}, + "PYTORCH_TUNABLEOP_HIPBLASLT_ENABLED": {"1": "0", "5": "1"}, + } + + def gemm_changed(gemm_display_val): + stored = rocm_mgr._dropdown_stored(str(gemm_display_val), rocm_vars.ROCM_ENV_VARS["MIOPEN_GEMM_ENFORCE_BACKEND"]["options"]) + cfg = rocm_mgr.load_config().copy() + cfg["MIOPEN_GEMM_ENFORCE_BACKEND"] = stored + for var, vals in _GEMM_COMPANIONS.items(): + cfg[var] = vals.get(stored, cfg.get(var, "")) + rocm_mgr.save_config(cfg) + rocm_mgr.apply_env(cfg) + arch = cfg.get(rocm_mgr._ARCH_KEY, "") + unavailable = rocm_profiles.UNAVAILABLE.get(arch, set()) + result = [gr.update(value=_build_style(unavailable, stored == "1"))] + for pname in var_names: + if pname in _GEMM_COMPANIONS: + meta = rocm_vars.ROCM_ENV_VARS[pname] + val = _GEMM_COMPANIONS[pname].get(stored, cfg.get(pname, "")) + result.append(gr.update(value=rocm_mgr._dropdown_display(val, meta["options"]))) + else: + result.append(gr.update()) + return result + + gemm_comp = components[var_names.index("MIOPEN_GEMM_ENFORCE_BACKEND")] + gemm_comp.change(fn=gemm_changed, inputs=[gemm_comp], outputs=[style_out] + components, show_progress='hidden') + def apply_fn(*values): rocm_mgr.apply_all(var_names, list(values)) saved = rocm_mgr.load_config() - result = [gr.update(value="")] + arch = saved.get(rocm_mgr._ARCH_KEY, "") + unavailable = rocm_profiles.UNAVAILABLE.get(arch, set()) + gemm_val = saved.get("MIOPEN_GEMM_ENFORCE_BACKEND", "1") + result = [gr.update(value=_build_style(unavailable, gemm_val == "1"))] for name in var_names: meta = rocm_vars.ROCM_ENV_VARS[name] val = saved.get(name, meta["default"]) @@ -124,19 +174,13 @@ class ROCmScript(scripts_manager.Script): result.append(gr.update(value=rocm_mgr._expand_venv(val))) return result - def _build_style(unavailable): - if not unavailable: - return "" - rules = " ".join( - f"#rocm_var_{v.lower()} label {{ text-decoration: line-through; opacity: 0.5; }}" - for v in unavailable - ) - return f"" - def reset_fn(): rocm_mgr.reset_defaults() updated = rocm_mgr.load_config() - result = [gr.update(value="")] + arch = updated.get(rocm_mgr._ARCH_KEY, "") + unavailable = rocm_profiles.UNAVAILABLE.get(arch, set()) + gemm_val = updated.get("MIOPEN_GEMM_ENFORCE_BACKEND", "1") + result = [gr.update(value=_build_style(unavailable, gemm_val == "1"))] for name in var_names: meta = rocm_vars.ROCM_ENV_VARS[name] val = updated.get(name, meta["default"]) @@ -150,7 +194,9 @@ class ROCmScript(scripts_manager.Script): def clear_fn(): rocm_mgr.clear_env() - result = [gr.update(value="")] + cfg = rocm_mgr.load_config() + gemm_val = cfg.get("MIOPEN_GEMM_ENFORCE_BACKEND", "1") + result = [gr.update(value=_build_style(None, gemm_val == "1"))] for name in var_names: meta = rocm_vars.ROCM_ENV_VARS[name] if meta["widget"] == "checkbox": @@ -163,7 +209,8 @@ class ROCmScript(scripts_manager.Script): def delete_fn(): rocm_mgr.delete_config() - result = [gr.update(value="")] + gemm_default = rocm_vars.ROCM_ENV_VARS.get("MIOPEN_GEMM_ENFORCE_BACKEND", {}).get("default", "1") + result = [gr.update(value=_build_style(None, gemm_default == "1"))] for name in var_names: meta = rocm_vars.ROCM_ENV_VARS[name] if meta["widget"] == "checkbox": @@ -175,11 +222,11 @@ class ROCmScript(scripts_manager.Script): return result def profile_fn(arch): - from scripts.rocm import rocm_profiles # pylint: disable=no-name-in-module rocm_mgr.apply_profile(arch) updated = rocm_mgr.load_config() unavailable = rocm_profiles.UNAVAILABLE.get(arch, set()) - result = [gr.update(value=_build_style(unavailable))] + gemm_val = updated.get("MIOPEN_GEMM_ENFORCE_BACKEND", "1") + result = [gr.update(value=_build_style(unavailable, gemm_val == "1"))] for pname in var_names: meta = rocm_vars.ROCM_ENV_VARS[pname] val = updated.get(pname, meta["default"]) From 1ed2811c8025087adb08850de0d247d6ef4c6340 Mon Sep 17 00:00:00 2001 From: resonantsky Date: Thu, 2 Apr 2026 15:21:54 +0200 Subject: [PATCH 02/12] rocm_mgr: use modules.paths.script_path for app root --- scripts/rocm/rocm_mgr.py | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/scripts/rocm/rocm_mgr.py b/scripts/rocm/rocm_mgr.py index 9bbc0199f..27b3b36fc 100644 --- a/scripts/rocm/rocm_mgr.py +++ b/scripts/rocm/rocm_mgr.py @@ -123,8 +123,12 @@ def _get_venv() -> str: def _get_root() -> str: - """App root — one level above the venv folder (e.g. E:\\Sd.Next).""" - return str(Path(_get_venv()).parent) + """App root — from modules.paths.script_path.""" + try: + from modules.paths import script_path # pylint: disable=import-outside-toplevel + return str(script_path) + except Exception: + return str(Path(_get_venv()).parent) def _expand_venv(value: str) -> str: From fdc2f464579a9d036b1f87f21abc86e8e097ed26 Mon Sep 17 00:00:00 2001 From: resonantsky Date: Thu, 2 Apr 2026 16:20:30 +0200 Subject: [PATCH 03/12] further rocblas default settings --- scripts/rocm/rocm_vars.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/scripts/rocm/rocm_vars.py b/scripts/rocm/rocm_vars.py index 235aef0d0..0e2235015 100644 --- a/scripts/rocm/rocm_vars.py +++ b/scripts/rocm/rocm_vars.py @@ -12,7 +12,7 @@ GENERAL_VARS: Dict[str, Dict[str, Any]] = { "restart_required": False, }, "PYTORCH_ROCM_USE_ROCBLAS": { - "default": "1", + "default": "0", "desc": "PyTorch ROCm: prioritise rocBLAS for linear algebra", "widget": "dropdown", "options": [("0 - Off", "0"), ("1 - On", "1")], @@ -116,7 +116,7 @@ GENERAL_VARS: Dict[str, Dict[str, Any]] = { "restart_required": False, }, "PYTORCH_TUNABLEOP_ROCBLAS_ENABLED": { - "default": "1", + "default": "0", "desc": "TunableOp: wrap and optimise rocBLAS GEMM calls", "widget": "dropdown", "options": [("0 - Off", "0"), ("1 - On", "1")], From b659a06c60f3f5ad43f21bb3e02f9694e00ae42b Mon Sep 17 00:00:00 2001 From: resonantsky Date: Thu, 2 Apr 2026 21:13:22 +0200 Subject: [PATCH 04/12] rocm_mgr: easy fixes - simplify _get_venv to sys.prefix, _get_root to direct script_path import --- scripts/rocm/rocm_mgr.py | 10 +++------- 1 file changed, 3 insertions(+), 7 deletions(-) diff --git a/scripts/rocm/rocm_mgr.py b/scripts/rocm/rocm_mgr.py index 27b3b36fc..ccfb90de9 100644 --- a/scripts/rocm/rocm_mgr.py +++ b/scripts/rocm/rocm_mgr.py @@ -119,16 +119,12 @@ def _resolve_dtype() -> str: # --- venv helpers --- def _get_venv() -> str: - return os.environ.get("VIRTUAL_ENV", "") or sys.prefix + return sys.prefix def _get_root() -> str: - """App root — from modules.paths.script_path.""" - try: - from modules.paths import script_path # pylint: disable=import-outside-toplevel - return str(script_path) - except Exception: - return str(Path(_get_venv()).parent) + from modules.paths import script_path # pylint: disable=import-outside-toplevel + return str(script_path) def _expand_venv(value: str) -> str: From 01d53edb25b73cc2e761f96a31fc2b64ac72e1c3 Mon Sep 17 00:00:00 2001 From: resonantsky Date: Fri, 3 Apr 2026 10:12:05 +0200 Subject: [PATCH 05/12] code quality and layout fixes --- scripts/rocm/rocm_mgr.py | 42 +++++++------- scripts/rocm/rocm_profiles.py | 82 +++++++++++++------------- scripts/rocm/rocm_vars.py | 105 +++++++++++++++++----------------- 3 files changed, 114 insertions(+), 115 deletions(-) diff --git a/scripts/rocm/rocm_mgr.py b/scripts/rocm/rocm_mgr.py index ccfb90de9..ea987a421 100644 --- a/scripts/rocm/rocm_mgr.py +++ b/scripts/rocm/rocm_mgr.py @@ -29,17 +29,17 @@ CONFIG = Path(os.path.abspath(os.path.join('data', 'rocm.json'))) _cache: Optional[Dict[str, str]] = None # loaded once, invalidated on save # Metadata key written into rocm.json to record which architecture profile is active. -# Not an environment variable — always skipped during env application but preserved in the +# Not an environment variable - always skipped during env application but preserved in the # saved config so that arch-safety enforcement is consistent across restarts. _ARCH_KEY = "_rocm_arch" # Vars that must never appear in the process environment. # -# _DTYPE_UNSAFE: alter FP16 inference dtype — must be cleared regardless of config -# MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL — DEBUG alias: routes all FP16 convs through BF16 exponent math -# MIOPEN_CONVOLUTION_ATTRIB_FP16_ALT_IMPL — API-level alias: same BF16-exponent effect -# MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_EXPEREMENTAL_FP16_TRANSFORM — unstable experimental FP16 path -# MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_PK_ATOMIC_ADD_FP16 — changes FP16 WrW atomic accumulation +# _DTYPE_UNSAFE: alter FP16 inference dtype - must be cleared regardless of config +# MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL - DEBUG alias: routes all FP16 convs through BF16 exponent math +# MIOPEN_CONVOLUTION_ATTRIB_FP16_ALT_IMPL - API-level alias: same BF16-exponent effect +# MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_EXPEREMENTAL_FP16_TRANSFORM - unstable experimental FP16 path +# MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_PK_ATOMIC_ADD_FP16 - changes FP16 WrW atomic accumulation # # SOLVER_DISABLED_BY_DEFAULT: every solver known to be incompatible with this runtime # (FP32-only, training-only WrW/BWD, fixed-geometry mismatches, XDLOPS/CDNA-only, arch-specific). @@ -54,18 +54,18 @@ _DTYPE_UNSAFE = { # regardless of saved config. Limited to dtype-corrupting vars only. # IMPORTANT: SOLVER_DISABLED_BY_DEFAULT is intentionally NOT included here. # When a solver var is absent (unset) MIOpen still calls IsApplicable() on every -# conv-find — wasted probing overhead. When a var is explicitly "0" MIOpen skips +# conv-find - wasted probing overhead. When a var is explicitly "0" MIOpen skips # IsApplicable() immediately. Solver defaults flow through the config loop as "0" # (their ROCM_ENV_VARS default is "0") so they are explicitly set to "0" in the env. _UNSET_VARS = _DTYPE_UNSAFE # Additional environment vars that must be removed from the process before MIOpen loads. # These are not MIOpen solver toggles but can corrupt MIOpen's runtime behaviour: -# HIP_PATH / HIP_PATH_71 — point to the system AMD ROCm install; override the venv-bundled +# HIP_PATH / HIP_PATH_71 - point to the system AMD ROCm install; override the venv-bundled # _rocm_sdk_devel DLLs with a potentially mismatched system version -# QML_*/QT_* — QtQuick shader/disk-cache flags leaked from Qt tools; harmless for +# QML_*/QT_* - QtQuick shader/disk-cache flags leaked from Qt tools; harmless for # PyTorch but can conflict with Gradio's embedded Qt helpers -# PYENV_VIRTUALENV_DISABLE_PROMPT — pyenv noise that confuses venv detection +# PYENV_VIRTUALENV_DISABLE_PROMPT - pyenv noise that confuses venv detection _EXTRA_CLEAR_VARS = { "HIP_PATH", "HIP_PATH_71", @@ -73,7 +73,7 @@ _EXTRA_CLEAR_VARS = { "QML_DISABLE_DISK_CACHE", "QML_FORCE_DISK_CACHE", "QT_DISABLE_SHADER_DISK_CACHE", - # PERF_VALS vars are NOT boolean toggles — MIOpen reads them as perf-config strings. + # PERF_VALS vars are NOT boolean toggles - MIOpen reads them as perf-config strings. # If inherited from a parent shell with value "1", MIOpen's GetPerfConfFromEnv parses # "1" as a degenerate config and can return dtype=float32 output from FP16 tensors. "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_PERF_VALS", @@ -82,12 +82,12 @@ _EXTRA_CLEAR_VARS = { # Solvers whose MIOpen IsApplicable() explicitly rejects non-FP32 tensors. # They are safe to leave enabled in FP32 mode. When the active dtype is FP16 or BF16 -# we force them OFF so MIOpen skips the IsApplicable probe entirely — avoids overhead on +# we force them OFF so MIOpen skips the IsApplicable probe entirely - avoids overhead on # every conv shape find. These are NOT in _UNSET_VARS because they are valid in FP32. _FP32_ONLY_SOLVERS = { - "MIOPEN_DEBUG_CONV_FFT", # FFT convolution — FP32 only (MIOpen source: IsFp32 check) - "MIOPEN_DEBUG_AMD_WINOGRAD_3X3", # Winograd 3x3 — FP32 only - "MIOPEN_DEBUG_AMD_FUSED_WINOGRAD", # Fused Winograd — FP32 only + "MIOPEN_DEBUG_CONV_FFT", # FFT convolution - FP32 only (MIOpen source: IsFp32 check) + "MIOPEN_DEBUG_AMD_WINOGRAD_3X3", # Winograd 3x3 - FP32 only + "MIOPEN_DEBUG_AMD_FUSED_WINOGRAD", # Fused Winograd - FP32 only } @@ -172,7 +172,7 @@ def load_config() -> Dict[str, str]: _cache = data if data else {k: v["default"] for k, v in ROCM_ENV_VARS.items()} # Purge unsafe vars from a stale saved config and re-persist only if the file existed. # When running without a saved config (first run / after Delete), load_config() must - # never create the file — that only happens via save_config() on Apply or Apply Profile. + # never create the file - that only happens via save_config() on Apply or Apply Profile. dirty = {k for k in _cache if k in _UNSET_VARS or (k != _ARCH_KEY and k not in ROCM_ENV_VARS)} if dirty: _cache = {k: v for k, v in _cache.items() if k not in dirty} @@ -221,7 +221,7 @@ def apply_env(config: Optional[Dict[str, str]] = None) -> None: os.environ[var] = expanded # Arch safety net: hard-force all hardware-incompatible vars to "0" in the env. # This runs *after* the config loop so it overrides any stale "1" that survived in the JSON. - # Source of truth: rocm_profiles.UNAVAILABLE[arch] — vars with no supporting hardware. + # Source of truth: rocm_profiles.UNAVAILABLE[arch] - vars with no supporting hardware. arch = config.get(_ARCH_KEY, "") unavailable = rocm_profiles.UNAVAILABLE.get(arch, set()) if unavailable: @@ -249,7 +249,7 @@ def apply_all(names: list, values: list) -> None: meta = ROCM_ENV_VARS[name] if meta["widget"] == "checkbox": if value is None: - pass # Gradio passed None (component not interacted with) — leave config unchanged + pass # Gradio passed None (component not interacted with) - leave config unchanged else: config[name] = "1" if value else "0" elif meta["widget"] == "radio": @@ -257,7 +257,7 @@ def apply_all(names: list, values: list) -> None: valid = {v for _, v in meta["options"]} if meta["options"] and isinstance(meta["options"][0], tuple) else set(meta["options"] or []) if stored in valid: config[name] = stored - # else: value was None/invalid — leave the existing saved value untouched + # else: value was None/invalid - leave the existing saved value untouched else: if meta.get("options"): value = _dropdown_stored(str(value), meta["options"]) @@ -300,7 +300,7 @@ def delete_config() -> None: CONFIG.unlink() log.info(f'ROCm delete_config: deleted {CONFIG}') _cache = None - # Delete the MIOpen user DB (~/.miopen/db) — stale entries can cause solver mismatches + # Delete the MIOpen user DB (~/.miopen/db) - stale entries can cause solver mismatches miopen_db = Path(os.path.expanduser('~')) / '.miopen' / 'db' if miopen_db.exists(): shutil.rmtree(miopen_db, ignore_errors=True) @@ -458,7 +458,7 @@ def info() -> dict: if ufiles: udb["files"] = ufiles - # --- User cache (~/.miopen/cache/) --- + # User cache (~/.miopen/cache/) cache_base = Path.home() / ".miopen" / "cache" db_hash = _extract_db_hash(user_db_path) if user_db_path.exists() else "" cache_path = cache_base / db_hash if db_hash else cache_base diff --git a/scripts/rocm/rocm_profiles.py b/scripts/rocm/rocm_profiles.py index eeba0628e..62a401de8 100644 --- a/scripts/rocm/rocm_profiles.py +++ b/scripts/rocm/rocm_profiles.py @@ -1,4 +1,4 @@ -""" +""" Architecture-specific MIOpen solver profiles for AMD GCN/RDNA GPUs. Sources: @@ -6,8 +6,8 @@ Sources: Key axis: consumer RDNA GPUs have NO XDLOPS hardware (that's CDNA/Instinct only). RDNA2 (gfx1030): RX 6000 series - RDNA3 (gfx1100): RX 7000 series — adds Fury Winograd, wider MPASS - RDNA4 (gfx1200): RX 9000 series — adds Rage Winograd, wider MPASS + RDNA3 (gfx1100): RX 7000 series - adds Fury Winograd, wider MPASS + RDNA4 (gfx1200): RX 9000 series - adds Rage Winograd, wider MPASS Each profile is a dict of {var: value} that will be MERGED on top of the current config (general vars like DB path / log level are preserved). @@ -15,9 +15,9 @@ current config (general vars like DB path / log level are preserved). from typing import Dict -# --------------------------------------------------------------------------- + # Shared: everything that must be OFF on ALL consumer RDNA (no XDLOPS hw) -# --------------------------------------------------------------------------- + _XDLOPS_OFF: Dict[str, str] = { # GTC XDLOPS (CDNA-only) "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS": "0", @@ -55,7 +55,7 @@ _XDLOPS_OFF: Dict[str, str] = { # MLIR (CDNA-only in practice) "MIOPEN_DEBUG_CONV_MLIR_IGEMM_WRW_XDLOPS": "0", "MIOPEN_DEBUG_CONV_MLIR_IGEMM_BWD_XDLOPS": "0", - # MP BD Winograd (Multi-pass Block-Decomposed — CDNA / high-end only) + # MP BD Winograd (Multi-pass Block-Decomposed - CDNA / high-end only) "MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F2X3": "0", "MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F3X3": "0", "MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F4X3": "0", @@ -68,17 +68,17 @@ _XDLOPS_OFF: Dict[str, str] = { "MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F6X3": "0", } -# --------------------------------------------------------------------------- -# RDNA2 — gfx1030 (RX 6000 series) + +# RDNA2 - gfx1030 (RX 6000 series) # No XDLOPS, no Fury/Rage Winograd, MPASS limited to F3x2/F3x3 # ASM IGEMM: V4R1 variants only; HIP IGEMM: non-XDLOPS V4R1/R4 only -# --------------------------------------------------------------------------- + RDNA2: Dict[str, str] = { **_XDLOPS_OFF, # General settings (architecture-independent; set here so all profiles cover them) "MIOPEN_SEARCH_CUTOFF": "0", "MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC": "0", - # Core algo enables — FFT is FP32-only but harmless (IsApplicable rejects it for fp16 tensors) + # Core algo enables - FFT is FP32-only but harmless (IsApplicable rejects it for fp16 tensors) "MIOPEN_DEBUG_CONV_FFT": "1", "MIOPEN_DEBUG_CONV_DIRECT": "1", "MIOPEN_DEBUG_CONV_GEMM": "1", @@ -93,16 +93,16 @@ RDNA2: Dict[str, str] = { "MIOPEN_DEBUG_OPENCL_CONVOLUTIONS": "1", "MIOPEN_DEBUG_OPENCL_WAVE64_NOWGP": "1", "MIOPEN_DEBUG_ATTN_SOFTMAX": "1", - # Direct ASM — dtype notes - # 3X3U / 1X1U / 1X1UV2: FP32/FP16 forward — enabled + # Direct ASM - dtype notes + # 3X3U / 1X1U / 1X1UV2: FP32/FP16 forward - enabled "MIOPEN_DEBUG_CONV_DIRECT_ASM_3X3U": "1", "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U": "1", "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1UV2": "1", - # 5X10U2V2: fixed geometry (5*10 stride-2), no SD conv matches — disabled + # 5X10U2V2: fixed geometry (5*10 stride-2), no SD conv matches - disabled "MIOPEN_DEBUG_CONV_DIRECT_ASM_5X10U2V2": "0", - # 7X7C3H224W224: hard-coded ImageNet stem (C=3, H=W=224, K=64) — never matches SD — disabled + # 7X7C3H224W224: hard-coded ImageNet stem (C=3, H=W=224, K=64) - never matches SD - disabled "MIOPEN_DEBUG_CONV_DIRECT_ASM_7X7C3H224W224": "0", - # WRW3X3 / WRW1X1: FP32-only weight-gradient (training only) — disabled for inference + # WRW3X3 / WRW1X1: FP32-only weight-gradient (training only) - disabled for inference "MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW3X3": "0", "MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW1X1": "0", # PERF_VALS intentionally blank: MIOpen reads this as a config string not a boolean; @@ -110,30 +110,30 @@ RDNA2: Dict[str, str] = { "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_PERF_VALS": "", "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_SEARCH_OPTIMIZED": "1", "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_AI_HEUR": "1", - # NAIVE_CONV_FWD: scalar FP32 reference solver — IsApplicable does NOT reliably filter for FP16; + # NAIVE_CONV_FWD: scalar FP32 reference solver - IsApplicable does NOT reliably filter for FP16; # can be selected for unusual shapes (e.g. VAE decoder 3-ch output) and returns dtype=float32 "MIOPEN_DEBUG_CONV_DIRECT_NAIVE_CONV_FWD": "0", - # Direct OCL — dtype notes - # FWD / FWD1X1: FP32/FP16 forward — enabled + # Direct OCL - dtype notes + # FWD / FWD1X1: FP32/FP16 forward - enabled "MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD": "1", "MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1": "1", - # FWD11X11: requires 11*11 kernel — no SD match — disabled + # FWD11X11: requires 11*11 kernel - no SD match - disabled "MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD11X11": "0", - # FWDGEN: FP32 generic OCL fallback — IsApplicable does NOT reliably reject for FP16; - # can produce dtype=float32 output for FP16 inputs — disabled + # FWDGEN: FP32 generic OCL fallback - IsApplicable does NOT reliably reject for FP16; + # can produce dtype=float32 output for FP16 inputs - disabled "MIOPEN_DEBUG_CONV_DIRECT_OCL_FWDGEN": "0", - # WRW2 / WRW53 / WRW1X1: training-only weight-gradient — disabled + # WRW2 / WRW53 / WRW1X1: training-only weight-gradient - disabled "MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW2": "0", "MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW53": "0", "MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW1X1": "0", - # Winograd RxS — dtype per MIOpen docs - # WINOGRAD_3X3: FP32-only — harmless (IsApplicable rejects for fp16); enabled + # Winograd RxS - dtype per MIOpen docs + # WINOGRAD_3X3: FP32-only - harmless (IsApplicable rejects for fp16); enabled "MIOPEN_DEBUG_AMD_WINOGRAD_3X3": "1", - # RXS: covers FP32/FP16 F(3,3) Fwd/Bwd + FP32 F(3,2) WrW — keep enabled (fp16 fwd/bwd path exists) + # RXS: covers FP32/FP16 F(3,3) Fwd/Bwd + FP32 F(3,2) WrW - keep enabled (fp16 fwd/bwd path exists) "MIOPEN_DEBUG_AMD_WINOGRAD_RXS": "1", - # RXS_FWD_BWD: FP32/FP16 — explicitly the fp16-capable subset + # RXS_FWD_BWD: FP32/FP16 - explicitly the fp16-capable subset "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_FWD_BWD": "1", - # RXS_WRW: FP32 WrW only — training-only, disabled for inference fp16 profile + # RXS_WRW: FP32 WrW only - training-only, disabled for inference fp16 profile "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_WRW": "0", # RXS_F3X2: FP32/FP16 Fwd/Bwd "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F3X2": "1", @@ -141,15 +141,15 @@ RDNA2: Dict[str, str] = { "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3": "1", # RXS_F2X3_G1: FP32/FP16 Fwd/Bwd (non-group convolutions) "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_G1": "1", - # FUSED_WINOGRAD: FP32-only — harmless (IsApplicable rejects for fp16); enabled + # FUSED_WINOGRAD: FP32-only - harmless (IsApplicable rejects for fp16); enabled "MIOPEN_DEBUG_AMD_FUSED_WINOGRAD": "1", - # PERF_VALS intentionally blank: same reason as ASM_1X1U — not a boolean, config string + # PERF_VALS intentionally blank: same reason as ASM_1X1U - not a boolean, config string "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_PERF_VALS": "", - # Fury/Rage Winograd — NOT available on RDNA2 + # Fury/Rage Winograd - NOT available on RDNA2 "MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F2X3": "0", "MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F3X2": "0", "MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3": "0", - # MPASS — only F3x2 and F3x3 are safe on RDNA2 + # MPASS - only F3x2 and F3x3 are safe on RDNA2 "MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X2": "1", "MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X3": "1", "MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X4": "0", @@ -159,50 +159,50 @@ RDNA2: Dict[str, str] = { "MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F5X4": "0", "MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X2": "0", "MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X3": "0", - # ASM Implicit GEMM — forward V4R1 only; no GTC/XDLOPS on RDNA2 - # BWD (backward data-gradient) and WrW (weight-gradient) are training-only — disabled + # ASM Implicit GEMM - forward V4R1 only; no GTC/XDLOPS on RDNA2 + # BWD (backward data-gradient) and WrW (weight-gradient) are training-only - disabled "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_V4R1": "1", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_V4R1_1X1": "1", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_BWD_V4R1": "0", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_V4R1": "0", - # HIP Implicit GEMM — non-XDLOPS V4R1/R4 forward only - # BWD (backward data-gradient) and WrW (weight-gradient) are training-only — disabled + # HIP Implicit GEMM - non-XDLOPS V4R1/R4 forward only + # BWD (backward data-gradient) and WrW (weight-gradient) are training-only - disabled "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1": "1", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4": "1", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1": "0", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1": "0", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R1": "0", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4": "0", - # Group Conv XDLOPS / CK default kernels — RDNA3/4 only, not available on RDNA2 + # Group Conv XDLOPS / CK default kernels - RDNA3/4 only, not available on RDNA2 "MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS": "0", "MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS_AI_HEUR": "0", "MIOPEN_DEBUG_CK_DEFAULT_KERNELS": "0", } # --------------------------------------------------------------------------- -# RDNA3 — gfx1100 (RX 7000 series) +# RDNA3 - gfx1100 (RX 7000 series) # Fury Winograd added; MPASS F3x4 enabled; Group Conv XDLOPS + CK default kernels enabled # --------------------------------------------------------------------------- RDNA3: Dict[str, str] = { **RDNA2, - # Fury Winograd — introduced for gfx1100 (RDNA3) + # Fury Winograd - introduced for gfx1100 (RDNA3) "MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F2X3": "1", "MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F3X2": "1", # Wider MPASS on RDNA3 "MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X4": "1", - # Group Conv XDLOPS / CK — available from gfx1100 (RDNA3) onwards + # Group Conv XDLOPS / CK - available from gfx1100 (RDNA3) onwards "MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS": "1", "MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS_AI_HEUR": "1", "MIOPEN_DEBUG_CK_DEFAULT_KERNELS": "1", } # --------------------------------------------------------------------------- -# RDNA4 — gfx1200 (RX 9000 series) +# RDNA4 - gfx1200 (RX 9000 series) # Rage Winograd added; MPASS F3x5 enabled # --------------------------------------------------------------------------- RDNA4: Dict[str, str] = { **RDNA3, - # Rage Winograd — introduced for gfx1200 (RDNA4) + # Rage Winograd - introduced for gfx1200 (RDNA4) "MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3": "1", # Wider MPASS on RDNA4 "MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X5": "1", diff --git a/scripts/rocm/rocm_vars.py b/scripts/rocm/rocm_vars.py index 0e2235015..1cf30231f 100644 --- a/scripts/rocm/rocm_vars.py +++ b/scripts/rocm/rocm_vars.py @@ -1,39 +1,51 @@ -from typing import Dict, Any, List, Tuple +from typing import Dict, Any, List, Tuple # --- General MIOpen/rocBLAS variables (dropdown/textbox/checkbox) --- GENERAL_VARS: Dict[str, Dict[str, Any]] = { - - # ── GEMM backend selector + companion toggles ────────────────────────── + "MIOPEN_SYSTEM_DB_PATH": { + "default": "{VIRTUAL_ENV}\\Lib\\site-packages\\_rocm_sdk_devel\\bin\\", + "desc": "MIOpen system DB path", + "widget": "textbox", + "options": None, + "restart_required": True, + }, + "ROCBLAS_TENSILE_LIBPATH": { + "default": "{VIRTUAL_ENV}\\Lib\\site-packages\\_rocm_sdk_devel\\bin\\rocblas\\library", + "desc": "rocBLAS Tensile library path", + "widget": "textbox", + "options": None, + "restart_required": True, + }, + # -- GEMM backend selector + companion toggles -------------------------- "MIOPEN_GEMM_ENFORCE_BACKEND": { "default": "1", - "desc": "Enforce GEMM backend", + "desc": "GEMM backend", "widget": "dropdown", "options": [("1 - rocBLAS", "1"), ("5 - hipBLASLt", "5")], "restart_required": False, }, "PYTORCH_ROCM_USE_ROCBLAS": { "default": "0", - "desc": "PyTorch ROCm: prioritise rocBLAS for linear algebra", + "desc": "PyTorch: Use rocBLAS.", "widget": "dropdown", "options": [("0 - Off", "0"), ("1 - On", "1")], "restart_required": True, }, "PYTORCH_HIPBLASLT_DISABLE": { "default": "1", - "desc": "Disable PyTorch hipBLASLt dispatcher", + "desc": "PyTorch: Use hipBLASLt.", "widget": "dropdown", "options": [("0 - Allow hipBLASLt", "0"), ("1 - Disable hipBLASLt", "1")], "restart_required": True, }, "ROCBLAS_USE_HIPBLASLT": { "default": "0", - "desc": "rocBLAS: use hipBLASLt backend (0 = Tensile)", + "desc": "rocBLAS: use hipBLASLt backend.", "widget": "dropdown", "options": [("0 - Tensile (rocBLAS)", "0"), ("1 - hipBLASLt", "1")], "restart_required": True, }, - - # ── MIOpen behavioural settings ──────────────────────────────────────── + # -- MIOpen behavioural settings ---------------------------------------- "MIOPEN_FIND_MODE": { "default": "2", "desc": "MIOpen Find Mode", @@ -57,34 +69,21 @@ GENERAL_VARS: Dict[str, Dict[str, Any]] = { }, "MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC": { "default": "0", - "desc": "Deterministic convolution (reproducible results, may be slower)", + "desc": "Deterministic convolutions", "widget": "dropdown", "options": [("0 - Off", "0"), ("1 - On", "1")], "restart_required": False, }, - # ── Paths / sizes ────────────────────────────────────────────────────── - "MIOPEN_SYSTEM_DB_PATH": { - "default": "{VIRTUAL_ENV}\\Lib\\site-packages\\_rocm_sdk_devel\\bin\\", - "desc": "MIOpen system DB path", - "widget": "textbox", - "options": None, - "restart_required": True, - }, + # -- Paths / sizes ------------------------------------------------------ + "MIOPEN_CONVOLUTION_MAX_WORKSPACE": { "default": "1073741824", - "desc": "MIOpen convolution max workspace (bytes; 1 GB default)", + "desc": "MIOpen convolutions: max workspace (bytes; 1 GB)", "widget": "textbox", "options": None, "restart_required": False, }, - "ROCBLAS_TENSILE_LIBPATH": { - "default": "{VIRTUAL_ENV}\\Lib\\site-packages\\_rocm_sdk_devel\\bin\\rocblas\\library", - "desc": "rocBLAS Tensile library path", - "widget": "textbox", - "options": None, - "restart_required": True, - }, "ROCBLAS_DEVICE_MEMORY_SIZE": { "default": "", "desc": "rocBLAS workspace size in bytes (empty = dynamic)", @@ -94,13 +93,13 @@ GENERAL_VARS: Dict[str, Dict[str, Any]] = { }, "PYTORCH_TUNABLEOP_CACHE_DIR": { "default": "{ROOT}\\models\\tunable", - "desc": "TunableOp: kernel profile cache directory", + "desc": "TunableOp cache directory", "widget": "textbox", "options": None, "restart_required": False, }, - # ── rocBLAS settings ─────────────────────────────────────────────────── + # -- rocBLAS settings --------------------------------------------------- "ROCBLAS_STREAM_ORDER_ALLOC": { "default": "1", "desc": "rocBLAS stream-ordered memory allocation", @@ -110,27 +109,27 @@ GENERAL_VARS: Dict[str, Dict[str, Any]] = { }, "ROCBLAS_DEFAULT_ATOMICS_MODE": { "default": "1", - "desc": "rocBLAS default atomics mode (1 = allow non-deterministic for performance)", + "desc": "rocBLAS allow atomics", "widget": "dropdown", "options": [("0 - Off (deterministic)", "0"), ("1 - On (performance)", "1")], "restart_required": False, }, "PYTORCH_TUNABLEOP_ROCBLAS_ENABLED": { "default": "0", - "desc": "TunableOp: wrap and optimise rocBLAS GEMM calls", + "desc": "TunableOp: Enable tuning", "widget": "dropdown", "options": [("0 - Off", "0"), ("1 - On", "1")], "restart_required": False, }, "PYTORCH_TUNABLEOP_TUNING": { "default": "0", - "desc": "TunableOp: tuning mode (1 = benchmark; 0 = use saved CSV)", + "desc": "TunableOp: Tuning mode", "widget": "dropdown", - "options": [("0 - Use saved CSV", "0"), ("1 - Benchmark new shapes", "1")], + "options": [("0 - Use Cache", "0"), ("1 - Benchmark new shapes", "1")], "restart_required": False, }, - # ── hipBLASLt settings ───────────────────────────────────────────────── + # -- hipBLASLt settings ------------------------------------------------- "PYTORCH_TUNABLEOP_HIPBLASLT_ENABLED": { "default": "0", "desc": "TunableOp: benchmark hipBLASLt kernels", @@ -139,7 +138,7 @@ GENERAL_VARS: Dict[str, Dict[str, Any]] = { "restart_required": False, }, - # ── Logging: MIOpen → rocBLAS → hipBLASLt ───────────────────────────── + # -- Logging: MIOpen -> rocBLAS -> hipBLASLt ----------------------------- "MIOPEN_LOG_LEVEL": { "default": "0", "desc": "MIOpen log verbosity level", @@ -171,13 +170,13 @@ GENERAL_VARS: Dict[str, Dict[str, Any]] = { } # --- Solver toggles (inference/FWD only, RDNA2/3/4 compatible) --- -# Removed entirely — not representable in the UI, cannot be set by users: -# WRW (weight-gradient) and BWD (data-gradient) — training passes only, never run during inference -# XDLOPS/CK CDNA-exclusive (MI100/MI200/MI300 matrix engine variants) — not on any RDNA -# Fixed-geometry (5x10, 7x7-ImageNet, 11x11) — shapes never appear in SD/video inference -# FP32-reference (NAIVE_CONV_FWD, FWDGEN) — IsApplicable() unreliable for FP16/BF16 -# Wide MPASS (F3x4..F7x3) — kernel sizes that cannot match any SD convolution shape -# Disabled by default (added but off): RDNA3/4-only — Group Conv XDLOPS, CK default kernels +# Removed entirely - not representable in the UI, cannot be set by users: +# WRW (weight-gradient) and BWD (data-gradient) - training passes only, never run during inference +# XDLOPS/CK CDNA-exclusive (MI100/MI200/MI300 matrix engine variants) - not on any RDNA +# Fixed-geometry (5x10, 7x7-ImageNet, 11x11) - shapes never appear in SD/video inference +# FP32-reference (NAIVE_CONV_FWD, FWDGEN) - IsApplicable() unreliable for FP16/BF16 +# Wide MPASS (F3x4..F7x3) - kernel sizes that cannot match any SD convolution shape +# Disabled by default (added but off): RDNA3/4-only - Group Conv XDLOPS, CK default kernels _SOLVER_DESCS: Dict[str, str] = {} _SOLVER_DESCS.update({ @@ -200,7 +199,7 @@ _SOLVER_DESCS.update({ "MIOPEN_DEBUG_ATTN_SOFTMAX": "Enable Attention Softmax", }) _SOLVER_DESCS.update({ - # Direct ASM — FWD inference only (WRW, fixed-geometry, FP32-reference removed) + # Direct ASM - FWD inference only (WRW, fixed-geometry, FP32-reference removed) "MIOPEN_DEBUG_CONV_DIRECT_ASM_3X3U": "Enable Direct ASM 3x3U", "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U": "Enable Direct ASM 1x1U", "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1UV2": "Enable Direct ASM 1x1UV2", @@ -208,12 +207,12 @@ _SOLVER_DESCS.update({ "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_AI_HEUR": "Enable Direct ASM 1x1U AI Heuristic", }) _SOLVER_DESCS.update({ - # Direct OCL — FWD inference only (WRW, FWD11X11 fixed-geom, FWDGEN FP32-ref removed) + # Direct OCL - FWD inference only (WRW, FWD11X11 fixed-geom, FWDGEN FP32-ref removed) "MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD": "Enable Direct OCL FWD", "MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1": "Enable Direct OCL FWD1X1", }) _SOLVER_DESCS.update({ - # Winograd FWD — WRW removed; Fury/Rage kept as RDNA3/4 inference (off by default) + # Winograd FWD - WRW removed; Fury/Rage kept as RDNA3/4 inference (off by default) "MIOPEN_DEBUG_AMD_WINOGRAD_3X3": "Enable AMD Winograd 3x3", "MIOPEN_DEBUG_AMD_WINOGRAD_RXS": "Enable AMD Winograd RxS", "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_FWD_BWD": "Enable AMD Winograd RxS FWD", @@ -226,32 +225,32 @@ _SOLVER_DESCS.update({ "MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3": "Enable AMD Winograd Rage RxS F2x3", }) _SOLVER_DESCS.update({ - # Multi-pass Winograd — only F3x2/F3x3 match typical 3x3 SD shapes; wider kernels removed + # Multi-pass Winograd - only F3x2/F3x3 match typical 3x3 SD shapes; wider kernels removed "MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X2": "Enable AMD Winograd MPASS F3x2", "MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X3": "Enable AMD Winograd MPASS F3x3", }) _SOLVER_DESCS.update({ - # Implicit GEMM FWD — BWD/WRW (training), CDNA-exclusive XDLOPS variants removed + # Implicit GEMM FWD - BWD/WRW (training), CDNA-exclusive XDLOPS variants removed "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_V4R1": "Enable ASM Implicit GEMM FWD V4R1", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_V4R1_1X1": "Enable ASM Implicit GEMM FWD V4R1 1x1", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1": "Enable HIP Implicit GEMM FWD V4R1", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4": "Enable HIP Implicit GEMM FWD V4R4", }) _SOLVER_DESCS.update({ - # Group Conv XDLOPS FWD — RDNA3/4 (gfx1100+) only; disabled by default + # Group Conv XDLOPS FWD - RDNA3/4 (gfx1100+) only; disabled by default "MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS": "Enable Group Conv Implicit GEMM XDLOPS FWD", "MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS_AI_HEUR": "Enable Group Conv Implicit GEMM XDLOPS FWD AI Heuristic", - # CK (Composable Kernel) default kernels — RDNA3/4 (gfx1100+); disabled by default + # CK (Composable Kernel) default kernels - RDNA3/4 (gfx1100+); disabled by default "MIOPEN_DEBUG_CK_DEFAULT_KERNELS": "Enable CK (Composable Kernel) default kernels", }) # Solvers still in the registry but disabled by default. -# FORCE_IMMED_MODE_FALLBACK — overrides FIND_MODE entirely, defeats tuning DB -# Fury RxS F2x3/F3x2 — RDNA3/4-only; harmless on RDNA2 but won't select -# Rage RxS F2x3 — RDNA4-only -# Group Conv XDLOPS — RDNA3/4-only (gfx1100+) -# CK_DEFAULT_KERNELS — RDNA3/4-only (gfx1100+) +# FORCE_IMMED_MODE_FALLBACK - overrides FIND_MODE entirely, defeats tuning DB +# Fury RxS F2x3/F3x2 - RDNA3/4-only; harmless on RDNA2 but won't select +# Rage RxS F2x3 - RDNA4-only +# Group Conv XDLOPS - RDNA3/4-only (gfx1100+) +# CK_DEFAULT_KERNELS - RDNA3/4-only (gfx1100+) SOLVER_DISABLED_BY_DEFAULT = { "MIOPEN_DEBUG_FORCE_IMMED_MODE_FALLBACK", "MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F2X3", From d2a47ee0eda97ffbc966a07427ec3b50ecdae257 Mon Sep 17 00:00:00 2001 From: resonantsky Date: Fri, 3 Apr 2026 10:29:53 +0200 Subject: [PATCH 06/12] code quality and layout fixes --- scripts/rocm/rocm_vars.py | 51 ++++++++++++--------------------------- 1 file changed, 16 insertions(+), 35 deletions(-) diff --git a/scripts/rocm/rocm_vars.py b/scripts/rocm/rocm_vars.py index 1cf30231f..0ecdf8298 100644 --- a/scripts/rocm/rocm_vars.py +++ b/scripts/rocm/rocm_vars.py @@ -1,6 +1,5 @@ -from typing import Dict, Any, List, Tuple +from typing import Dict, Any, List, Tuple -# --- General MIOpen/rocBLAS variables (dropdown/textbox/checkbox) --- GENERAL_VARS: Dict[str, Dict[str, Any]] = { "MIOPEN_SYSTEM_DB_PATH": { "default": "{VIRTUAL_ENV}\\Lib\\site-packages\\_rocm_sdk_devel\\bin\\", @@ -16,7 +15,6 @@ GENERAL_VARS: Dict[str, Dict[str, Any]] = { "options": None, "restart_required": True, }, - # -- GEMM backend selector + companion toggles -------------------------- "MIOPEN_GEMM_ENFORCE_BACKEND": { "default": "1", "desc": "GEMM backend", @@ -26,26 +24,25 @@ GENERAL_VARS: Dict[str, Dict[str, Any]] = { }, "PYTORCH_ROCM_USE_ROCBLAS": { "default": "0", - "desc": "PyTorch: Use rocBLAS.", + "desc": "PyTorch: Use rocBLAS", "widget": "dropdown", "options": [("0 - Off", "0"), ("1 - On", "1")], "restart_required": True, }, "PYTORCH_HIPBLASLT_DISABLE": { "default": "1", - "desc": "PyTorch: Use hipBLASLt.", + "desc": "PyTorch: Use hipBLASLt", "widget": "dropdown", "options": [("0 - Allow hipBLASLt", "0"), ("1 - Disable hipBLASLt", "1")], "restart_required": True, }, "ROCBLAS_USE_HIPBLASLT": { "default": "0", - "desc": "rocBLAS: use hipBLASLt backend.", + "desc": "rocBLAS: use hipBLASLt backend", "widget": "dropdown", "options": [("0 - Tensile (rocBLAS)", "0"), ("1 - hipBLASLt", "1")], "restart_required": True, }, - # -- MIOpen behavioural settings ---------------------------------------- "MIOPEN_FIND_MODE": { "default": "2", "desc": "MIOpen Find Mode", @@ -74,9 +71,6 @@ GENERAL_VARS: Dict[str, Dict[str, Any]] = { "options": [("0 - Off", "0"), ("1 - On", "1")], "restart_required": False, }, - - # -- Paths / sizes ------------------------------------------------------ - "MIOPEN_CONVOLUTION_MAX_WORKSPACE": { "default": "1073741824", "desc": "MIOpen convolutions: max workspace (bytes; 1 GB)", @@ -99,7 +93,6 @@ GENERAL_VARS: Dict[str, Dict[str, Any]] = { "restart_required": False, }, - # -- rocBLAS settings --------------------------------------------------- "ROCBLAS_STREAM_ORDER_ALLOC": { "default": "1", "desc": "rocBLAS stream-ordered memory allocation", @@ -128,8 +121,6 @@ GENERAL_VARS: Dict[str, Dict[str, Any]] = { "options": [("0 - Use Cache", "0"), ("1 - Benchmark new shapes", "1")], "restart_required": False, }, - - # -- hipBLASLt settings ------------------------------------------------- "PYTORCH_TUNABLEOP_HIPBLASLT_ENABLED": { "default": "0", "desc": "TunableOp: benchmark hipBLASLt kernels", @@ -137,8 +128,6 @@ GENERAL_VARS: Dict[str, Dict[str, Any]] = { "options": [("0 - Off", "0"), ("1 - On", "1")], "restart_required": False, }, - - # -- Logging: MIOpen -> rocBLAS -> hipBLASLt ----------------------------- "MIOPEN_LOG_LEVEL": { "default": "0", "desc": "MIOpen log verbosity level", @@ -169,14 +158,6 @@ GENERAL_VARS: Dict[str, Dict[str, Any]] = { }, } -# --- Solver toggles (inference/FWD only, RDNA2/3/4 compatible) --- -# Removed entirely - not representable in the UI, cannot be set by users: -# WRW (weight-gradient) and BWD (data-gradient) - training passes only, never run during inference -# XDLOPS/CK CDNA-exclusive (MI100/MI200/MI300 matrix engine variants) - not on any RDNA -# Fixed-geometry (5x10, 7x7-ImageNet, 11x11) - shapes never appear in SD/video inference -# FP32-reference (NAIVE_CONV_FWD, FWDGEN) - IsApplicable() unreliable for FP16/BF16 -# Wide MPASS (F3x4..F7x3) - kernel sizes that cannot match any SD convolution shape -# Disabled by default (added but off): RDNA3/4-only - Group Conv XDLOPS, CK default kernels _SOLVER_DESCS: Dict[str, str] = {} _SOLVER_DESCS.update({ @@ -199,7 +180,7 @@ _SOLVER_DESCS.update({ "MIOPEN_DEBUG_ATTN_SOFTMAX": "Enable Attention Softmax", }) _SOLVER_DESCS.update({ - # Direct ASM - FWD inference only (WRW, fixed-geometry, FP32-reference removed) + # Direct ASM — FWD inference only (WRW, fixed-geometry, FP32-reference removed) "MIOPEN_DEBUG_CONV_DIRECT_ASM_3X3U": "Enable Direct ASM 3x3U", "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U": "Enable Direct ASM 1x1U", "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1UV2": "Enable Direct ASM 1x1UV2", @@ -207,12 +188,12 @@ _SOLVER_DESCS.update({ "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_AI_HEUR": "Enable Direct ASM 1x1U AI Heuristic", }) _SOLVER_DESCS.update({ - # Direct OCL - FWD inference only (WRW, FWD11X11 fixed-geom, FWDGEN FP32-ref removed) + # Direct OCL — FWD inference only (WRW, FWD11X11 fixed-geom, FWDGEN FP32-ref removed) "MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD": "Enable Direct OCL FWD", "MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1": "Enable Direct OCL FWD1X1", }) _SOLVER_DESCS.update({ - # Winograd FWD - WRW removed; Fury/Rage kept as RDNA3/4 inference (off by default) + # Winograd FWD — WRW removed; Fury/Rage kept as RDNA3/4 inference (off by default) "MIOPEN_DEBUG_AMD_WINOGRAD_3X3": "Enable AMD Winograd 3x3", "MIOPEN_DEBUG_AMD_WINOGRAD_RXS": "Enable AMD Winograd RxS", "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_FWD_BWD": "Enable AMD Winograd RxS FWD", @@ -225,32 +206,32 @@ _SOLVER_DESCS.update({ "MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3": "Enable AMD Winograd Rage RxS F2x3", }) _SOLVER_DESCS.update({ - # Multi-pass Winograd - only F3x2/F3x3 match typical 3x3 SD shapes; wider kernels removed + # Multi-pass Winograd — only F3x2/F3x3 match typical 3x3 SD shapes; wider kernels removed "MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X2": "Enable AMD Winograd MPASS F3x2", "MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X3": "Enable AMD Winograd MPASS F3x3", }) _SOLVER_DESCS.update({ - # Implicit GEMM FWD - BWD/WRW (training), CDNA-exclusive XDLOPS variants removed + # Implicit GEMM FWD — BWD/WRW (training), CDNA-exclusive XDLOPS variants removed "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_V4R1": "Enable ASM Implicit GEMM FWD V4R1", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_V4R1_1X1": "Enable ASM Implicit GEMM FWD V4R1 1x1", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1": "Enable HIP Implicit GEMM FWD V4R1", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4": "Enable HIP Implicit GEMM FWD V4R4", }) _SOLVER_DESCS.update({ - # Group Conv XDLOPS FWD - RDNA3/4 (gfx1100+) only; disabled by default + # Group Conv XDLOPS FWD — RDNA3/4 (gfx1100+) only; disabled by default "MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS": "Enable Group Conv Implicit GEMM XDLOPS FWD", "MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS_AI_HEUR": "Enable Group Conv Implicit GEMM XDLOPS FWD AI Heuristic", - # CK (Composable Kernel) default kernels - RDNA3/4 (gfx1100+); disabled by default + # CK (Composable Kernel) default kernels — RDNA3/4 (gfx1100+); disabled by default "MIOPEN_DEBUG_CK_DEFAULT_KERNELS": "Enable CK (Composable Kernel) default kernels", }) # Solvers still in the registry but disabled by default. -# FORCE_IMMED_MODE_FALLBACK - overrides FIND_MODE entirely, defeats tuning DB -# Fury RxS F2x3/F3x2 - RDNA3/4-only; harmless on RDNA2 but won't select -# Rage RxS F2x3 - RDNA4-only -# Group Conv XDLOPS - RDNA3/4-only (gfx1100+) -# CK_DEFAULT_KERNELS - RDNA3/4-only (gfx1100+) +# FORCE_IMMED_MODE_FALLBACK — overrides FIND_MODE entirely, defeats tuning DB +# Fury RxS F2x3/F3x2 — RDNA3/4-only; harmless on RDNA2 but won't select +# Rage RxS F2x3 — RDNA4-only +# Group Conv XDLOPS — RDNA3/4-only (gfx1100+) +# CK_DEFAULT_KERNELS — RDNA3/4-only (gfx1100+) SOLVER_DISABLED_BY_DEFAULT = { "MIOPEN_DEBUG_FORCE_IMMED_MODE_FALLBACK", "MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F2X3", From 24f4490a59b6a542bcd83d1bccc6ca48d8765da8 Mon Sep 17 00:00:00 2001 From: resonantsky <90446525+resonantsky@users.noreply.github.com> Date: Fri, 3 Apr 2026 11:09:55 +0200 Subject: [PATCH 07/12] Apply suggestion from @awsr Co-authored-by: awsr <43862868+awsr@users.noreply.github.com> --- scripts/rocm/rocm_mgr.py | 2 ++ 1 file changed, 2 insertions(+) diff --git a/scripts/rocm/rocm_mgr.py b/scripts/rocm/rocm_mgr.py index ea987a421..913435f5f 100644 --- a/scripts/rocm/rocm_mgr.py +++ b/scripts/rocm/rocm_mgr.py @@ -7,6 +7,8 @@ from typing import Dict, Optional import installer from modules.logger import log from modules.json_helpers import readfile, writefile +from modules.shared import cmd_opts, opts +from modules.devices import has_rocm from scripts.rocm.rocm_vars import ROCM_ENV_VARS # pylint: disable=no-name-in-module from scripts.rocm import rocm_profiles # pylint: disable=no-name-in-module From ee3b141297e33725b98b39855e65b0312aa32333 Mon Sep 17 00:00:00 2001 From: resonantsky Date: Fri, 3 Apr 2026 11:21:09 +0200 Subject: [PATCH 08/12] fixes as requested by review --- scripts/rocm/rocm_mgr.py | 12 +++++++----- scripts/rocm_ext.py | 9 ++++----- 2 files changed, 11 insertions(+), 10 deletions(-) diff --git a/scripts/rocm/rocm_mgr.py b/scripts/rocm/rocm_mgr.py index 913435f5f..f9acd61a0 100644 --- a/scripts/rocm/rocm_mgr.py +++ b/scripts/rocm/rocm_mgr.py @@ -7,8 +7,10 @@ from typing import Dict, Optional import installer from modules.logger import log from modules.json_helpers import readfile, writefile -from modules.shared import cmd_opts, opts + +from modules.shared import cmd_opts from modules.devices import has_rocm + from scripts.rocm.rocm_vars import ROCM_ENV_VARS # pylint: disable=no-name-in-module from scripts.rocm import rocm_profiles # pylint: disable=no-name-in-module @@ -110,7 +112,7 @@ def _resolve_dtype() -> str: pass try: from modules import shared as _sh # pylint: disable=import-outside-toplevel - v = getattr(getattr(_sh, 'opts', None), 'cuda_dtype', None) + v = getattr(_sh.opts, 'cuda_dtype', None) if v in ('FP16', 'BF16', 'FP32'): return v except Exception: @@ -460,7 +462,7 @@ def info() -> dict: if ufiles: udb["files"] = ufiles - # User cache (~/.miopen/cache/) + # User cache (~/.miopen/cache/) cache_base = Path.home() / ".miopen" / "cache" db_hash = _extract_db_hash(user_db_path) if user_db_path.exists() else "" cache_path = cache_base / db_hash if db_hash else cache_base @@ -481,10 +483,10 @@ def info() -> dict: # Apply saved config to os.environ at import time (only when ROCm is present) -if is_rocm: +if has_rocm() and getattr(cmd_opts, 'use_rocm', False) and sys.platform == 'win32': try: apply_env() except Exception as _e: print(f"[rocm_mgr] Warning: failed to apply env at import: {_e}", file=sys.stderr) else: - log.debug('ROCm is not installed — skipping rocm_mgr env apply') + log.debug('ROCm is not installed - skipping rocm_mgr env apply') diff --git a/scripts/rocm_ext.py b/scripts/rocm_ext.py index e8ddf24ed..6ea9c0143 100644 --- a/scripts/rocm_ext.py +++ b/scripts/rocm_ext.py @@ -5,10 +5,9 @@ from modules import scripts_manager, shared # rocm_mgr exposes package-internal helpers (prefixed _) that are intentionally called here # pylint: disable=protected-access - class ROCmScript(scripts_manager.Script): def title(self): - return "ROCm: Advanced Config" + return "Windows ROCm: Advanced Config" def show(self, _is_img2img): if shared.cmd_opts.use_rocm or installer.torch_info.get('type') == 'rocm': @@ -75,9 +74,9 @@ class ROCmScript(scripts_manager.Script): rules.append(f"#rocm_var_{v.lower()} {{ opacity: 0.45; pointer-events: none; }}") return f"" if rules else "" - with gr.Accordion('ROCm: Advanced Config', open=False, elem_id='rocm_config'): + with gr.Accordion('Windows ROCm: Advanced Config', open=False, elem_id='rocm_config'): with gr.Row(): - gr.HTML("

Advanced configuration for ROCm users.


For best performance ensure that cudnn and torch tunable ops are set to default in Backend Settings.

") + gr.HTML("

Advanced configuration for Windows ROCm users.


This script aims to take the guesswork out of configuring MIOpen and rocBLAS on Windows ROCm, but also to expose the functioning switches of MIOpen for advanced configurations.


For best performance ensure that cuDNN and PyTorch tunable ops are set to default in Backend Settings.

") with gr.Row(): btn_info = gr.Button("Refresh Info", variant="primary", elem_id="rocm_btn_info", size="sm") btn_apply = gr.Button("Apply", variant="primary", elem_id="rocm_btn_apply", size="sm") @@ -96,7 +95,7 @@ class ROCmScript(scripts_manager.Script): # General vars (dropdowns, textboxes, checkboxes) with gr.Group(): - gr.HTML("

MIOpen Settings

") + gr.HTML("

MIOpen Settings

") for name, meta in rocm_vars.GENERAL_VARS.items(): comp = _make_component(name, meta, config) var_names.append(name) From 2177609e5443a5bcc6d494000f8af6038133ccfe Mon Sep 17 00:00:00 2001 From: resonantsky Date: Fri, 3 Apr 2026 11:28:25 +0200 Subject: [PATCH 09/12] further fixes requested by review --- scripts/rocm/rocm_mgr.py | 14 +++++--------- 1 file changed, 5 insertions(+), 9 deletions(-) diff --git a/scripts/rocm/rocm_mgr.py b/scripts/rocm/rocm_mgr.py index f9acd61a0..6ad3eebfc 100644 --- a/scripts/rocm/rocm_mgr.py +++ b/scripts/rocm/rocm_mgr.py @@ -7,8 +7,7 @@ from typing import Dict, Optional import installer from modules.logger import log from modules.json_helpers import readfile, writefile - -from modules.shared import cmd_opts +from modules.shared import cmd_opts, opts from modules.devices import has_rocm from scripts.rocm.rocm_vars import ROCM_ENV_VARS # pylint: disable=no-name-in-module @@ -16,13 +15,11 @@ from scripts.rocm import rocm_profiles # pylint: disable=no-name-in-module def _check_rocm() -> bool: - from modules import shared - if getattr(shared.cmd_opts, 'use_rocm', False): + if getattr(cmd_opts, 'use_rocm', False): return True if installer.torch_info.get('type') == 'rocm': return True - import torch # pylint: disable=import-outside-toplevel - return hasattr(torch.version, 'hip') and torch.version.hip is not None + return has_rocm() is_rocm = _check_rocm() @@ -111,8 +108,7 @@ def _resolve_dtype() -> str: except Exception: pass try: - from modules import shared as _sh # pylint: disable=import-outside-toplevel - v = getattr(_sh.opts, 'cuda_dtype', None) + v = getattr(opts, 'cuda_dtype', None) if v in ('FP16', 'BF16', 'FP32'): return v except Exception: @@ -483,7 +479,7 @@ def info() -> dict: # Apply saved config to os.environ at import time (only when ROCm is present) -if has_rocm() and getattr(cmd_opts, 'use_rocm', False) and sys.platform == 'win32': +if is_rocm and sys.platform == 'win32': try: apply_env() except Exception as _e: From ac9aacac662d968e47108208c6058caddf4789ed Mon Sep 17 00:00:00 2001 From: resonantsky Date: Fri, 3 Apr 2026 11:58:43 +0200 Subject: [PATCH 10/12] edited if installer.torch_info line 470 --- scripts/rocm/rocm_mgr.py | 29 +++++++++-------------------- 1 file changed, 9 insertions(+), 20 deletions(-) diff --git a/scripts/rocm/rocm_mgr.py b/scripts/rocm/rocm_mgr.py index 6ad3eebfc..b35b1ced1 100644 --- a/scripts/rocm/rocm_mgr.py +++ b/scripts/rocm/rocm_mgr.py @@ -7,24 +7,12 @@ from typing import Dict, Optional import installer from modules.logger import log from modules.json_helpers import readfile, writefile -from modules.shared import cmd_opts, opts -from modules.devices import has_rocm +from modules.shared import opts from scripts.rocm.rocm_vars import ROCM_ENV_VARS # pylint: disable=no-name-in-module from scripts.rocm import rocm_profiles # pylint: disable=no-name-in-module -def _check_rocm() -> bool: - if getattr(cmd_opts, 'use_rocm', False): - return True - if installer.torch_info.get('type') == 'rocm': - return True - return has_rocm() - - -is_rocm = _check_rocm() - - CONFIG = Path(os.path.abspath(os.path.join('data', 'rocm.json'))) _cache: Optional[Dict[str, str]] = None # loaded once, invalidated on save @@ -479,10 +467,11 @@ def info() -> dict: # Apply saved config to os.environ at import time (only when ROCm is present) -if is_rocm and sys.platform == 'win32': - try: - apply_env() - except Exception as _e: - print(f"[rocm_mgr] Warning: failed to apply env at import: {_e}", file=sys.stderr) -else: - log.debug('ROCm is not installed - skipping rocm_mgr env apply') +if installer.torch_info.get('type', None) == 'rocm': + if sys.platform == 'win32': + try: + apply_env() + except Exception as _e: + log.debug(f"[rocm_mgr] Warning: failed to apply env at import: {_e}") + else: + log.debug('Skipping ROCm Environment Manager: Currently only Windows is supported.') From e49d6262e903ea2e2b2b229dcb85fd1cd23b2159 Mon Sep 17 00:00:00 2001 From: resonantsky Date: Fri, 3 Apr 2026 12:25:06 +0200 Subject: [PATCH 11/12] UI edits, small corrections --- scripts/rocm_ext.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/scripts/rocm_ext.py b/scripts/rocm_ext.py index 6ea9c0143..d3a5f1467 100644 --- a/scripts/rocm_ext.py +++ b/scripts/rocm_ext.py @@ -74,9 +74,9 @@ class ROCmScript(scripts_manager.Script): rules.append(f"#rocm_var_{v.lower()} {{ opacity: 0.45; pointer-events: none; }}") return f"" if rules else "" - with gr.Accordion('Windows ROCm: Advanced Config', open=False, elem_id='rocm_config'): + with gr.Accordion('ROCm: Advanced Config', open=False, elem_id='rocm_config'): with gr.Row(): - gr.HTML("

Advanced configuration for Windows ROCm users.


This script aims to take the guesswork out of configuring MIOpen and rocBLAS on Windows ROCm, but also to expose the functioning switches of MIOpen for advanced configurations.


For best performance ensure that cuDNN and PyTorch tunable ops are set to default in Backend Settings.

") + gr.HTML("

Advanced configuration for ROCm users.


This script aims to take the guesswork out of configuring MIOpen and rocBLAS on Windows ROCm, but also to expose the functioning switches of MIOpen for advanced configurations.


For best performance ensure that cuDNN and PyTorch tunable ops are set to default in Backend Settings.


This script was written with the intent to support ROCm Windows users, it should however, function identically for Linux users.


") with gr.Row(): btn_info = gr.Button("Refresh Info", variant="primary", elem_id="rocm_btn_info", size="sm") btn_apply = gr.Button("Apply", variant="primary", elem_id="rocm_btn_apply", size="sm") From a809b616e69095b0390a351f29439c57b87f6188 Mon Sep 17 00:00:00 2001 From: resonantsky Date: Fri, 3 Apr 2026 15:46:15 +0200 Subject: [PATCH 12/12] Restoring platform agnosticism, Linux users report OK --- scripts/rocm/rocm_mgr.py | 11 ++++------- 1 file changed, 4 insertions(+), 7 deletions(-) diff --git a/scripts/rocm/rocm_mgr.py b/scripts/rocm/rocm_mgr.py index b35b1ced1..bb6b2c9a3 100644 --- a/scripts/rocm/rocm_mgr.py +++ b/scripts/rocm/rocm_mgr.py @@ -468,10 +468,7 @@ def info() -> dict: # Apply saved config to os.environ at import time (only when ROCm is present) if installer.torch_info.get('type', None) == 'rocm': - if sys.platform == 'win32': - try: - apply_env() - except Exception as _e: - log.debug(f"[rocm_mgr] Warning: failed to apply env at import: {_e}") - else: - log.debug('Skipping ROCm Environment Manager: Currently only Windows is supported.') + try: + apply_env() + except Exception as _e: + log.debug(f"[rocm_mgr] Warning: failed to apply env at import: {_e}")