Compare commits
4 commits
main
...
fix/amd-op
| Author | SHA1 | Date | |
|---|---|---|---|
|
|
c3d86a8983 | ||
|
|
3e606aefe5 | ||
|
|
e8d897466e | ||
|
|
7ecb049ac9 |
8 changed files with 97 additions and 23 deletions
|
|
@ -234,6 +234,14 @@ elif DEVICE_TYPE == "xpu":
|
||||||
# set SUPPORTS_BFLOAT16 as torch.xpu.is_bf16_supported()
|
# set SUPPORTS_BFLOAT16 as torch.xpu.is_bf16_supported()
|
||||||
SUPPORTS_BFLOAT16 = torch.xpu.is_bf16_supported()
|
SUPPORTS_BFLOAT16 = torch.xpu.is_bf16_supported()
|
||||||
|
|
||||||
|
# Backwards compatibility: some notebooks import `unsloth.is_bf16_supported`.
|
||||||
|
# Ensure it exists on all backends (HIP / XPU) and has a stable signature.
|
||||||
|
if "is_bf16_supported" not in globals():
|
||||||
|
|
||||||
|
def is_bf16_supported(including_emulation = False):
|
||||||
|
return SUPPORTS_BFLOAT16
|
||||||
|
|
||||||
|
|
||||||
# For Gradio HF Spaces?
|
# For Gradio HF Spaces?
|
||||||
# if "SPACE_AUTHOR_NAME" not in os.environ and "SPACE_REPO_NAME" not in os.environ:
|
# if "SPACE_AUTHOR_NAME" not in os.environ and "SPACE_REPO_NAME" not in os.environ:
|
||||||
import triton
|
import triton
|
||||||
|
|
|
||||||
|
|
@ -405,7 +405,10 @@ def fix_vllm_aimv2_issue():
|
||||||
spec = importlib.util.find_spec("vllm")
|
spec = importlib.util.find_spec("vllm")
|
||||||
if spec is None:
|
if spec is None:
|
||||||
return
|
return
|
||||||
vllm_version = importlib_version("vllm")
|
try:
|
||||||
|
vllm_version = importlib_version("vllm")
|
||||||
|
except Exception:
|
||||||
|
return
|
||||||
if Version(vllm_version) < Version("0.10.1"):
|
if Version(vllm_version) < Version("0.10.1"):
|
||||||
vllm_location = spec.origin
|
vllm_location = spec.origin
|
||||||
if vllm_location is None:
|
if vllm_location is None:
|
||||||
|
|
|
||||||
|
|
@ -136,11 +136,15 @@ def calculate_settings(
|
||||||
|
|
||||||
|
|
||||||
HAS_CUDA_STREAM = False
|
HAS_CUDA_STREAM = False
|
||||||
import bitsandbytes as bnb
|
try:
|
||||||
|
import bitsandbytes as bnb
|
||||||
|
|
||||||
# https://github.com/bitsandbytes-foundation/bitsandbytes/pull/1330/files
|
# https://github.com/bitsandbytes-foundation/bitsandbytes/pull/1330/files
|
||||||
HAS_CUDA_STREAM = Version(bnb.__version__) > Version("0.43.3")
|
HAS_CUDA_STREAM = Version(bnb.__version__) > Version("0.43.3")
|
||||||
get_ptr = bnb.functional.get_ptr
|
get_ptr = bnb.functional.get_ptr
|
||||||
|
except Exception:
|
||||||
|
bnb = None
|
||||||
|
get_ptr = None
|
||||||
|
|
||||||
if DEVICE_TYPE == "xpu":
|
if DEVICE_TYPE == "xpu":
|
||||||
HAS_XPU_STREAM = True
|
HAS_XPU_STREAM = True
|
||||||
|
|
@ -236,21 +240,32 @@ else:
|
||||||
WEIGHT_BUFFERS = []
|
WEIGHT_BUFFERS = []
|
||||||
ABSMAX_BUFFERS = []
|
ABSMAX_BUFFERS = []
|
||||||
|
|
||||||
# Bitsandbytes operations
|
# Bitsandbytes operations (optional)
|
||||||
ctypes_c_int = ctypes.c_int
|
ctypes_c_int = ctypes.c_int
|
||||||
ctypes_c_int32 = ctypes.c_int32
|
ctypes_c_int32 = ctypes.c_int32
|
||||||
cdequantize_blockwise_fp32 = bnb.functional.lib.cdequantize_blockwise_fp32
|
|
||||||
cdequantize_blockwise_fp16_nf4 = bnb.functional.lib.cdequantize_blockwise_fp16_nf4
|
|
||||||
cdequantize_blockwise_bf16_nf4 = bnb.functional.lib.cdequantize_blockwise_bf16_nf4
|
|
||||||
|
|
||||||
if DEVICE_TYPE == "xpu":
|
if bnb is not None:
|
||||||
# https://github.com/bitsandbytes-foundation/bitsandbytes/blob/c3b8de268fdb55a88f92feada23fc811a1e6877a/bitsandbytes/backends/xpu/ops.py#L115
|
cdequantize_blockwise_fp32 = bnb.functional.lib.cdequantize_blockwise_fp32
|
||||||
# for xpu, inference gemv using above link
|
cdequantize_blockwise_fp16_nf4 = bnb.functional.lib.cdequantize_blockwise_fp16_nf4
|
||||||
cgemm_4bit_inference_naive_fp16 = bnb.functional.lib.cgemv_4bit_inference_fp16
|
cdequantize_blockwise_bf16_nf4 = bnb.functional.lib.cdequantize_blockwise_bf16_nf4
|
||||||
cgemm_4bit_inference_naive_bf16 = bnb.functional.lib.cgemv_4bit_inference_bf16
|
|
||||||
|
if DEVICE_TYPE == "xpu":
|
||||||
|
# https://github.com/bitsandbytes-foundation/bitsandbytes/blob/c3b8de268fdb55a88f92feada23fc811a1e6877a/bitsandbytes/backends/xpu/ops.py#L115
|
||||||
|
cgemm_4bit_inference_naive_fp16 = bnb.functional.lib.cgemv_4bit_inference_fp16
|
||||||
|
cgemm_4bit_inference_naive_bf16 = bnb.functional.lib.cgemv_4bit_inference_bf16
|
||||||
|
else:
|
||||||
|
cgemm_4bit_inference_naive_fp16 = (
|
||||||
|
bnb.functional.lib.cgemm_4bit_inference_naive_fp16
|
||||||
|
)
|
||||||
|
cgemm_4bit_inference_naive_bf16 = (
|
||||||
|
bnb.functional.lib.cgemm_4bit_inference_naive_bf16
|
||||||
|
)
|
||||||
else:
|
else:
|
||||||
cgemm_4bit_inference_naive_fp16 = bnb.functional.lib.cgemm_4bit_inference_naive_fp16
|
cdequantize_blockwise_fp32 = None
|
||||||
cgemm_4bit_inference_naive_bf16 = bnb.functional.lib.cgemm_4bit_inference_naive_bf16
|
cdequantize_blockwise_fp16_nf4 = None
|
||||||
|
cdequantize_blockwise_bf16_nf4 = None
|
||||||
|
cgemm_4bit_inference_naive_fp16 = None
|
||||||
|
cgemm_4bit_inference_naive_bf16 = None
|
||||||
|
|
||||||
|
|
||||||
torch_device_stream = (
|
torch_device_stream = (
|
||||||
|
|
|
||||||
|
|
@ -2540,7 +2540,14 @@ def patch_tokenizer(model, tokenizer):
|
||||||
|
|
||||||
|
|
||||||
def patch_fast_lora():
|
def patch_fast_lora():
|
||||||
import peft.tuners.lora.bnb
|
try:
|
||||||
|
import peft.tuners.lora.bnb
|
||||||
|
except Exception as e:
|
||||||
|
print(
|
||||||
|
"Unsloth: bitsandbytes/peft bnb not available - skipping 4bit LoRA patch.",
|
||||||
|
repr(e),
|
||||||
|
)
|
||||||
|
return
|
||||||
from ..kernels.fast_lora import fast_lora_forward
|
from ..kernels.fast_lora import fast_lora_forward
|
||||||
|
|
||||||
peft.tuners.lora.bnb.Linear4bit.forward = fast_lora_forward
|
peft.tuners.lora.bnb.Linear4bit.forward = fast_lora_forward
|
||||||
|
|
|
||||||
|
|
@ -30,8 +30,18 @@ from .llama import (
|
||||||
LlamaLinearScalingRotaryEmbedding,
|
LlamaLinearScalingRotaryEmbedding,
|
||||||
)
|
)
|
||||||
from .mistral import *
|
from .mistral import *
|
||||||
from bitsandbytes.nn import Linear4bit as Bnb_Linear4bit
|
|
||||||
from peft.tuners.lora import Linear4bit as Peft_Linear4bit
|
try:
|
||||||
|
from bitsandbytes.nn import Linear4bit as Bnb_Linear4bit
|
||||||
|
except Exception:
|
||||||
|
Bnb_Linear4bit = None
|
||||||
|
|
||||||
|
try:
|
||||||
|
from peft.tuners.lora import Linear4bit as Peft_Linear4bit
|
||||||
|
except Exception:
|
||||||
|
Peft_Linear4bit = None
|
||||||
|
|
||||||
|
_BNB_LINEAR_TYPES = tuple(t for t in (Bnb_Linear4bit, Peft_Linear4bit) if t is not None)
|
||||||
|
|
||||||
try:
|
try:
|
||||||
from transformers.models.granite.modeling_granite import (
|
from transformers.models.granite.modeling_granite import (
|
||||||
|
|
@ -599,7 +609,7 @@ class FastGraniteModel(FastLlamaModel):
|
||||||
correct_dtype = lm_head.weight.dtype
|
correct_dtype = lm_head.weight.dtype
|
||||||
|
|
||||||
for name, module in model.named_modules():
|
for name, module in model.named_modules():
|
||||||
if isinstance(module, (Bnb_Linear4bit, Peft_Linear4bit)):
|
if _BNB_LINEAR_TYPES and isinstance(module, _BNB_LINEAR_TYPES):
|
||||||
weight = module.weight
|
weight = module.weight
|
||||||
quant_state = weight.quant_state
|
quant_state = weight.quant_state
|
||||||
|
|
||||||
|
|
|
||||||
|
|
@ -3208,6 +3208,18 @@ class FastLlamaModel:
|
||||||
if not SUPPORTS_RSLORA:
|
if not SUPPORTS_RSLORA:
|
||||||
del arguments["use_rslora"]
|
del arguments["use_rslora"]
|
||||||
|
|
||||||
|
# PEFT API compatibility: only pass kwargs supported by the installed peft version.
|
||||||
|
try:
|
||||||
|
import inspect as _inspect
|
||||||
|
|
||||||
|
if (
|
||||||
|
"ensure_weight_tying"
|
||||||
|
not in _inspect.signature(LoraConfig.__init__).parameters
|
||||||
|
):
|
||||||
|
arguments.pop("ensure_weight_tying", None)
|
||||||
|
except Exception:
|
||||||
|
arguments.pop("ensure_weight_tying", None)
|
||||||
|
|
||||||
_saved_temp_tokenizer = model._saved_temp_tokenizer
|
_saved_temp_tokenizer = model._saved_temp_tokenizer
|
||||||
|
|
||||||
lora_config = LoraConfig(**arguments)
|
lora_config = LoraConfig(**arguments)
|
||||||
|
|
|
||||||
|
|
@ -1889,6 +1889,13 @@ def grpo_trainer_compute_loss(function_name, function):
|
||||||
if x.shape[1] == 1: # when importance_sampling_level == "sequence"
|
if x.shape[1] == 1: # when importance_sampling_level == "sequence"
|
||||||
return x.mean()
|
return x.mean()
|
||||||
else:
|
else:
|
||||||
|
# Align mask/coef lengths when left-padding adds extra tokens.
|
||||||
|
if x.shape[1] != completion_mask.shape[1]:
|
||||||
|
min_len = min(x.shape[1], completion_mask.shape[1])
|
||||||
|
x = x[:, -min_len:]
|
||||||
|
cm = completion_mask[:, -min_len:]
|
||||||
|
denom = cm.sum().clamp(min = 1.0)
|
||||||
|
return (x * cm).sum() / denom
|
||||||
return (x * completion_mask).sum() / completion_token_count
|
return (x * completion_mask).sum() / completion_token_count
|
||||||
|
|
||||||
if advantages.dim() == 1:
|
if advantages.dim() == 1:
|
||||||
|
|
|
||||||
|
|
@ -33,8 +33,16 @@ except ImportError:
|
||||||
|
|
||||||
IS_WINDOWS = sys.platform == "win32"
|
IS_WINDOWS = sys.platform == "win32"
|
||||||
LLAMA_CPP_DEFAULT_DIR = "llama.cpp"
|
LLAMA_CPP_DEFAULT_DIR = "llama.cpp"
|
||||||
from bitsandbytes.nn import Linear4bit as Bnb_Linear4bit
|
|
||||||
from peft.tuners.lora import Linear4bit as Peft_Linear4bit
|
try:
|
||||||
|
from bitsandbytes.nn import Linear4bit as Bnb_Linear4bit
|
||||||
|
except Exception:
|
||||||
|
Bnb_Linear4bit = None
|
||||||
|
|
||||||
|
try:
|
||||||
|
from peft.tuners.lora import Linear4bit as Peft_Linear4bit
|
||||||
|
except Exception:
|
||||||
|
Peft_Linear4bit = None
|
||||||
from peft.tuners.lora import Linear as Peft_Linear
|
from peft.tuners.lora import Linear as Peft_Linear
|
||||||
from typing import Optional, Callable, Union, List
|
from typing import Optional, Callable, Union, List
|
||||||
import sys
|
import sys
|
||||||
|
|
@ -68,6 +76,10 @@ except:
|
||||||
from pathlib import Path
|
from pathlib import Path
|
||||||
from peft import PeftModelForCausalLM, PeftModel
|
from peft import PeftModelForCausalLM, PeftModel
|
||||||
|
|
||||||
|
_MERGE_LORA_LINEAR_TYPES = tuple(
|
||||||
|
t for t in (Bnb_Linear4bit, Peft_Linear4bit, Peft_Linear) if t is not None
|
||||||
|
)
|
||||||
|
|
||||||
__all__ = [
|
__all__ = [
|
||||||
"print_quantization_methods",
|
"print_quantization_methods",
|
||||||
"unsloth_save_model",
|
"unsloth_save_model",
|
||||||
|
|
@ -381,7 +393,7 @@ def _free_cached_model(model):
|
||||||
|
|
||||||
def _merge_lora(layer, name):
|
def _merge_lora(layer, name):
|
||||||
bias = getattr(layer, "bias", None)
|
bias = getattr(layer, "bias", None)
|
||||||
if isinstance(layer, (Bnb_Linear4bit, Peft_Linear4bit, Peft_Linear)):
|
if _MERGE_LORA_LINEAR_TYPES and isinstance(layer, _MERGE_LORA_LINEAR_TYPES):
|
||||||
# Is LoRA so we need to merge!
|
# Is LoRA so we need to merge!
|
||||||
W, quant_state, A, B, s, bias = get_lora_parameters_bias(layer)
|
W, quant_state, A, B, s, bias = get_lora_parameters_bias(layer)
|
||||||
if quant_state is not None:
|
if quant_state is not None:
|
||||||
|
|
|
||||||
Loading…
Add table
Add a link
Reference in a new issue