Skip to content

Commit c6ab72d

Browse files
jploskicebtenzzre
authored andcommitted
llm : add MPT support (ggml-org#3417)
Co-authored-by: Cebtenzzre <cebtenzzre@gmail.com> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> (cherry picked from commit f5f9121)
1 parent a63ffba commit c6ab72d

File tree

3 files changed

+700
-25
lines changed

3 files changed

+700
-25
lines changed

convert-mpt-hf-to-gguf.py

+216
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,216 @@
1+
#!/usr/bin/env python3
2+
# HF mpt--> gguf conversion
3+
4+
from __future__ import annotations
5+
6+
import argparse
7+
import json
8+
import os
9+
import struct
10+
import sys
11+
from pathlib import Path
12+
from typing import Any
13+
14+
import numpy as np
15+
import torch
16+
from transformers import AutoTokenizer # type: ignore[import]
17+
18+
if 'NO_LOCAL_GGUF' not in os.environ:
19+
sys.path.insert(1, str(Path(__file__).parent / 'gguf-py' / 'gguf'))
20+
import gguf
21+
22+
23+
def count_model_parts(dir_model: Path) -> int:
24+
num_parts = 0
25+
for filename in os.listdir(dir_model):
26+
if filename.startswith("pytorch_model-"):
27+
num_parts += 1
28+
29+
if num_parts > 0:
30+
print("gguf: found " + str(num_parts) + " model parts")
31+
return num_parts
32+
33+
34+
def parse_args() -> argparse.Namespace:
35+
parser = argparse.ArgumentParser(description="Convert an MPT model to a GGML compatible file")
36+
parser.add_argument(
37+
"--vocab-only", action="store_true",
38+
help="extract only the vocab",
39+
)
40+
parser.add_argument(
41+
"--outfile", type=Path,
42+
help="path to write to; default: based on input",
43+
)
44+
parser.add_argument(
45+
"model", type=Path,
46+
help="directory containing model file, or model file itself (*.bin)",
47+
)
48+
parser.add_argument(
49+
"ftype", type=int, choices=[0, 1], default=1, nargs='?',
50+
help="output format - use 0 for float32, 1 for float16",
51+
)
52+
return parser.parse_args()
53+
54+
args = parse_args()
55+
56+
dir_model = args.model
57+
ftype = args.ftype
58+
if not dir_model.is_dir():
59+
print(f'Error: {args.model} is not a directory', file = sys.stderr)
60+
sys.exit(1)
61+
62+
# possible tensor data types
63+
# ftype == 0 -> float32
64+
# ftype == 1 -> float16
65+
66+
# map from ftype to string
67+
ftype_str = ["f32", "f16"]
68+
69+
if args.outfile is not None:
70+
fname_out = args.outfile
71+
else:
72+
# output in the same directory as the model by default
73+
fname_out = dir_model / f'ggml-model-{ftype_str[ftype]}.gguf'
74+
75+
print("gguf: loading model "+dir_model.name)
76+
77+
with open(dir_model / "config.json", "r", encoding="utf-8") as f:
78+
hparams = json.load(f)
79+
80+
if hparams["architectures"][0] != "MPTForCausalLM":
81+
print("Model architecture not supported: " + hparams["architectures"][0])
82+
83+
sys.exit()
84+
85+
# get number of model parts
86+
num_parts = count_model_parts(dir_model)
87+
88+
ARCH=gguf.MODEL_ARCH.MPT
89+
gguf_writer = gguf.GGUFWriter(fname_out, gguf.MODEL_ARCH_NAMES[ARCH])
90+
91+
print("gguf: get model metadata")
92+
93+
block_count = hparams["n_layers"]
94+
95+
gguf_writer.add_name(dir_model.name)
96+
gguf_writer.add_context_length(hparams["max_seq_len"])
97+
gguf_writer.add_embedding_length(hparams["d_model"])
98+
gguf_writer.add_block_count(block_count)
99+
gguf_writer.add_feed_forward_length(4 * hparams["d_model"])
100+
gguf_writer.add_head_count(hparams["n_heads"])
101+
gguf_writer.add_layer_norm_eps(1e-05)
102+
if hparams["attn_config"]["clip_qkv"] is not None:
103+
gguf_writer.add_clamp_kqv(hparams["attn_config"]["clip_qkv"])
104+
gguf_writer.add_max_alibi_bias(hparams["attn_config"]["alibi_bias_max"])
105+
106+
# TOKENIZATION
107+
108+
print("gguf: get tokenizer metadata")
109+
110+
tokens: list[bytearray] = []
111+
scores: list[float] = []
112+
toktypes: list[int] = []
113+
114+
# gpt2 tokenizer
115+
gguf_writer.add_tokenizer_model("gpt2")
116+
117+
print("gguf: get gpt2 tokenizer vocab")
118+
119+
# MPT token embedding tensors have dimension 50432 (hparams["vocab_size"]), but
120+
# there are only 50254 (len(tokenizer.vocab)) tokens in the vocab, presumably to
121+
# accomodate some "reserved" tokens; this is causing problems down the line in
122+
# llama.cpp, so we pad the vocab with dummy tokens:
123+
124+
vocab_size = hparams["vocab_size"]
125+
126+
# ref: https://github.com/cmp-nct/ggllm.cpp/blob/master/falcon_convert.py
127+
tokenizer = AutoTokenizer.from_pretrained(dir_model)
128+
129+
reverse_vocab = {id: encoded_tok for encoded_tok, id in tokenizer.vocab.items()}
130+
131+
for i in range(vocab_size):
132+
tokens.append(reverse_vocab[i] if i in reverse_vocab else f"[PAD{i}]")
133+
scores.append(0.0) # dummy
134+
toktypes.append(gguf.TokenType.NORMAL)
135+
136+
gguf_writer.add_token_list(tokens)
137+
gguf_writer.add_token_scores(scores)
138+
gguf_writer.add_token_types(toktypes)
139+
140+
special_vocab = gguf.SpecialVocab(dir_model, load_merges = True)
141+
special_vocab.add_to_gguf(gguf_writer)
142+
143+
# TENSORS
144+
145+
tensor_map = gguf.get_tensor_name_map(ARCH,block_count)
146+
147+
# tensor info
148+
print("gguf: get tensor metadata")
149+
150+
if num_parts == 0:
151+
part_names = iter(("pytorch_model.bin",))
152+
else:
153+
part_names = (
154+
f"pytorch_model-{n:05}-of-{num_parts:05}.bin" for n in range(1, num_parts + 1)
155+
)
156+
157+
for part_name in part_names:
158+
if args.vocab_only:
159+
break
160+
print("gguf: loading model part '" + part_name + "'")
161+
model_part = torch.load(f"{dir_model}/{part_name}", map_location="cpu")
162+
163+
for name in model_part.keys():
164+
data = model_part[name]
165+
166+
old_dtype = data.dtype
167+
168+
# convert any unsupported data types to float32
169+
if data.dtype != torch.float16 and data.dtype != torch.float32:
170+
data = data.to(torch.float32)
171+
172+
data = data.squeeze().numpy()
173+
174+
# map tensor names
175+
new_name = tensor_map.get_name(name, try_suffixes = (".weight", ".bias"))
176+
if new_name is None:
177+
print("Cannot map tensor '" + name + "'")
178+
continue # for the sake of compatibility with some old published models, don't quit
179+
sys.exit()
180+
181+
n_dims = len(data.shape)
182+
data_dtype = data.dtype
183+
184+
# if f32 desired, convert any float16 to float32
185+
if ftype == 0 and data_dtype == np.float16:
186+
data = data.astype(np.float32)
187+
188+
# TODO: Why cant we use these float16 as-is? There should be not reason to store float16 as float32
189+
if ftype == 1 and data_dtype == np.float16 and n_dims == 1:
190+
data = data.astype(np.float32)
191+
192+
# if f16 desired, convert any float32 2-dim weight tensors to float16
193+
if ftype == 1 and data_dtype == np.float32 and name.endswith(".weight") and n_dims == 2:
194+
data = data.astype(np.float16)
195+
196+
print(new_name + ", n_dims = " + str(n_dims) + ", " + str(old_dtype) + " --> " + str(data.dtype))
197+
198+
gguf_writer.add_tensor(new_name, data)
199+
200+
# note: MPT output is tied to (same as) wte in original model;
201+
# for easier implementation in llama.cpp it's duplicated in GGUF, though :/
202+
if new_name == "token_embd.weight":
203+
gguf_writer.add_tensor("output.weight", data)
204+
205+
print("gguf: write header")
206+
gguf_writer.write_header_to_file()
207+
print("gguf: write metadata")
208+
gguf_writer.write_kv_data_to_file()
209+
if not args.vocab_only:
210+
print("gguf: write tensors")
211+
gguf_writer.write_tensors_to_file()
212+
213+
gguf_writer.close()
214+
215+
print(f"gguf: model successfully exported to '{fname_out}'")
216+
print("")

ggml-cuda.cu

+46-2
Original file line numberDiff line numberDiff line change
@@ -413,6 +413,7 @@ static_assert(sizeof(block_q6_K) == sizeof(ggml_fp16_t) + 13*QK_K/16, "wrong q6_
413413
#define CUDA_SILU_BLOCK_SIZE 256
414414
#define CUDA_CPY_BLOCK_SIZE 32
415415
#define CUDA_SCALE_BLOCK_SIZE 256
416+
#define CUDA_CLAMP_BLOCK_SIZE 256
416417
#define CUDA_ROPE_BLOCK_SIZE 256
417418
#define CUDA_ALIBI_BLOCK_SIZE 32
418419
#define CUDA_DIAG_MASK_INF_BLOCK_SIZE 32
@@ -4538,6 +4539,16 @@ static __global__ void scale_f32(const float * x, float * dst, const float scale
45384539
dst[i] = scale * x[i];
45394540
}
45404541

4542+
static __global__ void clamp_f32(const float * x, float * dst, const float min, const float max, const int k) {
4543+
const int i = blockDim.x*blockIdx.x + threadIdx.x;
4544+
4545+
if (i >= k) {
4546+
return;
4547+
}
4548+
4549+
dst[i] = x[i] < min ? min : (x[i] > max ? max : x[i]);
4550+
}
4551+
45414552
static void add_f32_cuda(const float * x, const float * y, float * dst, const int kx, const int ky, cudaStream_t stream) {
45424553
const int num_blocks = (kx + CUDA_ADD_BLOCK_SIZE - 1) / CUDA_ADD_BLOCK_SIZE;
45434554
add_f32<<<num_blocks, CUDA_ADD_BLOCK_SIZE, 0, stream>>>(x, y, dst, kx, ky);
@@ -5389,6 +5400,11 @@ static void scale_f32_cuda(const float * x, float * dst, const float scale, cons
53895400
scale_f32<<<num_blocks, CUDA_SCALE_BLOCK_SIZE, 0, stream>>>(x, dst, scale, k);
53905401
}
53915402

5403+
static void clamp_f32_cuda(const float * x, float * dst, const float min, const float max, const int k, cudaStream_t stream) {
5404+
const int num_blocks = (k + CUDA_CLAMP_BLOCK_SIZE - 1) / CUDA_CLAMP_BLOCK_SIZE;
5405+
clamp_f32<<<num_blocks, CUDA_CLAMP_BLOCK_SIZE, 0, stream>>>(x, dst, min, max, k);
5406+
}
5407+
53925408
static void rope_f32_cuda(const float * x, float * dst, const int ncols, const int nrows, const float p0,
53935409
const float p_delta, const int p_delta_rows, const float theta_scale, cudaStream_t stream) {
53945410
GGML_ASSERT(ncols % 2 == 0);
@@ -6186,12 +6202,12 @@ inline void ggml_cuda_op_alibi(
61866202
const int64_t ne02 = src0->ne[2];
61876203
const int64_t nrows = ggml_nrows(src0);
61886204

6189-
const int n_past = ((int32_t *) dst->op_params)[0];
6205+
//const int n_past = ((int32_t *) dst->op_params)[0];
61906206
const int n_head = ((int32_t *) dst->op_params)[1];
61916207
float max_bias;
61926208
memcpy(&max_bias, (int32_t *) dst->op_params + 2, sizeof(float));
61936209

6194-
GGML_ASSERT(ne01 + n_past == ne00);
6210+
//GGML_ASSERT(ne01 + n_past == ne00);
61956211
GGML_ASSERT(n_head == ne02);
61966212

61976213
const int n_heads_log2_floor = 1 << (int) floor(log2(n_head));
@@ -6260,6 +6276,24 @@ inline void ggml_cuda_op_scale(
62606276
(void) src1_dd;
62616277
}
62626278

6279+
inline void ggml_cuda_op_clamp(
6280+
const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst,
6281+
const float * src0_dd, const float * src1_dd, float * dst_dd, const cudaStream_t & main_stream) {
6282+
6283+
GGML_ASSERT(src0->type == GGML_TYPE_F32);
6284+
GGML_ASSERT( dst->type == GGML_TYPE_F32);
6285+
6286+
const float min = ((float *) dst->op_params)[0];
6287+
const float max = ((float *) dst->op_params)[1];
6288+
6289+
clamp_f32_cuda(src0_dd, dst_dd, min, max, ggml_nelements(src0), main_stream);
6290+
CUDA_CHECK(cudaGetLastError());
6291+
6292+
(void) src1;
6293+
(void) dst;
6294+
(void) src1_dd;
6295+
}
6296+
62636297
static void ggml_cuda_op_flatten(const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, const ggml_cuda_op_flatten_t op) {
62646298
const int64_t nrows0 = ggml_nrows(src0);
62656299

@@ -6817,6 +6851,10 @@ void ggml_cuda_scale(const ggml_tensor * src0, const ggml_tensor * src1, ggml_te
68176851
ggml_cuda_op_flatten(src0, src1, dst, ggml_cuda_op_scale);
68186852
}
68196853

6854+
static void ggml_cuda_clamp(const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
6855+
ggml_cuda_op_flatten(src0, src1, dst, ggml_cuda_op_clamp);
6856+
}
6857+
68206858
void ggml_cuda_cpy(const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
68216859
const int64_t ne = ggml_nelements(src0);
68226860
GGML_ASSERT(ne == ggml_nelements(src1));
@@ -7225,6 +7263,12 @@ bool ggml_cuda_compute_forward(struct ggml_compute_params * params, struct ggml_
72257263
}
72267264
func = ggml_cuda_scale;
72277265
break;
7266+
case GGML_OP_CLAMP:
7267+
if (!any_on_device) {
7268+
return false;
7269+
}
7270+
func = ggml_cuda_clamp;
7271+
break;
72287272
case GGML_OP_CPY:
72297273
if (!any_on_device) {
72307274
return false;

0 commit comments

Comments
 (0)