Add bitnet1.58 with custom metal kernel (#219)
* add bitnet * update activation to relu2 * working bitnet * remove artifacts * remove logging * add custom post quant * fix dtype and add compile * fixed weight unpack * add custom kernel to avoid memory overhead * compile relu2 * fix weight scale * remove unused * add tests and update tuner utils * update acknowledgements * add kernel caching * add act_quant and set float16 as default dtype * use mx.add and move scaling to kernel * remove act quant * move bitlinear layers to separate file * feat: add falcon-e and other bitnet support * refactor: address comments * add support for 1.58bit N-bit quants * 43.85% speedup in generation performance (M3 max) * refactor utils * remove masking (2% gen speed improvement) * add quantization config * test llama bitnet * refactor apply_hf_quant * default threadgroup: 64 -> 32 * add comment * fix prompt processing perf * remove modulo * compile kernel in the constructor * Improve the bitnet kernel * remove benchmark * refactor bitlinear swap * format * remove llama changes * revert utils * faster + cleanup * not trainable * fix tests --------- Co-authored-by: younesbelkada <younes.belkada@tii.ae> Co-authored-by: Angelos Katharopoulos <a_katharopoulos@apple.com> Co-authored-by: Awni Hannun <awni@apple.com>
This commit is contained in:
Prince Canuma
+1
-1
@@ -8,5 +8,5 @@ with a short description of your contribution(s) below. For example:
|
||||
MLX LM was developed with contributions from the following individuals:
|
||||
|
||||
- Shunta Saito: Added support for PLaMo models.
|
||||
- Prince Canuma: Helped add support for `Starcoder2` models.
|
||||
- Gökdeniz Gülmez: Added support for the following architectures: OpenBMB's `MiniCPM` and `MiniCPM3`, Kyutai's `Helium`, State-Space's`Mamba v1`, Z.ai & THUKEG's `GLM4`, Rednote `dots.llm1`, and Allenai's `OLMoE`; Added support for the following training algorithms: `full-fine-tuning`; Added support for the following other features: `Multiple Optimizers to choose for training`, and `reporting training metrics to WandB (Weights & Biases)`.
|
||||
- Prince Canuma: Helped add support for the following model architectures: HuggingFace's `Starcoder2`, Cohere's `Cohere (1 and 2)`, Alibaba Qwen's `Qwen (2, 3 and MoE)`, Microsoft's `Phi (3 and 3.5 MoE)`, `BitNet1.58`, Meta's `Llama (3 and 4)`, Google DeepMind's `Gemma 3`, and InterLM's `InternLM 2.5`.
|
||||
|
||||
@@ -0,0 +1,131 @@
|
||||
# Copyright © 2025 Apple Inc.
|
||||
|
||||
import mlx.core as mx
|
||||
import mlx.nn as nn
|
||||
from mlx.nn.layers.quantized import QuantizedLinear
|
||||
|
||||
|
||||
def make_bitlinear_kernel():
|
||||
"""
|
||||
Custom Metal kernel that performs matrix multiplication directly on
|
||||
packed weights and scales the output. This eliminates the need to
|
||||
store unpacked weights in memory.
|
||||
"""
|
||||
source = """
|
||||
constexpr int M = 4;
|
||||
constexpr int BLOCK = 32;
|
||||
|
||||
uint tid = thread_position_in_grid.y;
|
||||
uint in_offset = thread_position_in_grid.x;
|
||||
|
||||
uint batch_idx = tid / (out_features / 4);
|
||||
uint row_idx = tid % (out_features / 4);
|
||||
|
||||
float sum[4] = {0.0};
|
||||
|
||||
for (uint i = in_offset * M; i < in_features; i += BLOCK * M) {
|
||||
float v[M];
|
||||
for (int j=0; j<M; j++) {
|
||||
v[j] = x[batch_idx * in_features + i + j];
|
||||
}
|
||||
|
||||
for (int j=0; j<M; j++) {
|
||||
uint8_t w = packed_weights[row_idx * in_features + i + j];
|
||||
sum[0] += v[j] * ((w & 3) - 1);
|
||||
sum[1] += v[j] * (((w >> 2) & 3) - 1);
|
||||
sum[2] += v[j] * (((w >> 4) & 3) - 1);
|
||||
sum[3] += v[j] * (((w >> 6) & 3) - 1);
|
||||
}
|
||||
}
|
||||
|
||||
for (int j=0; j<4; j++) {
|
||||
sum[j] = simd_sum(sum[j]);
|
||||
}
|
||||
|
||||
// Apply weight scaling by diving them or multiplying them
|
||||
if (in_offset == 0) {
|
||||
float scale = invert_weight_scales ? 1 / weight_scale[0] : weight_scale[0];
|
||||
for (int i=0; i<4; i++) {
|
||||
out[batch_idx * out_features + row_idx + i * (out_features/4)] = static_cast<T>(sum[i] * scale);
|
||||
}
|
||||
}
|
||||
"""
|
||||
|
||||
return mx.fast.metal_kernel(
|
||||
name="bitlinear_matmul",
|
||||
input_names=["x", "packed_weights", "weight_scale"],
|
||||
output_names=["out"],
|
||||
source=source,
|
||||
)
|
||||
|
||||
|
||||
_bitlinear_kernel = make_bitlinear_kernel()
|
||||
|
||||
|
||||
class BitLinear(nn.Module):
|
||||
"""
|
||||
BitLinear module with memory-efficient weight handling.
|
||||
"""
|
||||
|
||||
def __init__(
|
||||
self,
|
||||
in_features,
|
||||
out_features,
|
||||
bias=True,
|
||||
invert_weight_scales=False,
|
||||
):
|
||||
super().__init__()
|
||||
self.in_features = in_features
|
||||
self.out_features = out_features
|
||||
# Calculate packed dimensions - the first dimension gets packed 4:1
|
||||
# The weights are ternary so can be represented with 2 bits, and they
|
||||
# are packed in uint8 tensors, hence the number of values per item is 4
|
||||
packed_out_features = (out_features + 3) // 4
|
||||
self.weight = mx.zeros((packed_out_features, in_features), dtype=mx.uint8)
|
||||
|
||||
self.invert_weight_scales = invert_weight_scales
|
||||
self.weight_scale = mx.array([1.0])
|
||||
|
||||
if bias:
|
||||
self.bias = mx.zeros((out_features,))
|
||||
else:
|
||||
self.bias = None
|
||||
|
||||
def execute_matmul_kernel(self, x, packed_weights):
|
||||
original_shape = x.shape
|
||||
if len(original_shape) > 2:
|
||||
x = x.reshape(-1, original_shape[-1])
|
||||
total_batch_elements, in_features = x.shape
|
||||
|
||||
out_features = self.out_features
|
||||
|
||||
dtype = self.weight_scale.dtype
|
||||
assert x.dtype == dtype, "Wrong type for input."
|
||||
out = _bitlinear_kernel(
|
||||
inputs=[
|
||||
x,
|
||||
packed_weights,
|
||||
self.weight_scale,
|
||||
],
|
||||
template=[
|
||||
("T", dtype),
|
||||
("invert_weight_scales", self.invert_weight_scales),
|
||||
("in_features", in_features),
|
||||
("out_features", out_features),
|
||||
],
|
||||
grid=(32, total_batch_elements * out_features // 4, 1),
|
||||
threadgroup=(32, 1, 1), # SIMD width is 32 threads
|
||||
output_shapes=[(total_batch_elements, out_features)],
|
||||
output_dtypes=[dtype],
|
||||
)[0]
|
||||
|
||||
if len(original_shape) > 2:
|
||||
out = out.reshape(*original_shape[:-1], out_features)
|
||||
return out
|
||||
|
||||
def __call__(self, x):
|
||||
y = self.execute_matmul_kernel(x, self.weight)
|
||||
|
||||
if self.bias is not None:
|
||||
y = mx.add(y, self.bias)
|
||||
return y
|
||||
@@ -0,0 +1,215 @@
|
||||
# Copyright © 2023-2024 Apple Inc.
|
||||
|
||||
from dataclasses import dataclass
|
||||
from functools import partial
|
||||
from typing import Any, Dict, Optional, Union
|
||||
|
||||
import mlx.core as mx
|
||||
import mlx.nn as nn
|
||||
|
||||
from .base import BaseModelArgs, create_attention_mask, scaled_dot_product_attention
|
||||
from .bitlinear_layers import BitLinear
|
||||
from .rope_utils import initialize_rope
|
||||
|
||||
|
||||
@dataclass
|
||||
class ModelArgs(BaseModelArgs):
|
||||
model_type: str
|
||||
hidden_size: int
|
||||
num_hidden_layers: int
|
||||
intermediate_size: int
|
||||
num_attention_heads: int
|
||||
num_key_value_heads: int
|
||||
rms_norm_eps: float
|
||||
vocab_size: int
|
||||
head_dim: Optional[int] = None
|
||||
max_position_embeddings: Optional[int] = None
|
||||
attention_bias: bool = False
|
||||
mlp_bias: bool = False
|
||||
rope_theta: float = 10000
|
||||
rope_traditional: bool = False
|
||||
rope_scaling: Optional[Dict[str, Union[float, str]]] = None
|
||||
tie_word_embeddings: bool = True
|
||||
|
||||
|
||||
class Attention(nn.Module):
|
||||
def __init__(self, args: ModelArgs):
|
||||
super().__init__()
|
||||
|
||||
dim = args.hidden_size
|
||||
self.n_heads = n_heads = args.num_attention_heads
|
||||
self.n_kv_heads = n_kv_heads = args.num_key_value_heads
|
||||
|
||||
self.head_dim = head_dim = args.head_dim or args.hidden_size // n_heads
|
||||
|
||||
self.scale = head_dim**-0.5
|
||||
attention_bias = args.attention_bias
|
||||
|
||||
self.q_proj = BitLinear(dim, n_heads * head_dim, bias=attention_bias)
|
||||
self.k_proj = BitLinear(dim, n_kv_heads * head_dim, bias=attention_bias)
|
||||
self.v_proj = BitLinear(dim, n_kv_heads * head_dim, bias=attention_bias)
|
||||
self.o_proj = BitLinear(n_heads * head_dim, dim, bias=attention_bias)
|
||||
|
||||
self.rope = initialize_rope(
|
||||
self.head_dim,
|
||||
args.rope_theta,
|
||||
args.rope_traditional,
|
||||
args.rope_scaling,
|
||||
args.max_position_embeddings,
|
||||
)
|
||||
self.attn_sub_norm = nn.RMSNorm(args.hidden_size, eps=args.rms_norm_eps)
|
||||
|
||||
def __call__(
|
||||
self,
|
||||
x: mx.array,
|
||||
mask: Optional[mx.array] = None,
|
||||
cache: Optional[Any] = None,
|
||||
) -> mx.array:
|
||||
B, L, D = x.shape
|
||||
|
||||
queries, keys, values = self.q_proj(x), self.k_proj(x), self.v_proj(x)
|
||||
|
||||
# Prepare the queries, keys and values for the attention computation
|
||||
queries = queries.reshape(B, L, self.n_heads, -1).transpose(0, 2, 1, 3)
|
||||
keys = keys.reshape(B, L, self.n_kv_heads, -1).transpose(0, 2, 1, 3)
|
||||
values = values.reshape(B, L, self.n_kv_heads, -1).transpose(0, 2, 1, 3)
|
||||
|
||||
if cache is not None:
|
||||
queries = self.rope(queries, offset=cache.offset)
|
||||
keys = self.rope(keys, offset=cache.offset)
|
||||
keys, values = cache.update_and_fetch(keys, values)
|
||||
else:
|
||||
queries = self.rope(queries)
|
||||
keys = self.rope(keys)
|
||||
|
||||
output = scaled_dot_product_attention(
|
||||
queries, keys, values, cache=cache, scale=self.scale, mask=mask
|
||||
)
|
||||
|
||||
output = output.transpose(0, 2, 1, 3).reshape(B, L, -1)
|
||||
output = self.attn_sub_norm(output)
|
||||
output = self.o_proj(output)
|
||||
|
||||
return output
|
||||
|
||||
|
||||
@partial(mx.compile, shapeless=True)
|
||||
def relu2(x):
|
||||
return mx.square(nn.relu(x))
|
||||
|
||||
|
||||
class MLP(nn.Module):
|
||||
def __init__(self, args: ModelArgs):
|
||||
super().__init__()
|
||||
|
||||
dim = args.hidden_size
|
||||
hidden_dim = args.intermediate_size
|
||||
if hasattr(args, "mlp_bias"):
|
||||
mlp_bias = args.mlp_bias
|
||||
else:
|
||||
mlp_bias = False
|
||||
|
||||
self.gate_proj = BitLinear(dim, hidden_dim, bias=mlp_bias)
|
||||
self.down_proj = BitLinear(hidden_dim, dim, bias=mlp_bias)
|
||||
self.up_proj = BitLinear(dim, hidden_dim, bias=mlp_bias)
|
||||
self.ffn_sub_norm = nn.RMSNorm(args.intermediate_size, eps=args.rms_norm_eps)
|
||||
|
||||
def __call__(self, x) -> mx.array:
|
||||
x = relu2(self.gate_proj(x)) * self.up_proj(x)
|
||||
x = self.ffn_sub_norm(x)
|
||||
x = self.down_proj(x)
|
||||
return x
|
||||
|
||||
|
||||
class TransformerBlock(nn.Module):
|
||||
def __init__(self, args: ModelArgs):
|
||||
super().__init__()
|
||||
self.num_attention_heads = args.num_attention_heads
|
||||
self.hidden_size = args.hidden_size
|
||||
self.self_attn = Attention(args)
|
||||
self.mlp = MLP(args)
|
||||
self.input_layernorm = nn.RMSNorm(args.hidden_size, eps=args.rms_norm_eps)
|
||||
self.post_attention_layernorm = nn.RMSNorm(
|
||||
args.hidden_size, eps=args.rms_norm_eps
|
||||
)
|
||||
|
||||
def __call__(
|
||||
self,
|
||||
x: mx.array,
|
||||
mask: Optional[mx.array] = None,
|
||||
cache: Optional[Any] = None,
|
||||
) -> mx.array:
|
||||
|
||||
r = self.self_attn(self.input_layernorm(x), mask, cache)
|
||||
h = x + r
|
||||
r = self.mlp(self.post_attention_layernorm(h))
|
||||
out = h + r
|
||||
return out
|
||||
|
||||
|
||||
class LlamaModel(nn.Module):
|
||||
def __init__(self, args: ModelArgs):
|
||||
super().__init__()
|
||||
self.args = args
|
||||
self.vocab_size = args.vocab_size
|
||||
self.num_hidden_layers = args.num_hidden_layers
|
||||
self.embed_tokens = nn.Embedding(args.vocab_size, args.hidden_size)
|
||||
self.layers = [
|
||||
TransformerBlock(args=args) for _ in range(args.num_hidden_layers)
|
||||
]
|
||||
self.norm = nn.RMSNorm(args.hidden_size, eps=args.rms_norm_eps)
|
||||
|
||||
def __call__(
|
||||
self,
|
||||
inputs: mx.array,
|
||||
mask: mx.array = None,
|
||||
cache=None,
|
||||
):
|
||||
h = self.embed_tokens(inputs)
|
||||
|
||||
if mask is None:
|
||||
mask = create_attention_mask(h, cache)
|
||||
|
||||
if cache is None:
|
||||
cache = [None] * len(self.layers)
|
||||
|
||||
for layer, c in zip(self.layers, cache):
|
||||
h = layer(h, mask, cache=c)
|
||||
|
||||
return self.norm(h)
|
||||
|
||||
|
||||
class Model(nn.Module):
|
||||
def __init__(self, args: ModelArgs):
|
||||
super().__init__()
|
||||
self.args = args
|
||||
self.model_type = args.model_type
|
||||
self.model = LlamaModel(args)
|
||||
if not args.tie_word_embeddings:
|
||||
self.lm_head = nn.Linear(args.hidden_size, args.vocab_size, bias=False)
|
||||
|
||||
def __call__(
|
||||
self,
|
||||
inputs: mx.array,
|
||||
mask: mx.array = None,
|
||||
cache=None,
|
||||
):
|
||||
out = self.model(inputs, mask, cache)
|
||||
if self.args.tie_word_embeddings:
|
||||
out = self.model.embed_tokens.as_linear(out)
|
||||
else:
|
||||
out = self.lm_head(out)
|
||||
return out
|
||||
|
||||
def sanitize(self, weights):
|
||||
# Remove unused precomputed rotary freqs
|
||||
weights = {
|
||||
k: v for k, v in weights.items() if "self_attn.rotary_emb.inv_freq" not in k
|
||||
}
|
||||
if self.args.tie_word_embeddings:
|
||||
weights.pop("lm_head.weight", None)
|
||||
return weights
|
||||
|
||||
@property
|
||||
def layers(self):
|
||||
return self.model.layers
|
||||
@@ -251,6 +251,24 @@ class TestModels(unittest.TestCase):
|
||||
model, args.model_type, args.vocab_size, args.num_hidden_layers
|
||||
)
|
||||
|
||||
def test_bitnet(self):
|
||||
from mlx_lm.models import bitnet
|
||||
|
||||
args = bitnet.ModelArgs(
|
||||
model_type="bitnet",
|
||||
hidden_size=1024,
|
||||
num_hidden_layers=4,
|
||||
intermediate_size=2048,
|
||||
num_attention_heads=4,
|
||||
num_key_value_heads=4,
|
||||
rms_norm_eps=1e-5,
|
||||
vocab_size=10_000,
|
||||
)
|
||||
model = bitnet.Model(args)
|
||||
self.model_test_runner(
|
||||
model, args.model_type, args.vocab_size, args.num_hidden_layers
|
||||
)
|
||||
|
||||
def test_phi2(self):
|
||||
from mlx_lm.models import phi
|
||||
|
||||
|
||||
Reference in New Issue
Block a user