Skip to content

Commit 5ff0500

Browse files
tdakhrantdakhran
committed
ASR with LFM2-Audio-1.5B
1 parent 96a181a commit 5ff0500

16 files changed

Lines changed: 677 additions & 29 deletions

File tree

‎common/arg.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1172,7 +1172,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
11721172
[](common_params & params, const std::string & value) {
11731173
params.system_prompt = value;
11741174
}
1175-
).set_examples({LLAMA_EXAMPLE_COMPLETION, LLAMA_EXAMPLE_CLI, LLAMA_EXAMPLE_DIFFUSION}));
1175+
).set_examples({LLAMA_EXAMPLE_COMPLETION, LLAMA_EXAMPLE_CLI, LLAMA_EXAMPLE_DIFFUSION, LLAMA_EXAMPLE_MTMD}));
11761176
add_opt(common_arg(
11771177
{"--perf"},
11781178
{"--no-perf"},

‎convert_hf_to_gguf.py‎

Lines changed: 77 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -719,6 +719,9 @@ def load_hparams(dir_model: Path, is_mistral_format: bool):
719719
if "thinker_config" in config:
720720
# rename for Qwen2.5-Omni
721721
config["text_config"] = config["thinker_config"]["text_config"]
722+
if "lfm" in config:
723+
# rename for LFM2-Audio
724+
config["text_config"] = config["lfm"]
722725
return config
723726

724727
@classmethod
@@ -9563,19 +9566,25 @@ def set_gguf_parameters(self):
95639566
self._add_feed_forward_length()
95649567

95659568
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
9566-
is_vision_tensor = "vision_tower" in name or "multi_modal_projector" in name
9567-
if is_vision_tensor:
9568-
# skip vision tensors
9569+
if self._is_vision_tensor(name) or self._is_audio_tensor(name):
9570+
# skip multimodal tensors
95699571
return []
95709572

9571-
name = name.replace("language_model.", "")
9573+
name = name.replace("language_model.", "") # vision
9574+
name = name.replace("lfm.", "model.") # audio
95729575

95739576
# conv op requires 2d tensor
95749577
if 'conv.conv' in name:
95759578
data_torch = data_torch.squeeze(1)
95769579

95779580
return [(self.map_tensor_name(name), data_torch)]
95789581

9582+
def _is_vision_tensor(self, name: str) -> bool:
9583+
return "vision_tower" in name or "multi_modal_projector" in name
9584+
9585+
def _is_audio_tensor(self, name: str):
9586+
return any(p in name for p in ["audio", "codebook", "conformer", "depth_embedding", "depthformer", "depth_linear"])
9587+
95799588

95809589
@ModelBase.register("Lfm2MoeForCausalLM")
95819590
class LFM2MoeModel(TextModel):
@@ -9681,6 +9690,70 @@ def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iter
96819690
return [] # skip other tensors
96829691

96839692

9693+
@ModelBase.register("Lfm2AudioForConditionalGeneration")
9694+
class LFM2AudioModel(MmprojModel):
9695+
has_vision_encoder = False
9696+
has_audio_encoder = True
9697+
model_name = "Lfm2AudioEncoder"
9698+
9699+
_batch_norm_tensors: list[dict[str, Tensor]] | None = None
9700+
9701+
def get_audio_config(self) -> dict[str, Any] | None:
9702+
return self.global_config.get("encoder")
9703+
9704+
def set_gguf_parameters(self):
9705+
self.hparams_audio["hidden_size"] = self.hparams_audio["d_model"]
9706+
self.hparams_audio["intermediate_size"] = self.hparams_audio["d_model"]
9707+
self.hparams_audio["num_attention_heads"] = self.hparams_audio["n_heads"]
9708+
super().set_gguf_parameters()
9709+
self.gguf_writer.add_clip_projector_type(gguf.VisionProjectorType.LFM2A)
9710+
self.gguf_writer.add_audio_num_mel_bins(self.hparams_audio["feat_in"])
9711+
self.gguf_writer.add_audio_attention_layernorm_eps(1e-5)
9712+
9713+
def tensor_force_quant(self, name, new_name, bid, n_dims):
9714+
if ".conv" in name and ".weight" in name:
9715+
return gguf.GGMLQuantizationType.F32
9716+
return super().tensor_force_quant(name, new_name, bid, n_dims)
9717+
9718+
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
9719+
# skip language model tensors
9720+
if name.startswith("lfm."):
9721+
return []
9722+
9723+
# for training only
9724+
if any(p in name for p in ["audio_loss_weight"]):
9725+
return []
9726+
9727+
# for audio output
9728+
if any(p in name for p in ["codebook_offsets", "depth_embeddings", "depth_linear", "depthformer"]):
9729+
return []
9730+
9731+
# fold running_mean, running_var and eps into weight and bias for batch_norm
9732+
if "batch_norm" in name:
9733+
if self._batch_norm_tensors is None:
9734+
self._batch_norm_tensors = [{} for _ in range(self.block_count)]
9735+
assert bid is not None
9736+
self._batch_norm_tensors[bid][name] = data_torch
9737+
9738+
if len(self._batch_norm_tensors[bid]) < 5:
9739+
return []
9740+
9741+
weight = self._batch_norm_tensors[bid][f"conformer.layers.{bid}.conv.batch_norm.weight"]
9742+
bias = self._batch_norm_tensors[bid][f"conformer.layers.{bid}.conv.batch_norm.bias"]
9743+
running_mean = self._batch_norm_tensors[bid][f"conformer.layers.{bid}.conv.batch_norm.running_mean"]
9744+
running_var = self._batch_norm_tensors[bid][f"conformer.layers.{bid}.conv.batch_norm.running_var"]
9745+
eps = 1e-5 # default value
9746+
9747+
a = weight / torch.sqrt(running_var + eps)
9748+
b = bias - running_mean * a
9749+
return [
9750+
(self.map_tensor_name(f"conformer.layers.{bid}.conv.batch_norm.weight"), a),
9751+
(self.map_tensor_name(f"conformer.layers.{bid}.conv.batch_norm.bias"), b),
9752+
]
9753+
9754+
return [(self.map_tensor_name(name), data_torch)]
9755+
9756+
96849757
@ModelBase.register("SmallThinkerForCausalLM")
96859758
class SmallThinkerModel(TextModel):
96869759
model_arch = gguf.MODEL_ARCH.SMALLTHINKER

‎ggml/src/ggml-cuda/ssm-conv.cu‎

Lines changed: 14 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -102,31 +102,25 @@ static void ssm_conv_f32_cuda(const float * src0, const float * src1, const int
102102
const int threads = 128;
103103
GGML_ASSERT(nr % threads == 0);
104104

105-
if (n_t <= 32) {
106-
const dim3 blocks(n_s, (nr + threads - 1) / threads, 1);
107-
if (nc == 4) {
108-
ssm_conv_f32<threads, 4><<<blocks, threads, 0, stream>>>(src0, src1, src0_nb0, src0_nb1, src0_nb2, src1_nb1,
109-
dst, dst_nb0, dst_nb1, dst_nb2, n_t);
110-
} else if (nc == 3) {
111-
ssm_conv_f32<threads, 3><<<blocks, threads, 0, stream>>>(src0, src1, src0_nb0, src0_nb1, src0_nb2, src1_nb1,
112-
dst, dst_nb0, dst_nb1, dst_nb2, n_t);
105+
auto launch_kernel = [&](auto NC) {
106+
constexpr int kNC = decltype(NC)::value;
107+
if (n_t <= 32) {
108+
const dim3 blocks(n_s, (nr + threads - 1) / threads, 1);
109+
ssm_conv_f32<threads, kNC><<<blocks, threads, 0, stream>>>(src0, src1, src0_nb0, src0_nb1, src0_nb2, src1_nb1,
110+
dst, dst_nb0, dst_nb1, dst_nb2, n_t);
113111
} else {
114-
GGML_ABORT("Only support kernel size = 3 or size = 4 right now.");
115-
}
116-
} else {
117-
if (nc == 4) {
118-
const int64_t split_n_t = 32;
119-
dim3 blocks(n_s, (nr + threads - 1) / threads, (n_t + split_n_t - 1) / split_n_t);
120-
ssm_conv_long_token_f32<threads, 4, split_n_t><<<blocks, threads, 0, stream>>>(
121-
src0, src1, src0_nb0, src0_nb1, src0_nb2, src1_nb1, dst, dst_nb0, dst_nb1, dst_nb2, n_t);
122-
} else if (nc == 3) {
123112
const int64_t split_n_t = 32;
124113
dim3 blocks(n_s, (nr + threads - 1) / threads, (n_t + split_n_t - 1) / split_n_t);
125-
ssm_conv_long_token_f32<threads, 3, split_n_t><<<blocks, threads, 0, stream>>>(
114+
ssm_conv_long_token_f32<threads, kNC, split_n_t><<<blocks, threads, 0, stream>>>(
126115
src0, src1, src0_nb0, src0_nb1, src0_nb2, src1_nb1, dst, dst_nb0, dst_nb1, dst_nb2, n_t);
127-
} else {
128-
GGML_ABORT("Only support kernel size = 3 or size = 4 right now.");
129116
}
117+
};
118+
119+
switch (nc) {
120+
case 3: launch_kernel(std::integral_constant<int, 3>{}); break;
121+
case 4: launch_kernel(std::integral_constant<int, 4>{}); break;
122+
case 9: launch_kernel(std::integral_constant<int, 9>{}); break;
123+
default: GGML_ABORT("Only support kernel sizes 3, 4, 9 right now.");
130124
}
131125
}
132126

‎gguf-py/gguf/constants.py‎

Lines changed: 43 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -687,6 +687,8 @@ class MODEL_TENSOR(IntEnum):
687687
V_TOK_EOI = auto() # cogvlm
688688
# audio (mtmd)
689689
A_ENC_EMBD_POS = auto()
690+
A_ENC_EMBD_NORM = auto()
691+
A_ENC_EMBD_TO_LOGITS = auto()
690692
A_ENC_CONV1D = auto()
691693
A_PRE_NORM = auto()
692694
A_POST_NORM = auto()
@@ -697,8 +699,13 @@ class MODEL_TENSOR(IntEnum):
697699
A_ENC_OUTPUT = auto()
698700
A_ENC_OUTPUT_NORM = auto()
699701
A_ENC_FFN_UP = auto()
702+
A_ENC_FFN_NORM = auto()
700703
A_ENC_FFN_GATE = auto()
701704
A_ENC_FFN_DOWN = auto()
705+
A_ENC_FFN_UP_1 = auto()
706+
A_ENC_FFN_NORM_1 = auto()
707+
A_ENC_FFN_GATE_1 = auto()
708+
A_ENC_FFN_DOWN_1 = auto()
702709
A_MMPROJ = auto()
703710
A_MMPROJ_FC = auto()
704711
A_MM_NORM_PRE = auto()
@@ -710,6 +717,12 @@ class MODEL_TENSOR(IntEnum):
710717
NEXTN_HNORM = auto()
711718
NEXTN_SHARED_HEAD_HEAD = auto()
712719
NEXTN_SHARED_HEAD_NORM = auto()
720+
# lfm2 audio
721+
A_ENC_NORM_CONV = auto()
722+
A_ENC_LINEAR_POS = auto()
723+
A_ENC_POS_BIAS_U = auto()
724+
A_ENC_POS_BIAS_V = auto()
725+
A_ENC_OUT = auto()
713726

714727

715728
MODEL_ARCH_NAMES: dict[MODEL_ARCH, str] = {
@@ -1059,6 +1072,8 @@ class MODEL_TENSOR(IntEnum):
10591072
MODEL_TENSOR.V_TOK_EOI: "v.eoi",
10601073
# audio (mtmd)
10611074
MODEL_TENSOR.A_ENC_EMBD_POS: "a.position_embd",
1075+
MODEL_TENSOR.A_ENC_EMBD_NORM: "a.position_embd_norm",
1076+
MODEL_TENSOR.A_ENC_EMBD_TO_LOGITS: "a.embd_to_logits",
10621077
MODEL_TENSOR.A_ENC_CONV1D: "a.conv1d.{bid}",
10631078
MODEL_TENSOR.A_PRE_NORM: "a.pre_ln",
10641079
MODEL_TENSOR.A_POST_NORM: "a.post_ln",
@@ -1068,9 +1083,14 @@ class MODEL_TENSOR(IntEnum):
10681083
MODEL_TENSOR.A_ENC_INPUT_NORM: "a.blk.{bid}.ln1",
10691084
MODEL_TENSOR.A_ENC_OUTPUT: "a.blk.{bid}.attn_out",
10701085
MODEL_TENSOR.A_ENC_OUTPUT_NORM: "a.blk.{bid}.ln2",
1086+
MODEL_TENSOR.A_ENC_FFN_NORM: "a.blk.{bid}.ffn_norm",
10711087
MODEL_TENSOR.A_ENC_FFN_UP: "a.blk.{bid}.ffn_up",
10721088
MODEL_TENSOR.A_ENC_FFN_GATE: "a.blk.{bid}.ffn_gate",
10731089
MODEL_TENSOR.A_ENC_FFN_DOWN: "a.blk.{bid}.ffn_down",
1090+
MODEL_TENSOR.A_ENC_FFN_NORM_1: "a.blk.{bid}.ffn_norm_1",
1091+
MODEL_TENSOR.A_ENC_FFN_UP_1: "a.blk.{bid}.ffn_up_1",
1092+
MODEL_TENSOR.A_ENC_FFN_GATE_1: "a.blk.{bid}.ffn_gate_1",
1093+
MODEL_TENSOR.A_ENC_FFN_DOWN_1: "a.blk.{bid}.ffn_down_1",
10741094
MODEL_TENSOR.A_MMPROJ: "mm.a.mlp.{bid}",
10751095
MODEL_TENSOR.A_MMPROJ_FC: "mm.a.fc",
10761096
MODEL_TENSOR.A_MM_NORM_PRE: "mm.a.norm_pre",
@@ -1082,6 +1102,12 @@ class MODEL_TENSOR(IntEnum):
10821102
MODEL_TENSOR.NEXTN_HNORM: "blk.{bid}.nextn.hnorm",
10831103
MODEL_TENSOR.NEXTN_SHARED_HEAD_HEAD: "blk.{bid}.nextn.shared_head_head",
10841104
MODEL_TENSOR.NEXTN_SHARED_HEAD_NORM: "blk.{bid}.nextn.shared_head_norm",
1105+
# lfm2
1106+
MODEL_TENSOR.A_ENC_NORM_CONV: "a.blk.{bid}.norm_conv",
1107+
MODEL_TENSOR.A_ENC_LINEAR_POS: "a.blk.{bid}.linear_pos",
1108+
MODEL_TENSOR.A_ENC_POS_BIAS_U: "a.blk.{bid}.pos_bias_u",
1109+
MODEL_TENSOR.A_ENC_POS_BIAS_V: "a.blk.{bid}.pos_bias_v",
1110+
MODEL_TENSOR.A_ENC_OUT: "a.pre_encode.out",
10851111
}
10861112

10871113
MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
@@ -1137,6 +1163,8 @@ class MODEL_TENSOR(IntEnum):
11371163
MODEL_TENSOR.V_TOK_EOI,
11381164
# audio
11391165
MODEL_TENSOR.A_ENC_EMBD_POS,
1166+
MODEL_TENSOR.A_ENC_EMBD_NORM,
1167+
MODEL_TENSOR.A_ENC_EMBD_TO_LOGITS,
11401168
MODEL_TENSOR.A_ENC_CONV1D,
11411169
MODEL_TENSOR.A_PRE_NORM,
11421170
MODEL_TENSOR.A_POST_NORM,
@@ -1146,13 +1174,27 @@ class MODEL_TENSOR(IntEnum):
11461174
MODEL_TENSOR.A_ENC_INPUT_NORM,
11471175
MODEL_TENSOR.A_ENC_OUTPUT,
11481176
MODEL_TENSOR.A_ENC_OUTPUT_NORM,
1177+
MODEL_TENSOR.A_ENC_FFN_NORM,
11491178
MODEL_TENSOR.A_ENC_FFN_UP,
11501179
MODEL_TENSOR.A_ENC_FFN_GATE,
11511180
MODEL_TENSOR.A_ENC_FFN_DOWN,
1181+
MODEL_TENSOR.A_ENC_FFN_NORM_1,
1182+
MODEL_TENSOR.A_ENC_FFN_UP_1,
1183+
MODEL_TENSOR.A_ENC_FFN_GATE_1,
1184+
MODEL_TENSOR.A_ENC_FFN_DOWN_1,
11521185
MODEL_TENSOR.A_MMPROJ,
11531186
MODEL_TENSOR.A_MMPROJ_FC,
11541187
MODEL_TENSOR.A_MM_NORM_PRE,
11551188
MODEL_TENSOR.A_MM_NORM_MID,
1189+
MODEL_TENSOR.CONVNEXT_DW,
1190+
MODEL_TENSOR.CONVNEXT_NORM,
1191+
MODEL_TENSOR.CONVNEXT_PW1,
1192+
MODEL_TENSOR.CONVNEXT_PW2,
1193+
MODEL_TENSOR.A_ENC_NORM_CONV,
1194+
MODEL_TENSOR.A_ENC_LINEAR_POS,
1195+
MODEL_TENSOR.A_ENC_POS_BIAS_U,
1196+
MODEL_TENSOR.A_ENC_POS_BIAS_V,
1197+
MODEL_TENSOR.A_ENC_OUT,
11561198
],
11571199
MODEL_ARCH.LLAMA: [
11581200
MODEL_TENSOR.TOKEN_EMBD,
@@ -3328,6 +3370,7 @@ class VisionProjectorType:
33283370
LIGHTONOCR = "lightonocr"
33293371
COGVLM = "cogvlm"
33303372
JANUS_PRO = "janus_pro"
3373+
LFM2A = "lfm2a" # audio
33313374

33323375

33333376
# Items here are (block size, type size)

0 commit comments

Comments
 (0)