Skip to content

Commit 4f4e628

Browse files
authored
Merge branch 'ggml-org:master' into master
2 parents 9232e6b + 6c442f4 commit 4f4e628

32 files changed

+409
-253
lines changed

common/arg.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -2555,15 +2555,15 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
25552555
{"--lora"}, "FNAME",
25562556
"path to LoRA adapter (can be repeated to use multiple adapters)",
25572557
[](common_params & params, const std::string & value) {
2558-
params.lora_adapters.push_back({ std::string(value), 1.0, nullptr });
2558+
params.lora_adapters.push_back({ std::string(value), 1.0, "", "", nullptr });
25592559
}
25602560
// we define this arg on both COMMON and EXPORT_LORA, so when showing help message of export-lora, it will be categorized as "example-specific" arg
25612561
).set_examples({LLAMA_EXAMPLE_COMMON, LLAMA_EXAMPLE_EXPORT_LORA}));
25622562
add_opt(common_arg(
25632563
{"--lora-scaled"}, "FNAME", "SCALE",
25642564
"path to LoRA adapter with user defined scaling (can be repeated to use multiple adapters)",
25652565
[](common_params & params, const std::string & fname, const std::string & scale) {
2566-
params.lora_adapters.push_back({ fname, std::stof(scale), nullptr });
2566+
params.lora_adapters.push_back({ fname, std::stof(scale), "", "", nullptr });
25672567
}
25682568
// we define this arg on both COMMON and EXPORT_LORA, so when showing help message of export-lora, it will be categorized as "example-specific" arg
25692569
).set_examples({LLAMA_EXAMPLE_COMMON, LLAMA_EXAMPLE_EXPORT_LORA}));

common/common.cpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -988,7 +988,12 @@ struct common_init_result common_init_from_params(common_params & params) {
988988
return iparams;
989989
}
990990

991+
char buf[1024];
991992
la.ptr = lora.get();
993+
llama_adapter_meta_val_str(la.ptr, "adapter.lora.task_name", buf, sizeof(buf));
994+
la.task_name = buf;
995+
llama_adapter_meta_val_str(la.ptr, "adapter.lora.prompt_prefix", buf, sizeof(buf));
996+
la.prompt_prefix = buf;
992997
iparams.lora.emplace_back(std::move(lora)); // copy to list of loaded adapters
993998
}
994999

common/common.h

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -34,6 +34,9 @@ struct common_adapter_lora_info {
3434
std::string path;
3535
float scale;
3636

37+
std::string task_name;
38+
std::string prompt_prefix;
39+
3740
struct llama_adapter_lora * ptr;
3841
};
3942

convert_hf_to_gguf.py

Lines changed: 77 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -72,6 +72,7 @@ class ModelBase:
7272
endianess: gguf.GGUFEndian
7373
use_temp_file: bool
7474
lazy: bool
75+
dry_run: bool
7576
part_names: list[str]
7677
is_safetensors: bool
7778
hparams: dict[str, Any]
@@ -111,6 +112,7 @@ def __init__(self, dir_model: Path, ftype: gguf.LlamaFileType, fname_out: Path,
111112
self.endianess = gguf.GGUFEndian.BIG if is_big_endian else gguf.GGUFEndian.LITTLE
112113
self.use_temp_file = use_temp_file
113114
self.lazy = not eager or (remote_hf_model_id is not None)
115+
self.dry_run = dry_run
114116
self.remote_hf_model_id = remote_hf_model_id
115117
if remote_hf_model_id is not None:
116118
self.is_safetensors = True
@@ -4871,11 +4873,35 @@ def modify_tensors(self, data_torch, name, bid):
48714873
@ModelBase.register("XLMRobertaModel", "XLMRobertaForSequenceClassification")
48724874
class XLMRobertaModel(BertModel):
48734875
model_arch = gguf.MODEL_ARCH.BERT
4876+
_lora_files = {}
4877+
_lora_names = []
48744878

4875-
def __init__(self, *args, **kwargs):
4876-
super().__init__(*args, **kwargs)
4879+
def __init__(self, dir_model: Path, ftype: gguf.LlamaFileType, fname_out: Path, **kwargs: Any):
4880+
hparams = kwargs.pop("hparams", None)
4881+
if hparams is None:
4882+
hparams = ModelBase.load_hparams(dir_model, False)
4883+
4884+
if lora_names := hparams.get("lora_adaptations"):
4885+
self._lora_names = lora_names
4886+
self.model_arch = gguf.MODEL_ARCH.JINA_BERT_V3
4887+
4888+
super().__init__(dir_model, ftype, fname_out, hparams=hparams, **kwargs)
48774889
self._xlmroberta_tokenizer_init()
48784890

4891+
def generate_extra_tensors(self) -> Iterable[tuple[str, Tensor]]:
4892+
if self._lora_names:
4893+
for name in self._lora_names:
4894+
fname = self.add_prefix_to_filename(self.fname_out, f"lora-{name}-")
4895+
self._lora_files[name] = gguf.GGUFWriter(fname, arch=gguf.MODEL_ARCH_NAMES[self.model_arch], endianess=self.endianess, use_temp_file=self.use_temp_file, dry_run=self.dry_run)
4896+
4897+
return super().generate_extra_tensors()
4898+
4899+
def set_type(self):
4900+
for lora_writer in self._lora_files.values():
4901+
lora_writer.add_type(gguf.GGUFType.ADAPTER)
4902+
lora_writer.add_string(gguf.Keys.Adapter.TYPE, "lora")
4903+
super().set_type()
4904+
48794905
def set_vocab(self):
48804906
self._xlmroberta_set_vocab()
48814907

@@ -4885,13 +4911,62 @@ def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iter
48854911
if name.startswith("roberta."):
48864912
name = name[8:]
48874913

4914+
# jina-embeddings-v3
4915+
if ".parametrizations." in name:
4916+
name = name.replace(".parametrizations.", ".")
4917+
if name.endswith(".original"):
4918+
name = name[:-9]
4919+
48884920
# position embeddings start at pad_token_id + 1, so just chop down the weight tensor
48894921
if name == "embeddings.position_embeddings.weight":
48904922
if self._position_offset is not None:
48914923
data_torch = data_torch[self._position_offset:,:]
48924924

4925+
if name.endswith(".0.lora_A") or name.endswith(".0.lora_B"):
4926+
if name.startswith("pooler.dense"):
4927+
return []
4928+
4929+
num_loras = data_torch.size(0)
4930+
assert num_loras == len(self._lora_names)
4931+
4932+
# Split out each LoRA in their own GGUF
4933+
for i, lora_writer in enumerate(self._lora_files.values()):
4934+
new_name = self.map_tensor_name(name[:-9]) + name[-7:].lower()
4935+
data = data_torch[i, :, :]
4936+
# Transpose/flip token_embd/types into correct shape
4937+
if new_name == "token_embd.weight.lora_b":
4938+
data = data.T
4939+
elif new_name.startswith("token_types.weight."):
4940+
new_name = new_name[:-1] + ("a" if new_name[-1:] == "b" else "b")
4941+
lora_writer.add_tensor(new_name, data.float().numpy(), raw_dtype=gguf.GGMLQuantizationType.F32)
4942+
4943+
return []
4944+
48934945
return super().modify_tensors(data_torch, name, bid)
48944946

4947+
def set_gguf_parameters(self):
4948+
super().set_gguf_parameters()
4949+
4950+
# jina-embeddings-v3
4951+
if rotary_emb_base := self.hparams.get("rotary_emb_base"):
4952+
self.gguf_writer.add_rope_freq_base(rotary_emb_base)
4953+
lora_alpha = self.hparams.get("lora_alpha")
4954+
if lora_prompt_prefixes := self.hparams.get("task_instructions"):
4955+
assert self._lora_files and all(lora_name in lora_prompt_prefixes for lora_name in self._lora_files.keys())
4956+
for lora_name, lora_writer in self._lora_files.items():
4957+
lora_writer.add_float32(gguf.Keys.Adapter.LORA_ALPHA, lora_alpha if lora_alpha is not None else 1.0)
4958+
lora_writer.add_string(gguf.Keys.Adapter.LORA_TASK_NAME, lora_name)
4959+
if lora_prompt_prefixes:
4960+
lora_writer.add_string(gguf.Keys.Adapter.LORA_PROMPT_PREFIX, lora_prompt_prefixes[lora_name])
4961+
4962+
def write(self):
4963+
super().write()
4964+
for lora_writer in self._lora_files.values():
4965+
lora_writer.write_header_to_file()
4966+
lora_writer.write_kv_data_to_file()
4967+
lora_writer.write_tensors_to_file(progress=True)
4968+
lora_writer.close()
4969+
48954970

48964971
@ModelBase.register("GemmaForCausalLM")
48974972
class GemmaModel(TextModel):

examples/model-conversion/Makefile

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -37,6 +37,20 @@ causal-convert-model:
3737
METADATA_OVERRIDE="$(METADATA_OVERRIDE)" \
3838
./scripts/causal/convert-model.sh
3939

40+
causal-convert-mm-model-bf16: OUTTYPE=bf16
41+
causal-convert-mm-model-bf16: MM_OUTTYPE=f16
42+
causal-convert-mm-model-bf16: causal-convert-mm-model
43+
44+
causal-convert-mm-model:
45+
$(call validate_model_path,causal-convert-mm-model)
46+
@MODEL_NAME="$(MODEL_NAME)" OUTTYPE="$(OUTTYPE)" MODEL_PATH="$(MODEL_PATH)" \
47+
METADATA_OVERRIDE="$(METADATA_OVERRIDE)" \
48+
./scripts/causal/convert-model.sh
49+
50+
@MODEL_NAME="$(MODEL_NAME)" OUTTYPE="$(MM_OUTTYPE)" MODEL_PATH="$(MODEL_PATH)" \
51+
METADATA_OVERRIDE="$(METADATA_OVERRIDE)" \
52+
./scripts/causal/convert-model.sh --mmproj
53+
4054
causal-run-original-model:
4155
$(call validate_model_path,causal-run-original-model)
4256
@MODEL_PATH="$(MODEL_PATH)" ./scripts/causal/run-org-model.py
Lines changed: 29 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,21 @@
11
#!/bin/bash
22

3+
set -e
4+
5+
# Parse command line arguments
6+
MMPROJ=""
7+
while [[ $# -gt 0 ]]; do
8+
case $1 in
9+
--mmproj)
10+
MMPROJ="--mmproj"
11+
shift
12+
;;
13+
*)
14+
shift
15+
;;
16+
esac
17+
done
18+
319
MODEL_NAME="${MODEL_NAME:-$(basename "$MODEL_PATH")}"
420
OUTPUT_DIR="${OUTPUT_DIR:-../../models}"
521
TYPE="${OUTTYPE:-f16}"
@@ -11,12 +27,20 @@ echo "Model name: ${MODEL_NAME}"
1127
echo "Data type: ${TYPE}"
1228
echo "Converted model path:: ${CONVERTED_MODEL}"
1329
echo "Metadata override: ${METADATA_OVERRIDE}"
14-
python ../../convert_hf_to_gguf.py --verbose \
15-
${MODEL_PATH} \
16-
--outfile ${CONVERTED_MODEL} \
17-
--outtype ${TYPE} \
18-
--metadata "${METADATA_OVERRIDE}"
30+
31+
CMD_ARGS=("python" "../../convert_hf_to_gguf.py" "--verbose")
32+
CMD_ARGS+=("${MODEL_PATH}")
33+
CMD_ARGS+=("--outfile" "${CONVERTED_MODEL}")
34+
CMD_ARGS+=("--outtype" "${TYPE}")
35+
[[ -n "$METADATA_OVERRIDE" ]] && CMD_ARGS+=("--metadata" "${METADATA_OVERRIDE}")
36+
[[ -n "$MMPROJ" ]] && CMD_ARGS+=("${MMPROJ}")
37+
38+
"${CMD_ARGS[@]}"
1939

2040
echo ""
2141
echo "The environment variable CONVERTED_MODEL can be set to this path using:"
2242
echo "export CONVERTED_MODEL=$(realpath ${CONVERTED_MODEL})"
43+
if [[ -n "$MMPROJ" ]]; then
44+
mmproj_file="${OUTPUT_DIR}/mmproj-$(basename "${CONVERTED_MODEL}")"
45+
echo "The mmproj model was created in $(realpath "$mmproj_file")"
46+
fi

ggml/src/ggml-cann/common.h

Lines changed: 0 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -374,7 +374,6 @@ struct ggml_backend_cann_context {
374374
#endif
375375
cann_task_queue task_queue;
376376
bool async_mode;
377-
bool support_set_rows;
378377
// Rope Cache
379378
void* rope_init_ptr = nullptr;
380379
void* rope_sin_ptr = nullptr;
@@ -400,14 +399,6 @@ struct ggml_backend_cann_context {
400399
async_mode = parse_bool(get_env("GGML_CANN_ASYNC_MODE").value_or(""));
401400
GGML_LOG_INFO("%s: device %d async operator submission is %s\n", __func__,
402401
device, async_mode ? "ON" : "OFF");
403-
404-
support_set_rows = parse_bool(get_env("LLAMA_SET_ROWS").value_or(""));
405-
GGML_LOG_INFO("%s: LLAMA_SET_ROWS is %s\n", __func__, support_set_rows ? "ON" : "OFF");
406-
407-
if (!support_set_rows) {
408-
GGML_LOG_INFO("%s: CANN Graph currently only supports execution when LLAMA_SET_ROWS is ON. "
409-
"Falling back to eager mode.\n", __func__);
410-
}
411402
}
412403

413404
/**

ggml/src/ggml-cann/ggml-cann.cpp

Lines changed: 0 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -2251,11 +2251,6 @@ static enum ggml_status ggml_backend_cann_graph_compute(
22512251
bool use_cann_graph = true;
22522252
bool cann_graph_update_required = false;
22532253

2254-
// check environment LLAMA_SET_ROWS
2255-
if (!cann_ctx->support_set_rows) {
2256-
use_cann_graph = false;
2257-
}
2258-
22592254
if (use_cann_graph) {
22602255
if (cann_ctx->cann_graph == nullptr) {
22612256
cann_ctx->cann_graph.reset(new ggml_cann_graph());

ggml/src/ggml-cpu/ggml-cpu-impl.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -489,7 +489,7 @@ inline static int16x8_t vec_padd_s16(int16x8_t a, int16x8_t b) {
489489
/**
490490
* @see https://github.com/ggml-org/llama.cpp/pull/14037
491491
*/
492-
inline float vec_hsum(float32x4_t v) {
492+
inline static float vec_hsum(float32x4_t v) {
493493
float32x4_t v_temp = v + vec_reve(v);
494494
return v_temp[0] + v_temp[1];
495495
}

ggml/src/ggml-cpu/ops.cpp

Lines changed: 11 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -9003,8 +9003,7 @@ static void ggml_compute_forward_ssm_scan_f32(
90039003
GGML_ASSERT(src4->nb[0] == sizeof(float));
90049004
GGML_ASSERT(src5->nb[0] == sizeof(float));
90059005
GGML_ASSERT(src6->nb[0] == sizeof(int32_t));
9006-
// allows optimizing the modulo since n_group should be a power of 2
9007-
GGML_ASSERT((ng & -ng) == ng);
9006+
GGML_ASSERT(nh % ng == 0);
90089007

90099008
// heads per thread
90109009
const int dh = (nh + nth - 1)/nth;
@@ -9035,6 +9034,7 @@ static void ggml_compute_forward_ssm_scan_f32(
90359034
// ref: https://github.com/state-spaces/mamba/blob/62db608da60f6fc790b8ed9f4b3225e95ca15fde/mamba_ssm/ops/triton/softplus.py#L16
90369035
const float dt_soft_plus = dt[h] <= 20.0f ? log1pf(expf(dt[h])) : dt[h];
90379036
const float dA = expf(dt_soft_plus * A[h]);
9037+
const int g = h / (nh / ng); // repeat_interleave
90389038

90399039
// dim
90409040
for (int i1 = 0; i1 < nr; ++i1) {
@@ -9057,8 +9057,8 @@ static void ggml_compute_forward_ssm_scan_f32(
90579057
// TODO: maybe unroll more?
90589058
for (int j = 0; j < 1; j++) {
90599059
GGML_F32_VEC t0 = GGML_F32_VEC_LOAD(s0 + i + j*ggml_f32_epr + ii*nc);
9060-
GGML_F32_VEC t1 = GGML_F32_VEC_LOAD(B + i + j*ggml_f32_epr + (h & (ng - 1))*nc);
9061-
GGML_F32_VEC t2 = GGML_F32_VEC_LOAD(C + i + j*ggml_f32_epr + (h & (ng - 1))*nc);
9060+
GGML_F32_VEC t1 = GGML_F32_VEC_LOAD(B + i + j*ggml_f32_epr + g*nc);
9061+
GGML_F32_VEC t2 = GGML_F32_VEC_LOAD(C + i + j*ggml_f32_epr + g*nc);
90629062

90639063
t0 = GGML_F32_VEC_MUL(t0, adA);
90649064
t1 = GGML_F32_VEC_MUL(t1, axdt);
@@ -9090,8 +9090,8 @@ static void ggml_compute_forward_ssm_scan_f32(
90909090
for (int i = 0; i < np; i += GGML_F32_STEP) {
90919091
for (int j = 0; j < GGML_F32_ARR; j++) {
90929092
ax[j] = GGML_F32_VEC_LOAD(s0 + i + j*GGML_F32_EPR + ii*nc);
9093-
ay[j] = GGML_F32_VEC_LOAD(B + i + j*GGML_F32_EPR + (h & (ng - 1))*nc);
9094-
az[j] = GGML_F32_VEC_LOAD(C + i + j*GGML_F32_EPR + (h & (ng - 1))*nc);
9093+
ay[j] = GGML_F32_VEC_LOAD(B + i + j*GGML_F32_EPR + g*nc);
9094+
az[j] = GGML_F32_VEC_LOAD(C + i + j*GGML_F32_EPR + g*nc);
90959095

90969096
ax[j] = GGML_F32_VEC_MUL(ax[j], adA);
90979097
ay[j] = GGML_F32_VEC_MUL(ay[j], axdt);
@@ -9113,7 +9113,7 @@ static void ggml_compute_forward_ssm_scan_f32(
91139113
// d_state
91149114
for (int i0 = np; i0 < nc; ++i0) {
91159115
const int i = i0 + ii*nc;
9116-
const int ig = i0 + (h & (ng - 1))*nc;
9116+
const int ig = i0 + g*nc;
91179117
// state = prev_state * dA + dB * x
91189118
const float state = (s0[i] * dA) + (B[ig] * x_dt);
91199119
// y = rowwise_dotprod(state, C)
@@ -9130,6 +9130,7 @@ static void ggml_compute_forward_ssm_scan_f32(
91309130
for (int h = ih0; h < ih1; ++h) {
91319131
// ref: https://github.com/state-spaces/mamba/blob/62db608da60f6fc790b8ed9f4b3225e95ca15fde/mamba_ssm/ops/triton/softplus.py#L16
91329132
const float dt_soft_plus = dt[h] <= 20.0f ? log1pf(expf(dt[h])) : dt[h];
9133+
const int g = h / (nh / ng); // repeat_interleave
91339134

91349135
// dim
91359136
for (int i1 = 0; i1 < nr; ++i1) {
@@ -9144,8 +9145,8 @@ static void ggml_compute_forward_ssm_scan_f32(
91449145
// TODO: what happens when (d_state % svcntw()) != 0?
91459146
for (int64_t k = 0; k < nc; k += svcntw()) {
91469147
svfloat32_t vA = GGML_F32_VEC_LOAD(&A[h*nc + k]);
9147-
svfloat32_t vB = GGML_F32_VEC_LOAD(&B[k + (h & (ng - 1))*nc]);
9148-
svfloat32_t vC = GGML_F32_VEC_LOAD(&C[k + (h & (ng - 1))*nc]);
9148+
svfloat32_t vB = GGML_F32_VEC_LOAD(&B[k + g*nc]);
9149+
svfloat32_t vC = GGML_F32_VEC_LOAD(&C[k + g*nc]);
91499150
svfloat32_t vs0 = GGML_F32_VEC_LOAD(&s0[ii*nc + k]);
91509151

91519152
svfloat32_t t1 = GGML_F32_VEC_MUL(vdt_soft_plus, vA);
@@ -9165,7 +9166,7 @@ static void ggml_compute_forward_ssm_scan_f32(
91659166
// d_state
91669167
for (int i0 = 0; i0 < nc; ++i0) {
91679168
const int i = i0 + ii*nc;
9168-
const int ig = i0 + (h & (ng - 1))*nc;
9169+
const int ig = i0 + g*nc;
91699170
// state = prev_state * dA + dB * x
91709171
const float state = (s0[i] * expf(dt_soft_plus * A[i0 + h*nc])) + (B[ig] * x_dt);
91719172
// y = rowwise_dotprod(state, C)

0 commit comments

Comments
 (0)