From 6c68713fc2cf3762b21973d37b86115685805288 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 24 Sep 2024 02:49:48 -0700 Subject: [PATCH 01/50] Layernorm --- unsloth/kernels/__init__.py | 1 + unsloth/kernels/layernorm.py | 160 +++++++++++++++++++++++++++++++++++ 2 files changed, 161 insertions(+) create mode 100644 unsloth/kernels/layernorm.py diff --git a/unsloth/kernels/__init__.py b/unsloth/kernels/__init__.py index cd1d90f262..2fb9d11f99 100644 --- a/unsloth/kernels/__init__.py +++ b/unsloth/kernels/__init__.py @@ -14,6 +14,7 @@ from .cross_entropy_loss import fast_cross_entropy_loss from .rms_layernorm import fast_rms_layernorm +from .layernorm import fast_layernorm from .rope_embedding import fast_rope_embedding, inplace_rope_embedding from .swiglu import swiglu_fg_kernel, swiglu_DWf_DW_dfg_kernel from .geglu import ( diff --git a/unsloth/kernels/layernorm.py b/unsloth/kernels/layernorm.py new file mode 100644 index 0000000000..c65b5412ff --- /dev/null +++ b/unsloth/kernels/layernorm.py @@ -0,0 +1,160 @@ +# Copyright 2023-present Daniel Han-Chen & the Unsloth team. All rights reserved. +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +import triton +import triton.language as tl +import torch +from .utils import calculate_settings + + +@triton.jit +def layernorm_forward( + Y, Y_row_stride, + X, X_row_stride, + W, + b, + r, + mu, + n_cols, eps, + BLOCK_SIZE : tl.constexpr +): + row_idx = tl.program_id(0) + col_offsets = tl.arange(0, BLOCK_SIZE) + mask = col_offsets < n_cols + + Y += row_idx * Y_row_stride + X += row_idx * X_row_stride + r += row_idx + mu += row_idx + + # According to https://pytorch.org/torchtune/stable/_modules/torchtune/modules/layer_norm.html#Fp32LayerNorm, all modules + # are in float32! + X_row = tl.load(X + col_offsets, mask = mask, other = 0).to(tl.float32) + W_row = tl.load(W + col_offsets, mask = mask, other = 0).to(tl.float32) + b_row = tl.load(b + col_offsets, mask = mask, other = 0).to(tl.float32) + + mean_X = tl.sum(X_row, axis = 0) / n_cols + XX = X_row - mean_X + row_var = tl.sum(XX * XX, axis = 0) / n_cols + inv_var = tl.math.rsqrt(row_var + eps) + tl.store (r, inv_var) + tl.store (mu, mean_X) + output = (XX * inv_var) * W_row + b_row + tl.store(Y + col_offsets, output, mask = mask) +pass + + +@triton.jit +def layernorm_backward( + dY, dY_row_stride, + X, X_row_stride, + W, + b, + r, + mu, + n_cols, eps, + BLOCK_SIZE : tl.constexpr +): + # Approximately follows https://github.com/karpathy/llm.c/blob/master/doc/layernorm/layernorm.md + row_idx = tl.program_id(0) + col_offsets = tl.arange(0, BLOCK_SIZE) + mask = col_offsets < n_cols + + dY += row_idx * dY_row_stride + X += row_idx * X_row_stride + r += row_idx + mu += row_idx + + # According to https://pytorch.org/torchtune/stable/_modules/torchtune/modules/layer_norm.html#Fp32LayerNorm, all modules + # are in float32! + dY_row = tl.load(dY + col_offsets, mask = mask, other = 0).to(tl.float32) + X_row = tl.load(X + col_offsets, mask = mask, other = 0).to(tl.float32) + W_row = tl.load(W + col_offsets, mask = mask, other = 0).to(tl.float32) + b_row = tl.load(b + col_offsets, mask = mask, other = 0).to(tl.float32) + + inv_var = tl.load(r) .to(tl.float32) + mean = tl.load(mu).to(tl.float32) + normed = (X_row - mean) * inv_var + dY_W = dY_row * W_row + dX_row = dY_W - tl.sum(dY_W, axis = 0) / n_cols - normed * tl.sum(dY_W * normed, axis = 0) / n_cols + dX_row = dX_row * inv_var + tl.store(dY + col_offsets, dX_row, mask = mask) +pass + + +class Fast_Layernorm(torch.autograd.Function): + @staticmethod + def forward(ctx, X, W, b, eps): + shape = X.shape + dim = shape[-1] + X = X.view(-1, dim) + n_rows, n_cols = X.shape + BLOCK_SIZE, num_warps = calculate_settings(n_cols) + + Y = torch.empty((n_rows, n_cols), dtype = X.dtype, device = "cuda:0") + r = torch.empty(n_rows, dtype = torch.float32, device = "cuda:0") + mu = torch.empty(n_rows, dtype = torch.float32, device = "cuda:0") + + layernorm_forward[(n_rows,)]( + Y, Y.stride(0), + X, X.stride(0), + W, + b, + r, + mu, + n_cols, eps, + BLOCK_SIZE = BLOCK_SIZE, + num_warps = num_warps, + ) + ctx.eps = eps + ctx.BLOCK_SIZE = BLOCK_SIZE + ctx.num_warps = num_warps + ctx.save_for_backward(X, W, r, mu) + return Y.view(*shape) + pass + + @staticmethod + def backward(ctx, dY): + shape = dY.shape + dim = shape[-1] + dY = dY.view(-1, dim) + X, W, r, mu = ctx.saved_tensors + n_rows, n_cols = dY.shape + + layernorm_backward[(n_rows,)]( + dY, dY.stride(0), + X, X .stride(0), + W, + b, + r, + mu, + n_cols, ctx.eps, + BLOCK_SIZE = ctx.BLOCK_SIZE, + num_warps = ctx.num_warps, + ) + dX = dY.view(*shape) + return dX, None, None, None, None + pass +pass + + +def fast_layernorm(layernorm, X): + W = layernorm.weight + bias = layernorm.bias + eps = layernorm.variance_epsilon if \ + hasattr(layernorm, "variance_epsilon") \ + else layernorm.eps + out = Fast_Layernorm.apply(X, W, bias, eps) + return out +pass From 1e0f8c7304c5395e3e5b831f258fadf3d590be7d Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 24 Sep 2024 16:29:54 -0700 Subject: [PATCH 02/50] Update layernorm.py --- unsloth/kernels/layernorm.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/unsloth/kernels/layernorm.py b/unsloth/kernels/layernorm.py index c65b5412ff..cfafa4a11b 100644 --- a/unsloth/kernels/layernorm.py +++ b/unsloth/kernels/layernorm.py @@ -120,7 +120,7 @@ class Fast_Layernorm(torch.autograd.Function): ctx.eps = eps ctx.BLOCK_SIZE = BLOCK_SIZE ctx.num_warps = num_warps - ctx.save_for_backward(X, W, r, mu) + ctx.save_for_backward(X, W, b, r, mu) return Y.view(*shape) pass @@ -129,7 +129,7 @@ class Fast_Layernorm(torch.autograd.Function): shape = dY.shape dim = shape[-1] dY = dY.view(-1, dim) - X, W, r, mu = ctx.saved_tensors + X, W, b, r, mu = ctx.saved_tensors n_rows, n_cols = dY.shape layernorm_backward[(n_rows,)]( From b48efb6a0be8130bd37d569a0c259dab0ebee129 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 24 Sep 2024 16:44:33 -0700 Subject: [PATCH 03/50] Update layernorm.py --- unsloth/kernels/layernorm.py | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/unsloth/kernels/layernorm.py b/unsloth/kernels/layernorm.py index cfafa4a11b..29c77462f5 100644 --- a/unsloth/kernels/layernorm.py +++ b/unsloth/kernels/layernorm.py @@ -41,8 +41,8 @@ def layernorm_forward( # According to https://pytorch.org/torchtune/stable/_modules/torchtune/modules/layer_norm.html#Fp32LayerNorm, all modules # are in float32! X_row = tl.load(X + col_offsets, mask = mask, other = 0).to(tl.float32) - W_row = tl.load(W + col_offsets, mask = mask, other = 0).to(tl.float32) - b_row = tl.load(b + col_offsets, mask = mask, other = 0).to(tl.float32) + W_row = tl.load(W + col_offsets, mask = mask, other = 0)#.to(tl.float32) + b_row = tl.load(b + col_offsets, mask = mask, other = 0)#.to(tl.float32) mean_X = tl.sum(X_row, axis = 0) / n_cols XX = X_row - mean_X @@ -50,7 +50,7 @@ def layernorm_forward( inv_var = tl.math.rsqrt(row_var + eps) tl.store (r, inv_var) tl.store (mu, mean_X) - output = (XX * inv_var) * W_row + b_row + output = (XX * inv_var).to(W_row.dtype) * W_row + b_row tl.store(Y + col_offsets, output, mask = mask) pass From 613116e19b05ea2680598309bdd5678603e695bd Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 24 Sep 2024 16:45:50 -0700 Subject: [PATCH 04/50] Update layernorm.py --- unsloth/kernels/layernorm.py | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/unsloth/kernels/layernorm.py b/unsloth/kernels/layernorm.py index 29c77462f5..cfafa4a11b 100644 --- a/unsloth/kernels/layernorm.py +++ b/unsloth/kernels/layernorm.py @@ -41,8 +41,8 @@ def layernorm_forward( # According to https://pytorch.org/torchtune/stable/_modules/torchtune/modules/layer_norm.html#Fp32LayerNorm, all modules # are in float32! X_row = tl.load(X + col_offsets, mask = mask, other = 0).to(tl.float32) - W_row = tl.load(W + col_offsets, mask = mask, other = 0)#.to(tl.float32) - b_row = tl.load(b + col_offsets, mask = mask, other = 0)#.to(tl.float32) + W_row = tl.load(W + col_offsets, mask = mask, other = 0).to(tl.float32) + b_row = tl.load(b + col_offsets, mask = mask, other = 0).to(tl.float32) mean_X = tl.sum(X_row, axis = 0) / n_cols XX = X_row - mean_X @@ -50,7 +50,7 @@ def layernorm_forward( inv_var = tl.math.rsqrt(row_var + eps) tl.store (r, inv_var) tl.store (mu, mean_X) - output = (XX * inv_var).to(W_row.dtype) * W_row + b_row + output = (XX * inv_var) * W_row + b_row tl.store(Y + col_offsets, output, mask = mask) pass From 987283e017a8db47d65749d8de2d3c913a79744a Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 24 Sep 2024 16:54:52 -0700 Subject: [PATCH 05/50] Update layernorm.py --- unsloth/kernels/layernorm.py | 45 ++++++++++++++++++++++++++++++++++++ 1 file changed, 45 insertions(+) diff --git a/unsloth/kernels/layernorm.py b/unsloth/kernels/layernorm.py index cfafa4a11b..058ec6a168 100644 --- a/unsloth/kernels/layernorm.py +++ b/unsloth/kernels/layernorm.py @@ -158,3 +158,48 @@ def fast_layernorm(layernorm, X): out = Fast_Layernorm.apply(X, W, bias, eps) return out pass + + +def test_layernorm( + dim = 1024, eps = 1e-5, dtype = torch.float16, + bsz = 21, random_state = 3407, seqlen = 3341, +): + from torch.nn import LayerNorm + layernorm = LayerNorm((dim,), eps = eps, device = "cuda", dtype = dtype) + torch.cuda.manual_seed(random_state) + torch.manual_seed(random_state) + torch.nn.init.uniform_(layernorm.weight) + torch.nn.init.uniform_(layernorm.bias) + X = torch.randn((bsz, seqlen, dim), dtype = dtype, device = "cuda") + XX = X.clone() + X .requires_grad_(True) + XX.requires_grad_(True) + Y = layernorm(X) + YY = torch.randn((bsz, seqlen, dim), dtype = dtype, device = "cuda", requires_grad = True) + Y.backward(YY) + correct_grad = X.grad.clone() + from unsloth.kernels import fast_layernorm + Y = fast_layernorm(layernorm, XX) + Y.backward(YY) + assert(torch.dist(correct_grad, XX.grad).item() <= 0.1) +pass + + +def testing_suite_layernorm(): + for dim in [512, 1024, 2048]: + for dtype in [torch.float16, torch.bfloat16]: + for seqlen in [3341, 2048, 349]: + for random_state in [3407, 42]: + test_layernorm( + dim = dim, + eps = 1e-5, + dtype = dtype, + bsz = 21, + random_state = random_state, + seqlen = seqlen, + ) + pass + pass + pass + pass +pass From dbcf225c32a3f934ad415495352bc71e5588e839 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 24 Sep 2024 17:00:23 -0700 Subject: [PATCH 06/50] Update layernorm.py --- unsloth/kernels/layernorm.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/unsloth/kernels/layernorm.py b/unsloth/kernels/layernorm.py index 058ec6a168..7939613804 100644 --- a/unsloth/kernels/layernorm.py +++ b/unsloth/kernels/layernorm.py @@ -178,7 +178,7 @@ def test_layernorm( YY = torch.randn((bsz, seqlen, dim), dtype = dtype, device = "cuda", requires_grad = True) Y.backward(YY) correct_grad = X.grad.clone() - from unsloth.kernels import fast_layernorm + # from unsloth.kernels import fast_layernorm Y = fast_layernorm(layernorm, XX) Y.backward(YY) assert(torch.dist(correct_grad, XX.grad).item() <= 0.1) From 616e97ef165215527aa448b08907373982269fef Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 24 Sep 2024 17:03:19 -0700 Subject: [PATCH 07/50] Update layernorm.py --- unsloth/kernels/layernorm.py | 22 ++++++++++++---------- 1 file changed, 12 insertions(+), 10 deletions(-) diff --git a/unsloth/kernels/layernorm.py b/unsloth/kernels/layernorm.py index 7939613804..c0ff4d7440 100644 --- a/unsloth/kernels/layernorm.py +++ b/unsloth/kernels/layernorm.py @@ -188,16 +188,18 @@ pass def testing_suite_layernorm(): for dim in [512, 1024, 2048]: for dtype in [torch.float16, torch.bfloat16]: - for seqlen in [3341, 2048, 349]: - for random_state in [3407, 42]: - test_layernorm( - dim = dim, - eps = 1e-5, - dtype = dtype, - bsz = 21, - random_state = random_state, - seqlen = seqlen, - ) + with torch.autocast(device_type = "cuda", dtype = dtype): + for seqlen in [3341, 2048, 349]: + for random_state in [3407, 42]: + test_layernorm( + dim = dim, + eps = 1e-5, + dtype = dtype, + bsz = 21, + random_state = random_state, + seqlen = seqlen, + ) + pass pass pass pass From a84b1dffc4a166c718048378fd2d35b76ab03635 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 24 Sep 2024 17:22:20 -0700 Subject: [PATCH 08/50] Patch layernorm --- unsloth/kernels/__init__.py | 6 +++++- unsloth/kernels/layernorm.py | 23 +++++++++++++++++++++++ 2 files changed, 28 insertions(+), 1 deletion(-) diff --git a/unsloth/kernels/__init__.py b/unsloth/kernels/__init__.py index 2fb9d11f99..841d0ce0f0 100644 --- a/unsloth/kernels/__init__.py +++ b/unsloth/kernels/__init__.py @@ -14,7 +14,11 @@ from .cross_entropy_loss import fast_cross_entropy_loss from .rms_layernorm import fast_rms_layernorm -from .layernorm import fast_layernorm +from .layernorm import ( + fast_layernorm, + patch_layernorm, + unpatch_layernorm, +) from .rope_embedding import fast_rope_embedding, inplace_rope_embedding from .swiglu import swiglu_fg_kernel, swiglu_DWf_DW_dfg_kernel from .geglu import ( diff --git a/unsloth/kernels/layernorm.py b/unsloth/kernels/layernorm.py index c0ff4d7440..0d456109ea 100644 --- a/unsloth/kernels/layernorm.py +++ b/unsloth/kernels/layernorm.py @@ -150,6 +150,7 @@ pass def fast_layernorm(layernorm, X): + assert(layernorm.elementwise_affine is True) W = layernorm.weight bias = layernorm.bias eps = layernorm.variance_epsilon if \ @@ -160,6 +161,28 @@ def fast_layernorm(layernorm, X): pass +from torch.nn import LayerNorm +class Fast_LayerNorm_Module(LayerNorm): + def forward(self, X): + return fast_layernorm(self, X) + pass +pass + + +def patch_layernorm(): + import torch.nn + torch.nn.LayerNorm = Fast_LayerNorm_Module + return +pass + + +def unpatch_layernorm(): + import torch.nn + torch.nn.LayerNorm = LayerNorm + return +pass + + def test_layernorm( dim = 1024, eps = 1e-5, dtype = torch.float16, bsz = 21, random_state = 3407, seqlen = 3341, From 03da2fe3ba0a6343814d8fe4213f2518260c3387 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 24 Sep 2024 17:24:39 -0700 Subject: [PATCH 09/50] Update layernorm.py --- unsloth/kernels/layernorm.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/unsloth/kernels/layernorm.py b/unsloth/kernels/layernorm.py index 0d456109ea..0546484f17 100644 --- a/unsloth/kernels/layernorm.py +++ b/unsloth/kernels/layernorm.py @@ -162,7 +162,7 @@ pass from torch.nn import LayerNorm -class Fast_LayerNorm_Module(LayerNorm): +class Unsloth_LayerNorm(LayerNorm): def forward(self, X): return fast_layernorm(self, X) pass @@ -171,7 +171,7 @@ pass def patch_layernorm(): import torch.nn - torch.nn.LayerNorm = Fast_LayerNorm_Module + torch.nn.LayerNorm = Unsloth_LayerNorm return pass From da6de6dcb84f99180ed191182eaedbf5c9c4f968 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 24 Sep 2024 22:50:33 -0700 Subject: [PATCH 10/50] RMS Layernorm --- unsloth/kernels/__init__.py | 6 +++++- unsloth/kernels/rms_layernorm.py | 22 ++++++++++++++++++++++ 2 files changed, 27 insertions(+), 1 deletion(-) diff --git a/unsloth/kernels/__init__.py b/unsloth/kernels/__init__.py index 841d0ce0f0..606adf80f7 100644 --- a/unsloth/kernels/__init__.py +++ b/unsloth/kernels/__init__.py @@ -13,7 +13,11 @@ # limitations under the License. from .cross_entropy_loss import fast_cross_entropy_loss -from .rms_layernorm import fast_rms_layernorm +from .rms_layernorm import ( + fast_rms_layernorm, + patch_rms_layernorm, + unpatch_rms_layernorm, +) from .layernorm import ( fast_layernorm, patch_layernorm, diff --git a/unsloth/kernels/rms_layernorm.py b/unsloth/kernels/rms_layernorm.py index ac5beb5ab1..75d491cb1a 100644 --- a/unsloth/kernels/rms_layernorm.py +++ b/unsloth/kernels/rms_layernorm.py @@ -192,3 +192,25 @@ def fast_rms_layernorm(layernorm, X, gemma = False): out = Fast_RMS_Layernorm.apply(X, W, eps, gemma) return out pass + + +from transformers.models.llama.modeling_llama import LlamaRMSNorm +class Unsloth_LlamaRMSNorm(LlamaRMSNorm): + def forward(self, X): + return fast_rms_layernorm(self, X, gemma = False) + pass +pass + + +def patch_rms_layernorm(): + import transformers.models.llama.modeling_llama + transformers.models.llama.modeling_llama.LlamaRMSNorm = Unsloth_LlamaRMSNorm + return +pass + + +def unpatch_rms_layernorm(): + import transformers.models.llama.modeling_llama + transformers.models.llama.modeling_llama.LlamaRMSNorm = LlamaRMSNorm + return +pass From 555082bf8484b39cbf04e2eabe2fc88c6c436b6e Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 24 Sep 2024 23:13:51 -0700 Subject: [PATCH 11/50] Update rms_layernorm.py --- unsloth/kernels/rms_layernorm.py | 46 ++++++++++++++++++++++++++++++++ 1 file changed, 46 insertions(+) diff --git a/unsloth/kernels/rms_layernorm.py b/unsloth/kernels/rms_layernorm.py index 75d491cb1a..43924d0ab3 100644 --- a/unsloth/kernels/rms_layernorm.py +++ b/unsloth/kernels/rms_layernorm.py @@ -214,3 +214,49 @@ def unpatch_rms_layernorm(): transformers.models.llama.modeling_llama.LlamaRMSNorm = LlamaRMSNorm return pass + + +def test_rms_layernorm( + dim = 1024, eps = 1e-5, dtype = torch.float16, + bsz = 21, random_state = 3407, seqlen = 3341, +): + from transformers.models.llama.modeling_llama import LlamaRMSNorm + layernorm = LlamaRMSNorm((dim,), eps = eps).to("cuda") + torch.cuda.manual_seed(random_state) + torch.manual_seed(random_state) + torch.nn.init.uniform_(layernorm.weight) + X = torch.randn((bsz, seqlen, dim), dtype = dtype, device = "cuda") + XX = X.clone() + X .requires_grad_(True) + XX.requires_grad_(True) + Y = layernorm(X) + YY = torch.randn((bsz, seqlen, dim), dtype = dtype, device = "cuda", requires_grad = True) + Y.backward(YY) + correct_grad = X.grad.clone() + # from unsloth.kernels import fast_rms_layernorm + Y = fast_rms_layernorm(layernorm, XX) + Y.backward(YY) + assert(torch.amax(correct_grad - XX.grad).item() <= 0.05) +pass + + +def testing_suite_layernorm(): + for dim in [512, 1024, 2048]: + for dtype in [torch.float16, torch.bfloat16]: + with torch.autocast(device_type = "cuda", dtype = dtype): + for seqlen in [3341, 2048, 349]: + for random_state in [3407, 42]: + test_rms_layernorm( + dim = dim, + eps = 1e-5, + dtype = dtype, + bsz = 21, + random_state = random_state, + seqlen = seqlen, + ) + pass + pass + pass + pass + pass +pass From f9516884f63162975e312456a9f7d9f8400c6658 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 24 Sep 2024 23:49:57 -0700 Subject: [PATCH 12/50] Causal LM --- unsloth/kernels/__init__.py | 6 +- unsloth/kernels/cross_entropy_loss.py | 79 +++++++++++++++++++++++++++ 2 files changed, 84 insertions(+), 1 deletion(-) diff --git a/unsloth/kernels/__init__.py b/unsloth/kernels/__init__.py index 606adf80f7..3e55332c80 100644 --- a/unsloth/kernels/__init__.py +++ b/unsloth/kernels/__init__.py @@ -12,7 +12,11 @@ # See the License for the specific language governing permissions and # limitations under the License. -from .cross_entropy_loss import fast_cross_entropy_loss +from .cross_entropy_loss import ( + fast_cross_entropy_loss, + patch_llama_for_causal_lm, + unpatch_llama_for_causal_lm, +) from .rms_layernorm import ( fast_rms_layernorm, patch_rms_layernorm, diff --git a/unsloth/kernels/cross_entropy_loss.py b/unsloth/kernels/cross_entropy_loss.py index 24e8002bec..74ee4ee66a 100644 --- a/unsloth/kernels/cross_entropy_loss.py +++ b/unsloth/kernels/cross_entropy_loss.py @@ -375,3 +375,82 @@ def fast_cross_entropy_loss( n_items = torch.count_nonzero(labels != -100) return loss.sum() / n_items pass + + +from transformers.models.llama.modeling_llama import LlamaForCausalLM +def patch_llama_for_causal_lm(): + import transformers.models.llama.modeling_llama + from transformers.models.llama.modeling_llama import ( + CausalLMOutputWithPast, + Optional, + Union, + Cache, + List, + Tuple, + ) + import inspect, re + function = inspect.getsource(transformers.models.llama.modeling_llama.LlamaForCausalLM.forward) + function = function.split("\n") + i = re.match(r"[ ]{1,}", function[0]).span(0)[1] + function = [x[i:] for x in function] + function = "\n".join(function) + function = function[function.find("def forward"):] + replacement = """ loss = None + logit_softcapping = getattr(self.config, "final_logit_softcapping", 0) + logit_scaling = getattr(self.config, "logit_scale", 0) + if labels is not None: + shift_logits = logits + if not hasattr(self, "extra_ignored_labels"): + # Fixes https://github.com/unslothai/unsloth/issues/10 + self.extra_ignored_labels = torch.full((self.max_seq_length, 1), -100, device = "cuda:0") + pass + + shift_labels = torch.hstack((labels[..., 1:], self.extra_ignored_labels[:labels.shape[0]])) + loss = fast_cross_entropy_loss( + logits = shift_logits, + labels = shift_labels, + logit_softcapping = logit_softcapping, + logit_scaling = logit_scaling, + ) + else: + if logit_scaling != 0: + if logits.requires_grad: + logits = logit_scaling * logits + else: + logits *= logit_scaling + pass + pass + if logit_softcapping != 0: + if logits.requires_grad: + logits = (1.0 / logit_softcapping) * logits + logits = torch.tanh(logits) + logits = logit_softcapping * logits + else: + logits *= (1.0 / logit_softcapping) + torch.tanh(logits, out = logits) + logits *= logit_softcapping + pass + pass + pass + """ + + function = \ + function[:function.find(" loss = None")] + \ + replacement + \ + function[ function.find(" if not return_dict"):] + function = function.replace("logits = logits.float()", "\n") + + patched_function = f"class Unsloth_LlamaForCausalLM(LlamaForCausalLM):\n"\ + f" {function}\n" + + exec(patched_function) + transformers.models.llama.modeling_llama.LlamaForCausalLM = Unsloth_LlamaForCausalLM + return +pass + + +def unpatch_llama_for_causal_lm(): + import transformers.models.llama.modeling_llama + transformers.models.llama.modeling_llama.LlamaForCausalLM = LlamaForCausalLM + return +pass From 26c4f88d8e111b1f4b6d1d05dc52a0d504d3fd61 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 24 Sep 2024 23:52:54 -0700 Subject: [PATCH 13/50] Update cross_entropy_loss.py --- 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 74ee4ee66a..59d0f3b694 100644 --- a/unsloth/kernels/cross_entropy_loss.py +++ b/unsloth/kernels/cross_entropy_loss.py @@ -443,6 +443,7 @@ def patch_llama_for_causal_lm(): patched_function = f"class Unsloth_LlamaForCausalLM(LlamaForCausalLM):\n"\ f" {function}\n" + print(patched_function) exec(patched_function) transformers.models.llama.modeling_llama.LlamaForCausalLM = Unsloth_LlamaForCausalLM return From a4659020467b878ea1be35ca35716347dafea255 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 24 Sep 2024 23:54:14 -0700 Subject: [PATCH 14/50] Update cross_entropy_loss.py --- unsloth/kernels/cross_entropy_loss.py | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/unsloth/kernels/cross_entropy_loss.py b/unsloth/kernels/cross_entropy_loss.py index 59d0f3b694..58104e4dfa 100644 --- a/unsloth/kernels/cross_entropy_loss.py +++ b/unsloth/kernels/cross_entropy_loss.py @@ -440,6 +440,10 @@ def patch_llama_for_causal_lm(): function[ function.find(" if not return_dict"):] function = function.replace("logits = logits.float()", "\n") + function = function.split("\n") + function = [" "*4 + x for x in function] + function = "\n".join(function) + patched_function = f"class Unsloth_LlamaForCausalLM(LlamaForCausalLM):\n"\ f" {function}\n" From cd42f4d772d11e1da374873399ebdb65fd42b6f2 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 24 Sep 2024 23:55:25 -0700 Subject: [PATCH 15/50] Update cross_entropy_loss.py --- unsloth/kernels/cross_entropy_loss.py | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/unsloth/kernels/cross_entropy_loss.py b/unsloth/kernels/cross_entropy_loss.py index 58104e4dfa..d06bf69913 100644 --- a/unsloth/kernels/cross_entropy_loss.py +++ b/unsloth/kernels/cross_entropy_loss.py @@ -440,10 +440,12 @@ def patch_llama_for_causal_lm(): function[ function.find(" if not return_dict"):] function = function.replace("logits = logits.float()", "\n") + # Missed spaces function = function.split("\n") - function = [" "*4 + x for x in function] + # Not the first one though! + function = [function[0]] + [" "*4 + x for x in function[1:]] function = "\n".join(function) - + patched_function = f"class Unsloth_LlamaForCausalLM(LlamaForCausalLM):\n"\ f" {function}\n" From f8a48090ec80d1e6ebbb43fb0cbb5d4ea8c18f07 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 24 Sep 2024 23:57:22 -0700 Subject: [PATCH 16/50] Update cross_entropy_loss.py --- unsloth/kernels/cross_entropy_loss.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/unsloth/kernels/cross_entropy_loss.py b/unsloth/kernels/cross_entropy_loss.py index d06bf69913..e791664a6f 100644 --- a/unsloth/kernels/cross_entropy_loss.py +++ b/unsloth/kernels/cross_entropy_loss.py @@ -437,7 +437,7 @@ def patch_llama_for_causal_lm(): function = \ function[:function.find(" loss = None")] + \ replacement + \ - function[ function.find(" if not return_dict"):] + function[ function.find("if not return_dict"):] function = function.replace("logits = logits.float()", "\n") # Missed spaces From 14936a94c15a47f6fb0e392a4f653575e7397b0e Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 24 Sep 2024 23:59:09 -0700 Subject: [PATCH 17/50] Update layernorm.py --- unsloth/kernels/layernorm.py | 1 + 1 file changed, 1 insertion(+) diff --git a/unsloth/kernels/layernorm.py b/unsloth/kernels/layernorm.py index 0546484f17..48ade6d5ec 100644 --- a/unsloth/kernels/layernorm.py +++ b/unsloth/kernels/layernorm.py @@ -1,4 +1,5 @@ # Copyright 2023-present Daniel Han-Chen & the Unsloth team. All rights reserved. +# Copyright 2024-present Andrej Karpathy & the llm.c team. All rights reserved. # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. From b30acca3b623ca6fe68d01e0d993ffad9596e3ed Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 00:00:51 -0700 Subject: [PATCH 18/50] Update cross_entropy_loss.py --- unsloth/kernels/cross_entropy_loss.py | 18 +++++++++--------- 1 file changed, 9 insertions(+), 9 deletions(-) diff --git a/unsloth/kernels/cross_entropy_loss.py b/unsloth/kernels/cross_entropy_loss.py index e791664a6f..e5ba8b55f1 100644 --- a/unsloth/kernels/cross_entropy_loss.py +++ b/unsloth/kernels/cross_entropy_loss.py @@ -377,17 +377,17 @@ def fast_cross_entropy_loss( pass -from transformers.models.llama.modeling_llama import LlamaForCausalLM +from transformers.models.llama.modeling_llama import ( + LlamaForCausalLM, + CausalLMOutputWithPast, + Optional, + Union, + Cache, + List, + Tuple, +) def patch_llama_for_causal_lm(): import transformers.models.llama.modeling_llama - from transformers.models.llama.modeling_llama import ( - CausalLMOutputWithPast, - Optional, - Union, - Cache, - List, - Tuple, - ) import inspect, re function = inspect.getsource(transformers.models.llama.modeling_llama.LlamaForCausalLM.forward) function = function.split("\n") From 336f3b9c24bafd376ca30dde4cd2f06f301b377a Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 00:01:56 -0700 Subject: [PATCH 19/50] Update cross_entropy_loss.py --- unsloth/kernels/cross_entropy_loss.py | 23 +++++++++++------------ 1 file changed, 11 insertions(+), 12 deletions(-) diff --git a/unsloth/kernels/cross_entropy_loss.py b/unsloth/kernels/cross_entropy_loss.py index e5ba8b55f1..a09f167db7 100644 --- a/unsloth/kernels/cross_entropy_loss.py +++ b/unsloth/kernels/cross_entropy_loss.py @@ -377,17 +377,17 @@ def fast_cross_entropy_loss( pass -from transformers.models.llama.modeling_llama import ( - LlamaForCausalLM, - CausalLMOutputWithPast, - Optional, - Union, - Cache, - List, - Tuple, -) +from transformers.models.llama.modeling_llama import LlamaForCausalLM def patch_llama_for_causal_lm(): import transformers.models.llama.modeling_llama + from transformers.models.llama.modeling_llama import ( + CausalLMOutputWithPast, + Optional, + Union, + Cache, + List, + Tuple, + ) import inspect, re function = inspect.getsource(transformers.models.llama.modeling_llama.LlamaForCausalLM.forward) function = function.split("\n") @@ -448,9 +448,8 @@ def patch_llama_for_causal_lm(): patched_function = f"class Unsloth_LlamaForCausalLM(LlamaForCausalLM):\n"\ f" {function}\n" - - print(patched_function) - exec(patched_function) + + exec(patched_function, globals()) transformers.models.llama.modeling_llama.LlamaForCausalLM = Unsloth_LlamaForCausalLM return pass From 3f811a5ef7942c764f882f72395ef07e7a6bed7f Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 00:04:58 -0700 Subject: [PATCH 20/50] Update cross_entropy_loss.py --- unsloth/kernels/cross_entropy_loss.py | 69 +++++++++++++-------------- 1 file changed, 34 insertions(+), 35 deletions(-) diff --git a/unsloth/kernels/cross_entropy_loss.py b/unsloth/kernels/cross_entropy_loss.py index a09f167db7..72bd3b92d6 100644 --- a/unsloth/kernels/cross_entropy_loss.py +++ b/unsloth/kernels/cross_entropy_loss.py @@ -377,25 +377,23 @@ def fast_cross_entropy_loss( pass -from transformers.models.llama.modeling_llama import LlamaForCausalLM -def patch_llama_for_causal_lm(): - import transformers.models.llama.modeling_llama - from transformers.models.llama.modeling_llama import ( - CausalLMOutputWithPast, - Optional, - Union, - Cache, - List, - Tuple, - ) - import inspect, re - function = inspect.getsource(transformers.models.llama.modeling_llama.LlamaForCausalLM.forward) - function = function.split("\n") - i = re.match(r"[ ]{1,}", function[0]).span(0)[1] - function = [x[i:] for x in function] - function = "\n".join(function) - function = function[function.find("def forward"):] - replacement = """ loss = None +from transformers.models.llama.modeling_llama import ( + LlamaForCausalLM, + CausalLMOutputWithPast, + Optional, + Union, + Cache, + List, + Tuple, +) +import inspect, re +function = inspect.getsource(transformers.models.llama.modeling_llama.LlamaForCausalLM.forward) +function = function.split("\n") +i = re.match(r"[ ]{1,}", function[0]).span(0)[1] +function = [x[i:] for x in function] +function = "\n".join(function) +function = function[function.find("def forward"):] +replacement = """ loss = None logit_softcapping = getattr(self.config, "final_logit_softcapping", 0) logit_scaling = getattr(self.config, "logit_scale", 0) if labels is not None: @@ -432,24 +430,25 @@ def patch_llama_for_causal_lm(): pass pass pass - """ +""" +function = \ + function[:function.find(" loss = None")] + \ + replacement + \ + function[ function.find("if not return_dict"):] +function = function.replace("logits = logits.float()", "\n") +# Missed spaces +function = function.split("\n") +# Not the first one though! +function = [function[0]] + [" "*4 + x for x in function[1:]] +function = "\n".join(function) +function = f"class Unsloth_LlamaForCausalLM(LlamaForCausalLM):\n"\ +f" {function}\n" +exec(function, globals()) +del function, replacement - function = \ - function[:function.find(" loss = None")] + \ - replacement + \ - function[ function.find("if not return_dict"):] - function = function.replace("logits = logits.float()", "\n") - # Missed spaces - function = function.split("\n") - # Not the first one though! - function = [function[0]] + [" "*4 + x for x in function[1:]] - function = "\n".join(function) - - patched_function = f"class Unsloth_LlamaForCausalLM(LlamaForCausalLM):\n"\ - f" {function}\n" - - exec(patched_function, globals()) +def patch_llama_for_causal_lm(): + import transformers.models.llama.modeling_llama transformers.models.llama.modeling_llama.LlamaForCausalLM = Unsloth_LlamaForCausalLM return pass From 06e06f922d2a19365b2cc6e2602993c15c1104dd Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 00:06:10 -0700 Subject: [PATCH 21/50] Update cross_entropy_loss.py --- unsloth/kernels/cross_entropy_loss.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/unsloth/kernels/cross_entropy_loss.py b/unsloth/kernels/cross_entropy_loss.py index 72bd3b92d6..1174ead7d1 100644 --- a/unsloth/kernels/cross_entropy_loss.py +++ b/unsloth/kernels/cross_entropy_loss.py @@ -387,7 +387,7 @@ from transformers.models.llama.modeling_llama import ( Tuple, ) import inspect, re -function = inspect.getsource(transformers.models.llama.modeling_llama.LlamaForCausalLM.forward) +function = inspect.getsource(LlamaForCausalLM.forward) function = function.split("\n") i = re.match(r"[ ]{1,}", function[0]).span(0)[1] function = [x[i:] for x in function] From b74a86a2bda201ac32748618237d3f93e49dbeca Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 00:07:49 -0700 Subject: [PATCH 22/50] Update cross_entropy_loss.py --- 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 1174ead7d1..49bdbd10ea 100644 --- a/unsloth/kernels/cross_entropy_loss.py +++ b/unsloth/kernels/cross_entropy_loss.py @@ -443,6 +443,7 @@ function = [function[0]] + [" "*4 + x for x in function[1:]] function = "\n".join(function) function = f"class Unsloth_LlamaForCausalLM(LlamaForCausalLM):\n"\ f" {function}\n" +print(function) exec(function, globals()) del function, replacement From 88fa0eb7f32e3372a648db0133842a7feabf25d5 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 00:09:19 -0700 Subject: [PATCH 23/50] Update cross_entropy_loss.py --- unsloth/kernels/cross_entropy_loss.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/unsloth/kernels/cross_entropy_loss.py b/unsloth/kernels/cross_entropy_loss.py index 49bdbd10ea..0034448be5 100644 --- a/unsloth/kernels/cross_entropy_loss.py +++ b/unsloth/kernels/cross_entropy_loss.py @@ -434,7 +434,7 @@ replacement = """ loss = None function = \ function[:function.find(" loss = None")] + \ replacement + \ - function[ function.find("if not return_dict"):] + function[ function.find(" if not return_dict"):] function = function.replace("logits = logits.float()", "\n") # Missed spaces function = function.split("\n") From 7f9ecd592f3e9add62c5a3dfd1b448128836caef Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 00:11:57 -0700 Subject: [PATCH 24/50] Update cross_entropy_loss.py --- 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 0034448be5..1fec5d7a85 100644 --- a/unsloth/kernels/cross_entropy_loss.py +++ b/unsloth/kernels/cross_entropy_loss.py @@ -443,9 +443,8 @@ function = [function[0]] + [" "*4 + x for x in function[1:]] function = "\n".join(function) function = f"class Unsloth_LlamaForCausalLM(LlamaForCausalLM):\n"\ f" {function}\n" -print(function) exec(function, globals()) -del function, replacement +del function, replacement, inspect, re def patch_llama_for_causal_lm(): From 020cbd2dd50204e46885cfa5b121805efedebb3b Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 00:22:54 -0700 Subject: [PATCH 25/50] Update _utils.py --- unsloth/models/_utils.py | 15 ++++++++++++++- 1 file changed, 14 insertions(+), 1 deletion(-) diff --git a/unsloth/models/_utils.py b/unsloth/models/_utils.py index f868c855bc..cd66825edf 100644 --- a/unsloth/models/_utils.py +++ b/unsloth/models/_utils.py @@ -41,6 +41,8 @@ __all__ = [ "torch_amp_custom_bwd", "accelerate_old_send_to_device", "accelerate_new_send_to_device", + "patch_gradient_checkpointing", + "unpatch_gradient_checkpointing", ] import torch @@ -791,7 +793,7 @@ class Unsloth_Offloaded_Gradient_Checkpointer(torch.autograd.Function): def backward(ctx, dY): (hidden_states,) = ctx.saved_tensors hidden_states = hidden_states.to("cuda:0", non_blocking = True).detach() - hidden_states.requires_grad = True + hidden_states.requires_grad_(True) with torch.enable_grad(): (output,) = ctx.forward_function(hidden_states, *ctx.args) torch.autograd.backward(output, dY) @@ -806,6 +808,17 @@ def unsloth_offloaded_gradient_checkpoint(function, *args, use_reentrant = None, pass +import torch.utils +old_checkpoint = torch.utils.checkpoint +def patch_gradient_checkpointing(): + torch.utils.checkpoint = unsloth_offloaded_gradient_checkpoint +pass + +def unpatch_gradient_checkpointing(): + torch.utils.checkpoint = old_checkpoint +pass + + # ============================================= # Fixes Bitsandbytes to remove missing warnings from transformers.utils.quantization_config import BitsAndBytesConfig, QuantizationMethod From 949a18dce364f5fd1d506d02164787d67d8a1aa8 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 00:23:23 -0700 Subject: [PATCH 26/50] Update _utils.py --- 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 cd66825edf..af7e1eb293 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__ = "2024.9.post2" +__version__ = "2024.9.post3" __all__ = [ "prepare_model_for_kbit_training", From 706f6ee18f93d3dfd3a18fb12a095005890a89a9 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 11:42:41 -0700 Subject: [PATCH 27/50] Llama 3.2 --- unsloth/kernels/rms_layernorm.py | 21 ++ unsloth/models/mapper.py | 16 + unsloth/models/vision.py | 592 +++++++++++++++++++++++++++++++ 3 files changed, 629 insertions(+) create mode 100644 unsloth/models/vision.py diff --git a/unsloth/kernels/rms_layernorm.py b/unsloth/kernels/rms_layernorm.py index 43924d0ab3..13faf08d6a 100644 --- a/unsloth/kernels/rms_layernorm.py +++ b/unsloth/kernels/rms_layernorm.py @@ -201,10 +201,25 @@ class Unsloth_LlamaRMSNorm(LlamaRMSNorm): pass pass +try: + from transformers.models.mllama.modeling_mllama import MllamaTextRMSNorm + class Unsloth_MllamaTextRMSNorm(MllamaTextRMSNorm): + def forward(self, X): + return fast_rms_layernorm(self, X, gemma = False) + pass + pass +except: + pass +pass def patch_rms_layernorm(): import transformers.models.llama.modeling_llama transformers.models.llama.modeling_llama.LlamaRMSNorm = Unsloth_LlamaRMSNorm + try: + import transformers.models.mllama.modeling_mllama + transformers.models.mllama.modeling_mllama.MllamaTextRMSNorm = Unsloth_MllamaTextRMSNorm + except: + pass return pass @@ -212,6 +227,12 @@ pass def unpatch_rms_layernorm(): import transformers.models.llama.modeling_llama transformers.models.llama.modeling_llama.LlamaRMSNorm = LlamaRMSNorm + try: + import transformers.models.mllama.modeling_mllama + transformers.models.mllama.modeling_mllama.MllamaTextRMSNorm = MllamaTextRMSNorm + except: + pass + return return pass diff --git a/unsloth/models/mapper.py b/unsloth/models/mapper.py index 50436a7a4f..7f27437904 100644 --- a/unsloth/models/mapper.py +++ b/unsloth/models/mapper.py @@ -400,6 +400,22 @@ __INT_TO_FLOAT_MAPPER = \ "unsloth/Qwen2.5-Coder-7B-Instruct", "Qwen/Qwen2.5-Coder-7B-Instruct", ), + "unsloth/Llama-3.2-1B-bnb-4bit" : ( + "unsloth/Llama-3.2-1B", + "meta-llama/Llama-3.2-1B", + ), + "unsloth/Llama-3.2-3B-bnb-4bit" : ( + "unsloth/Llama-3.2-3B", + "meta-llama/Llama-3.2-3B", + ), + "unsloth/Llama-3.2-1B-Instruct-bnb-4bit" : ( + "unsloth/Llama-3.2-1B-Instruct", + "meta-llama/Llama-3.2-1B-Instruct", + ), + "unsloth/Llama-3.2-3B-Instruct-bnb-4bit" : ( + "unsloth/Llama-3.2-3B-Instruct", + "meta-llama/Llama-3.2-3B-Instruct", + ), } INT_TO_FLOAT_MAPPER = {} diff --git a/unsloth/models/vision.py b/unsloth/models/vision.py new file mode 100644 index 0000000000..988fda4938 --- /dev/null +++ b/unsloth/models/vision.py @@ -0,0 +1,592 @@ +# Copyright 2023-present Daniel Han-Chen & the Unsloth team. All rights reserved. +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +from .llama import * +from ..kernels import patch_layernorm, unpatch_layernorm +from ..kernels import patch_rms_layernorm, unpatch_rms_layernorm +from ..kernels import patch_llama_for_causal_lm, unpatch_llama_for_causal_lm +from ._utils import patch_gradient_checkpointing + +from transformers import AutoProcessor, AutoModelForVision2Seq + + +class FastVisionModel: + + def pre_patch(self): + patch_gradient_checkpointing() + patch_layernorm() + patch_rms_layernorm() + patch_llama_for_causal_lm() + pass + + def post_unpatch(self): + unpatch_layernorm() + unpatch_rms_layernorm() + unpatch_llama_for_causal_lm() + pass + + + @staticmethod + def from_pretrained( + model_name = "llava-hf/llava-1.5-7b-hf", + max_seq_length = None, + dtype = None, + load_in_4bit = True, + token = None, + device_map = "sequential", + rope_scaling = None, + trust_remote_code = False, + **kwargs, + ): + if trust_remote_code: + print( + "Unsloth: WARNING `trust_remote_code` is True.\n"\ + "Are you certain you want to do remote code execution?" + ) + pass + if token is None: token = get_token() + if model_patcher is None: model_patcher = FastLlamaModel + SUPPORTS_BFLOAT16 = is_bfloat16_supported() + gpu_stats = torch.cuda.get_device_properties(0) + max_memory = round(gpu_stats.total_memory / 1024 / 1024 / 1024, 3) + + statistics = \ + f"==((====))== Unsloth {__version__}: Fast {model_patcher.__name__[4:-5]} patching. Transformers = {transformers_version}.\n"\ + f" \\\ /| GPU: {gpu_stats.name}. Max memory: {max_memory} GB. Platform = {platform_system}.\n"\ + f"O^O/ \_/ \\ Pytorch: {torch.__version__}. CUDA = {gpu_stats.major}.{gpu_stats.minor}. CUDA Toolkit = {torch.version.cuda}.\n"\ + f"\ / Bfloat16 = {str(SUPPORTS_BFLOAT16).upper()}. FA [Xformers = {xformers_version}. FA2 = {HAS_FLASH_ATTENTION}]\n"\ + f' "-____-" Free Apache license: http://github.com/unslothai/unsloth' + print(statistics) + + # Warn about fast transfers + old_hf_transfer = os.environ.get("HF_HUB_ENABLE_HF_TRANSFER", "0") + if os.environ.get("HF_HUB_ENABLE_HF_TRANSFER", "0") == "1": + print("Unsloth: Fast downloading is enabled - ignore downloading bars which are red colored!") + pass + # Return old flag + os.environ["HF_HUB_ENABLE_HF_TRANSFER"] = old_hf_transfer + + get_statistics() # For debugging - we use a download counter to see if environments are not breaking + + if dtype is None: + dtype = torch.float16 if not SUPPORTS_BFLOAT16 else torch.bfloat16 + elif dtype == torch.bfloat16 and not SUPPORTS_BFLOAT16: + logger.warning_once("Device does not support bfloat16. Will change to float16.") + dtype = torch.float16 + + assert(dtype == torch.float16 or dtype == torch.bfloat16 or dtype == torch.float32) + + # RoPE Scaling + model_config = AutoConfig.from_pretrained(model_name, token = token) + model_max_seq_length = model_config.max_position_embeddings + + # Check if RoPE Scaling is even allowed + model_function = MODEL_FOR_CAUSAL_LM_MAPPING[model_config.__class__] + has_rope_scaling = False + try: + with open(inspect.getfile(model_function), "r") as file: + has_rope_scaling = "self.config.rope_scaling" in file.read() + except: pass + has_rope_scaling = True + + # If max_seq_length is not specified, use maximum fron config + if max_seq_length is None: + max_seq_length = model_max_seq_length + pass + + if (rope_scaling is None) and (max_seq_length > model_max_seq_length): + + rope_scaling = max_seq_length / model_max_seq_length + + logger.warning_once( + f"Unsloth: {model_name} can only handle sequence lengths of at most "\ + f"{model_max_seq_length}.\nBut with kaiokendev's RoPE scaling of "\ + f"{round(rope_scaling, 3)}, it can be magically be extended to "\ + f"{max_seq_length}!" + ) + + # Warn RoPE scaling isn't allowed + if not has_rope_scaling: + raise RuntimeError( + "However, {model_name} doesn't support RoPE Scaling!\n"\ + "Please file a feature request at https://github.com/unslothai/unsloth." + ) + pass + + rope_scaling = {"type": "linear", "factor": rope_scaling,} + + # Add to kwargs + kwargs["rope_scaling"] = rope_scaling + pass + # We currently only support NVIDIA GPUs - AMD / Intel is a work in progress! + pre_check = check_nvidia() + + bnb_config = None + 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 = dtype, + ) + pass + + # https://huggingface.co/togethercomputer/LLaMA-2-7B-32K/discussions/12 + # RoPE Scaling's max_position_embeddings must be updated + max_position_embeddings = max(max_seq_length, model_max_seq_length) + kwargs.pop("attn_implementation", None); # No need since we auto call it + + # Cannot be None, since HF now checks for the config + if load_in_4bit: kwargs["quantization_config"] = bnb_config + + self.pre_patch() + model = AutoModelForVision2Seq.from_pretrained( + model_name, + device_map = device_map, + torch_dtype = dtype, + # quantization_config = bnb_config, + token = token, + max_position_embeddings = max_position_embeddings, + trust_remote_code = trust_remote_code, + attn_implementation = "eager", + **kwargs, + ) + self.post_unpatch() + + # Return old flag + os.environ["HF_HUB_ENABLE_HF_TRANSFER"] = old_hf_transfer + # We currently only support NVIDIA GPUs - AMD / Intel is a work in progress! + post_check = check_nvidia() + + # Counteract saved tokenizers + tokenizer = AutoProcessor.from_pretrained( + model_name, + ) + model = FastVisionModel.post_patch(model) + + # Patch Trainer + from transformers.trainer import Trainer + try: + if Trainer._inner_training_loop.__name__ != "_fast_inner_training_loop": + inner_training_loop = inspect.getsource(Trainer._inner_training_loop) + Trainer._original_training_loop = inner_training_loop + else: + inner_training_loop = Trainer._original_training_loop + except: + raise RuntimeError('Unsloth currently does not support multi GPU setups - but we are working on it!') + pass + + if ((post_check - pre_check) >= 1).sum() > 1: + raise RuntimeError('Unsloth currently does not support multi GPU setups - but we are working on it!') + + import transformers.trainer + items_in_trainer = dir(transformers.trainer) + good_items = [] + for item in items_in_trainer: + # TODO: Support Deepspeed + if item.startswith(("deepspeed", "xm", "met", "smp")): continue + if item in inner_training_loop: good_items.append(item) + pass + exec("from transformers.trainer import (" + ", ".join(x for x in good_items) + ")", globals()) + + start = re.search('logger\.info\([\"\'].+?Running training', inner_training_loop).span(0)[0] + end = inner_training_loop.find("\n\n", start) + original_debug = inner_training_loop[start:end] + spaces = re.search('\n([\s\t]{1,})', original_debug).group(0)[1:] + front_spaces = re.match('([\s\t]{1,})', inner_training_loop).group(0) + + debug_info = """debug_info = \\ + f"==((====))== Unsloth - 2x faster free finetuning | Num GPUs = {args.world_size}\\n"\\ + f" \\\\\\ /| Num examples = {num_examples:,} | Num Epochs = {num_train_epochs:,}\\n"\\ + f"O^O/ \\_/ \\ Batch size per device = {self._train_batch_size:,} | Gradient Accumulation steps = {args.gradient_accumulation_steps}\\n"\\ + f"\\ / Total batch size = {total_train_batch_size:,} | Total steps = {max_steps:,}\\n"\\ + f' "-____-" Number of trainable parameters = {get_model_param_count(model, trainable_only=True):,}' + logger.warning(debug_info) + import subprocess, re, gc, numpy as np + a = np.array([0,]) + try: + a = subprocess.check_output('nvidia-smi --query-gpu=memory.used --format=csv', shell = True) + a = re.findall(rb'([\\d]{1,})[\\s]{1,}M', a) + a = np.array([int(x.decode('utf-8'))/1024 for x in a]) + except: + if not torch.cuda.is_available(): + raise RuntimeError('Unsloth: We do not support AMD / Intel machines yet - it is a work in progress!') + if ((a - PRE_CHECK) >= 1).sum() > 1: + raise RuntimeError('Unsloth currently does not support multi GPU setups - but we are working on it!') + for _ in range(3): + gc.collect() + torch.cuda.empty_cache()""" + + debug_info = debug_info.split('\n') + debug_info = "\n".join([debug_info[0]] + [spaces + x[8:] for x in debug_info[1:]]) + inner_training_loop = inner_training_loop.replace(original_debug, debug_info) + + debug_info = """n_total_devices = total_train_batch_size // \\ + args.gradient_accumulation_steps // self._train_batch_size + if n_total_devices > 1: + logger.warning_once('Unsloth currently does not support multi GPU setups - but we are working on it!') + debug_info =""" + debug_info = debug_info.split('\n') + debug_info = "\n".join([debug_info[0]] + [spaces + x[8:] for x in debug_info[1:]]) + inner_training_loop = inner_training_loop.replace("debug_info =", debug_info, 1) + + front_spaces = re.match(r"[\t\s]{1,}", inner_training_loop).group(0) + inner_training_loop = re.sub(r"^" + front_spaces, "", inner_training_loop, flags = re.MULTILINE) + inner_training_loop = inner_training_loop.replace( + "train_dataloader = tpu_spmd_dataloader(train_dataloader)", + "raise RuntimeError('Unsloth: TPUs are not yet supported!')" + ) + inner_training_loop = inner_training_loop.replace( + "self.accelerator.free_memory()", + "self.accelerator.free_memory()\n" + \ + front_spaces + "if self.is_deepspeed_enabled:"\ + "raise RuntimeError('Unsloth: Deepspeed is not yet supported!')\n", 1, + ) + + check_batches = """train_dataloader = self.get_train_dataloader() + ga = args.gradient_accumulation_steps + bsz = self._train_batch_size + total_batches = bsz * ga * args.world_size + n_total_devices = total_batches // ga // bsz + if n_total_devices > 1: + logger.warning_once('Unsloth currently does not support multi GPU setups - but we are working on it!') + divisor = n_total_devices / 1 + bsz = self._train_batch_size = max(int(bsz / divisor), 1) + if total_batches // ga // bsz > 1: + divisor = n_total_devices / 1 + ga = args.gradient_accumulation_steps = max(int(ga / divisor), 1)""" + check_batches = check_batches.split('\n') + check_batches = "\n".join([check_batches[0]] + [front_spaces + x[8:] for x in check_batches[1:]]) + inner_training_loop = inner_training_loop.replace( + "train_dataloader = self.get_train_dataloader()", + check_batches, 1, + ) + inner_training_loop = inner_training_loop.replace( + "_inner_training_loop", + "_fast_inner_training_loop", 1, + ) + exec(inner_training_loop, globals()) + + Trainer._inner_training_loop = _fast_inner_training_loop + inner_training_loop = inner_training_loop.replace( + "is_torch_tpu_available()", + "False", + ) + if "n_total_devices >" not in inner_training_loop: + raise RuntimeError('Unsloth currently does not support multi GPU setups - but we are working on it!') + pass + inner_training_loop = inner_training_loop.replace( + "is_sagemaker_mp_enabled()", + "False", + ) + exec(inner_training_loop, globals()) + Trainer._inner_training_loop = _fast_inner_training_loop + + # Save max_seq_length + model.max_seq_length = max_position_embeddings + internal_model = model + while hasattr(internal_model, "model"): + internal_model.max_seq_length = max_position_embeddings + internal_model = internal_model.model + pass + internal_model.max_seq_length = max_position_embeddings + + # Fix up config for transformers uploading PEFT + # Not necessary anymore since we require transformers>=4.37! + if False: + name = model.config._name_or_path + if name.startswith("unsloth/") and name.endswith("-bnb-4bit"): + name = name[:len(name) - len("-bnb-4bit")] + model.config.update({"_name_or_path" : name}) + pass + pass + + # Log Unsloth version for future fastpaths for inference + model.config.update({"unsloth_version" : __version__}) + + # Add save modules + patch_saving_functions(model) + Trainer._inner_training_loop = _fast_inner_training_loop + + # Also fix torch_dtype + internal_model = model + while hasattr(internal_model, "model"): + if hasattr(internal_model, "config"): + if internal_model.config.torch_dtype == "float32": + internal_model.config.torch_dtype = torch.float32 + elif internal_model.config.torch_dtype == "bfloat16": + internal_model.config.torch_dtype = torch.bfloat16 + elif internal_model.config.torch_dtype == "float16": + internal_model.config.torch_dtype = torch.float16 + pass + pass + internal_model = internal_model.model + pass + if hasattr(internal_model, "config"): + if internal_model.config.torch_dtype == "float32": + internal_model.config.torch_dtype = torch.float32 + elif internal_model.config.torch_dtype == "bfloat16": + internal_model.config.torch_dtype = torch.bfloat16 + elif internal_model.config.torch_dtype == "float16": + internal_model.config.torch_dtype = torch.float16 + pass + pass + + return model, tokenizer + pass + + + @staticmethod + def post_patch(model): + # Patch model + layers = model.model.layers + lm_head = model.get_output_embeddings().weight + + # Also patch all dtypes - BnB seems to not allocate the correct type? + # BnB default dtype seems to be float16! + correct_dtype = lm_head.weight.dtype + + for name, module in model.named_modules(): + if isinstance(module, (Bnb_Linear4bit, Peft_Linear4bit)): + weight = module.weight + quant_state = weight.quant_state + + if type(quant_state) is list: + # BnB seems to have float16 as default! + module.weight.quant_state[2] = correct_dtype # Cast to correct dtype + else: + # https://github.com/TimDettmers/bitsandbytes/pull/763/files + quant_state.dtype = correct_dtype + pass + pass + pass + + # Clear deleted GPU items + for _ in range(3): + gc.collect() + torch.cuda.empty_cache() + return model + pass + + + @staticmethod + def get_peft_model( + model, + r = 16, + target_modules = "all-linear", + lora_alpha = 16, + lora_dropout = 0, + bias = "none", + layers_to_transform = None, + layers_pattern = None, + use_gradient_checkpointing = True, + random_state = 3407, + max_seq_length = 2048, # not used anymore + use_rslora = False, + modules_to_save = None, + init_lora_weights = True, + loftq_config = {}, + temporary_location = "_unsloth_temporary_saved_buffers", + **kwargs, + ): + transformers_set_seed(random_state) + + # Get LoRA + arguments = dict( + r = r, + lora_alpha = lora_alpha, + target_modules = target_modules, + lora_dropout = lora_dropout, + bias = bias, + layers_to_transform = layers_to_transform, + init_lora_weights = init_lora_weights, + # loftq_config = loftq_config, + # use_rslora = use_rslora, + modules_to_save = modules_to_save, + **kwargs, + ) + + lora_config = LoraConfig(**arguments) + + model = _get_peft_model(model, lora_config) + + model = FastVisionModel.patch_peft_model(model, use_gradient_checkpointing) + + # Clear deleted GPU items + for _ in range(3): + gc.collect() + torch.cuda.empty_cache() + pass + + return model + pass + + + @staticmethod + def patch_peft_model( + model, + use_gradient_checkpointing = True, + ): + + model = prepare_model_for_kbit_training( + model, + use_gradient_checkpointing = use_gradient_checkpointing, + use_reentrant = True, + ) + + # Fix up config for transformers uploading PEFT + for active_adapter in model.peft_config.keys(): + # Not necessary since we requires transformers >= 4.37 + if False: + name = model.peft_config[active_adapter].base_model_name_or_path + if name.startswith("unsloth/") and name.endswith("-bnb-4bit"): + name = name[:len(name) - len("-bnb-4bit")] + model.peft_config[active_adapter].base_model_name_or_path = name + pass + # Add revision to enable future fast inference paths + # [TODO] Bugs out!see https://github.com/unslothai/unsloth/issues/492 + # model.peft_config[active_adapter].revision = f"unsloth" + pass + + from transformers.trainer import Trainer + if Trainer._inner_training_loop.__name__ != "_fast_inner_training_loop": + raise RuntimeError( + 'Unsloth currently does not work on multi GPU setups - sadly we are a 2 brother team so '\ + 'enabling it will require much more work, so we have to prioritize. Please understand!\n'\ + 'We do have a separate beta version, which you can contact us about!\n'\ + 'Thank you for your understanding and we appreciate it immensely!' + ) + pass + + logger.warning_once( + f"Unsloth {__version__} patched {len(model.model.model.layers)} layers with "\ + f"{n_qkv} QKV layers, {n_o} O layers and {n_mlp} MLP layers.", + ) + patch_saving_functions(model) + + # Patch cross entropy loss labels + # Fixes https://github.com/unslothai/unsloth/issues/10 + max_seq_length = model.max_seq_length + extra_ignored_labels = torch.full((max_seq_length, 1), -100, device = "cuda:0") + model.model.extra_ignored_labels = extra_ignored_labels + internal_model = model + while hasattr(internal_model, "model"): + internal_model.max_seq_length = max_seq_length + internal_model = internal_model.model + pass + internal_model.max_seq_length = max_seq_length + + # Patch tokenizer to pad to the right + internal_model = model + while hasattr(internal_model, "model"): + if hasattr(internal_model, "_saved_temp_tokenizer"): + internal_model._saved_temp_tokenizer.padding_side = "right" + pass + internal_model = internal_model.model + pass + if hasattr(internal_model, "_saved_temp_tokenizer"): + internal_model._saved_temp_tokenizer.padding_side = "right" + pass + + # Clear deleted GPU items + for _ in range(3): + gc.collect() + torch.cuda.empty_cache() + pass + return model + pass + + + @staticmethod + def for_inference(model): + # if model.config.model_type == "qwen2": + # FastLlamaModel.for_training(model) + # return + # pass + + internal_model = model + internal_model.gradient_checkpointing = False + internal_model.training = False + + while hasattr(internal_model, "model"): + internal_model = internal_model.model + internal_model.gradient_checkpointing = False + internal_model.training = False + pass + if hasattr(internal_model, "training"): + internal_model.training = False + pass + + # Also check if lm_head / embeddings are trained + internal_model = model + while not hasattr(internal_model, "lm_head"): + internal_model = internal_model.model + pass + lm_head = internal_model.lm_head.weight + device_type = lm_head.device.type + dtype = model.config.torch_dtype + + if type(dtype) is str: + if dtype == "float16": dtype = torch.float16 + elif dtype == "bfloat16": dtype = torch.bfloat16 + pass + + # Also disable training for embeddings for NEFTune + if hasattr(model, "get_input_embeddings"): + embeddings = model.get_input_embeddings() + if hasattr(embeddings, "training"): embeddings.training = False + pass + if hasattr(model, "get_output_embeddings"): + embeddings = model.get_output_embeddings() + if hasattr(embeddings, "training"): embeddings.training = False + pass + + return model + pass + + + @staticmethod + def for_training(model, use_gradient_checkpointing = True): + internal_model = model + internal_model.gradient_checkpointing = use_gradient_checkpointing + internal_model.training = True + + # Delete all fast inference loras + for param in model.parameters(): + if hasattr(param, "_fast_lora"): + del param._fast_lora + pass + + while hasattr(internal_model, "model"): + internal_model = internal_model.model + internal_model.gradient_checkpointing = use_gradient_checkpointing + internal_model.training = True + pass + if hasattr(internal_model, "training"): + internal_model.training = True + pass + + # Also re-enable training for embeddings for NEFTune + if hasattr(model, "get_input_embeddings"): + embeddings = model.get_input_embeddings() + if hasattr(embeddings, "training"): embeddings.training = True + pass + if hasattr(model, "get_output_embeddings"): + embeddings = model.get_output_embeddings() + if hasattr(embeddings, "training"): embeddings.training = True + pass + + return model + pass +pass From 2bfa0aa6afda7322f081c1b58f31ec5b33cbd7ff Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 13:19:01 -0700 Subject: [PATCH 28/50] Update _utils.py --- unsloth/models/_utils.py | 34 +++++++++++++++++----------------- 1 file changed, 17 insertions(+), 17 deletions(-) diff --git a/unsloth/models/_utils.py b/unsloth/models/_utils.py index af7e1eb293..09b448a2f6 100644 --- a/unsloth/models/_utils.py +++ b/unsloth/models/_utils.py @@ -141,24 +141,24 @@ pass # ============================================= # Fix KeyError: 'Cache only has 0 layers, attempted to access layer with index 0' -import transformers.cache_utils -if hasattr(transformers.cache_utils, "DynamicCache") and \ - transformers.cache_utils.DynamicCache.__getitem__.__name__ != "__cache_utils_getitem__": +# import transformers.cache_utils +# if hasattr(transformers.cache_utils, "DynamicCache") and \ +# transformers.cache_utils.DynamicCache.__getitem__.__name__ != "__cache_utils_getitem__": - source = inspect.getsource(transformers.cache_utils.DynamicCache.__getitem__) - start = source.find("def") - spaces = start*" " - source = source.split("\n") - source = "\n".join(x[start:] for x in source) - where = source.find("raise KeyError") - source = source[:where] + \ - f"if len(self) == 0:\n{spaces}{spaces}"\ - " raise RuntimeError('Unsloth: You must call `FastLanguageModel.for_inference(model)` before doing inference for Unsloth models.')\n" + \ - f"{spaces}{spaces}else:\n{spaces}{spaces}{spaces}" + source[where:] - source = source.replace("__getitem__", "__cache_utils_getitem__", 1) - exec(source) - transformers.cache_utils.DynamicCache.__getitem__ = __cache_utils_getitem__ -pass +# source = inspect.getsource(transformers.cache_utils.DynamicCache.__getitem__) +# start = source.find("def") +# spaces = start*" " +# source = source.split("\n") +# source = "\n".join(x[start:] for x in source) +# where = source.find("raise KeyError") +# source = source[:where] + \ +# f"if len(self) == 0:\n{spaces}{spaces}"\ +# " raise RuntimeError('Unsloth: You must call `FastLanguageModel.for_inference(model)` before doing inference for Unsloth models.')\n" + \ +# f"{spaces}{spaces}else:\n{spaces}{spaces}{spaces}" + source[where:] +# source = source.replace("__getitem__", "__cache_utils_getitem__", 1) +# exec(source) +# transformers.cache_utils.DynamicCache.__getitem__ = __cache_utils_getitem__ +# pass # ============================================= # ============================================= From 0ee78d455c7b9131ef468848e31cab489f729306 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 13:28:05 -0700 Subject: [PATCH 29/50] Update _utils.py --- unsloth/models/_utils.py | 34 +++++++++++++++++----------------- 1 file changed, 17 insertions(+), 17 deletions(-) diff --git a/unsloth/models/_utils.py b/unsloth/models/_utils.py index 09b448a2f6..6144efe485 100644 --- a/unsloth/models/_utils.py +++ b/unsloth/models/_utils.py @@ -141,24 +141,24 @@ pass # ============================================= # Fix KeyError: 'Cache only has 0 layers, attempted to access layer with index 0' -# import transformers.cache_utils -# if hasattr(transformers.cache_utils, "DynamicCache") and \ -# transformers.cache_utils.DynamicCache.__getitem__.__name__ != "__cache_utils_getitem__": +import transformers.cache_utils +if hasattr(transformers.cache_utils, "DynamicCache") and \ + transformers.cache_utils.DynamicCache.__getitem__.__name__ != "__cache_utils_getitem__": -# source = inspect.getsource(transformers.cache_utils.DynamicCache.__getitem__) -# start = source.find("def") -# spaces = start*" " -# source = source.split("\n") -# source = "\n".join(x[start:] for x in source) -# where = source.find("raise KeyError") -# source = source[:where] + \ -# f"if len(self) == 0:\n{spaces}{spaces}"\ -# " raise RuntimeError('Unsloth: You must call `FastLanguageModel.for_inference(model)` before doing inference for Unsloth models.')\n" + \ -# f"{spaces}{spaces}else:\n{spaces}{spaces}{spaces}" + source[where:] -# source = source.replace("__getitem__", "__cache_utils_getitem__", 1) -# exec(source) -# transformers.cache_utils.DynamicCache.__getitem__ = __cache_utils_getitem__ -# pass + source = inspect.getsource(transformers.cache_utils.DynamicCache.__getitem__) + start = source.find("def") + spaces = start*" " + source = source.split("\n") + source = "\n".join(x[start:] for x in source) + where = source.find("raise KeyError") + # source = source[:where] + \ + # f"if len(self) == 0:\n{spaces}{spaces}"\ + # " raise RuntimeError('Unsloth: You must call `FastLanguageModel.for_inference(model)` before doing inference for Unsloth models.')\n" + \ + # f"{spaces}{spaces}else:\n{spaces}{spaces}{spaces}" + source[where:] + source = source.replace("__getitem__", "__cache_utils_getitem__", 1) + exec(source) + transformers.cache_utils.DynamicCache.__getitem__ = __cache_utils_getitem__ +pass # ============================================= # ============================================= From 5da313d8836e7625cbe148bab024ab651b3a0ea1 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 14:12:16 -0700 Subject: [PATCH 30/50] Update _utils.py --- unsloth/models/_utils.py | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/unsloth/models/_utils.py b/unsloth/models/_utils.py index 6144efe485..af7e1eb293 100644 --- a/unsloth/models/_utils.py +++ b/unsloth/models/_utils.py @@ -151,10 +151,10 @@ if hasattr(transformers.cache_utils, "DynamicCache") and \ source = source.split("\n") source = "\n".join(x[start:] for x in source) where = source.find("raise KeyError") - # source = source[:where] + \ - # f"if len(self) == 0:\n{spaces}{spaces}"\ - # " raise RuntimeError('Unsloth: You must call `FastLanguageModel.for_inference(model)` before doing inference for Unsloth models.')\n" + \ - # f"{spaces}{spaces}else:\n{spaces}{spaces}{spaces}" + source[where:] + source = source[:where] + \ + f"if len(self) == 0:\n{spaces}{spaces}"\ + " raise RuntimeError('Unsloth: You must call `FastLanguageModel.for_inference(model)` before doing inference for Unsloth models.')\n" + \ + f"{spaces}{spaces}else:\n{spaces}{spaces}{spaces}" + source[where:] source = source.replace("__getitem__", "__cache_utils_getitem__", 1) exec(source) transformers.cache_utils.DynamicCache.__getitem__ = __cache_utils_getitem__ From 25c6e1d8a77132bb662263e697f69266261ffef7 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 14:32:29 -0700 Subject: [PATCH 31/50] Update llama.py --- unsloth/models/llama.py | 6 +----- 1 file changed, 1 insertion(+), 5 deletions(-) diff --git a/unsloth/models/llama.py b/unsloth/models/llama.py index bae6d5b80d..2524ec1312 100644 --- a/unsloth/models/llama.py +++ b/unsloth/models/llama.py @@ -1386,11 +1386,7 @@ def _wrap_fast_inference(generate, device_type, dtype, model): pass # For newer HF - if SUPPORTS_LLAMA32: - # kwargs["cache_implementation"] = "hybrid" - pass - else: - kwargs["cache_implementation"] = "dynamic" + kwargs["cache_implementation"] = "dynamic" # For num_logits_to_keep kwargs["num_logits_to_keep"] = 1 From 4412dd7e9fd6f3b2807096ded42247e8a0c688e5 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 17:44:13 -0700 Subject: [PATCH 32/50] Update vision.py --- unsloth/models/vision.py | 62 +++++++--------------------------------- 1 file changed, 10 insertions(+), 52 deletions(-) diff --git a/unsloth/models/vision.py b/unsloth/models/vision.py index 988fda4938..0b8c08a371 100644 --- a/unsloth/models/vision.py +++ b/unsloth/models/vision.py @@ -18,8 +18,14 @@ from ..kernels import patch_rms_layernorm, unpatch_rms_layernorm from ..kernels import patch_llama_for_causal_lm, unpatch_llama_for_causal_lm from ._utils import patch_gradient_checkpointing -from transformers import AutoProcessor, AutoModelForVision2Seq - +from transformers import AutoProcessor +try: + from transformers import MllamaForConditionalGeneration +except: + raise ImportError( + "Unsloth: Please update your transformers version to 4.46.0 for Llama 3.2 support!" + ) +pass class FastVisionModel: @@ -56,7 +62,6 @@ class FastVisionModel: ) pass if token is None: token = get_token() - if model_patcher is None: model_patcher = FastLlamaModel SUPPORTS_BFLOAT16 = is_bfloat16_supported() gpu_stats = torch.cuda.get_device_properties(0) max_memory = round(gpu_stats.total_memory / 1024 / 1024 / 1024, 3) @@ -87,48 +92,6 @@ class FastVisionModel: assert(dtype == torch.float16 or dtype == torch.bfloat16 or dtype == torch.float32) - # RoPE Scaling - model_config = AutoConfig.from_pretrained(model_name, token = token) - model_max_seq_length = model_config.max_position_embeddings - - # Check if RoPE Scaling is even allowed - model_function = MODEL_FOR_CAUSAL_LM_MAPPING[model_config.__class__] - has_rope_scaling = False - try: - with open(inspect.getfile(model_function), "r") as file: - has_rope_scaling = "self.config.rope_scaling" in file.read() - except: pass - has_rope_scaling = True - - # If max_seq_length is not specified, use maximum fron config - if max_seq_length is None: - max_seq_length = model_max_seq_length - pass - - if (rope_scaling is None) and (max_seq_length > model_max_seq_length): - - rope_scaling = max_seq_length / model_max_seq_length - - logger.warning_once( - f"Unsloth: {model_name} can only handle sequence lengths of at most "\ - f"{model_max_seq_length}.\nBut with kaiokendev's RoPE scaling of "\ - f"{round(rope_scaling, 3)}, it can be magically be extended to "\ - f"{max_seq_length}!" - ) - - # Warn RoPE scaling isn't allowed - if not has_rope_scaling: - raise RuntimeError( - "However, {model_name} doesn't support RoPE Scaling!\n"\ - "Please file a feature request at https://github.com/unslothai/unsloth." - ) - pass - - rope_scaling = {"type": "linear", "factor": rope_scaling,} - - # Add to kwargs - kwargs["rope_scaling"] = rope_scaling - pass # We currently only support NVIDIA GPUs - AMD / Intel is a work in progress! pre_check = check_nvidia() @@ -142,16 +105,11 @@ class FastVisionModel: ) pass - # https://huggingface.co/togethercomputer/LLaMA-2-7B-32K/discussions/12 - # RoPE Scaling's max_position_embeddings must be updated - max_position_embeddings = max(max_seq_length, model_max_seq_length) - kwargs.pop("attn_implementation", None); # No need since we auto call it - # Cannot be None, since HF now checks for the config if load_in_4bit: kwargs["quantization_config"] = bnb_config self.pre_patch() - model = AutoModelForVision2Seq.from_pretrained( + model = MllamaForConditionalGeneration.from_pretrained( model_name, device_map = device_map, torch_dtype = dtype, @@ -159,7 +117,7 @@ class FastVisionModel: token = token, max_position_embeddings = max_position_embeddings, trust_remote_code = trust_remote_code, - attn_implementation = "eager", + attn_implementation = "sdpa", **kwargs, ) self.post_unpatch() From a4ee0bc8060b1439b50249909d42ea8afeac4831 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 17:46:29 -0700 Subject: [PATCH 33/50] Update llama.py --- unsloth/models/llama.py | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/unsloth/models/llama.py b/unsloth/models/llama.py index 2524ec1312..b7dc68586c 100644 --- a/unsloth/models/llama.py +++ b/unsloth/models/llama.py @@ -843,6 +843,9 @@ def LlamaModel_fast_forward( pass +global past_key_values_all +past_key_values_all = None + # https://github.com/huggingface/transformers/blob/main/src/transformers/models/llama/modeling_llama.py#L825 def LlamaModel_fast_forward_inference( self, @@ -855,6 +858,8 @@ def LlamaModel_fast_forward_inference( hidden_states = self.model.embed_tokens(input_ids) hidden_states = hidden_states.to(self.config.torch_dtype) bsz, q_len, hd = hidden_states.shape + global past_key_values_all + past_key_values_all = past_key_values seq_len = past_key_values[0][0].shape[-2] if bsz != 1: attention_mask = _prepare_4d_causal_attention_mask_for_sdpa( From 360234fbc7736f336c947507d0187715160dba9b Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 17:51:36 -0700 Subject: [PATCH 34/50] Update llama.py --- unsloth/models/llama.py | 7 ++----- 1 file changed, 2 insertions(+), 5 deletions(-) diff --git a/unsloth/models/llama.py b/unsloth/models/llama.py index b7dc68586c..07bb5505ef 100644 --- a/unsloth/models/llama.py +++ b/unsloth/models/llama.py @@ -843,9 +843,6 @@ def LlamaModel_fast_forward( pass -global past_key_values_all -past_key_values_all = None - # https://github.com/huggingface/transformers/blob/main/src/transformers/models/llama/modeling_llama.py#L825 def LlamaModel_fast_forward_inference( self, @@ -858,8 +855,8 @@ def LlamaModel_fast_forward_inference( hidden_states = self.model.embed_tokens(input_ids) hidden_states = hidden_states.to(self.config.torch_dtype) bsz, q_len, hd = hidden_states.shape - global past_key_values_all - past_key_values_all = past_key_values + import os + os.environ["past_key_values_all"] = past_key_values seq_len = past_key_values[0][0].shape[-2] if bsz != 1: attention_mask = _prepare_4d_causal_attention_mask_for_sdpa( From e19dee7474350a421e1a54062eee8d6d80968aac Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 17:55:42 -0700 Subject: [PATCH 35/50] Update llama.py --- unsloth/models/llama.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/unsloth/models/llama.py b/unsloth/models/llama.py index 07bb5505ef..c187449902 100644 --- a/unsloth/models/llama.py +++ b/unsloth/models/llama.py @@ -1388,7 +1388,7 @@ def _wrap_fast_inference(generate, device_type, dtype, model): pass # For newer HF - kwargs["cache_implementation"] = "dynamic" + # kwargs["cache_implementation"] = "dynamic" # For num_logits_to_keep kwargs["num_logits_to_keep"] = 1 From d4eed9807f094285342c26b037d1943b23df4d9a Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 18:07:21 -0700 Subject: [PATCH 36/50] Update llama.py --- unsloth/models/llama.py | 2 ++ 1 file changed, 2 insertions(+) diff --git a/unsloth/models/llama.py b/unsloth/models/llama.py index c187449902..48d1f7ec6e 100644 --- a/unsloth/models/llama.py +++ b/unsloth/models/llama.py @@ -1409,6 +1409,8 @@ def _wrap_fast_inference(generate, device_type, dtype, model): # Autocasted with torch.autocast(device_type = device_type, dtype = dtype): + print(args) + print(kwargs) output = generate(*args, **kwargs) pass From dc85077f1f790b35504808f64c1d47550be91541 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 18:10:46 -0700 Subject: [PATCH 37/50] Update llama.py --- unsloth/models/llama.py | 2 ++ 1 file changed, 2 insertions(+) diff --git a/unsloth/models/llama.py b/unsloth/models/llama.py index 48d1f7ec6e..a31a485676 100644 --- a/unsloth/models/llama.py +++ b/unsloth/models/llama.py @@ -1372,6 +1372,8 @@ def _wrap_fast_inference(generate, device_type, dtype, model): # Wraps inference with bfloat16 / float16 @torch.inference_mode def _fast_generate(*args, **kwargs): + print(args) + print(kwargs) # Set a flag for generation! internal_model = model From cbbff0697337d1dd8b5f03d8d69aa181d8b0d370 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 18:24:30 -0700 Subject: [PATCH 38/50] Update llama.py --- unsloth/models/llama.py | 10 ++++------ 1 file changed, 4 insertions(+), 6 deletions(-) diff --git a/unsloth/models/llama.py b/unsloth/models/llama.py index a31a485676..dd6d805ec4 100644 --- a/unsloth/models/llama.py +++ b/unsloth/models/llama.py @@ -1384,10 +1384,10 @@ def _wrap_fast_inference(generate, device_type, dtype, model): internal_model._flag_for_generation = True # Must patch accelerate for Xformers - if accelerate_new_send_to_device is not None: - import accelerate.utils.operations - accelerate.utils.operations.send_to_device = accelerate_new_send_to_device - pass + # if accelerate_new_send_to_device is not None: + # import accelerate.utils.operations + # accelerate.utils.operations.send_to_device = accelerate_new_send_to_device + # pass # For newer HF # kwargs["cache_implementation"] = "dynamic" @@ -1411,8 +1411,6 @@ def _wrap_fast_inference(generate, device_type, dtype, model): # Autocasted with torch.autocast(device_type = device_type, dtype = dtype): - print(args) - print(kwargs) output = generate(*args, **kwargs) pass From 72de37e3214c6b013c0ea6dc4557a9a9f322b25a Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 19:15:40 -0700 Subject: [PATCH 39/50] Update llama.py --- unsloth/models/llama.py | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/unsloth/models/llama.py b/unsloth/models/llama.py index dd6d805ec4..24f9942119 100644 --- a/unsloth/models/llama.py +++ b/unsloth/models/llama.py @@ -1392,17 +1392,17 @@ def _wrap_fast_inference(generate, device_type, dtype, model): # For newer HF # kwargs["cache_implementation"] = "dynamic" # For num_logits_to_keep - kwargs["num_logits_to_keep"] = 1 + # kwargs["num_logits_to_keep"] = 1 - # Remove token_type_ids - kwargs.pop("token_type_ids", None) + # # Remove token_type_ids + # kwargs.pop("token_type_ids", None) - # Check pad_token - model_eos_token_id = getattr(model.config, "eos_token_id", None) - if model_eos_token_id is not None and hasattr(model_eos_token_id, "__iter__"): - model_eos_token_id = model_eos_token_id[0] + # # Check pad_token + # model_eos_token_id = getattr(model.config, "eos_token_id", None) + # if model_eos_token_id is not None and hasattr(model_eos_token_id, "__iter__"): + # model_eos_token_id = model_eos_token_id[0] - kwargs["pad_token_id"] = kwargs.pop("pad_token_id", model_eos_token_id) + # kwargs["pad_token_id"] = kwargs.pop("pad_token_id", model_eos_token_id) # Set pad token # old_pad_token_id = getattr(model.config, "pad_token_id", None) From 7ca67313d0fd65765f61241af46f25dcd2a5b8ac Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 25 Sep 2024 19:38:01 -0700 Subject: [PATCH 40/50] Update llama.py --- unsloth/models/llama.py | 33 ++++++++++++++++----------------- 1 file changed, 16 insertions(+), 17 deletions(-) diff --git a/unsloth/models/llama.py b/unsloth/models/llama.py index 24f9942119..f5dc02704e 100644 --- a/unsloth/models/llama.py +++ b/unsloth/models/llama.py @@ -855,8 +855,6 @@ def LlamaModel_fast_forward_inference( hidden_states = self.model.embed_tokens(input_ids) hidden_states = hidden_states.to(self.config.torch_dtype) bsz, q_len, hd = hidden_states.shape - import os - os.environ["past_key_values_all"] = past_key_values seq_len = past_key_values[0][0].shape[-2] if bsz != 1: attention_mask = _prepare_4d_causal_attention_mask_for_sdpa( @@ -1372,8 +1370,6 @@ def _wrap_fast_inference(generate, device_type, dtype, model): # Wraps inference with bfloat16 / float16 @torch.inference_mode def _fast_generate(*args, **kwargs): - print(args) - print(kwargs) # Set a flag for generation! internal_model = model @@ -1384,25 +1380,28 @@ def _wrap_fast_inference(generate, device_type, dtype, model): internal_model._flag_for_generation = True # Must patch accelerate for Xformers - # if accelerate_new_send_to_device is not None: - # import accelerate.utils.operations - # accelerate.utils.operations.send_to_device = accelerate_new_send_to_device - # pass + if accelerate_new_send_to_device is not None: + import accelerate.utils.operations + accelerate.utils.operations.send_to_device = accelerate_new_send_to_device + pass # For newer HF - # kwargs["cache_implementation"] = "dynamic" + if SUPPORTS_LLAMA32: + kwargs["cache_implementation"] = "hybrid" + else: + kwargs["cache_implementation"] = "dynamic" # For num_logits_to_keep - # kwargs["num_logits_to_keep"] = 1 + kwargs["num_logits_to_keep"] = 1 - # # Remove token_type_ids - # kwargs.pop("token_type_ids", None) + # Remove token_type_ids + kwargs.pop("token_type_ids", None) - # # Check pad_token - # model_eos_token_id = getattr(model.config, "eos_token_id", None) - # if model_eos_token_id is not None and hasattr(model_eos_token_id, "__iter__"): - # model_eos_token_id = model_eos_token_id[0] + # Check pad_token + model_eos_token_id = getattr(model.config, "eos_token_id", None) + if model_eos_token_id is not None and hasattr(model_eos_token_id, "__iter__"): + model_eos_token_id = model_eos_token_id[0] - # kwargs["pad_token_id"] = kwargs.pop("pad_token_id", model_eos_token_id) + kwargs["pad_token_id"] = kwargs.pop("pad_token_id", model_eos_token_id) # Set pad token # old_pad_token_id = getattr(model.config, "pad_token_id", None) From 98a1e16f57158c89da5d993bcf6e00fecd678aa2 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sun, 29 Sep 2024 23:13:44 -0700 Subject: [PATCH 41/50] Update loader.py --- unsloth/models/loader.py | 19 +++++++++++++++++-- 1 file changed, 17 insertions(+), 2 deletions(-) diff --git a/unsloth/models/loader.py b/unsloth/models/loader.py index 0ac9b02743..2a89c06c21 100644 --- a/unsloth/models/loader.py +++ b/unsloth/models/loader.py @@ -23,6 +23,7 @@ from peft import PeftConfig, PeftModel from .mapper import INT_TO_FLOAT_MAPPER, FLOAT_TO_INT_MAPPER, MAP_TO_UNSLOTH_16bit import os from huggingface_hub.utils._token import get_token +from huggingface_hub import HfFileSystem # https://github.com/huggingface/transformers/pull/26037 allows 4 bit loading! from packaging.version import Version @@ -191,14 +192,28 @@ class FastLanguageModel(FastLlamaModel): is_peft = False pass - # Cannot be both! - if (is_model and is_peft) and not SUPPORTS_LLAMA32: + # Both config.json and adapter_config.json should not exist! + + # Old transformers versions check + both_exist = (is_model and is_peft) and not SUPPORTS_LLAMA32 + + if SUPPORTS_LLAMA32: + # New transformers need to check manually. + files = HfFileSystem(token = token).glob(os.path.join(model_name, "*.json")) + if sum(x.endswith(("adapter_config.json", "config.json")) for x in files) >= 2: + both_exist = True + pass + pass + + # Error out if both LoRA and normal model config exists. + if both_exist: raise RuntimeError( "Unsloth: Your repo has a LoRA adapter and a base model.\n"\ "You have 2 files `config.json` and `adapter_config.json`.\n"\ "We must only allow one config file.\n"\ "Please separate the LoRA and base models to 2 repos." ) + elif not is_model and not is_peft: error = autoconfig_error or peft_error # Old transformers version From fe9171d15ca123d201cf1185c62b8c3616ff1845 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sun, 29 Sep 2024 23:15:22 -0700 Subject: [PATCH 42/50] Update loader.py --- unsloth/models/loader.py | 2 ++ 1 file changed, 2 insertions(+) diff --git a/unsloth/models/loader.py b/unsloth/models/loader.py index 2a89c06c21..9ea6c308e3 100644 --- a/unsloth/models/loader.py +++ b/unsloth/models/loader.py @@ -197,12 +197,14 @@ class FastLanguageModel(FastLlamaModel): # Old transformers versions check both_exist = (is_model and is_peft) and not SUPPORTS_LLAMA32 + print(both_exist) if SUPPORTS_LLAMA32: # New transformers need to check manually. files = HfFileSystem(token = token).glob(os.path.join(model_name, "*.json")) if sum(x.endswith(("adapter_config.json", "config.json")) for x in files) >= 2: both_exist = True pass + print(both_exist) pass # Error out if both LoRA and normal model config exists. From fa3b87e013a44bcd99406c8fdffcbf87f63566c7 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sun, 29 Sep 2024 23:22:05 -0700 Subject: [PATCH 43/50] Update loader.py --- unsloth/models/loader.py | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/unsloth/models/loader.py b/unsloth/models/loader.py index 9ea6c308e3..61e8132731 100644 --- a/unsloth/models/loader.py +++ b/unsloth/models/loader.py @@ -196,15 +196,14 @@ class FastLanguageModel(FastLlamaModel): # Old transformers versions check both_exist = (is_model and is_peft) and not SUPPORTS_LLAMA32 - - print(both_exist) + if SUPPORTS_LLAMA32: # New transformers need to check manually. files = HfFileSystem(token = token).glob(os.path.join(model_name, "*.json")) - if sum(x.endswith(("adapter_config.json", "config.json")) for x in files) >= 2: + files = (os.path.split(x)[-1] for x in files) + if sum(x == "adapter_config.json" or x == "config.json" for x in files) >= 2: both_exist = True pass - print(both_exist) pass # Error out if both LoRA and normal model config exists. From c8f45c8def05262554d0a50e3d0a16123fc61c6b Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Mon, 30 Sep 2024 02:09:15 -0700 Subject: [PATCH 44/50] Dependencies --- pyproject.toml | 8 ++++---- unsloth/tokenizer_utils.py | 34 +++++++++++++++++++++++++++++++++- 2 files changed, 37 insertions(+), 5 deletions(-) diff --git a/pyproject.toml b/pyproject.toml index 59ef1b8aab..1e8fd7e236 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -42,8 +42,8 @@ huggingface = [ "psutil", "wheel>=0.42.0", "numpy", - "accelerate>=0.26.1", - "trl>=0.7.9,!=0.9.0,!=0.9.1,!=0.9.2,!=0.9.3", + "accelerate>=0.34.1", + "trl>=0.7.9,!=0.9.0,!=0.9.1,!=0.9.2,!=0.9.3,<=0.11.1", "peft>=0.7.1,!=0.11.0", "protobuf<4.0.0", "huggingface_hub", @@ -224,8 +224,8 @@ colab-new = [ "hf_transfer", ] colab-no-deps = [ - "accelerate>=0.26.1", - "trl>=0.7.9,!=0.9.0,!=0.9.1,!=0.9.2,!=0.9.3", + "accelerate>=0.34.1", + "trl>=0.7.9,!=0.9.0,!=0.9.1,!=0.9.2,!=0.9.3,<=0.11.1", "peft>=0.7.1", "xformers<0.0.27", "bitsandbytes>=0.43.3", diff --git a/unsloth/tokenizer_utils.py b/unsloth/tokenizer_utils.py index 04690d3566..cdce372b50 100644 --- a/unsloth/tokenizer_utils.py +++ b/unsloth/tokenizer_utils.py @@ -1135,7 +1135,39 @@ from inspect import getsource import trl.trainer.sft_trainer from trl.trainer.sft_trainer import * from transformers.trainer import * -from trl.trainer.sft_trainer import neftune_post_forward_hook +try: + from trl.trainer.sft_trainer import neftune_post_forward_hook +except: + def neftune_post_forward_hook(module, input, output): + """ + Implements the NEFTune forward pass for the model using forward hooks. Note this works only for + torch.nn.Embedding layers. This method is slightly adapted from the original source code + that can be found here: https://github.com/neelsjain/NEFTune + + Simply add it to your model as follows: + ```python + model = ... + model.embed_tokens.neftune_noise_alpha = 0.1 + model.embed_tokens.register_forward_hook(neftune_post_forward_hook) + ``` + + Args: + module (`torch.nn.Module`): + The embedding module where the hook is attached. Note that you need to set + `module.neftune_noise_alpha` to the desired noise alpha value. + input (`torch.Tensor`): + The input tensor to the model. + output (`torch.Tensor`): + The output tensor of the model (i.e. the embeddings). + """ + if module.training: + dims = torch.tensor(output.size(1) * output.size(2)) + mag_norm = module.neftune_noise_alpha / torch.sqrt(dims) + output = output + torch.zeros_like(output).uniform_(-mag_norm, mag_norm) + return output + pass +pass + def patch_sft_trainer_tokenizer(): """ From 51475011218286c27e5715288b9d4e108997a698 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Mon, 30 Sep 2024 02:48:31 -0700 Subject: [PATCH 45/50] Update pyproject.toml --- pyproject.toml | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/pyproject.toml b/pyproject.toml index 1e8fd7e236..9499d771c8 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -35,7 +35,7 @@ exclude = ["images*"] huggingface = [ "packaging", "tyro", - "transformers>=4.45.1", + "transformers<4.45.0", "datasets>=2.16.0", "sentencepiece>=0.2.0", "tqdm", @@ -212,7 +212,7 @@ colab-ampere-torch220 = [ colab-new = [ "packaging", "tyro", - "transformers>=4.45.1", + "transformers<4.45.0", "datasets>=2.16.0", "sentencepiece>=0.2.0", "tqdm", From cd7fe775fd229f5c5a288f71e8ee77fbf33b032b Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Mon, 30 Sep 2024 02:51:32 -0700 Subject: [PATCH 46/50] Update _utils.py --- 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 309c0d913e..b14bb39144 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__ = "2024.9.post3" +__version__ = "2024.9.post4" __all__ = [ "prepare_model_for_kbit_training", From 3ac34b50718b18d912541095602234b7f9fdd274 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 1 Oct 2024 00:14:52 -0700 Subject: [PATCH 47/50] Update tokenizer_utils.py --- unsloth/tokenizer_utils.py | 37 ++++++++++++++++++++++++++++++++++++- 1 file changed, 36 insertions(+), 1 deletion(-) diff --git a/unsloth/tokenizer_utils.py b/unsloth/tokenizer_utils.py index cdce372b50..df36552b4b 100644 --- a/unsloth/tokenizer_utils.py +++ b/unsloth/tokenizer_utils.py @@ -1116,6 +1116,40 @@ def add_new_tokens( pass +@torch.inference_mode +def fix_zero_training_loss(model, tokenizer, train_dataset): + """ + Sometimes the labels get masked by all -100s, causing the loss + to be 0. We check for this! + """ + if len(train_dataset) == 0: return + + row = train_dataset[0] + if type(row) is dict and "labels" in row: + + # Check the first 100 rows + seen_bad = 0 + seen_good = 0 + for i, row in enumerate(train_dataset): + try: check_tokens = list(set(row["labels"])) + except: continue + if len(check_tokens) == 1 and check_tokens[0] == -100: seen_bad += 1 + else: seen_good += 1 + if i >= 100: break + pass + + # Check ratio + if seen_bad / (seen_bad + seen_good) >= 0.9: + logger.warning( + "Unsloth: Most labels in your dataset are -100. Training losses will be 0.\n"\ + "Are you usre you used `train_on_responses_only` correctly?\n"\ + "Or did you mask our tokens incorrectly? Maybe this is intended?" + ) + pass + pass +pass + + def check_nvidia(): # Unsloth doesn't work yet on AMD devices - we're working on it! output = np.array([0,]) @@ -1228,7 +1262,8 @@ def patch_sft_trainer_tokenizer(): " torch.cuda.empty_cache()\n"\ "pass\n"\ "\n"\ - "fix_untrained_tokens(self.model, self.tokenizer, self.train_dataset, eps = 1e-16)\n\n" + "fix_untrained_tokens(self.model, self.tokenizer, self.train_dataset, eps = 1e-16)\n\n"\ + "fix_zero_training_loss(self.model, self.tokenizer, self.train_dataset)\n\n" # Add NEFTune since it doesn't seem to work?? We need to manually inject it check_text += \ From cd530425f020bf8d1861df634a14214a215868e8 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 1 Oct 2024 00:20:01 -0700 Subject: [PATCH 48/50] Update tokenizer_utils.py --- unsloth/tokenizer_utils.py | 12 +++++------- 1 file changed, 5 insertions(+), 7 deletions(-) diff --git a/unsloth/tokenizer_utils.py b/unsloth/tokenizer_utils.py index df36552b4b..5efe610a10 100644 --- a/unsloth/tokenizer_utils.py +++ b/unsloth/tokenizer_utils.py @@ -42,15 +42,13 @@ IGNORED_TOKENIZER_CHECKING = frozenset(( IGNORED_TOKENIZER_NAMES = [ - # "unsloth/Mistral-Nemo-Instruct-2407-bnb-4bit", - # "unsloth/Mistral-Nemo-Instruct-2407", - # "mistralai/Mistral-Nemo-Instruct-2407", - # "unsloth/Mistral-Nemo-Base-2407-bnb-4bit", - # "unsloth/Mistral-Nemo-Base-2407", - # "mistralai/Mistral-Nemo-Base-2407", + # Qwen Coder did not train on tool calling. Math did! + "unsloth/Qwen2.5-Coder-1.5B-Instruct", + "unsloth/Qwen2.5-Coder-7B-Instruct", ] IGNORED_TOKENIZER_NAMES = frozenset( - [x.lower() for x in IGNORED_TOKENIZER_NAMES] + [x.lower() for x in IGNORED_TOKENIZER_NAMES] + \ + [x.lower()+"-bnb-4bit" for x in IGNORED_TOKENIZER_NAMES] ) # Check environments From 682f16aa5a7003fc1704000c3a872854f3cedd8d Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 1 Oct 2024 00:35:54 -0700 Subject: [PATCH 49/50] Update tokenizer_utils.py --- unsloth/tokenizer_utils.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/unsloth/tokenizer_utils.py b/unsloth/tokenizer_utils.py index 5efe610a10..196e496188 100644 --- a/unsloth/tokenizer_utils.py +++ b/unsloth/tokenizer_utils.py @@ -1140,7 +1140,7 @@ def fix_zero_training_loss(model, tokenizer, train_dataset): if seen_bad / (seen_bad + seen_good) >= 0.9: logger.warning( "Unsloth: Most labels in your dataset are -100. Training losses will be 0.\n"\ - "Are you usre you used `train_on_responses_only` correctly?\n"\ + "For example, are you sure you used `train_on_responses_only` correctly?\n"\ "Or did you mask our tokens incorrectly? Maybe this is intended?" ) pass From 98a4a570adb18c696ca2ed952742b2ef8df85648 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Tue, 1 Oct 2024 00:40:17 -0700 Subject: [PATCH 50/50] Update README.md --- README.md | 1 + 1 file changed, 1 insertion(+) diff --git a/README.md b/README.md index 2b8e79f021..51342fdf27 100644 --- a/README.md +++ b/README.md @@ -22,6 +22,7 @@ All notebooks are **beginner friendly**! Add your dataset, click "Run All", and | Unsloth supports | Free Notebooks | Performance | Memory use | |-----------|---------|--------|----------| +| **Llama 3.2 (3B)** | [▶️ Start for free](https://colab.research.google.com/drive/1T5-zKWM_5OD21QHwXHiV9ixTRR7k3iB9?usp=sharing) | 2x faster | 60% less | | **Llama 3.1 (8B)** | [▶️ Start for free](https://colab.research.google.com/drive/1Ys44kVvmeZtnICzWz0xgpRnrIOjZAuxp?usp=sharing) | 2x faster | 60% less | | **Phi-3.5 (mini)** | [▶️ Start for free](https://colab.research.google.com/drive/1lN6hPQveB_mHSnTOYifygFcrO8C1bxq4?usp=sharing) | 2x faster | 50% less | | **Gemma 2 (9B)** | [▶️ Start for free](https://colab.research.google.com/drive/1vIrqH5uYDQwsJ4-OO3DErvuv4pBgVwk4?usp=sharing) | 2x faster | 63% less |