From 1e8922af2b9943fbba6e5592baf188936bd723bc Mon Sep 17 00:00:00 2001 From: Daniel Han-Chen Date: Sun, 10 Mar 2024 17:42:41 +1100 Subject: [PATCH] Revert --- unsloth/kernels/rms_layernorm.py | 2 +- unsloth/kernels/swiglu.py | 6 +++--- unsloth/save.py | 10 +++++----- 3 files changed, 9 insertions(+), 9 deletions(-) diff --git a/unsloth/kernels/rms_layernorm.py b/unsloth/kernels/rms_layernorm.py index 2adedff663..4db89b7816 100644 --- a/unsloth/kernels/rms_layernorm.py +++ b/unsloth/kernels/rms_layernorm.py @@ -41,7 +41,7 @@ def _rms_layernorm_forward( r += row_idx * r_row_stride 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) + W_row = tl.load(W + col_offsets, mask = mask, other = 0)#.to(tl.float32) row_var = tl.sum(X_row * X_row, axis = 0) / n_cols inv_var = tl.math.rsqrt(row_var + eps) diff --git a/unsloth/kernels/swiglu.py b/unsloth/kernels/swiglu.py index 4f4fbe9ee9..ff6b162680 100644 --- a/unsloth/kernels/swiglu.py +++ b/unsloth/kernels/swiglu.py @@ -25,7 +25,7 @@ def _fg_kernel(e, g, h, n_elements, BLOCK_SIZE : tl.constexpr,): mask = offsets < n_elements e_row = tl.load(e + offsets, mask = mask, other = 0).to(tl.float32) - g_row = tl.load(g + offsets, mask = mask, other = 0).to(tl.float32) + g_row = tl.load(g + offsets, mask = mask, other = 0)#.to(tl.float32) # f = e * sigmoid(e) f_row = e_row * tl.sigmoid(e_row) # e_row / (1 + tl.exp(-e_row)) @@ -63,9 +63,9 @@ def _DWf_DW_dfg_kernel(DW, e, g, n_elements, BLOCK_SIZE : tl.constexpr,): offsets = block_idx*BLOCK_SIZE + tl.arange(0, BLOCK_SIZE) mask = offsets < n_elements - DW_row = tl.load(DW + offsets, mask = mask, other = 0).to(tl.float32) + DW_row = tl.load(DW + offsets, mask = mask, other = 0)#.to(tl.float32) e_row = tl.load(e + offsets, mask = mask, other = 0).to(tl.float32) - g_row = tl.load(g + offsets, mask = mask, other = 0).to(tl.float32) + g_row = tl.load(g + offsets, mask = mask, other = 0)#.to(tl.float32) # e = e.float() # se = 1.0 / (1.0 + torch.exp(-e)) diff --git a/unsloth/save.py b/unsloth/save.py index 2a5cf79599..5971d76e6e 100644 --- a/unsloth/save.py +++ b/unsloth/save.py @@ -90,13 +90,13 @@ def _merge_lora(layer, name): W = fast_dequantize(W, quant_state) else: dtype = W.dtype - # W = W.to(torch.float32).t() - W = W.t() + W = W.to(torch.float32).t() + # W = W.t() if A is not None: - sAB = (A.t().to(W.dtype) @ (s * B.t().to(W.dtype))) - W += sAB - # W.addmm_(A.t().to(torch.float32), B.t().to(torch.float32), alpha = s) + # sAB = (A.t().to(torch.float32) @ (s * B.t().to(torch.float32))) + # W += sAB + W.addmm_(A.t().to(torch.float32), B.t().to(torch.float32), alpha = s) # W.addmm_(A.t().to(W.dtype), B.t().to(W.dtype), alpha = s) # if not torch.isfinite(W).all(): maximum_element = torch.max(W.min().abs(), W.max())