diff --git a/data/rocm-config.json b/data/rocm-config.json new file mode 100644 index 000000000..30bf96ea9 --- /dev/null +++ b/data/rocm-config.json @@ -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" +} \ No newline at end of file diff --git a/scripts/rocm/rocm_mgr.py b/scripts/rocm/rocm_mgr.py index 507a73018..9b98d5405 100644 --- a/scripts/rocm/rocm_mgr.py +++ b/scripts/rocm/rocm_mgr.py @@ -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') diff --git a/scripts/rocm/rocm_profiles.py b/scripts/rocm/rocm_profiles.py new file mode 100644 index 000000000..9e714d1f7 --- /dev/null +++ b/scripts/rocm/rocm_profiles.py @@ -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", + }, +} diff --git a/scripts/rocm/rocm_vars.py b/scripts/rocm/rocm_vars.py new file mode 100644 index 000000000..dc768f88e --- /dev/null +++ b/scripts/rocm/rocm_vars.py @@ -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", + ]), +] diff --git a/scripts/rocm_ext.py b/scripts/rocm_ext.py index 520aedb40..b899cc4db 100644 --- a/scripts/rocm_ext.py +++ b/scripts/rocm_ext.py @@ -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("

ROCm is not installed. This section is disabled.

") 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 = "" + rows = [] + def section(title): + rows.append(f"{title}") + def row(k, v): + rows.append(f"{k}{v}") + 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}   {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"{''.join(rows)}
" + + with gr.Accordion('ROCm: Advanced Config', open=False, elem_id='rocm_config'): + gr.HTML("""""") + 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.

") + 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("

MIOpen Settings

") + 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"

{group_name}

") + 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("
📄 MIOpen Environment Variables Reference

") + + 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"" + + 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