Public/automatic
1
0
Fork 0
mirror of https://github.com/vladmandic/automatic synced 2026-09-19 17:24:32 +02:00
Code Issues Packages Projects Releases Wiki Activity

ROCm: Advanced Config

Browse Source
This commit is contained in:
resonantsky
2026-03-22 23:18:22 +02:00
parent 41fddce2e9
commit 73b854e4cf
5 changed files with 1232 additions and 44 deletions
Download Patch File Download Diff File Expand all files Collapse all files
data/rocm-config.json
+115
View File
@@ -0,0 +1,115 @@
{
"MIOPEN_GEMM_ENFORCE_BACKEND": "1",
"MIOPEN_FIND_MODE": "2",
"MIOPEN_FIND_ENFORCE": "1",
"MIOPEN_SEARCH_CUTOFF": "0",
"MIOPEN_SYSTEM_DB_PATH": "{VIRTUAL_ENV}\\Lib\\site-packages\\_rocm_sdk_devel\\bin\\",
"MIOPEN_LOG_LEVEL": "0",
"MIOPEN_DEBUG_ENABLE": "0",
"ROCBLAS_LAYER": "0",
"HIPBLASLT_LOG_LEVEL": "0",
"MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC": "0",
"MIOPEN_DEBUG_CONV_FFT": "1",
"MIOPEN_DEBUG_CONV_DIRECT": "1",
"MIOPEN_DEBUG_CONV_GEMM": "1",
"MIOPEN_DEBUG_CONV_WINOGRAD": "1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM": "1",
"MIOPEN_DEBUG_CONV_IMMED_FALLBACK": "1",
"MIOPEN_DEBUG_ENABLE_AI_IMMED_MODE_FALLBACK": "1",
"MIOPEN_DEBUG_FORCE_IMMED_MODE_FALLBACK": "0",
"MIOPEN_DEBUG_GCN_ASM_KERNELS": "1",
"MIOPEN_DEBUG_HIP_KERNELS": "1",
"MIOPEN_DEBUG_OPENCL_CONVOLUTIONS": "1",
"MIOPEN_DEBUG_OPENCL_WAVE64_NOWGP": "1",
"MIOPEN_DEBUG_ATTN_SOFTMAX": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_3X3U": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1UV2": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_5X10U2V2": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_7X7C3H224W224": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW3X3": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW1X1": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_PERF_VALS": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_SEARCH_OPTIMIZED": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_AI_HEUR": "1",
"MIOPEN_DEBUG_CONV_DIRECT_NAIVE_CONV_FWD": "1",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD": "1",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1": "1",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD11X11": "1",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_FWDGEN": "1",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW2": "1",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW53": "1",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW1X1": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_3X3": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_FWD_BWD": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_WRW": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F3X2": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_G1": "1",
"MIOPEN_DEBUG_AMD_FUSED_WINOGRAD": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_PERF_VALS": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F2X3": "0",
"MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F3X2": "0",
"MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3": "0",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X2": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X3": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X4": "0",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X5": "0",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X6": "0",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F5X3": "0",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F5X4": "0",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X2": "0",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X3": "0",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F2X3": "0",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F3X3": "0",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F4X3": "0",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F5X3": "0",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F6X3": "0",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F2X3": "0",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F3X3": "0",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F4X3": "0",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F5X3": "0",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F6X3": "0",
"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": "1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_V4R1": "1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_BWD_GTC_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_GTC_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS_NHWC": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_BWD_GTC_XDLOPS_NHWC": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_GTC_XDLOPS_NHWC": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_DLOPS_NCHWC": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_GROUP_BWD_XDLOPS": "0",
"MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_BWD_XDLOPS_AI_HEUR": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_FWD_V4R4_XDLOPS_ADD_VECTOR_LOAD_GEMMN_TUNE_PARAM": "0",
"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": "1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1": "1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R1": "1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4": "1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R5_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4_PADDED_GEMM_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4_PADDED_GEMM_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_XDLOPS_EMULATE": "0",
"MIOPEN_DEBUG_IMPLICIT_GEMM_XDLOPS_INLINE_ASM": "0",
"MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS": "0",
"MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_BWD_XDLOPS": "0",
"MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_WRW_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_V6R1_DLOPS_NCHW": "0",
"MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_BIAS_ACTIV": "0",
"MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_BIAS_RES_ADD_ACTIV": "0",
"MIOPEN_DEBUG_CONV_MLIR_IGEMM_WRW_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_MLIR_IGEMM_BWD_XDLOPS": "0"
}
scripts/rocm/rocm_mgr.py
+308 -18
View File
@@ -1,32 +1,322 @@
import os
import sys
from pathlib import Path
from typing import Dict, Any, Optional
import installer
from modules.logger import log
from modules.json_helpers import readfile, writefile
from scripts.rocm.rocm_vars import ROCM_ENV_VARS, SOLVER_GROUPS # pylint: disable=no-name-in-module
from scripts.rocm import rocm_profiles # pylint: disable=no-name-in-module
CONFIG = os.path.join('data', 'rocm.json')
DATA = {}
def _check_rocm() -> bool:
try:
from modules import shared
if getattr(shared.cmd_opts, 'use_rocm', False):
return True
except Exception:
pass
try:
if installer.torch_info.get('type') == 'rocm':
return True
except Exception:
pass
try:
import torch
if hasattr(torch.version, 'hip') and torch.version.hip is not None:
return True
except Exception:
pass
return False
def load():
global DATA # pylint: disable=global-statement
DATA = readfile(CONFIG, silent=True) or {}
log.debug(f'ROCm load: config={CONFIG} items={len(DATA)}')
is_rocm = _check_rocm()
def reset():
log.info(f"ROCm reset: config={CONFIG}")
DATA.clear()
writefile(DATA, CONFIG)
CONFIG = Path(os.path.abspath(os.path.join('data', 'rocm-config.json')))
_cache: Optional[Dict[str, str]] = None # loaded once, invalidated on save
# Vars that must never be set — they interfere with PyTorch dtype handling
_UNSET_VARS = {
"MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_EXPEREMENTAL_FP16_TRANSFORM",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_PK_ATOMIC_ADD_FP16",
}
def apply(db_path):
DATA["MIOPEN_SYSTEM_DB_PATH"] = db_path
log.info(f'ROCm apply: config={CONFIG} items={len(DATA)}')
writefile(DATA, CONFIG)
# --- venv helpers ---
def _get_venv() -> str:
return os.environ.get("VIRTUAL_ENV", "") or sys.prefix
def info():
for gpu in installer.gpu_info:
gpu['db'] = DATA.get('MIOPEN_SYSTEM_DB_PATH', '')
return installer.gpu_info
def _expand_venv(value: str) -> str:
return value.replace("{VIRTUAL_ENV}", _get_venv())
def _collapse_venv(value: str) -> str:
venv = _get_venv()
if venv and value.startswith(venv):
return "{VIRTUAL_ENV}" + value[len(venv):]
return value
# --- dropdown helpers ---
def _dropdown_display(stored_val: str, options) -> str:
if options and isinstance(options[0], tuple):
return next((label for label, val in options if val == str(stored_val)), str(stored_val))
return str(stored_val)
def _dropdown_stored(display_val: str, options) -> str:
if options and isinstance(options[0], tuple):
return next((val for label, val in options if label == str(display_val)), str(display_val))
return str(display_val)
def _dropdown_choices(options):
if options and isinstance(options[0], tuple):
return [label for label, _ in options]
return options
# --- config I/O ---
def load_config() -> Dict[str, str]:
global _cache # pylint: disable=global-statement
if _cache is None:
if CONFIG.exists():
data = readfile(str(CONFIG), lock=True, as_type="dict")
_cache = data if data else {k: v["default"] for k, v in ROCM_ENV_VARS.items()}
else:
_cache = {k: v["default"] for k, v in ROCM_ENV_VARS.items()}
log.debug(f'ROCm load_config: path={CONFIG} items={len(_cache)}')
return _cache
def save_config(config: Dict[str, str]) -> None:
global _cache # pylint: disable=global-statement
writefile(config, str(CONFIG))
_cache = config
def apply_env(config: Optional[Dict[str, str]] = None) -> None:
if config is None:
config = load_config()
applied = 0
skipped = 0
for var in _UNSET_VARS:
if var in os.environ:
del os.environ[var]
for var, value in config.items():
if var in _UNSET_VARS:
skipped += 1
continue
expanded = _expand_venv(str(value))
if expanded == "":
skipped += 1
continue
os.environ[var] = expanded
applied += 1
def apply_all(names: list, values: list) -> None:
config = load_config().copy()
for name, value in zip(names, values):
if name not in ROCM_ENV_VARS:
log.warning(f'ROCm apply_all: unknown variable={name}')
continue
meta = ROCM_ENV_VARS[name]
if meta["widget"] == "checkbox":
if value is None:
pass # Gradio passed None (component not interacted with) — leave config unchanged
else:
config[name] = "1" if value else "0"
elif meta["widget"] == "radio":
stored = _dropdown_stored(str(value), meta["options"])
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:
config[name] = _collapse_venv(str(value))
save_config(config)
apply_env(config)
def reset_defaults() -> None:
defaults = {k: v["default"] for k, v in ROCM_ENV_VARS.items()}
save_config(defaults)
apply_env(defaults)
log.info('ROCm reset_defaults: config reset to defaults')
def clear_env() -> None:
"""Remove all managed ROCm vars from os.environ without writing to disk."""
cleared = 0
for var in ROCM_ENV_VARS:
if var in os.environ:
del os.environ[var]
cleared += 1
for var in _UNSET_VARS:
if var in os.environ:
del os.environ[var]
log.info(f'ROCm clear_env: cleared={cleared}')
def delete_config() -> None:
"""Delete the saved config file and clear all vars from the environment."""
global _cache # pylint: disable=global-statement
clear_env()
if CONFIG.exists():
CONFIG.unlink()
log.info(f'ROCm delete_config: deleted {CONFIG}')
_cache = None
def apply_profile(name: str) -> None:
"""Merge an architecture profile on top of the current config, then save and apply."""
profile = rocm_profiles.PROFILES.get(name)
if profile is None:
log.warning(f'ROCm apply_profile: unknown profile={name}')
return
config = load_config().copy()
config.update(profile)
save_config(config)
apply_env(config)
log.info(f'ROCm apply_profile: profile={name} overrides={len(profile)}')
def _hip_version_from_file(db_path: Path) -> str:
"""Parse HIP_VERSION_* keys from .hipVersion in the SDK bin folder."""
hip_ver_file = db_path / ".hipVersion"
if not hip_ver_file.exists():
return ""
kv = {}
for line in hip_ver_file.read_text(errors="ignore").splitlines():
if "=" in line and not line.startswith("#"):
k, _, v = line.partition("=")
kv[k.strip()] = v.strip()
major = kv.get("HIP_VERSION_MAJOR", "")
minor = kv.get("HIP_VERSION_MINOR", "")
patch = kv.get("HIP_VERSION_PATCH", "")
git = kv.get("HIP_VERSION_GITHASH", "")
if major:
return f"{major}.{minor}.{patch} ({git})"
return ""
def _pkg_version(name: str) -> str:
try:
import importlib.metadata as _m # pylint: disable=import-outside-toplevel
return _m.version(name)
except Exception:
return "n/a"
def _db_file_summary(path: Path, patterns: list) -> dict:
"""Return {filename: 'N KB'} for files matching any of the given glob patterns."""
out = {}
for pat in patterns:
for f in sorted(path.glob(pat)):
kb = f.stat().st_size // 1024
out[f.name] = f"{kb} KB"
return out
def _user_db_summary(path: Path) -> dict:
"""Return {filename: 'N KB, M entries'} for user MIOpen DB txt files."""
out = {}
for pat in ("*.udb.txt", "*.ufdb.txt"):
for f in sorted(path.glob(pat)):
kb = f.stat().st_size // 1024
try:
lines = sum(1 for _ in f.open(errors="ignore"))
except Exception:
lines = 0
out[f.name] = f"{kb} KB, {lines} entries"
return out
def info() -> dict:
config = load_config()
db_path = Path(_expand_venv(config.get("MIOPEN_SYSTEM_DB_PATH", "")))
# --- ROCm / HIP package versions ---
rocm_pkgs = {}
for pkg in ("rocm", "rocm-sdk-core", "rocm-sdk-devel"):
v = _pkg_version(pkg)
if v != "n/a":
rocm_pkgs[pkg] = v
libs_pkg = _pkg_version("rocm-sdk-libraries-gfx103x-dgpu")
if libs_pkg != "n/a":
rocm_pkgs["rocm-sdk-libraries (gfx103x)"] = libs_pkg
hip_ver = _hip_version_from_file(db_path)
if not hip_ver:
try:
import torch # pylint: disable=import-outside-toplevel
hip_ver = getattr(torch.version, "hip", "") or ""
except Exception:
pass
rocm_section = {}
if hip_ver:
rocm_section["hip_version"] = hip_ver
rocm_section.update(rocm_pkgs)
# --- Torch ---
torch_section = {}
try:
import torch # pylint: disable=import-outside-toplevel
torch_section["version"] = torch.__version__
torch_section["hip"] = getattr(torch.version, "hip", None) or "n/a"
except Exception:
pass
# --- GPU ---
gpu_section = [dict(g) for g in installer.gpu_info]
# --- System DB ---
sdb = {"path": str(db_path)}
if db_path.exists():
solver_db = _db_file_summary(db_path, ["*.db.txt"])
find_db = _db_file_summary(db_path, ["*.HIP.fdb.txt", "*.fdb.txt"])
kernel_db = _db_file_summary(db_path, ["*.kdb"])
if solver_db:
sdb["solver_db"] = solver_db
if find_db:
sdb["find_db"] = find_db
if kernel_db:
sdb["kernel_db"] = kernel_db
else:
sdb["exists"] = False
# --- User DB (~/.miopen/db) ---
user_db_path = Path.home() / ".miopen" / "db"
udb = {"path": str(user_db_path), "exists": user_db_path.exists()}
if user_db_path.exists():
ufiles = _user_db_summary(user_db_path)
if ufiles:
udb["files"] = ufiles
return {
"rocm": rocm_section,
"torch": torch_section,
"gpu": gpu_section,
"system_db": sdb,
"user_db": udb,
}
# Apply saved config to os.environ at import time (only when ROCm is present)
if is_rocm:
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')
scripts/rocm/rocm_profiles.py
+218
View File
@@ -0,0 +1,218 @@
"""
Architecture-specific MIOpen solver profiles for AMD GCN/RDNA GPUs.
Sources:
https://rocm.docs.amd.com/projects/MIOpen/en/develop/reference/env_variables.html
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
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).
"""
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",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_BWD_GTC_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_GTC_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS_NHWC": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_BWD_GTC_XDLOPS_NHWC": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_GTC_XDLOPS_NHWC": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_DLOPS_NCHWC": "0",
# HIP XDLOPS variants (CDNA-only)
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R5_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4_PADDED_GEMM_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4_PADDED_GEMM_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_XDLOPS": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_XDLOPS_EMULATE": "0",
"MIOPEN_DEBUG_IMPLICIT_GEMM_XDLOPS_INLINE_ASM": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_GROUP_BWD_XDLOPS": "0",
"MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_BWD_XDLOPS_AI_HEUR": "0",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_FWD_V4R4_XDLOPS_ADD_VECTOR_LOAD_GEMMN_TUNE_PARAM": "0",
# 3D XDLOPS (CDNA-only; no 3D conv XDLOPS on consumer RDNA)
"MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS": "0",
"MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_BWD_XDLOPS": "0",
"MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_WRW_XDLOPS": "0",
# Composable Kernel (requires XDLOPS / CDNA)
"MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_V6R1_DLOPS_NCHW": "0",
"MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_BIAS_ACTIV": "0",
"MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_BIAS_RES_ADD_ACTIV": "0",
# 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)
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F2X3": "0",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F3X3": "0",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F4X3": "0",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F5X3": "0",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F6X3": "0",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F2X3": "0",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F3X3": "0",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F4X3": "0",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F5X3": "0",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F6X3": "0",
}
# ---------------------------------------------------------------------------
# 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
"MIOPEN_DEBUG_CONV_FFT": "1",
"MIOPEN_DEBUG_CONV_DIRECT": "1",
"MIOPEN_DEBUG_CONV_GEMM": "1",
"MIOPEN_DEBUG_CONV_WINOGRAD": "1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM": "1",
"MIOPEN_DEBUG_CONV_IMMED_FALLBACK": "1",
"MIOPEN_DEBUG_ENABLE_AI_IMMED_MODE_FALLBACK": "1",
"MIOPEN_DEBUG_FORCE_IMMED_MODE_FALLBACK": "0",
# Kernel backends
"MIOPEN_DEBUG_GCN_ASM_KERNELS": "1",
"MIOPEN_DEBUG_HIP_KERNELS": "1",
"MIOPEN_DEBUG_OPENCL_CONVOLUTIONS": "1",
"MIOPEN_DEBUG_OPENCL_WAVE64_NOWGP": "1",
"MIOPEN_DEBUG_ATTN_SOFTMAX": "1",
# Direct ASM — all work on RDNA2
"MIOPEN_DEBUG_CONV_DIRECT_ASM_3X3U": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1UV2": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_5X10U2V2": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_7X7C3H224W224": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW3X3": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW1X1": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_PERF_VALS": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_SEARCH_OPTIMIZED": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_AI_HEUR": "1",
"MIOPEN_DEBUG_CONV_DIRECT_NAIVE_CONV_FWD": "1",
# Direct OCL — all work on RDNA2
"MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD": "1",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1": "1",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD11X11": "1",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_FWDGEN": "1",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW2": "1",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW53": "1",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW1X1": "1",
# Winograd RxS — all base variants work on RDNA2
"MIOPEN_DEBUG_AMD_WINOGRAD_3X3": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_FWD_BWD": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_WRW": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F3X2": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_G1": "1",
"MIOPEN_DEBUG_AMD_FUSED_WINOGRAD": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_PERF_VALS": "1",
# 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
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X2": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X3": "1",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X4": "0",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X5": "0",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X6": "0",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F5X3": "0",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F5X4": "0",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X2": "0",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X3": "0",
# ASM Implicit GEMM — V4R1 only; no GTC/XDLOPS on RDNA2
"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": "1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_V4R1": "1",
# HIP Implicit GEMM — non-XDLOPS V4R1/R4 only
"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": "1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1": "1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R1": "1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4": "1",
}
# ---------------------------------------------------------------------------
# RDNA3 — gfx1100 (RX 7000 series)
# Fury Winograd added; MPASS F3x4 enabled
# ---------------------------------------------------------------------------
RDNA3: Dict[str, str] = {
**RDNA2,
# 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",
}
# ---------------------------------------------------------------------------
# RDNA4 — gfx1200 (RX 9000 series)
# Rage Winograd added; MPASS F3x5 enabled
# ---------------------------------------------------------------------------
RDNA4: Dict[str, str] = {
**RDNA3,
# 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",
}
PROFILES: Dict[str, Dict[str, str]] = {
"RDNA2": RDNA2,
"RDNA3": RDNA3,
"RDNA4": RDNA4,
}
# Vars that are architecturally unavailable (no supporting hardware) per arch.
# These will be visually marked in the UI with strikethrough.
_UNAVAILABLE_ALL_RDNA = set(_XDLOPS_OFF.keys())
UNAVAILABLE: Dict[str, set] = {
"RDNA2": _UNAVAILABLE_ALL_RDNA | {
"MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F2X3",
"MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F3X2",
"MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X4",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X5",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X6",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F5X3",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F5X4",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X2",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X3",
},
"RDNA3": _UNAVAILABLE_ALL_RDNA | {
"MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X5",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X6",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F5X3",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F5X4",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X2",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X3",
},
"RDNA4": _UNAVAILABLE_ALL_RDNA | {
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X6",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F5X3",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F5X4",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X2",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X3",
},
}
scripts/rocm/rocm_vars.py
+374
View File
@@ -0,0 +1,374 @@
from typing import Dict, Any, List, Tuple
# --- General MIOpen/rocBLAS variables (dropdown/textbox/checkbox) ---
GENERAL_VARS: Dict[str, Dict[str, Any]] = {
"MIOPEN_GEMM_ENFORCE_BACKEND": {
"default": "1",
"desc": "Enforce GEMM backend",
"widget": "dropdown",
"options": [("1 - rocBLAS", "1"), ("5 - hipBLASLt", "5")],
"restart_required": False,
},
"MIOPEN_FIND_MODE": {
"default": "2",
"desc": "MIOpen Find Mode",
"widget": "dropdown",
"options": [("1 - NORMAL", "1"), ("2 - FAST", "2"), ("3 - HYBRID", "3"), ("5 - DYNAMIC_HYBRID", "5"), ("6 - TRUST_VERIFY", "6"), ("7 - TRUST_VERIFY_FULL", "7")],
"restart_required": True,
},
"MIOPEN_FIND_ENFORCE": {
"default": "1",
"desc": "MIOpen Find Enforce",
"widget": "dropdown",
"options": [("1 - NONE", "1"), ("2 - DB_UPDATE", "2"), ("3 - SEARCH", "3"), ("4 - SEARCH_DB_UPDATE", "4"), ("5 - DB_CLEAN", "5")],
"restart_required": True,
},
"MIOPEN_SEARCH_CUTOFF": {
"default": "0",
"desc": "Enable early termination of suboptimal searches",
"widget": "dropdown",
"options": [("0 - Off", "0"), ("1 - On", "1")],
"restart_required": True,
},
"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,
},
"MIOPEN_LOG_LEVEL": {
"default": "0",
"desc": "MIOpen log verbosity level",
"widget": "dropdown",
"options": [("0 - Default", "0"), ("1 - Quiet", "1"), ("3 - Error", "3"), ("4 - Warning", "4"), ("5 - Info", "5"), ("6 - Detail", "6"), ("7 - Trace", "7")],
"restart_required": False,
},
"MIOPEN_DEBUG_ENABLE": {
"default": "0",
"desc": "Enable MIOpen logging",
"widget": "dropdown",
"options": [("0 - Off", "0"), ("1 - On", "1")],
"restart_required": False,
},
"ROCBLAS_LAYER": {
"default": "0",
"desc": "rocBLAS logging",
"widget": "dropdown",
"options": [("0 - Off", "0"), ("1 - Trace", "1"), ("2 - Bench", "2"), ("3 - Trace+Bench", "3"), ("4 - Profile", "4"), ("5 - Trace+Profile", "5"), ("6 - Bench+Profile", "6"), ("7 - All", "7")],
"restart_required": False,
},
"HIPBLASLT_LOG_LEVEL": {
"default": "0",
"desc": "hipBLASLt logging",
"widget": "dropdown",
"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: name -> description string ---
# These are all checkbox widgets; default "1" unless in _SOLVER_DISABLED_BY_DEFAULT.
_SOLVER_DESCS: Dict[str, str] = {}
_SOLVER_DESCS.update({
"MIOPEN_DEBUG_CONV_FFT": "Enable FFT solver",
"MIOPEN_DEBUG_CONV_DIRECT": "Enable Direct solver",
"MIOPEN_DEBUG_CONV_GEMM": "Enable GEMM solver",
"MIOPEN_DEBUG_CONV_WINOGRAD": "Enable Winograd solver",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM": "Enable Implicit GEMM solver",
})
_SOLVER_DESCS.update({
"MIOPEN_DEBUG_CONV_IMMED_FALLBACK": "Enable Immediate Fallback",
"MIOPEN_DEBUG_ENABLE_AI_IMMED_MODE_FALLBACK": "Enable AI Immediate Mode Fallback",
"MIOPEN_DEBUG_FORCE_IMMED_MODE_FALLBACK": "Force Immediate Mode Fallback",
})
_SOLVER_DESCS.update({
"MIOPEN_DEBUG_GCN_ASM_KERNELS": "Enable GCN ASM kernels",
"MIOPEN_DEBUG_HIP_KERNELS": "Enable HIP kernels",
"MIOPEN_DEBUG_OPENCL_CONVOLUTIONS": "Enable OpenCL convolutions",
"MIOPEN_DEBUG_OPENCL_WAVE64_NOWGP": "Enable OpenCL Wave64 NOWGP",
"MIOPEN_DEBUG_ATTN_SOFTMAX": "Enable Attention Softmax",
})
_SOLVER_DESCS.update({
"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",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_5X10U2V2": "Enable Direct ASM 5x10U2V2",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_7X7C3H224W224": "Enable Direct ASM 7x7C3H224W224",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW3X3": "Enable Direct ASM WRW3X3",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW1X1": "Enable Direct ASM WRW1X1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_PERF_VALS": "Enable Direct ASM 1x1U Perf Vals",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_SEARCH_OPTIMIZED": "Enable Direct ASM 1x1U Search Optimized",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_AI_HEUR": "Enable Direct ASM 1x1U AI Heuristic",
"MIOPEN_DEBUG_CONV_DIRECT_NAIVE_CONV_FWD": "Enable Direct Naive Conv FWD",
})
_SOLVER_DESCS.update({
"MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD": "Enable Direct OCL FWD",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1": "Enable Direct OCL FWD1X1",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD11X11": "Enable Direct OCL FWD11X11",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_FWDGEN": "Enable Direct OCL FWDGEN",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW2": "Enable Direct OCL WRW2",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW53": "Enable Direct OCL WRW53",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW1X1": "Enable Direct OCL WRW1X1",
})
_SOLVER_DESCS.update({
"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/BWD",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_WRW": "Enable AMD Winograd RxS WRW",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F3X2": "Enable AMD Winograd RxS F3x2",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3": "Enable AMD Winograd RxS F2x3",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_G1": "Enable AMD Winograd RxS F2x3 G1",
"MIOPEN_DEBUG_AMD_FUSED_WINOGRAD": "Enable AMD Fused Winograd",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_PERF_VALS":"Enable AMD Winograd RxS F2x3 Perf Vals",
"MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F2X3": "Enable AMD Winograd Fury RxS F2x3",
"MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F3X2": "Enable AMD Winograd Fury RxS F3x2",
"MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3": "Enable AMD Winograd Rage RxS F2x3",
})
_SOLVER_DESCS.update({
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X2": "Enable AMD Winograd MPASS F3x2",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X3": "Enable AMD Winograd MPASS F3x3",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X4": "Enable AMD Winograd MPASS F3x4",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X5": "Enable AMD Winograd MPASS F3x5",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X6": "Enable AMD Winograd MPASS F3x6",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F5X3": "Enable AMD Winograd MPASS F5x3",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F5X4": "Enable AMD Winograd MPASS F5x4",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X2": "Enable AMD Winograd MPASS F7x2",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X3": "Enable AMD Winograd MPASS F7x3",
})
_SOLVER_DESCS.update({
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F2X3": "Enable AMD MP BD Winograd F2x3",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F3X3": "Enable AMD MP BD Winograd F3x3",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F4X3": "Enable AMD MP BD Winograd F4x3",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F5X3": "Enable AMD MP BD Winograd F5x3",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F6X3": "Enable AMD MP BD Winograd F6x3",
})
_SOLVER_DESCS.update({
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F2X3": "Enable AMD MP BD XDLOPS Winograd F2x3",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F3X3": "Enable AMD MP BD XDLOPS Winograd F3x3",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F4X3": "Enable AMD MP BD XDLOPS Winograd F4x3",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F5X3": "Enable AMD MP BD XDLOPS Winograd F5x3",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F6X3": "Enable AMD MP BD XDLOPS Winograd F6x3",
})
_SOLVER_DESCS.update({
"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_ASM_BWD_V4R1": "Enable ASM Implicit GEMM BWD V4R1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_V4R1": "Enable ASM Implicit GEMM WRW V4R1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS": "Enable ASM Implicit GEMM FWD GTC XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_BWD_GTC_XDLOPS": "Enable ASM Implicit GEMM BWD GTC XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_GTC_XDLOPS": "Enable ASM Implicit GEMM WRW GTC XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS_NHWC": "Enable ASM Implicit GEMM FWD GTC XDLOPS NHWC",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_BWD_GTC_XDLOPS_NHWC": "Enable ASM Implicit GEMM BWD GTC XDLOPS NHWC",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_GTC_XDLOPS_NHWC": "Enable ASM Implicit GEMM WRW GTC XDLOPS NHWC",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_DLOPS_NCHWC": "Enable ASM Implicit GEMM FWD GTC DLOPS NCHWC",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_GROUP_BWD_XDLOPS": "Enable HIP Group BWD XDLOPS",
"MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_BWD_XDLOPS_AI_HEUR": "Enable Group HIP BWD XDLOPS AI Heuristic",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_FWD_V4R4_XDLOPS_ADD_VECTOR_LOAD_GEMMN_TUNE_PARAM": "Enable FWD V4R4 XDLOPS Add Vector Load GEMMN Tune Param",
})
_SOLVER_DESCS.update({
"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",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1": "Enable HIP Implicit GEMM BWD V1R1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1": "Enable HIP Implicit GEMM BWD V4R1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R1": "Enable HIP Implicit GEMM WRW V4R1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4": "Enable HIP Implicit GEMM WRW V4R4",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4_XDLOPS": "Enable HIP Implicit GEMM FWD V4R4 XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R5_XDLOPS": "Enable HIP Implicit GEMM FWD V4R5 XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1_XDLOPS": "Enable HIP Implicit GEMM BWD V1R1 XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1_XDLOPS": "Enable HIP Implicit GEMM BWD V4R1 XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4_XDLOPS": "Enable HIP Implicit GEMM WRW V4R4 XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4_PADDED_GEMM_XDLOPS": "Enable HIP Implicit GEMM FWD V4R4 Padded GEMM XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4_PADDED_GEMM_XDLOPS": "Enable HIP Implicit GEMM WRW V4R4 Padded GEMM XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS": "Enable HIP Implicit GEMM FWD XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_XDLOPS": "Enable HIP Implicit GEMM BWD XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_XDLOPS": "Enable HIP Implicit GEMM WRW XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_XDLOPS": "Enable Implicit GEMM XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_XDLOPS_EMULATE": "Enable Implicit GEMM XDLOPS Emulation",
"MIOPEN_DEBUG_IMPLICIT_GEMM_XDLOPS_INLINE_ASM": "Enable Implicit GEMM XDLOPS Inline ASM",
})
_SOLVER_DESCS.update({
"MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS": "Enable 3D HIP Implicit GEMM FWD XDLOPS",
"MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_BWD_XDLOPS": "Enable 3D HIP Implicit GEMM BWD XDLOPS",
"MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_WRW_XDLOPS": "Enable 3D HIP Implicit GEMM WRW XDLOPS",
})
_SOLVER_DESCS.update({
"MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_V6R1_DLOPS_NCHW": "Enable CK iGEMM FWD V6R1 DLOPS NCHW",
"MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_BIAS_ACTIV": "Enable CK iGEMM FWD Bias Activation",
"MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_BIAS_RES_ADD_ACTIV": "Enable CK iGEMM FWD Bias Res Add Activation",
})
_SOLVER_DESCS.update({
"MIOPEN_DEBUG_CONV_MLIR_IGEMM_WRW_XDLOPS": "Enable MLIR iGEMM WRW XDLOPS",
"MIOPEN_DEBUG_CONV_MLIR_IGEMM_BWD_XDLOPS": "Enable MLIR iGEMM BWD XDLOPS",
})
# Vars whose solver default is "0" (disabled) — CDNA-only, RDNA3/4-only, experimental, or not used in SD inference
_SOLVER_DISABLED_BY_DEFAULT = {
"MIOPEN_DEBUG_FORCE_IMMED_MODE_FALLBACK",
"MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F2X3",
"MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F3X2",
"MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X4",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X5",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X6",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F5X3",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F5X4",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X2",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X3",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F2X3",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F3X3",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F4X3",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F5X3",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F6X3",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F2X3",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F3X3",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F4X3",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F5X3",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F6X3",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_BWD_GTC_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_GTC_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS_NHWC",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_BWD_GTC_XDLOPS_NHWC",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_GTC_XDLOPS_NHWC",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_DLOPS_NCHWC",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_GROUP_BWD_XDLOPS",
"MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_BWD_XDLOPS_AI_HEUR",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_FWD_V4R4_XDLOPS_ADD_VECTOR_LOAD_GEMMN_TUNE_PARAM",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R5_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4_PADDED_GEMM_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4_PADDED_GEMM_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_XDLOPS_EMULATE",
"MIOPEN_DEBUG_IMPLICIT_GEMM_XDLOPS_INLINE_ASM",
"MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS",
"MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_BWD_XDLOPS",
"MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_WRW_XDLOPS",
"MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_V6R1_DLOPS_NCHW",
"MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_BIAS_ACTIV",
"MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_BIAS_RES_ADD_ACTIV",
"MIOPEN_DEBUG_CONV_MLIR_IGEMM_WRW_XDLOPS",
"MIOPEN_DEBUG_CONV_MLIR_IGEMM_BWD_XDLOPS",
}
# Build full merged var registry
ROCM_ENV_VARS: Dict[str, Dict[str, Any]] = {}
ROCM_ENV_VARS.update(GENERAL_VARS)
for _var, _desc in _SOLVER_DESCS.items():
ROCM_ENV_VARS[_var] = {
"default": "0" if _var in _SOLVER_DISABLED_BY_DEFAULT else "1",
"desc": _desc,
"widget": "checkbox",
"options": None,
"restart_required": False,
}
# UI group ordering for solver sections
SOLVER_GROUPS: List[Tuple[str, List[str]]] = [
("Algorithm/Solver Group Enables", [
"MIOPEN_DEBUG_CONV_FFT", "MIOPEN_DEBUG_CONV_DIRECT", "MIOPEN_DEBUG_CONV_GEMM",
"MIOPEN_DEBUG_CONV_WINOGRAD", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM",
]),
("Immediate Fallback Mode", [
"MIOPEN_DEBUG_CONV_IMMED_FALLBACK", "MIOPEN_DEBUG_ENABLE_AI_IMMED_MODE_FALLBACK",
"MIOPEN_DEBUG_FORCE_IMMED_MODE_FALLBACK",
]),
("Build Method Toggles", [
"MIOPEN_DEBUG_GCN_ASM_KERNELS", "MIOPEN_DEBUG_HIP_KERNELS",
"MIOPEN_DEBUG_OPENCL_CONVOLUTIONS", "MIOPEN_DEBUG_OPENCL_WAVE64_NOWGP",
"MIOPEN_DEBUG_ATTN_SOFTMAX",
]),
("Direct ASM Solver Toggles", [
"MIOPEN_DEBUG_CONV_DIRECT_ASM_3X3U", "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1UV2", "MIOPEN_DEBUG_CONV_DIRECT_ASM_5X10U2V2",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_7X7C3H224W224", "MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW3X3",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW1X1", "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_PERF_VALS",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_SEARCH_OPTIMIZED", "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_AI_HEUR",
"MIOPEN_DEBUG_CONV_DIRECT_NAIVE_CONV_FWD",
]),
("Direct OpenCL Solver Toggles", [
"MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD", "MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD11X11", "MIOPEN_DEBUG_CONV_DIRECT_OCL_FWDGEN",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW2", "MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW53",
"MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW1X1",
]),
("Winograd Solver Toggles", [
"MIOPEN_DEBUG_AMD_WINOGRAD_3X3", "MIOPEN_DEBUG_AMD_WINOGRAD_RXS",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_FWD_BWD", "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_WRW",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F3X2", "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_G1", "MIOPEN_DEBUG_AMD_FUSED_WINOGRAD",
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_PERF_VALS", "MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F2X3",
"MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F3X2", "MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3",
]),
("Multi-pass Winograd Toggles", [
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X2", "MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X3",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X4", "MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X5",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X6", "MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F5X3",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F5X4", "MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X2",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X3",
]),
("MP BD Winograd Toggles", [
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F2X3", "MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F3X3",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F4X3", "MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F5X3",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F6X3",
]),
("MP BD XDLOPS Winograd Toggles", [
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F2X3", "MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F3X3",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F4X3", "MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F5X3",
"MIOPEN_DEBUG_AMD_MP_BD_XDLOPS_WINOGRAD_F6X3",
]),
("ASM Implicit GEMM Toggles", [
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_V4R1", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_V4R1_1X1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_BWD_V4R1", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_V4R1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_BWD_GTC_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_GTC_XDLOPS", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS_NHWC",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_BWD_GTC_XDLOPS_NHWC", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_GTC_XDLOPS_NHWC",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_DLOPS_NCHWC",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_GROUP_BWD_XDLOPS",
"MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_BWD_XDLOPS_AI_HEUR",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_FWD_V4R4_XDLOPS_ADD_VECTOR_LOAD_GEMMN_TUNE_PARAM",
]),
("HIP Implicit GEMM Toggles", [
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R1", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4_XDLOPS", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R5_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1_XDLOPS", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4_PADDED_GEMM_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4_PADDED_GEMM_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_XDLOPS", "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_XDLOPS",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_XDLOPS_EMULATE", "MIOPEN_DEBUG_IMPLICIT_GEMM_XDLOPS_INLINE_ASM",
]),
("3D Implicit GEMM Toggles", [
"MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS",
"MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_BWD_XDLOPS",
"MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_WRW_XDLOPS",
]),
("CK Implicit GEMM Toggles", [
"MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_V6R1_DLOPS_NCHW",
"MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_BIAS_ACTIV",
"MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_BIAS_RES_ADD_ACTIV",
]),
("MLIR Implicit GEMM Toggles", [
"MIOPEN_DEBUG_CONV_MLIR_IGEMM_WRW_XDLOPS",
"MIOPEN_DEBUG_CONV_MLIR_IGEMM_BWD_XDLOPS",
]),
]
scripts/rocm_ext.py
+217 -26
View File
@@ -1,6 +1,14 @@
import gradio as gr
import installer
from modules import scripts_manager, shared
from modules import scripts_manager
def _is_rocm() -> bool:
try:
from scripts.rocm import rocm_mgr # pylint: disable=no-name-in-module
return rocm_mgr.is_rocm
except Exception:
return False
class Script(scripts_manager.Script):
@@ -8,35 +16,218 @@ class Script(scripts_manager.Script):
return "ROCm: Advanced Config"
def show(self, _is_img2img):
if shared.cmd_opts.use_rocm or installer.torch_info.get('type') == 'rocm':
return scripts_manager.AlwaysVisible # script should be visible only if rocm is detected or forced
if _is_rocm():
return scripts_manager.AlwaysVisible
return False
def ui(self, _is_img2img):
# all ui elements go here
# logic should go into scripts/rocm/rocm_mgr.py and be called from here
if not shared.cmd_opts.use_rocm and not installer.torch_info.get('type') == 'rocm': # skip ui creation if not rocm
from scripts.rocm import rocm_mgr, rocm_vars # pylint: disable=no-name-in-module
if not rocm_mgr.is_rocm:
with gr.Accordion('ROCm: Advanced Config', open=False, elem_id='rocm_config'):
gr.HTML("<p><b>ROCm is not installed.</b> This section is disabled.</p>")
return []
from scripts.rocm import rocm_mgr # pylint: disable=no-name-in-module
rocm_mgr.load() # load config before ui creation so we can populate fields with current values
with gr.Accordion('ROCM: Advanced Config', open=False, elem_id='rocm_config'): # set all ui in this accordion
with gr.Row():
gr.HTML('Advanced configuration for ROCm users')
with gr.Row():
btn_info = gr.Button("Refresh") # refresh info and populate with additional fields
btn_apply = gr.Button("Apply") # apply changes and save config
btn_reset = gr.Button("Reset") # reset config and save config
with gr.Row():
info = gr.JSON(installer.gpu_info, label="ROCm Info") # just an example how to display info we already have
with gr.Row():
db_path = gr.Textbox(label="MiOpen SystemDB Path", value=rocm_mgr.DATA.get('MIOPEN_SYSTEM_DB_PATH', ''), lines=1)
args = [ # list all params here and note that it *must* match with scripts/rocm/rocm_mgr.apply() args
db_path,
]
config = rocm_mgr.load_config()
var_names = []
components = []
btn_info.click(fn=rocm_mgr.info, inputs=[], outputs=[info])
btn_apply.click(fn=rocm_mgr.apply, inputs=args, outputs=[])
btn_reset.click(fn=rocm_mgr.reset, inputs=[], outputs=[])
def _make_component(name, meta, cfg):
val = cfg.get(name, meta["default"])
if meta["widget"] == "checkbox":
return gr.Checkbox(label=meta["desc"], value=(val == "1"), elem_id=f"rocm_var_{name.lower()}")
elif meta["widget"] == "dropdown":
choices = rocm_mgr._dropdown_choices(meta["options"])
display = rocm_mgr._dropdown_display(val, meta["options"])
if display not in choices: # corrupted stored value — fall back to default
display = rocm_mgr._dropdown_display(meta["default"], meta["options"])
return gr.Dropdown(label=meta["desc"], choices=choices, value=display, elem_id=f"rocm_var_{name.lower()}")
else: # textbox
return gr.Textbox(label=meta["desc"], value=rocm_mgr._expand_venv(val), lines=1)
return args
def _info_html():
d = rocm_mgr.info()
ts = "<style>#rocm_info_table table{width:100%;border-collapse:collapse;font-size:12px} #rocm_info_table td,#rocm_info_table th{padding:2px 8px;border-bottom:1px solid var(--sd-panel-border-color)} #rocm_info_table th{text-align:left;color:var(--sd-muted-color);font-weight:normal} #rocm_info_table td:first-child{color:var(--sd-muted-color);width:38%}</style>"
rows = []
def section(title):
rows.append(f"<tr><th colspan='2' style='padding-top:6px;color:var(--highlight-color)'>{title}</th></tr>")
def row(k, v):
rows.append(f"<tr><td>{k}</td><td>{v}</td></tr>")
section("ROCm / HIP")
for k, v in d.get("rocm", {}).items():
row(k, v)
section("System DB")
sdb = d.get("system_db", {})
row("path", sdb.get("path", ""))
for sub in ("solver_db", "find_db", "kernel_db"):
for fname, sz in sdb.get(sub, {}).items():
row(sub.replace("_", " "), f"{fname} &nbsp; {sz}")
section("User DB (~/.miopen/db)")
udb = d.get("user_db", {})
row("path", udb.get("path", ""))
for fname, finfo in udb.get("files", {}).items():
row(fname, finfo)
return ts + f"<table>{''.join(rows)}</table>"
with gr.Accordion('ROCm: Advanced Config', open=False, elem_id='rocm_config'):
gr.HTML("""<style>
#rocm_btn_apply button { background: var(--button-primary-background-fill) !important; color: var(--button-primary-text-color) !important; border-color: var(--button-primary-border-color) !important; }
#rocm_btn_apply button:hover { background: var(--button-primary-background-fill-hover) !important; border-color: var(--button-primary-border-color) !important; }
#rocm_btn_delete button { color: var(--color-error) !important; border-color: var(--color-error) !important; background: transparent !important; }
#rocm_btn_delete button:hover { background: color-mix(in srgb, var(--color-error) 12%, transparent) !important; }
/* Strip all block/group backgrounds and borders inside rocm_config */
#rocm_config .block,
#rocm_config .form,
#rocm_config .gradio-group {
background: transparent !important;
box-shadow: none !important;
border: none !important;
padding: 0 !important;
margin: 0 !important;
gap: 0 !important;
}
#rocm_config fieldset { border: none !important; padding: 0 !important; margin: 0 !important; box-shadow: none !important; background: transparent !important; }
</style>""")
with gr.Row():
gr.HTML("<p>Advanced configuration for ROCm users.</p><br><p>Set Your database and solver selections based on GPU profile or individually.</p><br><p>Enable cuDNN in Backend Settings to activate MIOpen.</p>")
with gr.Row():
btn_info = gr.Button("Refresh Info", variant="primary", elem_id="rocm_btn_info")
btn_apply = gr.Button("Apply", elem_id="rocm_btn_apply")
btn_reset = gr.Button("Defaults")
btn_clear = gr.Button("Clear Runtime")
btn_delete = gr.Button("Delete", elem_id="rocm_btn_delete")
with gr.Row():
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("")
info_out = gr.HTML(value=_info_html, elem_id="rocm_info_table")
# General vars (dropdowns, textboxes, checkboxes)
with gr.Group():
gr.HTML("<h3>MIOpen Settings</h3><hr>")
for name, meta in rocm_vars.GENERAL_VARS.items():
comp = _make_component(name, meta, config)
var_names.append(name)
components.append(comp)
# Solver groups (all checkboxes, grouped by section)
for group_name, varlist in rocm_vars.SOLVER_GROUPS:
with gr.Group():
gr.HTML(f"<h3>{group_name}</h3><hr>")
for name in varlist:
meta = rocm_vars.ROCM_ENV_VARS[name]
comp = _make_component(name, meta, config)
var_names.append(name)
components.append(comp)
gr.HTML("<br><center><div style='margin:0 Auto'><a href='https://rocm.docs.amd.com/projects/MIOpen/en/develop/reference/env_variables.html' target='_blank'>&#128196; MIOpen Environment Variables Reference</a></div></center><br>")
def _autosave_dropdown(name, value):
meta = rocm_vars.ROCM_ENV_VARS[name]
if meta["widget"] == "dropdown":
stored = rocm_mgr._dropdown_stored(str(value), meta["options"])
config = rocm_mgr.load_config()
config[name] = stored
rocm_mgr.save_config(config)
rocm_mgr.apply_env(config)
for name, comp in zip(var_names, components):
meta = rocm_vars.ROCM_ENV_VARS[name]
if meta["widget"] == "dropdown":
comp.change(fn=lambda v, n=name: _autosave_dropdown(n, v), inputs=[comp], outputs=[])
def apply_fn(*values):
rocm_mgr.apply_all(var_names, list(values))
saved = rocm_mgr.load_config()
result = [gr.update(value="")]
for name in var_names:
meta = rocm_vars.ROCM_ENV_VARS[name]
val = saved.get(name, meta["default"])
if meta["widget"] == "checkbox":
result.append(gr.update(value=(val == "1")))
elif meta["widget"] == "dropdown":
result.append(gr.update(value=rocm_mgr._dropdown_display(val, meta["options"])))
else:
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"<style>{rules}</style>"
def reset_fn():
rocm_mgr.reset_defaults()
updated = rocm_mgr.load_config()
result = [gr.update(value="")]
for name in var_names:
meta = rocm_vars.ROCM_ENV_VARS[name]
val = updated.get(name, meta["default"])
if meta["widget"] == "checkbox":
result.append(gr.update(value=(val == "1")))
elif meta["widget"] == "radio":
result.append(gr.update(value=rocm_mgr._dropdown_display(val, meta["options"])))
else:
result.append(gr.update(value=rocm_mgr._expand_venv(val)))
return result
def clear_fn():
rocm_mgr.clear_env()
result = [gr.update(value="")]
for name in var_names:
meta = rocm_vars.ROCM_ENV_VARS[name]
if meta["widget"] == "checkbox":
result.append(gr.update(value=False))
elif meta["widget"] == "radio":
choices = rocm_mgr._dropdown_choices(meta["options"])
result.append(gr.update(value=choices[0] if choices else None))
else:
result.append(gr.update(value=""))
return result
def delete_fn():
rocm_mgr.delete_config()
result = [gr.update(value="")]
for name in var_names:
meta = rocm_vars.ROCM_ENV_VARS[name]
if meta["widget"] == "checkbox":
result.append(gr.update(value=False))
elif meta["widget"] == "radio":
choices = rocm_mgr._dropdown_choices(meta["options"])
result.append(gr.update(value=choices[0] if choices else None))
else:
result.append(gr.update(value=""))
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))]
for pname in var_names:
meta = rocm_vars.ROCM_ENV_VARS[pname]
val = updated.get(pname, meta["default"])
if meta["widget"] == "checkbox":
result.append(gr.update(value=(val == "1")))
elif meta["widget"] == "radio":
result.append(gr.update(value=rocm_mgr._dropdown_display(val, meta["options"])))
else:
result.append(gr.update(value=rocm_mgr._expand_venv(val)))
return result
btn_info.click(fn=_info_html, inputs=[], outputs=[info_out])
btn_apply.click(fn=apply_fn, inputs=components, outputs=[style_out] + components)
btn_reset.click(fn=reset_fn, inputs=[], outputs=[style_out] + components)
btn_clear.click(fn=clear_fn, inputs=[], outputs=[style_out] + components)
btn_delete.click(fn=delete_fn, inputs=[], outputs=[style_out] + components)
btn_rdna2.click(fn=lambda: profile_fn("RDNA2"), inputs=[], outputs=[style_out] + components)
btn_rdna3.click(fn=lambda: profile_fn("RDNA3"), inputs=[], outputs=[style_out] + components)
btn_rdna4.click(fn=lambda: profile_fn("RDNA4"), inputs=[], outputs=[style_out] + components)
return components