From 56158454b3651e580e81a2db5563fc4c10dd3195 Mon Sep 17 00:00:00 2001 From: Lei Zhenyuan Date: Wed, 11 Feb 2026 09:50:33 +0800 Subject: [PATCH 01/63] fix for tma (#4023) --- unsloth/kernels/moe/grouped_gemm/interface.py | 3 +++ 1 file changed, 3 insertions(+) diff --git a/unsloth/kernels/moe/grouped_gemm/interface.py b/unsloth/kernels/moe/grouped_gemm/interface.py index 572b02702b..554e5fcc03 100644 --- a/unsloth/kernels/moe/grouped_gemm/interface.py +++ b/unsloth/kernels/moe/grouped_gemm/interface.py @@ -4,6 +4,7 @@ import logging import warnings from dataclasses import asdict +from unsloth import DEVICE_TYPE import torch import triton @@ -41,6 +42,8 @@ logger.addHandler(ch) # 1. GPU capability >= 9 (Hopper+) # 2. Triton version with TMA API (make_tensor_descriptor or _experimental_make_tensor_descriptor) def _check_tma_support(): + if DEVICE_TYPE == "xpu": + return False import triton.language as tl gpu_supports_tma = torch.cuda.get_device_capability()[0] >= 9 From 25cae890e6bdb56bb3e72c2a734220e7128efbc7 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Thu, 12 Feb 2026 02:56:13 -0800 Subject: [PATCH 02/63] Create CODEOWNERS (#4039) --- .github/CODEOWNERS | 46 ++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 46 insertions(+) create mode 100644 .github/CODEOWNERS diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS new file mode 100644 index 0000000000..f4d18645f2 --- /dev/null +++ b/.github/CODEOWNERS @@ -0,0 +1,46 @@ +# Inspired from https://github.com/vllm-project/vllm/blob/main/.github/CODEOWNERS + +/unsloth/models/loader.py @danielhanchen @mmathew23 +/unsloth/models/llama.py @Datta0 @danielhanchen @mmathew23 +/unsloth/models/rl.py @Datta0 @pluesclues @danielhanchen +/unsloth/models/rl_replacements.py @Datta0 @pluesclues @danielhanchen +/unsloth/trainer.py @danielhanchen +/unsloth/models/sentence_transformer.py @Etherll @danielhanchen +/unsloth/save.py @rolandtannous @danielhanchen +/unsloth/tokenizer_utils.py @mmathew23 @danielhanchen +/unsloth/chat_templates.py @rolandtannous @danielhanchen +/unsloth/ollama_template_mappers.py @rolandtannous @danielhanchen +/unsloth/kernels/moe/*.py @Datta0 +/unsloth/import_fixes.py @danielhanchen +/unsloth/device_type.py @danielhanchen +/unsloth/_auto_install.py @danielhanchen +/unsloth/dataprep/*.py @danielhanchen +/unsloth/kernels/cross_entropy_loss.py @danielhanchen +/unsloth/kernels/fast_lora.py @danielhanchen +/unsloth/kernels/flex_attention.py @danielhanchen +/unsloth/kernels/fp8.py @Datta0 +/unsloth/kernels/geglu.py @danielhanchen +/unsloth/kernels/layernorm.py @danielhanchen +/unsloth/kernels/rms_layernorm.py @danielhanchen +/unsloth/kernels/rope_embedding.py @danielhanchen +/unsloth/kernels/swiglu.py @danielhanchen +/unsloth/kernels/utils.py @danielhanchen @Datta0 +/unsloth/models/_utils.py @danielhanchen @mmathew23 +/unsloth/models/cohere.py @danielhanchen +/unsloth/models/dpo.py @danielhanchen +/unsloth/models/falcon_h1.py @danielhanchen +/unsloth/models/gemma.py @danielhanchen +/unsloth/models/gemma2.py @danielhanchen +/unsloth/models/glm4_moe.py @Datta0 +/unsloth/models/granite.py @danielhanchen +/unsloth/models/llama4.py @danielhanchen +/unsloth/models/loader_utils.py @Datta0 @danielhanchen +/unsloth/models/mapper.py @danielhanchen +/unsloth/models/mistral.py @danielhanchen +/unsloth/models/qwen2.py @danielhanchen +/unsloth/models/qwen3.py @Datta0 +/unsloth/models/qwen3_moe.py @Datta0 +/unsloth/models/vision.py @mmathew23 @danielhanchen +/unsloth/utils/attention_dispatch.py @mmathew23 +/unsloth/utils/hf_hub.py @mmathew23 +/unsloth/utils/packing.py @mmathew23 From 9b2a44290a45c0ac6a9dfc24028c1b31909090c9 Mon Sep 17 00:00:00 2001 From: Michael Han <107991372+shimmyshimmer@users.noreply.github.com> Date: Fri, 13 Feb 2026 19:38:23 -0800 Subject: [PATCH 03/63] Update README with faster MoE.md Adding MoE --- README.md | 1 + 1 file changed, 1 insertion(+) diff --git a/README.md b/README.md index bfe7f2255b..88f8c77ffa 100644 --- a/README.md +++ b/README.md @@ -53,6 +53,7 @@ Use our official [Unsloth Docker image](https://hub.docker.com/r/unsloth/unsloth For RTX 50x, B200, 6000 GPUs: `pip install unsloth`. Read our [Blackwell Guide](https://unsloth.ai/docs/blog/fine-tuning-llms-with-blackwell-rtx-50-series-and-unsloth) and [DGX Spark Guide](https://unsloth.ai/docs/blog/fine-tuning-llms-with-nvidia-dgx-spark-and-unsloth) for more details. ## 🦥 Unsloth News +- Train **MoE LLMs 12x faster** with 35% less VRAM - DeepSeek, GLM, Qwen and gpt-oss. [Blog](https://unsloth.ai/docs/new/faster-moe) - **Embedding models**: Unsloth now supports ~1.8-3.3x faster embedding fine-tuning. [Blog](https://unsloth.ai/docs/new/embedding-finetuning) • [Notebooks](https://unsloth.ai/docs/get-started/unsloth-notebooks#embedding-models) - New **7x longer context RL** vs. all other setups, via our new batching algorithms. [Blog](https://unsloth.ai/docs/new/grpo-long-context) - New RoPE & MLP **Triton Kernels** & **Padding Free + Packing**: 3x faster training & 30% less VRAM. [Blog](https://unsloth.ai/docs/new/3x-faster-training-packing) From 7254d47da1cf152b591fe532d9cb98d9ade3e31e Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Fri, 13 Feb 2026 21:20:25 -0800 Subject: [PATCH 04/63] Handle broken causal_conv1d at import time (#4047) * Handle broken causal_conv1d import at runtime Add a startup import-time probe for causal_conv1d and disable the fast path when the shared library is ABI broken. This keeps Falcon H1/model loading resilient without requiring env flags. - Add disable_broken_causal_conv1d in import_fixes. - Invoke it early from unsloth/__init__ during package init. - Make Falcon H1 optional imports in loader and models/__init__ soft-fail instead of failing hard. * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Enforce unavailable semantics for broken causal_conv1d * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Remove Falcon H1 import swallowing * Restore optional Falcon H1 import guard * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Remove causal_conv1d regression tests * Trim FA2 fallback messaging --------- Co-authored-by: Daniel Hanchen Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- unsloth/__init__.py | 3 + unsloth/import_fixes.py | 157 +++++++++++++++++++++++++++++++++++++++ unsloth/models/_utils.py | 6 +- 3 files changed, 162 insertions(+), 4 deletions(-) diff --git a/unsloth/__init__.py b/unsloth/__init__.py index b068d6a5fc..a505e89ad4 100644 --- a/unsloth/__init__.py +++ b/unsloth/__init__.py @@ -28,16 +28,19 @@ already_imported = [mod for mod in critical_modules if mod in sys.modules] from .import_fixes import ( fix_message_factory_issue, check_fbgemm_gpu_version, + disable_broken_causal_conv1d, torchvision_compatibility_check, fix_diffusers_warnings, fix_huggingface_hub, ) +disable_broken_causal_conv1d() fix_message_factory_issue() check_fbgemm_gpu_version() torchvision_compatibility_check() fix_diffusers_warnings() fix_huggingface_hub() +del disable_broken_causal_conv1d del fix_message_factory_issue del check_fbgemm_gpu_version del torchvision_compatibility_check diff --git a/unsloth/import_fixes.py b/unsloth/import_fixes.py index 4a2211f9a0..8d00b793c9 100644 --- a/unsloth/import_fixes.py +++ b/unsloth/import_fixes.py @@ -13,6 +13,8 @@ # limitations under the License. import os +import importlib.abc +import importlib.machinery import importlib.util from pathlib import Path from importlib.metadata import version as importlib_version @@ -21,6 +23,7 @@ import re import logging import textwrap import warnings +import sys # We cannot do from unsloth_zoo.log import logger since FBGEMM might cause seg faults. UNSLOTH_ENABLE_LOGGING = os.environ.get("UNSLOTH_ENABLE_LOGGING", "0") in ( @@ -1163,3 +1166,157 @@ def disable_torchcodec_if_broken(): tf_import_utils._torchcodec_available = False except (ImportError, AttributeError): pass + + +CAUSAL_CONV1D_BROKEN = False +_CAUSAL_CONV1D_PREFIX = "causal_conv1d" +_CAUSAL_CONV1D_BLOCKER_SENTINEL = "_unsloth_causal_conv1d_blocker" + + +def _is_causal_conv1d_name(module_name: str) -> bool: + return module_name == _CAUSAL_CONV1D_PREFIX or module_name.startswith( + _CAUSAL_CONV1D_PREFIX + "." + ) + + +def _resolve_module_name(module_name, package): + if not isinstance(module_name, str): + return module_name + if module_name.startswith("."): + try: + return importlib.util.resolve_name(module_name, package) + except Exception: + return module_name + return module_name + + +def _is_broken_causal_conv1d_error(error) -> bool: + checked = set() + current = error + while current is not None and id(current) not in checked: + checked.add(id(current)) + message = str(current).lower() + if ( + ("causal_conv1d_cuda" in message and "undefined symbol" in message) + or ("_zn3c103hip28c10_hip_check_implementation" in message) + or ("causal_conv1d" in message and "undefined symbol" in message) + ): + return True + current = getattr(current, "__cause__", None) or getattr( + current, "__context__", None + ) + return False + + +class _CausalConv1dImportBlockerLoader(importlib.abc.Loader): + __slots__ = ("module_name",) + + def __init__(self, module_name): + self.module_name = module_name + + def create_module(self, spec): + return None + + def exec_module(self, module): + raise ModuleNotFoundError(f"No module named '{self.module_name}'") + + +class _CausalConv1dImportBlockerFinder(importlib.abc.MetaPathFinder): + __slots__ = (_CAUSAL_CONV1D_BLOCKER_SENTINEL,) + + def __init__(self): + setattr(self, _CAUSAL_CONV1D_BLOCKER_SENTINEL, True) + + def find_spec(self, fullname, path = None, target = None): + if not CAUSAL_CONV1D_BROKEN or not _is_causal_conv1d_name(fullname): + return None + return importlib.machinery.ModuleSpec( + name = fullname, + loader = _CausalConv1dImportBlockerLoader(fullname), + is_package = fullname == _CAUSAL_CONV1D_PREFIX, + ) + + +def _patch_find_spec_for_causal_conv1d(): + current_find_spec = importlib.util.find_spec + if getattr(current_find_spec, "_unsloth_causal_conv1d_find_spec_patch", False): + return + + def _blocked_find_spec(name, package = None): + resolved_name = _resolve_module_name(name, package) + if CAUSAL_CONV1D_BROKEN and isinstance(resolved_name, str): + if _is_causal_conv1d_name(resolved_name): + return None + return current_find_spec(name, package) + + _blocked_find_spec._unsloth_causal_conv1d_find_spec_patch = True + _blocked_find_spec._unsloth_original_find_spec = current_find_spec + importlib.util.find_spec = _blocked_find_spec + + +def _install_causal_conv1d_blocker(): + _patch_find_spec_for_causal_conv1d() + for finder in sys.meta_path: + if getattr(finder, _CAUSAL_CONV1D_BLOCKER_SENTINEL, False): + return + sys.meta_path.insert(0, _CausalConv1dImportBlockerFinder()) + + +def _clear_causal_conv1d_modules(): + for module_name in list(sys.modules): + if _is_causal_conv1d_name(module_name): + sys.modules.pop(module_name, None) + + +def _disable_transformers_causal_conv1d(): + try: + import transformers.utils.import_utils as tf_import_utils + except Exception: + return + + if hasattr(tf_import_utils, "is_causal_conv1d_available"): + tf_import_utils.is_causal_conv1d_available = lambda: False + + for attr_name in ( + "_causal_conv1d_available", + "_is_causal_conv1d_available", + ): + if hasattr(tf_import_utils, attr_name): + setattr(tf_import_utils, attr_name, False) + + +def disable_broken_causal_conv1d(): + """Disable causal_conv1d dynamically when its shared library is ABI-broken. + + This mirrors Unsloth's FlashAttention fallback behavior: if importing causal_conv1d + fails with a known binary symbol error, we disable it at startup so model imports do + not hard-fail. + """ + global CAUSAL_CONV1D_BROKEN + if CAUSAL_CONV1D_BROKEN: + _install_causal_conv1d_blocker() + _disable_transformers_causal_conv1d() + return + + try: + if importlib.util.find_spec("causal_conv1d") is None: + return + except Exception: + return + + try: + import causal_conv1d # noqa: F401 + + return + except Exception as error: + if not _is_broken_causal_conv1d_error(error): + return + + CAUSAL_CONV1D_BROKEN = True + _clear_causal_conv1d_modules() + _install_causal_conv1d_blocker() + _disable_transformers_causal_conv1d() + print( + "Unsloth: Detected broken causal_conv1d binary; " + "disabling causal_conv1d fast path and continuing import." + ) diff --git a/unsloth/models/_utils.py b/unsloth/models/_utils.py index 70d71b3e06..9c4e8742d2 100644 --- a/unsloth/models/_utils.py +++ b/unsloth/models/_utils.py @@ -845,8 +845,7 @@ if DEVICE_TYPE == "cuda": "Unsloth: Your Flash Attention 2 installation seems to be broken?\n" "A possible explanation is you have a new CUDA version which isn't\n" "yet compatible with FA2? Please file a ticket to Unsloth or FA2.\n" - "We shall now use Xformers instead, which does not have any performance hits!\n" - "We found this negligible impact by benchmarking on 1x A100." + "We shall now use Xformers instead, which does not have any performance hits!" ) # Stop Flash Attention from importing! @@ -897,8 +896,7 @@ elif DEVICE_TYPE == "hip": "Unsloth: Your Flash Attention 2 installation seems to be broken?\n" "A possible explanation is you have a new CUDA version which isn't\n" "yet compatible with FA2? Please file a ticket to Unsloth or FA2.\n" - "We shall now use Xformers instead, which does not have any performance hits!\n" - "We found this negligible impact by benchmarking on 1x A100." + "We shall now use Xformers instead, which does not have any performance hits!" ) # Stop Flash Attention from importing! From a0f4b078c6f021840bc27d638850deb0003779a2 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Fri, 13 Feb 2026 21:32:34 -0800 Subject: [PATCH 05/63] Improve HIP GPU name detection in startup banner (#4048) * Improve HIP GPU name reporting in startup banner * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --------- Co-authored-by: Daniel Hanchen Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- unsloth/models/_utils.py | 32 ++++++++++++++++++++++++++++++++ unsloth/models/llama.py | 4 +--- unsloth/models/vision.py | 4 +--- 3 files changed, 34 insertions(+), 6 deletions(-) diff --git a/unsloth/models/_utils.py b/unsloth/models/_utils.py index 9c4e8742d2..e5b73db079 100644 --- a/unsloth/models/_utils.py +++ b/unsloth/models/_utils.py @@ -28,6 +28,7 @@ __all__ = [ "HAS_FLASH_ATTENTION_SOFTCAPPING", "USE_MODELSCOPE", "platform_system", + "resolve_hip_gpu_stats_name", "patch_tokenizer", "get_statistics", "Unsloth_Offloaded_Gradient_Checkpointer", @@ -149,6 +150,37 @@ from unsloth_zoo.compiler import ( from unsloth_zoo.training_utils import ( prepare_model_for_training, ) + + +def resolve_hip_gpu_stats_name(gpu_stats): + name = str(getattr(gpu_stats, "name", "") or "").strip() + normalized_name = name.lower().strip(". ") + if normalized_name and normalized_name not in ("amd radeon graphics",): + return name + ". " + + try: + torch_name = str(torch.cuda.get_device_name(0) or "").strip() + except Exception: + torch_name = "" + normalized_torch_name = torch_name.lower().strip(". ") + if normalized_torch_name and normalized_torch_name not in ("amd radeon graphics",): + return torch_name + ". " + + arch_name = "" + for key in ("gcnArchName", "gcn_arch_name", "arch_name", "gfx_arch_name"): + value = getattr(gpu_stats, key, None) + if value is not None and str(value).strip(): + arch_name = str(value).strip() + break + + if arch_name: + # gfx942 maps to MI300X on current ROCm naming. + if arch_name.lower().startswith("gfx942"): + return f"AMD Instinct MI300X ({arch_name}). " + return f"AMD GPU ({arch_name}). " + return "AMD GPU Device. " + + from unsloth_zoo.temporary_patches import ( TEMPORARY_PATCHES, ) diff --git a/unsloth/models/llama.py b/unsloth/models/llama.py index 043d2363c1..dcaa770baa 100644 --- a/unsloth/models/llama.py +++ b/unsloth/models/llama.py @@ -2245,9 +2245,7 @@ class FastLlamaModel: vllm_version = "" elif DEVICE_TYPE == "hip": gpu_stats = torch.cuda.get_device_properties(0) - gpu_stats_name = ( - gpu_stats.name + ". " if gpu_stats.name != "" else "AMD GPU Device. " - ) + gpu_stats_name = resolve_hip_gpu_stats_name(gpu_stats) gpu_version = torch.version.hip gpu_stats_snippet = f"ROCm Toolkit: {gpu_version}." try: diff --git a/unsloth/models/vision.py b/unsloth/models/vision.py index c294dbdb0b..9e292a2849 100644 --- a/unsloth/models/vision.py +++ b/unsloth/models/vision.py @@ -497,9 +497,7 @@ class FastBaseModel: vllm_version = "" elif DEVICE_TYPE == "hip": gpu_stats = torch.cuda.get_device_properties(0) - gpu_stats_name = ( - gpu_stats.name + ". " if gpu_stats.name != "" else "AMD GPU Device. " - ) + gpu_stats_name = resolve_hip_gpu_stats_name(gpu_stats) gpu_version = torch.version.hip gpu_stats_snippet = f"ROCm Toolkit: {gpu_version}." try: From 596683e7ca43a300b447a2ba575de8ea4ff7e227 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sat, 14 Feb 2026 02:24:03 -0800 Subject: [PATCH 06/63] Simplify MI300X startup banner name (#4049) * Improve HIP GPU name reporting in startup banner * Drop MI300X arch suffix in banner name * Normalize _utils.py file mode * Simplify FA2 fallback text and filter AMD ids noise * Strip trailing GPU arch suffix via regex * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Use gfx lookup default and normalize Ryzen AI naming * Remove name-path Ryzen AI normalization * Expand ROCm gfx map to full documented GPU name aliases * Simplify HIP fallback naming to AMD gfx token * Remove Ryzen Al torch_name normalization --------- Co-authored-by: Daniel Hanchen Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- unsloth/import_fixes.py | 2 ++ unsloth/models/_utils.py | 24 +++++++++++------------- 2 files changed, 13 insertions(+), 13 deletions(-) diff --git a/unsloth/import_fixes.py b/unsloth/import_fixes.py index 8d00b793c9..d037903bec 100644 --- a/unsloth/import_fixes.py +++ b/unsloth/import_fixes.py @@ -199,6 +199,8 @@ if os.environ.get("UNSLOTH_ENABLE_LOGGING", "0") != "1": # Triton "df: No such file or directory" stderr noise sys.stderr.add_filter("df: No such file") + # ROCm/libdrm missing ids table stderr noise on some AMD setups + sys.stderr.add_filter("amdgpu.ids: No such file or directory") # Fix up AttributeError: 'MessageFactory' object has no attribute 'GetPrototype' diff --git a/unsloth/models/_utils.py b/unsloth/models/_utils.py index e5b73db079..451f60091f 100644 --- a/unsloth/models/_utils.py +++ b/unsloth/models/_utils.py @@ -154,12 +154,14 @@ from unsloth_zoo.training_utils import ( def resolve_hip_gpu_stats_name(gpu_stats): name = str(getattr(gpu_stats, "name", "") or "").strip() + name = re.sub(r"\s*\([^)]*\)\s*$", "", name).strip() normalized_name = name.lower().strip(". ") if normalized_name and normalized_name not in ("amd radeon graphics",): return name + ". " try: torch_name = str(torch.cuda.get_device_name(0) or "").strip() + torch_name = re.sub(r"\s*\([^)]*\)\s*$", "", torch_name).strip() except Exception: torch_name = "" normalized_torch_name = torch_name.lower().strip(". ") @@ -174,11 +176,11 @@ def resolve_hip_gpu_stats_name(gpu_stats): break if arch_name: - # gfx942 maps to MI300X on current ROCm naming. - if arch_name.lower().startswith("gfx942"): - return f"AMD Instinct MI300X ({arch_name}). " - return f"AMD GPU ({arch_name}). " - return "AMD GPU Device. " + arch_name = arch_name.strip() + match = re.search(r"(gfx[0-9a-z]+)", arch_name, flags = re.I) + if match: + return f"AMD {match.group(1).lower()} GPU. " + return "AMD GPU. " from unsloth_zoo.temporary_patches import ( @@ -874,10 +876,8 @@ if DEVICE_TYPE == "cuda": ) except: print( - "Unsloth: Your Flash Attention 2 installation seems to be broken?\n" - "A possible explanation is you have a new CUDA version which isn't\n" - "yet compatible with FA2? Please file a ticket to Unsloth or FA2.\n" - "We shall now use Xformers instead, which does not have any performance hits!" + "Unsloth: Your Flash Attention 2 installation seems to be broken. " + "Using Xformers instead. No performance changes will be seen." ) # Stop Flash Attention from importing! @@ -925,10 +925,8 @@ elif DEVICE_TYPE == "hip": ) except: print( - "Unsloth: Your Flash Attention 2 installation seems to be broken?\n" - "A possible explanation is you have a new CUDA version which isn't\n" - "yet compatible with FA2? Please file a ticket to Unsloth or FA2.\n" - "We shall now use Xformers instead, which does not have any performance hits!" + "Unsloth: Your Flash Attention 2 installation seems to be broken. " + "Using Xformers instead. No performance changes will be seen." ) # Stop Flash Attention from importing! From 483c6fe055fbb10ead69666075b0628ac43ea73f Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sat, 14 Feb 2026 02:44:34 -0800 Subject: [PATCH 07/63] Suppress HIP amdgpu.ids stderr noise during causal_conv1d check (#4052) * Suppress HIP libdrm stderr noise in causal_conv1d probe * Broaden HIP libdrm stderr suppression for early ROCm startup --------- Co-authored-by: Daniel Hanchen --- unsloth/__init__.py | 21 ++++++++++------- unsloth/import_fixes.py | 52 ++++++++++++++++++++++++++++++++++++++++- 2 files changed, 63 insertions(+), 10 deletions(-) diff --git a/unsloth/__init__.py b/unsloth/__init__.py index a505e89ad4..990e4b915f 100644 --- a/unsloth/__init__.py +++ b/unsloth/__init__.py @@ -29,6 +29,7 @@ from .import_fixes import ( fix_message_factory_issue, check_fbgemm_gpu_version, disable_broken_causal_conv1d, + _suppress_hip_libdrm_ids_noise, torchvision_compatibility_check, fix_diffusers_warnings, fix_huggingface_hub, @@ -106,7 +107,8 @@ del PackageNotFoundError, importlib_version # Try importing PyTorch and check version try: - import torch + with _suppress_hip_libdrm_ids_noise(): + import torch except ModuleNotFoundError: raise ImportError( "Unsloth: Pytorch is not installed. Go to https://pytorch.org/.\n" @@ -115,14 +117,15 @@ except ModuleNotFoundError: except: raise -from unsloth_zoo.device_type import ( - is_hip, - get_device_type, - DEVICE_TYPE, - DEVICE_TYPE_TORCH, - DEVICE_COUNT, - ALLOW_PREQUANTIZED_MODELS, -) +with _suppress_hip_libdrm_ids_noise(): + from unsloth_zoo.device_type import ( + is_hip, + get_device_type, + DEVICE_TYPE, + DEVICE_TYPE_TORCH, + DEVICE_COUNT, + ALLOW_PREQUANTIZED_MODELS, + ) # Fix other issues from .import_fixes import ( diff --git a/unsloth/import_fixes.py b/unsloth/import_fixes.py index d037903bec..0cac361888 100644 --- a/unsloth/import_fixes.py +++ b/unsloth/import_fixes.py @@ -16,6 +16,7 @@ import os import importlib.abc import importlib.machinery import importlib.util +import contextlib from pathlib import Path from importlib.metadata import version as importlib_version from packaging.version import Version as TrueVersion @@ -1175,6 +1176,54 @@ _CAUSAL_CONV1D_PREFIX = "causal_conv1d" _CAUSAL_CONV1D_BLOCKER_SENTINEL = "_unsloth_causal_conv1d_blocker" +def _is_rocm_torch_build() -> bool: + try: + return "rocm" in str(importlib_version("torch")).lower() + except Exception: + return False + + +@contextlib.contextmanager +def _suppress_stderr_fd(): + saved_stderr_fd = None + devnull_fd = None + redirected = False + try: + saved_stderr_fd = os.dup(2) + devnull_fd = os.open(os.devnull, os.O_WRONLY) + os.dup2(devnull_fd, 2) + redirected = True + except Exception: + redirected = False + + try: + yield + finally: + if redirected and saved_stderr_fd is not None: + try: + os.dup2(saved_stderr_fd, 2) + except Exception: + pass + if devnull_fd is not None: + try: + os.close(devnull_fd) + except Exception: + pass + if saved_stderr_fd is not None: + try: + os.close(saved_stderr_fd) + except Exception: + pass + + +def _suppress_hip_libdrm_ids_noise(): + # ROCm/libdrm can emit amdgpu.ids missing errors via low-level fd=2 writes. + # Python-level stderr filters cannot intercept those writes. + if not _is_rocm_torch_build(): + return contextlib.nullcontext() + return _suppress_stderr_fd() + + def _is_causal_conv1d_name(module_name: str) -> bool: return module_name == _CAUSAL_CONV1D_PREFIX or module_name.startswith( _CAUSAL_CONV1D_PREFIX + "." @@ -1307,7 +1356,8 @@ def disable_broken_causal_conv1d(): return try: - import causal_conv1d # noqa: F401 + with _suppress_hip_libdrm_ids_noise(): + import causal_conv1d # noqa: F401 return except Exception as error: From ed1cd2716cc25b8a544caa0352b2e0dec899248f Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E9=87=91=E9=BB=84=E8=89=B2=E8=91=A1=E8=90=84=E7=90=83?= =?UTF-8?q?=E5=90=9B=E5=90=9B?= Date: Sat, 14 Feb 2026 18:50:05 +0800 Subject: [PATCH 08/63] ROCm: Add gfx950 (MI355X/CDNA4) to is_cdna() (#4051) MI355X (gfx950) has the same 1024-thread workgroup limit as MI300X (gfx942), but was missing from is_cdna(), causing all Triton kernels to use num_warps=32 (2048 threads) instead of 16 (1024 threads), resulting in OutOfResources crash. Tested on: 8x AMD Instinct MI355X (gfx950), ROCm 7.1 --- unsloth/kernels/utils.py | 1 + 1 file changed, 1 insertion(+) diff --git a/unsloth/kernels/utils.py b/unsloth/kernels/utils.py index 5dcc7c232c..0b13c04532 100644 --- a/unsloth/kernels/utils.py +++ b/unsloth/kernels/utils.py @@ -82,6 +82,7 @@ def is_cdna(): "gfx940", "gfx941", "gfx942", + "gfx950", # CDNA4 (MI350/MI355X) ) From f6b49063c199ce2f5c73795b9111d3390f0c977c Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sat, 14 Feb 2026 02:59:49 -0800 Subject: [PATCH 09/63] Make ROCm suppression detection robust for custom torch builds (#4053) * Make ROCm suppression detection robust for custom torch builds * Add ROCm detection debug logging behind UNSLOTH_ENABLE_LOGGING --------- Co-authored-by: Daniel Hanchen --- unsloth/import_fixes.py | 49 +++++++++++++++++++++++++++++++++++++++-- 1 file changed, 47 insertions(+), 2 deletions(-) diff --git a/unsloth/import_fixes.py b/unsloth/import_fixes.py index 0cac361888..63e99386d8 100644 --- a/unsloth/import_fixes.py +++ b/unsloth/import_fixes.py @@ -25,6 +25,7 @@ import logging import textwrap import warnings import sys +import functools # We cannot do from unsloth_zoo.log import logger since FBGEMM might cause seg faults. UNSLOTH_ENABLE_LOGGING = os.environ.get("UNSLOTH_ENABLE_LOGGING", "0") in ( @@ -1176,11 +1177,55 @@ _CAUSAL_CONV1D_PREFIX = "causal_conv1d" _CAUSAL_CONV1D_BLOCKER_SENTINEL = "_unsloth_causal_conv1d_blocker" +@functools.lru_cache(1) def _is_rocm_torch_build() -> bool: + # Most official ROCm wheels include a local version suffix like +rocmX.Y. + # Some custom/source builds do not, so we fall back to runtime hints. try: - return "rocm" in str(importlib_version("torch")).lower() + torch_version_raw = str(importlib_version("torch")).lower() + if "rocm" in torch_version_raw: + if UNSLOTH_ENABLE_LOGGING: + logger.info( + "Unsloth: ROCm detection matched torch version tag (+rocm)." + ) + return True except Exception: - return False + pass + + # Environment hints commonly present on ROCm runtimes. + for key in ( + "ROCM_PATH", + "ROCM_HOME", + "HIP_PATH", + "HSA_PATH", + "HIP_VISIBLE_DEVICES", + "ROCR_VISIBLE_DEVICES", + ): + value = os.environ.get(key, "") + if isinstance(value, str) and value.strip(): + if UNSLOTH_ENABLE_LOGGING: + logger.info(f"Unsloth: ROCm detection matched environment key `{key}`.") + return True + + # Filesystem / driver hints for ROCm stacks. + for path in ( + Path("/opt/rocm"), + Path("/dev/kfd"), + Path("/sys/module/amdgpu"), + ): + try: + if path.exists(): + if UNSLOTH_ENABLE_LOGGING: + logger.info( + f"Unsloth: ROCm detection matched filesystem hint `{path}`." + ) + return True + except Exception: + continue + + if UNSLOTH_ENABLE_LOGGING: + logger.info("Unsloth: ROCm detection did not match any known hints.") + return False @contextlib.contextmanager From a2e3a8ee544d358d5f945c0af6c02d7ff9172e0f Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sat, 14 Feb 2026 03:35:41 -0800 Subject: [PATCH 10/63] Filter only amdgpu.ids fd2 noise during ROCm startup (#4054) Co-authored-by: Daniel Hanchen --- unsloth/import_fixes.py | 38 +++++++++++++++++++++++++++++++------- 1 file changed, 31 insertions(+), 7 deletions(-) diff --git a/unsloth/import_fixes.py b/unsloth/import_fixes.py index 63e99386d8..2dad8da30c 100644 --- a/unsloth/import_fixes.py +++ b/unsloth/import_fixes.py @@ -26,6 +26,7 @@ import textwrap import warnings import sys import functools +import tempfile # We cannot do from unsloth_zoo.log import logger since FBGEMM might cause seg faults. UNSLOTH_ENABLE_LOGGING = os.environ.get("UNSLOTH_ENABLE_LOGGING", "0") in ( @@ -1229,14 +1230,20 @@ def _is_rocm_torch_build() -> bool: @contextlib.contextmanager -def _suppress_stderr_fd(): +def _filter_stderr_fd( + suppressed_substrings = ("amdgpu.ids: No such file or directory",), +): + """ + Capture low-level fd=2 writes, drop only known noisy substrings, and replay + everything else after the protected block. + """ saved_stderr_fd = None - devnull_fd = None + temp_file = None redirected = False try: saved_stderr_fd = os.dup(2) - devnull_fd = os.open(os.devnull, os.O_WRONLY) - os.dup2(devnull_fd, 2) + temp_file = tempfile.TemporaryFile(mode = "w+b") + os.dup2(temp_file.fileno(), 2) redirected = True except Exception: redirected = False @@ -1244,14 +1251,31 @@ def _suppress_stderr_fd(): try: yield finally: + captured = b"" + if redirected and temp_file is not None: + try: + temp_file.flush() + temp_file.seek(0) + captured = temp_file.read() + except Exception: + captured = b"" if redirected and saved_stderr_fd is not None: try: os.dup2(saved_stderr_fd, 2) except Exception: pass - if devnull_fd is not None: + if captured and saved_stderr_fd is not None: try: - os.close(devnull_fd) + for raw_line in captured.splitlines(keepends = True): + line = raw_line.decode("utf-8", errors = "ignore") + if any(s in line for s in suppressed_substrings): + continue + os.write(saved_stderr_fd, raw_line) + except Exception: + pass + if temp_file is not None: + try: + temp_file.close() except Exception: pass if saved_stderr_fd is not None: @@ -1266,7 +1290,7 @@ def _suppress_hip_libdrm_ids_noise(): # Python-level stderr filters cannot intercept those writes. if not _is_rocm_torch_build(): return contextlib.nullcontext() - return _suppress_stderr_fd() + return _filter_stderr_fd() def _is_causal_conv1d_name(module_name: str) -> bool: From c577835d976c793ca8edc59156e1739fbdcec5cb Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sat, 14 Feb 2026 03:59:57 -0800 Subject: [PATCH 11/63] Wrap unsloth_zoo import with HIP amdgpu.ids filter (#4056) * Wrap unsloth_zoo import with HIP amdgpu.ids filter * Refactor ROCm ids filter helpers for readability * Rename ROCm ids filter helper and annotate call sites * Remove obsolete amdgpu ids filter alias * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --------- Co-authored-by: Daniel Hanchen Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- unsloth/__init__.py | 12 +++++--- unsloth/import_fixes.py | 66 +++++++++++++++++++++++------------------ 2 files changed, 45 insertions(+), 33 deletions(-) diff --git a/unsloth/__init__.py b/unsloth/__init__.py index 990e4b915f..e3ced412a5 100644 --- a/unsloth/__init__.py +++ b/unsloth/__init__.py @@ -29,7 +29,7 @@ from .import_fixes import ( fix_message_factory_issue, check_fbgemm_gpu_version, disable_broken_causal_conv1d, - _suppress_hip_libdrm_ids_noise, + _filter_rocm_amdgpu_ids_fd2_noise, torchvision_compatibility_check, fix_diffusers_warnings, fix_huggingface_hub, @@ -96,7 +96,9 @@ try: # os.system("pip install --upgrade --no-cache-dir --no-deps --user unsloth_zoo") # except: # raise ImportError("Unsloth: Please update unsloth_zoo via `pip install --upgrade --no-cache-dir --no-deps unsloth_zoo`") - import unsloth_zoo + # Filter native fd=2 amdgpu.ids noise during early unsloth_zoo import. + with _filter_rocm_amdgpu_ids_fd2_noise(): + import unsloth_zoo except PackageNotFoundError: raise ImportError( f"Unsloth: Please install unsloth_zoo via `pip install unsloth_zoo` then retry!" @@ -107,7 +109,8 @@ del PackageNotFoundError, importlib_version # Try importing PyTorch and check version try: - with _suppress_hip_libdrm_ids_noise(): + # Filter native fd=2 amdgpu.ids noise during torch import on ROCm. + with _filter_rocm_amdgpu_ids_fd2_noise(): import torch except ModuleNotFoundError: raise ImportError( @@ -117,7 +120,8 @@ except ModuleNotFoundError: except: raise -with _suppress_hip_libdrm_ids_noise(): +# Filter native fd=2 amdgpu.ids noise during early device detection import. +with _filter_rocm_amdgpu_ids_fd2_noise(): from unsloth_zoo.device_type import ( is_hip, get_device_type, diff --git a/unsloth/import_fixes.py b/unsloth/import_fixes.py index 2dad8da30c..f6e66a27f7 100644 --- a/unsloth/import_fixes.py +++ b/unsloth/import_fixes.py @@ -46,6 +46,8 @@ else: ) logger.setLevel(logging.WARNING) +_AMDGPU_IDS_MISSING_TEXT = "amdgpu.ids: No such file or directory" + def Version(version): try: @@ -203,7 +205,7 @@ if os.environ.get("UNSLOTH_ENABLE_LOGGING", "0") != "1": # Triton "df: No such file or directory" stderr noise sys.stderr.add_filter("df: No such file") # ROCm/libdrm missing ids table stderr noise on some AMD setups - sys.stderr.add_filter("amdgpu.ids: No such file or directory") + sys.stderr.add_filter(_AMDGPU_IDS_MISSING_TEXT) # Fix up AttributeError: 'MessageFactory' object has no attribute 'GetPrototype' @@ -1176,6 +1178,24 @@ def disable_torchcodec_if_broken(): CAUSAL_CONV1D_BROKEN = False _CAUSAL_CONV1D_PREFIX = "causal_conv1d" _CAUSAL_CONV1D_BLOCKER_SENTINEL = "_unsloth_causal_conv1d_blocker" +_ROCM_ENV_HINT_KEYS = ( + "ROCM_PATH", + "ROCM_HOME", + "HIP_PATH", + "HSA_PATH", + "HIP_VISIBLE_DEVICES", + "ROCR_VISIBLE_DEVICES", +) +_ROCM_PATH_HINTS = ( + Path("/opt/rocm"), + Path("/dev/kfd"), + Path("/sys/module/amdgpu"), +) + + +def _log_rocm_detection(message): + if UNSLOTH_ENABLE_LOGGING: + logger.info(message) @functools.lru_cache(1) @@ -1185,53 +1205,40 @@ def _is_rocm_torch_build() -> bool: try: torch_version_raw = str(importlib_version("torch")).lower() if "rocm" in torch_version_raw: - if UNSLOTH_ENABLE_LOGGING: - logger.info( - "Unsloth: ROCm detection matched torch version tag (+rocm)." - ) + _log_rocm_detection( + "Unsloth: ROCm detection matched torch version tag (+rocm)." + ) return True except Exception: pass # Environment hints commonly present on ROCm runtimes. - for key in ( - "ROCM_PATH", - "ROCM_HOME", - "HIP_PATH", - "HSA_PATH", - "HIP_VISIBLE_DEVICES", - "ROCR_VISIBLE_DEVICES", - ): + for key in _ROCM_ENV_HINT_KEYS: value = os.environ.get(key, "") if isinstance(value, str) and value.strip(): - if UNSLOTH_ENABLE_LOGGING: - logger.info(f"Unsloth: ROCm detection matched environment key `{key}`.") + _log_rocm_detection( + f"Unsloth: ROCm detection matched environment key `{key}`." + ) return True # Filesystem / driver hints for ROCm stacks. - for path in ( - Path("/opt/rocm"), - Path("/dev/kfd"), - Path("/sys/module/amdgpu"), - ): + for path in _ROCM_PATH_HINTS: try: if path.exists(): - if UNSLOTH_ENABLE_LOGGING: - logger.info( - f"Unsloth: ROCm detection matched filesystem hint `{path}`." - ) + _log_rocm_detection( + f"Unsloth: ROCm detection matched filesystem hint `{path}`." + ) return True except Exception: continue - if UNSLOTH_ENABLE_LOGGING: - logger.info("Unsloth: ROCm detection did not match any known hints.") + _log_rocm_detection("Unsloth: ROCm detection did not match any known hints.") return False @contextlib.contextmanager def _filter_stderr_fd( - suppressed_substrings = ("amdgpu.ids: No such file or directory",), + suppressed_substrings = (_AMDGPU_IDS_MISSING_TEXT,), ): """ Capture low-level fd=2 writes, drop only known noisy substrings, and replay @@ -1285,7 +1292,7 @@ def _filter_stderr_fd( pass -def _suppress_hip_libdrm_ids_noise(): +def _filter_rocm_amdgpu_ids_fd2_noise(): # ROCm/libdrm can emit amdgpu.ids missing errors via low-level fd=2 writes. # Python-level stderr filters cannot intercept those writes. if not _is_rocm_torch_build(): @@ -1425,7 +1432,8 @@ def disable_broken_causal_conv1d(): return try: - with _suppress_hip_libdrm_ids_noise(): + # Suppress only native fd=2 amdgpu.ids noise during causal_conv1d probe. + with _filter_rocm_amdgpu_ids_fd2_noise(): import causal_conv1d # noqa: F401 return From 9b83a1da4e98d90d8d4dd34c34a35d2b5d502d83 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sat, 14 Feb 2026 04:13:25 -0800 Subject: [PATCH 12/63] Wrap models import with ROCm amdgpu ids fd2 filter (#4057) Co-authored-by: Daniel Hanchen --- unsloth/__init__.py | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/unsloth/__init__.py b/unsloth/__init__.py index e3ced412a5..5587aa8c8e 100644 --- a/unsloth/__init__.py +++ b/unsloth/__init__.py @@ -305,8 +305,10 @@ elif DEVICE_TYPE == "xpu": # TODO: check triton for intel installed properly. pass -from .models import * -from .models import __version__ +# Filter native fd=2 amdgpu.ids noise during model import startup. +with _filter_rocm_amdgpu_ids_fd2_noise(): + from .models import * + from .models import __version__ from .save import * from .chat_templates import * from .tokenizer_utils import * From fd1034629ea3155dbc7deab0738ee25edd13cd17 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sat, 14 Feb 2026 21:52:31 -0800 Subject: [PATCH 13/63] Auto-configure AMDGPU_ASIC_ID_TABLE_PATH on ROCm startup (#4060) * Auto-configure AMDGPU_ASIC_ID_TABLE_PATH on ROCm startup * Remove ROCm fd2 amdgpu.ids noise filter wrappers * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Use PyPI bitsandbytes for amd extra to avoid malformed wheel URL * Add amd-preview extra for bitsandbytes continuous wheel channel * Keep amd extra on bitsandbytes>=0.49.1 and remove amd-preview --------- Co-authored-by: Daniel Hanchen Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- pyproject.toml | 5 +- unsloth/__init__.py | 37 +++++------- unsloth/import_fixes.py | 127 +++++++++++++++++++++------------------- 3 files changed, 84 insertions(+), 85 deletions(-) diff --git a/pyproject.toml b/pyproject.toml index f67f9ad121..bd6107626b 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -994,9 +994,8 @@ intel = [ ] amd = [ "unsloth[huggingfacenotorch]", - "bitsandbytes @ https://github.com/bitsandbytes-foundation/bitsandbytes/releases/download/continuous-release_main/bitsandbytes-1.33.7.preview-py3-none-manylinux_2_24_x86_64.whl ; ('linux' in sys_platform) and (platform_machine == 'AMD64' or platform_machine == 'x86_64')", - "bitsandbytes @ https://github.com/bitsandbytes-foundation/bitsandbytes/releases/download/continuous-release_main/bitsandbytes-1.33.7.preview-py3-none-win_amd64.whl ; (sys_platform == 'win32') and (platform_machine == 'AMD64' or platform_machine == 'x86_64')", - "bitsandbytes @ https://github.com/bitsandbytes-foundation/bitsandbytes/releases/download/continuous-release_main/bitsandbytes-1.33.7.preview-py3-none-manylinux_2_24_aarch64.whl ; ('linux' in sys_platform) and (platform_machine == 'aarch64')", + "bitsandbytes>=0.49.1 ; ('linux' in sys_platform) and (platform_machine == 'AMD64' or platform_machine == 'x86_64' or platform_machine == 'aarch64')", + "bitsandbytes>=0.49.1 ; (sys_platform == 'win32') and (platform_machine == 'AMD64' or platform_machine == 'x86_64')", ] [project.urls] diff --git a/unsloth/__init__.py b/unsloth/__init__.py index 5587aa8c8e..605df5a930 100644 --- a/unsloth/__init__.py +++ b/unsloth/__init__.py @@ -29,18 +29,21 @@ from .import_fixes import ( fix_message_factory_issue, check_fbgemm_gpu_version, disable_broken_causal_conv1d, - _filter_rocm_amdgpu_ids_fd2_noise, + configure_amdgpu_asic_id_table_path, torchvision_compatibility_check, fix_diffusers_warnings, fix_huggingface_hub, ) +# Configure libdrm ids table path early so ROCm can resolve AMD GPU names. +configure_amdgpu_asic_id_table_path() disable_broken_causal_conv1d() fix_message_factory_issue() check_fbgemm_gpu_version() torchvision_compatibility_check() fix_diffusers_warnings() fix_huggingface_hub() +del configure_amdgpu_asic_id_table_path del disable_broken_causal_conv1d del fix_message_factory_issue del check_fbgemm_gpu_version @@ -96,9 +99,7 @@ try: # os.system("pip install --upgrade --no-cache-dir --no-deps --user unsloth_zoo") # except: # raise ImportError("Unsloth: Please update unsloth_zoo via `pip install --upgrade --no-cache-dir --no-deps unsloth_zoo`") - # Filter native fd=2 amdgpu.ids noise during early unsloth_zoo import. - with _filter_rocm_amdgpu_ids_fd2_noise(): - import unsloth_zoo + import unsloth_zoo except PackageNotFoundError: raise ImportError( f"Unsloth: Please install unsloth_zoo via `pip install unsloth_zoo` then retry!" @@ -109,9 +110,7 @@ del PackageNotFoundError, importlib_version # Try importing PyTorch and check version try: - # Filter native fd=2 amdgpu.ids noise during torch import on ROCm. - with _filter_rocm_amdgpu_ids_fd2_noise(): - import torch + import torch except ModuleNotFoundError: raise ImportError( "Unsloth: Pytorch is not installed. Go to https://pytorch.org/.\n" @@ -120,16 +119,14 @@ except ModuleNotFoundError: except: raise -# Filter native fd=2 amdgpu.ids noise during early device detection import. -with _filter_rocm_amdgpu_ids_fd2_noise(): - from unsloth_zoo.device_type import ( - is_hip, - get_device_type, - DEVICE_TYPE, - DEVICE_TYPE_TORCH, - DEVICE_COUNT, - ALLOW_PREQUANTIZED_MODELS, - ) +from unsloth_zoo.device_type import ( + is_hip, + get_device_type, + DEVICE_TYPE, + DEVICE_TYPE_TORCH, + DEVICE_COUNT, + ALLOW_PREQUANTIZED_MODELS, +) # Fix other issues from .import_fixes import ( @@ -305,10 +302,8 @@ elif DEVICE_TYPE == "xpu": # TODO: check triton for intel installed properly. pass -# Filter native fd=2 amdgpu.ids noise during model import startup. -with _filter_rocm_amdgpu_ids_fd2_noise(): - from .models import * - from .models import __version__ +from .models import * +from .models import __version__ from .save import * from .chat_templates import * from .tokenizer_utils import * diff --git a/unsloth/import_fixes.py b/unsloth/import_fixes.py index f6e66a27f7..f11328f7e9 100644 --- a/unsloth/import_fixes.py +++ b/unsloth/import_fixes.py @@ -16,7 +16,6 @@ import os import importlib.abc import importlib.machinery import importlib.util -import contextlib from pathlib import Path from importlib.metadata import version as importlib_version from packaging.version import Version as TrueVersion @@ -26,7 +25,6 @@ import textwrap import warnings import sys import functools -import tempfile # We cannot do from unsloth_zoo.log import logger since FBGEMM might cause seg faults. UNSLOTH_ENABLE_LOGGING = os.environ.get("UNSLOTH_ENABLE_LOGGING", "0") in ( @@ -1191,6 +1189,13 @@ _ROCM_PATH_HINTS = ( Path("/dev/kfd"), Path("/sys/module/amdgpu"), ) +_AMDGPU_ASIC_ID_TABLE_PATH_ENV = "AMDGPU_ASIC_ID_TABLE_PATH" +_AMDGPU_ASIC_ID_CANDIDATE_PATHS = ( + Path("/usr/share/libdrm/amdgpu.ids"), + Path("/usr/local/share/libdrm/amdgpu.ids"), + Path("/opt/rocm/share/libdrm/amdgpu.ids"), + Path("/opt/amdgpu/share/libdrm/amdgpu.ids"), +) def _log_rocm_detection(message): @@ -1236,68 +1241,70 @@ def _is_rocm_torch_build() -> bool: return False -@contextlib.contextmanager -def _filter_stderr_fd( - suppressed_substrings = (_AMDGPU_IDS_MISSING_TEXT,), -): - """ - Capture low-level fd=2 writes, drop only known noisy substrings, and replay - everything else after the protected block. - """ - saved_stderr_fd = None - temp_file = None - redirected = False +def _iter_amdgpu_asic_id_table_candidates(): + # Try torch-adjacent ids table paths first without importing torch. try: - saved_stderr_fd = os.dup(2) - temp_file = tempfile.TemporaryFile(mode = "w+b") - os.dup2(temp_file.fileno(), 2) - redirected = True + torch_spec = importlib.util.find_spec("torch") except Exception: - redirected = False + torch_spec = None - try: - yield - finally: - captured = b"" - if redirected and temp_file is not None: - try: - temp_file.flush() - temp_file.seek(0) - captured = temp_file.read() - except Exception: - captured = b"" - if redirected and saved_stderr_fd is not None: - try: - os.dup2(saved_stderr_fd, 2) - except Exception: - pass - if captured and saved_stderr_fd is not None: - try: - for raw_line in captured.splitlines(keepends = True): - line = raw_line.decode("utf-8", errors = "ignore") - if any(s in line for s in suppressed_substrings): - continue - os.write(saved_stderr_fd, raw_line) - except Exception: - pass - if temp_file is not None: - try: - temp_file.close() - except Exception: - pass - if saved_stderr_fd is not None: - try: - os.close(saved_stderr_fd) - except Exception: - pass + roots = [] + if torch_spec is not None: + if torch_spec.origin: + roots.append(Path(torch_spec.origin).resolve().parent) + if torch_spec.submodule_search_locations: + for location in torch_spec.submodule_search_locations: + roots.append(Path(location).resolve()) + + seen = set() + for root in roots: + for candidate in ( + root / "share" / "libdrm" / "amdgpu.ids", + root.parent / "share" / "libdrm" / "amdgpu.ids", + root.parent.parent / "share" / "libdrm" / "amdgpu.ids", + ): + candidate_str = str(candidate) + if candidate_str in seen: + continue + seen.add(candidate_str) + yield candidate + + for candidate in _AMDGPU_ASIC_ID_CANDIDATE_PATHS: + candidate_str = str(candidate) + if candidate_str in seen: + continue + seen.add(candidate_str) + yield candidate -def _filter_rocm_amdgpu_ids_fd2_noise(): - # ROCm/libdrm can emit amdgpu.ids missing errors via low-level fd=2 writes. - # Python-level stderr filters cannot intercept those writes. +def configure_amdgpu_asic_id_table_path(): + # Honor an existing valid user-provided path. + configured = os.environ.get(_AMDGPU_ASIC_ID_TABLE_PATH_ENV, "").strip() + if configured: + configured_path = Path(configured) + try: + if configured_path.is_file(): + return str(configured_path) + except Exception: + pass + + # Only attempt this on ROCm-like environments. if not _is_rocm_torch_build(): - return contextlib.nullcontext() - return _filter_stderr_fd() + return None + + for candidate in _iter_amdgpu_asic_id_table_candidates(): + try: + if candidate.is_file(): + os.environ[_AMDGPU_ASIC_ID_TABLE_PATH_ENV] = str(candidate) + if UNSLOTH_ENABLE_LOGGING: + logger.info( + f"Unsloth: Set {_AMDGPU_ASIC_ID_TABLE_PATH_ENV}={candidate}" + ) + return str(candidate) + except Exception: + continue + + return None def _is_causal_conv1d_name(module_name: str) -> bool: @@ -1432,9 +1439,7 @@ def disable_broken_causal_conv1d(): return try: - # Suppress only native fd=2 amdgpu.ids noise during causal_conv1d probe. - with _filter_rocm_amdgpu_ids_fd2_noise(): - import causal_conv1d # noqa: F401 + import causal_conv1d # noqa: F401 return except Exception as error: From 04fb97aee2bc921d48b302600dfe4928a9b82741 Mon Sep 17 00:00:00 2001 From: anonymous dev <44989609+devchilll@users.noreply.github.com> Date: Sat, 14 Feb 2026 22:13:07 -0800 Subject: [PATCH 14/63] [FIX] Move labels to logits device in cross-entropy loss for multi-GPU support (#4041) (#4059) When using device_map='balanced' with multiple GPUs, the labels tensor may reside on a different device than the logits/losses tensors. This causes a RuntimeError at the masked_fill_ call in the chunked cross-entropy forward path. Fix: explicitly move labels to the same device as logits at the start of Fast_CrossEntropyLoss.forward(). This is a no-op on single-GPU setups. Fixes #4041 --- unsloth/kernels/cross_entropy_loss.py | 1 + 1 file changed, 1 insertion(+) diff --git a/unsloth/kernels/cross_entropy_loss.py b/unsloth/kernels/cross_entropy_loss.py index fbb14013ff..cdd5b1cf12 100644 --- a/unsloth/kernels/cross_entropy_loss.py +++ b/unsloth/kernels/cross_entropy_loss.py @@ -301,6 +301,7 @@ class Fast_CrossEntropyLoss(torch.autograd.Function): vocab_size: int n_rows, vocab_size = logits.shape device = logits.device + labels = labels.to(device) div, mod = divmod(vocab_size, MAX_FUSED_SIZE) n_chunks: int = div + (mod != 0) From 529436a007ccbe8144f0167070a69298d72ae962 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sat, 14 Feb 2026 22:14:05 -0800 Subject: [PATCH 15/63] Silence Apex Aiter RoPE warning unless logging is enabled (#4058) * Silence Apex Aiter RoPE warning unless logging is enabled * Update unsloth/import_fixes.py Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> --------- Co-authored-by: Daniel Hanchen Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> --- unsloth/import_fixes.py | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/unsloth/import_fixes.py b/unsloth/import_fixes.py index f11328f7e9..6d3df0b470 100644 --- a/unsloth/import_fixes.py +++ b/unsloth/import_fixes.py @@ -97,7 +97,7 @@ class HidePrintMessage: return getattr(self._original_stream, name) -if os.environ.get("UNSLOTH_ENABLE_LOGGING", "0") != "1": +if not UNSLOTH_ENABLE_LOGGING: import sys # Apply to stderr for FBGEMM and CUTLASS errors @@ -204,6 +204,13 @@ if os.environ.get("UNSLOTH_ENABLE_LOGGING", "0") != "1": sys.stderr.add_filter("df: No such file") # ROCm/libdrm missing ids table stderr noise on some AMD setups sys.stderr.add_filter(_AMDGPU_IDS_MISSING_TEXT) + # Apex ROCm fused RoPE backend selection warning when Aiter is enabled. + warnings.filterwarnings( + "ignore", + message = r"^Aiter backend is selected for fused RoPE\.?", + category = UserWarning, + module = r"^apex\.transformer\.functional\.fused_rope$", + ) # Fix up AttributeError: 'MessageFactory' object has no attribute 'GetPrototype' From e9e8ad2c5d68797e35431852d7d896785c3e8310 Mon Sep 17 00:00:00 2001 From: nole69 <89107458+nole69@users.noreply.github.com> Date: Sun, 15 Feb 2026 01:09:40 -0800 Subject: [PATCH 16/63] [FIX] Move loss and n_items to logits device in fast_cross_entropy_loss loss for multi-GPU support (#4063) * bug fix for multi-GPU * Apply suggestion from @gemini-code-assist[bot] Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> --------- Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> --- unsloth/kernels/cross_entropy_loss.py | 3 +++ 1 file changed, 3 insertions(+) diff --git a/unsloth/kernels/cross_entropy_loss.py b/unsloth/kernels/cross_entropy_loss.py index cdd5b1cf12..d92229314f 100644 --- a/unsloth/kernels/cross_entropy_loss.py +++ b/unsloth/kernels/cross_entropy_loss.py @@ -438,6 +438,7 @@ def fast_cross_entropy_loss( batch, seq_len, d = logits.shape assert labels.shape == (batch, seq_len) + device = logits.device loss = Fast_CrossEntropyLoss.apply( logits.view(batch * seq_len, d), labels.view(-1), @@ -446,6 +447,8 @@ def fast_cross_entropy_loss( ) if n_items is None: n_items = torch.count_nonzero(labels != -100) + if torch.is_tensor(n_items): + n_items = n_items.to(device) return loss.sum() / n_items From 9a61677123faf258a34da0af8de339038807e662 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sun, 15 Feb 2026 19:14:19 -0800 Subject: [PATCH 17/63] Raise ImportError on stable torch/torchvision mismatch (#4065) * Raise ImportError for stable torchvision mismatches * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Remove torchvision compatibility tests from PR scope --------- Co-authored-by: Daniel Hanchen Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- unsloth/import_fixes.py | 14 ++++---------- 1 file changed, 4 insertions(+), 10 deletions(-) diff --git a/unsloth/import_fixes.py b/unsloth/import_fixes.py index 6d3df0b470..c7914f39a1 100644 --- a/unsloth/import_fixes.py +++ b/unsloth/import_fixes.py @@ -668,7 +668,6 @@ def torchvision_compatibility_check(): # Try known table first, then fall back to formula for forward compatibility required = TORCH_TORCHVISION_COMPAT.get((torch_major, torch_minor)) - is_in_known_table = required is not None if required is None: required = _infer_required_torchvision(torch_major, torch_minor) @@ -705,15 +704,10 @@ def torchvision_compatibility_check(): t in torchvision_version_raw for t in _pre_tags ) - # Downgrade to warning for custom/source/pre-release builds or formula-predicted - if is_custom or is_prerelease or not is_in_known_table: - reason = ( - "custom/source build" - if is_custom - else "pre-release build" - if is_prerelease - else "newer torch version" - ) + # Only downgrade to warning for custom/source or prerelease builds. + # Stable mismatches should fail fast to prevent runtime operator errors. + if is_custom or is_prerelease: + reason = "custom/source build" if is_custom else "pre-release build" logger.warning( f"{message}\n" f"Detected a {reason}. " From 971539e88f8ddceb7cb998bdfa20789f327a2619 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sun, 15 Feb 2026 19:36:16 -0800 Subject: [PATCH 18/63] Add torchvision upgrade hint to mismatch ImportError (#4067) Co-authored-by: Daniel Hanchen --- unsloth/import_fixes.py | 1 + 1 file changed, 1 insertion(+) diff --git a/unsloth/import_fixes.py b/unsloth/import_fixes.py index c7914f39a1..a86ad84779 100644 --- a/unsloth/import_fixes.py +++ b/unsloth/import_fixes.py @@ -689,6 +689,7 @@ def torchvision_compatibility_check(): f"Unsloth: torch=={torch_version_raw} requires " f"torchvision>={required_tv_str}, " f"but found torchvision=={torchvision_version_raw}. " + f'Try updating torchvision via `pip install --upgrade "torchvision>={required_tv_str}"`. ' f"Please refer to https://pytorch.org/get-started/previous-versions/ " f"for more information." ) From 37258cba82c913da32ec75662e0da5b7cf8d6d8d Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sun, 15 Feb 2026 22:09:29 -0800 Subject: [PATCH 19/63] Guard optional vLLM imports when extension is broken (#4068) * Guard optional vLLM imports when extension is broken * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Remove vLLM import guard tests from PR scope * Block broken vLLM imports like causal_conv1d --------- Co-authored-by: Daniel Hanchen Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- unsloth/__init__.py | 3 + unsloth/dataprep/synthetic.py | 20 +++-- unsloth/import_fixes.py | 133 ++++++++++++++++++++++++++++++++- unsloth/models/loader_utils.py | 8 +- unsloth/models/vision.py | 5 -- 5 files changed, 154 insertions(+), 15 deletions(-) diff --git a/unsloth/__init__.py b/unsloth/__init__.py index 605df5a930..7932fe966d 100644 --- a/unsloth/__init__.py +++ b/unsloth/__init__.py @@ -29,6 +29,7 @@ from .import_fixes import ( fix_message_factory_issue, check_fbgemm_gpu_version, disable_broken_causal_conv1d, + disable_broken_vllm, configure_amdgpu_asic_id_table_path, torchvision_compatibility_check, fix_diffusers_warnings, @@ -38,6 +39,7 @@ from .import_fixes import ( # Configure libdrm ids table path early so ROCm can resolve AMD GPU names. configure_amdgpu_asic_id_table_path() disable_broken_causal_conv1d() +disable_broken_vllm() fix_message_factory_issue() check_fbgemm_gpu_version() torchvision_compatibility_check() @@ -45,6 +47,7 @@ fix_diffusers_warnings() fix_huggingface_hub() del configure_amdgpu_asic_id_table_path del disable_broken_causal_conv1d +del disable_broken_vllm del fix_message_factory_issue del check_fbgemm_gpu_version del torchvision_compatibility_check diff --git a/unsloth/dataprep/synthetic.py b/unsloth/dataprep/synthetic.py index 7735b1ecd2..612c531f47 100644 --- a/unsloth/dataprep/synthetic.py +++ b/unsloth/dataprep/synthetic.py @@ -27,11 +27,6 @@ import torch import gc import time import re -from unsloth_zoo.vllm_utils import ( - load_vllm, - patch_vllm, - delete_vllm, -) from unsloth_zoo.log import logger import numpy as np @@ -40,6 +35,16 @@ from .synthetic_configs import ( ) +def _load_vllm_utils(): + from unsloth_zoo.vllm_utils import ( + load_vllm, + patch_vllm, + delete_vllm, + ) + + return load_vllm, patch_vllm, delete_vllm + + def terminate_tree(proc: subprocess.Popen, timeout = 15): if proc is None or proc.poll() is not None: return @@ -182,6 +187,8 @@ class SyntheticDataKit: model_name, token = token, ) + load_vllm, patch_vllm, delete_vllm = _load_vllm_utils() + self._delete_vllm = delete_vllm patch_vllm(debug = False) engine_args = load_vllm( model_name = model_name, @@ -364,7 +371,8 @@ class SyntheticDataKit: gc.collect() # Delete vLLM module as well - delete_vllm(llm = None) + if hasattr(self, "_delete_vllm"): + self._delete_vllm(llm = None) def __enter__(self): return self diff --git a/unsloth/import_fixes.py b/unsloth/import_fixes.py index a86ad84779..0581cf7bdf 100644 --- a/unsloth/import_fixes.py +++ b/unsloth/import_fixes.py @@ -421,14 +421,18 @@ def fix_vllm_guided_decoding_params(): # trl still wants to use GuidedDecodingParams. This is a temporary patch till trl updates try: import vllm - except ImportError as e: + except (ImportError, OSError) as e: _maybe_raise_vllm_transformers_mismatch(e) + if disable_broken_vllm(e): + return raise try: from vllm.sampling_params import GuidedDecodingParams - except ImportError as e: + except (ImportError, OSError) as e: _maybe_raise_vllm_transformers_mismatch(e) + if disable_broken_vllm(e): + return if not hasattr(vllm, "sampling_params") or not hasattr( vllm.sampling_params, "StructuredOutputsParams" ): @@ -1028,7 +1032,7 @@ def fix_vllm_pdl_blackwell(): def _spec_exists(name): try: return importlib.util.find_spec(name) is not None - except (ModuleNotFoundError, ValueError): + except (ImportError, OSError, ModuleNotFoundError, ValueError): return False # Check if vLLM has the PDL-related modules before doing internet check @@ -1178,6 +1182,9 @@ def disable_torchcodec_if_broken(): CAUSAL_CONV1D_BROKEN = False _CAUSAL_CONV1D_PREFIX = "causal_conv1d" _CAUSAL_CONV1D_BLOCKER_SENTINEL = "_unsloth_causal_conv1d_blocker" +VLLM_BROKEN = False +_VLLM_PREFIX = "vllm" +_VLLM_BLOCKER_SENTINEL = "_unsloth_vllm_blocker" _ROCM_ENV_HINT_KEYS = ( "ROCM_PATH", "ROCM_HOME", @@ -1315,6 +1322,10 @@ def _is_causal_conv1d_name(module_name: str) -> bool: ) +def _is_vllm_name(module_name: str) -> bool: + return module_name == _VLLM_PREFIX or module_name.startswith(_VLLM_PREFIX + ".") + + def _resolve_module_name(module_name, package): if not isinstance(module_name, str): return module_name @@ -1344,6 +1355,27 @@ def _is_broken_causal_conv1d_error(error) -> bool: return False +def _is_broken_vllm_error(error) -> bool: + checked = set() + current = error + while current is not None and id(current) not in checked: + checked.add(id(current)) + message = str(current).lower() + if ( + ("vllm/_c" in message or "vllm._c" in message) + and ( + "undefined symbol" in message + or "cannot open shared object file" in message + or ".so:" in message + ) + ) or ("vllm" in message and "undefined symbol" in message): + return True + current = getattr(current, "__cause__", None) or getattr( + current, "__context__", None + ) + return False + + class _CausalConv1dImportBlockerLoader(importlib.abc.Loader): __slots__ = ("module_name",) @@ -1373,6 +1405,35 @@ class _CausalConv1dImportBlockerFinder(importlib.abc.MetaPathFinder): ) +class _VllmImportBlockerLoader(importlib.abc.Loader): + __slots__ = ("module_name",) + + def __init__(self, module_name): + self.module_name = module_name + + def create_module(self, spec): + return None + + def exec_module(self, module): + raise ModuleNotFoundError(f"No module named '{self.module_name}'") + + +class _VllmImportBlockerFinder(importlib.abc.MetaPathFinder): + __slots__ = (_VLLM_BLOCKER_SENTINEL,) + + def __init__(self): + setattr(self, _VLLM_BLOCKER_SENTINEL, True) + + def find_spec(self, fullname, path = None, target = None): + if not VLLM_BROKEN or not _is_vllm_name(fullname): + return None + return importlib.machinery.ModuleSpec( + name = fullname, + loader = _VllmImportBlockerLoader(fullname), + is_package = fullname == _VLLM_PREFIX, + ) + + def _patch_find_spec_for_causal_conv1d(): current_find_spec = importlib.util.find_spec if getattr(current_find_spec, "_unsloth_causal_conv1d_find_spec_patch", False): @@ -1390,6 +1451,23 @@ def _patch_find_spec_for_causal_conv1d(): importlib.util.find_spec = _blocked_find_spec +def _patch_find_spec_for_vllm(): + current_find_spec = importlib.util.find_spec + if getattr(current_find_spec, "_unsloth_vllm_find_spec_patch", False): + return + + def _blocked_find_spec(name, package = None): + resolved_name = _resolve_module_name(name, package) + if VLLM_BROKEN and isinstance(resolved_name, str): + if _is_vllm_name(resolved_name): + return None + return current_find_spec(name, package) + + _blocked_find_spec._unsloth_vllm_find_spec_patch = True + _blocked_find_spec._unsloth_original_find_spec = current_find_spec + importlib.util.find_spec = _blocked_find_spec + + def _install_causal_conv1d_blocker(): _patch_find_spec_for_causal_conv1d() for finder in sys.meta_path: @@ -1398,12 +1476,61 @@ def _install_causal_conv1d_blocker(): sys.meta_path.insert(0, _CausalConv1dImportBlockerFinder()) +def _install_vllm_blocker(): + _patch_find_spec_for_vllm() + for finder in sys.meta_path: + if getattr(finder, _VLLM_BLOCKER_SENTINEL, False): + return + sys.meta_path.insert(0, _VllmImportBlockerFinder()) + + def _clear_causal_conv1d_modules(): for module_name in list(sys.modules): if _is_causal_conv1d_name(module_name): sys.modules.pop(module_name, None) +def _clear_vllm_modules(): + for module_name in list(sys.modules): + if _is_vllm_name(module_name): + sys.modules.pop(module_name, None) + + +def disable_broken_vllm(error = None): + """Disable vLLM dynamically when its shared library is ABI-broken.""" + global VLLM_BROKEN + if VLLM_BROKEN: + _install_vllm_blocker() + return True + + failure = error + if failure is None: + try: + if importlib.util.find_spec("vllm") is None: + return False + except Exception: + return False + + try: + import vllm # noqa: F401 + + return False + except Exception as import_error: + failure = import_error + + if not _is_broken_vllm_error(failure): + return False + + VLLM_BROKEN = True + _clear_vllm_modules() + _install_vllm_blocker() + logger.warning( + "Unsloth: Detected broken vLLM binary extension; " + "disabling vLLM imports and continuing import." + ) + return True + + def _disable_transformers_causal_conv1d(): try: import transformers.utils.import_utils as tf_import_utils diff --git a/unsloth/models/loader_utils.py b/unsloth/models/loader_utils.py index 01d221c725..f17ebc53e3 100644 --- a/unsloth/models/loader_utils.py +++ b/unsloth/models/loader_utils.py @@ -31,7 +31,6 @@ from .mapper import ( from transformers import __version__ as transformers_version from unsloth.models._utils import TorchAOConfig from unsloth_zoo.utils import Version -from unsloth_zoo.vllm_utils import _get_torchao_fp8_config import gc transformers_version = Version(transformers_version) @@ -49,6 +48,13 @@ BAD_MAPPINGS = { } +def _get_torchao_fp8_config(fp8_mode): + # Import lazily so an optional, broken vLLM install does not break plain `import unsloth`. + from unsloth_zoo.vllm_utils import _get_torchao_fp8_config as _impl + + return _impl(fp8_mode) + + def _get_env_int(keys): for key in keys: value = os.environ.get(key) diff --git a/unsloth/models/vision.py b/unsloth/models/vision.py index 9e292a2849..00d8a3b7f2 100644 --- a/unsloth/models/vision.py +++ b/unsloth/models/vision.py @@ -126,11 +126,6 @@ _compile_config = CompileConfig( ) _compile_config.disable = True # Must set manually -from unsloth_zoo.vllm_utils import ( - convert_lora_modules, - return_lora_modules, -) - try: torch_compiler_set_stance = torch.compiler.set_stance except: From ee311a649832e759f9df325ce9a3253b23f71221 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sun, 15 Feb 2026 23:02:12 -0800 Subject: [PATCH 20/63] Add reinstall command to broken vLLM warning (#4070) * Add vLLM reinstall command to broken-extension warning * Apply suggestion from @danielhanchen --------- Co-authored-by: Daniel Hanchen --- unsloth/import_fixes.py | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/unsloth/import_fixes.py b/unsloth/import_fixes.py index 0581cf7bdf..6e5826f6b8 100644 --- a/unsloth/import_fixes.py +++ b/unsloth/import_fixes.py @@ -1526,7 +1526,9 @@ def disable_broken_vllm(error = None): _install_vllm_blocker() logger.warning( "Unsloth: Detected broken vLLM binary extension; " - "disabling vLLM imports and continuing import." + "disabling vLLM imports and continuing import.\n" + "Please reinstall via `uv pip install unsloth vllm torchvision torchaudio " + "--torch-backend=auto`." ) return True From 2a56922f881f0a30ac2d222a1be344e5fd927133 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sun, 15 Feb 2026 23:16:17 -0800 Subject: [PATCH 21/63] Fix regressions from security PRs #4042, #4044, and #4045 (#4062) * Fix security-regression fallout in chat templates and PDL patching * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Drop security regression test files from PR scope * Apply suggestion from @danielhanchen --------- Co-authored-by: Daniel Hanchen Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- unsloth/chat_templates.py | 128 +++++++++++++++++++++++--------------- unsloth/import_fixes.py | 23 +++++-- 2 files changed, 98 insertions(+), 53 deletions(-) diff --git a/unsloth/chat_templates.py b/unsloth/chat_templates.py index a17a6f6299..35eb871529 100644 --- a/unsloth/chat_templates.py +++ b/unsloth/chat_templates.py @@ -1959,41 +1959,61 @@ def _parse_combined_prompt(combined_prompt, dataset): def _create_formatter(possible_columns, final_optional_prompts, user_column_name): - # Start final prompt! - function = ["def __combined_prompt_processor__(examples):"] - columns = list(set(possible_columns)) - for column in columns: - function.append(f"{' '*4}{column}__ = examples['{column}']") - function.append(f"{' '*4}texts = []") - function.append(f"{' '*4}for ({', '.join(columns)}) in zip({', '.join(f'{x}__' for x in columns)}):") - - # Add optional tags as well! - final_prompt = "" - formatter = [] + columns = list(dict.fromkeys(possible_columns)) + merged_prompt_parts = [] + formatter_templates = [] for j, optional_prompt in enumerate(final_optional_prompts): if type(optional_prompt) is str: - columns = re.findall(r"\{(.+?)\}", optional_prompt) - formatter += columns - # Must escape \n \r - final_prompt += optional_prompt.encode("unicode-escape").decode("utf-8").replace("'", "\\'").replace('"', '\\"') - else: - where, prompt = optional_prompt - # Strip [[...]] - # Must escape \n \r - prompt = prompt[2:-2].encode("unicode-escape").decode("utf-8").replace("'", "\\'").replace('"', '\\"') - columns = re.findall(r"\{(.+?)\}", prompt) - x = f"__optional_{j}__" - prompt = f"{' '*8}{x} = '{prompt}'.format({', '.join(f'{x} = {x}' for x in columns)}) if {columns[0]} else ''" - function.append(prompt) - formatter.append(x) - final_prompt += "{" + x + "}" + needed_columns = re.findall(r"\{(.+?)\}", optional_prompt) + formatter_templates.append(("required", optional_prompt, needed_columns)) + merged_prompt_parts.append(optional_prompt) + continue - function.insert(1, f"{' '*4}__combined_prompt__ = '{final_prompt}'") - function.append(f"{' '*8}texts.append("\ - f"__combined_prompt__.format({', '.join(f'{x} = {x}' for x in formatter)}))") - function.append(f"{' '*4}return " + "{ " + f"'{user_column_name}' : texts" + " }") - return "\n".join(function) + _, prompt = optional_prompt + prompt = prompt[2:-2] + needed_columns = re.findall(r"\{(.+?)\}", prompt) + if len(needed_columns) == 0: + raise IndexError("Unsloth: Optional [[...]] blocks must contain at least 1 {column}.") + optional_name = f"__optional_{j}__" + formatter_templates.append(("optional", optional_name, prompt, needed_columns)) + merged_prompt_parts.append("{" + optional_name + "}") + + merged_prompt = "".join(merged_prompt_parts) + + def __combined_prompt_processor__(examples): + if len(examples) == 0: + return {user_column_name: []} + + first_key = next(iter(examples.keys()), None) + if first_key is None: + return {user_column_name: []} + n_rows = len(examples[first_key]) + + texts = [] + for row_idx in range(n_rows): + row_values = {column: examples[column][row_idx] for column in columns} + formatter_values = {} + + for formatter_template in formatter_templates: + if formatter_template[0] == "required": + _, _, needed_columns = formatter_template + for column in needed_columns: + formatter_values[column] = row_values[column] + continue + + _, optional_name, prompt, needed_columns = formatter_template + if row_values[needed_columns[0]] not in (None, ""): + prompt_values = {column: row_values[column] for column in needed_columns} + formatter_values[optional_name] = prompt.format(**prompt_values) + else: + formatter_values[optional_name] = "" + + texts.append(merged_prompt.format(**formatter_values)) + + return {user_column_name: texts} + + return __combined_prompt_processor__ def to_sharegpt( @@ -2025,13 +2045,17 @@ def to_sharegpt( raise TypeError("Unsloth: Your dataset is probably already in ShareGPT format!") possible_columns, final_optional_prompts = _parse_combined_prompt(merged_prompt, dataset) - function = _create_formatter(possible_columns, final_optional_prompts, merged_column_name) - exec(function, globals()) - dataset = dataset.map(__combined_prompt_processor__, batched = True, desc = "Merging columns") + formatter = _create_formatter(possible_columns, final_optional_prompts, merged_column_name) + dataset = dataset.map(formatter, batched = True, desc = "Merging columns") def __convert_to_sharegpt__(examples): users = examples[merged_column_name] assistants = examples[output_column_name] + if len(users) != len(assistants): + raise ValueError( + "Unsloth: Input and output columns must have matching batch lengths. " + f"Got {len(users)} {merged_column_name} rows and {len(assistants)} {output_column_name} rows." + ) texts = [ [ {"from" : "human", "value" : str(user) }, @@ -2062,19 +2086,18 @@ def to_sharegpt( dataset = concatenate_datasets(all_shuffled, axis = 1) # Combine them into 1 - function = "def __combine_conversations__(examples):\n" n_extensions += 1 - for j in range(n_extensions): - function += f"{' '*4}conversations{j}__ = examples['conversations{j}']\n" - function += f"{' '*4}convos = []\n" - function += f"{' '*4}for ({', '.join(f'conversations{j}' for j in range(n_extensions))}) "\ - f"in zip({', '.join(f'conversations{j}__' for j in range(n_extensions))}):\n" - function += f"{' '*8}convos.append("\ - f"{'+'.join(f'conversations{j}' for j in range(n_extensions))})\n" - function += f"{' '*4}return " + "{ " + "'conversations' : convos" + " }" + conversation_columns = [f"conversations{j}" for j in range(n_extensions)] + def __combine_conversations__(examples): + columns = [examples[column] for column in conversation_columns] + convos = [] + for conversations in zip(*columns): + merged_conversation = [] + for conversation in conversations: + merged_conversation.extend(conversation) + convos.append(merged_conversation) + return {"conversations" : convos} - # Map function - exec(function, globals()) dataset = dataset.map( __combine_conversations__, batched = True, @@ -2682,16 +2705,23 @@ def test_hf_gguf_equivalence(tokenizer, gguf_model = "./model-unsloth.F16.gguf") if tokenizer.chat_template is not None: prompt = tokenizer.apply_chat_template(messages, tokenize = False, add_generation_prompt = True) - prompt = prompt.replace("'", "") # Subprocess does not like '' prompt = remove_special_tokens(tokenizer, prompt) prompts.append(prompt) for prompt in prompts: - command = f"./llama.cpp/llama-cli -m {gguf_model} -n 0 --temp 0.0 --verbose-prompt "\ - f"--check-tensors -p '{prompt}'" + # Use a list of args with shell=False so prompt content is passed literally. + command = [ + "./llama.cpp/llama-cli", + "-m", gguf_model, + "-n", "0", + "--temp", "0.0", + "--verbose-prompt", + "--check-tensors", + "-p", prompt, + ] datas = [] - with subprocess.Popen(command, shell = True, stdout = subprocess.PIPE, stderr = subprocess.STDOUT, bufsize = 1) as sp: + with subprocess.Popen(command, shell = False, stdout = subprocess.PIPE, stderr = subprocess.STDOUT, bufsize = 1) as sp: for line in sp.stdout: datas.append(line.decode("utf-8", errors = "replace")) gguf_tokens = "".join(datas) diff --git a/unsloth/import_fixes.py b/unsloth/import_fixes.py index 6e5826f6b8..67e4a26ffb 100644 --- a/unsloth/import_fixes.py +++ b/unsloth/import_fixes.py @@ -1045,10 +1045,10 @@ def fix_vllm_pdl_blackwell(): return # Check if vLLM version includes the fix - VLLM_PDL_FIX_VERSION = "0.13.2" + VLLM_PDL_FIX_VERSION = "0.15.0" try: vllm_version = Version(importlib_version("vllm")) - if vllm_version > Version(VLLM_PDL_FIX_VERSION): + if vllm_version >= Version(VLLM_PDL_FIX_VERSION): logger.info( f"Unsloth: SM100 ({sm100_gpu_name}) detected but vLLM {vllm_version} " f"should include PDL fix - skipping workaround" @@ -1066,6 +1066,12 @@ def fix_vllm_pdl_blackwell(): return False patched = [] + patched_names = set() + + def _record_patch(name): + if name not in patched_names: + patched.append(name) + patched_names.add(name) # First, patch the source module (utils.py) where supports_pdl is defined. # This is critical because supports_pdl uses @lru_cache - we must clear the @@ -1077,7 +1083,7 @@ def fix_vllm_pdl_blackwell(): if hasattr(original_fn, "cache_clear"): original_fn.cache_clear() utils_module.supports_pdl = fake_supports_pdl - patched.append("utils") + _record_patch("utils") except (ImportError, ModuleNotFoundError, AttributeError): pass @@ -1094,10 +1100,19 @@ def fix_vllm_pdl_blackwell(): module = importlib.import_module(path) if hasattr(module, "supports_pdl"): module.supports_pdl = fake_supports_pdl - patched.append(name) + _record_patch(name) except (ImportError, ModuleNotFoundError, AttributeError): pass + # Patch any additional already-loaded triton ops consumers that expose supports_pdl. + for module_name, module in tuple(sys.modules.items()): + if not module_name.startswith("vllm.lora.ops.triton_ops."): + continue + if module is None or not hasattr(module, "supports_pdl"): + continue + module.supports_pdl = fake_supports_pdl + _record_patch(module_name.rsplit(".", 1)[-1]) + if patched: logger.info( f"Unsloth: Applied PDL fix for SM100 ({sm100_gpu_name}) - " From e43f6e96c7f53840553cb5afe4915e1a7b4029a8 Mon Sep 17 00:00:00 2001 From: Datta Nimmaturi Date: Mon, 16 Feb 2026 14:51:30 +0530 Subject: [PATCH 22/63] [Feat] FP8 per tensor quant support (#4043) * FP8 per tensor quant support * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --------- Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- unsloth/kernels/fp8.py | 106 +++++++++++++++++++++-------------------- 1 file changed, 54 insertions(+), 52 deletions(-) diff --git a/unsloth/kernels/fp8.py b/unsloth/kernels/fp8.py index e9f9161709..0b073d59d6 100644 --- a/unsloth/kernels/fp8.py +++ b/unsloth/kernels/fp8.py @@ -95,8 +95,11 @@ def weight_dequant_block( def weight_dequant(x: torch.Tensor, s: torch.Tensor, dtype = torch.bfloat16): - if s.shape[1] == 1: - # this is row quantized weight, just simple multiplication suffices + # Per-tensor scale: single value for entire weight matrix + if s.numel() == 1: + return x.to(dtype) * s.view(1, 1).to(dtype) + # Row quantized weight: scale shape is (m, 1) or (n, 1) + elif s.ndim == 2 and s.shape[1] == 1: if x.shape[0] == s.shape[0]: y = x.to(dtype) * s.to(dtype) elif x.shape[1] == s.shape[0]: @@ -106,8 +109,8 @@ def weight_dequant(x: torch.Tensor, s: torch.Tensor, dtype = torch.bfloat16): else: raise ValueError(f"Incompatible shapes {x.shape = }, {s.shape = }") return y + # Block quantized weight: scale shape is (ceil(m/block_m), ceil(n/block_n)) else: - # this is block quantized weight return weight_dequant_block(x, s, dtype = dtype) @@ -238,44 +241,29 @@ def w8a8_block_fp8_matmul_triton( block_size: list[int], output_dtype: torch.dtype = torch.float32, ) -> torch.Tensor: - """This function performs matrix multiplication with block-wise - quantization. - It takes two input tensors `A` and `B` with scales `As` and `Bs`. - The output is returned in the specified `output_dtype`. - Args: - A: The input tensor, e.g., activation. - B: The input tensor, e.g., weight. - As: The per-token-group quantization scale for `A`. - Bs: The per-block quantization scale for `B`. - block_size: The block size for per-block quantization. It should - be 2-dim, e.g., [128, 128]. - output_dytpe: The dtype of the returned tensor. - Returns: - torch.Tensor: The result of matmul. - """ - assert len(block_size) == 2 - block_n, block_k = block_size[0], block_size[1] + """Block-wise FP8 matmul.""" + if block_size is None: + block_n, block_k = 128, 128 + else: + assert len(block_size) == 2 + block_n, block_k = block_size[0], block_size[1] + N, K = B.shape assert A.shape[-1] == B.shape[-1] assert A.shape[:-1] == As.shape[:-1] and A.is_contiguous() assert triton.cdiv(A.shape[-1], block_k) == As.shape[-1] - M = A.numel() // A.shape[-1] - assert B.ndim == 2 and B.is_contiguous() and Bs.ndim == 2 - N, K = B.shape assert triton.cdiv(N, block_n) == Bs.shape[0] assert triton.cdiv(K, block_k) == Bs.shape[1] + M = A.numel() // A.shape[-1] C_shape = A.shape[:-1] + (N,) C = A.new_empty(C_shape, dtype = output_dtype) BLOCK_SIZE_M = 128 if M < BLOCK_SIZE_M: - BLOCK_SIZE_M = triton.next_power_of_2(M) - BLOCK_SIZE_M = max(BLOCK_SIZE_M, 16) - BLOCK_SIZE_K = block_k - assert block_k % BLOCK_SIZE_K == 0 - BLOCK_SIZE_N = block_n + BLOCK_SIZE_M = max(triton.next_power_of_2(M), 16) + BLOCK_SIZE_K, BLOCK_SIZE_N = block_k, block_n def grid(META): return ( @@ -342,29 +330,41 @@ fp8_block_matmul = ( class FP8BlockQuantLinear(torch.autograd.Function): @staticmethod def forward(ctx, X, weight, weight_scale): - # block_size = getattr(weight, 'block_size', [128,128]) m, n = weight.shape - p, q = weight_scale.shape - block_size = getattr(weight, "block_size", None) or getattr( - weight_scale, "block_size", [128, 128] - ) - assert block_size is not None, "block_size is not set" - if triton.cdiv(m, block_size[0]) != p or triton.cdiv(n, block_size[1]) != q: - if ( - triton.cdiv(m, block_size[0]) == q - and triton.cdiv(n, block_size[1]) == p - ): - # weights are transposed during backward pass for training :) - # We transpose weight scale to counter that. Note that transposing weight would cause issues with matmul with input X - weight_scale = weight_scale.T - else: - raise ValueError( - f"Weight shape {weight.shape} and scales shape {weight_scale.shape} is not compatible with block size {block_size}" - ) + + # Save original scale for backward (before any transformation) + original_weight_scale = weight_scale + + # Handle per-tensor quantization: expand scalar to block scale shape + if weight_scale.numel() == 1: + block_size = [128, 128] + # Expand scalar to (ceil(m/128), ceil(n/128)) - same value for all blocks + num_blocks_m = triton.cdiv(m, block_size[0]) + num_blocks_n = triton.cdiv(n, block_size[1]) + weight_scale = weight_scale.expand(num_blocks_m, num_blocks_n).contiguous() + else: + # Block quantization path + p, q = weight_scale.shape + block_size = getattr(weight, "block_size", None) or getattr( + weight_scale, "block_size", [128, 128] + ) + assert block_size is not None, "block_size is not set" + if triton.cdiv(m, block_size[0]) != p or triton.cdiv(n, block_size[1]) != q: + if ( + triton.cdiv(m, block_size[0]) == q + and triton.cdiv(n, block_size[1]) == p + ): + weight_scale = weight_scale.T + original_weight_scale = weight_scale # Update for transposed case + else: + raise ValueError( + f"Weight shape {weight.shape} and scales shape {weight_scale.shape} is not compatible with block size {block_size}" + ) if not weight.is_contiguous(): weight = weight.contiguous() - # this is replica of https://github.com/huggingface/transformers/blob/01c9e1ba683b3e50d7c76bf92f2d470759fd5e81/src/transformers/integrations/finegrained_fp8.py#L331-L353 + + # Quantize input and run FP8 matmul qinput, scale = act_quant(X, block_size[1]) output = fp8_block_matmul( qinput, @@ -375,8 +375,7 @@ class FP8BlockQuantLinear(torch.autograd.Function): output_dtype = X.dtype, ) ctx.weight = weight - ctx.weight_scale = weight_scale - ctx.block_size = block_size + ctx.weight_scale = original_weight_scale # Save original for backward return output.to(X.dtype) @staticmethod @@ -592,11 +591,14 @@ except: @torch_compile def fp8_linear(X, weight, weight_scale, bias = None): - if weight_scale.ndim == 2 and weight_scale.shape[1] > 1: - # This is block quantized FP8 matmul + # Per-tensor quantization: single scalar scale for entire weight + # Block quantized FP8: 2D scale tensor with multiple columns + if weight_scale.numel() == 1 or ( + weight_scale.ndim == 2 and weight_scale.shape[1] > 1 + ): out = fp8_block_quant_linear(X, weight, weight_scale) + # Row/channel quantized FP8: 2D scale with shape (n, 1) else: - # Row quantized FP8 out = fbgemm_fp8_linear(X, weight, weight_scale, bias) return out From d55a768456c92dd456d89d7e6580e158103b05cf Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Mon, 16 Feb 2026 21:19:45 -0800 Subject: [PATCH 23/63] [pre-commit.ci] pre-commit autoupdate (#4072) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit updates: - [github.com/astral-sh/ruff-pre-commit: v0.15.0 → v0.15.1](https://github.com/astral-sh/ruff-pre-commit/compare/v0.15.0...v0.15.1) Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- .pre-commit-config.yaml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index 47c2dd8010..92d8e80201 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -1,6 +1,6 @@ repos: - repo: https://github.com/astral-sh/ruff-pre-commit - rev: v0.15.0 + rev: v0.15.1 hooks: - id: ruff args: From 753312fe7b77f0e60040d8f5265ecf007a4db407 Mon Sep 17 00:00:00 2001 From: Michael Han <107991372+shimmyshimmer@users.noreply.github.com> Date: Tue, 17 Feb 2026 07:23:31 -0800 Subject: [PATCH 24/63] Update README Install.md Updating to include new installation links --- README.md | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/README.md b/README.md index 88f8c77ffa..4a9d4d01c8 100644 --- a/README.md +++ b/README.md @@ -49,8 +49,9 @@ For Windows, `pip install unsloth` works only if you have Pytorch installed. Rea ### Docker Use our official [Unsloth Docker image](https://hub.docker.com/r/unsloth/unsloth) ```unsloth/unsloth``` container. Read our [Docker Guide](https://unsloth.ai/docs/get-started/install/docker). -### Blackwell & DGX Spark -For RTX 50x, B200, 6000 GPUs: `pip install unsloth`. Read our [Blackwell Guide](https://unsloth.ai/docs/blog/fine-tuning-llms-with-blackwell-rtx-50-series-and-unsloth) and [DGX Spark Guide](https://unsloth.ai/docs/blog/fine-tuning-llms-with-nvidia-dgx-spark-and-unsloth) for more details. +### AMD, Intel, Blackwell & DGX Spark +For RTX 50x, B200, 6000 GPUs: `pip install unsloth`. Read our guides for: [Blackwell](https://unsloth.ai/docs/blog/fine-tuning-llms-with-blackwell-rtx-50-series-and-unsloth) and [DGX Spark](https://unsloth.ai/docs/blog/fine-tuning-llms-with-nvidia-dgx-spark-and-unsloth).
+To install Unsloth on **AMD** and **Intel** GPUs, follow our [AMD Guide](https://unsloth.ai/docs/get-started/install/amd) and [Intel Guide](https://unsloth.ai/docs/get-started/install/intel). ## 🦥 Unsloth News - Train **MoE LLMs 12x faster** with 35% less VRAM - DeepSeek, GLM, Qwen and gpt-oss. [Blog](https://unsloth.ai/docs/new/faster-moe) @@ -59,7 +60,6 @@ For RTX 50x, B200, 6000 GPUs: `pip install unsloth`. Read our [Blackwell Guide]( - New RoPE & MLP **Triton Kernels** & **Padding Free + Packing**: 3x faster training & 30% less VRAM. [Blog](https://unsloth.ai/docs/new/3x-faster-training-packing) - **500K Context**: Training a 20B model with >500K context is now possible on an 80GB GPU. [Blog](https://unsloth.ai/docs/blog/500k-context-length-fine-tuning) - **FP8 Reinforcement Learning**: You can now do FP8 GRPO on consumer GPUs. [Blog](https://unsloth.ai/docs/get-started/reinforcement-learning-rl-guide/fp8-reinforcement-learning) • [Notebook](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Qwen3_8B_FP8_GRPO.ipynb) -- **DeepSeek-OCR**: Fine-tune to improve language understanding by 89%. [Guide](https://unsloth.ai/docs/models/tutorials/deepseek-ocr-how-to-run-and-fine-tune) • [Notebook](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Deepseek_OCR_(3B).ipynb) - **Docker**: Use Unsloth with no setup & environment issues with our new image. [Guide](https://unsloth.ai/docs/blog/how-to-fine-tune-llms-with-unsloth-and-docker) • [Docker image](https://hub.docker.com/r/unsloth/unsloth) - **Vision RL**: You can now train VLMs with GRPO or GSPO in Unsloth! [Read guide](https://unsloth.ai/docs/get-started/reinforcement-learning-rl-guide/vision-reinforcement-learning-vlm-rl) - **gpt-oss** by OpenAI: Read our [RL blog](https://unsloth.ai/docs/models/gpt-oss-how-to-run-and-fine-tune/gpt-oss-reinforcement-learning), [Flex Attention](https://unsloth.ai/docs/models/gpt-oss-how-to-run-and-fine-tune/long-context-gpt-oss-training) blog and [gpt-oss Guide](https://unsloth.ai/docs/models/gpt-oss-how-to-run-and-fine-tune). 20B works on 14GB VRAM. 120B on 65GB. @@ -103,9 +103,9 @@ For RTX 50x, B200, 6000 GPUs: `pip install unsloth`. Read our [Blackwell Guide]( * Supports **all models** including [TTS](https://unsloth.ai/docs/basics/text-to-speech-tts-fine-tuning), multimodal, [embedding](https://unsloth.ai/docs/new/embedding-finetuning) and more! Any model that works in transformers, works in Unsloth. * The most efficient library for [Reinforcement Learning (RL)](https://unsloth.ai/docs/get-started/reinforcement-learning-rl-guide), using 80% less VRAM. Supports GRPO, GSPO, DrGRPO, DAPO etc. * **0% loss in accuracy** - no approximation methods - all exact. -* Export and [deploy your model](https://unsloth.ai/docs/basics/inference-and-deployment) to GGUF, llama.cpp, vLLM, SGLang and Hugging Face. -* Supports NVIDIA (since 2018), [AMD](https://unsloth.ai/docs/get-started/install/amd) and Intel GPUs. Minimum CUDA Capability 7.0 (V100, T4, Titan V, RTX 20, 30, 40x, A100, H100, L40 etc) -* Works on **Linux**, WSL and **Windows** +* Export and [deploy your model](https://unsloth.ai/docs/basics/inference-and-deployment) to [GGUF](https://unsloth.ai/docs/basics/inference-and-deployment/saving-to-gguf) llama.cpp, [vLLM](https://unsloth.ai/docs/basics/inference-and-deployment/vllm-guide), [SGLang](https://unsloth.ai/docs/basics/inference-and-deployment/sglang-guide) and Hugging Face. +* Supports NVIDIA (since 2018), [AMD](https://unsloth.ai/docs/get-started/install/amd) and [Intel](https://unsloth.ai/docs/get-started/install/intel) GPUs. Minimum CUDA Capability 7.0 (V100, T4, Titan V, RTX 20, 30, 40x, A100, H100, L40 etc) +* Works on **Linux**, WSL and **[Windows](https://unsloth.ai/docs/get-started/install/windows-installation)** * All kernels written in OpenAI's Triton language. Manual backprop engine. * If you trained a model with 🦥Unsloth, you can use this cool sticker!   From a5e0867f89e04c78dd99be20e2963c09279b777d Mon Sep 17 00:00:00 2001 From: Kaitao Yang <21039614+ykaitao@users.noreply.github.com> Date: Wed, 18 Feb 2026 17:13:33 -0800 Subject: [PATCH 25/63] reduce code duplicaton by inheritting from LlamaRotaryEmbedding (#3878) * simplify_code_using_apply_time_scaling * modify LlamaRotaryEmbedding for better inheritance * reduce_code_duplication_LlamaExtendedRotaryEmbedding --- unsloth/models/llama.py | 157 +++++++++------------------------------- 1 file changed, 33 insertions(+), 124 deletions(-) diff --git a/unsloth/models/llama.py b/unsloth/models/llama.py index dcaa770baa..a3fc9ab49a 100644 --- a/unsloth/models/llama.py +++ b/unsloth/models/llama.py @@ -1640,6 +1640,17 @@ class LlamaRotaryEmbedding(torch.nn.Module): self.multi_gpu_cos_cached = [None] * DEVICE_COUNT self.multi_gpu_sin_cached = [None] * DEVICE_COUNT + # Normal Llama-3 RoPE + inv_freq = 1.0 / ( + self.base + ** ( + torch.arange(0, self.dim, 2, dtype = torch.int64, device = "cpu").float() + / self.dim + ) + ) + inv_freq = self._apply_inv_freq_scaling(inv_freq) + self.register_buffer("inv_freq", inv_freq, persistent = False) + # Build here to make `torch.jit.trace` work. for device_idx in range(DEVICE_COUNT): self._set_cos_sin_cache( @@ -1656,22 +1667,24 @@ class LlamaRotaryEmbedding(torch.nn.Module): 1, device = get_current_device(), dtype = torch.get_default_dtype() ) + def _apply_inv_freq_scaling(self, inv_freq): + """Override to apply custom inv_freq scaling (e.g., extended RoPE).""" + return inv_freq + + def _apply_time_scaling(self, t): + """Override to apply custom time scaling (e.g., linear scaling).""" + return t + def _set_cos_sin_cache(self, seq_len, device, dtype): # Note: on the original Llama codebase, these tensors are created on the target device (and not on CPU) and # in FP32. They are applied (multiplied) in FP32 as well. self.current_rope_size = seq_len - inv_freq = 1.0 / ( - self.base - ** ( - torch.arange(0, self.dim, 2, dtype = torch.int64, device = "cpu").float() - / self.dim - ) - ) t = torch.arange( - self.current_rope_size, device = "cpu", dtype = torch.int64 + self.current_rope_size, device = self.inv_freq.device, dtype = torch.int64 ).float() + t = self._apply_time_scaling(t) - freqs = torch.outer(t, inv_freq) + freqs = torch.outer(t, self.inv_freq) # Different from paper, but it uses a different permutation in order to obtain the same calculation emb = torch.cat((freqs, freqs), dim = -1) cos = emb.cos().to(dtype = dtype, device = device, non_blocking = True) @@ -1733,33 +1746,14 @@ class LlamaLinearScalingRotaryEmbedding(LlamaRotaryEmbedding): config = config, ) - def _set_cos_sin_cache(self, seq_len, device, dtype): - self.current_rope_size = seq_len - inv_freq = 1.0 / ( - self.base - ** ( - torch.arange(0, self.dim, 2, dtype = torch.int64, device = "cpu").float() - / self.dim - ) - ) - t = torch.arange( - self.current_rope_size, device = "cpu", dtype = torch.int64 - ).float() - t = t / self.scaling_factor - - freqs = torch.outer(t, inv_freq) - # Different from paper, but it uses a different permutation in order to obtain the same calculation - emb = torch.cat((freqs, freqs), dim = -1) - cos = emb.cos().to(dtype = dtype, device = device, non_blocking = True) - sin = emb.sin().to(dtype = dtype, device = device, non_blocking = True) - self.multi_gpu_cos_cached[device.index] = cos - self.multi_gpu_sin_cached[device.index] = sin - return cos, sin + def _apply_time_scaling(self, t): + """Apply linear scaling to time indices.""" + return t / self.scaling_factor # See https://github.com/vllm-project/vllm/blob/main/vllm/model_executor/layers/rotary_embedding.py#L736 # For Llama 3.1 -class LlamaExtendedRotaryEmbedding(torch.nn.Module): +class LlamaExtendedRotaryEmbedding(LlamaRotaryEmbedding): def __init__( self, dim = None, @@ -1768,101 +1762,16 @@ class LlamaExtendedRotaryEmbedding(torch.nn.Module): device = None, config = None, # [TODO] Hack to pass in config - need to remove later ): - super().__init__() - if config is not None: - # [TODO] Hack to pass in config - need to remove later - base = _get_rope_theta(config, default = base) - partial_rotary_factor = ( - config.partial_rotary_factor - if hasattr(config, "partial_rotary_factor") - else 1.0 - ) - dim = int((config.hidden_size // config.num_attention_heads)) - device = DEVICE_TYPE_TORCH - max_position_embeddings = config.max_position_embeddings - - self.dim = dim - self.max_position_embeddings = max_position_embeddings - self.base = base - # Dynamic RoPE we first set it to a max of 4 * 8192 tokens then we iteratively grow this - self.current_rope_size = min(4 * 8192, self.max_position_embeddings) - self.multi_gpu_cos_cached = [None] * DEVICE_COUNT - self.multi_gpu_sin_cached = [None] * DEVICE_COUNT - - # Normal Llama-3 RoPE - inv_freq = 1.0 / ( - self.base - ** ( - torch.arange(0, self.dim, 2, dtype = torch.int64, device = "cpu").float() - / self.dim - ) + super().__init__( + dim = dim, + max_position_embeddings = max_position_embeddings, + base = base, + device = device, + config = config, ) - inv_freq = self.apply_scaling(inv_freq) - self.register_buffer("inv_freq", inv_freq, persistent = False) - - # Build here to make `torch.jit.trace` work. - for device_idx in range(DEVICE_COUNT): - self._set_cos_sin_cache( - seq_len = self.current_rope_size, - device = torch.device(device_idx), - dtype = torch.get_default_dtype(), - ) - - # dummy so that patch_utils doesn't fail for now - self.cos_cached = torch.empty( - 1, device = get_current_device(), dtype = torch.get_default_dtype() - ) - self.sin_cached = torch.empty( - 1, device = get_current_device(), dtype = torch.get_default_dtype() - ) - - def _set_cos_sin_cache(self, seq_len, device, dtype): - # Note: on the original Llama codebase, these tensors are created on the target device (and not on CPU) and - # in FP32. They are applied (multiplied) in FP32 as well. - self.current_rope_size = seq_len - - t = torch.arange( - self.current_rope_size, device = self.inv_freq.device, dtype = torch.int64 - ).float() - - freqs = torch.outer(t, self.inv_freq) - # Different from paper, but it uses a different permutation in order to obtain the same calculation - emb = torch.cat((freqs, freqs), dim = -1) - cos = emb.cos().to(dtype = dtype, device = device, non_blocking = True) - sin = emb.sin().to(dtype = dtype, device = device, non_blocking = True) - self.multi_gpu_cos_cached[device.index] = cos - self.multi_gpu_sin_cached[device.index] = sin - return cos, sin - - def forward(self, x, position_ids = None, seq_len = None): - # x: [bs, num_attention_heads, seq_len, head_size] - if seq_len is not None and seq_len > self.current_rope_size: - self._set_cos_sin_cache(seq_len = seq_len, device = x.device, dtype = x.dtype) - device_index = x.device.index - return ( - self.multi_gpu_cos_cached[device_index][:seq_len], - self.multi_gpu_sin_cached[device_index][:seq_len], - ) - - def get_cached(self, seq_len = None, device_index = None): - if device_index is None: - device_index = get_current_device() - return self.multi_gpu_cos_cached[device_index], self.multi_gpu_sin_cached[ - device_index - ] - - def extend_rope_embedding(self, x, seq_len): - if seq_len <= self.current_rope_size: - return - # Iteratively grow by increments of 8192 - self.current_rope_size = ((seq_len // 8192) + ((seq_len % 8192) != 0)) * 8192 - for device_idx in range(DEVICE_COUNT): - self._set_cos_sin_cache( - self.current_rope_size, device = torch.device(device_idx), dtype = x.dtype - ) # From https://github.com/meta-llama/llama-models/blob/main/models/llama3_1/api/model.py#L41 - def apply_scaling(self, freqs: torch.Tensor): + def _apply_inv_freq_scaling(self, freqs: torch.Tensor): # Values obtained from grid search scale_factor = 8 low_freq_factor = 1 From 431fa22891c52334b61b1d84cd1bed294728f6c1 Mon Sep 17 00:00:00 2001 From: Daniel van Strien Date: Thu, 19 Feb 2026 11:53:35 +0000 Subject: [PATCH 26/63] Add optional datasets metadata support to save/push functions (#4076) * Add `datasets` metadata support to model cards Add an optional `datasets` parameter to all save/push functions so users can specify which datasets were used for training. The metadata is set via `ModelCard.data.datasets` for standard paths and via `metadata_update` for GGUF and generic save paths. Co-Authored-By: Claude Opus 4.6 * Fix datasets metadata for existing repos, add token, improve errors - Add metadata_update fallback in create_huggingface_repo and upload_to_huggingface so datasets metadata is set even when the repo already exists (previously only worked on first creation). - Pass token=token to all metadata_update calls so they work without a global HF login. - Replace silent except:pass with logger.warning_once for metadata failures so users know if something went wrong. Co-Authored-By: Claude Opus 4.6 * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Fix generic datasets metadata repo resolution for PR #4076 * Fix create_huggingface_repo username resolution for PR #4076 --------- Co-authored-by: Claude Opus 4.6 Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> Co-authored-by: Daniel Hanchen --- unsloth/save.py | 71 ++++++++++++++++++++++++++++++++++++++++++++++--- 1 file changed, 68 insertions(+), 3 deletions(-) diff --git a/unsloth/save.py b/unsloth/save.py index fc3b7b8771..a07575c35d 100644 --- a/unsloth/save.py +++ b/unsloth/save.py @@ -257,6 +257,7 @@ def unsloth_save_model( # Our functions temporary_location: str = "_unsloth_temporary_saved_buffers", maximum_memory_usage: float = 0.9, + datasets: Optional[List[str]] = None, ): if token is None: token = get_token() @@ -289,6 +290,7 @@ def unsloth_save_model( "save_method", "temporary_location", "maximum_memory_usage", + "datasets", ): del save_pretrained_settings[deletion] @@ -366,6 +368,7 @@ def unsloth_save_model( file_location = None, old_username = None, private = private, + datasets = datasets, ) getattr(model, "original_push_to_hub", model.push_to_hub)( @@ -475,6 +478,7 @@ def unsloth_save_model( file_location = None, old_username = None, private = private, + datasets = datasets, ) if tokenizer is not None: @@ -737,6 +741,7 @@ def unsloth_save_model( file_location = None, old_username = username, private = private, + datasets = datasets, ) # First check if we're pushing to an organization! @@ -1362,6 +1367,7 @@ def unsloth_save_pretrained_merged( tags: List[str] = None, temporary_location: str = "_unsloth_temporary_saved_buffers", maximum_memory_usage: float = 0.75, + datasets: Optional[List[str]] = None, ): """ Same as .save_pretrained(...) except 4bit weights are auto @@ -1403,6 +1409,7 @@ def unsloth_push_to_hub_merged( tags: Optional[List[str]] = None, temporary_location: str = "_unsloth_temporary_saved_buffers", maximum_memory_usage: float = 0.75, + datasets: Optional[List[str]] = None, ): """ Same as .push_to_hub(...) except 4bit weights are auto @@ -1480,10 +1487,11 @@ def create_huggingface_repo( save_directory, token = None, private = False, + datasets = None, ): if token is None: token = get_token() - save_directory, username = _determine_username(save_directory, "", token) + save_directory, username = _determine_username(save_directory, None, token) from huggingface_hub import create_repo @@ -1507,9 +1515,22 @@ def create_huggingface_repo( extra = "unsloth", ) card = ModelCard(content) + if datasets: + card.data.datasets = datasets card.push_to_hub(save_directory, token = token) except: - pass + # Repo already exists — update datasets metadata separately + if datasets: + try: + from huggingface_hub import metadata_update + + metadata_update( + save_directory, {"datasets": datasets}, overwrite = True, token = token + ) + except Exception as e: + logger.warning_once( + f"Unsloth: Could not update datasets metadata for {save_directory}: {e}" + ) hf_api = HfApi(token = token) return save_directory, hf_api @@ -1524,6 +1545,7 @@ def upload_to_huggingface( old_username = None, private = None, create_config = True, + datasets = None, ): save_directory, username = _determine_username(save_directory, old_username, token) @@ -1549,9 +1571,22 @@ def upload_to_huggingface( extra = extra, ) card = ModelCard(content) + if datasets: + card.data.datasets = datasets card.push_to_hub(save_directory, token = token) except: - pass + # Repo already exists — update datasets metadata separately + if datasets: + try: + from huggingface_hub import metadata_update + + metadata_update( + save_directory, {"datasets": datasets}, overwrite = True, token = token + ) + except Exception as e: + logger.warning_once( + f"Unsloth: Could not update datasets metadata for {save_directory}: {e}" + ) if file_location is not None: # Now upload file @@ -2083,6 +2118,7 @@ def unsloth_push_to_hub_gguf( tags: Optional[List[str]] = None, temporary_location: str = "_unsloth_temporary_saved_buffers", maximum_memory_usage: float = 0.85, + datasets: Optional[List[str]] = None, ): """ Same as .push_to_hub(...) except 4bit weights are auto @@ -2338,6 +2374,18 @@ This model was finetuned and converted to GGUF format using [Unsloth](https://gi except: pass + if datasets: + try: + from huggingface_hub import metadata_update + + metadata_update( + full_repo_id, {"datasets": datasets}, overwrite = True, token = token + ) + except Exception as e: + logger.warning_once( + f"Unsloth: Could not update datasets metadata for {full_repo_id}: {e}" + ) + except Exception as e: raise RuntimeError(f"Failed to upload to Hugging Face Hub: {e}") @@ -2645,6 +2693,7 @@ def unsloth_generic_save( # Our functions temporary_location: str = "_unsloth_temporary_saved_buffers", maximum_memory_usage: float = 0.9, + datasets: Optional[List[str]] = None, ): if token is None and push_to_hub: token = get_token() @@ -2672,6 +2721,20 @@ def unsloth_generic_save( low_disk_space_usage = True, use_temp_file = False, ) + + if push_to_hub and datasets: + try: + from huggingface_hub import metadata_update + + save_dir, _ = _determine_username(save_directory, None, token) + metadata_update( + save_dir, {"datasets": datasets}, overwrite = True, token = token + ) + except Exception as e: + logger.warning_once( + f"Unsloth: Could not update datasets metadata for {save_directory}: {e}" + ) + return @@ -2692,6 +2755,7 @@ def unsloth_generic_save_pretrained_merged( tags: List[str] = None, temporary_location: str = "_unsloth_temporary_saved_buffers", maximum_memory_usage: float = 0.75, + datasets: Optional[List[str]] = None, ): """ Same as .push_to_hub(...) except 4bit weights are auto @@ -2733,6 +2797,7 @@ def unsloth_generic_push_to_hub_merged( tags: Optional[List[str]] = None, temporary_location: str = "_unsloth_temporary_saved_buffers", maximum_memory_usage: float = 0.75, + datasets: Optional[List[str]] = None, ): """ Same as .push_to_hub(...) except 4bit weights are auto From 5e2e75243fade29aacc215996ad17f3bc12a21d5 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Thu, 19 Feb 2026 04:40:14 -0800 Subject: [PATCH 27/63] Patch trunc_normal_ for low-precision stability (#4027) * Fix low-precision trunc_normal initialization instability * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Document TorchTitan trunc_normal low-precision failure mode * Fix trunc_normal generator positional compatibility * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Fix trunc_normal generator TypeError fallback --------- Co-authored-by: Daniel Hanchen Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- tests/utils/test_trunc_normal_patch.py | 114 +++++++++++++++++++++++++ unsloth/__init__.py | 3 + unsloth/import_fixes.py | 82 ++++++++++++++++++ 3 files changed, 199 insertions(+) create mode 100644 tests/utils/test_trunc_normal_patch.py diff --git a/tests/utils/test_trunc_normal_patch.py b/tests/utils/test_trunc_normal_patch.py new file mode 100644 index 0000000000..b84a0772d8 --- /dev/null +++ b/tests/utils/test_trunc_normal_patch.py @@ -0,0 +1,114 @@ +# Copyright 2023-present Daniel Han-Chen, Michael Han-Chen & the Unsloth team. All rights reserved. +# +# This program is free software: you can redistribute it and/or modify +# it under the terms of the GNU Lesser General Public License as published by +# the Free Software Foundation, either version 3 of the License, or +# (at your option) any later version. +# +# This program is distributed in the hope that it will be useful, +# but WITHOUT ANY WARRANTY; without even the implied warranty of +# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the +# GNU General Public License for more details. +# +# You should have received a copy of the GNU Lesser General Public License +# along with this program. If not, see . + +"""Tests for trunc_normal low-precision patch compatibility.""" + +import importlib.util +import inspect +from pathlib import Path + +import pytest +import torch + + +_MISSING = object() + + +def _load_import_fixes_module(): + repo_root = Path(__file__).resolve().parents[2] + import_fixes_path = repo_root / "unsloth" / "import_fixes.py" + spec = importlib.util.spec_from_file_location( + "unsloth_import_fixes_local", import_fixes_path + ) + assert spec is not None and spec.loader is not None + module = importlib.util.module_from_spec(spec) + spec.loader.exec_module(module) + return module + + +def _getattr_or_missing(obj, name): + return getattr(obj, name) if hasattr(obj, name) else _MISSING + + +def _restore_attr(obj, name, value): + if value is _MISSING: + if hasattr(obj, name): + delattr(obj, name) + return + setattr(obj, name, value) + + +def test_trunc_normal_patch_accepts_positional_generator(): + import_fixes = _load_import_fixes_module() + patch_fn = import_fixes.patch_trunc_normal_precision_issue + + init_mod = torch.nn.init + old_fn = init_mod.trunc_normal_ + old_patched = _getattr_or_missing(init_mod, "_unsloth_trunc_normal_patched") + old_original = _getattr_or_missing(init_mod, "_unsloth_trunc_normal_original") + try: + # Normalize to an unpatched baseline before applying the patch. + if old_original is not _MISSING: + init_mod.trunc_normal_ = old_original + if hasattr(init_mod, "_unsloth_trunc_normal_patched"): + delattr(init_mod, "_unsloth_trunc_normal_patched") + if hasattr(init_mod, "_unsloth_trunc_normal_original"): + delattr(init_mod, "_unsloth_trunc_normal_original") + + patch_fn() + sig = inspect.signature(init_mod.trunc_normal_) + assert "generator" in sig.parameters + assert sig.parameters["generator"].kind is not inspect.Parameter.KEYWORD_ONLY + + tensor = torch.empty(1024, dtype = torch.float32) + gen = torch.Generator() + gen.manual_seed(3407) + + init_mod.trunc_normal_(tensor, 0.0, 1.0, -2.0, 2.0, gen) + init_mod.trunc_normal_(tensor, mean = 0.0, std = 1.0, a = -2.0, b = 2.0, generator = gen) + finally: + init_mod.trunc_normal_ = old_fn + _restore_attr(init_mod, "_unsloth_trunc_normal_patched", old_patched) + _restore_attr(init_mod, "_unsloth_trunc_normal_original", old_original) + + +def test_trunc_normal_patch_rejects_invalid_generator(): + import_fixes = _load_import_fixes_module() + patch_fn = import_fixes.patch_trunc_normal_precision_issue + + init_mod = torch.nn.init + old_fn = init_mod.trunc_normal_ + old_patched = _getattr_or_missing(init_mod, "_unsloth_trunc_normal_patched") + old_original = _getattr_or_missing(init_mod, "_unsloth_trunc_normal_original") + try: + if old_original is not _MISSING: + init_mod.trunc_normal_ = old_original + if hasattr(init_mod, "_unsloth_trunc_normal_patched"): + delattr(init_mod, "_unsloth_trunc_normal_patched") + if hasattr(init_mod, "_unsloth_trunc_normal_original"): + delattr(init_mod, "_unsloth_trunc_normal_original") + + patch_fn() + sig = inspect.signature(init_mod.trunc_normal_) + if "generator" not in sig.parameters: + pytest.skip("torch.nn.init.trunc_normal_ lacks a generator parameter") + + tensor = torch.empty(16, dtype = torch.float32) + with pytest.raises(TypeError): + init_mod.trunc_normal_(tensor, generator = 123) + finally: + init_mod.trunc_normal_ = old_fn + _restore_attr(init_mod, "_unsloth_trunc_normal_patched", old_patched) + _restore_attr(init_mod, "_unsloth_trunc_normal_original", old_original) diff --git a/unsloth/__init__.py b/unsloth/__init__.py index 7932fe966d..466fbf0628 100644 --- a/unsloth/__init__.py +++ b/unsloth/__init__.py @@ -140,6 +140,7 @@ from .import_fixes import ( fix_vllm_pdl_blackwell, fix_triton_compiled_kernel_missing_attrs, fix_rocm_triton_key_error, + patch_trunc_normal_precision_issue, ignore_logger_messages, patch_ipykernel_hf_xet, patch_trackio, @@ -161,6 +162,7 @@ fix_vllm_guided_decoding_params() fix_vllm_pdl_blackwell() fix_triton_compiled_kernel_missing_attrs() fix_rocm_triton_key_error() +patch_trunc_normal_precision_issue() ignore_logger_messages() patch_ipykernel_hf_xet() patch_trackio() @@ -180,6 +182,7 @@ del fix_vllm_guided_decoding_params del fix_vllm_pdl_blackwell del fix_triton_compiled_kernel_missing_attrs del fix_rocm_triton_key_error +del patch_trunc_normal_precision_issue del ignore_logger_messages del patch_ipykernel_hf_xet del patch_trackio diff --git a/unsloth/import_fixes.py b/unsloth/import_fixes.py index 67e4a26ffb..b3200baa05 100644 --- a/unsloth/import_fixes.py +++ b/unsloth/import_fixes.py @@ -927,6 +927,88 @@ def fix_rocm_triton_key_error(): ) +def patch_trunc_normal_precision_issue(): + """ + Patch torch.nn.init.trunc_normal_ for low precision tensors to run init in fp32. + + torch.nn.init.trunc_normal_ can saturate at truncation bounds in fp16/bf16 on + some versions/backends. This was observed in TorchTitan investigations where + low-precision truncation produced boundary-heavy initialization behavior: + https://github.com/pytorch/torchtitan/pull/2342 + + To avoid that failure mode, initialize into a temporary fp32 tensor, then copy + back to the original dtype. + """ + try: + import torch + except (ImportError, ModuleNotFoundError): + return + + if getattr(torch.nn.init, "_unsloth_trunc_normal_patched", False): + return + + original_trunc_normal = torch.nn.init.trunc_normal_ + if getattr(original_trunc_normal, "__unsloth_trunc_normal_patched__", False): + torch.nn.init._unsloth_trunc_normal_patched = True + return + + low_precision_dtypes = {torch.float16, torch.bfloat16} + + def _call_original(target, mean, std, a, b, generator): + if generator is None: + return original_trunc_normal(target, mean = mean, std = std, a = a, b = b) + try: + return original_trunc_normal( + target, mean = mean, std = std, a = a, b = b, generator = generator + ) + except TypeError as exc: + # Older torch versions may not accept a generator keyword argument. + msg = str(exc).lower() + if "unexpected keyword argument" in msg and "generator" in msg: + return original_trunc_normal(target, mean = mean, std = std, a = a, b = b) + raise + + try: + from torch.distributed._tensor import DTensor + except Exception: + DTensor = None + + @torch.no_grad() + def _patched_trunc_normal_( + tensor, + mean: float = 0.0, + std: float = 1.0, + a: float = -2.0, + b: float = 2.0, + generator = None, + ): + if DTensor is not None and isinstance(tensor, DTensor): + local_tensor = getattr(tensor, "_local_tensor", None) + if local_tensor is None: + return _call_original(tensor, mean, std, a, b, generator) + if local_tensor.dtype in low_precision_dtypes: + local_fp32 = local_tensor.float() + _call_original(local_fp32, mean, std, a, b, generator) + local_tensor.copy_(local_fp32.to(dtype = local_tensor.dtype)) + return tensor + return _call_original(tensor, mean, std, a, b, generator) + + if tensor.dtype in low_precision_dtypes: + tensor_fp32 = tensor.float() + _call_original(tensor_fp32, mean, std, a, b, generator) + tensor.copy_(tensor_fp32.to(dtype = tensor.dtype)) + return tensor + + return _call_original(tensor, mean, std, a, b, generator) + + _patched_trunc_normal_.__unsloth_trunc_normal_patched__ = True + _patched_trunc_normal_._unsloth_original = original_trunc_normal + torch.nn.init._unsloth_trunc_normal_original = original_trunc_normal + torch.nn.init.trunc_normal_ = _patched_trunc_normal_ + torch.nn.init._unsloth_trunc_normal_patched = True + logger.info("Unsloth: Patched torch.nn.init.trunc_normal_ for fp16/bf16 stability.") + + def check_vllm_torch_sm100_compatibility(): """ Check for incompatible vLLM + torch < 2.9.0 + SM100 (Blackwell) combination. From 74e7af649a4d6e7e3902fa62c598ecc3a57f7e6c Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Mon, 23 Feb 2026 01:06:53 -0800 Subject: [PATCH 28/63] Fix VLM processor load degradation and vLLM CUDA version detection (#4091) * Fix VLM processor load degradation and vLLM CUDA version detection vision.py - Fix VLM processor load for issue #4085: - Before loading the processor, scan local config files and strip the _Unsloth_Patched_ prefix. AutoProcessor.from_pretrained silently degrades to a text-only tokenizer instead of raising an exception when it encounters the unrecognized class name, so the existing get_auto_processor fallback never triggers. Sanitizing the configs before loading fixes backwards compat for old corrupted saves. - After loading, detect when AutoProcessor returned a text-only tokenizer for a VLM model (has no image_processor attribute) and trigger the manual fallback constructor. import_fixes.py - Fix vLLM CUDA version mismatch detection: - _is_broken_vllm_error now also matches CUDA shared library errors (libcudart, libcublas, libnvrtc) with "cannot open shared object file". Previously it only matched errors containing "vllm._c" in the message text, which missed cases where the error message was about the missing CUDA library itself (e.g. vllm built for CUDA 12 on a CUDA 13 system). - New _get_vllm_cuda_mismatch_message function extracts the CUDA version from the error, compares to the system CUDA version via torch.version.cuda, and returns a targeted install command using the correct GitHub releases wheel URL. - disable_broken_vllm uses the targeted message when a CUDA mismatch is detected, falling back to the existing generic message otherwise. * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --------- Co-authored-by: Ubuntu Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- unsloth/import_fixes.py | 82 +++++++++++++++++++++++++++++++++++++--- unsloth/models/vision.py | 39 ++++++++++++++++++- 2 files changed, 113 insertions(+), 8 deletions(-) diff --git a/unsloth/import_fixes.py b/unsloth/import_fixes.py index b3200baa05..172ba36f63 100644 --- a/unsloth/import_fixes.py +++ b/unsloth/import_fixes.py @@ -1467,12 +1467,78 @@ def _is_broken_vllm_error(error) -> bool: ) ) or ("vllm" in message and "undefined symbol" in message): return True + # Also catch CUDA shared library mismatches during vllm import + # e.g. "libcudart.so.12: cannot open shared object file" + if ( + "libcudart" in message or "libcublas" in message or "libnvrtc" in message + ) and "cannot open shared object file" in message: + return True current = getattr(current, "__cause__", None) or getattr( current, "__context__", None ) return False +def _get_vllm_cuda_mismatch_message(error): + """If the error is a CUDA version mismatch, return a helpful install message.""" + import re as _re + + checked = set() + current = error + wanted_cuda = None + while current is not None and id(current) not in checked: + checked.add(id(current)) + message = str(current) + # Extract the CUDA version vllm was built for, e.g. "libcudart.so.12" + match = _re.search(r"libcudart\.so\.(\d+)", message) + if match: + wanted_cuda = match.group(1) + break + current = getattr(current, "__cause__", None) or getattr( + current, "__context__", None + ) + if wanted_cuda is None: + return None + + # Detect what CUDA version is actually available on the system + system_cuda_display = None # Human-readable, e.g. "13.0" + system_cuda_tag = None # For wheel URL, e.g. "130" + try: + import torch + + cuda_version = torch.version.cuda # e.g. "13.0" or "12.8" + if cuda_version: + system_cuda_display = cuda_version + system_cuda_tag = cuda_version.replace(".", "")[:3] # "130" or "128" + except Exception: + pass + + if system_cuda_tag is None or system_cuda_tag.startswith(wanted_cuda): + return None # Not a mismatch or can't determine + + try: + vllm_version = importlib_version("vllm").split("+")[0] + except Exception: + vllm_version = "VLLM_VERSION" + + cpu_arch = "x86_64" + try: + import platform + + cpu_arch = platform.machine() + except Exception: + pass + + return ( + f"Unsloth: vLLM was built for CUDA {wanted_cuda} but this system has " + f"CUDA {system_cuda_display}. Please reinstall vLLM with the correct CUDA version:\n" + f"\n" + f" uv pip install https://github.com/vllm-project/vllm/releases/download/" + f"v{vllm_version}/vllm-{vllm_version}+cu{system_cuda_tag}-cp38-abi3-" + f"manylinux_2_35_{cpu_arch}.whl" + ) + + class _CausalConv1dImportBlockerLoader(importlib.abc.Loader): __slots__ = ("module_name",) @@ -1621,12 +1687,16 @@ def disable_broken_vllm(error = None): VLLM_BROKEN = True _clear_vllm_modules() _install_vllm_blocker() - logger.warning( - "Unsloth: Detected broken vLLM binary extension; " - "disabling vLLM imports and continuing import.\n" - "Please reinstall via `uv pip install unsloth vllm torchvision torchaudio " - "--torch-backend=auto`." - ) + cuda_msg = _get_vllm_cuda_mismatch_message(failure) + if cuda_msg: + logger.warning(cuda_msg) + else: + logger.warning( + "Unsloth: Detected broken vLLM binary extension; " + "disabling vLLM imports and continuing import.\n" + "Please reinstall via `uv pip install unsloth vllm torchvision torchaudio " + "--torch-backend=auto`." + ) return True diff --git a/unsloth/models/vision.py b/unsloth/models/vision.py index 00d8a3b7f2..2d6c7ac93f 100644 --- a/unsloth/models/vision.py +++ b/unsloth/models/vision.py @@ -916,6 +916,32 @@ class FastBaseModel: # Counteract saved tokenizers tokenizer_name = model_name if tokenizer_name is None else tokenizer_name + + # Fix _Unsloth_Patched_ prefix in local config files from old saves (issue #4085) + if os.path.isdir(tokenizer_name): + import json as _json + + for _cfg_name in ( + "processor_config.json", + "preprocessor_config.json", + "tokenizer_config.json", + ): + _cfg_path = os.path.join(tokenizer_name, _cfg_name) + if os.path.exists(_cfg_path): + try: + with open(_cfg_path, "r", encoding = "utf-8") as _f: + _cfg = _json.load(_f) + if _cfg.get("processor_class", "").startswith( + "_Unsloth_Patched_" + ): + _cfg["processor_class"] = _cfg["processor_class"][ + len("_Unsloth_Patched_") : + ] + with open(_cfg_path, "w", encoding = "utf-8") as _f: + _json.dump(_cfg, _f, indent = 2, ensure_ascii = False) + except Exception: + pass + if (whisper_language and whisper_task) or auto_model.__name__.endswith( "ForConditionalGeneration" ): @@ -947,14 +973,23 @@ class FastBaseModel: ) # If processor loading failed (e.g., tokenizer class not found), + # or if AutoProcessor silently degraded to a text-only tokenizer + # instead of returning a full VLM processor (issue #4085), # try constructing the processor manually from separate components. - if tokenizer is None and is_vlm: - tokenizer = _construct_vlm_processor_fallback( + _processor_is_degraded = ( + is_vlm + and tokenizer is not None + and not hasattr(tokenizer, "image_processor") + ) + if (tokenizer is None or _processor_is_degraded) and is_vlm: + _fallback = _construct_vlm_processor_fallback( tokenizer_name, model_type_arch, token, trust_remote_code, ) + if _fallback is not None: + tokenizer = _fallback if tokenizer is None: import sys From a2fd00d727392e12e34e55bcafb48c3df2383fec Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Mon, 23 Feb 2026 01:27:10 -0800 Subject: [PATCH 29/63] Suppress FBGEMM CUTLASS stdout spam on Blackwell GPUs (#4092) * Suppress FBGEMM CUTLASS "Arch conditional MMA" stdout spam on Blackwell GPUs On Blackwell GPUs (B200/B100, SM100), FBGEMM's f8f8bf16_blockwise kernel is hardcoded to cutlass::arch::Sm90 with no SM100 code path. When test_has_fbgemm() probes this kernel, it fires 2304 "ERROR : Arch conditional MMA instruction used without targeting appropriate compute capability" lines before aborting and returning zeros. The existing HidePrintMessage filter on sys.stderr (line 109) does not catch these because CUDA device-side printf writes to stdout fd 1 at the C level, bypassing Python's sys.stdout/sys.stderr entirely. Fix: add suppress_cuda_printf() context manager in import_fixes.py that redirects fd 1 and fd 2 to /dev/null at the OS level, with torch.cuda.synchronize() and libc fflush before restoring. Wrap the test_has_fbgemm() call in fp8.py with this context manager. Tested on B200 with fbgemm-gpu-genai 1.4.0+cu130 and 1.5.0+cu130: - Before: 2304 warning lines on every import - After: 0 warning lines - UNSLOTH_HAS_FBGEMM correctly set to 0 (Triton fallback works) - Works with both UNSLOTH_ENABLE_LOGGING=0 and =1 * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Guard _libc init and fflush to prevent fd leak on failure --------- Co-authored-by: Ubuntu Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- unsloth/import_fixes.py | 47 +++++++++++++++++++++++++++++++++++++++++ unsloth/kernels/fp8.py | 9 +++++++- 2 files changed, 55 insertions(+), 1 deletion(-) diff --git a/unsloth/import_fixes.py b/unsloth/import_fixes.py index 172ba36f63..ee82715306 100644 --- a/unsloth/import_fixes.py +++ b/unsloth/import_fixes.py @@ -97,6 +97,53 @@ class HidePrintMessage: return getattr(self._original_stream, name) +import contextlib +import ctypes + +try: + _libc = ctypes.CDLL(None) +except Exception: + _libc = None + + +@contextlib.contextmanager +def suppress_cuda_printf(): + """Suppress CUDA device-side printf by redirecting stdout/stderr fds to /dev/null. + + CUDA device printf (eg CUTLASS "Arch conditional MMA" errors on Blackwell) + writes to stdout fd 1 at the C level, bypassing Python sys.stdout entirely. + The existing HidePrintMessage filter on sys.stderr cannot catch these since + they go to a different fd at a different layer. This context manager redirects + both fd 1 and fd 2 at the OS level, syncs CUDA, then restores them. + """ + sys.stdout.flush() + sys.stderr.flush() + saved_fds = {} + try: + for fd in (1, 2): + saved_fds[fd] = os.dup(fd) + devnull = os.open(os.devnull, os.O_WRONLY) + os.dup2(devnull, fd) + os.close(devnull) + yield + finally: + try: + import torch + + if torch.cuda.is_available(): + torch.cuda.synchronize() + except Exception: + pass + if _libc is not None: + try: + _libc.fflush(None) + except Exception: + pass + for fd, saved in saved_fds.items(): + os.dup2(saved, fd) + os.close(saved) + + if not UNSLOTH_ENABLE_LOGGING: import sys diff --git a/unsloth/kernels/fp8.py b/unsloth/kernels/fp8.py index 0b073d59d6..a57f4ffb64 100644 --- a/unsloth/kernels/fp8.py +++ b/unsloth/kernels/fp8.py @@ -579,7 +579,14 @@ try: if Version(fbgemm_gpu.__version__) >= Version("1.4.0"): # We must manually confirm if blockwise FBGEMM works! # This check is a must for consumer grade GPUs which fail - if test_has_fbgemm(): + # Suppress CUDA device printf during probe -- on Blackwell (SM100) GPUs, + # FBGEMM's CUTLASS blockwise kernel (hardcoded SM90) fires thousands of + # "Arch conditional MMA" lines to stdout fd 1 before aborting. + from unsloth.import_fixes import suppress_cuda_printf + + with suppress_cuda_printf(): + _has_fbgemm = test_has_fbgemm() + if _has_fbgemm: os.environ["UNSLOTH_HAS_FBGEMM"] = "1" logger.info(f"Using fbgemm_gpu block quantized FP8 matmul") fp8_block_quant_linear = fp8_fbgemm_block_linear From c20cd03f9527d41e4cf9323f362a4ba3f7907372 Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Mon, 23 Feb 2026 17:04:34 -0800 Subject: [PATCH 30/63] [pre-commit.ci] pre-commit autoupdate (#4096) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit updates: - [github.com/astral-sh/ruff-pre-commit: v0.15.1 → v0.15.2](https://github.com/astral-sh/ruff-pre-commit/compare/v0.15.1...v0.15.2) Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- .pre-commit-config.yaml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index 92d8e80201..245c636fe7 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -1,6 +1,6 @@ repos: - repo: https://github.com/astral-sh/ruff-pre-commit - rev: v0.15.1 + rev: v0.15.2 hooks: - id: ruff args: From 738176772eb6a2148c1daba0e630220df4ad7393 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 24 Feb 2026 05:56:07 -0800 Subject: [PATCH 31/63] Fix FP8 model loading: redirect to BF16 sibling for BNB/16-bit (#4095) * Fix FP8 model loading for BNB/16-bit: redirect to BF16 sibling Models like Ministral-3-3B-Instruct-2512 ship with FP8 weights and an FP8 quantization_config in their config.json. Loading these with BNB 4-bit/8-bit fails because BNB cannot quantize FP8 tensors. Loading with 16-bit also fails because the FP8 quantization config has activation_scheme=static which is unsupported by transformers' FineGrainedFP8Config. When an FP8 model is detected and the user is not explicitly requesting FP8 loading, check if a BF16 sibling repo exists (model_name + "-BF16") and redirect to it. This happens early in the loading flow before any quantization config processing. Also pass the modified model_config to auto_model.from_pretrained to avoid transformers re-reading the original config from the model repo. Tested with Ministral-3-3B in 4-bit and 16-bit modes. Both now load and train correctly. * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Simplify FP8 condition and narrow exception handling Simplify the load_in_fp8 check (works for bool and string values). Narrow inner except to KeyError and add comment for outer except. * Warn user when FP8 model has no BF16 sibling for redirect Previously the except block silently fell through with `pass`, so users would get a confusing BNB dtype error later. Now prints a clear message explaining the FP8 situation and suggesting load_in_fp8=True or uploading a BF16 version. * Fix FP8 redirect state corruption and add fbgemm_fp8 support - Fix state corruption: model_name was reassigned before AutoConfig.from_pretrained, so if config fetch failed, model_name pointed to BF16 repo while auto_config still had FP8. Now only updates state after both checks succeed. - Save original model_name so warning message is correct even on failure. - Handle fbgemm_fp8 quant method in addition to fp8. * Extract FP8 redirect to shared _redirect_fp8_to_bf16() in _utils.py Addresses reviewer feedback: - Move FP8 redirect logic to a shared function callable from both vision.py (FastBaseModel) and llama.py (FastLlamaModel) - Raise RuntimeError instead of warning when BF16 sibling not found - Add FP8 redirect to llama.py for text-only model loading path * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Add Ministral 3B/8B/14B mapper entries Adds all 9 Ministral model variants to the mapper: - Instruct (3B, 8B, 14B) with FP8 variant mappings - Base (3B, 8B, 14B) - Reasoning (3B, 8B, 14B) This routes mistralai/Ministral-* to unsloth/Ministral-* repos (BF16 weights), which also avoids the FP8 config issue for the standard loading path through loader.py. * Add FP8 mapper entries for Mistral-Small-3.2 and Magistral-Small-2509 --------- Co-authored-by: Ubuntu Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- unsloth/models/_utils.py | 54 +++++++++++++++++++++ unsloth/models/llama.py | 10 ++++ unsloth/models/mapper.py | 101 +++++++++++++++++++++++++++++++++++---- unsloth/models/vision.py | 21 ++++++++ 4 files changed, 176 insertions(+), 10 deletions(-) diff --git a/unsloth/models/_utils.py b/unsloth/models/_utils.py index 451f60091f..e7cc5ff1ef 100644 --- a/unsloth/models/_utils.py +++ b/unsloth/models/_utils.py @@ -74,6 +74,7 @@ __all__ = [ "dequantize_module_weight", "patch_hf_quantizer", "verify_fp8_support_if_applicable", + "_redirect_fp8_to_bf16", "_get_inference_mode_context_manager", "hf_login", "is_moe_model", @@ -2570,6 +2571,59 @@ def patch_hf_quantizer(): patch_hf_quantizer() +def _redirect_fp8_to_bf16( + model_name, auto_config, load_in_fp8, token, trust_remote_code +): + """ + Detect FP8 quantization in model config and redirect to BF16 sibling. + + Models shipping FP8 as default (e.g. mistralai/Ministral-3-*B-Instruct) + cannot be loaded with BNB 4-bit/8-bit or 16-bit mode. This detects + quant_method in ("fp8", "fbgemm_fp8") and redirects to {model_name}-BF16. + + Redirect is SKIPPED when load_in_fp8 is truthy (True or 'block'), + meaning the user explicitly wants FP8 loading. + + Returns (model_name, auto_config) -- possibly updated. + """ + if not hasattr(auto_config, "quantization_config"): + return model_name, auto_config + + _qc = auto_config.quantization_config + _qm = ( + _qc.get("quant_method", "") + if isinstance(_qc, dict) + else getattr(_qc, "quant_method", "") + ) + if _qm not in ("fp8", "fbgemm_fp8") or load_in_fp8: + return model_name, auto_config + + _bf16_name = model_name.rstrip("/") + "-BF16" + _original_name = model_name + try: + from huggingface_hub import model_info as _hf_model_info + from transformers import AutoConfig + + _hf_model_info(_bf16_name, token = token) + _bf16_config = AutoConfig.from_pretrained( + _bf16_name, + token = token, + trust_remote_code = trust_remote_code, + ) + print( + f"Unsloth: {_original_name} uses FP8 weights. " + f"Redirecting to {_bf16_name}." + ) + return _bf16_name, _bf16_config + except Exception: + raise RuntimeError( + f"Unsloth: {_original_name} uses FP8 weights but no BF16 version " + f"was found at {_bf16_name}.\n" + f"Loading FP8 weights with BitsAndBytes or in 16-bit will fail.\n" + f"Set load_in_fp8=True to use FP8 mode, or upload a BF16 version." + ) + + def verify_fp8_support_if_applicable(model_config): quant_method = get_quant_type(model_config) if quant_method in ["fbgemm_fp8", "fp8"] and DEVICE_TYPE != "cuda": diff --git a/unsloth/models/llama.py b/unsloth/models/llama.py index a3fc9ab49a..342366f02e 100644 --- a/unsloth/models/llama.py +++ b/unsloth/models/llama.py @@ -25,6 +25,7 @@ from ._utils import move_to_device from ._utils import ( _get_inference_mode_context_manager, _prepare_model_for_qat, + _redirect_fp8_to_bf16, ) from .loader_utils import _get_fp8_mode_and_check_settings from ..utils.packing import ( @@ -2228,6 +2229,15 @@ class FastLlamaModel: token = token, attn_implementation = "sdpa", ) + # Handle FP8 models: redirect to BF16 sibling when the model ships with + # FP8 weights. Redirect is skipped when load_in_fp8 is truthy (True or 'block'). + model_name, model_config = _redirect_fp8_to_bf16( + model_name, + model_config, + load_in_fp8, + token, + trust_remote_code, + ) model_config.model_name = model_name model_max_seq_length = model_config.max_position_embeddings diff --git a/unsloth/models/mapper.py b/unsloth/models/mapper.py index e7296291f1..ec7a7a8046 100644 --- a/unsloth/models/mapper.py +++ b/unsloth/models/mapper.py @@ -965,11 +965,18 @@ __INT_TO_FLOAT_MAPPER = \ "mistralai/Magistral-Small-2506", "unsloth/Magistral-Small-2506-bnb-4bit", ), - "unsloth/Mistral-Small-3.2-24B-Instruct-2506-unsloth-bnb-4bit" : ( - "unsloth/Mistral-Small-3.2-24B-Instruct-2506", - "mistralai/Mistral-Small-3.2-24B-Instruct-2506", - "unsloth/Mistral-Small-3.2-24B-Instruct-2506-bnb-4bit", - ), + "unsloth/Mistral-Small-3.2-24B-Instruct-2506-unsloth-bnb-4bit" : { + "8" : ( + "mistralai/Mistral-Small-3.2-24B-Instruct-2506", + "unsloth/Mistral-Small-3.2-24B-Instruct-2506-FP8", + "unsloth/Mistral-Small-3.2-24B-Instruct-2506-FP8", + ), + "16" : ( + "unsloth/Mistral-Small-3.2-24B-Instruct-2506", + "mistralai/Mistral-Small-3.2-24B-Instruct-2506", + "unsloth/Mistral-Small-3.2-24B-Instruct-2506-bnb-4bit", + ), + }, "unsloth/gemma-3n-E4B-it-unsloth-bnb-4bit" : ( "unsloth/gemma-3n-E4B-it", "google/gemma-3n-E4B-it", @@ -1056,11 +1063,18 @@ __INT_TO_FLOAT_MAPPER = \ "mistralai/Magistral-Small-2507", "unsloth/Magistral-Small-2507-bnb-4bit", ), - "unsloth/Magistral-Small-2509-unsloth-bnb-4bit" : ( - "unsloth/Magistral-Small-2509", - "mistralai/Magistral-Small-2509", - "unsloth/Magistral-Small-2509-bnb-4bit", - ), + "unsloth/Magistral-Small-2509-unsloth-bnb-4bit" : { + "8" : ( + "mistralai/Magistral-Small-2509", + "unsloth/Magistral-Small-2509-FP8-Dynamic", + "unsloth/Magistral-Small-2509-FP8-Dynamic", + ), + "16" : ( + "unsloth/Magistral-Small-2509", + "mistralai/Magistral-Small-2509", + "unsloth/Magistral-Small-2509-bnb-4bit", + ), + }, "unsloth/Apertus-70B-Instruct-2509-unsloth-bnb-4bit" : ( "unsloth/Apertus-70B-Instruct-2509", "swiss-ai/Apertus-70B-2509", @@ -1256,6 +1270,73 @@ __INT_TO_FLOAT_MAPPER = \ "google/functiongemma-270m-it", "unsloth/functiongemma-270m-it-unsloth-bnb-4bit", ), + # Ministral 3 models + "unsloth/Ministral-3-3B-Instruct-2512-unsloth-bnb-4bit" : { + "8" : ( + "mistralai/Ministral-3-3B-Instruct-2512", + "unsloth/Ministral-3-3B-Instruct-2512-FP8", + "unsloth/Ministral-3-3B-Instruct-2512-FP8", + ), + "16" : ( + "unsloth/Ministral-3-3B-Instruct-2512", + "mistralai/Ministral-3-3B-Instruct-2512", + "unsloth/Ministral-3-3B-Instruct-2512-bnb-4bit", + ), + }, + "unsloth/Ministral-3-3B-Base-2512-unsloth-bnb-4bit" : ( + "unsloth/Ministral-3-3B-Base-2512", + "mistralai/Ministral-3-3B-Base-2512", + "unsloth/Ministral-3-3B-Base-2512-bnb-4bit", + ), + "unsloth/Ministral-3-3B-Reasoning-2512-unsloth-bnb-4bit" : ( + "unsloth/Ministral-3-3B-Reasoning-2512", + "mistralai/Ministral-3-3B-Reasoning-2512", + "unsloth/Ministral-3-3B-Reasoning-2512-bnb-4bit", + ), + "unsloth/Ministral-3-8B-Instruct-2512-unsloth-bnb-4bit" : { + "8" : ( + "mistralai/Ministral-3-8B-Instruct-2512", + "unsloth/Ministral-3-8B-Instruct-2512-FP8", + "unsloth/Ministral-3-8B-Instruct-2512-FP8", + ), + "16" : ( + "unsloth/Ministral-3-8B-Instruct-2512", + "mistralai/Ministral-3-8B-Instruct-2512", + "unsloth/Ministral-3-8B-Instruct-2512-bnb-4bit", + ), + }, + "unsloth/Ministral-3-8B-Base-2512-unsloth-bnb-4bit" : ( + "unsloth/Ministral-3-8B-Base-2512", + "mistralai/Ministral-3-8B-Base-2512", + "unsloth/Ministral-3-8B-Base-2512-bnb-4bit", + ), + "unsloth/Ministral-3-8B-Reasoning-2512-unsloth-bnb-4bit" : ( + "unsloth/Ministral-3-8B-Reasoning-2512", + "mistralai/Ministral-3-8B-Reasoning-2512", + "unsloth/Ministral-3-8B-Reasoning-2512-bnb-4bit", + ), + "unsloth/Ministral-3-14B-Instruct-2512-unsloth-bnb-4bit" : { + "8" : ( + "mistralai/Ministral-3-14B-Instruct-2512", + "unsloth/Ministral-3-14B-Instruct-2512-FP8", + "unsloth/Ministral-3-14B-Instruct-2512-FP8", + ), + "16" : ( + "unsloth/Ministral-3-14B-Instruct-2512", + "mistralai/Ministral-3-14B-Instruct-2512", + "unsloth/Ministral-3-14B-Instruct-2512-bnb-4bit", + ), + }, + "unsloth/Ministral-3-14B-Base-2512-unsloth-bnb-4bit" : ( + "unsloth/Ministral-3-14B-Base-2512", + "mistralai/Ministral-3-14B-Base-2512", + "unsloth/Ministral-3-14B-Base-2512-bnb-4bit", + ), + "unsloth/Ministral-3-14B-Reasoning-2512-unsloth-bnb-4bit" : ( + "unsloth/Ministral-3-14B-Reasoning-2512", + "mistralai/Ministral-3-14B-Reasoning-2512", + "unsloth/Ministral-3-14B-Reasoning-2512-bnb-4bit", + ), } INT_TO_FLOAT_MAPPER = {} diff --git a/unsloth/models/vision.py b/unsloth/models/vision.py index 2d6c7ac93f..56cd615a49 100644 --- a/unsloth/models/vision.py +++ b/unsloth/models/vision.py @@ -30,6 +30,7 @@ from ..kernels import ( post_patch_loss_function, ) from ._utils import __version__, importlib_version, _prepare_model_for_qat +from ._utils import _redirect_fp8_to_bf16 from ._utils import * from .loader_utils import _get_fp8_mode_and_check_settings from ..save import patch_saving_functions @@ -611,6 +612,24 @@ class FastBaseModel: model_class = None flex_attn_impl = prefer_flex_attn_if_supported(model_class, auto_config) + # Handle FP8 models: redirect to BF16 sibling when the model ships with + # FP8 weights (e.g. Ministral-3-3B-Instruct-2512). FP8 weights cannot be + # directly loaded by BNB, and the FP8 quantization config can cause issues + # even for 16-bit loading. + # Redirect is skipped when load_in_fp8 is truthy (True or 'block'). + model_name, auto_config = _redirect_fp8_to_bf16( + model_name, + auto_config, + load_in_fp8, + token, + trust_remote_code, + ) + # Re-resolve model_class after potential config change + try: + model_class = auto_model._model_mapping[auto_config.__class__] + except KeyError: + pass + default_attn_impl = "flex_attention" if flex_attn_impl else "sdpa" if not ("attn_implementation" in kwargs): kwargs["attn_implementation"] = default_attn_impl @@ -759,6 +778,7 @@ class FastBaseModel: if hasattr(auto_config, "attn_implementation"): setattr(auto_config, "attn_implementation", config_attn_impl) model_config = auto_config + verify_fp8_support_if_applicable(model_config) raise_handler = RaiseUninitialized() @@ -767,6 +787,7 @@ class FastBaseModel: load_in_fp8 = kwargs.pop("load_in_fp8", None) model = auto_model.from_pretrained( model_name, + config = model_config, device_map = device_map, # torch_dtype = torch_dtype, # Transformers removed torch_dtype # quantization_config = bnb_config, From 1e08d111807f8718ea2d3adebac328504da75370 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Feb 2026 03:14:12 -0800 Subject: [PATCH 32/63] Fix Nemotron-H and Nemotron-VL model support (#4105) * Fix Nemotron-H and Nemotron-VL model support - Add Mamba kernel precision settings for Nemotron-H hybrid models - Fix VL model auto_model selection for models that only register AutoModelForCausalLM in their auto_map - Skip quantization of out_proj for Nemotron-H Mamba layers * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Simplify VLM auto_model selection logic Reduce three branches to two since the first and third both assign AutoModelForVision2Seq. The simplified condition checks whether the auto_map exclusively registers AutoModelForCausalLM without the VLM class, and defaults to AutoModelForVision2Seq otherwise. * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --------- Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- unsloth/models/loader.py | 25 ++++++++++++++++++++++++- unsloth/models/vision.py | 11 +++++++++-- 2 files changed, 33 insertions(+), 3 deletions(-) diff --git a/unsloth/models/loader.py b/unsloth/models/loader.py index 4054f1b7f5..04c15bb2c3 100644 --- a/unsloth/models/loader.py +++ b/unsloth/models/loader.py @@ -1084,6 +1084,15 @@ class FastModel(FastBaseModel): ";" "os.environ['TRITON_F32_DEFAULT'] = 'ieee'" ) + elif "nemotron_h" in model_types_all: + # NemotronH (hybrid Mamba-2 + Transformer) uses same Mamba kernels as Falcon-H1 + # Mamba kernels need float32 Triton precision + os.environ["UNSLOTH_FORCE_CUSTOM_DTYPE"] = ( + "float16;torch.float32;torch.float16;" + "if name.endswith(('q_proj', 'k_proj', 'v_proj', 'o_proj', 'gate_proj', 'up_proj', 'down_proj', 'head')): module.to(torch.float16)" + ";" + "os.environ['TRITON_F32_DEFAULT'] = 'ieee'" + ) elif "gpt_oss" in model_types_all: os.environ["UNSLOTH_DISABLE_STATIC_GENERATION"] = "1" if not load_in_4bit: @@ -1281,7 +1290,21 @@ class FastModel(FastBaseModel): is_vlm = any(x.endswith("ForConditionalGeneration") for x in architectures) is_vlm = is_vlm or hasattr(model_config, "vision_config") if auto_model is None: - auto_model = AutoModelForVision2Seq if is_vlm else AutoModelForCausalLM + if is_vlm: + # Check if the model's auto_map supports the VLM auto class. + # Some VL models (e.g. Nemotron-VL) only register AutoModelForCausalLM + # in their auto_map, not AutoModelForImageTextToText/AutoModelForVision2Seq. + _auto_map = getattr(model_config, "auto_map", {}) or {} + _vlm_class_name = AutoModelForVision2Seq.__name__ + if ( + "AutoModelForCausalLM" in _auto_map + and _vlm_class_name not in _auto_map + ): + auto_model = AutoModelForCausalLM + else: + auto_model = AutoModelForVision2Seq + else: + auto_model = AutoModelForCausalLM load_in_4bit_kwargs = load_in_4bit load_in_8bit_kwargs = load_in_8bit diff --git a/unsloth/models/vision.py b/unsloth/models/vision.py index 56cd615a49..24e72e1535 100644 --- a/unsloth/models/vision.py +++ b/unsloth/models/vision.py @@ -654,18 +654,25 @@ class FastBaseModel: raise RuntimeError( "Unsloth: Can only load in 4bit or 8bit or 16bit, not a combination!" ) + _skip_modules = SKIP_QUANTIZATION_MODULES.copy() + # Nemotron-H uses 'mixer' (not 'mamba') for Mamba layers. + # Mamba fused kernels pass out_proj.weight directly to F.linear, + # which fails with quantized Params4bit. Skip out_proj from quantization. + if any(mt == "nemotron_h" for mt in (model_types or [])): + _skip_modules.append("out_proj") + if load_in_4bit: bnb_config = BitsAndBytesConfig( load_in_4bit = True, bnb_4bit_use_double_quant = True, bnb_4bit_quant_type = "nf4", bnb_4bit_compute_dtype = bnb_compute_dtype, - llm_int8_skip_modules = SKIP_QUANTIZATION_MODULES.copy(), + llm_int8_skip_modules = _skip_modules, ) elif load_in_8bit: bnb_config = BitsAndBytesConfig( load_in_8bit = True, - llm_int8_skip_modules = SKIP_QUANTIZATION_MODULES.copy(), + llm_int8_skip_modules = _skip_modules, ) elif load_in_16bit: bnb_config = None From d9fdada470f5e9409db6abf89b33ebd3898dfafa Mon Sep 17 00:00:00 2001 From: Irfan Ali Date: Wed, 25 Feb 2026 13:39:16 +0100 Subject: [PATCH 33/63] =?UTF-8?q?fix:=20correct=20gpt-oss=20Ollama=20gener?= =?UTF-8?q?ation=20prompt=20and=20add=20quantization=20wa=E2=80=A6=20(#408?= =?UTF-8?q?7)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * Warn when save_pretrained_gguf overrides quantization to MXFP4 for GPT-OSS GPT-OSS only supports MXFP4 format. If the user passes a different quantization_method, log a warning via logger.warning_once before overriding. Pass quantization_method=None to suppress the warning. * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --------- Co-authored-by: Daniel Han Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- unsloth/save.py | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/unsloth/save.py b/unsloth/save.py index a07575c35d..66d5ae2d60 100644 --- a/unsloth/save.py +++ b/unsloth/save.py @@ -1925,6 +1925,20 @@ def unsloth_save_pretrained_gguf( arguments["push_to_hub"] = False # We handle upload ourselves # GPT-OSS needs mxfp4 save method if is_gpt_oss: + if quantization_method is not None: + _qm = ( + quantization_method + if isinstance(quantization_method, (list, tuple)) + else [quantization_method] + ) + _ignored = [q for q in _qm if str(q).lower() != "mxfp4"] + if _ignored: + logger.warning_once( + f"Unsloth: GPT-OSS does not support GGUF quantization " + f"(requested: {', '.join(str(q) for q in _ignored)}). " + f"Overriding to MXFP4 format. " + f"Pass quantization_method=None to suppress this warning." + ) arguments["save_method"] = "mxfp4" else: arguments["save_method"] = "merged_16bit" From a409c943b66aee8436208f54e19dc555b0848028 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Feb 2026 06:34:21 -0800 Subject: [PATCH 34/63] Add resilience to TRL internal API reclassification (#4111) * Add resilience to TRL internal API reclassification TRL is moving toward v1.0 and will reclassify several currently-importable symbols as internal with no stability guarantees. This adds try/except cascading imports with local fallbacks so Unsloth keeps working regardless of whether TRL removes, moves, or restructures these symbols. Changes: - rl.py: Add try/except cascade for unwrap_model_for_generation with local contextmanager fallback. Wire sanitize_logprob from RL_REPLACEMENTS into the compiled trainer template (same pipeline as selective_log_softmax and other global functions). Add import math and import logging to the template header. - rl_replacements.py: Remove inline import of sanitize_logprob from trl.scripts.vllm_serve in the regex replacement. The function is now a module-level global in the compiled file. - tokenizer_utils.py: Wrap dynamic exec import with per-item fallback so a single removed symbol does not break the entire bulk import. Depends on unslothai/unsloth-zoo#516. Tested across all TRL versions from 0.22.2 through 0.29.0.dev0 (git main). Training losses and grad norms are bit-identical to unpatched runs. * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --------- Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- unsloth/models/rl.py | 41 ++++++++++++++++++++++++++++++- unsloth/models/rl_replacements.py | 3 ++- unsloth/tokenizer_utils.py | 9 ++++++- 3 files changed, 50 insertions(+), 3 deletions(-) diff --git a/unsloth/models/rl.py b/unsloth/models/rl.py index 181e9479df..fbb32b2c50 100755 --- a/unsloth/models/rl.py +++ b/unsloth/models/rl.py @@ -94,7 +94,40 @@ def vLLMSamplingParams(**kwargs): def PatchRL(FastLanguageModel): - from trl.models.utils import unwrap_model_for_generation + try: + from trl.models.utils import unwrap_model_for_generation + except ImportError: + try: + from trl.models import unwrap_model_for_generation + except ImportError: + # Local fallback -- TRL removed or moved this symbol + from contextlib import contextmanager as _cm + + @_cm + def unwrap_model_for_generation( + model, accelerator, gather_deepspeed3_params = True + ): + unwrapped_model = accelerator.unwrap_model(model) + is_gc = getattr(unwrapped_model, "is_gradient_checkpointing", False) + if is_gc: + unwrapped_model.gradient_checkpointing_disable() + if ( + getattr(accelerator, "state", None) is not None + and getattr(accelerator.state, "deepspeed_plugin", None) is not None + and accelerator.state.deepspeed_plugin.zero_stage == 3 + ): + if not gather_deepspeed3_params: + yield accelerator.unwrap_model(model) + else: + import deepspeed + + with deepspeed.zero.GatheredParameters(model.parameters()): + yield accelerator.unwrap_model(model) + else: + yield unwrapped_model + if is_gc: + unwrapped_model.gradient_checkpointing_enable() + from contextlib import contextmanager @contextmanager @@ -253,9 +286,12 @@ create_completion_attention_mask = RL_REPLACEMENTS["create_completion_attention_ left_pack_padding = RL_REPLACEMENTS["left_pack_padding"] align_logprobs_with_mask = RL_REPLACEMENTS["align_logprobs_with_mask"] autotune_batch_and_chunks = RL_REPLACEMENTS["grpo_autotune_batch_and_chunks"] +sanitize_logprob = RL_REPLACEMENTS["sanitize_logprob"] RLTrainer_replacement = ''' import os +import math +import logging from typing import * from dataclasses import dataclass, field from packaging.version import Version @@ -324,6 +360,7 @@ torch_compile_options = {{ {left_pack_padding_code} {align_logprobs_with_mask_code} {autotune_batch_and_chunks_code} +{sanitize_logprob_code} {RL_pre} @@ -1228,6 +1265,7 @@ def _patch_trl_rl_trainers(trainer_file = "grpo_trainer"): left_pack_padding_code = inspect.getsource(left_pack_padding) align_logprobs_with_mask_code = inspect.getsource(align_logprobs_with_mask) autotune_batch_and_chunks_code = inspect.getsource(autotune_batch_and_chunks) + sanitize_logprob_code = inspect.getsource(sanitize_logprob) # Get final source code RLTrainer_source = RLTrainer_replacement.format( RLTrainer_name = RLTrainer_name, @@ -1256,6 +1294,7 @@ def _patch_trl_rl_trainers(trainer_file = "grpo_trainer"): autotune_batch_and_chunks_code = autotune_batch_and_chunks_code, left_pack_padding_code = left_pack_padding_code, align_logprobs_with_mask_code = align_logprobs_with_mask_code, + sanitize_logprob_code = sanitize_logprob_code, ) if RLTrainer_name == "GRPOTrainer": diff --git a/unsloth/models/rl_replacements.py b/unsloth/models/rl_replacements.py index 27f00f10f3..b3a55440f9 100755 --- a/unsloth/models/rl_replacements.py +++ b/unsloth/models/rl_replacements.py @@ -355,8 +355,9 @@ def grpo_trainer__generate_and_score_completions(function_name, function): re.DOTALL | re.MULTILINE, ) + # sanitize_logprob is injected as a module-level function via RLTrainer_replacement + # template in rl.py (from RL_REPLACEMENTS), so just reference it directly here. replacement_text = ( - r"\1from trl.scripts.vllm_serve import sanitize_logprob\n" r"\1all_logprobs = [\n" r"\1 [sanitize_logprob(next(iter(logprob.values()))) for logprob in output.logprobs]\n" r"\1 for outputs in all_outputs\n" diff --git a/unsloth/tokenizer_utils.py b/unsloth/tokenizer_utils.py index 1c107bd84f..c445879df7 100644 --- a/unsloth/tokenizer_utils.py +++ b/unsloth/tokenizer_utils.py @@ -1007,7 +1007,14 @@ def patch_sft_trainer_tokenizer(): function = function.replace(replacer, check_text + replacer) x = [x for x in all_imports if x in function] - exec(f"from trl.trainer.sft_trainer import ({','.join(x)})", locals()) + try: + exec(f"from trl.trainer.sft_trainer import ({','.join(x)})", locals()) + except ImportError: + for _item in x: + try: + exec(f"from trl.trainer.sft_trainer import {_item}", locals()) + except ImportError: + pass exec(function, locals(), globals()) exec( f"trl.trainer.sft_trainer.SFTTrainer.{function_name} = {function_name}", From f04404972bc1d0678b45cf9ada71f303255154b4 Mon Sep 17 00:00:00 2001 From: Datta Nimmaturi Date: Wed, 25 Feb 2026 20:22:18 +0530 Subject: [PATCH 35/63] Allow fp8 for non fast inference (#3904) * Allow fp8 for non fast inference * Extensive fp8 alow and quantizer patch * Clean up commented-out code, duplicate import, and revert unnecessary Version() changes - Delete commented-out FP8 fast_inference guard in FastModel (loader.py) instead of leaving it commented -- matches FastLanguageModel which was properly deleted - Delete commented-out fast_inference guard in loader_utils.py - Remove duplicate `from transformers import GenerationConfig, CompileConfig` in vision.py (line 112 already imports both plus AutoConfig) - Revert Version(trl.__version__) back to Version(trl) in trainer.py -- trainer.py imports Version from unsloth_zoo.utils which already handles module objects --------- Co-authored-by: Daniel Han --- unsloth/models/_utils.py | 8 ++++++++ unsloth/models/loader.py | 13 ------------- unsloth/models/loader_utils.py | 4 ---- 3 files changed, 8 insertions(+), 17 deletions(-) diff --git a/unsloth/models/_utils.py b/unsloth/models/_utils.py index e7cc5ff1ef..5f4927b92f 100644 --- a/unsloth/models/_utils.py +++ b/unsloth/models/_utils.py @@ -2567,6 +2567,14 @@ def patch_hf_quantizer(): except Exception as e: logger.warning(f"Failed to patch FbgemmFp8HfQuantizer. Error {e}") + try: + from transformers.quantizers.quantizer_torchao import TorchAoHfQuantizer + + TorchAoHfQuantizer.is_trainable = property(make_trainable) + TorchAoHfQuantizer.is_qat_trainable = property(make_trainable) + except Exception as e: + logger.warning(f"Failed to patch TorchAoHfQuantizer. Error {e}") + patch_hf_quantizer() diff --git a/unsloth/models/loader.py b/unsloth/models/loader.py index 04c15bb2c3..72a6d7d10f 100644 --- a/unsloth/models/loader.py +++ b/unsloth/models/loader.py @@ -261,12 +261,6 @@ class FastLanguageModel(FastLlamaModel): fast_inference = False break - # [TODO] For now fast_inference only works with fast_inference ie vLLM - if load_in_fp8 != False: - if not fast_inference: - raise NotImplementedError( - "Unsloth: set `fast_inference = True` when doing `load_in_fp8`." - ) # Check if 4bit is allowed specifically for AMD if not ALLOW_BITSANDBYTES and not use_exact_model_name: if load_in_4bit or load_in_8bit or model_name.lower().endswith("-bnb-4bit"): @@ -876,13 +870,6 @@ class FastModel(FastBaseModel): fast_inference = False break - # [TODO] For now fast_inference only works with fast_inference ie vLLM - if load_in_fp8 != False: - if not fast_inference: - raise NotImplementedError( - "Unsloth: set `fast_inference = True` when doing `load_in_fp8`." - ) - # Find FP8, BnB 4bit, other mapped names old_model_name = model_name fp8_mode = None diff --git a/unsloth/models/loader_utils.py b/unsloth/models/loader_utils.py index f17ebc53e3..40ac49ca78 100644 --- a/unsloth/models/loader_utils.py +++ b/unsloth/models/loader_utils.py @@ -350,10 +350,6 @@ def _get_fp8_mode_and_check_settings( raise ValueError( f"Unsloth: `load_in_fp8` can only be 'row' or 'block', got '{fp8_mode}'" ) - if not fast_inference: - raise ValueError( - "Unsloth: `load_in_fp8` is only supported for `fast_inference` for now" - ) if full_finetuning: raise ValueError( "Unsloth: `load_in_fp8` is not compatible with full finetuning" From 73e2ac372aded01a8bd3eb0bb1849556f5c038b4 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Feb 2026 07:17:16 -0800 Subject: [PATCH 36/63] Support Python 3.14 in package metadata (#4113) --- pyproject.toml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/pyproject.toml b/pyproject.toml index bd6107626b..d21ceb490f 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -7,7 +7,7 @@ name = "unsloth" dynamic = ["version"] description = "2-5X faster training, reinforcement learning & finetuning" readme = "README.md" -requires-python = ">=3.9,<3.14" +requires-python = ">=3.9,<3.15" license = "Apache-2.0" keywords = ["ai", "llm", "reinforcement learning", "machine learning", "artificial intelligence", "pytorch"] authors = [ From 9bad70b0c39efa3da4c9a8ceaf74439ebdd397e1 Mon Sep 17 00:00:00 2001 From: DoubleMathew Date: Wed, 25 Feb 2026 09:21:04 -0600 Subject: [PATCH 37/63] Fix/pr 3699 leftpad prefill main (#4100) * Fix left-padding masks and positions in batched decode/prefill * Fix batched generation with left padding * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Fix attention mask handling, padding_idx zeroing, and Mistral batched generation 1. attention_dispatch.py: Fall back from flash/xformers to SDPA when an attention_mask is present, since flash attention only supports causal masking via flag and cannot consume arbitrary padding masks. 2. gemma2.py: Apply attention_mask during decode inference for bsz > 1. Guard against boolean SWA/GA flags with isinstance check. Slice mask to match K/V length when sliding window is active. Remove dead commented-out SDPA branch (SDPA does not support softcapping). 3. granite.py: Apply attention_mask during decode inference for bsz > 1. Remove dead commented-out SDPA branch and misleading comment. 4. mistral.py: Fix 2D-to-4D padding mask conversion -- convert 0/1 mask to additive format (0 for keep, -inf for mask) before combining with the causal mask. Force SDPA backend when attention_mask is present. 5. llama.py: Skip zeroing embed_tokens.weight[padding_idx] when the embedding is weight-tied to lm_head, since zeroing the shared weight forces logit(pad) = 0 which is higher than real token logits in models like Gemma, causing the decoder to emit pad tokens as gibberish. Also add eos != pad guard, clean up unused _seq_length variable, and fix get_max_cache_shape handling. 6. vision.py: Same padding_idx fix as llama.py for the vision model loading path. Tested on gemma-2b-it, gemma-2-2b-it, Llama-3.2-1B, Mistral-7B-v0.3, Qwen2.5-0.5B, Qwen3-0.6B with flash-attn 2.8.3 active. All outputs coherent, zero crashes, zero resize warnings. * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Inference path optimizations: eliminate per-layer GPU-CPU sync, cache inspect.signature, add Granite SDPA split * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * More inference path optimizations across model files - gemma: hoist rotary_seq_len computation to model level (eliminates N per-layer GPU-CPU syncs from position_ids.max().item()), pre-convert attention mask to bool once for all layers, use scalar float multiply instead of torch.tensor allocation for embedding scaling - gemma2: use in-place tanh_() for softcap attention, use scalar float multiply for embedding scaling - granite: pre-convert attention mask to bool once for all layers - cohere: use in-place neg_() for rotary embedding (consistent with all other model files) - falcon_h1: use in-place mul_() for key_multiplier scaling - llama: use in-place tanh_() for logit softcapping * Revert scalar multiply for Gemma/Gemma2 embedding scaling The original torch.tensor(..., dtype=hidden_states.dtype) is intentional: sqrt(3072) rounds to 55.5 in bfloat16 vs 55.4256 in float32. A plain scalar multiply may compute at higher precision internally, producing different results. Restore the explicit dtype-cast tensor to match the training path in LlamaModel_fast_forward. * Fix hardcoded cuda:0 device strings and add Cohere .eq(0) bool mask Replace 15 hardcoded "cuda:0" with f"{DEVICE_TYPE_TORCH}:0" across gemma.py, gemma2.py, cohere.py, and falcon_h1.py to support multi-GPU and non-CUDA devices (XPU, etc.). Add .eq(0) bool mask pre-conversion in CohereModel_fast_forward_inference for batched inference consistency with llama.py, granite.py, and gemma.py. * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Disable flex_attention for Mllama (Llama 3.2 Vision) Mllama's _update_causal_mask uses the deprecated make_flex_block_causal_mask which creates a BlockMask with Q_LEN=KV_LEN=total_seq_len. During decode with KV cache, q_len=1 but the block_mask still has Q_LEN=total_seq_len, causing a ValueError. This is an upstream transformers issue -- newer models use flex_attention_mask from masking_utils which handles decode correctly via cache_position, but mllama has not been updated yet. Add mllama to the exclusion list in prefer_flex_attn_if_supported alongside gpt_oss so it falls back to sdpa, which works correctly for both training and inference. * Fix off-by-one in sliding window K/V slicing for gemma2, qwen3, falcon_h1, cohere The old formula `slicing_tokens = 1 - sliding_window` uses negative indexing that keeps `sliding_window - 1` tokens instead of `sliding_window`. For example with sliding_window=32 and kv_seq_len=100, `1-32 = -31` keeps indices 69..99 (31 tokens) instead of the correct 68..99 (32 tokens). Replace with `start = kv_seq_len - sliding_window` to match the fix already applied in llama.py and the canonical definition in transformers masking_utils (sliding_window_overlay: kv_idx > q_idx - W, which keeps exactly W tokens). Also add attention_mask slicing after K/V trim in qwen3, falcon_h1, and cohere to prevent mask/K dimension mismatch during batched SDPA inference, matching the pattern already used in llama.py. Currently only gemma2 (sliding_window=4096) is actively affected. The other three models have sliding_window=None in their configs so the code path is not triggered, but this keeps it correct for any future models that set it. * Fix Gemma2 softcapping order: apply mask after softcap, not before The attention mask must be applied AFTER logit softcapping, not before. Both the Google DeepMind reference implementation (google-deepmind/gemma, gm/nn/_modules.py lines 254-277) and transformers' eager_attention_forward (gemma2/modeling_gemma2.py lines 187-193) use this order: 1. logits = Q @ K^T * scale 2. logits = tanh(logits / softcap) * softcap # softcap first 3. logits = logits + mask # mask after 4. probs = softmax(logits) The PR had the mask addition before softcapping, which causes tanh to clamp the -inf mask values to -softcap instead of preserving them as -inf for softmax. While the practical impact is small (masked positions get ~1e-23 probability instead of exact zero), this should match upstream. * Clarify GQA condition precedence and remove stale comments Add explicit parentheses to grouped query attention conditions in llama.py, qwen3.py, granite.py to make operator precedence clear. The expression `bsz == 1 or not X and Y` relies on Python binding `not` > `and` > `or` which is correct but easy to misread. Remove dead commented-out code (`# else: # Knn, Vnn = Knn, Vnn`) and stale mask comments (`# if attention_mask ...`) from the bsz==1 fast path in llama, qwen3, cohere, falcon_h1, gemma2 inference functions. These were leftover from the pre-batched-inference structure and no longer apply. --------- Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> Co-authored-by: Daniel Han --- unsloth/models/_utils.py | 11 +- unsloth/models/cohere.py | 47 +++--- unsloth/models/falcon_h1.py | 34 +++-- unsloth/models/gemma.py | 12 +- unsloth/models/gemma2.py | 39 ++--- unsloth/models/granite.py | 57 ++++--- unsloth/models/llama.py | 225 +++++++++++++++++++++------- unsloth/models/mistral.py | 22 ++- unsloth/models/qwen3.py | 42 ++++-- unsloth/models/vision.py | 15 +- unsloth/utils/attention_dispatch.py | 83 +++++++++- 11 files changed, 436 insertions(+), 151 deletions(-) diff --git a/unsloth/models/_utils.py b/unsloth/models/_utils.py index 5f4927b92f..dafdec4e1e 100644 --- a/unsloth/models/_utils.py +++ b/unsloth/models/_utils.py @@ -235,11 +235,14 @@ def prefer_flex_attn_if_supported(model_class, config): model_class, "_supports_flex_attn", False ): return None - # GPT-OSS uses eager attention during inference since flex attention - # returns incorrect results (likely due to left padding issues). - # Skip setting flex_attention to avoid BlockMask type errors. + # GPT-OSS and Mllama use eager/sdpa attention during inference since + # flex attention returns incorrect results or errors out. + # GPT-OSS: left padding issues cause incorrect outputs. + # Mllama: _update_causal_mask uses make_flex_block_causal_mask which + # creates BlockMask with Q_LEN=KV_LEN=total_seq_len, but during + # decode q_len=1, causing ValueError. Needs transformers update. model_type = getattr(config, "model_type", "") if config else "" - if model_type == "gpt_oss": + if model_type in ("gpt_oss", "mllama"): return None if config is not None: setattr(config, "_attn_implementation", "flex_attention") diff --git a/unsloth/models/cohere.py b/unsloth/models/cohere.py index c33317ee02..4251f3acd9 100644 --- a/unsloth/models/cohere.py +++ b/unsloth/models/cohere.py @@ -188,7 +188,9 @@ def CohereDecoderLayer_fast_forward( self, "_flag_for_generation" ): # past_key_value is not None: out_weight = torch.empty( - self.input_layernorm.weight.shape, dtype = torch.float32, device = "cuda:0" + self.input_layernorm.weight.shape, + dtype = torch.float32, + device = f"{DEVICE_TYPE_TORCH}:0", ) # Self Attention @@ -254,6 +256,7 @@ def CohereAttention_fast_forward_inference( position_ids, do_prefill = False, attention_mask = None, + **kwargs, ): Xn = hidden_states bsz, _, hd = hidden_states.size() @@ -277,26 +280,28 @@ def CohereAttention_fast_forward_inference( self.paged_attention = torch.empty( (KV_CACHE_INCREMENT + seq_len + 1, 2, bsz, n_kv_heads, head_dim), dtype = dtype, - device = "cuda:0", + device = f"{DEVICE_TYPE_TORCH}:0", ) self.paged_attention_K = self.paged_attention[:, 0] self.paged_attention_V = self.paged_attention[:, 1] self.paged_attention_K[:seq_len] = K1.permute(2, 0, 1, 3) self.paged_attention_V[:seq_len] = V1.permute(2, 0, 1, 3) self.temp_QA = torch.empty( - (2, bsz, 1, attention_size), dtype = dtype, device = "cuda:0" + (2, bsz, 1, attention_size), dtype = dtype, device = f"{DEVICE_TYPE_TORCH}:0" ) self.temp_KV = torch.empty( - (2, bsz, 1, n_kv_heads * head_dim), dtype = dtype, device = "cuda:0" + (2, bsz, 1, n_kv_heads * head_dim), + dtype = dtype, + device = f"{DEVICE_TYPE_TORCH}:0", ) self.RH_Q = torch.empty( - (bsz, n_heads, 1, head_dim), dtype = dtype, device = "cuda:0" + (bsz, n_heads, 1, head_dim), dtype = dtype, device = f"{DEVICE_TYPE_TORCH}:0" ) # Mistral Nemo 12b has weird dimensions if attention_size != hidden_size: self.temp_O = torch.empty( - (1, bsz, hidden_size), dtype = dtype, device = "cuda:0" + (bsz, 1, hidden_size), dtype = dtype, device = f"{DEVICE_TYPE_TORCH}:0" ) else: self.temp_O = self.temp_QA[1][:, :, :hidden_size] @@ -304,17 +309,21 @@ def CohereAttention_fast_forward_inference( self.attention = torch.empty( (bsz, n_heads, 1, KV_CACHE_INCREMENT + seq_len), dtype = dtype, - device = "cuda:0", + device = f"{DEVICE_TYPE_TORCH}:0", ) self.scalar = 1.0 / math_sqrt(self.head_dim) self.half_head_dim = head_dim // 2 # Cohere has QK layernorms if self.use_qk_norm: self.q_norm_out_weight = torch.empty( - self.q_norm.weight.shape, dtype = torch.float32, device = "cuda:0" + self.q_norm.weight.shape, + dtype = torch.float32, + device = f"{DEVICE_TYPE_TORCH}:0", ) self.k_norm_out_weight = torch.empty( - self.k_norm.weight.shape, dtype = torch.float32, device = "cuda:0" + self.k_norm.weight.shape, + dtype = torch.float32, + device = f"{DEVICE_TYPE_TORCH}:0", ) else: self.q_norm_out_weight = None @@ -355,7 +364,7 @@ def CohereAttention_fast_forward_inference( RH_Q = self.RH_Q RH_Q[:, :, :, :h] = Qn[:, :, :, h:] RH_Q[:, :, :, h:] = Qn[:, :, :, :h] - torch.neg(RH_Q[:, :, :, :h], out = RH_Q[:, :, :, :h]) + RH_Q[:, :, :, :h].neg_() Qn *= cos Qn.addcmul_(RH_Q, sin) @@ -364,7 +373,7 @@ def CohereAttention_fast_forward_inference( ] # torch.empty((n_kv_heads, 1, head_dim), dtype = dtype, device = "cuda:0") RH_K[:, :, :, :h] = Kn[:, :, :, h:] RH_K[:, :, :, h:] = Kn[:, :, :, :h] - torch.neg(RH_K[:, :, :, :h], out = RH_K[:, :, :, :h]) + RH_K[:, :, :, :h].neg_() Kn *= cos Kn.addcmul_(RH_K, sin) @@ -379,10 +388,11 @@ def CohereAttention_fast_forward_inference( # Handle sliding windows sliding_window = getattr(self.config, "sliding_window", None) if sliding_window is not None and kv_seq_len > sliding_window: - # From https://github.com/huggingface/transformers/blob/main/src/transformers/models/mistral/modeling_mistral.py#L193 - slicing_tokens = 1 - sliding_window - Knn = Kn[:, :, slicing_tokens:, :] # .contiguous() - Vnn = Vn[:, :, slicing_tokens:, :] # .contiguous() + start = kv_seq_len - sliding_window + Knn = Kn[:, :, start:, :] # .contiguous() + Vnn = Vn[:, :, start:, :] # .contiguous() + if attention_mask is not None: + attention_mask = attention_mask[..., start:] else: Knn, Vnn = Kn, Vn @@ -397,9 +407,6 @@ def CohereAttention_fast_forward_inference( ) Knn = Knn.reshape(bsz, n_heads, cached_len, head_dim) Vnn = Vnn.reshape(bsz, n_heads, cached_len, head_dim) - # else: - # Knn, Vnn = Knn, Vnn - # pass # Attention if bsz == 1: @@ -408,7 +415,6 @@ def CohereAttention_fast_forward_inference( A = torch_matmul( Qn, Knn.transpose(2, 3), out = self.attention[:, :, :, :cached_len] ) - # if attention_mask is not None: A += attention_mask # Must add attention_mask for batched A[:] = torch_nn_functional_softmax( A, dim = -1, dtype = torch.float32 ) # .to(A.dtype) @@ -453,6 +459,9 @@ def CohereModel_fast_forward_inference( seq_len, sliding_window = getattr(self.config, "sliding_window", None), ) + # Pre-convert to bool once for all layers (avoids per-layer .eq(0)) + if attention_mask is not None and attention_mask.dtype != torch.bool: + attention_mask = attention_mask.eq(0) else: attention_mask = None diff --git a/unsloth/models/falcon_h1.py b/unsloth/models/falcon_h1.py index 428f49d727..6e3b16b21b 100644 --- a/unsloth/models/falcon_h1.py +++ b/unsloth/models/falcon_h1.py @@ -197,6 +197,7 @@ def FalconH1Attention_fast_forward_inference( position_ids, do_prefill = False, attention_mask = None, + **kwargs, ): """ https://github.com/huggingface/transformers/blob/main/src/transformers/models/llama/modeling_llama.py#L406 @@ -265,7 +266,7 @@ def FalconH1Attention_fast_forward_inference( # Mistral Nemo 12b has weird dimensions if attention_size != hidden_size: - self.temp_O = torch.empty((1, bsz, hidden_size), dtype = dtype, device = device) + self.temp_O = torch.empty((bsz, 1, hidden_size), dtype = dtype, device = device) else: self.temp_O = self.temp_QA[1][:, :, :hidden_size] @@ -292,7 +293,7 @@ def FalconH1Attention_fast_forward_inference( Qn = fast_linear_forward(self.q_proj, Xn, out = self.temp_QA[0]) Kn = fast_linear_forward(self.k_proj, Xn, out = self.temp_KV[0]) - Kn = Kn * self.config.key_multiplier + Kn.mul_(self.config.key_multiplier) Vn = fast_linear_forward(self.v_proj, Xn, out = self.temp_KV[1]) Qn = Qn.view( bsz, 1, n_heads, head_dim @@ -343,10 +344,11 @@ def FalconH1Attention_fast_forward_inference( # Handle sliding windows sliding_window = getattr(self.config, "sliding_window", None) if sliding_window is not None and kv_seq_len > sliding_window: - # From https://github.com/huggingface/transformers/blob/main/src/transformers/models/mistral/modeling_mistral.py#L193 - slicing_tokens = 1 - sliding_window - Knn = Kn[:, :, slicing_tokens:, :] # .contiguous() - Vnn = Vn[:, :, slicing_tokens:, :] # .contiguous() + start = kv_seq_len - sliding_window + Knn = Kn[:, :, start:, :] # .contiguous() + Vnn = Vn[:, :, start:, :] # .contiguous() + if attention_mask is not None: + attention_mask = attention_mask[..., start:] else: Knn, Vnn = Kn, Vn @@ -361,9 +363,6 @@ def FalconH1Attention_fast_forward_inference( ) Knn = Knn.reshape(bsz, n_heads, cached_len, head_dim) Vnn = Vnn.reshape(bsz, n_heads, cached_len, head_dim) - # else: - # Knn, Vnn = Knn, Vnn - # pass # Attention if bsz == 1: @@ -372,7 +371,6 @@ def FalconH1Attention_fast_forward_inference( A = torch_matmul( Qn, Knn.transpose(2, 3), out = self.attention[:, :, :, :cached_len] ) - # if attention_mask is not None: A += attention_mask # Must add attention_mask for batched A[:] = torch_nn_functional_softmax( A, dim = -1, dtype = torch.float32 ) # .to(A.dtype) @@ -533,11 +531,19 @@ def _FalconH1_fast_forward_inference( bsz, q_len, hd = X.shape assert q_len == 1 # Get saved buffers to reduce memory movement - residual = torch.empty((bsz, q_len, hd), dtype = torch.float32, device = "cuda:0") - _XX = torch.empty((2, bsz, q_len, hd), dtype = torch.float32, device = "cuda:0") + residual = torch.empty( + (bsz, q_len, hd), dtype = torch.float32, device = f"{DEVICE_TYPE_TORCH}:0" + ) + _XX = torch.empty( + (2, bsz, q_len, hd), dtype = torch.float32, device = f"{DEVICE_TYPE_TORCH}:0" + ) XX, XX2 = _XX[0], _XX[1] - variance = torch.empty((bsz, q_len, 1), dtype = torch.float32, device = "cuda:0") - temp_mlp = torch.empty((2, bsz, 1, mlp_size), dtype = X.dtype, device = "cuda:0") + variance = torch.empty( + (bsz, q_len, 1), dtype = torch.float32, device = f"{DEVICE_TYPE_TORCH}:0" + ) + temp_mlp = torch.empty( + (2, bsz, 1, mlp_size), dtype = X.dtype, device = f"{DEVICE_TYPE_TORCH}:0" + ) temp_gate, temp_up = temp_mlp[0], temp_mlp[1] seq_len = past_key_values[0][0].shape[-2] if bsz != 1: diff --git a/unsloth/models/gemma.py b/unsloth/models/gemma.py index 55a8c8697f..cf543ae094 100644 --- a/unsloth/models/gemma.py +++ b/unsloth/models/gemma.py @@ -97,7 +97,9 @@ def GemmaDecoderLayer_fast_forward( self, "_flag_for_generation" ): # past_key_value is not None: out_weight = torch.empty( - self.input_layernorm.weight.shape, dtype = torch.float32, device = "cuda:0" + self.input_layernorm.weight.shape, + dtype = torch.float32, + device = f"{DEVICE_TYPE_TORCH}:0", ) # Self Attention @@ -191,6 +193,7 @@ def GemmaModel_fast_forward_inference( bsz, q_len, hd = hidden_states.shape seq_len = past_key_values[0][0].shape[-2] + kv_seq_len = seq_len + 1 if bsz != 1: attention_mask = _prepare_4d_causal_attention_mask_for_sdpa( attention_mask, @@ -198,6 +201,12 @@ def GemmaModel_fast_forward_inference( hidden_states, seq_len, ) + # Pre-convert to bool once for all layers (avoids per-layer .eq(0)) + if attention_mask is not None and attention_mask.dtype != torch.bool: + attention_mask = attention_mask.eq(0) + + # Compute rotary_seq_len once to avoid per-layer GPU-CPU sync from .item() + rotary_seq_len = max(kv_seq_len, int(position_ids.max().item()) + 1) next_decoder_cache = [] for idx, decoder_layer in enumerate(self.model.layers): @@ -217,6 +226,7 @@ def GemmaModel_fast_forward_inference( position_ids = position_ids, attention_mask = attention_mask, do_prefill = not hasattr(decoder_layer.self_attn, "paged_attention"), + rotary_seq_len = rotary_seq_len, ) hidden_states += residual diff --git a/unsloth/models/gemma2.py b/unsloth/models/gemma2.py index 03e77f6504..e59b8d5ebd 100644 --- a/unsloth/models/gemma2.py +++ b/unsloth/models/gemma2.py @@ -222,7 +222,9 @@ def Gemma2DecoderLayer_fast_forward( self, "_flag_for_generation" ): # past_key_value is not None: out_weight = torch.empty( - self.input_layernorm.weight.shape, dtype = torch.float32, device = "cuda:0" + self.input_layernorm.weight.shape, + dtype = torch.float32, + device = f"{DEVICE_TYPE_TORCH}:0", ) # Self Attention @@ -352,7 +354,7 @@ def Gemma2Attention_fast_forward_inference( ) self.RH_Q = torch.empty((bsz, n_heads, 1, head_dim), dtype = dtype, device = device) # Only for Gemma2 - self.temp_O = torch.empty((1, bsz, hidden_size), dtype = dtype, device = device) + self.temp_O = torch.empty((bsz, 1, hidden_size), dtype = dtype, device = device) self.attention = torch.empty( (bsz, n_heads, 1, KV_CACHE_INCREMENT + seq_len), dtype = dtype, device = device ) @@ -399,7 +401,7 @@ def Gemma2Attention_fast_forward_inference( RH_Q = self.RH_Q RH_Q[:, :, :, :h] = Qn[:, :, :, h:] RH_Q[:, :, :, h:] = Qn[:, :, :, :h] - torch.neg(RH_Q[:, :, :, :h], out = RH_Q[:, :, :, :h]) + RH_Q[:, :, :, :h].neg_() Qn *= cos Qn.addcmul_(RH_Q, sin) @@ -408,7 +410,7 @@ def Gemma2Attention_fast_forward_inference( ] # torch.empty((n_kv_heads, 1, head_dim), dtype = dtype, device = "cuda:0") RH_K[:, :, :, :h] = Kn[:, :, :, h:] RH_K[:, :, :, h:] = Kn[:, :, :, :h] - torch.neg(RH_K[:, :, :, :h], out = RH_K[:, :, :, :h]) + RH_K[:, :, :, :h].neg_() Kn *= cos Kn.addcmul_(RH_K, sin) @@ -423,10 +425,9 @@ def Gemma2Attention_fast_forward_inference( # Handle sliding windows sliding_window = self.config.sliding_window if use_sliding_window and kv_seq_len > sliding_window: - # From https://github.com/huggingface/transformers/blob/main/src/transformers/models/mistral/modeling_mistral.py#L193 - slicing_tokens = 1 - sliding_window - Knn = Kn[:, :, slicing_tokens:, :] # .contiguous() - Vnn = Vn[:, :, slicing_tokens:, :] # .contiguous() + start = kv_seq_len - sliding_window + Knn = Kn[:, :, start:, :] # .contiguous() + Vnn = Vn[:, :, start:, :] # .contiguous() else: Knn, Vnn = Kn, Vn @@ -441,28 +442,32 @@ def Gemma2Attention_fast_forward_inference( ) Knn = Knn.reshape(bsz, n_heads, cached_len, head_dim) Vnn = Vnn.reshape(bsz, n_heads, cached_len, head_dim) - # else: - # Knn, Vnn = Knn, Vnn - # pass # Attention - # if bsz == 1: + # [TODO] Gemma2 uses manual matmul for all batch sizes because SDPA does + # not support softcapping (tanh logit scaling). If a future PyTorch adds + # a softcap param to scaled_dot_product_attention, consider using SDPA + # for bsz > 1 to match the llama/qwen3 pattern. Qn *= ( self.scalar ) # See https://github.com/ggerganov/llama.cpp/issues/7805#issuecomment-2153349963 # It seems like doing (Q * scalar) @ K is better than (Q @ K) * scalar to stop overflows A = torch_matmul(Qn, Knn.transpose(2, 3), out = self.attention[:, :, :, :cached_len]) - # if attention_mask is not None: A += attention_mask # Must add attention_mask for batched + # Softcapping must happen BEFORE the mask is applied. + # Reference: google-deepmind/gemma _modules.py and transformers gemma2 eager_attention_forward A *= self.reciprocal_t - torch_tanh(A, out = A) + A.tanh_() A *= self.t # Logit softcapping + if attention_mask is not None and isinstance(attention_mask, torch.Tensor): + # Slice mask to match K/V when sliding window is active + if attention_mask.shape[-1] != A.shape[-1]: + attention_mask = attention_mask[:, :, :, -A.shape[-1] :] + A += attention_mask + A[:] = torch_nn_functional_softmax(A, dim = -1, dtype = torch.float32) # .to(A.dtype) A = torch_matmul(A, Vnn, out = Qn) - # else: - # A = scaled_dot_product_attention(Qn, Knn, Vnn, attn_mask = attention_mask, is_causal = False) - # pass A = A.transpose(1, 2) A = A.reshape(bsz, 1, attention_size) A = fast_linear_forward(self.o_proj, A, out = self.temp_O) diff --git a/unsloth/models/granite.py b/unsloth/models/granite.py index 168df90f4c..79ac41c43f 100644 --- a/unsloth/models/granite.py +++ b/unsloth/models/granite.py @@ -323,8 +323,7 @@ def GraniteAttention_fast_forward_inference( (2, bsz, 1, n_kv_heads * head_dim), dtype = dtype, device = device ) self.RH_Q = torch.empty((bsz, n_heads, 1, head_dim), dtype = dtype, device = device) - # Only for Gemma2 - self.temp_O = torch.empty((1, bsz, hidden_size), dtype = dtype, device = device) + self.temp_O = torch.empty((bsz, 1, hidden_size), dtype = dtype, device = device) self.attention = torch.empty( (bsz, n_heads, 1, KV_CACHE_INCREMENT + seq_len), dtype = dtype, device = device ) @@ -362,7 +361,7 @@ def GraniteAttention_fast_forward_inference( RH_Q = self.RH_Q RH_Q[:, :, :, :h] = Qn[:, :, :, h:] RH_Q[:, :, :, h:] = Qn[:, :, :, :h] - torch.neg(RH_Q[:, :, :, :h], out = RH_Q[:, :, :, :h]) + RH_Q[:, :, :, :h].neg_() Qn *= cos Qn.addcmul_(RH_Q, sin) @@ -371,7 +370,7 @@ def GraniteAttention_fast_forward_inference( ] # torch.empty((n_kv_heads, 1, head_dim), dtype = dtype, device = "cuda:0") RH_K[:, :, :, :h] = Kn[:, :, :, h:] RH_K[:, :, :, h:] = Kn[:, :, :, :h] - torch.neg(RH_K[:, :, :, :h], out = RH_K[:, :, :, :h]) + RH_K[:, :, :, :h].neg_() Kn *= cos Kn.addcmul_(RH_K, sin) @@ -385,7 +384,7 @@ def GraniteAttention_fast_forward_inference( # Grouped query attention _, _, cached_len, _ = Kn.shape - if n_groups != 1: + if bsz == 1 or ((not SDPA_HAS_GQA) and n_groups != 1): Kn = Kn[:, :, None, :, :].expand( bsz, n_kv_heads, n_groups, cached_len, head_dim ) @@ -394,20 +393,39 @@ def GraniteAttention_fast_forward_inference( ) Kn = Kn.reshape(bsz, n_heads, cached_len, head_dim) Vn = Vn.reshape(bsz, n_heads, cached_len, head_dim) - # else: - # Kn, Vn = Kn, Vn - # pass - Qn *= self.scaling - A = torch_matmul(Qn, Kn.transpose(2, 3), out = self.attention[:, :, :, :cached_len]) - - # if attention_mask is not None: A += attention_mask # Must add attention_mask for batched - - A[:] = torch_nn_functional_softmax(A, dim = -1, dtype = torch.float32) # .to(A.dtype) - A = torch_matmul(A, Vn, out = Qn) - # else: - # A = scaled_dot_product_attention(Qn, Kn, Vn, attn_mask = attention_mask, is_causal = False) - # pass + # Attention + if bsz == 1: + Qn *= self.scaling + A = torch_matmul( + Qn, Kn.transpose(2, 3), out = self.attention[:, :, :, :cached_len] + ) + A[:] = torch_nn_functional_softmax(A, dim = -1, dtype = torch.float32) + A = torch_matmul(A, Vn, out = Qn) + else: + if ( + attention_mask is not None + and attention_mask.dim() == 4 + and attention_mask.dtype != torch.bool + ): + attention_mask = attention_mask.eq(0) + if SDPA_HAS_GQA: + A = scaled_dot_product_attention( + Qn, + Kn, + Vn, + attn_mask = attention_mask, + scale = self.scaling, + enable_gqa = True, + ) + else: + A = scaled_dot_product_attention( + Qn, + Kn, + Vn, + attn_mask = attention_mask, + scale = self.scaling, + ) A = A.transpose(1, 2) A = A.reshape(bsz, 1, attention_size) A = fast_linear_forward(self.o_proj, A, out = self.temp_O) @@ -442,6 +460,9 @@ def GraniteModel_fast_forward_inference( hidden_states, seq_len, ) + # Pre-convert to bool once for all layers (avoids per-layer .eq(0)) + if attention_mask is not None and attention_mask.dtype != torch.bool: + attention_mask = attention_mask.eq(0) else: attention_mask = None diff --git a/unsloth/models/llama.py b/unsloth/models/llama.py index 342366f02e..f80a55fdd2 100644 --- a/unsloth/models/llama.py +++ b/unsloth/models/llama.py @@ -36,6 +36,7 @@ from ..utils.attention_dispatch import ( AttentionConfig, AttentionContext, run_attention, + SDPA, select_attention_backend, ) from torch.nn.functional import scaled_dot_product_attention @@ -213,11 +214,22 @@ def _fast_prepare_inputs_for_generation( **kwargs, ): past_key_values = kwargs.get("past_key_values", None) + original_attention_mask = attention_mask # Handle inputs_embeds - only use on FIRST generation step (no cache) # This fixes GitHub issue #3798: inputs_embeds was ignored use_inputs_embeds = inputs_embeds is not None and past_key_values is None + if input_ids is not None and input_ids.numel() > 0: + bs, seq_length = input_ids.shape + device = input_ids.device + elif inputs_embeds is not None: + bs, seq_length, _ = inputs_embeds.shape + device = inputs_embeds.device + else: + bs, seq_length = 1, 0 + device = "cuda" if torch.cuda.is_available() else "cpu" + if past_key_values is not None: # Check for uninitialized DynamicCache if len(past_key_values) == 0: @@ -234,16 +246,47 @@ def _fast_prepare_inputs_for_generation( use_inputs_embeds = inputs_embeds is not None else: if input_ids is not None and input_ids.numel() > 0: - bs, cache_length = input_ids.shape + bs = input_ids.shape[0] input_ids = input_ids[:, [-1]] device = input_ids.device + seq_length = 1 elif inputs_embeds is not None: - bs, cache_length, _ = inputs_embeds.shape + bs, seq_length, _ = inputs_embeds.shape device = inputs_embeds.device else: - bs, cache_length = 1, 0 + bs, seq_length = 1, 0 device = "cuda" if torch.cuda.is_available() else "cpu" + if hasattr(past_key_values, "get_seq_length"): + past_len = int(past_key_values.get_seq_length()) + else: + # legacy tuple cache: (layer, (K,V)) + past_len = int(past_key_values[0][0].shape[-2]) + + max_cache_len = None + if hasattr(past_key_values, "get_max_cache_shape"): + m = past_key_values.get_max_cache_shape() + max_cache_len = int(m) if m is not None and m > 0 else None + elif hasattr(past_key_values, "get_max_length"): + m = past_key_values.get_max_length() + max_cache_len = int(m) if m is not None else None + + # ensure cache_position + cache_position = kwargs.get("cache_position", None) + if cache_position is None: + kwargs["cache_position"] = torch.arange( + past_len, + past_len + seq_length, + device = device, + dtype = torch.long, + ) + else: + if ( + hasattr(cache_position, "device") + and cache_position.device != device + ): + kwargs["cache_position"] = cache_position.to(device) + # Get to the base model base_model = self if hasattr(base_model, "base_model_prefix"): @@ -252,45 +295,49 @@ def _fast_prepare_inputs_for_generation( if hasattr( base_model, "_prepare_4d_causal_attention_mask_with_cache_position" ): + if not hasattr(base_model, "_unsloth_mask_needs_device"): - def needs_device_kw(fn) -> bool: - try: - sig = inspect.signature(inspect.unwrap(fn)) - return "device" in sig.parameters - except: - # transformers <= 4.51.3 includes device arg but > 4.51.3 does not - return transformers_version < Version("4.52.0") + def _check_needs_device(fn) -> bool: + try: + sig = inspect.signature(inspect.unwrap(fn)) + return "device" in sig.parameters + except: + # transformers <= 4.51.3 includes device arg but > 4.51.3 does not + return transformers_version < Version("4.52.0") - kwargs = { - "sequence_length": 1, - "target_length": cache_length, + base_model._unsloth_mask_needs_device = _check_needs_device( + base_model._prepare_4d_causal_attention_mask_with_cache_position + ) + + if max_cache_len is not None: + target_length = max_cache_len + elif ( + original_attention_mask is not None + and original_attention_mask.dim() == 2 + ): + target_length = original_attention_mask.shape[-1] + else: + target_length = past_len + seq_length + + mask_kwargs = { + "sequence_length": seq_length, + "target_length": target_length, "dtype": self.dtype, - "cache_position": torch.arange( - cache_length, cache_length + 1, device = device - ), + "cache_position": kwargs["cache_position"], "batch_size": bs, "config": self.config, "past_key_values": past_key_values, } - try: - if needs_device_kw( - base_model._prepare_4d_causal_attention_mask_with_cache_position - ): - kwargs["device"] = device - except: - print( - f"Unsloth: Could not inspect signature of {base_model._prepare_4d_causal_attention_mask_with_cache_position}" - ) + if base_model._unsloth_mask_needs_device: + mask_kwargs["device"] = device attention_mask = ( base_model._prepare_4d_causal_attention_mask_with_cache_position( attention_mask, - **kwargs, + **mask_kwargs, ) ) else: - if attention_mask is not None: - attention_mask = attention_mask[:, [-1]] if transformers_version <= Version("4.52.4"): logger.warning_once( f"{self.__class__.__name__} has no `_prepare_4d_causal_attention_mask_with_cache_position` method " @@ -299,8 +346,17 @@ def _fast_prepare_inputs_for_generation( "issue on GitHub." ) - if "cache_position" in kwargs: - kwargs["position_ids"] = kwargs["cache_position"] + if kwargs.get("position_ids", None) is None: + if original_attention_mask is not None and original_attention_mask.dim() == 2: + position_ids = original_attention_mask.long().cumsum(-1) - 1 + position_ids.masked_fill_(original_attention_mask == 0, 1) + position_ids = position_ids[:, -seq_length:] + kwargs["position_ids"] = position_ids + elif kwargs.get("cache_position", None) is not None: + cp = kwargs["cache_position"] + if cp.dim() == 1: + cp = cp.unsqueeze(0).expand(bs, -1) + kwargs["position_ids"] = cp result = { "attention_mask": attention_mask, @@ -330,6 +386,7 @@ def LlamaAttention_fast_forward_inference( position_ids, do_prefill = False, attention_mask = None, + rotary_seq_len = None, ): """ https://github.com/huggingface/transformers/blob/main/src/transformers/models/llama/modeling_llama.py#L406 @@ -398,7 +455,7 @@ def LlamaAttention_fast_forward_inference( # Mistral Nemo 12b has weird dimensions if attention_size != hidden_size: - self.temp_O = torch.empty((1, bsz, hidden_size), dtype = dtype, device = device) + self.temp_O = torch.empty((bsz, 1, hidden_size), dtype = dtype, device = device) else: self.temp_O = self.temp_QA[1][:, :, :hidden_size] @@ -435,10 +492,19 @@ def LlamaAttention_fast_forward_inference( # Need to do it prior 2 steps before hitting full on short KV cache # or else error - self.rotary_emb.extend_rope_embedding(Vn, seq_len + 2) - cos, sin = self.rotary_emb.get_cached(kv_seq_len, Qn.device.index) - cos = cos[position_ids].unsqueeze(1) - sin = sin[position_ids].unsqueeze(1) + # ensure correct shape + if position_ids.dim() == 1: + position_ids = position_ids[:, None] + position_ids = position_ids.to(Qn.device) + + if rotary_seq_len is None: + rotary_seq_len = max(kv_seq_len, int(position_ids.max().item()) + 1) + self.rotary_emb.extend_rope_embedding(Vn, rotary_seq_len + 1) # +1 slack + cos, sin = self.rotary_emb.get_cached(rotary_seq_len, Qn.device.index or 0) + + cos = cos[position_ids].unsqueeze(1).to(device = Qn.device, dtype = Qn.dtype) + sin = sin[position_ids].unsqueeze(1).to(device = Qn.device, dtype = Qn.dtype) + h = self.half_head_dim RH_Q = self.RH_Q @@ -469,15 +535,17 @@ def LlamaAttention_fast_forward_inference( sliding_window = getattr(self.config, "sliding_window", None) if sliding_window is not None and kv_seq_len > sliding_window: # From https://github.com/huggingface/transformers/blob/main/src/transformers/models/mistral/modeling_mistral.py#L193 - slicing_tokens = 1 - sliding_window - Knn = Kn[:, :, slicing_tokens:, :] # .contiguous() - Vnn = Vn[:, :, slicing_tokens:, :] # .contiguous() + start = kv_seq_len - sliding_window + Knn = Kn[:, :, start:, :] # .contiguous() + Vnn = Vn[:, :, start:, :] # .contiguous() + if attention_mask is not None: + attention_mask = attention_mask[..., start:] else: Knn, Vnn = Kn, Vn # Grouped query attention _, _, cached_len, _ = Knn.shape - if bsz == 1 or not SDPA_HAS_GQA and n_groups != 1: + if bsz == 1 or ((not SDPA_HAS_GQA) and n_groups != 1): Knn = Knn[:, :, None, :, :].expand( bsz, n_kv_heads, n_groups, cached_len, head_dim ) @@ -486,9 +554,6 @@ def LlamaAttention_fast_forward_inference( ) Knn = Knn.reshape(bsz, n_heads, cached_len, head_dim) Vnn = Vnn.reshape(bsz, n_heads, cached_len, head_dim) - # else: - # Knn, Vnn = Knn, Vnn - # pass # when qlen==vlen and attn_mask is None, we should use causal attention Q_len = Qn.shape[-2] @@ -504,12 +569,23 @@ def LlamaAttention_fast_forward_inference( A = torch_matmul( Qn, Knn.transpose(2, 3), out = self.attention[:, :, :, :cached_len] ) - # if attention_mask is not None: A += attention_mask # Must add attention_mask for batched A[:] = torch_nn_functional_softmax( A, dim = -1, dtype = torch.float32 ) # .to(A.dtype) A = torch_matmul(A, Vnn, out = Qn) + # --- attention_mask fixup for SDPA if user passes 2D padding mask else: + if attention_mask is not None and attention_mask.dim() == 2: + attention_mask = attention_mask[:, None, None, :].to(torch.bool) + # is it more appropriate to use _prepare_4d_causal_attention_mask_for_sdpa? + elif ( + attention_mask is not None + and attention_mask.dim() == 4 + and attention_mask.dtype != torch.bool + ): + # Decode is more stable with boolean keep masks than additive bf16 masks. + attention_mask = attention_mask.eq(0) + if SDPA_HAS_GQA: A = scaled_dot_product_attention( Qn, @@ -663,6 +739,8 @@ def LlamaAttention_fast_forward( rotary_emb = self.rotary_emb rotary_emb.extend_rope_embedding(V, seq_len = kv_seq_len) cos, sin = rotary_emb.get_cached(kv_seq_len, Q.device.index) + cos = cos.to(device = Q.device, dtype = Q.dtype) + sin = sin.to(device = Q.device, dtype = Q.dtype) rope_position_ids = position_ids if rope_position_ids is None and seq_info is not None: @@ -682,7 +760,11 @@ def LlamaAttention_fast_forward( # Attention module use_varlen = seq_info is not None and past_key_value is None - backend = select_attention_backend(use_varlen) + backend = ( + SDPA if attention_mask is not None else select_attention_backend(use_varlen) + ) + + # should dropout be hardcoded to 0.0? config = AttentionConfig( backend = backend, n_kv_heads = n_kv_heads, @@ -1257,7 +1339,8 @@ def _LlamaModel_fast_forward_inference( ) seq_len = past_key_values[0][0].shape[-2] - if bsz != 1: + kv_seq_len = seq_len + 1 + if attention_mask is not None: attention_mask = _prepare_4d_causal_attention_mask_for_sdpa( attention_mask, (bsz, q_len), @@ -1265,9 +1348,15 @@ def _LlamaModel_fast_forward_inference( seq_len, sliding_window = getattr(self.config, "sliding_window", None), ) + # Pre-convert to bool once for all layers (avoids per-layer .eq(0)) + if attention_mask is not None and attention_mask.dtype != torch.bool: + attention_mask = attention_mask.eq(0) else: attention_mask = None + # Compute rotary_seq_len once to avoid per-layer GPU-CPU sync from .item() + rotary_seq_len = max(kv_seq_len, int(position_ids.max().item()) + 1) + next_decoder_cache = [] for idx, decoder_layer in enumerate(self.model.layers): @@ -1290,6 +1379,7 @@ def _LlamaModel_fast_forward_inference( position_ids = position_ids, attention_mask = attention_mask, do_prefill = not hasattr(decoder_layer.self_attn, "paged_attention"), + rotary_seq_len = rotary_seq_len, ) X += residual @@ -1528,7 +1618,7 @@ def CausalLM_fast_forward(fast_forward_inference): logits = logit_softcapping * logits else: logits *= 1.0 / logit_softcapping - torch.tanh(logits, out = logits) + logits.tanh_() logits *= logit_softcapping if not return_dict: @@ -2586,16 +2676,39 @@ class FastLlamaModel: model._old_generate = model.generate unsloth_fast_generate.__doc__ = model._old_generate.__doc__ model.generate = types.MethodType(unsloth_fast_generate, model) - # Set weight[padding_idx] = 0 - with torch.no_grad(): - for name, module in model.named_modules(): - if type(module) is torch.nn.Embedding: - if ( - getattr(module, "weight", None) is not None - and getattr(module, "padding_idx", None) is not None - ): - if module.padding_idx < module.weight.shape[0]: - module.weight[module.padding_idx] = 0 + # Set weight[padding_idx] = 0 for embeddings that are NOT tied with the + # lm_head. When weights are tied, zeroing the padding row also zeros + # the corresponding lm_head row, forcing logit = 0 for the pad token. + # This is higher than the (negative) logits for real tokens in models + # like Gemma, causing the decoder to emit and produce gibberish. + # Skip entirely if eos_token == pad_token to avoid zeroing EOS embedding. + eos_token_id = ( + getattr(tokenizer, "eos_token_id", None) if tokenizer is not None else None + ) + pad_token_id = ( + getattr(tokenizer, "pad_token_id", None) if tokenizer is not None else None + ) + if tokenizer is not None and eos_token_id != pad_token_id: + lm_head = getattr(model, "lm_head", None) + lm_head_weight = ( + getattr(lm_head, "weight", None) if lm_head is not None else None + ) + with torch.no_grad(): + for name, module in model.named_modules(): + if type(module) is torch.nn.Embedding: + if ( + getattr(module, "weight", None) is not None + and getattr(module, "padding_idx", None) is not None + ): + if module.padding_idx < module.weight.shape[0]: + # Skip if tied to lm_head + if ( + lm_head_weight is not None + and module.weight.data_ptr() + == lm_head_weight.data_ptr() + ): + continue + module.weight[module.padding_idx] = 0 return model, tokenizer @staticmethod diff --git a/unsloth/models/mistral.py b/unsloth/models/mistral.py index 5e893d2b6f..83e9ab9486 100644 --- a/unsloth/models/mistral.py +++ b/unsloth/models/mistral.py @@ -25,6 +25,7 @@ from ..utils.attention_dispatch import ( AttentionConfig, AttentionContext, run_attention, + SDPA, select_attention_backend, ) from .llama import ( @@ -115,7 +116,9 @@ def MistralAttention_fast_forward( use_varlen = ( seq_info is not None and past_key_value is None and window_size == (-1, -1) ) - backend = select_attention_backend(use_varlen) + backend = ( + SDPA if attention_mask is not None else select_attention_backend(use_varlen) + ) attention_config = AttentionConfig( backend = backend, n_kv_heads = n_kv_heads, @@ -216,13 +219,18 @@ def MistralForCausalLM_fast_forward( bsz, 1, q_len, q_len ) else: - # attention_mask should be [bsz, 1, q_len, q_len] or broadcastable - # Add causal mask to existing attention mask if attention_mask.dim() == 2: - # [bsz, seq_len] -> [bsz, 1, 1, seq_len] - attention_mask = attention_mask[:, None, None, :] - attention_mask = attention_mask.expand(bsz, 1, q_len, q_len) - attention_mask = attention_mask + causal_mask_values[None, None, :, :] + # Convert 0/1 padding mask to additive format: 1->0 (keep), 0->-inf (mask) + padding_mask = torch.where( + attention_mask[:, None, None, :].bool(), + 0.0, + -torch.inf, + ) + attention_mask = causal_mask_values[None, None, :, :] + padding_mask + else: + attention_mask = ( + attention_mask + causal_mask_values[None, None, :, :] + ) attention_mask = attention_mask.to( dtype = _get_dtype(dtype_from_config(self.config)) diff --git a/unsloth/models/qwen3.py b/unsloth/models/qwen3.py index ea06016d72..b93dddb186 100644 --- a/unsloth/models/qwen3.py +++ b/unsloth/models/qwen3.py @@ -21,6 +21,7 @@ from ..utils.attention_dispatch import ( AttentionConfig, AttentionContext, run_attention, + SDPA, select_attention_backend, ) from .llama import ( @@ -139,7 +140,9 @@ def Qwen3Attention_fast_forward( # Attention module use_varlen = seq_info is not None and past_key_value is None - backend = select_attention_backend(use_varlen) + backend = ( + SDPA if attention_mask is not None else select_attention_backend(use_varlen) + ) attention_config = AttentionConfig( backend = backend, n_kv_heads = n_kv_heads, @@ -181,6 +184,7 @@ def Qwen3Attention_fast_forward_inference( position_ids, do_prefill = False, attention_mask = None, + **kwargs, ): """ https://github.com/huggingface/transformers/blob/main/src/transformers/models/llama/modeling_llama.py#L406 @@ -249,7 +253,7 @@ def Qwen3Attention_fast_forward_inference( # Mistral Nemo 12b has weird dimensions if attention_size != hidden_size: - self.temp_O = torch.empty((1, bsz, hidden_size), dtype = dtype, device = device) + self.temp_O = torch.empty((bsz, 1, hidden_size), dtype = dtype, device = device) else: self.temp_O = self.temp_QA[1][:, :, :hidden_size] @@ -329,24 +333,42 @@ def Qwen3Attention_fast_forward_inference( # Handle sliding windows sliding_window = getattr(self.config, "sliding_window", None) if sliding_window is not None and kv_seq_len > sliding_window: - # From https://github.com/huggingface/transformers/blob/main/src/transformers/models/mistral/modeling_mistral.py#L193 - slicing_tokens = 1 - sliding_window - Knn = Kn[:, :, slicing_tokens:, :] # .contiguous() - Vnn = Vn[:, :, slicing_tokens:, :] # .contiguous() + start = kv_seq_len - sliding_window + Knn = Kn[:, :, start:, :] # .contiguous() + Vnn = Vn[:, :, start:, :] # .contiguous() + if attention_mask is not None: + attention_mask = attention_mask[..., start:] else: Knn, Vnn = Kn, Vn # when qlen==vlen and attn_mask is None, we should use causal attention Q_len = Qn.shape[-2] K_len = Knn.shape[-2] + if attention_mask is not None and attention_mask.dim() == 2: + attention_mask = attention_mask[:, None, None, :].to(torch.bool) + elif ( + attention_mask is not None + and attention_mask.dim() == 4 + and attention_mask.dtype != torch.bool + ): + attention_mask = attention_mask.eq(0) if attention_mask is None and Q_len == K_len: is_causal = True else: is_causal = False + use_sdpa_gqa = SDPA_HAS_GQA + if ( + use_sdpa_gqa + and isinstance(attention_mask, torch.Tensor) + and attention_mask.dim() >= 3 + and attention_mask.shape[0] > 1 + ): + # Avoid SDPA GQA drift for batched masked decode. + use_sdpa_gqa = False # Grouped query attention _, _, cached_len, _ = Knn.shape - if bsz == 1 or not SDPA_HAS_GQA and n_groups != 1: + if bsz == 1 or ((not use_sdpa_gqa) and n_groups != 1): Knn = Knn[:, :, None, :, :].expand( bsz, n_kv_heads, n_groups, cached_len, head_dim ) @@ -355,9 +377,6 @@ def Qwen3Attention_fast_forward_inference( ) Knn = Knn.reshape(bsz, n_heads, cached_len, head_dim) Vnn = Vnn.reshape(bsz, n_heads, cached_len, head_dim) - # else: - # Knn, Vnn = Knn, Vnn - # pass # Attention if bsz == 1: @@ -366,13 +385,12 @@ def Qwen3Attention_fast_forward_inference( A = torch_matmul( Qn, Knn.transpose(2, 3), out = self.attention[:, :, :, :cached_len] ) - # if attention_mask is not None: A += attention_mask # Must add attention_mask for batched A[:] = torch_nn_functional_softmax( A, dim = -1, dtype = torch.float32 ) # .to(A.dtype) A = torch_matmul(A, Vnn, out = Qn) else: - if SDPA_HAS_GQA: + if use_sdpa_gqa: A = scaled_dot_product_attention( Qn, Knn, diff --git a/unsloth/models/vision.py b/unsloth/models/vision.py index 24e72e1535..e6c859d44e 100644 --- a/unsloth/models/vision.py +++ b/unsloth/models/vision.py @@ -1410,9 +1410,15 @@ class FastBaseModel: m.for_training = functools.partial(FastBaseModel.for_training, m) m.for_inference = functools.partial(FastBaseModel.for_inference, m) m = m.model - # Set weight[padding_idx] = 0 + # Set weight[padding_idx] = 0 for embeddings that are NOT tied with the + # lm_head. When weights are tied, zeroing the padding row also zeros + # the corresponding lm_head row, forcing logit = 0 for the pad token. # Only do this if tokenizer is defined since eos_token == pad_token sometimes! pad_token_id = getattr(tokenizer, "pad_token_id", None) + lm_head = getattr(model, "lm_head", None) + lm_head_weight = ( + getattr(lm_head, "weight", None) if lm_head is not None else None + ) if ( tokenizer is not None and getattr(tokenizer, "eos_token_id", None) != pad_token_id @@ -1428,6 +1434,13 @@ class FastBaseModel: module.padding_idx == pad_token_id and module.padding_idx < module.weight.shape[0] ): + # Skip if tied to lm_head + if ( + lm_head_weight is not None + and module.weight.data_ptr() + == lm_head_weight.data_ptr() + ): + continue module.weight[module.padding_idx] = 0 return model diff --git a/unsloth/utils/attention_dispatch.py b/unsloth/utils/attention_dispatch.py index a7620549be..72d52ab376 100644 --- a/unsloth/utils/attention_dispatch.py +++ b/unsloth/utils/attention_dispatch.py @@ -20,6 +20,7 @@ from __future__ import annotations from dataclasses import dataclass from typing import Any, Optional, Tuple +import torch from torch import Tensor from torch.nn.functional import scaled_dot_product_attention @@ -119,6 +120,19 @@ def run_attention( backend = config.backend if backend == FLASH_VARLEN and context.seq_info is None: backend = FLASH_DENSE if HAS_FLASH_ATTENTION else SDPA + + # [TODO] Flash attention does not support arbitrary attention masks (only + # causal via flag). When a padding mask is present (e.g. left-padded + # batched generation), fall back to SDPA which consumes attn_mask. + # xFormers also does not thread context.attention_mask through, so the + # same fallback applies. + if context.attention_mask is not None and backend in ( + FLASH_DENSE, + FLASH_VARLEN, + XFORMERS, + ): + backend = SDPA + flash_dense_kwargs = config.flash_dense_kwargs or {} flash_varlen_kwargs = config.flash_varlen_kwargs or {} sdpa_kwargs = config.sdpa_kwargs or {} @@ -234,14 +248,79 @@ def run_attention( else: q_len_local = Q.shape[-2] k_len_local = K.shape[-2] + # ---- SDPA mask normalization for left padding / 2D masks ---- + if local_mask is not None and isinstance(local_mask, torch.Tensor): + local_mask = local_mask.to(device = Q.device) + + if local_mask.dim() == 2: + # key padding keep mask: (bsz, k_len), 1/True = real token + if local_mask.dtype == torch.bool: + key_keep = local_mask + else: + # tokenizer attention_mask is typically int 0/1 + key_keep = local_mask != 0 + + past_len = ( + k_len_local - q_len_local + ) # works for prefill (0) and decode + q_pos = torch.arange( + past_len, past_len + q_len_local, device = Q.device + ) + k_pos = torch.arange(k_len_local, device = Q.device) + + causal_keep = ( + k_pos[None, :] <= q_pos[:, None] + ) # True = allowed (SDPA) + if sliding_window is not None: + causal_keep &= k_pos[None, :] >= ( + q_pos[:, None] - (sliding_window - 1) + ) + + # (bsz, 1, q_len, k_len) boolean keep mask + local_mask = ( + causal_keep[None, None, :, :] & key_keep[:, None, None, :] + ) + + elif local_mask.dim() == 3: + # (bsz, q_len, k_len) -> (bsz, 1, q_len, k_len) + local_mask = local_mask[:, None, :, :] + + elif local_mask.dim() == 4: + if local_mask.dtype != torch.bool: + # Use boolean keep masks for better SDPA stability. + local_mask = local_mask.eq(0) + else: + raise ValueError( + f"Unsupported SDPA attention_mask rank: {local_mask.dim()}" + ) + + # Avoid NaNs from fully-masked rows (common with left padding). + if local_mask.dtype == torch.bool: + no_allowed = ~local_mask.any( + dim = -1, keepdim = True + ) # (bsz,1,q_len,1) + local_mask = local_mask | no_allowed + is_causal_local = local_mask is None and q_len_local == k_len_local kwargs = dict(sdpa_kwargs) kwargs.setdefault("attn_mask", local_mask) kwargs.setdefault("is_causal", is_causal_local) - if SDPA_HAS_GQA: - kwargs.setdefault("enable_gqa", config.n_groups != 1) + use_sdpa_gqa = SDPA_HAS_GQA and config.n_groups != 1 + if ( + use_sdpa_gqa + and (not requires_grad) + and isinstance(local_mask, torch.Tensor) + and local_mask.dim() >= 3 + and local_mask.shape[0] > 1 + ): + # Batched masked inference has shown row-coupled drift with SDPA GQA. + # Fall back to explicit KV expansion for deterministic row-wise behavior. + use_sdpa_gqa = False + + if use_sdpa_gqa: + kwargs.setdefault("enable_gqa", True) out = scaled_dot_product_attention(Q, K, V, **kwargs) return out.transpose(1, 2) From 0e64336ac9a09f72f3b2b1525e85f13ded5bfccc Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Feb 2026 08:18:45 -0800 Subject: [PATCH 38/63] Fix transformers v5 RoPE inv_freq corruption and generate() BatchEncoding compat (#4112) * Fix transformers v5 RoPE inv_freq corruption during model loading Transformers v5 initializes models on the meta device, then _move_missing_keys_from_meta_to_device() replaces all non-persistent buffers with torch.empty_like() (uninitialized memory). Vanilla transformers restores inv_freq via _init_weights() checking for original_inv_freq, but Unsloth's LlamaRotaryEmbedding subclasses lack this attribute, so inv_freq stays corrupted with garbage values. This caused 5-11x higher training loss on transformers v5 for all models using Unsloth's rope (Llama 3.x, Qwen3, Mistral, TinyLlama, Granite). Models using native transformers rope (Gemma, Phi-4, Falcon-H1) were unaffected. The fix recomputes inv_freq from the stored base/dim after model loading, applies model-specific scaling via _apply_inv_freq_scaling(), and rebuilds cos/sin caches. Also handles LongRopeRotaryEmbedding (Phi-3.5 style short/long inv_freq). Guarded by transformers >= 5.0.0 so it is a no-op on v4. Tested on: Llama 3.1 8B, Llama 3.2 3B, Qwen3 14B, Qwen3 4B, Phi-4, TinyLlama, Mistral 7B, Gemma2 2B, Falcon-H1 -- all v5 losses now match v4 baselines to < 0.004 absolute difference. * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Unpack BatchEncoding in generate() for v4/v5 backwards compatibility Old notebooks pass the full tokenizer output as input_ids: inputs = tokenizer(..., return_tensors="pt").to("cuda") model.generate(input_ids=inputs, ...) This worked on transformers v4 because generate() internally extracted the tensor. Transformers v5 calls .shape on input_ids directly, which crashes since BatchEncoding has no .shape attribute. Fix: in unsloth_fast_generate(), detect when input_ids is a dict-like object (BatchEncoding) and unpack its contents into separate kwargs before forwarding to the underlying generate(). This makes both old and new notebook patterns work on both v4 and v5. * Remove redundant seen_ids dedup in _fix_rope_inv_freq named_modules() already deduplicates with remove_duplicate=True (default). Also clarify that native v5 rotary classes (Gemma3 etc.) have original_inv_freq which transformers v5's _init_weights() uses to restore inv_freq, so they do not need this fix. --------- Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- unsloth/models/llama.py | 18 ++++++-- unsloth/models/loader.py | 99 ++++++++++++++++++++++++++++++++++++++++ 2 files changed, 114 insertions(+), 3 deletions(-) diff --git a/unsloth/models/llama.py b/unsloth/models/llama.py index f80a55fdd2..6fe60cf940 100644 --- a/unsloth/models/llama.py +++ b/unsloth/models/llama.py @@ -2062,6 +2062,21 @@ def unsloth_fast_generate( FastLlamaModel.for_inference(self) + # Unpack BatchEncoding passed as input_ids for backwards compatibility. + # Old notebooks do model.generate(input_ids=tokenizer(...)) where the tokenizer + # output is a BatchEncoding (dict-like). Transformers v5 generate() calls + # .shape on it directly and crashes. Unpack into separate kwargs so both + # v4 and v5 work transparently. + _maybe_encoding = kwargs.get("input_ids", None) + if ( + _maybe_encoding is not None + and not isinstance(_maybe_encoding, torch.Tensor) + and hasattr(_maybe_encoding, "items") + ): + batch_data = kwargs.pop("input_ids") + for key, val in batch_data.items(): + kwargs.setdefault(key, val) + dtype = _get_dtype(dtype_from_config(self.config)) if hasattr(self, "config") and hasattr(self.config, "max_position_embeddings"): @@ -2071,9 +2086,6 @@ def unsloth_fast_generate( and "max_new_tokens" in kwargs ): _ids = kwargs["input_ids"] - # Handle BatchEncoding from transformers 5.0+ (no .shape attribute) - if hasattr(_ids, "input_ids"): - _ids = _ids["input_ids"] if hasattr(_ids, "shape") and ( _ids.shape[-1] + kwargs["max_new_tokens"] > self.config.max_position_embeddings diff --git a/unsloth/models/loader.py b/unsloth/models/loader.py index 72a6d7d10f..711476b759 100644 --- a/unsloth/models/loader.py +++ b/unsloth/models/loader.py @@ -78,6 +78,9 @@ SUPPORTS_QWEN3_MOE = transformers_version >= Version("4.50.3") SUPPORTS_FALCON_H1 = transformers_version >= Version("4.53.0") SUPPORTS_GEMMA3N = transformers_version >= Version("4.53.0") SUPPORTS_GPTOSS = transformers_version >= Version("4.55.0") +# Transformers v5 meta-device loading corrupts non-persistent buffers (inv_freq). +# See _fix_rope_inv_freq() below for details. +_NEEDS_ROPE_FIX = transformers_version >= Version("5.0.0") if SUPPORTS_GEMMA: from .gemma import FastGemmaModel if SUPPORTS_GEMMA2: @@ -121,6 +124,100 @@ DISABLE_SDPA_MODEL_NAMES = [ ] +def _fix_rope_inv_freq(model): + """Fix inv_freq corruption caused by transformers v5 meta-device loading. + + Transformers v5 initializes models on the meta device, then + _move_missing_keys_from_meta_to_device() (modeling_utils.py) replaces ALL + non-persistent buffers with torch.empty_like() -- uninitialized memory. + + Vanilla transformers restores inv_freq via _init_weights() which checks for + hasattr(module, "original_inv_freq"). Unsloth's LlamaRotaryEmbedding and + subclasses do not have this attribute, so inv_freq stays corrupted. This + produces wrong positional encodings and causes 5-11x higher training loss. + + This function recomputes inv_freq from the stored base and dim, applies + any model-specific scaling, and rebuilds the cos/sin caches. + + Only runs on transformers >= 5.0.0. No-op on v4. + """ + if not _NEEDS_ROPE_FIX: + return model + + for name, module in model.named_modules(): + # Unsloth's LlamaRotaryEmbedding and subclasses (Extended, LinearScaling, + # Granite). Native v5 rotary classes (Gemma3, etc.) have original_inv_freq + # which v5's _init_weights() uses to restore inv_freq, so they are fine. + if ( + hasattr(module, "inv_freq") + and hasattr(module, "base") + and hasattr(module, "dim") + and hasattr(module, "_apply_inv_freq_scaling") + and hasattr(module, "multi_gpu_cos_cached") + ): + inv_freq = 1.0 / ( + module.base + ** ( + torch.arange( + 0, module.dim, 2, dtype = torch.int64, device = "cpu" + ).float() + / module.dim + ) + ) + inv_freq = module._apply_inv_freq_scaling(inv_freq) + module.inv_freq = inv_freq + for device_idx in range(len(module.multi_gpu_cos_cached)): + if module.multi_gpu_cos_cached[device_idx] is not None: + module._set_cos_sin_cache( + seq_len = module.current_rope_size, + device = torch.device(device_idx), + dtype = torch.get_default_dtype(), + ) + + # LongRopeRotaryEmbedding (Phi-3.5 style with short_inv_freq + long_inv_freq) + elif ( + hasattr(module, "short_inv_freq") + and hasattr(module, "long_inv_freq") + and hasattr(module, "base") + and hasattr(module, "dim") + ): + config = getattr(model, "config", None) + rope_scaling = getattr(config, "rope_scaling", None) if config else None + if rope_scaling is not None: + short_factor = rope_scaling.get("short_factor", None) + long_factor = rope_scaling.get("long_factor", None) + if short_factor is not None and long_factor is not None: + inv_freq_shape = ( + torch.arange( + 0, module.dim, 2, dtype = torch.int64, device = "cpu" + ).float() + / module.dim + ) + sf = torch.tensor(short_factor, device = "cpu", dtype = torch.float32) + lf = torch.tensor(long_factor, device = "cpu", dtype = torch.float32) + module.short_inv_freq = 1.0 / (sf * module.base**inv_freq_shape) + module.long_inv_freq = 1.0 / (lf * module.base**inv_freq_shape) + + dtype = torch.bfloat16 if is_bfloat16_supported() else torch.float16 + t = torch.arange( + module.original_max_position_embeddings, + device = module.short_inv_freq.device, + dtype = torch.int64, + ).float() + freqs = torch.outer(t, module.short_inv_freq) + emb = torch.cat((freqs, freqs), dim = -1) + for device_idx in range(len(module.multi_gpu_short_cos_cached)): + if module.multi_gpu_short_cos_cached[device_idx] is not None: + device_obj = torch.device(device_idx) + module.multi_gpu_short_cos_cached[device_idx] = ( + emb.cos() * module.scaling_factor + ).to(dtype = dtype, device = device_obj, non_blocking = True) + module.multi_gpu_short_sin_cached[device_idx] = ( + emb.sin() * module.scaling_factor + ).to(dtype = dtype, device = device_obj, non_blocking = True) + return model + + class FastLanguageModel(FastLlamaModel): @staticmethod def from_pretrained( @@ -685,6 +782,7 @@ class FastLanguageModel(FastLlamaModel): if patch_tiled_mlp_choice != "0" or unsloth_tiled_mlp: patch_tiled_mlp(model, patch_options_str = patch_tiled_mlp_choice) + model = _fix_rope_inv_freq(model) return model, tokenizer @@ -1408,6 +1506,7 @@ class FastModel(FastBaseModel): if patch_tiled_mlp_choice != "0" or unsloth_tiled_mlp: patch_tiled_mlp(model, patch_options_str = patch_tiled_mlp_choice) + model = _fix_rope_inv_freq(model) return model, tokenizer From 885b4899446f158b092b62ee3b9bac586258d976 Mon Sep 17 00:00:00 2001 From: Michael Han <107991372+shimmyshimmer@users.noreply.github.com> Date: Wed, 25 Feb 2026 23:56:48 -0800 Subject: [PATCH 39/63] Qwen3.5 update.md --- README.md | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/README.md b/README.md index 4a9d4d01c8..b7fc6b6a88 100644 --- a/README.md +++ b/README.md @@ -54,15 +54,15 @@ For RTX 50x, B200, 6000 GPUs: `pip install unsloth`. Read our guides for: [Black To install Unsloth on **AMD** and **Intel** GPUs, follow our [AMD Guide](https://unsloth.ai/docs/get-started/install/amd) and [Intel Guide](https://unsloth.ai/docs/get-started/install/intel). ## 🦥 Unsloth News +- **Qwen3.5** including 35-A3B, 27B, 112B-A10B is now supported. [Guide + notebooks](https://unsloth.ai/docs/new/faster-moe) - Train **MoE LLMs 12x faster** with 35% less VRAM - DeepSeek, GLM, Qwen and gpt-oss. [Blog](https://unsloth.ai/docs/new/faster-moe) - **Embedding models**: Unsloth now supports ~1.8-3.3x faster embedding fine-tuning. [Blog](https://unsloth.ai/docs/new/embedding-finetuning) • [Notebooks](https://unsloth.ai/docs/get-started/unsloth-notebooks#embedding-models) - New **7x longer context RL** vs. all other setups, via our new batching algorithms. [Blog](https://unsloth.ai/docs/new/grpo-long-context) - New RoPE & MLP **Triton Kernels** & **Padding Free + Packing**: 3x faster training & 30% less VRAM. [Blog](https://unsloth.ai/docs/new/3x-faster-training-packing) - **500K Context**: Training a 20B model with >500K context is now possible on an 80GB GPU. [Blog](https://unsloth.ai/docs/blog/500k-context-length-fine-tuning) -- **FP8 Reinforcement Learning**: You can now do FP8 GRPO on consumer GPUs. [Blog](https://unsloth.ai/docs/get-started/reinforcement-learning-rl-guide/fp8-reinforcement-learning) • [Notebook](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Qwen3_8B_FP8_GRPO.ipynb) +- **FP8 & Vision RL**: You can now do FP8 & VLM GRPO on consumer GPUs. [FP8 Blog](https://unsloth.ai/docs/get-started/reinforcement-learning-rl-guide/fp8-reinforcement-learning) • [Vision RL](https://unsloth.ai/docs/get-started/reinforcement-learning-rl-guide/vision-reinforcement-learning-vlm-rl) - **Docker**: Use Unsloth with no setup & environment issues with our new image. [Guide](https://unsloth.ai/docs/blog/how-to-fine-tune-llms-with-unsloth-and-docker) • [Docker image](https://hub.docker.com/r/unsloth/unsloth) -- **Vision RL**: You can now train VLMs with GRPO or GSPO in Unsloth! [Read guide](https://unsloth.ai/docs/get-started/reinforcement-learning-rl-guide/vision-reinforcement-learning-vlm-rl) -- **gpt-oss** by OpenAI: Read our [RL blog](https://unsloth.ai/docs/models/gpt-oss-how-to-run-and-fine-tune/gpt-oss-reinforcement-learning), [Flex Attention](https://unsloth.ai/docs/models/gpt-oss-how-to-run-and-fine-tune/long-context-gpt-oss-training) blog and [gpt-oss Guide](https://unsloth.ai/docs/models/gpt-oss-how-to-run-and-fine-tune). 20B works on 14GB VRAM. 120B on 65GB. +- **gpt-oss** by OpenAI: Read our [RL blog](https://unsloth.ai/docs/models/gpt-oss-how-to-run-and-fine-tune/gpt-oss-reinforcement-learning), [Flex Attention](https://unsloth.ai/docs/models/gpt-oss-how-to-run-and-fine-tune/long-context-gpt-oss-training) blog and [Guide](https://unsloth.ai/docs/models/gpt-oss-how-to-run-and-fine-tune).
Click for more news From 8b0f1731e0c34fb12a36b36695838b9418b1094e Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Thu, 26 Feb 2026 08:06:21 -0800 Subject: [PATCH 40/63] Update README.md (#4118) --- README.md | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/README.md b/README.md index b7fc6b6a88..b0b6a2f52a 100644 --- a/README.md +++ b/README.md @@ -6,7 +6,7 @@ unsloth logo - + @@ -22,8 +22,6 @@ Notebooks are beginner friendly. Read our [guide](https://unsloth.ai/docs/get-st | Model | Free Notebooks | Performance | Memory use | |-----------|---------|--------|----------| -| **gpt-oss (20B)** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/gpt-oss-(20B)-Fine-tuning.ipynb) | 1.5x faster | 70% less | -| **gpt-oss (20B): GRPO** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/gpt-oss-(20B)-GRPO.ipynb) | 2x faster | 80% less | | **Qwen3: Advanced GRPO** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Qwen3_(4B)-GRPO.ipynb) | 2x faster | 50% less | | **Qwen3-VL (8B): GSPO** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Qwen3_VL_(8B)-Vision-GRPO.ipynb) | 1.5x faster | 80% less | | **Gemma 3 (4B) Vision** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Gemma3_(4B)-Vision.ipynb) | 1.7x faster | 60% less | @@ -33,6 +31,8 @@ Notebooks are beginner friendly. Read our [guide](https://unsloth.ai/docs/get-st | **Llama 3.1 (8B) Alpaca** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Llama3.1_(8B)-Alpaca.ipynb) | 2x faster | 70% less | | **Llama 3.2 Conversational** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Llama3.2_(1B_and_3B)-Conversational.ipynb) | 2x faster | 70% less | | **Orpheus-TTS (3B)** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Orpheus_(3B)-TTS.ipynb) | 1.5x faster | 50% less | +| **gpt-oss (20B)** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/gpt-oss-(20B)-Fine-tuning.ipynb) | 1.5x faster | 70% less | +| **gpt-oss (20B): GRPO** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/gpt-oss-(20B)-GRPO.ipynb) | 2x faster | 80% less | - See all our notebooks for: [Kaggle](https://github.com/unslothai/notebooks?tab=readme-ov-file#-kaggle-notebooks), [GRPO](https://unsloth.ai/docs/get-started/unsloth-notebooks#grpo-reasoning-rl-notebooks), [TTS](https://unsloth.ai/docs/get-started/unsloth-notebooks#text-to-speech-tts-notebooks), [embedding](https://unsloth.ai/docs/new/embedding-finetuning) & [Vision](https://unsloth.ai/docs/get-started/unsloth-notebooks#vision-multimodal-notebooks) - See [all our models](https://unsloth.ai/docs/get-started/unsloth-model-catalog) and [all our notebooks](https://unsloth.ai/docs/get-started/unsloth-notebooks) From 4870fc5fb4767424667df25c925556c439cbd653 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Thu, 26 Feb 2026 09:18:29 -0800 Subject: [PATCH 41/63] Update README.md (#4119) --- README.md | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/README.md b/README.md index b0b6a2f52a..b7fc6b6a88 100644 --- a/README.md +++ b/README.md @@ -6,7 +6,7 @@ unsloth logo - + @@ -22,6 +22,8 @@ Notebooks are beginner friendly. Read our [guide](https://unsloth.ai/docs/get-st | Model | Free Notebooks | Performance | Memory use | |-----------|---------|--------|----------| +| **gpt-oss (20B)** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/gpt-oss-(20B)-Fine-tuning.ipynb) | 1.5x faster | 70% less | +| **gpt-oss (20B): GRPO** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/gpt-oss-(20B)-GRPO.ipynb) | 2x faster | 80% less | | **Qwen3: Advanced GRPO** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Qwen3_(4B)-GRPO.ipynb) | 2x faster | 50% less | | **Qwen3-VL (8B): GSPO** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Qwen3_VL_(8B)-Vision-GRPO.ipynb) | 1.5x faster | 80% less | | **Gemma 3 (4B) Vision** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Gemma3_(4B)-Vision.ipynb) | 1.7x faster | 60% less | @@ -31,8 +33,6 @@ Notebooks are beginner friendly. Read our [guide](https://unsloth.ai/docs/get-st | **Llama 3.1 (8B) Alpaca** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Llama3.1_(8B)-Alpaca.ipynb) | 2x faster | 70% less | | **Llama 3.2 Conversational** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Llama3.2_(1B_and_3B)-Conversational.ipynb) | 2x faster | 70% less | | **Orpheus-TTS (3B)** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Orpheus_(3B)-TTS.ipynb) | 1.5x faster | 50% less | -| **gpt-oss (20B)** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/gpt-oss-(20B)-Fine-tuning.ipynb) | 1.5x faster | 70% less | -| **gpt-oss (20B): GRPO** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/gpt-oss-(20B)-GRPO.ipynb) | 2x faster | 80% less | - See all our notebooks for: [Kaggle](https://github.com/unslothai/notebooks?tab=readme-ov-file#-kaggle-notebooks), [GRPO](https://unsloth.ai/docs/get-started/unsloth-notebooks#grpo-reasoning-rl-notebooks), [TTS](https://unsloth.ai/docs/get-started/unsloth-notebooks#text-to-speech-tts-notebooks), [embedding](https://unsloth.ai/docs/new/embedding-finetuning) & [Vision](https://unsloth.ai/docs/get-started/unsloth-notebooks#vision-multimodal-notebooks) - See [all our models](https://unsloth.ai/docs/get-started/unsloth-model-catalog) and [all our notebooks](https://unsloth.ai/docs/get-started/unsloth-notebooks) From d650c15983e6812d5ef2062eb1d318fbee4caa72 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Thu, 26 Feb 2026 17:48:38 -0800 Subject: [PATCH 42/63] Guard Gemma3N variants from flex attention defaults (#4116) --- unsloth/models/_utils.py | 8 +++++--- unsloth/models/vision.py | 8 +++++++- 2 files changed, 12 insertions(+), 4 deletions(-) diff --git a/unsloth/models/_utils.py b/unsloth/models/_utils.py index dafdec4e1e..f883d466f0 100644 --- a/unsloth/models/_utils.py +++ b/unsloth/models/_utils.py @@ -235,14 +235,16 @@ def prefer_flex_attn_if_supported(model_class, config): model_class, "_supports_flex_attn", False ): return None - # GPT-OSS and Mllama use eager/sdpa attention during inference since - # flex attention returns incorrect results or errors out. + # GPT-OSS, Mllama and Gemma3N use eager/sdpa attention during + # inference since flex attention returns incorrect results or errors out. # GPT-OSS: left padding issues cause incorrect outputs. # Mllama: _update_causal_mask uses make_flex_block_causal_mask which # creates BlockMask with Q_LEN=KV_LEN=total_seq_len, but during # decode q_len=1, causing ValueError. Needs transformers update. + # Gemma3N: timm vision wrappers (eg Gemma3nVisionConfig) do not + # support flex_attention. model_type = getattr(config, "model_type", "") if config else "" - if model_type in ("gpt_oss", "mllama"): + if model_type in ("gpt_oss", "mllama") or str(model_type).startswith("gemma3n"): return None if config is not None: setattr(config, "_attn_implementation", "flex_attention") diff --git a/unsloth/models/vision.py b/unsloth/models/vision.py index e6c859d44e..1f6b240a7d 100644 --- a/unsloth/models/vision.py +++ b/unsloth/models/vision.py @@ -630,7 +630,13 @@ class FastBaseModel: except KeyError: pass - default_attn_impl = "flex_attention" if flex_attn_impl else "sdpa" + model_type = str(getattr(auto_config, "model_type", "")).lower() + if model_type.startswith("gemma3n"): + # Gemma3N variants initialize timm-based vision towers which do + # not support flex_attention, so default to eager unless overridden. + default_attn_impl = "eager" + else: + default_attn_impl = "flex_attention" if flex_attn_impl else "sdpa" if not ("attn_implementation" in kwargs): kwargs["attn_implementation"] = default_attn_impl if not supports_sdpa and kwargs.get("attn_implementation") == "sdpa": From 9dca5d8e76580ab1b12b8122c6a1d6e72d62ec3d Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Thu, 26 Feb 2026 23:43:57 -0800 Subject: [PATCH 43/63] Fix Whisper auto_model mapping fallback for concrete model classes (#4115) --- unsloth/models/vision.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/unsloth/models/vision.py b/unsloth/models/vision.py index 1f6b240a7d..5aea471263 100644 --- a/unsloth/models/vision.py +++ b/unsloth/models/vision.py @@ -627,8 +627,8 @@ class FastBaseModel: # Re-resolve model_class after potential config change try: model_class = auto_model._model_mapping[auto_config.__class__] - except KeyError: - pass + except Exception: + model_class = None model_type = str(getattr(auto_config, "model_type", "")).lower() if model_type.startswith("gemma3n"): From cd37a10a2a8d86772edbc863c2e64bab927765f9 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E9=87=91=E9=BB=84=E8=89=B2=E8=91=A1=E8=90=84=E7=90=83?= =?UTF-8?q?=E5=90=9B=E5=90=9B?= Date: Sun, 1 Mar 2026 15:59:12 +0800 Subject: [PATCH 44/63] =?UTF-8?q?fix(ROCm):=20remove=20fix=5Frocm=5Ftriton?= =?UTF-8?q?=5Fkey=5Ferror=20=E2=80=94=20based=20on=20a=20false=20premise?= =?UTF-8?q?=20(#4125)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The function (introduced in #3923) assumed that the absence of `triton.runtime.triton_key` on ROCm means torch.compile will crash. Investigation shows this is incorrect: 1. `triton.runtime.triton_key` was renamed/removed in the ROCm Triton fork — it does not exist at that path. However, `triton.compiler.compiler.triton_key` (the path torch._inductor actually imports) EXISTS and works correctly on ROCm. 2. Both call-sites in torch._inductor (codecache.py and async_compile.py) already wrap the import in try/except, so even a genuinely missing triton_key would be handled gracefully. 3. Comprehensive testing on ROCm 7.1 + Triton 3.4.0 + gfx1100 confirms torch.compile works correctly for matmul, cross-entropy, RMSNorm, multi-layer transformer forward+backward, and LoRA — all without triton.runtime.triton_key. The original code was also ineffective (environment variables set after torch import have no effect on torch._dynamo config), so removing it has zero behavioral change on existing installations. Supersedes the compile-disable portion of #3923. --- unsloth/__init__.py | 3 --- unsloth/import_fixes.py | 33 --------------------------------- 2 files changed, 36 deletions(-) diff --git a/unsloth/__init__.py b/unsloth/__init__.py index 466fbf0628..60c5b3bba6 100644 --- a/unsloth/__init__.py +++ b/unsloth/__init__.py @@ -139,7 +139,6 @@ from .import_fixes import ( fix_vllm_guided_decoding_params, fix_vllm_pdl_blackwell, fix_triton_compiled_kernel_missing_attrs, - fix_rocm_triton_key_error, patch_trunc_normal_precision_issue, ignore_logger_messages, patch_ipykernel_hf_xet, @@ -161,7 +160,6 @@ check_vllm_torch_sm100_compatibility() fix_vllm_guided_decoding_params() fix_vllm_pdl_blackwell() fix_triton_compiled_kernel_missing_attrs() -fix_rocm_triton_key_error() patch_trunc_normal_precision_issue() ignore_logger_messages() patch_ipykernel_hf_xet() @@ -181,7 +179,6 @@ del check_vllm_torch_sm100_compatibility del fix_vllm_guided_decoding_params del fix_vllm_pdl_blackwell del fix_triton_compiled_kernel_missing_attrs -del fix_rocm_triton_key_error del patch_trunc_normal_precision_issue del ignore_logger_messages del patch_ipykernel_hf_xet diff --git a/unsloth/import_fixes.py b/unsloth/import_fixes.py index ee82715306..ebd81f9568 100644 --- a/unsloth/import_fixes.py +++ b/unsloth/import_fixes.py @@ -941,39 +941,6 @@ def fix_triton_compiled_kernel_missing_attrs(): ) -def fix_rocm_triton_key_error(): - """ - ROCm + torch.compile can fail if Triton lacks `triton_key`. - Disable Inductor/compile only on ROCm when that symbol is missing. - """ - try: - import torch - except (ImportError, ModuleNotFoundError): - return - - if not getattr(torch.version, "hip", None): - return - - try: - import triton - except (ImportError, ModuleNotFoundError): - return - - try: - from triton.runtime import triton_key # noqa: F401 - - return - except ImportError: - pass - - os.environ.setdefault("TORCHINDUCTOR_DISABLE", "1") - os.environ.setdefault("TORCH_COMPILE_DISABLE", "1") - logger.info( - "Unsloth: ROCm detected and Triton lacks triton_key; " - "disabling torch.compile/Inductor to avoid backend crash." - ) - - def patch_trunc_normal_precision_issue(): """ Patch torch.nn.init.trunc_normal_ for low precision tensors to run init in fp32. From dc75d00d1465e3c3251877c9a02ad66cb00bbf83 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E9=87=91=E9=BB=84=E8=89=B2=E8=91=A1=E8=90=84=E7=90=83?= =?UTF-8?q?=E5=90=9B=E5=90=9B?= Date: Sun, 1 Mar 2026 15:59:17 +0800 Subject: [PATCH 45/63] fix(ROCm): Comprehensive RDNA GPU support - fix Gemma3 NaN & add is_rdna() (#4109) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * fix(ROCm): comprehensive RDNA GPU support - fix Gemma3 NaN & add is_rdna() - Add is_rdna() detection for RDNA3/3.5/RDNA4 consumer GPUs (gfx11xx, gfx1151, gfx12xx) - Disable torch.compile for Gemma3 on HIP to fix NaN loss (fixes #3385, #4029) - Export is_cdna/is_rdna from kernels for downstream use - Import is_rdna into cross_entropy_loss for future RDNA-specific tuning Tested on AMD Radeon PRO W7900 (gfx1100) with ROCm 7.1: ✓ Gemma3-1B: loss 3.37→3.25 (no NaN) ✓ Llama-3.2-1B: loss 2.44→2.37 (no NaN) ✓ Qwen2.5-1.5B: loss 1.89→1.85 (no NaN) ✓ RMS LayerNorm Triton kernel: bf16/fp16 PASSED ✓ Cross Entropy Loss Triton kernel: 32K/256K vocab PASSED * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Address review: scope compile disable to RDNA only, use partial mode, remove unused import Changes based on Daniel's review: 1. (HIGH) Replace DEVICE_TYPE=='hip' with is_rdna() to avoid disabling torch.compile on CDNA GPUs (MI250X/MI300X/MI350) where it works fine 2. (MEDIUM) Use 'partial' instead of '1' for UNSLOTH_COMPILE_DISABLE to only disable model forward compilation while keeping loss compilation, matching the existing Sesame pattern 3. (LOW) Remove unused is_rdna import from cross_entropy_loss.py (F401) * Remove redundant is_cdna/is_rdna exports from kernels/__init__.py These functions are imported directly from .utils where needed (e.g. cross_entropy_loss.py, loader.py). No external code imports them from the unsloth.kernels namespace. * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --------- Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- unsloth/kernels/utils.py | 9 +++++++++ unsloth/models/loader.py | 8 ++++++++ 2 files changed, 17 insertions(+) diff --git a/unsloth/kernels/utils.py b/unsloth/kernels/utils.py index 0b13c04532..b1211ae345 100644 --- a/unsloth/kernels/utils.py +++ b/unsloth/kernels/utils.py @@ -86,6 +86,15 @@ def is_cdna(): ) +@functools.lru_cache(1) +def is_rdna(): + """Detect RDNA consumer/workstation GPUs (RDNA3, RDNA3.5, RDNA4).""" + if not is_hip(): + return False + arch = triton.runtime.driver.active.get_current_target().arch + return arch.startswith("gfx1") and not is_cdna() + + def calculate_settings( n: int, ) -> ( diff --git a/unsloth/models/loader.py b/unsloth/models/loader.py index 711476b759..7bb2a24e00 100644 --- a/unsloth/models/loader.py +++ b/unsloth/models/loader.py @@ -1131,6 +1131,14 @@ class FastModel(FastBaseModel): # Set norms to float32 since anyways they get upcasted to float32 # common in both gemma-3 and gemma-3n os.environ["UNSLOTH_HIGH_PRECISION_LAYERNORM"] = "1" + # ROCm/HIP: Gemma3 compiled forward produces NaN on RDNA GPUs + # (gfx1100, gfx1101, gfx1102, gfx1150, gfx1151, etc.). + # Disable torch.compile for model forward; loss compilation is fine. + # See https://github.com/unslothai/unsloth/issues/3385 + from unsloth.kernels.utils import is_rdna + + if is_rdna(): + os.environ["UNSLOTH_COMPILE_DISABLE"] = "partial" # Cohere elif "cohere2" in model_types_all and transformers_version < Version( "4.50.0.dev0" From e4daae62d99365fce520754be9e7fb514fc5d9fc Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E9=87=91=E9=BB=84=E8=89=B2=E8=91=A1=E8=90=84=E7=90=83?= =?UTF-8?q?=E5=90=9B=E5=90=9B?= Date: Sun, 1 Mar 2026 15:59:22 +0800 Subject: [PATCH 46/63] fix(Triton): ensure float32 eps in RMS LayerNorm rsqrt for HIP/ROCm (#4110) * fix(Triton): ensure float32 eps in RMS LayerNorm rsqrt for HIP/ROCm On HIP (AMD ROCm), Triton constexpr eps may not promote to float32 in rsqrt, causing numerical instability (NaN/Inf) on RDNA GPUs (gfx1100, gfx1151 Strix Halo, etc.). Use tl.full((), eps, tl.float32) to explicitly create a float32 scalar before adding to row_var in rsqrt. Applied to both standard and Gemma RMS LayerNorm forward kernels. Tested on W7900 (gfx1100): full test suite passed (dim 512-2048, bf16/fp16, various seqlen). Related: #3385, #3588 * Apply same float32 eps fix to layernorm.py for PR #4110 layernorm.py has the identical tl.constexpr eps pattern in layernorm_forward that can misfire on HIP/ROCm. Apply the same tl.full((), eps, tl.float32) fix for consistency. Both testing_suite_layernorm (standard LayerNorm) and testing_suite_layernorm (RMS LayerNorm) pass on NVIDIA after this change. --------- Co-authored-by: Daniel Han --- unsloth/kernels/layernorm.py | 4 +++- unsloth/kernels/rms_layernorm.py | 8 ++++++-- 2 files changed, 9 insertions(+), 3 deletions(-) diff --git a/unsloth/kernels/layernorm.py b/unsloth/kernels/layernorm.py index 5e2e3af2f8..9e64c3d341 100644 --- a/unsloth/kernels/layernorm.py +++ b/unsloth/kernels/layernorm.py @@ -55,7 +55,9 @@ def layernorm_forward( # (X[0] - mean) == -mean so we need to mask it out XX = tl.where(mask, X_row - mean_X, 0) row_var = tl.sum(XX * XX, axis = 0) / n_cols - inv_var = tl.math.rsqrt(row_var + eps) + # Explicit float32 scalar to ensure correct type promotion on HIP/ROCm + eps_f32 = tl.full((), eps, tl.float32) + inv_var = tl.math.rsqrt(row_var + eps_f32) tl.store(r, inv_var) tl.store(mu, mean_X) output = (XX * inv_var) * W_row + b_row diff --git a/unsloth/kernels/rms_layernorm.py b/unsloth/kernels/rms_layernorm.py index 82e0cd0e9b..74c16c1e63 100644 --- a/unsloth/kernels/rms_layernorm.py +++ b/unsloth/kernels/rms_layernorm.py @@ -49,7 +49,9 @@ def _rms_layernorm_forward( W_row = tl.load(W + col_offsets, mask = mask, other = 0) # .to(tl.float32) row_var = tl.sum(X_row * X_row, axis = 0) / n_cols - inv_var = tl.math.rsqrt(row_var + eps) + # Explicit float32 scalar to ensure correct type promotion on HIP/ROCm + eps_f32 = tl.full((), eps, tl.float32) + inv_var = tl.math.rsqrt(row_var + eps_f32) tl.store(r, inv_var) normed = X_row * inv_var normed = normed.to(W_row.dtype) # Exact copy from HF @@ -147,7 +149,9 @@ def _gemma_rms_layernorm_forward( W_row = tl.load(W + col_offsets, mask = mask, other = 0).to(tl.float32) row_var = tl.sum(X_row * X_row, axis = 0) / n_cols - inv_var = tl.math.rsqrt(row_var + eps) + # Explicit float32 scalar to ensure correct type promotion on HIP/ROCm + eps_f32 = tl.full((), eps, tl.float32) + inv_var = tl.math.rsqrt(row_var + eps_f32) tl.store(r, inv_var) normed = X_row * inv_var output = normed * (W_row + 1.0) From b8b88edc7f13ab6e818d60568f67af3c25bd448d Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E9=87=91=E9=BB=84=E8=89=B2=E8=91=A1=E8=90=84=E7=90=83?= =?UTF-8?q?=E5=90=9B=E5=90=9B?= Date: Sun, 1 Mar 2026 15:59:27 +0800 Subject: [PATCH 47/63] fix(ROCm): prevent false TMA support detection on AMD GPUs (#4126) TMA (Tensor Memory Accelerator) is an NVIDIA Hopper+ feature that does not exist on AMD GPUs. However, _check_tma_support() incorrectly returns True on ROCm because: 1. torch.cuda.get_device_capability() returns (11, 0) for gfx1100, satisfying the >= 9 check intended for Hopper (sm_90). 2. ROCm Triton exports tl.make_tensor_descriptor (the symbol exists even though the hardware does not support TMA). This would cause MoE grouped_gemm to attempt TMA operations on AMD GPUs, leading to runtime failures. Fix: early-return False for HIP devices, matching the existing XPU guard. --- unsloth/kernels/moe/grouped_gemm/interface.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/unsloth/kernels/moe/grouped_gemm/interface.py b/unsloth/kernels/moe/grouped_gemm/interface.py index 554e5fcc03..5588458973 100644 --- a/unsloth/kernels/moe/grouped_gemm/interface.py +++ b/unsloth/kernels/moe/grouped_gemm/interface.py @@ -39,10 +39,10 @@ logger.addHandler(ch) # Precompute TMA support to avoid graph breaks # TMA requires both: -# 1. GPU capability >= 9 (Hopper+) +# 1. NVIDIA GPU with capability >= 9 (Hopper+) # 2. Triton version with TMA API (make_tensor_descriptor or _experimental_make_tensor_descriptor) def _check_tma_support(): - if DEVICE_TYPE == "xpu": + if DEVICE_TYPE in ("xpu", "hip"): return False import triton.language as tl From 6e499c421b0a0af1caf991334bb47d1d3dc5c471 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E9=87=91=E9=BB=84=E8=89=B2=E8=91=A1=E8=90=84=E7=90=83?= =?UTF-8?q?=E5=90=9B=E5=90=9B?= Date: Sun, 1 Mar 2026 15:59:34 +0800 Subject: [PATCH 48/63] perf(ROCm): optimize chunked CE loss num_warps for RDNA GPUs (#4123) Use 16 warps for RDNA in the chunked cross-entropy forward kernel (large vocab > 65536), matching the existing CDNA optimization. Benchmarked on W7900 (gfx1100) with actual unsloth kernels (5 trials, median): - Chunked CE forward (BS=65536): 16 warps = 2.4-2.6x faster than 32 - All other kernels (LayerNorm, RoPE, SwiGLU): default heuristic is already optimal for RDNA; no modification needed. Depends on: #4109 (provides is_rdna() detection) --- unsloth/kernels/cross_entropy_loss.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/unsloth/kernels/cross_entropy_loss.py b/unsloth/kernels/cross_entropy_loss.py index d92229314f..1bb1691776 100644 --- a/unsloth/kernels/cross_entropy_loss.py +++ b/unsloth/kernels/cross_entropy_loss.py @@ -22,6 +22,7 @@ from .utils import ( triton_cast, torch_gpu_device, is_cdna, + is_rdna, ) from transformers.models.llama.modeling_llama import logger from unsloth_zoo.utils import Version @@ -364,7 +365,7 @@ class Fast_CrossEntropyLoss(torch.autograd.Function): SOFTCAP = logit_softcapping, DO_LOGIT_SCALING = DO_LOGIT_SCALING, LOGIT_SCALE = logit_scaling, - num_warps = 32 if not is_cdna() else 16, + num_warps = 16 if is_cdna() or is_rdna() else 32, ) # logsumexp(chunked_logsumexp) - x # Do the -x separately From 3ddb683d6be83a17cd2ba8ce99b6798a743e703e Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E9=87=91=E9=BB=84=E8=89=B2=E8=91=A1=E8=90=84=E7=90=83?= =?UTF-8?q?=E5=90=9B=E5=90=9B?= Date: Sun, 1 Mar 2026 16:15:47 +0800 Subject: [PATCH 49/63] Fix global dequantize buffer dtype mismatch across mixed-precision loads (#4026) Fix global dequantize buffer dtype mismatch when loading multiple 4-bit models with different dtypes in the same process. Adds dtype check alongside existing None check for WEIGHT_BUFFER in both CUDA/HIP and XPU paths. --- unsloth/kernels/utils.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/unsloth/kernels/utils.py b/unsloth/kernels/utils.py index b1211ae345..3122282b38 100644 --- a/unsloth/kernels/utils.py +++ b/unsloth/kernels/utils.py @@ -398,7 +398,7 @@ if DEVICE_TYPE == "xpu" and HAS_XPU_STREAM: global ABSMAX_BUFFERS WEIGHT_BUFFER = WEIGHT_BUFFERS[device_index] ABSMAX_BUFFER = ABSMAX_BUFFERS[device_index] - if WEIGHT_BUFFER is None: + if WEIGHT_BUFFER is None or WEIGHT_BUFFER.dtype != dtype: WEIGHT_BUFFERS[device_index] = WEIGHT_BUFFER = torch_empty( size, dtype = dtype, device = device, requires_grad = False ) @@ -508,7 +508,7 @@ elif DEVICE_TYPE in ("cuda", "hip") and HAS_CUDA_STREAM: global ABSMAX_BUFFERS WEIGHT_BUFFER = WEIGHT_BUFFERS[device_index] ABSMAX_BUFFER = ABSMAX_BUFFERS[device_index] - if WEIGHT_BUFFER is None: + if WEIGHT_BUFFER is None or WEIGHT_BUFFER.dtype != dtype: WEIGHT_BUFFERS[device_index] = WEIGHT_BUFFER = torch_empty( size, dtype = dtype, device = device, requires_grad = False ) From fc63b7ec4d2d7ba50dcf641503482f641779f78d Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sun, 1 Mar 2026 00:42:37 -0800 Subject: [PATCH 50/63] rl: guard warnings_issued before TRL estimate_tokens write (#4034) Co-authored-by: Daniel Hanchen --- unsloth/models/rl.py | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/unsloth/models/rl.py b/unsloth/models/rl.py index fbb32b2c50..eb5d4be171 100755 --- a/unsloth/models/rl.py +++ b/unsloth/models/rl.py @@ -810,6 +810,18 @@ def _patch_trl_rl_trainers(trainer_file = "grpo_trainer"): " os.environ['UNSLOTH_RETURN_LOGITS'] = '1'\n" ) extra_args += logits_check + warnings_issued_check = ( + "if model is not None:\n" + " _warnings_issued = getattr(model, 'warnings_issued', None)\n" + " if _warnings_issued is None:\n" + " model.warnings_issued = {}\n" + " elif not isinstance(_warnings_issued, dict):\n" + " try:\n" + " model.warnings_issued = dict(_warnings_issued)\n" + " except Exception:\n" + " model.warnings_issued = {}\n" + ) + extra_args += warnings_issued_check # Check max_seq_length if "model" in call_args: From 6f73648eadc2c8ddbcd7db44ef0d79277a847411 Mon Sep 17 00:00:00 2001 From: DoubleMathew Date: Sun, 1 Mar 2026 21:30:47 -0600 Subject: [PATCH 51/63] Fix auto padding free logic to respect user passed False (#4128) * Fix auto padding free logic to respect user passed * Update unsloth/trainer.py Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> --------- Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> --- unsloth/models/rl.py | 1 + unsloth/trainer.py | 3 ++- 2 files changed, 3 insertions(+), 1 deletion(-) diff --git a/unsloth/models/rl.py b/unsloth/models/rl.py index eb5d4be171..e4f34c908e 100755 --- a/unsloth/models/rl.py +++ b/unsloth/models/rl.py @@ -1037,6 +1037,7 @@ def _patch_trl_rl_trainers(trainer_file = "grpo_trainer"): "include_num_input_tokens_seen": False, "auto_find_batch_size": False, # Auto /2 batch size - too many people complained so removing "dataloader_pin_memory": True, + "padding_free": None, # None = user didn't set it, allows auto-enable detection # Might fail so disable for now # "dataloader_persistent_workers" : True, # Keeps dataloader in RAM # "dataloader_prefetch_factor" : 2, diff --git a/unsloth/trainer.py b/unsloth/trainer.py index cb36b8639d..65abe6801f 100644 --- a/unsloth/trainer.py +++ b/unsloth/trainer.py @@ -73,7 +73,7 @@ def _should_auto_padding_free(config) -> bool: or getattr(config, "packing", False) ): return False - return not getattr(config, "padding_free", False) + return getattr(config, "padding_free", None) is None def _disable_sample_packing(config): @@ -392,6 +392,7 @@ def _patch_sft_trainer_auto_packing(trl_module): packing_active = True logger.info("Unsloth: Sample packing enabled for SFTTrainer instance.") + # Resolve padding_free: None (default) = auto-enable unless env-disabled or packing auto_padding_free_active = False padding_free_requested = getattr(config_arg, "padding_free", None) is True if not blocked: From fa0ff304ade5cda76123496866d68d33397eb00b Mon Sep 17 00:00:00 2001 From: Etherll <61019402+Etherll@users.noreply.github.com> Date: Mon, 2 Mar 2026 21:36:28 +0200 Subject: [PATCH 52/63] Add Qwen 3.5 to FORCE_FLOAT32 (#4134) * Add Qwen3.5 to FORCE_FLOAT32 * fix vision encoder dtype mismatch * revert vision cast changes --- unsloth/models/loader.py | 1 + unsloth/models/vision.py | 1 + 2 files changed, 2 insertions(+) diff --git a/unsloth/models/loader.py b/unsloth/models/loader.py index 7bb2a24e00..b32c8303d8 100644 --- a/unsloth/models/loader.py +++ b/unsloth/models/loader.py @@ -106,6 +106,7 @@ FORCE_FLOAT32 = [ "gemma3text", # Gemma3TextModel (EmbeddingGemma, standalone text-only Gemma3) "gemma3n", "gpt_oss", + "qwen3_5", # Qwen3.5 RMSNorm uses (1+w) pattern like Gemma3, overflows float16 ] global DISABLE_COMPILE_MODEL_NAMES diff --git a/unsloth/models/vision.py b/unsloth/models/vision.py index 5aea471263..28c624c213 100644 --- a/unsloth/models/vision.py +++ b/unsloth/models/vision.py @@ -1054,6 +1054,7 @@ class FastBaseModel: do_forced_float32 = do_forced_float32, correct_dtype = correct_dtype, ) + try: model, tokenizer = patch_tokenizer(model, tokenizer) except Exception as _patch_err: From 55ba7c56402431f1f8bb06aea6720eb1d47709d6 Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Mon, 2 Mar 2026 21:48:36 -0800 Subject: [PATCH 53/63] [pre-commit.ci] pre-commit autoupdate (#4141) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit updates: - [github.com/astral-sh/ruff-pre-commit: v0.15.2 → v0.15.4](https://github.com/astral-sh/ruff-pre-commit/compare/v0.15.2...v0.15.4) Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- .pre-commit-config.yaml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index 245c636fe7..5c822d9b38 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -1,6 +1,6 @@ repos: - repo: https://github.com/astral-sh/ruff-pre-commit - rev: v0.15.2 + rev: v0.15.4 hooks: - id: ruff args: From b422b7336f4ebf720d74bec1d6d4c297a6c4d9fb Mon Sep 17 00:00:00 2001 From: Michael Han <107991372+shimmyshimmer@users.noreply.github.com> Date: Mon, 2 Mar 2026 23:33:22 -0800 Subject: [PATCH 54/63] Qwen3.5 Update.md Updated with Qwen3.5 Small models --- README.md | 43 +++++++++++++++++++++++++------------------ 1 file changed, 25 insertions(+), 18 deletions(-) diff --git a/README.md b/README.md index b7fc6b6a88..1314cb1c59 100644 --- a/README.md +++ b/README.md @@ -22,12 +22,11 @@ Notebooks are beginner friendly. Read our [guide](https://unsloth.ai/docs/get-st | Model | Free Notebooks | Performance | Memory use | |-----------|---------|--------|----------| -| **gpt-oss (20B)** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/gpt-oss-(20B)-Fine-tuning.ipynb) | 1.5x faster | 70% less | +| **Qwen3.5 (4B)** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Qwen3_5_(4B)_Vision.ipynb) | 1.5x faster | 60% less | +| **gpt-oss (20B)** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/gpt-oss-(20B)-Fine-tuning.ipynb) | 2x faster | 70% less | | **gpt-oss (20B): GRPO** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/gpt-oss-(20B)-GRPO.ipynb) | 2x faster | 80% less | | **Qwen3: Advanced GRPO** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Qwen3_(4B)-GRPO.ipynb) | 2x faster | 50% less | -| **Qwen3-VL (8B): GSPO** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Qwen3_VL_(8B)-Vision-GRPO.ipynb) | 1.5x faster | 80% less | | **Gemma 3 (4B) Vision** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Gemma3_(4B)-Vision.ipynb) | 1.7x faster | 60% less | -| **Gemma 3n (e4B)** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Gemma3N_(4B)-Conversational.ipynb) | 1.5x faster | 50% less | | **embeddinggemma (300M)** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/EmbeddingGemma_(300M).ipynb) | 2x faster | 20% less | | **Mistral Ministral 3 (3B)** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Ministral_3_VL_(3B)_Vision.ipynb) | 1.5x faster | 60% less | | **Llama 3.1 (8B) Alpaca** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Llama3.1_(8B)-Alpaca.ipynb) | 2x faster | 70% less | @@ -54,7 +53,7 @@ For RTX 50x, B200, 6000 GPUs: `pip install unsloth`. Read our guides for: [Black To install Unsloth on **AMD** and **Intel** GPUs, follow our [AMD Guide](https://unsloth.ai/docs/get-started/install/amd) and [Intel Guide](https://unsloth.ai/docs/get-started/install/intel). ## 🦥 Unsloth News -- **Qwen3.5** including 35-A3B, 27B, 112B-A10B is now supported. [Guide + notebooks](https://unsloth.ai/docs/new/faster-moe) +- **Qwen3.5** - 0.8B, 2B, 4B, 9B, 27B, 35-A3B, 112B-A10B are now supported. [Guide + notebooks](https://unsloth.ai/docs/models/qwen3.5/fine-tune) - Train **MoE LLMs 12x faster** with 35% less VRAM - DeepSeek, GLM, Qwen and gpt-oss. [Blog](https://unsloth.ai/docs/new/faster-moe) - **Embedding models**: Unsloth now supports ~1.8-3.3x faster embedding fine-tuning. [Blog](https://unsloth.ai/docs/new/embedding-finetuning) • [Notebooks](https://unsloth.ai/docs/get-started/unsloth-notebooks#embedding-models) - New **7x longer context RL** vs. all other setups, via our new batching algorithms. [Blog](https://unsloth.ai/docs/new/grpo-long-context) @@ -125,25 +124,33 @@ pip install --upgrade --force-reinstall --no-cache-dir unsloth unsloth_zoo ``` See [here](#advanced-pip-installation) for advanced pip install instructions. ### Windows Installation +For this method, we will be utilizing Anaconda. You can view the [full guide with screenshots here](https://unsloth.ai/docs/get-started/install/windows-installation). +1. **Install Miniconda (or Anaconda):** Miniconda is recommended. Install [Miniconda](https://www.anaconda.com/docs/getting-started/miniconda/install) or [Anaconda](https://www.anaconda.com/download), then open Anaconda PowerShell Prompt to continue. -1. **Install NVIDIA Video Driver:** - You should install the latest driver for your GPU. Download drivers here: [NVIDIA GPU Driver](https://www.nvidia.com/Download/index.aspx). +2. **Create a Conda Environment:** Create and activate a fresh Python 3.12 environment for Unsloth. -2. **Install Visual Studio C++:** - You will need Visual Studio, with C++ installed. By default, C++ is not installed with [Visual Studio](https://visualstudio.microsoft.com/vs/community/), so make sure you select all of the C++ options. Also select options for Windows 10/11 SDK. For detailed instructions with options, see [here](https://unsloth.ai/docs/get-started/install/windows-installation#method-3-windows-directly). + ```bash + conda create --name unsloth_env python==3.12 -y + conda activate unsloth_env + ``` -3. **Install CUDA Toolkit:** - Follow the instructions to install [CUDA Toolkit](https://developer.nvidia.com/cuda-toolkit-archive). +3. **Check Your GPU and CUDA Version:** Run `nvidia-smi` to confirm that your NVIDIA GPU is detected and note the CUDA version shown in the output. If `nvidia-smi` does not work, reinstall the latest [NVIDIA drivers](https://www.nvidia.com/en-us/drivers/). -4. **Install PyTorch:** - You will need the correct version of PyTorch that is compatible with your CUDA drivers, so make sure to select them carefully. - [Install PyTorch](https://pytorch.org/get-started/locally/). +4. **Install PyTorch:** Install the Windows pip build of PyTorch that matches your CUDA version. Use [Install PyTorch](https://pytorch.org/get-started/locally/) to select the correct command for your system, then verify that PyTorch can see your GPU. -5. **Install Unsloth:** - -```python -pip install unsloth -``` + ```python + import torch + print(torch.cuda.is_available()) + A = torch.ones((10, 10), device="cuda") + B = torch.ones((10, 10), device="cuda") + A @ B + ``` + +5. **Install Unsloth:** Only install Unsloth after PyTorch is working correctly. + + ```bash + pip install unsloth + ``` #### Advanced/Troubleshooting For **advanced installation instructions** or if you see weird errors during installations: From bdb98a96ee8de6d4f14630b7344565d5d33cb814 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E9=87=91=E9=BB=84=E8=89=B2=E8=91=A1=E8=90=84=E7=90=83?= =?UTF-8?q?=E5=90=9B=E5=90=9B?= Date: Tue, 3 Mar 2026 19:05:32 +0800 Subject: [PATCH 55/63] Revert "perf(ROCm): optimize chunked CE loss num_warps for RDNA GPUs (#4123)" (#4139) This reverts commit 6e499c421b0a0af1caf991334bb47d1d3dc5c471. --- unsloth/kernels/cross_entropy_loss.py | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/unsloth/kernels/cross_entropy_loss.py b/unsloth/kernels/cross_entropy_loss.py index 1bb1691776..d92229314f 100644 --- a/unsloth/kernels/cross_entropy_loss.py +++ b/unsloth/kernels/cross_entropy_loss.py @@ -22,7 +22,6 @@ from .utils import ( triton_cast, torch_gpu_device, is_cdna, - is_rdna, ) from transformers.models.llama.modeling_llama import logger from unsloth_zoo.utils import Version @@ -365,7 +364,7 @@ class Fast_CrossEntropyLoss(torch.autograd.Function): SOFTCAP = logit_softcapping, DO_LOGIT_SCALING = DO_LOGIT_SCALING, LOGIT_SCALE = logit_scaling, - num_warps = 16 if is_cdna() or is_rdna() else 32, + num_warps = 32 if not is_cdna() else 16, ) # logsumexp(chunked_logsumexp) - x # Do the -x separately From f1c847ca8271a974077ae0264af302f9ffaadfd3 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E9=87=91=E9=BB=84=E8=89=B2=E8=91=A1=E8=90=84=E7=90=83?= =?UTF-8?q?=E5=90=9B=E5=90=9B?= Date: Tue, 3 Mar 2026 19:05:38 +0800 Subject: [PATCH 56/63] fix(ROCm): restrict is_rdna() to ROCm-officially-supported RDNA GPUs (#4136) Current arch.startswith("gfx1") incorrectly matches: - RDNA1 (gfx10xx) and RDNA2 (gfx103x): not ROCm supported - gfx1102 (RX 7600), gfx1103 (Phoenix APU): not in ROCm support matrix - gfx1150/1151/1152 (RDNA3.5 APUs): not in ROCm support matrix Replace with explicit whitelist aligned to the ROCm Linux support matrix: https://rocm.docs.amd.com/projects/install-on-linux/en/latest/reference/system-requirements.html gfx1100 - RDNA3 discrete (RX 7900 series, PRO W7900/W7800) gfx1101 - RDNA3 discrete (RX 7800/7700 series, PRO W7700) gfx1200 - RDNA4 discrete (RX 9060 series) gfx1201 - RDNA4 discrete (RX 9070 series, AI PRO R9700) Mirrors the existing is_cdna() pattern. Avoids silently applying unverified Triton kernel tuning to unsupported hardware. --- unsloth/kernels/utils.py | 12 +++++++----- 1 file changed, 7 insertions(+), 5 deletions(-) diff --git a/unsloth/kernels/utils.py b/unsloth/kernels/utils.py index 3122282b38..90f2d5d238 100644 --- a/unsloth/kernels/utils.py +++ b/unsloth/kernels/utils.py @@ -88,11 +88,13 @@ def is_cdna(): @functools.lru_cache(1) def is_rdna(): - """Detect RDNA consumer/workstation GPUs (RDNA3, RDNA3.5, RDNA4).""" - if not is_hip(): - return False - arch = triton.runtime.driver.active.get_current_target().arch - return arch.startswith("gfx1") and not is_cdna() + """Detect ROCm-supported RDNA consumer/workstation GPUs (RDNA3, RDNA4).""" + return is_hip() and triton.runtime.driver.active.get_current_target().arch in ( + "gfx1100", + "gfx1101", + "gfx1200", + "gfx1201", + ) def calculate_settings( From bc3b390feac4a7f1c6460e4bb3f535a8e1ca8476 Mon Sep 17 00:00:00 2001 From: Datta Nimmaturi Date: Tue, 3 Mar 2026 20:00:13 +0530 Subject: [PATCH 57/63] [Fix] lm_head lora save (#4106) * Fix lm_head lora save * Fix _need_to_train_embeddings guard for lm_head LoRA targets When lm_head is already in final_modules as a LoRA target, the _need_to_train_embeddings block should not also add it to modules_to_save. This prevents dual-wrapping (LoRA + modules_to_save on the same module) which causes assertion failures downstream. Check if embed_tokens/lm_head are already being trained as LoRA targets before adding them to modules_to_save. Also prevents duplicate entries with elif guards. * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --------- Co-authored-by: Daniel Han Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- unsloth/models/llama.py | 45 ++++++++++++++++++++--------------------- 1 file changed, 22 insertions(+), 23 deletions(-) diff --git a/unsloth/models/llama.py b/unsloth/models/llama.py index 6fe60cf940..6e2fdb57ca 100644 --- a/unsloth/models/llama.py +++ b/unsloth/models/llama.py @@ -2957,6 +2957,7 @@ class FastLlamaModel: accepted_modules = frozenset( ( + "lm_head", "q_proj", "k_proj", "v_proj", @@ -2975,18 +2976,7 @@ class FastLlamaModel: train_embed_tokens = False final_modules = [] for module in target_modules: - if module == "lm_head": - # logger.warning_once( - # "Unsloth: `lm_head` should be placed in `modules_to_save` and not `target_modules`. "\ - # "Luckily, we shall do it for you!" - # ) - train_lm_head = True - if modules_to_save is None: - modules_to_save = ["lm_head"] - else: - modules_to_save.append("lm_head") - - elif module == "embed_tokens": + if module == "embed_tokens": # logger.warning_once( # "Unsloth: `embed_tokens` should be placed in `modules_to_save` and not `target_modules`. "\ # "Luckily, we shall do it for you!" @@ -3011,23 +3001,32 @@ class FastLlamaModel: # Check if we added new tokens! if hasattr(model, "_need_to_train_embeddings"): - if not train_lm_head or not train_embed_tokens: + # Check if embed_tokens/lm_head are already being trained + # (either as LoRA targets in final_modules or via modules_to_save) + _embed_already_trained = ( + train_embed_tokens or "embed_tokens" in final_modules + ) + _lm_head_already_trained = train_lm_head or "lm_head" in final_modules + if not _lm_head_already_trained or not _embed_already_trained: print( "Unsloth: You added new tokens but did not specify if you wanted to " "train the lm_head and embed_tokens.\nWe must turn it on for you." ) - train_lm_head = True - train_embed_tokens = True - if modules_to_save is None: - modules_to_save = ["embed_tokens"] - else: - modules_to_save.append("embed_tokens") + # Only add to modules_to_save if not already a LoRA target + if not _embed_already_trained: + train_embed_tokens = True + if modules_to_save is None: + modules_to_save = ["embed_tokens"] + elif "embed_tokens" not in modules_to_save: + modules_to_save.append("embed_tokens") - if modules_to_save is None: - modules_to_save = ["lm_head"] - else: - modules_to_save.append("lm_head") + if not _lm_head_already_trained: + train_lm_head = True + if modules_to_save is None: + modules_to_save = ["lm_head"] + elif "lm_head" not in modules_to_save: + modules_to_save.append("lm_head") # Check for Llama-3 # if hasattr(model._saved_temp_tokenizer, "_using_llama3_template"): From f1a79488bc781c91c826b84a1a520163a559844e Mon Sep 17 00:00:00 2001 From: Lei Zhenyuan Date: Tue, 3 Mar 2026 22:33:45 +0800 Subject: [PATCH 58/63] add intel support for torch210 within pyproject.toml (#4144) * add intel support for torch210 * fix for typo --- pyproject.toml | 37 +++++++++++++++++++++++++++++++++++++ 1 file changed, 37 insertions(+) diff --git a/pyproject.toml b/pyproject.toml index d21ceb490f..26a1499e16 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -989,6 +989,43 @@ intelgputorch290 = [ intel-gpu-torch290 = [ "unsloth[intelgputorch290]" ] +intelgputorch210 = [ + "unsloth_zoo[intelgpu]", + "unsloth[huggingfacenotorch]", + + "pytorch_triton_xpu @ https://download.pytorch.org/whl/pytorch_triton_xpu-3.5.0-cp310-cp310-manylinux_2_27_x86_64.manylinux_2_28_x86_64.whl#sha256=c169a1de14c19673b17c751290d467fa282fc90fa5da4314b2e5cdab1f553146 ; platform_system == 'Linux' and python_version == '3.10' and platform_machine == 'x86_64'", + "pytorch_triton_xpu @ https://download.pytorch.org/whl/pytorch_triton_xpu-3.5.0-cp311-cp311-manylinux_2_27_x86_64.manylinux_2_28_x86_64.whl#sha256=013d9dd5d6479bd22983161f462e61c8dbe1d82e6730624a7a8d5945507eaa61 ; platform_system == 'Linux' and python_version == '3.11' and platform_machine == 'x86_64'", + "pytorch_triton_xpu @ https://download.pytorch.org/whl/pytorch_triton_xpu-3.5.0-cp312-cp312-manylinux_2_27_x86_64.manylinux_2_28_x86_64.whl#sha256=afc8cabfbf7ed51fd278d1e0f88d6afc157b0201bad4b99d681e4d542f9e66d4 ; platform_system == 'Linux' and python_version == '3.12' and platform_machine == 'x86_64'", + "pytorch_triton_xpu @ https://download.pytorch.org/whl/pytorch_triton_xpu-3.5.0-cp313-cp313-manylinux_2_27_x86_64.manylinux_2_28_x86_64.whl#sha256=0d24c1716088f2764d0d24c64227732195b6a42706c3c5fc89eeb4904bfa0818 ; platform_system == 'Linux' and python_version == '3.13' and platform_machine == 'x86_64'", + "pytorch_triton_xpu @ https://download.pytorch.org/whl/pytorch_triton_xpu-3.5.0-cp310-cp310-win_amd64.whl#sha256=c83ab007311d9cfb6e809ee5a4587d99a9eef4be720b90da4f1aaa68b45139a0 ; sys_platform == 'win32' and python_version == '3.10' and (platform_machine == 'AMD64' or platform_machine == 'x86_64')", + "pytorch_triton_xpu @ https://download.pytorch.org/whl/pytorch_triton_xpu-3.5.0-cp311-cp311-win_amd64.whl#sha256=debf75348da8e8c7166b4d4a9b91d1508bb8d6581e339f79f7604b2e6746bacd ; sys_platform == 'win32' and python_version == '3.11' and (platform_machine == 'AMD64' or platform_machine == 'x86_64')", + "pytorch_triton_xpu @ https://download.pytorch.org/whl/pytorch_triton_xpu-3.5.0-cp312-cp312-win_amd64.whl#sha256=97337a47425f1963a723475bd61037460e84ba01db4f87a1d662c3718ff6c47e ; sys_platform == 'win32' and python_version == '3.12' and (platform_machine == 'AMD64' or platform_machine == 'x86_64')", + "pytorch_triton_xpu @ https://download.pytorch.org/whl/pytorch_triton_xpu-3.5.0-cp313-cp313-win_amd64.whl#sha256=2caf8138695f6abb023ecd02031a2611ba1bf8fff2f19802567cb2fadefe9e87 ; sys_platform == 'win32' and python_version == '3.13' and (platform_machine == 'AMD64' or platform_machine == 'x86_64')", + + "torch @ https://download.pytorch.org/whl/xpu/torch-2.10.0%2Bxpu-cp310-cp310-linux_x86_64.whl#sha256=abb1d1ec1ac672bac0ff35420c965f2df0c636ef9d94e2a830e34578489d0a57 ; platform_system == 'Linux' and python_version == '3.10' and platform_machine == 'x86_64'", + "torch @ https://download.pytorch.org/whl/xpu/torch-2.10.0%2Bxpu-cp311-cp311-linux_x86_64.whl#sha256=71ad2f82da0f41eaec159f39fc85854e27c2391efa91b373e550648a6f4aaad3 ; platform_system == 'Linux' and python_version == '3.11' and platform_machine == 'x86_64'", + "torch @ https://download.pytorch.org/whl/xpu/torch-2.10.0%2Bxpu-cp312-cp312-linux_x86_64.whl#sha256=b473571d478912f92881cc13f15fa18f8463fb0fb8a068c96ed47a7d45a4da0a ; platform_system == 'Linux' and python_version == '3.12' and platform_machine == 'x86_64'", + "torch @ https://download.pytorch.org/whl/xpu/torch-2.10.0%2Bxpu-cp313-cp313-linux_x86_64.whl#sha256=3bc64a746ff25a93de140902c60c9e819d7413f5cea1e88d80999c27a5901e9c ; platform_system == 'Linux' and python_version == '3.13' and platform_machine == 'x86_64'", + "torch @ https://download.pytorch.org/whl/xpu/torch-2.10.0%2Bxpu-cp310-cp310-win_amd64.whl#sha256=ce50691ab3fb6301d9b7bb8b3834cf5fa7152a2b5f91fd24c5efdc601a25b780 ; sys_platform == 'win32' and python_version == '3.10' and (platform_machine == 'AMD64' or platform_machine == 'x86_64')", + "torch @ https://download.pytorch.org/whl/xpu/torch-2.10.0%2Bxpu-cp311-cp311-win_amd64.whl#sha256=cb9d37f21cb9fb7df67d62863f021c3144e8d8832b9ea8e8523ac308bc620ea1 ; sys_platform == 'win32' and python_version == '3.11' and (platform_machine == 'AMD64' or platform_machine == 'x86_64')", + "torch @ https://download.pytorch.org/whl/xpu/torch-2.10.0%2Bxpu-cp312-cp312-win_amd64.whl#sha256=3ad605be4728b6d3a28a44d07dd794b1a9e45551b0057815bf25eb2a6d6a56a7 ; sys_platform == 'win32' and python_version == '3.12' and (platform_machine == 'AMD64' or platform_machine == 'x86_64')", + "torch @ https://download.pytorch.org/whl/xpu/torch-2.10.0%2Bxpu-cp313-cp313-win_amd64.whl#sha256=2b4b56dd6c792aef82006904fa888692e3782e4ae5da27526801bad4898f05a5 ; sys_platform == 'win32' and python_version == '3.13' and (platform_machine == 'AMD64' or platform_machine == 'x86_64')", + + "bitsandbytes @ https://github.com/bitsandbytes-foundation/bitsandbytes/releases/download/continuous-release_main/bitsandbytes-1.33.7.preview-py3-none-manylinux_2_24_x86_64.whl ; ('linux' in sys_platform) and (platform_machine == 'AMD64' or platform_machine == 'x86_64')", + "bitsandbytes @ https://github.com/bitsandbytes-foundation/bitsandbytes/releases/download/continuous-release_main/bitsandbytes-1.33.7.preview-py3-none-win_amd64.whl ; (sys_platform == 'win32') and (platform_machine == 'AMD64' or platform_machine == 'x86_64')", + + "torchvision @ https://download.pytorch.org/whl/xpu/torchvision-0.25.0%2Bxpu-cp310-cp310-manylinux_2_28_x86_64.whl#sha256=7e1e7b170fcf7161c8499b67156c5a05462243626dc0974010791a0bab4378d3 ; platform_system == 'Linux' and python_version == '3.10' and platform_machine == 'x86_64'", + "torchvision @ https://download.pytorch.org/whl/xpu/torchvision-0.25.0%2Bxpu-cp311-cp311-manylinux_2_28_x86_64.whl#sha256=bd6add201bd7628af70437292e1447abb368e0b5f4ff9abd334ae435efd44792 ; platform_system == 'Linux' and python_version == '3.11' and platform_machine == 'x86_64'", + "torchvision @ https://download.pytorch.org/whl/xpu/torchvision-0.25.0%2Bxpu-cp312-cp312-manylinux_2_28_x86_64.whl#sha256=6ad2543496bc29e59d3dd614a94d09aa9870318aedb66045344fffddfedd2cf8 ; platform_system == 'Linux' and python_version == '3.12' and platform_machine == 'x86_64'", + "torchvision @ https://download.pytorch.org/whl/xpu/torchvision-0.25.0%2Bxpu-cp313-cp313-manylinux_2_28_x86_64.whl#sha256=80269f37865fcd8b57f20e4786efae2200bfa2b2727926c3c7acc82f0e7d3548 ; platform_system == 'Linux' and python_version == '3.13' and platform_machine == 'x86_64'", + "torchvision @ https://download.pytorch.org/whl/xpu/torchvision-0.25.0%2Bxpu-cp310-cp310-win_amd64.whl#sha256=6b9485ba85dcba4d196d6134d9c3332fb228fb2556416bf0450a64e8a472fcba ; sys_platform == 'win32' and python_version == '3.10' and (platform_machine == 'AMD64' or platform_machine == 'x86_64')", + "torchvision @ https://download.pytorch.org/whl/xpu/torchvision-0.25.0%2Bxpu-cp311-cp311-win_amd64.whl#sha256=36cbaedf10f6412af5c89afd9aeea474e6a56a0050348ada8fabe1ecaf6b879e ; sys_platform == 'win32' and python_version == '3.11' and (platform_machine == 'AMD64' or platform_machine == 'x86_64')", + "torchvision @ https://download.pytorch.org/whl/xpu/torchvision-0.25.0%2Bxpu-cp312-cp312-win_amd64.whl#sha256=738357d97468d75fe3d510ac37e65130f2787f81d9bbc1518898f7396dc3403f ; sys_platform == 'win32' and python_version == '3.12' and (platform_machine == 'AMD64' or platform_machine == 'x86_64')", + "torchvision @ https://download.pytorch.org/whl/xpu/torchvision-0.25.0%2Bxpu-cp313-cp313-win_amd64.whl#sha256=1c4b44b36a557f7381e3076fb8843366742238648441d607c8d049c6da0f8886 ; sys_platform == 'win32' and python_version == '3.13' and (platform_machine == 'AMD64' or platform_machine == 'x86_64')", +] +intel-gpu-torch210 = [ + "unsloth[intelgputorch210]" +] intel = [ "unsloth[intelgputorch280]", ] From 363cc6884149e06151a8ac4f6bdcb996fbda076a Mon Sep 17 00:00:00 2001 From: Roland Tannous <115670425+rolandtannous@users.noreply.github.com> Date: Tue, 3 Mar 2026 18:34:09 +0400 Subject: [PATCH 59/63] fix: update GGUF save paths to use ~/.unsloth/llama.cpp with Windows support (#4138) * fix: update GGUF save paths to use ~/.unsloth/llama.cpp with Windows support * fix: quote LLAMA_CPP_DEFAULT_DIR in fallback shell commands to handle paths with spaces * refactor: deduplicate platform-specific build instructions in quantization error message * chore: remove accidentally committed PR description file * Fix import safety and f-string bugs in save.py - H4: Add defensive try/except for LLAMA_CPP_DEFAULT_DIR and IS_WINDOWS imports with fallback defaults, so save.py works even if zoo PR #526 is not merged yet - H5: Fix Kaggle error path using plain "Error: {e}" instead of f"Error: {e}", so the actual exception is shown to users * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --------- Co-authored-by: Datta Nimmaturi Co-authored-by: Daniel Han Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- unsloth/save.py | 42 +++++++++++++++++++++++++++++++++--------- 1 file changed, 33 insertions(+), 9 deletions(-) diff --git a/unsloth/save.py b/unsloth/save.py index 66d5ae2d60..6e38d1e952 100644 --- a/unsloth/save.py +++ b/unsloth/save.py @@ -23,6 +23,16 @@ from unsloth_zoo.llama_cpp import ( check_llama_cpp, _download_convert_hf_to_gguf, ) + +# H4: Defensive imports -- these were added in unsloth-zoo PR #526 +# and may not exist on older versions +try: + from unsloth_zoo.llama_cpp import LLAMA_CPP_DEFAULT_DIR, IS_WINDOWS +except ImportError: + import sys + + IS_WINDOWS = sys.platform == "win32" + LLAMA_CPP_DEFAULT_DIR = "llama.cpp" from bitsandbytes.nn import Linear4bit as Bnb_Linear4bit from peft.tuners.lora import Linear4bit as Peft_Linear4bit from peft.tuners.lora import Linear as Peft_Linear @@ -1316,18 +1326,26 @@ def save_to_gguf( "`model.{save_pretrained/push_to_hub}_gguf will use too much disk space.\n" "You can try saving it to the `/tmp` directory for larger disk space.\n" "I suggest you to save the 16bit model first, then use manual llama.cpp conversion.\n" - "Error: {e}" + f"Error: {e}" ) else: + if IS_WINDOWS: + build_instructions = ( + f'cd "{LLAMA_CPP_DEFAULT_DIR}"\n' + f"cmake -S . -B build -DBUILD_SHARED_LIBS=OFF\n" + f"cmake --build build --config Release" + ) + else: + build_instructions = f'cd "{LLAMA_CPP_DEFAULT_DIR}" && make clean && make all -j' + raise RuntimeError( f"Unsloth: Quantization failed for {output_location}\n" "You might have to compile llama.cpp yourself, then run this again.\n" "You do not need to close this Python program. Run the following commands in a new terminal:\n" - "You must run this in the same folder as you're saving your model.\n" - "git clone --recursive https://github.com/ggerganov/llama.cpp\n" - "cd llama.cpp && make clean && make all -j\n" + f'git clone --recursive https://github.com/ggerganov/llama.cpp "{LLAMA_CPP_DEFAULT_DIR}"\n' + f"{build_instructions}\n" "Once that's done, redo the quantization.\n" - "Error: {e}" + f"Error: {e}" ) print("Unsloth: Model files cleanup...") if quants_created: @@ -2084,16 +2102,22 @@ def unsloth_save_pretrained_gguf( "Unsloth: ##### We removed it in GGUF's chat template for you." ) + _exe = ".exe" if IS_WINDOWS else "" + if IS_WINDOWS: + _bin_dir = os.path.join(LLAMA_CPP_DEFAULT_DIR, "build", "bin", "Release") + else: + _bin_dir = LLAMA_CPP_DEFAULT_DIR + if is_vlm_update: print("\n") print( - f"Unsloth: example usage for Multimodal LLMs: llama.cpp/llama-mtmd-cli -m {all_file_locations[0]} --mmproj {all_file_locations[-1]}" + f"Unsloth: example usage for Multimodal LLMs: {os.path.join(_bin_dir, 'llama-mtmd-cli' + _exe)} -m {all_file_locations[0]} --mmproj {all_file_locations[-1]}" ) print("Unsloth: load image inside llama.cpp runner: /image test_image.jpg") print("Unsloth: Prompt model to describe the image") else: print( - f'Unsloth: example usage for text only LLMs: llama.cpp/llama-cli --model {all_file_locations[0]} -p "why is the sky blue?"' + f'Unsloth: example usage for text only LLMs: {os.path.join(_bin_dir, "llama-cli" + _exe)} --model {all_file_locations[0]} -p "why is the sky blue?"' ) if ollama_success: @@ -2307,8 +2331,8 @@ tags: This model was finetuned and converted to GGUF format using [Unsloth](https://github.com/unslothai/unsloth). **Example usage**: -- For text only LLMs: `./llama.cpp/llama-cli -hf {repo_id} --jinja` -- For multimodal models: `./llama.cpp/llama-mtmd-cli -hf {repo_id} --jinja` +- For text only LLMs: `llama-cli -hf {repo_id} --jinja` +- For multimodal models: `llama-mtmd-cli -hf {repo_id} --jinja` ## Available Model files: """ From 2afd686c621795a0204283b8c09d2d1c5f2fd7a1 Mon Sep 17 00:00:00 2001 From: Mustafa Eyceoz Date: Tue, 3 Mar 2026 09:45:41 -0500 Subject: [PATCH 60/63] Fix multi-node distributed training with single GPU per node (#4143) --- unsloth/models/_utils.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/unsloth/models/_utils.py b/unsloth/models/_utils.py index f883d466f0..8684991431 100644 --- a/unsloth/models/_utils.py +++ b/unsloth/models/_utils.py @@ -1473,7 +1473,7 @@ BitsAndBytesConfig__init__ = BitsAndBytesConfig__init__.replace( ) exec(BitsAndBytesConfig__init__, globals()) -if DEVICE_COUNT == 1: +if DEVICE_COUNT == 1 and int(os.environ.get("WORLD_SIZE", "1")) <= 1: from accelerate.utils.dataclasses import DistributedType def _prepare_backend(self, *args, **kwargs): From deca813464fc3e01310c05c13e1c64daf06f3b21 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 3 Mar 2026 06:53:58 -0800 Subject: [PATCH 61/63] Update --- pyproject.toml | 8 ++++---- unsloth/models/_utils.py | 2 +- 2 files changed, 5 insertions(+), 5 deletions(-) diff --git a/pyproject.toml b/pyproject.toml index 26a1499e16..9d410b4033 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -55,13 +55,13 @@ huggingfacenotorch = [ "huggingface_hub>=0.34.0", "hf_transfer", "diffusers", - "transformers>=4.51.3,!=4.52.0,!=4.52.1,!=4.52.2,!=4.52.3,!=4.53.0,!=4.54.0,!=4.55.0,!=4.55.1,!=4.57.0,!=4.57.4,!=4.57.5,<=4.57.6", + "transformers>=4.51.3,!=4.52.0,!=4.52.1,!=4.52.2,!=4.52.3,!=4.53.0,!=4.54.0,!=4.55.0,!=4.55.1,!=4.57.0,!=4.57.4,!=4.57.5,!=5.0.0,!=5.1.0,<=5.2.0", "trl>=0.18.2,!=0.19.0,<=0.24.0", "sentence-transformers", ] huggingface = [ "unsloth[huggingfacenotorch]", - "unsloth_zoo>=2026.2.1", + "unsloth_zoo>=2026.3.1", "torchvision", "unsloth[triton]", ] @@ -551,10 +551,10 @@ colab-ampere-torch220 = [ "flash-attn>=2.6.3 ; ('linux' in sys_platform)", ] colab-new = [ - "unsloth_zoo>=2026.2.1", + "unsloth_zoo>=2026.3.1", "packaging", "tyro", - "transformers>=4.51.3,!=4.52.0,!=4.52.1,!=4.52.2,!=4.52.3,!=4.53.0,!=4.54.0,!=4.55.0,!=4.55.1,!=4.57.0,!=4.57.4,!=4.57.5,<=4.57.6", + "transformers>=4.51.3,!=4.52.0,!=4.52.1,!=4.52.2,!=4.52.3,!=4.53.0,!=4.54.0,!=4.55.0,!=4.55.1,!=4.57.0,!=4.57.4,!=4.57.5,!=5.0.0,!=5.1.0,<=5.2.0", "datasets>=3.4.1,!=4.0.*,!=4.1.0,<4.4.0", "sentencepiece>=0.2.0", "tqdm", diff --git a/unsloth/models/_utils.py b/unsloth/models/_utils.py index 8684991431..d0d74ab5bd 100644 --- a/unsloth/models/_utils.py +++ b/unsloth/models/_utils.py @@ -12,7 +12,7 @@ # See the License for the specific language governing permissions and # limitations under the License. -__version__ = "2026.2.1" +__version__ = "2026.3.1" __all__ = [ "SUPPORTS_BFLOAT16", From f636b198b536f9e21a6e221bfb2c14b1eef49da4 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 3 Mar 2026 06:55:08 -0800 Subject: [PATCH 62/63] Update __init__.py --- unsloth/__init__.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/unsloth/__init__.py b/unsloth/__init__.py index 60c5b3bba6..aff9e3d186 100644 --- a/unsloth/__init__.py +++ b/unsloth/__init__.py @@ -89,7 +89,7 @@ from importlib.metadata import PackageNotFoundError # Check for unsloth_zoo try: unsloth_zoo_version = importlib_version("unsloth_zoo") - if Version(unsloth_zoo_version) < Version("2026.2.1"): + if Version(unsloth_zoo_version) < Version("2026.3.1"): print( "Unsloth: Please update Unsloth and Unsloth-Zoo to the latest version!\n" "Do this via `pip install --upgrade --force-reinstall --no-cache-dir --no-deps unsloth unsloth_zoo`" From ef316549eac87ef069d4f0994d14c4a3de15392d Mon Sep 17 00:00:00 2001 From: Datta Nimmaturi Date: Tue, 3 Mar 2026 20:27:25 +0530 Subject: [PATCH 63/63] Fixup mapper issues and resolve properly (#4124) * Fixup mapper issues and resolve properly * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --------- Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- tests/test_get_model_name.py | 127 +++++++++++++++++++++++++++++++++ unsloth/models/_utils.py | 54 -------------- unsloth/models/llama.py | 12 +--- unsloth/models/loader.py | 14 +++- unsloth/models/loader_utils.py | 55 +++++++++----- unsloth/models/mapper.py | 74 ++++++++++--------- unsloth/models/vision.py | 16 +---- 7 files changed, 224 insertions(+), 128 deletions(-) create mode 100644 tests/test_get_model_name.py diff --git a/tests/test_get_model_name.py b/tests/test_get_model_name.py new file mode 100644 index 0000000000..ad89f595f0 --- /dev/null +++ b/tests/test_get_model_name.py @@ -0,0 +1,127 @@ +import unittest +from unittest.mock import patch +from unsloth.models.loader_utils import get_model_name +from unsloth.models import loader_utils +from unsloth.models.mapper import FLOAT_TO_INT_MAPPER, MAP_TO_UNSLOTH_16bit + + +def _no_remote_mapper(): + return {}, {}, {} + + +class TestGetModelName(unittest.TestCase): + def _assert_mapping(self, model_name, load_in_4bit, expected, should_change): + mapped = get_model_name(model_name, load_in_4bit = load_in_4bit) + self.assertEqual(mapped.lower(), expected.lower()) + if should_change: + self.assertNotEqual(mapped.lower(), model_name.lower()) + else: + self.assertEqual(mapped.lower(), model_name.lower()) + + @patch.object(loader_utils, "_get_new_mapper", _no_remote_mapper) + def test_resolution_matrix(self): + cases = [ + # Core mappings + ("meta-llama/Llama-2-7b-hf", True, "unsloth/llama-2-7b-bnb-4bit", True), + ("meta-llama/Llama-2-7b-hf", False, "unsloth/llama-2-7b", True), + ( + "mistralai/Ministral-8B-Instruct-2410", + True, + "mistralai/Ministral-8B-Instruct-2410", + False, + ), + ( + "meta-llama/Llama-3.2-1B-Instruct", + False, + "unsloth/Llama-3.2-1B-Instruct", + True, + ), + ( + "meta-llama/Llama-2-7b-chat-hf", + True, + "unsloth/llama-2-7b-chat-bnb-4bit", + True, + ), + ( + "meta-llama/Llama-3.3-70B-Instruct", + True, + "unsloth/llama-3.3-70b-instruct-unsloth-bnb-4bit", + True, + ), + ("Qwen/Qwen3-8B", True, "unsloth/Qwen3-8B-unsloth-bnb-4bit", True), + ("Qwen/Qwen3-8B", False, "unsloth/Qwen3-8B", True), + ("Qwen/Qwen3-8B-FP8", False, "unsloth/Qwen3-8B-FP8", True), + ("Qwen/Qwen3-8B-FP8", True, "unsloth/Qwen3-8B-unsloth-bnb-4bit", True), + ( + "mistralai/Ministral-3-3B-Instruct-2512", + True, + "unsloth/Ministral-3-3B-Instruct-2512-unsloth-bnb-4bit", + True, + ), + ( + "mistralai/Ministral-3-3B-Instruct-2512", + False, + "unsloth/Ministral-3-3B-Instruct-2512", + True, + ), + ("unsloth/Kimi-K2-Instruct", True, "unsloth/Kimi-K2-Instruct-BF16", True), + ("unsloth/Kimi-K2-Instruct", False, "unsloth/Kimi-K2-Instruct", False), + # Fallback-to-original behavior + "nonexistent-user/nonexistent-model-123", + "google/gemma-3-random-prototype-123", + "imdatta0/nanoqwen-fp8", + "imdatta0/nanoqwen-bf16", + # Backward compatibility for legacy 4bit names + ("unsloth/llama-2-7b-bnb-4bit", True, "unsloth/llama-2-7b-bnb-4bit", False), + ("unsloth/llama-2-7b-bnb-4bit", False, "unsloth/llama-2-7b", True), + ("google/gemma-2-9b", True, "unsloth/gemma-2-9b-bnb-4bit", True), + # GPT-OSS behavior + ("openai/gpt-oss-20b", False, "unsloth/gpt-oss-20b", True), + ("openai/gpt-oss-20b", True, "unsloth/gpt-oss-20b-unsloth-bnb-4bit", True), + ("unsloth/gpt-oss-20b", True, "unsloth/gpt-oss-20b-unsloth-bnb-4bit", True), + ("unsloth/gpt-oss-20b-bf16", True, "unsloth/gpt-oss-20b-bf16", False), + ( + "unsloth/gpt-oss-20b-unsloth-bnb-4bit", + False, + "unsloth/gpt-oss-20b", + True, + ), + ( + "unsloth/gpt-oss-20b-bnb-4bit", + True, + "unsloth/gpt-oss-20b-bnb-4bit", + False, + ), + ] + for case in cases: + if isinstance(case, str): + model_name = case + with self.subTest(model_name = model_name, load_in_4bit = True): + self._assert_mapping(model_name, True, model_name, False) + else: + model_name, load_in_4bit, expected, should_change = case + with self.subTest(model_name = model_name, load_in_4bit = load_in_4bit): + self._assert_mapping( + model_name, load_in_4bit, expected, should_change + ) + + def test_static_mapper_contract(self): + contracts = [ + ("qwen/qwen3-8b", "unsloth/qwen3-8b-unsloth-bnb-4bit"), + ("qwen/qwen3-8b-fp8", "unsloth/qwen3-8b-unsloth-bnb-4bit"), + ( + "mistralai/ministral-3-3b-instruct-2512", + "unsloth/ministral-3-3b-instruct-2512-unsloth-bnb-4bit", + ), + ("unsloth/kimi-k2-instruct", "unsloth/kimi-k2-instruct-bf16"), + ] + for src, expected in contracts: + with self.subTest(src = src): + self.assertEqual(FLOAT_TO_INT_MAPPER[src], expected) + self.assertEqual( + MAP_TO_UNSLOTH_16bit["qwen/qwen3-8b-fp8"], "unsloth/Qwen3-8B-FP8" + ) + + +if __name__ == "__main__": + unittest.main() diff --git a/unsloth/models/_utils.py b/unsloth/models/_utils.py index d0d74ab5bd..b74743d724 100644 --- a/unsloth/models/_utils.py +++ b/unsloth/models/_utils.py @@ -74,7 +74,6 @@ __all__ = [ "dequantize_module_weight", "patch_hf_quantizer", "verify_fp8_support_if_applicable", - "_redirect_fp8_to_bf16", "_get_inference_mode_context_manager", "hf_login", "is_moe_model", @@ -2584,59 +2583,6 @@ def patch_hf_quantizer(): patch_hf_quantizer() -def _redirect_fp8_to_bf16( - model_name, auto_config, load_in_fp8, token, trust_remote_code -): - """ - Detect FP8 quantization in model config and redirect to BF16 sibling. - - Models shipping FP8 as default (e.g. mistralai/Ministral-3-*B-Instruct) - cannot be loaded with BNB 4-bit/8-bit or 16-bit mode. This detects - quant_method in ("fp8", "fbgemm_fp8") and redirects to {model_name}-BF16. - - Redirect is SKIPPED when load_in_fp8 is truthy (True or 'block'), - meaning the user explicitly wants FP8 loading. - - Returns (model_name, auto_config) -- possibly updated. - """ - if not hasattr(auto_config, "quantization_config"): - return model_name, auto_config - - _qc = auto_config.quantization_config - _qm = ( - _qc.get("quant_method", "") - if isinstance(_qc, dict) - else getattr(_qc, "quant_method", "") - ) - if _qm not in ("fp8", "fbgemm_fp8") or load_in_fp8: - return model_name, auto_config - - _bf16_name = model_name.rstrip("/") + "-BF16" - _original_name = model_name - try: - from huggingface_hub import model_info as _hf_model_info - from transformers import AutoConfig - - _hf_model_info(_bf16_name, token = token) - _bf16_config = AutoConfig.from_pretrained( - _bf16_name, - token = token, - trust_remote_code = trust_remote_code, - ) - print( - f"Unsloth: {_original_name} uses FP8 weights. " - f"Redirecting to {_bf16_name}." - ) - return _bf16_name, _bf16_config - except Exception: - raise RuntimeError( - f"Unsloth: {_original_name} uses FP8 weights but no BF16 version " - f"was found at {_bf16_name}.\n" - f"Loading FP8 weights with BitsAndBytes or in 16-bit will fail.\n" - f"Set load_in_fp8=True to use FP8 mode, or upload a BF16 version." - ) - - def verify_fp8_support_if_applicable(model_config): quant_method = get_quant_type(model_config) if quant_method in ["fbgemm_fp8", "fp8"] and DEVICE_TYPE != "cuda": diff --git a/unsloth/models/llama.py b/unsloth/models/llama.py index 6e2fdb57ca..93d93e26d6 100644 --- a/unsloth/models/llama.py +++ b/unsloth/models/llama.py @@ -25,7 +25,8 @@ from ._utils import move_to_device from ._utils import ( _get_inference_mode_context_manager, _prepare_model_for_qat, - _redirect_fp8_to_bf16, + is_bfloat16_supported, + get_quant_type, ) from .loader_utils import _get_fp8_mode_and_check_settings from ..utils.packing import ( @@ -2331,15 +2332,6 @@ class FastLlamaModel: token = token, attn_implementation = "sdpa", ) - # Handle FP8 models: redirect to BF16 sibling when the model ships with - # FP8 weights. Redirect is skipped when load_in_fp8 is truthy (True or 'block'). - model_name, model_config = _redirect_fp8_to_bf16( - model_name, - model_config, - load_in_fp8, - token, - trust_remote_code, - ) model_config.model_name = model_name model_max_seq_length = model_config.max_position_embeddings diff --git a/unsloth/models/loader.py b/unsloth/models/loader.py index b32c8303d8..e7b975ceea 100644 --- a/unsloth/models/loader.py +++ b/unsloth/models/loader.py @@ -372,7 +372,11 @@ class FastLanguageModel(FastLlamaModel): fp8_mode = None if not use_exact_model_name: new_model_name = get_model_name( - model_name, load_in_4bit = load_in_4bit, load_in_fp8 = load_in_fp8 + model_name, + load_in_4bit = load_in_4bit, + load_in_fp8 = load_in_fp8, + token = token, + trust_remote_code = trust_remote_code, ) if new_model_name is None and load_in_fp8 != False: fp8_mode = _get_fp8_mode_and_check_settings( @@ -525,7 +529,13 @@ class FastLanguageModel(FastLlamaModel): # Check base model again for PEFT model_name = peft_config.base_model_name_or_path if not use_exact_model_name: - model_name = get_model_name(model_name, load_in_4bit) + model_name = get_model_name( + model_name, + load_in_4bit = load_in_4bit, + load_in_fp8 = load_in_fp8, + token = token, + trust_remote_code = trust_remote_code, + ) # Check if pre-quantized models are allowed # For eg AMD Instinct GPUs need blocksize = 128, but our pre-quants are blocksize = 64 if not ALLOW_PREQUANTIZED_MODELS and model_name.lower().endswith( diff --git a/unsloth/models/loader_utils.py b/unsloth/models/loader_utils.py index 40ac49ca78..cf5af983a6 100644 --- a/unsloth/models/loader_utils.py +++ b/unsloth/models/loader_utils.py @@ -198,18 +198,42 @@ def _get_new_mapper(): return {}, {}, {} -def get_model_name(model_name, load_in_4bit = True, load_in_fp8 = False): - assert load_in_fp8 in (True, False, "block") - new_model_name = __get_model_name( +def _resolve_with_mappers( + model_name, + load_in_4bit, + load_in_fp8, + int_to_float, + float_to_int, + map_to_unsloth_16bit, +): + return __get_model_name( model_name = model_name, load_in_4bit = load_in_4bit, - INT_TO_FLOAT_MAPPER = INT_TO_FLOAT_MAPPER, - FLOAT_TO_INT_MAPPER = FLOAT_TO_INT_MAPPER, - MAP_TO_UNSLOTH_16bit = MAP_TO_UNSLOTH_16bit, + INT_TO_FLOAT_MAPPER = int_to_float, + FLOAT_TO_INT_MAPPER = float_to_int, + MAP_TO_UNSLOTH_16bit = map_to_unsloth_16bit, load_in_fp8 = load_in_fp8, FLOAT_TO_FP8_BLOCK_MAPPER = FLOAT_TO_FP8_BLOCK_MAPPER, FLOAT_TO_FP8_ROW_MAPPER = FLOAT_TO_FP8_ROW_MAPPER, ) + + +def get_model_name( + model_name, + load_in_4bit = True, + load_in_fp8 = False, + token = None, + trust_remote_code = False, +): + assert load_in_fp8 in (True, False, "block") + new_model_name = _resolve_with_mappers( + model_name = model_name, + load_in_4bit = load_in_4bit, + load_in_fp8 = load_in_fp8, + int_to_float = INT_TO_FLOAT_MAPPER, + float_to_int = FLOAT_TO_INT_MAPPER, + map_to_unsloth_16bit = MAP_TO_UNSLOTH_16bit, + ) # In the rare case, we convert bad model names to other names # For eg too large dynamic quants or MoEs if ( @@ -228,15 +252,13 @@ def get_model_name(model_name, load_in_4bit = True, load_in_fp8 = False): NEW_INT_TO_FLOAT_MAPPER, NEW_FLOAT_TO_INT_MAPPER, NEW_MAP_TO_UNSLOTH_16bit = ( _get_new_mapper() ) - upgraded_model_name = __get_model_name( + upgraded_model_name = _resolve_with_mappers( model_name = model_name, load_in_4bit = load_in_4bit, - INT_TO_FLOAT_MAPPER = NEW_INT_TO_FLOAT_MAPPER, - FLOAT_TO_INT_MAPPER = NEW_FLOAT_TO_INT_MAPPER, - MAP_TO_UNSLOTH_16bit = NEW_MAP_TO_UNSLOTH_16bit, load_in_fp8 = load_in_fp8, - FLOAT_TO_FP8_BLOCK_MAPPER = FLOAT_TO_FP8_BLOCK_MAPPER, - FLOAT_TO_FP8_ROW_MAPPER = FLOAT_TO_FP8_ROW_MAPPER, + int_to_float = NEW_INT_TO_FLOAT_MAPPER, + float_to_int = NEW_FLOAT_TO_INT_MAPPER, + map_to_unsloth_16bit = NEW_MAP_TO_UNSLOTH_16bit, ) if upgraded_model_name is not None: raise NotImplementedError( @@ -245,10 +267,11 @@ def get_model_name(model_name, load_in_4bit = True, load_in_fp8 = False): 'pip install --upgrade --no-cache-dir "unsloth[colab-new] @ git+https://github.com/unslothai/unsloth.git"\n' 'pip install --upgrade --no-cache-dir "git+https://github.com/unslothai/unsloth-zoo.git"\n' ) - if load_in_fp8 != False: - # Handle on the fly TorchAO FP8 quantization - return new_model_name - return new_model_name if new_model_name is not None else model_name + + if new_model_name is None: + new_model_name = model_name + + return new_model_name def _offline_quantize_to_fp8(model_name: str, fp8_mode: str) -> str: diff --git a/unsloth/models/mapper.py b/unsloth/models/mapper.py index ec7a7a8046..f0f430eb7e 100644 --- a/unsloth/models/mapper.py +++ b/unsloth/models/mapper.py @@ -1337,6 +1337,9 @@ __INT_TO_FLOAT_MAPPER = \ "mistralai/Ministral-3-14B-Reasoning-2512", "unsloth/Ministral-3-14B-Reasoning-2512-bnb-4bit", ), + "unsloth/Kimi-K2-Instruct-BF16" : ( + "unsloth/Kimi-K2-Instruct", + ), } INT_TO_FLOAT_MAPPER = {} @@ -1345,6 +1348,19 @@ MAP_TO_UNSLOTH_16bit = {} FLOAT_TO_FP8_BLOCK_MAPPER = {} FLOAT_TO_FP8_ROW_MAPPER = {} + +def _add_with_lower(mapper, key, value): + if key is None: + return + mapper[key] = value + mapper[key.lower()] = value + + +def _add_lower_only(mapper, key, value): + if key is None: + return + mapper[key.lower()] = value + for key, values in __INT_TO_FLOAT_MAPPER.items(): block, row = None, None if type(values) is dict: @@ -1355,21 +1371,24 @@ for key, values in __INT_TO_FLOAT_MAPPER.items(): float8_values = values["8"] assert len(float8_values) == 3 official, block, row = float8_values - FLOAT_TO_FP8_BLOCK_MAPPER[key.lower()] = block - FLOAT_TO_FP8_ROW_MAPPER[key.lower()] = row - FLOAT_TO_FP8_BLOCK_MAPPER[official.lower() + "-dynamic"] = block - FLOAT_TO_FP8_ROW_MAPPER[official.lower()] = row - FLOAT_TO_FP8_ROW_MAPPER[official.lower() + "-dynamic"] = row - FLOAT_TO_FP8_BLOCK_MAPPER[float16_values[0]] = block - FLOAT_TO_FP8_BLOCK_MAPPER[float16_values[0].lower()] = block - FLOAT_TO_FP8_ROW_MAPPER[float16_values[0]] = block - FLOAT_TO_FP8_ROW_MAPPER[float16_values[0].lower()] = block - for k in float8_values: - FLOAT_TO_FP8_BLOCK_MAPPER[k.lower()] = block - FLOAT_TO_FP8_ROW_MAPPER[k.lower()] = row - for k in float16_values: - FLOAT_TO_FP8_BLOCK_MAPPER[k.lower()] = block - FLOAT_TO_FP8_ROW_MAPPER[k.lower()] = row + _add_lower_only(FLOAT_TO_FP8_BLOCK_MAPPER, key, block) + _add_lower_only(FLOAT_TO_FP8_ROW_MAPPER, key, row) + _add_lower_only(FLOAT_TO_FP8_BLOCK_MAPPER, official + "-dynamic", block) + _add_lower_only(FLOAT_TO_FP8_ROW_MAPPER, official, row) + _add_lower_only(FLOAT_TO_FP8_ROW_MAPPER, official + "-dynamic", row) + for k in float8_values + float16_values: + _add_lower_only(FLOAT_TO_FP8_BLOCK_MAPPER, k, block) + _add_lower_only(FLOAT_TO_FP8_ROW_MAPPER, k, row) + + if float8_values[1] is not None and float8_values[1].startswith("unsloth"): + for value in float8_values: + if value is not None: + _add_with_lower(MAP_TO_UNSLOTH_16bit, value, float8_values[1]) + + for value in float8_values: + if value is not None: + FLOAT_TO_INT_MAPPER[value] = key + FLOAT_TO_INT_MAPPER[value.lower()] = key.lower() values = float16_values INT_TO_FLOAT_MAPPER[key] = values[0] @@ -1379,27 +1398,16 @@ for key, values in __INT_TO_FLOAT_MAPPER.items(): # Map to Unsloth version for 16bit versions if len(values) == 2: if values[0].startswith("unsloth"): - MAP_TO_UNSLOTH_16bit[values[1]] = values[0] - MAP_TO_UNSLOTH_16bit[values[1].lower()] = values[0] - if block is not None: - MAP_TO_UNSLOTH_16bit[block] = values[0] - MAP_TO_UNSLOTH_16bit[block.lower()] = values[0] - if row is not None: - MAP_TO_UNSLOTH_16bit[row] = values[0] - MAP_TO_UNSLOTH_16bit[row.lower()] = values[0] + _add_with_lower(MAP_TO_UNSLOTH_16bit, values[1], values[0]) + _add_with_lower(MAP_TO_UNSLOTH_16bit, block, values[0]) + _add_with_lower(MAP_TO_UNSLOTH_16bit, row, values[0]) elif len(values) == 3: # Dynamic Unsloth quantization if values[0].startswith("unsloth"): - MAP_TO_UNSLOTH_16bit[values[1]] = values[0] - MAP_TO_UNSLOTH_16bit[values[1].lower()] = values[0] - MAP_TO_UNSLOTH_16bit[values[2]] = values[0] - MAP_TO_UNSLOTH_16bit[values[2].lower()] = values[0] - if block is not None: - MAP_TO_UNSLOTH_16bit[block] = values[0] - MAP_TO_UNSLOTH_16bit[block.lower()] = values[0] - if row is not None: - MAP_TO_UNSLOTH_16bit[row] = values[0] - MAP_TO_UNSLOTH_16bit[row.lower()] = values[0] + _add_with_lower(MAP_TO_UNSLOTH_16bit, values[1], values[0]) + _add_with_lower(MAP_TO_UNSLOTH_16bit, values[2], values[0]) + _add_with_lower(MAP_TO_UNSLOTH_16bit, block, values[0]) + _add_with_lower(MAP_TO_UNSLOTH_16bit, row, values[0]) pass # Get lowercased diff --git a/unsloth/models/vision.py b/unsloth/models/vision.py index 28c624c213..a8adba99e7 100644 --- a/unsloth/models/vision.py +++ b/unsloth/models/vision.py @@ -30,7 +30,6 @@ from ..kernels import ( post_patch_loss_function, ) from ._utils import __version__, importlib_version, _prepare_model_for_qat -from ._utils import _redirect_fp8_to_bf16 from ._utils import * from .loader_utils import _get_fp8_mode_and_check_settings from ..save import patch_saving_functions @@ -612,18 +611,9 @@ class FastBaseModel: model_class = None flex_attn_impl = prefer_flex_attn_if_supported(model_class, auto_config) - # Handle FP8 models: redirect to BF16 sibling when the model ships with - # FP8 weights (e.g. Ministral-3-3B-Instruct-2512). FP8 weights cannot be - # directly loaded by BNB, and the FP8 quantization config can cause issues - # even for 16-bit loading. - # Redirect is skipped when load_in_fp8 is truthy (True or 'block'). - model_name, auto_config = _redirect_fp8_to_bf16( - model_name, - auto_config, - load_in_fp8, - token, - trust_remote_code, - ) + # Handle FP8 models: get_model_name has already redirected this to BF16 sibling if the model ships with + # FP8 weights. We just need to update it here for sanity. + auto_config.model_name = model_name # Re-resolve model_class after potential config change try: model_class = auto_model._model_mapping[auto_config.__class__]