diff --git a/data/rocm-config.json b/data/rocm-config.json
deleted file mode 100644
index 30bf96ea9..000000000
--- a/data/rocm-config.json
+++ /dev/null
@@ -1,115 +0,0 @@
-{
- "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 9b98d5405..af07396a3 100644
--- a/scripts/rocm/rocm_mgr.py
+++ b/scripts/rocm/rocm_mgr.py
@@ -6,7 +6,7 @@ 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.rocm_vars import ROCM_ENV_VARS, SOLVER_GROUPS, SOLVER_DISABLED_BY_DEFAULT # pylint: disable=no-name-in-module
from scripts.rocm import rocm_profiles # pylint: disable=no-name-in-module
@@ -34,16 +34,96 @@ def _check_rocm() -> bool:
is_rocm = _check_rocm()
-CONFIG = Path(os.path.abspath(os.path.join('data', 'rocm-config.json')))
+CONFIG = Path(os.path.abspath(os.path.join('data', 'rocm.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 = {
+# Metadata key written into rocm.json to record which architecture profile is active.
+# Not an environment variable — always skipped during env application but preserved in the
+# saved config so that arch-safety enforcement is consistent across restarts.
+_ARCH_KEY = "_rocm_arch"
+
+# Vars that must never appear in the process environment.
+#
+# _DTYPE_UNSAFE: alter FP16 inference dtype — must be cleared regardless of config
+# MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL — DEBUG alias: routes all FP16 convs through BF16 exponent math
+# MIOPEN_CONVOLUTION_ATTRIB_FP16_ALT_IMPL — API-level alias: same BF16-exponent effect
+# MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_EXPEREMENTAL_FP16_TRANSFORM — unstable experimental FP16 path
+# MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_PK_ATOMIC_ADD_FP16 — changes FP16 WrW atomic accumulation
+#
+# SOLVER_DISABLED_BY_DEFAULT: every solver known to be incompatible with this runtime
+# (FP32-only, training-only WrW/BWD, fixed-geometry mismatches, XDLOPS/CDNA-only, arch-specific).
+# Actively unsetting these ensures no inherited shell value can re-enable them.
+_DTYPE_UNSAFE = {
"MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL",
+ "MIOPEN_CONVOLUTION_ATTRIB_FP16_ALT_IMPL",
"MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_EXPEREMENTAL_FP16_TRANSFORM",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_PK_ATOMIC_ADD_FP16",
}
+# _UNSET_VARS: hard-blocked vars that are DELETED from the process env and never written,
+# regardless of saved config. Limited to dtype-corrupting vars only.
+# IMPORTANT: SOLVER_DISABLED_BY_DEFAULT is intentionally NOT included here.
+# When a solver var is absent (unset) MIOpen still calls IsApplicable() on every
+# conv-find — wasted probing overhead. When a var is explicitly "0" MIOpen skips
+# IsApplicable() immediately. Solver defaults flow through the config loop as "0"
+# (their ROCM_ENV_VARS default is "0") so they are explicitly set to "0" in the env.
+_UNSET_VARS = _DTYPE_UNSAFE
+
+# Additional environment vars that must be removed from the process before MIOpen loads.
+# These are not MIOpen solver toggles but can corrupt MIOpen's runtime behaviour:
+# HIP_PATH / HIP_PATH_71 — point to the system AMD ROCm install; override the venv-bundled
+# _rocm_sdk_devel DLLs with a potentially mismatched system version
+# QML_*/QT_* — QtQuick shader/disk-cache flags leaked from Qt tools; harmless for
+# PyTorch but can conflict with Gradio's embedded Qt helpers
+# PYENV_VIRTUALENV_DISABLE_PROMPT — pyenv noise that confuses venv detection
+_EXTRA_CLEAR_VARS = {
+ "HIP_PATH",
+ "HIP_PATH_71",
+ "PYENV_VIRTUALENV_DISABLE_PROMPT",
+ "QML_DISABLE_DISK_CACHE",
+ "QML_FORCE_DISK_CACHE",
+ "QT_DISABLE_SHADER_DISK_CACHE",
+ # PERF_VALS vars are NOT boolean toggles — MIOpen reads them as perf-config strings.
+ # If inherited from a parent shell with value "1", MIOpen's GetPerfConfFromEnv parses
+ # "1" as a degenerate config and can return dtype=float32 output from FP16 tensors.
+ "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_PERF_VALS",
+ "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_PERF_VALS",
+}
+
+# Solvers whose MIOpen IsApplicable() explicitly rejects non-FP32 tensors.
+# They are safe to leave enabled in FP32 mode. When the active dtype is FP16 or BF16
+# we force them OFF so MIOpen skips the IsApplicable probe entirely — avoids overhead on
+# every conv shape find. These are NOT in _UNSET_VARS because they are valid in FP32.
+_FP32_ONLY_SOLVERS = {
+ "MIOPEN_DEBUG_CONV_FFT", # FFT convolution — FP32 only (MIOpen source: IsFp32 check)
+ "MIOPEN_DEBUG_AMD_WINOGRAD_3X3", # Winograd 3x3 — FP32 only
+ "MIOPEN_DEBUG_AMD_FUSED_WINOGRAD", # Fused Winograd — FP32 only
+}
+
+
+def _resolve_dtype() -> str:
+ """Return the resolved active compute dtype: 'FP16', 'BF16', 'FP32', or '' (not yet known).
+ Prefers the resolved devices.dtype (post test_fp16/bf16) over the raw opts string."""
+ try:
+ import torch # pylint: disable=import-outside-toplevel
+ from modules import devices as _dev # pylint: disable=import-outside-toplevel
+ if _dev.dtype is not None:
+ if _dev.dtype == torch.float16:
+ return 'FP16'
+ if _dev.dtype == torch.bfloat16:
+ return 'BF16'
+ if _dev.dtype == torch.float32:
+ return 'FP32'
+ except Exception:
+ pass
+ try:
+ from modules import shared as _sh # pylint: disable=import-outside-toplevel
+ v = getattr(getattr(_sh, 'opts', None), 'cuda_dtype', None)
+ if v in ('FP16', 'BF16', 'FP32'):
+ return v
+ except Exception:
+ pass
+ return ''
# --- venv helpers ---
@@ -88,47 +168,86 @@ def _dropdown_choices(options):
def load_config() -> Dict[str, str]:
global _cache # pylint: disable=global-statement
if _cache is None:
- if CONFIG.exists():
+ file_existed = CONFIG.exists()
+ if file_existed:
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()}
+ # Purge unsafe vars from a stale saved config and re-persist only if the file existed.
+ # When running without a saved config (first run / after Delete), load_config() must
+ # never create the file — that only happens via save_config() on Apply or Apply Profile.
+ dirty = {k for k in _cache if k in _UNSET_VARS or (k != _ARCH_KEY and k not in ROCM_ENV_VARS)}
+ if dirty:
+ _cache = {k: v for k, v in _cache.items() if k not in dirty}
+ writefile(_cache, str(CONFIG))
+ log.debug(f'ROCm load_config: purged {len(dirty)} stale/unsafe var(s) from saved config')
else:
_cache = {k: v["default"] for k, v in ROCM_ENV_VARS.items()}
- log.debug(f'ROCm load_config: path={CONFIG} items={len(_cache)}')
+ log.debug(f'ROCm load_config: path={CONFIG} existed={file_existed} 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
+ sanitized = {k: v for k, v in config.items() if k not in _UNSET_VARS}
+ # Enforce arch-incompatible solvers to "0" before writing.
+ # Prevents malformed edits (UI or JSON hand-edit) from persisting incompatible "1" values.
+ arch = sanitized.get(_ARCH_KEY, "")
+ unavailable = rocm_profiles.UNAVAILABLE.get(arch, set())
+ for var in unavailable:
+ if var in sanitized and sanitized[var] != "0":
+ sanitized[var] = "0"
+ log.debug(f'ROCm save_config: clamped arch-incompatible var={var} arch={arch}')
+ writefile(sanitized, str(CONFIG))
+ _cache = sanitized
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
+ if var == _ARCH_KEY:
continue
+ if var in _UNSET_VARS:
+ continue
+ if var not in ROCM_ENV_VARS:
+ continue
+ meta = ROCM_ENV_VARS.get(var, {})
+ if meta.get("options"):
+ value = _dropdown_stored(str(value), meta["options"])
expanded = _expand_venv(str(value))
if expanded == "":
- skipped += 1
continue
os.environ[var] = expanded
- applied += 1
+ # Arch safety net: hard-force all hardware-incompatible vars to "0" in the env.
+ # This runs *after* the config loop so it overrides any stale "1" that survived in the JSON.
+ # Source of truth: rocm_profiles.UNAVAILABLE[arch] — vars with no supporting hardware.
+ arch = config.get(_ARCH_KEY, "")
+ unavailable = rocm_profiles.UNAVAILABLE.get(arch, set())
+ if unavailable:
+ for var in unavailable:
+ os.environ[var] = "0"
+ dtype_str = _resolve_dtype()
+ if dtype_str in ('FP16', 'BF16'):
+ for var in _FP32_ONLY_SOLVERS:
+ os.environ[var] = "0"
def apply_all(names: list, values: list) -> None:
config = load_config().copy()
+ arch = config.get(_ARCH_KEY, "")
+ unavailable = rocm_profiles.UNAVAILABLE.get(arch, set())
for name, value in zip(names, values):
if name not in ROCM_ENV_VARS:
log.warning(f'ROCm apply_all: unknown variable={name}')
continue
+ # Arch safety net: silently clamp incompatible solvers back to "0".
+ # The UI may send the current checkbox state even for greyed-out vars.
+ if name in unavailable:
+ config[name] = "0"
+ continue
meta = ROCM_ENV_VARS[name]
if meta["widget"] == "checkbox":
if value is None:
@@ -142,6 +261,8 @@ def apply_all(names: list, values: list) -> None:
config[name] = stored
# else: value was None/invalid — leave the existing saved value untouched
else:
+ if meta.get("options"):
+ value = _dropdown_stored(str(value), meta["options"])
config[name] = _collapse_venv(str(value))
save_config(config)
apply_env(config)
@@ -149,32 +270,45 @@ def apply_all(names: list, values: list) -> None:
def reset_defaults() -> None:
defaults = {k: v["default"] for k, v in ROCM_ENV_VARS.items()}
+ # Preserve the active arch key so safety nets survive a defaults reset.
+ arch = load_config().get(_ARCH_KEY, "")
+ if arch:
+ defaults[_ARCH_KEY] = arch
save_config(defaults)
apply_env(defaults)
- log.info('ROCm reset_defaults: config reset to defaults')
+ log.info(f'ROCm reset_defaults: config reset to defaults arch={arch or "(none)"}')
def clear_env() -> None:
- """Remove all managed ROCm vars from os.environ without writing to disk."""
+ """Remove all managed ROCm vars and known noise 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:
+ for var in _UNSET_VARS | _EXTRA_CLEAR_VARS:
if var in os.environ:
del os.environ[var]
+ cleared += 1
log.info(f'ROCm clear_env: cleared={cleared}')
def delete_config() -> None:
- """Delete the saved config file and clear all vars from the environment."""
+ """Delete the saved config file, clear all vars, and wipe the MIOpen user DB cache."""
+ import shutil # pylint: disable=import-outside-toplevel
global _cache # pylint: disable=global-statement
clear_env()
if CONFIG.exists():
CONFIG.unlink()
log.info(f'ROCm delete_config: deleted {CONFIG}')
_cache = None
+ # Delete the MIOpen user DB (~/.miopen/db) — stale entries can cause solver mismatches
+ miopen_db = Path(os.path.expanduser('~')) / '.miopen' / 'db'
+ if miopen_db.exists():
+ shutil.rmtree(miopen_db, ignore_errors=True)
+ log.info(f'ROCm delete_config: wiped MIOpen user DB at {miopen_db}')
+ else:
+ log.debug(f'ROCm delete_config: MIOpen user DB not found at {miopen_db} — nothing to wipe')
def apply_profile(name: str) -> None:
@@ -185,6 +319,7 @@ def apply_profile(name: str) -> None:
return
config = load_config().copy()
config.update(profile)
+ config[_ARCH_KEY] = name # stamp the active arch so safety nets survive restarts
save_config(config)
apply_env(config)
log.info(f'ROCm apply_profile: profile={name} overrides={len(profile)}')
diff --git a/scripts/rocm/rocm_profiles.py b/scripts/rocm/rocm_profiles.py
index 9e714d1f7..3c9a52b7f 100644
--- a/scripts/rocm/rocm_profiles.py
+++ b/scripts/rocm/rocm_profiles.py
@@ -78,7 +78,7 @@ RDNA2: Dict[str, str] = {
# General settings (architecture-independent; set here so all profiles cover them)
"MIOPEN_SEARCH_CUTOFF": "0",
"MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC": "0",
- # Core algo enables
+ # Core algo enables — FFT is FP32-only but harmless (IsApplicable rejects it for fp16 tensors)
"MIOPEN_DEBUG_CONV_FFT": "1",
"MIOPEN_DEBUG_CONV_DIRECT": "1",
"MIOPEN_DEBUG_CONV_GEMM": "1",
@@ -93,36 +93,58 @@ RDNA2: Dict[str, str] = {
"MIOPEN_DEBUG_OPENCL_CONVOLUTIONS": "1",
"MIOPEN_DEBUG_OPENCL_WAVE64_NOWGP": "1",
"MIOPEN_DEBUG_ATTN_SOFTMAX": "1",
- # Direct ASM — all work on RDNA2
+ # Direct ASM — dtype notes
+ # 3X3U / 1X1U / 1X1UV2: FP32/FP16 forward — enabled
"MIOPEN_DEBUG_CONV_DIRECT_ASM_3X3U": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1UV2": "1",
- "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",
+ # 5X10U2V2: fixed geometry (5×10 stride-2), no SD conv matches — disabled
+ "MIOPEN_DEBUG_CONV_DIRECT_ASM_5X10U2V2": "0",
+ # 7X7C3H224W224: hard-coded ImageNet stem (C=3, H=W=224, K=64) — never matches SD — disabled
+ "MIOPEN_DEBUG_CONV_DIRECT_ASM_7X7C3H224W224": "0",
+ # WRW3X3 / WRW1X1: FP32-only weight-gradient (training only) — disabled for inference
+ "MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW3X3": "0",
+ "MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW1X1": "0",
+ # PERF_VALS intentionally blank: MIOpen reads this as a config string not a boolean;
+ # setting to "1" causes GetPerfConfFromEnv to use a degenerate config and return float32
+ "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_PERF_VALS": "",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_SEARCH_OPTIMIZED": "1",
"MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_AI_HEUR": "1",
- "MIOPEN_DEBUG_CONV_DIRECT_NAIVE_CONV_FWD": "1",
- # Direct OCL — all work on RDNA2
+ # NAIVE_CONV_FWD: scalar FP32 reference solver — IsApplicable does NOT reliably filter for FP16;
+ # can be selected for unusual shapes (e.g. VAE decoder 3-ch output) and returns dtype=float32
+ "MIOPEN_DEBUG_CONV_DIRECT_NAIVE_CONV_FWD": "0",
+ # Direct OCL — dtype notes
+ # FWD / FWD1X1: FP32/FP16 forward — enabled
"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
+ # FWD11X11: requires 11×11 kernel — no SD match — disabled
+ "MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD11X11": "0",
+ # FWDGEN: FP32 generic OCL fallback — IsApplicable does NOT reliably reject for FP16;
+ # can produce dtype=float32 output for FP16 inputs — disabled
+ "MIOPEN_DEBUG_CONV_DIRECT_OCL_FWDGEN": "0",
+ # WRW2 / WRW53 / WRW1X1: training-only weight-gradient — disabled
+ "MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW2": "0",
+ "MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW53": "0",
+ "MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW1X1": "0",
+ # Winograd RxS — dtype per MIOpen docs
+ # WINOGRAD_3X3: FP32-only — harmless (IsApplicable rejects for fp16); enabled
"MIOPEN_DEBUG_AMD_WINOGRAD_3X3": "1",
+ # RXS: covers FP32/FP16 F(3,3) Fwd/Bwd + FP32 F(3,2) WrW — keep enabled (fp16 fwd/bwd path exists)
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS": "1",
+ # RXS_FWD_BWD: FP32/FP16 — explicitly the fp16-capable subset
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_FWD_BWD": "1",
- "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_WRW": "1",
+ # RXS_WRW: FP32 WrW only — training-only, disabled for inference fp16 profile
+ "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_WRW": "0",
+ # RXS_F3X2: FP32/FP16 Fwd/Bwd
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F3X2": "1",
+ # RXS_F2X3: FP32/FP16 Fwd/Bwd (group convolutions)
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3": "1",
+ # RXS_F2X3_G1: FP32/FP16 Fwd/Bwd (non-group convolutions)
"MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_G1": "1",
+ # FUSED_WINOGRAD: FP32-only — harmless (IsApplicable rejects for fp16); enabled
"MIOPEN_DEBUG_AMD_FUSED_WINOGRAD": "1",
- "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_PERF_VALS": "1",
+ # PERF_VALS intentionally blank: same reason as ASM_1X1U — not a boolean, config string
+ "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_PERF_VALS": "",
# Fury/Rage Winograd — NOT available on RDNA2
"MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F2X3": "0",
"MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F3X2": "0",
@@ -137,18 +159,20 @@ RDNA2: Dict[str, str] = {
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F5X4": "0",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X2": "0",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F7X3": "0",
- # ASM Implicit GEMM — V4R1 only; no GTC/XDLOPS on RDNA2
+ # ASM Implicit GEMM — forward V4R1 only; no GTC/XDLOPS on RDNA2
+ # BWD (backward data-gradient) and WrW (weight-gradient) are training-only — disabled
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_V4R1": "1",
"MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_V4R1_1X1": "1",
- "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_BWD_V4R1": "1",
- "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_V4R1": "1",
- # HIP Implicit GEMM — non-XDLOPS V4R1/R4 only
+ "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_BWD_V4R1": "0",
+ "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_V4R1": "0",
+ # HIP Implicit GEMM — non-XDLOPS V4R1/R4 forward only
+ # BWD (backward data-gradient) and WrW (weight-gradient) are training-only — disabled
"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_BWD_V1R1": "0",
+ "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1": "0",
+ "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R1": "0",
+ "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4": "0",
}
# ---------------------------------------------------------------------------
diff --git a/scripts/rocm/rocm_vars.py b/scripts/rocm/rocm_vars.py
index dc768f88e..b6fce4faa 100644
--- a/scripts/rocm/rocm_vars.py
+++ b/scripts/rocm/rocm_vars.py
@@ -75,8 +75,13 @@ GENERAL_VARS: Dict[str, Dict[str, Any]] = {
},
}
-# --- Solver toggles: name -> description string ---
-# These are all checkbox widgets; default "1" unless in _SOLVER_DISABLED_BY_DEFAULT.
+# --- Solver toggles (inference/FWD only, RDNA2/3/4 compatible) ---
+# Removed entirely — not representable in the UI, cannot be set by users:
+# WRW (weight-gradient) and BWD (data-gradient) — training passes only, never run during inference
+# XDLOPS, CK (Composable Kernel), MLIR, MP BD — CDNA (MI100/MI200/MI300) matrix engine; not on RDNA
+# Fixed-geometry (5x10, 7x7-ImageNet, 11x11) — shapes never appear in SD/video inference
+# FP32-reference (NAIVE_CONV_FWD, FWDGEN) — IsApplicable() unreliable for FP16/BF16
+# Wide MPASS (F3x4..F7x3) — kernel sizes that cannot match any SD convolution shape
_SOLVER_DESCS: Dict[str, str] = {}
_SOLVER_DESCS.update({
@@ -87,184 +92,90 @@ _SOLVER_DESCS.update({
"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",
+ "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",
+ "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",
+ # Direct ASM — FWD inference only (WRW, fixed-geometry, FP32-reference removed)
+ "MIOPEN_DEBUG_CONV_DIRECT_ASM_3X3U": "Enable Direct ASM 3x3U",
+ "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U": "Enable Direct ASM 1x1U",
+ "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1UV2": "Enable Direct ASM 1x1UV2",
+ "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",
})
_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",
+ # Direct OCL — FWD inference only (WRW, FWD11X11 fixed-geom, FWDGEN FP32-ref removed)
+ "MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD": "Enable Direct OCL FWD",
+ "MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1": "Enable Direct OCL FWD1X1",
})
_SOLVER_DESCS.update({
- "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",
+ # Winograd FWD — WRW removed; Fury/Rage kept as RDNA3/4 inference (off by default)
+ "MIOPEN_DEBUG_AMD_WINOGRAD_3X3": "Enable AMD Winograd 3x3",
+ "MIOPEN_DEBUG_AMD_WINOGRAD_RXS": "Enable AMD Winograd RxS",
+ "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_FWD_BWD": "Enable AMD Winograd RxS FWD",
+ "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_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({
+ # Multi-pass Winograd — only F3x2/F3x3 match typical 3x3 SD shapes; wider kernels removed
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X2": "Enable AMD Winograd MPASS F3x2",
"MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X3": "Enable AMD Winograd MPASS F3x3",
- "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",
+ # Implicit GEMM FWD — BWD/WRW (training), all XDLOPS/CDNA variants removed
+ "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_V4R1": "Enable ASM Implicit GEMM FWD V4R1",
+ "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_V4R1_1X1": "Enable ASM Implicit GEMM FWD V4R1 1x1",
+ "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1": "Enable HIP Implicit GEMM FWD V4R1",
+ "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4": "Enable HIP Implicit GEMM FWD V4R4",
})
-# Vars whose solver default is "0" (disabled) — CDNA-only, RDNA3/4-only, experimental, or not used in SD inference
-_SOLVER_DISABLED_BY_DEFAULT = {
+
+# Solvers still in the registry but disabled by default.
+# FORCE_IMMED_MODE_FALLBACK — overrides FIND_MODE entirely, defeats tuning DB
+# Fury RxS F2x3/F3x2 — RDNA3/4-only; harmless on RDNA2 but won't select
+# Rage RxS F2x3 — RDNA4-only
+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",
+}
+
+SOLVER_DTYPE_TAGS: Dict[str, str] = {
+ "MIOPEN_DEBUG_CONV_DIRECT_ASM_3X3U": "FP16/FP32",
+ "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U": "FP16/FP32",
+ "MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1UV2": "FP16/FP32",
+ "MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD": "FP16/FP32",
+ "MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1": "FP16/FP32",
+ "MIOPEN_DEBUG_AMD_WINOGRAD_3X3": "FP32",
+ "MIOPEN_DEBUG_AMD_FUSED_WINOGRAD": "FP32",
+ "MIOPEN_DEBUG_AMD_WINOGRAD_RXS": "FP16/FP32",
+ "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_FWD_BWD": "FP16/FP32",
+ "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F3X2": "FP16/FP32",
+ "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3": "FP16/FP32",
+ "MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_G1": "FP16/FP32",
+ "MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F2X3": "FP16/FP32",
+ "MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F3X2": "FP16/FP32",
+ "MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3": "FP16/FP32",
+ "MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X2": "FP16/FP32",
+ "MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X3": "FP16/FP32",
+ "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_V4R1": "FP16/FP32",
+ "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_V4R1_1X1": "FP16/FP32",
+ "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1": "FP16/FP32",
+ "MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4": "FP16/FP32",
}
# Build full merged var registry
@@ -272,10 +183,11 @@ 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",
+ "default": "0" if _var in SOLVER_DISABLED_BY_DEFAULT else "1",
"desc": _desc,
"widget": "checkbox",
"options": None,
+ "dtype": SOLVER_DTYPE_TAGS.get(_var),
"restart_required": False,
}
@@ -296,79 +208,26 @@ SOLVER_GROUPS: List[Tuple[str, List[str]]] = [
]),
("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_1X1UV2",
"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_FWD_BWD",
"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_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", [
+ ("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 b899cc4db..307df918a 100644
--- a/scripts/rocm_ext.py
+++ b/scripts/rocm_ext.py
@@ -35,12 +35,12 @@ class Script(scripts_manager.Script):
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()}")
+ dtype_tag = meta.get("dtype")
+ label = f"[{dtype_tag}] {meta['desc']}" if dtype_tag else meta["desc"]
+ return gr.Checkbox(label=label, 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)
@@ -88,15 +88,45 @@ class Script(scripts_manager.Script):
gap: 0 !important;
}
#rocm_config fieldset { border: none !important; padding: 0 !important; margin: 0 !important; box-shadow: none !important; background: transparent !important; }
+ /* No horizontal overflow on the container */
+ #rocm_config { overflow-x: hidden !important; }
+ /* Smaller text for dense solver checkbox lists */
+ #rocm_config .gradio-checkbox label { font-size: 11px !important; line-height: 1.4 !important; }
+
+ /* CSS-only tooltips — script tags are sanitized by Gradio */
+ #rocm_btn_info, #rocm_btn_apply, #rocm_btn_reset, #rocm_btn_clear, #rocm_btn_delete,
+ #rocm_btn_rdna2, #rocm_btn_rdna3, #rocm_btn_rdna4 { position: relative !important; overflow: visible !important; }
+ #rocm_btn_info::after { content: "Reload DB and device version info"; }
+ #rocm_btn_apply::after { content: "Save and apply current settings to runtime"; }
+ #rocm_btn_reset::after { content: "Reset all settings to built-in defaults"; }
+ #rocm_btn_clear::after { content: "Remove all MIOpen/HIP vars from runtime env"; }
+ #rocm_btn_delete::after { content: "Delete saved config + wipe MIOpen user DB cache"; }
+ #rocm_btn_rdna2::after { content: "Apply RDNA2 (RX 6000 series) solver profile"; }
+ #rocm_btn_rdna3::after { content: "Apply RDNA3 (RX 7000 series) solver profile"; }
+ #rocm_btn_rdna4::after { content: "Apply RDNA4 (RX 9000 series) solver profile"; }
+ #rocm_btn_info::after, #rocm_btn_apply::after, #rocm_btn_reset::after,
+ #rocm_btn_clear::after, #rocm_btn_delete::after,
+ #rocm_btn_rdna2::after, #rocm_btn_rdna3::after, #rocm_btn_rdna4::after {
+ display: block; position: absolute; top: calc(100% + 4px); left: 0;
+ background: var(--background-fill-secondary, #2a2a2a);
+ color: var(--body-text-color, #ddd);
+ border: 1px solid var(--border-color-primary, #555);
+ padding: 3px 8px; border-radius: 4px; font-size: 11px;
+ white-space: nowrap; pointer-events: none;
+ opacity: 0; transition: opacity 0.15s; z-index: 999;
+ }
+ #rocm_btn_info:hover::after, #rocm_btn_apply:hover::after, #rocm_btn_reset:hover::after,
+ #rocm_btn_clear:hover::after, #rocm_btn_delete:hover::after,
+ #rocm_btn_rdna2:hover::after, #rocm_btn_rdna3:hover::after, #rocm_btn_rdna4:hover::after { opacity: 1; }
""")
with gr.Row():
- gr.HTML("Advanced configuration for ROCm users.
Set Your database and solver selections based on GPU profile or individually.
Enable cuDNN in Backend Settings to activate MIOpen.
")
+ gr.HTML("Advanced configuration for ROCm users.
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")
+ btn_info = gr.Button("Refresh Info", variant="primary", elem_id="rocm_btn_info", size="sm")
+ btn_apply = gr.Button("Apply", elem_id="rocm_btn_apply", size="sm")
+ btn_reset = gr.Button("Defaults", elem_id="rocm_btn_reset", size="sm")
+ btn_clear = gr.Button("Clear vEnv", elem_id="rocm_btn_clear", size="sm")
+ btn_delete = gr.Button("Delete UserDb", elem_id="rocm_btn_delete", size="sm")
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")
@@ -123,19 +153,18 @@ class Script(scripts_manager.Script):
components.append(comp)
gr.HTML("
")
- def _autosave_dropdown(name, value):
+ def _autosave_field(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)
+ stored = rocm_mgr._dropdown_stored(str(value), meta["options"])
+ cfg = rocm_mgr.load_config()
+ cfg[name] = stored
+ rocm_mgr.save_config(cfg)
+ rocm_mgr.apply_env(cfg)
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=[])
+ comp.change(fn=lambda v, n=name: _autosave_field(n, v), inputs=[comp], outputs=[], show_progress='hidden')
def apply_fn(*values):
rocm_mgr.apply_all(var_names, list(values))
@@ -170,7 +199,7 @@ class Script(scripts_manager.Script):
val = updated.get(name, meta["default"])
if meta["widget"] == "checkbox":
result.append(gr.update(value=(val == "1")))
- elif meta["widget"] == "radio":
+ 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)))
@@ -183,9 +212,8 @@ class Script(scripts_manager.Script):
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))
+ elif meta["widget"] == "dropdown":
+ result.append(gr.update(value=rocm_mgr._dropdown_display(meta["default"], meta["options"])))
else:
result.append(gr.update(value=""))
return result
@@ -197,9 +225,8 @@ class Script(scripts_manager.Script):
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))
+ elif meta["widget"] == "dropdown":
+ result.append(gr.update(value=rocm_mgr._dropdown_display(meta["default"], meta["options"])))
else:
result.append(gr.update(value=""))
return result
@@ -215,19 +242,19 @@ class Script(scripts_manager.Script):
val = updated.get(pname, meta["default"])
if meta["widget"] == "checkbox":
result.append(gr.update(value=(val == "1")))
- elif meta["widget"] == "radio":
+ 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
- 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)
+ btn_info.click(fn=_info_html, inputs=[], outputs=[info_out], show_progress='hidden')
+ btn_apply.click(fn=apply_fn, inputs=components, outputs=[style_out] + components, show_progress='hidden')
+ btn_reset.click(fn=reset_fn, inputs=[], outputs=[style_out] + components, show_progress='hidden')
+ btn_clear.click(fn=clear_fn, inputs=[], outputs=[style_out] + components, show_progress='hidden')
+ btn_delete.click(fn=delete_fn, inputs=[], outputs=[style_out] + components, show_progress='hidden')
+ btn_rdna2.click(fn=lambda: profile_fn("RDNA2"), inputs=[], outputs=[style_out] + components, show_progress='hidden')
+ btn_rdna3.click(fn=lambda: profile_fn("RDNA3"), inputs=[], outputs=[style_out] + components, show_progress='hidden')
+ btn_rdna4.click(fn=lambda: profile_fn("RDNA4"), inputs=[], outputs=[style_out] + components, show_progress='hidden')
return components