Merge branch 'main' into nightly
This commit is contained in:
Daniel Han
commit
8f95022daa
5 changed files with 315 additions and 61 deletions
22
README.md
22
README.md
|
|
@ -23,7 +23,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/github/unslothai/notebooks/blob/main/nb/Llama3.2_(1B_and_3B)-Conversational.ipynb) | 2x faster | 60% less |
|
||||
| **Phi-4** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Phi_4-Conversational.ipynb) | 2x faster | 50% less |
|
||||
| **Phi-4 (14B)** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Phi_4-Conversational.ipynb) | 2x faster | 50% less |
|
||||
| **Llama 3.2 Vision (11B)** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Llama3.2_(11B)-Vision.ipynb) | 2x faster | 40% less |
|
||||
| **Llama 3.1 (8B)** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Llama3.1_(8B)-Alpaca.ipynb) | 2x faster | 60% less |
|
||||
| **Gemma 2 (9B)** | [▶️ Start for free](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Gemma2_(9B)-Alpaca.ipynb) | 2x faster | 63% less |
|
||||
|
|
@ -36,27 +36,27 @@ All notebooks are **beginner friendly**! Add your dataset, click "Run All", and
|
|||
- See [all our notebooks](https://docs.unsloth.ai/get-started/unsloth-notebooks) and [all our models](https://docs.unsloth.ai/get-started/all-our-models)
|
||||
- **Kaggle Notebooks** for [Llama 3.2 Kaggle notebook](https://www.kaggle.com/danielhanchen/kaggle-llama-3-2-1b-3b-unsloth-notebook), [Llama 3.1 (8B)](https://www.kaggle.com/danielhanchen/kaggle-llama-3-1-8b-unsloth-notebook), [Gemma 2 (9B)](https://www.kaggle.com/code/danielhanchen/kaggle-gemma-7b-unsloth-notebook/), [Mistral (7B)](https://www.kaggle.com/code/danielhanchen/kaggle-mistral-7b-unsloth-notebook)
|
||||
- Run notebooks for [Llama 3.2 conversational](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Llama3.2_(1B_and_3B)-Conversational.ipynb), [Llama 3.1 conversational](https://colab.research.google.com/drive/15OyFkGoCImV9dSsewU1wa2JuKB4-mDE_?usp=sharing) and [Mistral v0.3 ChatML](https://colab.research.google.com/drive/15F1xyn8497_dUbxZP4zWmPZ3PJx1Oymv?usp=sharing)
|
||||
- This [text completion notebook](https://colab.research.google.com/drive/1ef-tab5bhkvWmBOObepl1WgJvfvSzn5Q?usp=sharing) is for continued pretraining / raw text
|
||||
- This [text completion notebook](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Mistral_(7B)-Text_Completion.ipynb) is for continued pretraining / raw text
|
||||
- This [continued pretraining notebook](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Mistral_v0.3_(7B)-CPT.ipynb) is for learning another language
|
||||
- Click [here](https://docs.unsloth.ai/) for detailed documentation for Unsloth.
|
||||
|
||||
## 🦥 Unsloth.ai News
|
||||
- 📣 NEW! [DeepSeek-R1](https://unsloth.ai/blog/deepseek-r1) - the most powerful open reasoning models with Llama & Qwen distillations. Run or fine-tune them now! More details: [unsloth.ai/blog/deepseek-r1](https://unsloth.ai/blog/deepseek-r1). All model uploads: [here](https://huggingface.co/collections/unsloth/deepseek-r1-all-versions-678e1c48f5d2fce87892ace5).
|
||||
- 📣 NEW! [Phi-4](https://unsloth.ai/blog/phi4) by Microsoft is now supported. We also [fixed bugs](https://unsloth.ai/blog/phi4) in Phi-4 and [uploaded GGUFs, 4-bit](https://huggingface.co/collections/unsloth/phi-4-all-versions-677eecf93784e61afe762afa). Try the [Phi-4 Colab notebook](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Phi_4-Conversational.ipynb)
|
||||
- 📣 NEW! [Llama 3.3 (70B)](https://huggingface.co/collections/unsloth/llama-33-all-versions-67535d7d994794b9d7cf5e9f), Meta's latest model is supported.
|
||||
- 📣 NEW! We worked with Apple to add [Cut Cross Entropy](https://arxiv.org/abs/2411.09009). Unsloth now supports 89K context for Meta's Llama 3.3 (70B) on a 80GB GPU - 13x longer than HF+FA2. For Llama 3.1 (8B), Unsloth enables 342K context, surpassing its native 128K support.
|
||||
- 📣 NEW! Introducing Unsloth [Dynamic 4-bit Quantization](https://unsloth.ai/blog/dynamic-4bit)! We dynamically opt not to quantize certain parameters and this greatly increases accuracy while only using <10% more VRAM than BnB 4-bit. See our collection on [Hugging Face here.](https://huggingface.co/collections/unsloth/unsloth-4-bit-dynamic-quants-67503bb873f89e15276c44e7)
|
||||
- 📣 NEW! [Vision models](https://unsloth.ai/blog/vision) now supported! [Llama 3.2 Vision (11B)](https://colab.research.google.com/drive/1j0N4XTY1zXXy7mPAhOC1_gMYZ2F2EBlk?usp=sharing), [Qwen 2.5 VL (7B)](https://colab.research.google.com/drive/1whHb54GNZMrNxIsi2wm2EY_-Pvo2QyKh?usp=sharing) and [Pixtral (12B) 2409](https://colab.research.google.com/drive/1K9ZrdwvZRE96qGkCq_e88FgV3MLnymQq?usp=sharing)
|
||||
- 📣 NEW! Qwen-2.5 including [Coder](https://colab.research.google.com/drive/18sN803sU23XuJV9Q8On2xgqHSer6-UZF?usp=sharing) models are now supported with bugfixes. 14b fits in a Colab GPU! [Qwen 2.5 conversational notebook](https://colab.research.google.com/drive/1qN1CEalC70EO1wGKhNxs1go1W9So61R5?usp=sharing)
|
||||
- 📣 NEW! We found and helped fix a [gradient accumulation bug](https://unsloth.ai/blog/gradient)! Please update Unsloth and transformers.
|
||||
- 📣 Introducing Unsloth [Dynamic 4-bit Quantization](https://unsloth.ai/blog/dynamic-4bit)! We dynamically opt not to quantize certain parameters and this greatly increases accuracy while only using <10% more VRAM than BnB 4-bit. See our collection on [Hugging Face here.](https://huggingface.co/collections/unsloth/unsloth-4-bit-dynamic-quants-67503bb873f89e15276c44e7)
|
||||
- 📣 [Vision models](https://unsloth.ai/blog/vision) now supported! [Llama 3.2 Vision (11B)](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Llama3.2_(11B)-Vision.ipynb), [Qwen 2.5 VL (7B)](https://colab.research.google.com/drive/1whHb54GNZMrNxIsi2wm2EY_-Pvo2QyKh?usp=sharing) and [Pixtral (12B) 2409](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Qwen2_VL_(7B)-Vision.ipynb)
|
||||
<details>
|
||||
<summary>Click for more news</summary>
|
||||
|
||||
- 📣 Try out [Chat interface](https://colab.research.google.com/drive/1i-8ESvtLRGNkkUQQr_-z_rcSAIo9c3lM?usp=sharing)!
|
||||
- 📣 NEW! [Mistral Small 22b notebook](https://colab.research.google.com/drive/1oCEHcED15DzL8xXGU1VTx5ZfOJM8WY01?usp=sharing) finetuning fits in under 16GB of VRAM!
|
||||
- 📣 NEW! [Llama 3.1 8b, 70b](https://colab.research.google.com/drive/1Ys44kVvmeZtnICzWz0xgpRnrIOjZAuxp?usp=sharing) & [Mistral Nemo-12b](https://colab.research.google.com/drive/17d3U-CAIwzmbDRqbZ9NnpHxCkmXB6LZ0?usp=sharing) both Base and Instruct are now supported
|
||||
- 📣 We found and helped fix a [gradient accumulation bug](https://unsloth.ai/blog/gradient)! Please update Unsloth and transformers.
|
||||
- 📣 Try out [Chat interface](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Unsloth_Studio.ipynb)!
|
||||
- 📣 NEW! Qwen-2.5 including [Coder](https://unsloth.ai/blog/qwen-coder) models are now supported with bugfixes. 14b fits in a Colab GPU! [Qwen 2.5 conversational notebook](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Qwen2.5_Coder_(14B)-Conversational.ipynb)
|
||||
- 📣 NEW! [Mistral Small 22b notebook](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Mistral_Small_(22B)-Alpaca.ipynb) finetuning fits in under 16GB of VRAM!
|
||||
- 📣 NEW! `pip install unsloth` now works! Head over to [pypi](https://pypi.org/project/unsloth/) to check it out! This allows non git pull installs. Use `pip install unsloth[colab-new]` for non dependency installs.
|
||||
- 📣 NEW! Continued Pretraining [notebook](https://colab.research.google.com/drive/1tEd1FrOXWMnCU9UIvdYhs61tkxdMuKZu?usp=sharing) for other languages like Korean!
|
||||
- 📣 [2x faster inference](https://colab.research.google.com/drive/1aqlNQi7MMJbynFDyOQteD2t0yVfjb9Zh?usp=sharing) added for all our models
|
||||
- 📣 NEW! Continued Pretraining [notebook](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Mistral_v0.3_(7B)-CPT.ipynb) for other languages like Korean!
|
||||
- 📣 [2x faster inference](https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Llama3.1_(8B)-Inference.ipynb) added for all our models
|
||||
- 📣 We cut memory usage by a [further 30%](https://unsloth.ai/blog/long-context) and now support [4x longer context windows](https://unsloth.ai/blog/long-context)!
|
||||
</details>
|
||||
|
||||
|
|
|
|||
|
|
@ -134,65 +134,80 @@ else:
|
|||
torch.cuda.is_bf16_supported = is_bf16_supported
|
||||
pass
|
||||
|
||||
# For Gradio HF Spaces?
|
||||
# if "SPACE_AUTHOR_NAME" not in os.environ and "SPACE_REPO_NAME" not in os.environ:
|
||||
import triton
|
||||
libcuda_dirs = lambda: None
|
||||
if Version(triton.__version__) >= Version("3.0.0"):
|
||||
try: from triton.backends.nvidia.driver import libcuda_dirs
|
||||
except: pass
|
||||
else: from triton.common.build import libcuda_dirs
|
||||
|
||||
def fix_triton_ops():
|
||||
# Check if triton.ops exists
|
||||
try:
|
||||
import triton.ops
|
||||
except:
|
||||
# Triton 3.2 removed triton.ops
|
||||
from .kernels.matmul_perf_model import (
|
||||
early_config_prune,
|
||||
estimate_matmul_time,
|
||||
)
|
||||
triton.ops.early_config_prune = early_config_prune
|
||||
triton.ops.estimate_matmul_time = estimate_matmul_time
|
||||
pass
|
||||
pass
|
||||
fix_triton_ops()
|
||||
|
||||
# Try loading bitsandbytes and triton
|
||||
import bitsandbytes as bnb
|
||||
try:
|
||||
cdequantize_blockwise_fp32 = bnb.functional.lib.cdequantize_blockwise_fp32
|
||||
libcuda_dirs()
|
||||
except:
|
||||
warnings.warn(
|
||||
"Unsloth: Running `ldconfig /usr/lib64-nvidia` to link CUDA."\
|
||||
)
|
||||
|
||||
if "SPACE_AUTHOR_NAME" not in os.environ and "SPACE_REPO_NAME" not in os.environ:
|
||||
if os.path.exists("/usr/lib64-nvidia"):
|
||||
os.system("ldconfig /usr/lib64-nvidia")
|
||||
elif os.path.exists("/usr/local"):
|
||||
# Sometimes bitsandbytes cannot be linked properly in Runpod for example
|
||||
possible_cudas = subprocess.check_output(["ls", "-al", "/usr/local"]).decode("utf-8").split("\n")
|
||||
find_cuda = re.compile(r"[\s](cuda\-[\d\.]{2,})$")
|
||||
possible_cudas = [find_cuda.search(x) for x in possible_cudas]
|
||||
possible_cudas = [x.group(1) for x in possible_cudas if x is not None]
|
||||
|
||||
import triton
|
||||
libcuda_dirs = lambda: None
|
||||
if Version(triton.__version__) >= Version("3.0.0"):
|
||||
try: from triton.backends.nvidia.driver import libcuda_dirs
|
||||
except: pass
|
||||
else: from triton.common.build import libcuda_dirs
|
||||
# Try linking cuda folder, or everything in local
|
||||
if len(possible_cudas) == 0:
|
||||
os.system("ldconfig /usr/local/")
|
||||
else:
|
||||
find_number = re.compile(r"([\d\.]{2,})")
|
||||
latest_cuda = np.argsort([float(find_number.search(x).group(1)) for x in possible_cudas])[::-1][0]
|
||||
latest_cuda = possible_cudas[latest_cuda]
|
||||
os.system(f"ldconfig /usr/local/{latest_cuda}")
|
||||
pass
|
||||
|
||||
importlib.reload(bnb)
|
||||
importlib.reload(triton)
|
||||
try:
|
||||
libcuda_dirs = lambda: None
|
||||
if Version(triton.__version__) >= Version("3.0.0"):
|
||||
try: from triton.backends.nvidia.driver import libcuda_dirs
|
||||
except: pass
|
||||
else: from triton.common.build import libcuda_dirs
|
||||
cdequantize_blockwise_fp32 = bnb.functional.lib.cdequantize_blockwise_fp32
|
||||
libcuda_dirs()
|
||||
fix_triton_ops()
|
||||
except:
|
||||
warnings.warn(
|
||||
"Unsloth: Running `ldconfig /usr/lib64-nvidia` to link CUDA."\
|
||||
"Unsloth: CUDA is not linked properly.\n"\
|
||||
"Try running `python -m bitsandbytes` then `python -m xformers.info`\n"\
|
||||
"We tried running `ldconfig /usr/lib64-nvidia` ourselves, but it didn't work.\n"\
|
||||
"You need to run in your terminal `sudo ldconfig /usr/lib64-nvidia` yourself, then import Unsloth.\n"\
|
||||
"Also try `sudo ldconfig /usr/local/cuda-xx.x` - find the latest cuda version.\n"\
|
||||
"Unsloth will still run for now, but maybe it might crash - let's hope it works!"
|
||||
)
|
||||
|
||||
if os.path.exists("/usr/lib64-nvidia"):
|
||||
os.system("ldconfig /usr/lib64-nvidia")
|
||||
elif os.path.exists("/usr/local"):
|
||||
# Sometimes bitsandbytes cannot be linked properly in Runpod for example
|
||||
possible_cudas = subprocess.check_output(["ls", "-al", "/usr/local"]).decode("utf-8").split("\n")
|
||||
find_cuda = re.compile(r"[\s](cuda\-[\d\.]{2,})$")
|
||||
possible_cudas = [find_cuda.search(x) for x in possible_cudas]
|
||||
possible_cudas = [x.group(1) for x in possible_cudas if x is not None]
|
||||
|
||||
# Try linking cuda folder, or everything in local
|
||||
if len(possible_cudas) == 0:
|
||||
os.system("ldconfig /usr/local/")
|
||||
else:
|
||||
find_number = re.compile(r"([\d\.]{2,})")
|
||||
latest_cuda = np.argsort([float(find_number.search(x).group(1)) for x in possible_cudas])[::-1][0]
|
||||
latest_cuda = possible_cudas[latest_cuda]
|
||||
os.system(f"ldconfig /usr/local/{latest_cuda}")
|
||||
pass
|
||||
|
||||
importlib.reload(bnb)
|
||||
importlib.reload(triton)
|
||||
try:
|
||||
libcuda_dirs = lambda: None
|
||||
if Version(triton.__version__) >= Version("3.0.0"):
|
||||
try: from triton.backends.nvidia.driver import libcuda_dirs
|
||||
except: pass
|
||||
else: from triton.common.build import libcuda_dirs
|
||||
cdequantize_blockwise_fp32 = bnb.functional.lib.cdequantize_blockwise_fp32
|
||||
libcuda_dirs()
|
||||
except:
|
||||
warnings.warn(
|
||||
"Unsloth: CUDA is not linked properly.\n"\
|
||||
"Try running `python -m bitsandbytes` then `python -m xformers.info`\n"\
|
||||
"We tried running `ldconfig /usr/lib64-nvidia` ourselves, but it didn't work.\n"\
|
||||
"You need to run in your terminal `sudo ldconfig /usr/lib64-nvidia` yourself, then import Unsloth.\n"\
|
||||
"Also try `sudo ldconfig /usr/local/cuda-xx.x` - find the latest cuda version.\n"\
|
||||
"Unsloth will still run for now, but maybe it might crash - let's hope it works!"
|
||||
)
|
||||
pass
|
||||
pass
|
||||
|
||||
# Check for unsloth_zoo
|
||||
|
|
|
|||
211
unsloth/kernels/matmul_perf_model.py
Normal file
211
unsloth/kernels/matmul_perf_model.py
Normal file
|
|
@ -0,0 +1,211 @@
|
|||
# Adapted from https://github.com/triton-lang/kernels/blob/eeeebdd8be7d13629de22d600621e6234057eed3/kernels/matmul_perf_model.py
|
||||
# https://github.com/triton-lang/kernels is licensed under the MIT License.
|
||||
|
||||
import functools
|
||||
import heapq
|
||||
|
||||
import torch
|
||||
|
||||
from triton import cdiv
|
||||
from triton.runtime import driver
|
||||
from triton.testing import (
|
||||
get_dram_gbps,
|
||||
get_max_simd_tflops,
|
||||
get_max_tensorcore_tflops,
|
||||
nvsmi,
|
||||
)
|
||||
|
||||
|
||||
@functools.lru_cache
|
||||
def get_clock_rate_in_khz():
|
||||
try:
|
||||
return nvsmi(["clocks.max.sm"])[0] * 1e3
|
||||
except FileNotFoundError:
|
||||
import pynvml
|
||||
|
||||
pynvml.nvmlInit()
|
||||
handle = pynvml.nvmlDeviceGetHandleByIndex(0)
|
||||
return pynvml.nvmlDeviceGetMaxClockInfo(handle, pynvml.NVML_CLOCK_SM) * 1e3
|
||||
|
||||
|
||||
def get_tensorcore_tflops(device, num_ctas, num_warps, dtype):
|
||||
"""return compute throughput in TOPS"""
|
||||
total_warps = num_ctas * min(num_warps, 4)
|
||||
num_subcores = driver.active.utils.get_device_properties(device)["multiprocessor_count"] * 4 # on recent GPUs
|
||||
tflops = (
|
||||
min(num_subcores, total_warps)
|
||||
/ num_subcores
|
||||
* get_max_tensorcore_tflops(dtype, get_clock_rate_in_khz(), device)
|
||||
)
|
||||
return tflops
|
||||
|
||||
|
||||
def get_simd_tflops(device, num_ctas, num_warps, dtype):
|
||||
"""return compute throughput in TOPS"""
|
||||
total_warps = num_ctas * min(num_warps, 4)
|
||||
num_subcores = driver.active.utils.get_device_properties(device)["multiprocessor_count"] * 4 # on recent GPUs
|
||||
tflops = (
|
||||
min(num_subcores, total_warps) / num_subcores * get_max_simd_tflops(dtype, get_clock_rate_in_khz(), device)
|
||||
)
|
||||
return tflops
|
||||
|
||||
|
||||
def get_tflops(device, num_ctas, num_warps, dtype):
|
||||
capability = torch.cuda.get_device_capability(device)
|
||||
if capability[0] < 8 and dtype == torch.float32:
|
||||
return get_simd_tflops(device, num_ctas, num_warps, dtype)
|
||||
return get_tensorcore_tflops(device, num_ctas, num_warps, dtype)
|
||||
|
||||
|
||||
def estimate_matmul_time(
|
||||
# backend, device,
|
||||
num_warps,
|
||||
num_stages, #
|
||||
A,
|
||||
B,
|
||||
C, #
|
||||
M,
|
||||
N,
|
||||
K, #
|
||||
BLOCK_M,
|
||||
BLOCK_N,
|
||||
BLOCK_K,
|
||||
SPLIT_K, #
|
||||
debug=False,
|
||||
**kwargs, #
|
||||
):
|
||||
"""return estimated running time in ms
|
||||
= max(compute, loading) + store"""
|
||||
device = torch.cuda.current_device()
|
||||
dtype = A.dtype
|
||||
dtsize = A.element_size()
|
||||
|
||||
num_cta_m = cdiv(M, BLOCK_M)
|
||||
num_cta_n = cdiv(N, BLOCK_N)
|
||||
num_cta_k = SPLIT_K
|
||||
num_ctas = num_cta_m * num_cta_n * num_cta_k
|
||||
|
||||
# If the input is smaller than the block size
|
||||
M, N = max(M, BLOCK_M), max(N, BLOCK_N)
|
||||
|
||||
# time to compute
|
||||
total_ops = 2 * M * N * K / (1024 * 1024 * 1024) # GOPS
|
||||
tput = get_tflops(device, num_ctas, num_warps, dtype)
|
||||
compute_ms = total_ops / tput
|
||||
|
||||
# time to load data
|
||||
num_sm = driver.active.utils.get_device_properties(device)["multiprocessor_count"]
|
||||
active_cta_ratio = min(1, num_ctas / num_sm)
|
||||
active_cta_ratio_bw1 = min(1, num_ctas / 32) # 32 active ctas are enough to saturate
|
||||
active_cta_ratio_bw2 = max(min(1, (num_ctas - 32) / (108 - 32)), 0) # 32-108, remaining 5%
|
||||
dram_bw = get_dram_gbps(device) * (active_cta_ratio_bw1 * 0.95 + active_cta_ratio_bw2 * 0.05) # in GB/s
|
||||
l2_bw = dram_bw * 4 # rough estimation (should be 4.7 for A100?)
|
||||
# assume 80% of (following) loads are in L2 cache
|
||||
load_a_dram = M * K * dtsize * (1 + 0.2 * (num_cta_n - 1))
|
||||
load_a_l2 = M * K * dtsize * 0.8 * (num_cta_n - 1)
|
||||
load_b_dram = N * K * dtsize * (1 + 0.2 * (num_cta_m - 1))
|
||||
load_b_l2 = N * K * dtsize * 0.8 * (num_cta_m - 1)
|
||||
# total
|
||||
total_dram = (load_a_dram + load_b_dram) / (1024 * 1024) # MB
|
||||
total_l2 = (load_a_l2 + load_b_l2) / (1024 * 1024)
|
||||
# loading time in ms
|
||||
load_ms = total_dram / dram_bw + total_l2 / l2_bw
|
||||
|
||||
# estimate storing time
|
||||
store_bw = dram_bw * 0.6 # :o
|
||||
store_c_dram = M * N * dtsize * SPLIT_K / (1024 * 1024) # MB
|
||||
if SPLIT_K == 1:
|
||||
store_ms = store_c_dram / store_bw
|
||||
else:
|
||||
reduce_bw = store_bw
|
||||
store_ms = store_c_dram / reduce_bw
|
||||
# c.zero_()
|
||||
zero_ms = M * N * 2 / (1024 * 1024) / store_bw
|
||||
store_ms += zero_ms
|
||||
|
||||
total_time_ms = max(compute_ms, load_ms) + store_ms
|
||||
if debug:
|
||||
print(
|
||||
f"Total time: {total_time_ms}ms, compute time: {compute_ms}ms, "
|
||||
f"loading time: {load_ms}ms, store time: {store_ms}ms, "
|
||||
f"Activate CTAs: {active_cta_ratio*100}%"
|
||||
)
|
||||
return total_time_ms
|
||||
|
||||
|
||||
def early_config_prune(configs, named_args, **kwargs):
|
||||
device = torch.cuda.current_device()
|
||||
capability = torch.cuda.get_device_capability()
|
||||
# BLOCK_M, BLOCK_N, BLOCK_K, SPLIT_K, num_warps, num_stages
|
||||
dtsize = named_args["A"].element_size()
|
||||
dtype = named_args["A"].dtype
|
||||
|
||||
# 1. make sure we have enough smem
|
||||
pruned_configs = []
|
||||
for config in configs:
|
||||
kw = config.kwargs
|
||||
BLOCK_M, BLOCK_N, BLOCK_K, num_stages = (
|
||||
kw["BLOCK_M"],
|
||||
kw["BLOCK_N"],
|
||||
kw["BLOCK_K"],
|
||||
config.num_stages,
|
||||
)
|
||||
|
||||
max_shared_memory = driver.active.utils.get_device_properties(device)["max_shared_mem"]
|
||||
required_shared_memory = (BLOCK_M + BLOCK_N) * BLOCK_K * num_stages * dtsize
|
||||
if required_shared_memory <= max_shared_memory:
|
||||
pruned_configs.append(config)
|
||||
configs = pruned_configs
|
||||
|
||||
# Some dtypes do not allow atomic_add
|
||||
if dtype not in [torch.float16, torch.float32]:
|
||||
configs = [config for config in configs if config.kwargs["SPLIT_K"] == 1]
|
||||
|
||||
# group configs by (BLOCK_M,_N,_K, SPLIT_K, num_warps)
|
||||
configs_map = {}
|
||||
for config in configs:
|
||||
kw = config.kwargs
|
||||
BLOCK_M, BLOCK_N, BLOCK_K, SPLIT_K, num_warps, num_stages = (
|
||||
kw["BLOCK_M"],
|
||||
kw["BLOCK_N"],
|
||||
kw["BLOCK_K"],
|
||||
kw["SPLIT_K"],
|
||||
config.num_warps,
|
||||
config.num_stages,
|
||||
)
|
||||
|
||||
key = (BLOCK_M, BLOCK_N, BLOCK_K, SPLIT_K, num_warps)
|
||||
if key in configs_map:
|
||||
configs_map[key].append((config, num_stages))
|
||||
else:
|
||||
configs_map[key] = [(config, num_stages)]
|
||||
|
||||
pruned_configs = []
|
||||
for k, v in configs_map.items():
|
||||
BLOCK_M, BLOCK_N, BLOCK_K, SPLIT_K, num_warps = k
|
||||
if capability[0] >= 8:
|
||||
# compute cycles (only works for ampere GPUs)
|
||||
mmas = BLOCK_M * BLOCK_N * BLOCK_K / (16 * 8 * 16)
|
||||
mma_cycles = mmas / min(4, num_warps) * 8
|
||||
|
||||
ldgsts_latency = 300 # Does this matter?
|
||||
optimal_num_stages = ldgsts_latency / mma_cycles
|
||||
|
||||
# nearest stages, prefer large #stages
|
||||
nearest = heapq.nsmallest(
|
||||
2,
|
||||
v,
|
||||
key=lambda x: (
|
||||
10 + abs(x[1] - optimal_num_stages)
|
||||
if (x[1] - optimal_num_stages) < 0
|
||||
else x[1] - optimal_num_stages
|
||||
),
|
||||
)
|
||||
|
||||
for n in nearest:
|
||||
pruned_configs.append(n[0])
|
||||
else: # Volta & Turing only supports num_stages <= 2
|
||||
random_config = v[0][0]
|
||||
random_config.num_stages = 2
|
||||
pruned_configs.append(random_config)
|
||||
return pruned_configs
|
||||
|
|
@ -12,7 +12,7 @@
|
|||
# See the License for the specific language governing permissions and
|
||||
# limitations under the License.
|
||||
|
||||
__version__ = "2025.1.6"
|
||||
__version__ = "2025.1.7"
|
||||
|
||||
__all__ = [
|
||||
"SUPPORTS_BFLOAT16",
|
||||
|
|
|
|||
|
|
@ -522,6 +522,34 @@ __INT_TO_FLOAT_MAPPER = \
|
|||
"microsoft/phi-4",
|
||||
"unsloth/phi-4-bnb-4bit",
|
||||
),
|
||||
"unsloth/DeepSeek-R1-Distill-Qwen-32B-bnb-4bit" : (
|
||||
"unsloth/DeepSeek-R1-Distill-Qwen-32B",
|
||||
"deepseek-ai/DeepSeek-R1-Distill-Qwen-32B",
|
||||
),
|
||||
"unsloth/DeepSeek-R1-Distill-Qwen-14B-unsloth-bnb-4bit" : (
|
||||
"unsloth/DeepSeek-R1-Distill-Qwen-14B",
|
||||
"deepseek-ai/DeepSeek-R1-Distill-Qwen-14B",
|
||||
"unsloth/DeepSeek-R1-Distill-Qwen-14B-bnb-4bit",
|
||||
),
|
||||
"unsloth/DeepSeek-R1-Distill-Qwen-7B-unsloth-bnb-4bit" : (
|
||||
"unsloth/DeepSeek-R1-Distill-Qwen-7B",
|
||||
"deepseek-ai/DeepSeek-R1-Distill-Qwen-7B",
|
||||
"unsloth/DeepSeek-R1-Distill-Qwen-7B-bnb-4bit",
|
||||
),
|
||||
"unsloth/DeepSeek-R1-Distill-Qwen-1.5B-unsloth-bnb-4bit" : (
|
||||
"unsloth/DeepSeek-R1-Distill-Qwen-1.5B",
|
||||
"deepseek-ai/DeepSeek-R1-Distill-Qwen-1.5B",
|
||||
"unsloth/DeepSeek-R1-Distill-Qwen-1.5B-bnb-4bit",
|
||||
),
|
||||
"unsloth/DeepSeek-R1-Distill-Llama-8B-unsloth-bnb-4bit" : (
|
||||
"unsloth/DeepSeek-R1-Distill-Llama-8B",
|
||||
"deepseek-ai/DeepSeek-R1-Distill-Llama-8B",
|
||||
"unsloth/DeepSeek-R1-Distill-Llama-8B-bnb-4bit",
|
||||
),
|
||||
"unsloth/DeepSeek-R1-Distill-Llama-70B-bnb-4bit" : (
|
||||
"unsloth/DeepSeek-R1-Distill-Llama-70B",
|
||||
"deepseek-ai/DeepSeek-R1-Distill-Llama-70B",
|
||||
),
|
||||
}
|
||||
|
||||
INT_TO_FLOAT_MAPPER = {}
|
||||
|
|
|
|||
Loading…
Add table
Add a link
Reference in a new issue