mirror of
https://github.com/vladmandic/automatic
synced 2026-09-19 09:14:35 +02:00
ROCm: Advanced Configuration Script
This commit is contained in:
resonantsky
@@ -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"
|
||||
}
|
||||
+153
-18
@@ -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)}')
|
||||
|
||||
@@ -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",
|
||||
}
|
||||
|
||||
# ---------------------------------------------------------------------------
|
||||
|
||||
+78
-219
@@ -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",
|
||||
]),
|
||||
]
|
||||
|
||||
|
||||
+60
-33
@@ -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; }
|
||||
</style>""")
|
||||
with gr.Row():
|
||||
gr.HTML("<p>Advanced configuration for ROCm users.</p><br><p>Set Your database and solver selections based on GPU profile or individually.</p><br><p>Enable cuDNN in Backend Settings to activate MIOpen.</p>")
|
||||
gr.HTML("<p>Advanced configuration for ROCm users.</p><br><p>Set your database and solver selections based on GPU profile or individually.</p><br><p>Enable cuDNN in Backend Settings to activate MIOpen.</p>")
|
||||
with gr.Row():
|
||||
btn_info = gr.Button("Refresh Info", variant="primary", elem_id="rocm_btn_info")
|
||||
btn_apply = gr.Button("Apply", elem_id="rocm_btn_apply")
|
||||
btn_reset = gr.Button("Defaults")
|
||||
btn_clear = gr.Button("Clear Runtime")
|
||||
btn_delete = gr.Button("Delete", elem_id="rocm_btn_delete")
|
||||
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("<br><center><div style='margin:0 Auto'><a href='https://rocm.docs.amd.com/projects/MIOpen/en/develop/reference/env_variables.html' target='_blank'>📄 MIOpen Environment Variables Reference</a></div></center><br>")
|
||||
|
||||
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
|
||||
|
||||
Reference in New Issue
Block a user