大模型推理的"第一性原理" 从Deepseek V3到Kimi K2 无论模型如何变化，当前主流大模型的核心架构都是基于transformer。其本质是一个由多层相同结构堆叠而成的序列到序列的模型，每一层中有一个最核心的模块：Self-Attention（自注意力），让模型理解“哪些词和那些词相关”。 核心：Self-Attention直观理解 Self-Attention本质上就是回答一个问题：让计算机能够理解这些词的含义，对于当前这个词，序列里其他哪些词对它最重要？重要多少？（在一个句子中，根据上下文来进行理解每个词的含义） 直观理解，“苹果公司发布了一款新的手机，它的股价可能也会随之上涨”。 提问（Query）：“它”，指的是什么？ 关键词（Key）: 在上下文中扫描关键信息，每个词都自带一个“身份标签”（Key），比如“苹果”可能是水果或公司，“公司”是实体，“手机”是产品。 匹配相关性，将你的问题（Query）和每个词的“身份标签”（Key）比对：“它”指代物 -> 和“公司”“手机”更相关，和“发布”（动作）不太相关。“股价”通常属于公司 ->所以“公司”的标签最匹配。匹配度越高，你分配的关注度（Attention Score）就越大。 提取信息（Value）：确定“它”指“苹果公司”后，你从“公司”这个词背后提取的信息（Value）是：这是一个企业实体，有股价，属于科技行业……，最终，你综合所有词的信息（但主要权重给了“公司”），得出理解：“它” = 苹果公司，其股价上涨了。 总结一下：计算机怎么理解每个词的意思？词跟句子中的每个词去做相关性计算（矩阵相乘）得到一个注意力分数，然后通过这个分数再去跟句子的中每个词做提取信息，最后得到的一个归一信息就是计算机对该词的理解。 $$Attention(Q, K, V) = softmax\left(\frac{QK^T}{\sqrt{d_k}}\right) V$$ 其中： Q (Query)：当前位置的"提问"向量，代表"我在找什么" K (Key)：所有位置的"标签"向量，代表"我是什么" V (Value)：所有位置的"内容"向量，代表"我能提供什么" d_k：Key 向量的维度，用于缩放防止点积过大 步骤1: 线性投影 Q = X × Wq # [seq_len, d_model] × [d_model, d_k] → [seq_len, d_k] K = X × Wk # 同上 V = X × Wv # 同上 步骤2: 计算注意力分数 S = Q × K^T # [seq_len, d_k] × [d_k, seq_len] → [seq_len, seq_len] # ⚠ 这个 N×N 矩阵是内存瓶颈的根源 步骤3: 缩放 + Softmax P = softmax(S / √d_k) # 每行归一化为概率分布 步骤4: 加权求和 O = P × V # [seq_len, seq_len] × [seq_len, d_v] → [seq_len, d_v] transformer架构 Multi-Head Attention 单组Q/K/V只能捕获一种“关注模式”。但语言中的关系是多维度的，同一个词可能同时需要关注语法结构，语义关联，位置关系等不同维度的信息。而Multi-Head Attention的做法就是：把一个大的注意力拆成多个小的"头"（head），每个头独立学习一种关注模式，最多把所有的结果拼接起来。还是以“苹果公司发布了一款新的手机，它的股价可能也会随之上涨”为例，当模型处理"它"时，不同的头可能各自关注不同的维度： Head 0 (指代关系): "它" → 重点关注 "公司"(0.45) "苹果"(0.30) → 解析代词指代 Head 1 (主谓关系): "它" → 重点关注 "股价"(0.40) "上涨"(0.35) → 理解主语-谓语 Head 2 (修饰关系): "它" → 重点关注 "新的"(0.30) "手机"(0.25) → 捕捉形容词-名词 Head 3 (位置邻近): "它" → 重点关注 "，"(0.35) "的"(0.30) → 关注局部上下文 每个头只看 d_dmole / h 维的子空间（如 4096/32 = 128 维），计算量和单头 Attention 一样，但能并行捕捉多种关系。 MultiHead(Q, K, V) = Concat(head₁, head₂, ..., headₕ) × Wo 其中 headᵢ = Attention(X × Wqᵢ, X × Wkᵢ, X × Wvᵢ) Attention变体演进 随着模型规模的增长和序列变长，标准的MHA面临两大瓶颈：KV Cache（下文会阐述）显存占用和N * N注意力矩阵的计算/内存开销。为此，业界提出了多种优化方案。 GQA 由 Google 在 2023 年提出（论文：GQA: Training Generalized Multi-Query Attention），核心思想是让多个 Query 头共享同一组 Key/Value 头，从而减少 KV Cache 的大小。GQA 是 MHA 和 MQA 的折中——既显著减少了 KV Cache（LLaMA 3 使用 8 组 GQA，缓存压缩 4 倍），又保持了接近 MHA 的模型质量。LLaMA 2 70B、LLaMA 3、Mistral 等主流模型均采用 GQA。 MLA 由 DeepSeek 在 DeepSeek-V2（2024）中提出，是目前 KV Cache 压缩最激进的方案。核心思想：不再缓存完整的 K 和 V，而是将它们压缩为一个低维的潜变量（latent variable） ckvckv，推理时再通过上投影矩阵恢复。KV Cache 只需存储 ckvckv（维度 d_c），而非完整的 K 和 V（维度 h×d_k + h×d_v）。DeepSeek-V2 中 d_c=512，而 h×d_k=12288，压缩比高达 24 倍。 Flash Attention（Tri Dao, 2022）解决的不是"用更少的头"，而是"如何高效计算注意力本身"。标准 Attention 需要在 GPU 显存（HBM）中存储完整的 N×N 注意力矩阵，当序列长度 N 很大时，显存和带宽都成为瓶颈。而flash attention解决的是"单次Attention计算中的N x N矩阵太大"的问题，分块计算，不存完整的注意力矩阵，这样可以省显存带宽不改变计算量。类比：一次性搬运10吨大石头过桥（承重桥不够，还慢），分20次搬运20公斤，反而更快。 Flash Attention 的核心技巧是分块计算（tiling）， 不改变模型结构，是纯粹的工程优化，可以和 MHA/GQA/MLA 任意组合使用。llama.cpp 中通过 llama_set_flash_attn()启用。 将 Q/K/V 分成小块，每次只在 GPU 的片上缓存（SRAM）中计算一个小块的注意力 使用在线 softmax 算法（Milakov & Gimelshein, 2018），无需看到完整的一行就能增量更新 softmax 永远不在 HBM 中存储完整的 N×N 矩阵 应用示例: 翻译英文 Tokenization（分词）：将原始文本（我有一个苹果）切分成模型可识别的词元（如[“我”, “有”, “一”, “个”, “苹”, “果”]），并映射为Token ID（如[2769, 2150, 671, ...]）。这是文本进入模型的第一步。 Word Embedding（词嵌入）：将每个Token ID转换为一个稠密的、可学习的向量，这个向量能捕捉词语的语义信息。 Positional Encoding（位置编码）：为词嵌入向量添加位置信息，让模型知道单词的顺序（因为Self-Attention本身是顺序无关的）。 Encoder Stack（编码器堆叠）：由N个相同的Encoder层堆叠而成。每层通过自注意力（Self-Attention） 和前馈网络（FFN），让输入序列的每个词元都能充分“看到”并融合其他所有词元的信息，最终输出包含完整上下文语义的编码表示。 Decoder Stack（解码器堆栈）：由N个相同的Decoder层堆叠而成。除了自注意力和FFN，它还包含交叉注意力，用于“聚焦”Encoder的输出，从而基于源语言信息生成目标语言。推理时采用自回归方式，逐个生成Token。 Linear Projection（线性投影）：将Decoder顶层的输出向量（维度d_model）映射到词表大小的向量。这一步通常在图中Decoder之后，用“Linear + Softmax”表示。 Softmax & Sampling（概率化与采样）：对上一步的输出做Softmax，得到每个词元在词表中的概率分布。然后根据策略（贪心、束搜索、采样等）选择下一个Token。 Detokenization（反分词）：将生成的Token ID序列转换回人类可读的文本。这是模型输出的最后一步。 大模型推理的软件实现 llama.cpp的软件框架图 simple.cpp为例：推理流程 模型加载：不只是读取，更是映射 模型加载可以分为4个阶段：GGUF文件解析、架构识别与超参数、张量创建与设备分配、权重数据加载。 阶段一：GGUF文件解析 GGUF（GGML Universal Format）是 llama.cpp 的模型文件格式。gguf_init_from_file()按顺序解析其二进制结构。上图展示了 GGUF 文件的二进制布局，从上到下依次是： Header：Magic Number（"GGUF"）、版本号、张量数量、KV 元数据数量。 KV Metadata：模型超参数（架构类型、embedding 维度、层数、注意力头数等）和完整的词表信息（分词算法、token 列表、合并规则、特殊 token ID 等），每个 KV 对由 key + value type + value data 组成 Tensor Info：每个权重张量的"目录"——包含张量名称（如 blk.0.attn_q.weight）、维度数组（如 [4096, 4096]）、数据类型（Q4_0/Q8_0/F16/F32 等）以及该张量在数据区中的字节偏移量。这里只记录元信息，不含实际权重数据，类似于书的目录页。 Alignment Padding：填充到 16 字节边界，确保数据区内存对齐。 Tensor Data：所有权重的实际二进制数据，按偏移量排列。 GGUF 文件只存储数据，不存储计算逻辑。矩阵乘法、softmax、残差连接等计算图操作并不在文件中，而是由 llama.cpp 代码根据模型架构（general.architecture）在推理时动态构建（详见下文）。换句话说，GGUF 回答的是"参数是什么"，代码回答的是"怎么算"。 阶段二：架构识别与超参数加载 阶段一解析出的KV元数据是"原始数据"，阶段二的任务是将这些原始数据转换为代码可以直接使用的结构化参数（填充到结构体中），主要分为以下3步。 // src/llama-model.cpp — 调用链 model.load_arch(ml); // 从 "general.architecture" 识别模型类型 model.load_hparams(ml); // 提取 n_embd, n_layer, n_head, n_vocab 等 model.load_vocab(ml); // 加载词表和特殊 Token (BOS/EOS/PAD) load_arch()：读取 general.architecture字段（如 "llama"），映射到内部枚举 LLM_ARCH_LLAMA。这个枚举决定了后续推理时使用哪套计算图构建逻辑——不同架构（LLaMA、GPT-2、Falcon 等）的注意力机制、归一化方式、FFN 结构各不相同，代码需要知道"这是什么模型"才能正确计算。 load_hparams()：根据架构类型，从 KV 元数据中提取对应的超参数。 参数 KV Key 值 含义 n_embd llama.embedding_length 4096 隐藏层维度 n_layer llama.block_count 32 Transformer 层数 n_head llama.attention.head_count 32 注意力头数 n_head_kv llama.attention.head_count_kv 32 KV 头数（GQA 时可能更少） n_vocab llama.vocab_size 32000 词表大小 n_ff llama.feed_forward_length 11008 FFN 中间层维度 n_ctx_train llama.context_length 4096 训练时的上下文长度 这些参数直接决定了每个权重张量的形状。例如 attn_q.weight 的形状是 [n_embd, n_embd] = [4096, 4096]，ffn_gate.weight 的形状是 [n_embd, n_ff] = [4096, 11008]。此外还会解析 RoPE 相关参数（rope_type、rope_freq_base）和归一化参数（norm_eps）等。 load_vocab()：从 KV 元数据中加载完整的词表信息，包括 token 字符串列表、分数/优先级、BPE 合并规则、特殊 token ID（BOS/EOS/PAD/UNK）等。这些信息在后续的分词阶段（下文）会被直接使用。 阶段三：张量创建与设备分配 load_tensors()（src/llama-model.cpp:2576）是整个加载过程中最精巧的阶段。它的核心任务是：根据阶段二确定的模型结构，为每个权重创建 ggml_tensor 描述符，并决定每个张量应该放在哪个计算设备上。 load_tensors()（src/llama-model.cpp:2576）是整个加载过程中最精巧的阶段。它的核心任务是：根据阶段二确定的模型结构，为每个权重创建 ggml_tensor 描述符，并决定每个张量应该放在哪个计算设备上。 张量描述符创建：对于 LLaMA-7B，需要创建约 290+ 个张量描述符（1 个 token embedding + 32 层 × 9 个权重/层 + 输出层）。每个描述符只是一个轻量级的元数据结构（约 400 字节），记录张量的名称、形状、数据类型和设备归属，不包含实际的权重数据。这里使用了 no_alloc = true 策略——先规划，后分配，避免在设备分配确定之前浪费内存。 设备分配策略：n_gpu_layers 参数控制了 CPU/GPU 的分工： 层类型 设备分配策略 输入层（token_embd） 始终在 CPU（embedding 查表操作 CPU 更高效） 中间层 0 ~ N-1 根据 n_gpu_layers 从第 0 层开始依次分配到 GPU 输出层（output） 跟随最后一个中间层的设备 以 LLaMA-7B（32 层）为例： n_gpu_layers=0：纯 CPU 推理，所有层在 CPU n_gpu_layers=20：layer 0~19 在 GPU，layer 20~31 在 CPU（混合推理） n_gpu_layers=33：全部 32 层 + 输出层都在 GPU（全 GPU 推理） Buffer 分配：设备分配确定后，为每个设备创建对应的 ggml_backend_buffer。CPU 使用 ggml_backend_cpu_buffer，GPU 使用对应的 CUDA/Metal buffer。Buffer 是一块连续的设备内存，同一设备上的所有张量共享同一个 buffer，每个张量在 buffer 中占据一段连续区域。这种设计减少了内存碎片，也方便后续的批量数据传输。 阶段四：权重数据的加载 ml.load_all_data()（src/llama-model-loader.cpp:1359）是最后一步——将 GGUF 文件中的实际权重数据填充到阶段三创建的张量描述符中。这个阶段遍历每个张量，根据其目标设备和系统能力，选择三条加载路径之一。 路径 A — mmap 绑定（零拷贝，默认）：直接将张量的 data 指针指向 mmap_base + offset，不发生任何数据拷贝。操作系统通过虚拟内存机制，在首次访问时才从磁盘读取对应的页面（4KB/页）。这是最快的"加载"方式——对于 3.5GB 的模型，mmap 映射本身只需不到 0.1 秒。 路径 B — fread 文件读取（备选）：通过 fread() 将权重数据直接读入 CPU buffer。用于不支持 mmap 的平台（如某些嵌入式系统），或用户显式禁用 mmap 的场景。这种方式启动较慢（需要实际读取全部数据），但首次推理不会有 page fault 开销。 路径 C — GPU 异步上传：对于分配到 GPU 的张量，先通过 fread 读入 CPU 端的 pinned memory（页锁定内存），然后通过 GPU backend 的异步传输接口上传到显存。llama.cpp 使用 event 机制实现流水线化——在上传第 N 个张量的同时，可以并行读取第 N+1 个张量的数据，最大化 PCIe 带宽利用率。 此外，加载过程中还可能发生量化转换：如果目标设备要求的数据类型与文件中存储的不同（例如某些 backend 不支持 Q4_0 但支持 Q8_0），会在加载时自动进行 dequant/requant 转换。mmap 是默认且推荐的方式——一个 7B Q4_0 模型约 3.5GB，mmap 让进程几乎瞬间"加载"完毕，实际的磁盘读取延迟到首次访问时由操作系统透明处理。 语言转数字：Tokenizer 的实现逻辑 llama.cpp 支持四种主流分词算法，具体使用哪种取决于模型的词表类型，无论哪种算法，流程都是三步：先提取特殊 Token（BOS/EOS 等），再对普通文本执行分词算法，最后在序列首尾添加特殊 Token。 算法 代表模型 核心思想 SPM (SentencePiece) LLaMA 1/2 将文本拆为字符，贪心合并频率最高的相邻对 BPE (Byte-Pair Encoding) LLaMA 3, GPT 正则预分词后，按 rank 排序合并 bigram WPM (WordPiece) BERT 按空格分词，对每个词做贪心最长子串匹配 UGM (Unigram) T5 Viterbi 动态规划，选择全局概率最优的分词方案 上下文初始化：llama_model与llama_conext 到目前为止，我们完成了两件事：加载模型权重（3.1 节）和将用户输入转化为 token 序列（3.2 节）。但要真正开始推理，还缺少一个关键角色——推理上下文。 模型权重是"知识"，token 序列是"问题"，而推理上下文就是"工作台"：它提供了推理过程中所需的全部运行时资源，包括 KV Cache（存储注意力历史）、计算图缓冲区（存放中间计算结果）、后端调度器（管理 CPU/GPU 协作）等。没有这个工作台，模型虽然"知道"所有知识，却没有地方展开计算。 llama.cpp 将这个工作台抽象为两个核心结构体：llama_model（模型本身）和 llama_context（推理上下文）。这种分离设计是整个推理框架的架构基石——理解它们的关系，是理解后续所有推理步骤的前提。 Model contex的关系 llama_model 和 llama_context 的关系类似于"类"和"实例"——model 是只读的共享知识库，context 是每次推理会话的独立工作空间： 概念 llama_model llama_context 生命周期 加载一次，多次复用 每次会话创建 包含内容 权重张量、架构参数、词表 KV Cache、计算图、后端调度器 内存占用 大（模型权重 ~3.5GB） 中（KV Cache + 计算缓冲区 ~1-2GB） 线程安全 只读，可多线程共享 非线程安全，每线程一个 这种分离允许一个模型服务多个并发请求。想象一个 API 服务器：模型权重只加载一次（3.5GB）常驻内存，每个用户请求创建独立的 context（~1.5GB），10 个并发请求只需 3.5GB + 10×1.5GB = 18.5GB，而不是 10×5GB = 50GB。同时，model 只读可安全共享，context 包含可变状态各自独立，多线程实现更简单。 Context 创建过程 llama_context_params ctx_params = llama_context_default_params(); ctx_params.n_ctx = n_prompt + n_predict - 1; // 上下文窗口大小 ctx_params.n_batch = n_prompt; // 单次 decode 最大 token 数 ctx_params.no_perf = false; // 启用性能计数 llama_context * ctx = llama_init_from_model(model, ctx_params); llama_init_from_model() 内部会依次创建三个核心组件： KV cache：缓存历史Key和Value KV Cache 是 Transformer 推理加速的关键——没有它，每生成一个新 token 都要重新计算所有历史 token 的 K/V，复杂度从 O(n) 变成 O(n²)。它的大小由 n_ctx 参数决定： KV Cache 大小 = n_ctx × n_layer × 2 × n_embd_gqa × sizeof(type) 各参数含义：n_ctx 是上下文窗口大小（用户指定），n_layer 是模型层数，2 代表 K 和 V 各一份，n_embd_gqa 是每层 KV 的维度（GQA 模型会压缩），sizeof(type) 是数据类型大小（F16 为 2 字节）。 模型 n_ctx n_layer n_embd_gqa type KV Cache 大小 LLaMA-7B 2048 32 4096 F16 2048×32×2×4096×2 = 1.0 GB LLaMA-7B 4096 32 4096 F16 4096×32×2×4096×2 = 2.0 GB LLaMA-70B 2048 80 1024 F16 2048×80×2×1024×2 = 0.64 GB 注意 LLaMA-70B 使用了 GQA（8 组），n_embd_gqa = 8192/8 = 1024，KV Cache 反而比 7B 小。n_ctx 的选择需要权衡：太小无法处理长文本，太大浪费内存且注意力计算复杂度是 O(n_ctx²)。 后端调度器与计算缓冲 后端调度器（ggml_backend_sched）是多设备协作的核心，负责三件事：根据算子类型和输入张量的设备决定在哪个后端执行（CPU/CUDA/Metal）；当算子输入在不同设备时自动插入数据传输操作；通过 event 机制确保跨设备操作的正确时序。例如在 CPU+GPU 混合推理中（n_gpu_layers=20），前 20 层在 GPU 执行，后 12 层在 CPU，调度器在边界自动插入 GPU→CPU 数据搬运。 计算缓冲区存储推理过程中每层的中间激活值，大小取决于 n_batch 和模型的中间层维度。llama.cpp 支持计算图复用——如果连续两次推理的 batch 形状相同（如 Decode 阶段每次都是 1 个 token），可以跳过重新构建计算图，直接复用上次的图结构。 n_batch和n_ubatch：批次拆分 代码中设置的 n_batch 是用户指定的批次大小（Prefill 阶段等于 prompt 长度，Decode 阶段为 1），但 llama.cpp 内部会进一步拆分为微批次（ubatch，默认 512）。例如 n_batch=1024 会被拆成 2 个 ubatch 分别处理。 拆分的核心动机是控制 Self-Attention 的内存峰值。计算 Q×K^T 时，Q 矩阵大小取决于当前处理的 token 数，而 K 矩阵始终是完整的（包含所有历史 token）。拆分 ubatch 只缩小 Q，不影响 K： 场景 Q 矩阵大小 K 矩阵大小 注意力矩阵大小 内存占用 (FP32) 不拆分 (n_batch=1024) [1024, 128] [1536, 128] [1024, 1536] 6 MB 拆分 ubatch[0] [512, 128] [1536, 128] [512, 1536] 3 MB 拆分 ubatch[1] [512, 128] [1536, 128] [512, 1536] 3 MB 拆分后峰值内存降低 50%。此外，相同大小的 ubatch 可以复用计算图结构（Decode 阶段每次 1 个 token，图只需构建一次），也能避免显存不足时的 OOM。默认值 512 是内存与性能的经验折中，可通ctx_params.n_ubatch 调整。 循环生成：Token的"轮回" // examples/simple/simple.cpp:149-203 llama_batch batch = llama_batch_get_one(prompt_tokens.data(), prompt_tokens.size()); for (int n_pos = 0; n_pos + batch.n_tokens < n_prompt + n_predict; ) { // 核心：将 batch 送入模型计算 llama_decode(ctx, batch); n_pos += batch.n_tokens; // 从 logits 中采样下一个 token new_token_id = llama_sampler_sample(smpl, ctx, -1); // 检查是否生成了结束符 if (llama_vocab_is_eog(vocab, new_token_id)) break; // 输出文本 char buf[128]; int n = llama_token_to_piece(vocab, new_token_id, buf, sizeof(buf), 0, true); printf("%s", std::string(buf, n).c_str()); // 用新 token 构建下一轮的 batch batch = llama_batch_get_one(&new_token_id, 1); } prefill阶段：第一次迭代，Prefill 就是把整个 prompt（比如 "今天天气"）一次性送进 Decoder-only 模型，每个 token 都走完完整的decode （Attention + FFN × 32 层）。它的产出是两个东西：所有 prompt token 的 K、V 向量，写入 KV Cache（这是最重要的产出，后续 Decode 阶段要复用）最后一个 token 位置的 logits，用于采样第一个生成字。 decode阶段：每次只送入 1 个新 token（上一步刚生成的），同样走完 32 层 Transformer。区别是 Self-Attention计算时，只需要为这 1 个 token 算新的 Q，而 K 和 V 直接从 KV Cache中读取历史的，不用重算。算完后把这个新 token 的 K、V 追加到 KV Cache，采样出下一个token，循环往复直到遇到 EOS。 推理引擎基础：GGML 前面第3章从用户视角看了推理流程：加载模型 ->分词 -> 初始化 context -> 循环生成。但这些高层 API 背后，GGML 引擎是如何组织数据和计算的？ 想象一家餐厅要做一道"Transformer炒饭"：需要工作台（内存池）、食材（张量）、菜谱（计算图）、预处理食材包（量化）。理解这些数据结构，才能真正看懂 llama.cpp 的推理引擎。 内存池：静态内存规划 推理过程中需要创建大量张量描述符（一个 7B 模型约 290+ 个权重张量，加上每次推理的中间结果张量）。如果每个张量都单独 malloc/free，不仅有分配开销，还会产生内存碎片。ggml_context 的解决方案是：预分配一块连续内存，所有张量从中顺序分配。 struct ggml_context { size_t mem_size; // 内存池总大小 void * mem_buffer; // 内存池起始地址 bool mem_buffer_owned; // 是否由 context 自己分配 bool no_alloc; // true = 只分配描述符，不分配数据 int n_objects; // 对象数量 struct ggml_object * objects_begin; // 对象链表头 struct ggml_object * objects_end; // 对象链表尾 }; ggml_context 采用线性分配器（bump allocator）模式： 内存池在创建时一次性分配，之后只做指针递增，不做单独释放 no_alloc = true 时只分配 tensor 描述符（约 400 字节/个），数据由后端 buffer 管理 这种设计避免了频繁的 malloc/free，对推理场景的确定性内存需求非常友好 张量：计算的原子 ggml_tensor 是 GGML 中最核心的数据结构，既描述数据的形状和类型，也记录了计算图中的依赖关系。每个张量同时扮演两个角色：数据容器（通过 type/data 存储实际数值）和计算图节点（通过 op/src 描述运算和输入依赖）。 struct ggml_tensor { enum ggml_type type; // 数据类型：F32, F16, Q4_0, Q8_0... struct ggml_backend_buffer * buffer; // 所属的后端缓冲区 int64_t ne[4]; // 各维度元素数: ne[0]=列, ne[1]=行, ne[2]=batch... size_t nb[4]; // 各维度步长(字节): nb[0]=sizeof(type), nb[1]=ne[0]*nb[0]... enum ggml_op op; // 算子类型: MUL_MAT, ADD, SOFT_MAX... int32_t op_params[16]; // 算子参数 struct ggml_tensor * src[10]; // 输入张量（计算图的边） void * data; // 数据指针（指向 buffer 中的实际数据） char name[64]; // 张量名称，如 "blk.0.attn_q.weight" }; ne/nb 分离：ne 描述逻辑形状，nb 描述物理布局。这使得转置、切片等操作只需修改 nb 而无需拷贝数据（零拷贝视图） src 数组：最多 10 个输入，构成计算图的有向边。例如 MUL_MAT 操作的 src[0] 是权重，src[1] 是输入 op + data 双重身份：一个 tensor 既是数据容器（通过 data 指针），也是计算图中的节点（通过 op 和 src） 计算图：指令的蓝图 struct ggml_tensor * ggml_mul_mat( struct ggml_context * ctx, struct ggml_tensor * a, struct ggml_tensor * b) { GGML_ASSERT(ggml_can_mul_mat(a, b)); GGML_ASSERT(!ggml_is_transposed(a)); const int64_t ne[4] = { a->ne[1], b->ne[1], b->ne[2], b->ne[3] }; struct ggml_tensor * result = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, ne); result->op = GGML_OP_MUL_MAT; result->src[0] = a; result->src[1] = b; return result; } 上图展示了 Transformer 单层的计算图结构。计算图是如何构建和执行的了？ 计算图的构建，nodes 数组天然满足拓扑序：每个节点的依赖一定排在它前面。 遇到 op ≠ 无 的张量 → 加入 nodes 数组（计算节点） 遇到 op = 无 的张量 → 加入 leafs 数组（叶子节点） 使用 visited_hash_set 去重，防止同一个张量被重复添加 递归顺序：先访问 src[0]，再访问 src[1]（从左到右） 计算图的运行： 第1步：执行 nodes[0] = mul_mat_0 ├─ 看 op：MUL_MAT（矩阵乘法） ├─ 看 src：src[0]=weight_0（已有数据），src[1]=input（已有数据） └─ 执行：weight_0 × input → 结果写入 mul_mat_0->data 第2步：执行 nodes[1] = add_0 ├─ 看 op：ADD（加法） ├─ 看 src：src[0]=mul_mat_0（第1步刚算好），src[1]=weight_1（已有数据） └─ 执行：mul_mat_0 + weight_1 → 结果写入 add_0->data ... 后续节点同理，直到 add_1（最终输出）完成 量化格式：用更少的位数表示权重 量化就像把高清菜谱照片压缩成缩略图。原图（FP32）每个像素都精确记录，文件很大；缩略图（Q4_0）只保留大致的色块信息，文件小了 7 倍，远看效果差不多，但放大就能看到模糊和色差。量化也是一样——用更少的位数近似表示权重，省了内存、算得更快，但每个权重值都不再精确，会引入误差。位数越少（Q8->Q4 ->Q2），压缩越狠，误差越大，模型输出质量也会逐渐下降。 量化是 llama.cpp 能在消费级硬件上运行大模型的关键技术。核心思想是将 FP32/FP16 权重压缩为更低位宽的整数表示。 以最常用的 Q4_0 为例，每 32 个权重值组成一个 block： 计算这 32 个值的最大绝对值，得到 scale = max_abs / 7 每个值量化为 4-bit 整数：q = round(val / scale) + 8 反量化：val = (q - 8) × scale 一个 Q4_0 block 只占 18 字节（2B scale + 16B data），而 FP32 需要 128 字节，压缩率约 7 倍。 llama.cpp 支持多种量化格式，精度和压缩率各有取舍： 格式 位宽 每32权重占用 压缩率(vs FP32) 精度损失 适用场景 F32 32-bit 128 B 1× 无 调试/基准 F16 16-bit 64 B 2× 极小 GPU 推理 Q8_0 8-bit 34 B 3.8× 很小 精度敏感场景 Q5_K_M 5-bit ~22 B 5.8× 小 精度与大小的平衡点 Q4_0 4-bit 18 B 7.1× 中等 最常用，性价比高 Q4_K_M 4-bit ~20 B 6.4× 较小 Q4 中精度最好 Q2_K 2-bit ~12 B 10.7× 大 极端内存受限 其中 K 系列（Q4_K_M、Q5_K_M 等）是改进版量化，对不同层使用不同的量化精度——注意力层用更高精度，FFN 层用更低精度，在相同位宽下获得更好的输出质量。量化对实际推理的影响： 内存：LLaMA-7B 从 FP16 的 13GB 降到 Q4_0 的 3.5GB，可以在 8GB 显存的消费级 GPU 上运行 速度：量化权重更小，内存带宽占用更低，Decode 阶段（内存带宽瓶颈）速度反而更快 质量：Q8_0 几乎无损，Q4_K_M 在大多数任务上表现接近 FP16，Q2_K 在复杂推理任务上会有明显退化 llama.cpp关键技术补充 计算图的构建：llm_build_graph如何动态生成算子依赖 整个流程可以分为3步：架构识别、构建函数分发，逐层构建计算图。 架构识别：模型加载时，load_arch()（src/llama-model.cpp）从 GGUF 的 KV 元数据中读取 general.architecture 字段（如 "llama"），并映射为枚举值 LLM_ARCH_LLAMA。llama.cpp 在 src/llama-arch.h 中定义了 123+ 种架构枚举，涵盖主流的开源模型。 架构家族 代表枚举 说明 LLaMA 系列 LLM_ARCH_LLAMA, LLM_ARCH_LLAMA4 Meta LLaMA 1/2/3/4 Qwen 系列 LLM_ARCH_QWEN2, LLM_ARCH_QWEN3, LLM_ARCH_QWEN35MOE 阿里通义千问 Gemma 系列 LLM_ARCH_GEMMA, LLM_ARCH_GEMMA2, LLM_ARCH_GEMMA3 Google Gemma 状态空间模型 LLM_ARCH_MAMBA, LLM_ARCH_MAMBA2, LLM_ARCH_RWKV6 非 Transformer 架构 BERT 系列 LLM_ARCH_BERT, LLM_ARCH_MODERN_BERT 编码器模型 其他 LLM_ARCH_FALCON, LLM_ARCH_BAICHUAN, LLM_ARCH_DEEPSEEK_V2 60+ 种其他架构 构建函数分发：推理时，build_graph()（src/llama-model.cpp:8134）根据架构枚举，通过一个巨大的 switch 语句（500+ 行）分发到对应的构建函数 ggml_cgraph * llama_model::build_graph(const llm_graph_params & params) const { std::unique_ptr<llm_graph_context> llm; switch (arch) { case LLM_ARCH_LLAMA: llm = std::make_unique<llm_build_llama<false>>(*this, params); break; case LLM_ARCH_QWEN2: llm = std::make_unique<llm_build_qwen2>(*this, params); break; case LLM_ARCH_GEMMA: llm = std::make_unique<llm_build_gemma>(*this, params); break; // ... 120+ 个 case default: GGML_ABORT("fatal error"); } llm->build_pooling(...); // 后处理：池化、采样、输出层 llm->build_sampling(); llm->build_dense_out(...); return llm->res->get_gf(); // 返回完整计算图 } 每个架构对应一个独立的构建类（如 llm_build_llama），实现在 src/models/ 目录下（共 113 个文件）。这些类都继承自 llm_graph_context 基类，在构造函数中完成计算图的构建。 逐层构建计算图：以 LLaMA 架构为例，llm_build_llama 的构造函数按照 Transformer 的标准流程，逐层调用 GGML 的算子创建函数。 // src/models/llama.cpp — llm_build_llama 构造函数（简化） struct llm_build_llama : public llm_graph_context { llm_build_llama(const llama_model & model, const llm_graph_params & params) { // 1. 输入 embedding ggml_tensor * inpL = build_inp_embd(model.tok_embd); ggml_tensor * inp_pos = build_inp_pos(); // 2. 逐层构建 for (int il = 0; il < n_layer; ++il) { ggml_tensor * cur; // Attention 子图 cur = build_norm(inpL, model.layers[il].attn_norm, LLM_NORM_RMS, il); ggml_tensor * Qcur = build_lora_mm(model.layers[il].wq, cur); // Q = Wq × x ggml_tensor * Kcur = build_lora_mm(model.layers[il].wk, cur); // K = Wk × x ggml_tensor * Vcur = build_lora_mm(model.layers[il].wv, cur); // V = Wv × x Qcur = ggml_rope_ext(ctx0, Qcur, inp_pos, ...); // RoPE 位置编码 Kcur = ggml_rope_ext(ctx0, Kcur, inp_pos, ...); cur = build_attn(inp_attn, Qcur, Kcur, Vcur, ...); // Self-Attention cur = build_lora_mm(model.layers[il].wo, cur); // 输出投影 // 残差连接 + FFN cur = ggml_add(ctx0, cur, inpL); // 残差① ggml_tensor * ffn_inp = cur; cur = build_norm(cur, model.layers[il].ffn_norm, LLM_NORM_RMS, il); cur = build_ffn(cur, ...); // SwiGLU FFN inpL = ggml_add(ctx0, cur, ffn_inp); // 残差② } // 3. 最终输出 cur = build_norm(inpL, model.output_norm, LLM_NORM_RMS, -1); cur = build_lora_mm(model.output, cur); // LM Head: [n_vocab, n_embd] × [n_embd, 1] ggml_build_forward_expand(gf, cur); // 展开为完整计算图 } }; 每个 build_* / ggml_* 调用并不立即执行计算，而是创建一个新的 ggml_tensor 节点并记录其 op 和 src 依赖。最终 ggml_build_forward_expand() 从输出节点反向遍历，将所有依赖的节点按拓扑序收集到 ggml_cgraph 中。 问题1：为什么不把计算图存入GGUF文件中？ 牺牲通用性，换取极致的简单和快速。可以做到代码随时优化，无需重新转换模型，方便更新计算图。 问题2：权重张量如何与计算图关联？ 答案是张量命名规范。GGUF文件通过张量名称提供权重数据，代码逻辑通过构建函数定义计算图结构。两者在load_tensors中通过名称匹配"焊接"在一起。 GGUF 文件中: "blk.0.attn_q.weight" → offset=0x1A00, dims=[4096,4096], type=Q4_0 ↓ load_tensors() 按名称匹配 代码中: model.layers[0].wq → tensor.data = mmap_base + 0x1A00 ↓ build_graph() 构建计算图 计算图: ggml_mul_mat(wq, cur) src[0] = wq (权重, 来自GGUF) src[1] = cur (激活值, 来自上一步计算) ↓ graph_compute() 执行 结果: Qcur = Wq × cur KV cache的管理：逻辑槽位分配、物理内存复用 KV Cache 是推理性能的关键。它缓存了已计算的 Key 和 Value 向量，避免 Decode 阶段的重复计算。llama.cpp 的 KV Cache 采用 Ring Buffer（环形缓冲区） 设计。 槽位分配流程（find_slot()，src/llama-kv-cache.cpp:704）： 从 head 位置开始扫描空闲槽位 如果 head 已经远超已用区域，重置到 0（环形复用） 对每个 token 分配一个槽位，记录其 pos（位置）和 seq_id（序列ID） 分配失败时回滚所有修改（事务性设计） // src/llama-kv-cache.h — 核心数据结构（简化） struct llama_kv_cache { // 每层的 K、V 张量 struct kv_layer { ggml_tensor * k; // [n_embd_k_gqa, kv_size] ggml_tensor * v; // [n_embd_v_gqa, kv_size] }; std::vector<kv_layer> layers; // 槽位管理 std::vector<llama_kv_cells> v_cells; // 每个槽位的状态（pos, seq_id） std::vector<int32_t> v_heads; // 环形缓冲区头指针 }; 算子是如何跑到硬件上的 调用llama_decode调用链如下，具体代码的调用链路： llama_decode(ctx, batch) // [src/llama-context.cpp] 应用入口 └→ llama_context::decode(batch) // [src/llama-context.cpp:1511] 拆分 batch 为多个 ubatch └→ process_ubatch(ubatch) // [src/llama-context.cpp:1159] 处理每个微批次 │ ├→ model.build_graph(params) // [src/llama-model.cpp:8134] Step 1: 构建计算图 │ └→ switch(arch) // 根据模型架构分发（500+ 行 switch） │ ├→ llm_build_llama(model, params) // [src/models/llama.cpp] 逐层调用 ggml_* │ ├→ llm_build_qwen2(model, params) // [src/models/qwen2.cpp] │ └→ ...（113种架构，113个文件） │ ├→ ggml_backend_sched_alloc_graph(sched,gf) // [ggml/src/ggml-backend.cpp] 分配张量内存 │ └→ graph_compute(gf) // [src/llama-context.cpp:2145] Step 2: 执行计算图 └→ ggml_backend_sched_graph_compute_async(sched, gf) // [ggml/src/ggml-backend.cpp:1797] ├→ 绑定 backend_id // 遍历节点，根据 supports_op() 分配后端 ├→ 分割子图 (splits) // 相同后端的连续节点归为一组 └→ ggml_backend_sched_compute_splits(sched) // [ggml/src/ggml-backend.cpp:1445] └→ for each split: // --->逐个子图执行（确定后端） └→ backend->iface.graph_compute(split) // [ggml/src/ggml-backend.cpp:364] │ ├→ CUDA 后端: ggml_cuda_graph_compute() // [ggml/src/ggml-cuda/ggml-cuda.cu] │ └→ for each node: │ ggml_cuda_compute_forward(node) // switch(op) → ggml_cuda_mul_mat() 等 │ ├→ Metal / SYCL / CANN / Vulkan / ... // 各自的 graph_compute() │ └→ CPU 后端（下面展开） CPU 后端的执行路径（以 SpacemiT RISC-V 为例）： backend->iface.graph_compute(split) └→ ggml_backend_cpu_graph_compute() // [ggml/src/ggml-cpu/ggml-cpu.cpp:170] ├→ ggml_graph_plan(cgraph, n_threads) // 计算工作缓冲区大小、线程分配 └→ ggml_graph_compute(cgraph, &cplan) // [ggml/src/ggml-cpu/ggml-cpu.c:3223] └→ ggml_graph_compute_thread() // 多线程，每个线程遍历 nodes └→ for each node: ggml_compute_forward(params, node) // [ggml/src/ggml-cpu/ggml-cpu.c:1694] │ │ //关键分叉点：先尝试硬件加速插件 ├→ ggml_cpu_extra_compute_forward(params, node) // [ggml/src/ggml-cpu/traits.cpp:12] │ └→ 遍历注册的 extra_buffer_types: │ ├→ SpacemiT IME (GGML_USE_CPU_RISCV64_SPACEMIT) │ │ └→ tensor_traits::compute_forward() │ │ └→ forward_mul_mat_q4() // [spacemit/ime.cpp] │ │ └→ IME 内核 (vmadot) // [spacemit/ime1_kernels.cpp] │ ├→ Intel AMX (__AMX_INT8__) │ ├→ ARM KleidiAI (GGML_USE_CPU_KLEIDIAI) │ └→ CPU Repack (GGML_USE_CPU_REPACK) │ │ // 如果没有插件处理，走标准 CPU 路径 └→ switch (node->op) ├→ MUL_MAT: ggml_compute_forward_mul_mat() // 标准 GEMM │ └→ llamafile_sgemm() 或 vec_dot // SIMD 加速 ├→ RMS_NORM: ggml_compute_forward_rms_norm() ├→ ROPE: ggml_compute_forward_rope() ├→ SOFT_MAX: ggml_compute_forward_soft_max() ├→ ADD: ggml_compute_forward_add() └→ ...（100+ 种算子） 核心思路：构建计算图 → 调度器分配后端 → 遍历节点逐个执行算子。CPU 后端有一个插件机制（extra_buffer_type），允许各硬件平台（SpacemiT、AMX、KleidiAI）注册自己的加速实现，优先于标准 CPU 路径执行。 LLM 推理的核心算子与加速策略 ggml 定义了 100+ 种算子（enum ggml_op，见 ggml/include/ggml.h），但跑一个 Transformer 大语言模型，真正用到的只有十几种。按计算量排序如下。MUL_MAT 一个算子就占了 90% 以上的计算量。这就是为什么各硬件平台的加速工作几乎都集中在矩阵乘法上。 算子 作用 计算量占比 说明 MUL_MAT 矩阵乘法 ~90% Q/K/V 投影、FFN、LM Head，几乎所有"权重 x 输入"都是它 FLASH_ATTN_EXT 融合注意力 ~5% 将 QK^T、softmax、V 融合为一个算子，减少显存访问 ROPE 旋转位置编码 <1% 给 Q、K 注入位置信息 RMS_NORM RMS 归一化 <1% 每层 Attention 和 FFN 前的归一化 SOFT_MAX Softmax <1% 注意力分数归一化（未融合时使用） ADD 逐元素加法 <1% 残差连接 MUL 逐元素乘法 <1% SwiGLU 中的门控乘法 SILU / GELU 激活函数 <1% FFN 中的非线性激活 GET_ROWS 行查找 <1% Embedding 查表：token_id → 向量 CPY / CONT / VIEW / PERMUTE 数据搬运/变形 ~0% 不做计算，只改变数据的排列方式 llama.cpp 通过后端接口对接了 10+ 种硬件平台，每个平台重点加速的算子不同： 平台 后端文件 重点加速算子 加速手段 x86 CPU ggml-cpu/arch/x86/ MUL_MAT AVX2/AVX-512 SIMD，量化内核 ARM CPU ggml-cpu/arch/arm/ MUL_MAT NEON/SVE 向量指令 RISC-V CPU (SpacemiT) ggml-cpu/spacemit/ MUL_MAT, NORM, RMS_NORM IME 矩阵扩展（Q4_0/Q4_1/Q4_K），向量扩展 (RVV) NVIDIA GPU ggml-cuda/ (61个.cu文件) MUL_MAT, FLASH_ATTN, ROPE, SOFT_MAX 等 CUDA Kernel，Tensor Core，cuBLAS Apple GPU ggml-metal/ MUL_MAT, FLASH_ATTN 等 Metal Shader Intel GPU ggml-sycl/ MUL_MAT 等 SYCL/oneAPI 华为 NPU ggml-cann/ MUL_MAT 等 CANN 算子库 通用 GPU ggml-vulkan/ MUL_MAT 等 Vulkan Compute Shader CPU 后端通常只加速 MUL_MAT（因为其他算子计算量太小，优化收益不大），但 SpacemiT 是个例外——除了 MUL_MAT，还通过 IME 加速了 NORM 和 RMS_NORM。GPU 后端则会加速更多算子，原因不是这些小算子在GPU 上跑得更快，而是为了避免来回搬运：如果 GPU 只加速 MUL_MAT，一个 Transformer 层的执行路径会变成 GPU 做矩阵乘->搬到 CPU 做 RMS_NORM ->搬回 GPU 做矩阵乘 ->搬到 CPU 做 ROPE ->搬回 GPU...每次搬运都要走 PCIe 总线，延迟高、带宽有限。把小算子也留在 GPU 上，整层计算都在显存里完成，省掉的搬运开销远大于小算子本身的计算开销。这也是调度器 Pass 2 让 GPU "吞噬"相邻节点的原因——尽量减少 split 数量和跨设备传输。 CUDA 后端是最完整的，61 个 .cu 文件覆盖了几乎所有算子。其他后端通常只实现高频算子，不支持的算子会自动 fallback 到 CPU 执行。 了解了整体架构和各平台的加速策略后，接下来深入调度层的工作机制。计算图构建完毕后，调度器通过四个阶段将其送上硬件： 后端抽象层：屏蔽硬件差异的接口 GGML 通过一套 C 接口（虚函数表）将硬件差异封装在后端实现中，上层代码完全不感知底层是 CPU、CUDA 还是 Metal。 // ggml/src/ggml-backend-impl.h:87 — 后端接口 struct ggml_backend_i { const char * (*get_name)(ggml_backend_t backend); void (*free)(ggml_backend_t backend); // 异步数据传输 void (*set_tensor_async)(ggml_backend_t, ggml_tensor *, const void *, size_t, size_t); void (*get_tensor_async)(ggml_backend_t, const ggml_tensor *, void *, size_t, size_t); bool (*cpy_tensor_async)(ggml_backend_t src, ggml_backend_t dst, const ggml_tensor *, ggml_tensor *); void (*synchronize)(ggml_backend_t); // 核心：计算图执行 enum ggml_status (*graph_compute)(ggml_backend_t, struct ggml_cgraph *); // 事件同步（流水线并行） void (*event_record)(ggml_backend_t backend, ggml_backend_event_t event); void (*event_wait)(ggml_backend_t backend, ggml_backend_event_t event); }; // 后端实例 struct ggml_backend { ggml_guid_t guid; struct ggml_backend_i iface; // ★ 函数指针表（C语言的"虚函数"） ggml_backend_dev_t device; // 所属设备 void * context; // 后端私有数据（如 CUDA stream） }; 这是经典的 C 语言多态模式：通过函数指针表实现接口抽象。上层代码调用 backend->iface.graph_compute(backend, cgraph) 时，实际执行的是 CPU/CUDA/Metal 各自的实现。不同后端填入各自的函数指针： // CPU后端 — ggml/src/ggml-cpu/ggml-cpu.cpp:193 static const struct ggml_backend_i ggml_backend_cpu_i = { /* .graph_compute = */ ggml_backend_cpu_graph_compute, // ★ 核心 /* .set_tensor_async = */ NULL, // CPU不需要异步传输 /* .synchronize = */ NULL, // CPU不需要同步 /* .event_record = */ NULL, // CPU不支持事件 }; // CUDA后端（对比） static const struct ggml_backend_i ggml_backend_cuda_i = { /* .graph_compute = */ ggml_cuda_graph_compute, /* .event_record = */ ggml_backend_cuda_event_record, // GPU支持事件 /* .event_wait = */ ggml_backend_cuda_event_wait, // ... 支持异步传输等 }; CPU 后端的大部分可选接口都是 NULL——因为 CPU 上所有操作天然是同步的，不需要异步传输和事件机制。每个后端还需要实现设备接口，声明自己支持哪些操作。supports_op 和 offload_op 是调度决策的基础——调度器通过询问每个后端"是否能跑这个算子？ struct ggml_backend_device_i { const char * (*get_name)(ggml_backend_dev_t dev); void (*get_memory)(ggml_backend_dev_t dev, size_t * free, size_t * total); // ★ 调度器用这两个方法决定算子分配 bool (*supports_op)(ggml_backend_dev_t, const struct ggml_tensor * op); // 是否支持该算子 bool (*offload_op)(ggml_backend_dev_t, const struct ggml_tensor * op); // 是否建议卸载到此设备 ggml_backend_buffer_type_t (*get_buffer_type)(ggml_backend_dev_t); // 获取缓冲区类型 }; 子图分割与后端调度 ggml_backend_sched 是多后端协作的核心调度器。它的职责是：将计算图中的每个算子分配到最合适的后端，按后端边界切分子图，并处理跨后端的数据搬运。 一个计算图里可能混合了不同后端的算子——前20层权重在 GPU，后12层在 CPU。但每个后端的 graph_compute() 只能执行自己设备上的节点，不能跨设备访问。所以在执行之前，必须把一张大图按后端边界切分成多个子图（split），每个 split 内的节点都在同一个后端上，跨 split 的数据通过拷贝节点搬运。这就是 ggml_backend_sched（调度器）的核心工作：先决定每个节点跑在哪个后端（调度），再按后端边界切分子图（分割），最后为跨后端的张量插入拷贝节点。整个过程通过5轮 Pass 完成。 顶层：顶部是计算图中的两类张量：权重张量（Wq、Wk 等）在 load_tensors() 阶段已有确定的 buffer，n_gpu_layers决定了它们的设备归属；中间张量（Qcur、FFN_out 等）在 build_graph() 时创建，buffer 为 NULL，后端待定。 中间：中间是 split_graph() 的5轮 Pass：Pass 1 根据已有 buffer 确定后端，中间张量跟随权重；Pass 2 让 GPU 向相邻未分配节点扩展，CPU 兜底；Pass 3 尝试升级到更高优先级后端；Pass 4 确保无遗漏；Pass 5 按后端边界切分 split 并插入跨后端的 cpy_tensor。 底层：底部是输出的 splits 数组——以 n_gpu_layers=20 的32层模型为例，切出3个 split，相邻 split 之间通过 cpy 完成跨设备数据搬运。 以 n_gpu_layers=20、模型32层为例： Pass 1-4 后的 nodes 后端分配: [embd(CPU), layer0_norm(GPU), layer0_attn(GPU), ..., layer19_ffn(GPU), layer20_norm(CPU), ..., lm_head(CPU)] Pass 5 遍历 nodes: i=0: embd → backend=CPU, cur=CPU → split[0] 开始 i=1: layer0_norm → backend=GPU, cur=CPU → ★ 后端变了! 切分! split[0].i_end=1, split[1] 开始, inputs=[embd] i=2~N: layer0~19 → 全是GPU → 不切 i=N+1: layer20_norm → backend=CPU → ★ 后端变了! 切分! split[1].i_end=N+1, split[2] 开始, inputs=[layer19_output] ...直到结束 结果: 3个splits split[0]: CPU [embd] inputs=[] split[1]: GPU [layer0~19] inputs=[embd] split[2]: CPU [layer20~31 + lm_head] inputs=[layer19_output] 对于跨后端的输入张量，Pass 5 会创建拷贝张量并替换原始图中的 src 指针： Pass 4 后的后端分配: nodes: [A, B, C, D, E, F, G] 后端ID: [GPU,GPU,GPU,CPU,CPU,GPU,GPU] Pass 5 切分结果: split[0]: backend=GPU, nodes=[A, B, C] inputs: [] (无跨后端输入) split[1]: backend=CPU, nodes=[D, E] inputs: [C_output] ← 需要从GPU拷贝到CPU 操作: D.src[0] 被替换为 C_output 的CPU拷贝 split[2]: backend=GPU, nodes=[F, G] inputs: [E_output] ← 需要从CPU拷贝到GPU 操作: F.src[0] 被替换为 E_output 的GPU拷贝 切分的目的是让每个 split 内部的所有节点都在同一个后端上执行，避免单个 graph_compute调用中出现跨设备访问。 CPU后端执行模型：节点间串行、节点内并行 当 compute_splits() 调用 backend->iface.graph_compute() 进入 CPU 后端时，执行分为三层： ggml_backend_cpu_graph_compute(backend, cgraph) // 后端入口 ├→ ggml_graph_plan(cgraph, n_threads, threadpool) // 规划：计算work buffer大小 └→ ggml_graph_compute(cgraph, &cplan) // 执行：启动多线程 └→ ggml_graph_compute_thread(worker) // 每个线程 └→ for each node: ggml_compute_forward(params, node) // 算子分发 ggml_barrier(threadpool) // 线程同步屏障 多线程执行 CPU 后端启动 N 个线程，每个线程执行相同的 ggml_graph_compute_thread() 函数： // ggml/src/ggml-cpu/ggml-cpu.c:2948 static thread_ret_t ggml_graph_compute_thread(void * data) { struct ggml_compute_params params = { .ith = state->ith, // 我是第几个线程 .nth = n_threads, // 总共几个线程 .wdata = cplan->work_data, // 共享work buffer指针 }; // ★ 所有线程遍历相同的节点序列 for (int node_n = 0; node_n < cgraph->n_nodes; node_n++) { struct ggml_tensor * node = cgraph->nodes[node_n]; // 每个线程都调用同一个算子，内部根据 ith/nth 分工 ggml_compute_forward(¶ms, node); // ★ 节点间屏障：所有线程必须完成当前节点才能进入下一个 ggml_barrier(state->threadpool); } } 理解了"节点间串行、节点内并行"的执行模型后，就能理解图分割调度的设计逻辑： 节点间串行意味着：一个 split 内的所有节点必须在同一个后端上顺序执行。如果计算图中的节点分属不同后端（如 GPU 和 CPU），就必须把图切成多个 split，每个 split 内部是同一后端的连续节点序列。 原始计算图（拓扑序）: node0[GPU] → node1[GPU] → node2[CPU] → node3[CPU] → node4[GPU] → node5[GPU] 分割后: split[0]: backend=GPU, nodes=[node0, node1] split[1]: backend=CPU, nodes=[node2, node3] ← node1→node2 需要 GPU→CPU 拷贝 split[2]: backend=GPU, nodes=[node4, node5] ← node3→node4 需要 CPU→GPU 拷贝 split 之间也是串行的。compute_splits() 逐个执行 split，每个 split 完成后才执行下一个： // ggml/src/ggml-cpu/ggml-cpu.c (简化) for (int i = 0; i < sched->n_splits; i++) { // 1. 拷贝跨后端输入 for (int j = 0; j < split->n_inputs; j++) { ggml_backend_tensor_copy_async(input_backend, split_backend, input, copy); } // 2. 在该后端上执行整个子图（节点间串行、节点内并行） ggml_backend_graph_compute_async(split_backend, &split->graph); // 3. 记录事件（用于流水线优化） ggml_backend_event_record(split->event, split_backend); } 因此，整体执行模型是三层串行嵌套： [图片] 每次 split 切换都有代价：跨后端数据拷贝（split 边界处的张量需要从一个后端搬到另一个）、同步开销（上一个 split 必须完全结束下一个才能开始）、线程池重启（不同后端可能有不同的线程池配置）。这就是 Pass 2（扩展相邻节点到同一后端）和 Pass 3（升级到更高优先级后端）的核心动机——让尽可能多的连续节点在同一后端执行，减少 split 切换。 节点内并行的3种模式 不同算子采用不同的并行策略： 模式一：行交错（Stride） — SOFT_MAX、RMS_NORM 等使用。每个线程处理第 ith, ith+nth, ith+2*nth, ... 行： 4线程处理12行: 线程0: 行0, 行4, 行8 线程1: 行1, 行5, 行9 线程2: 行2, 行6, 行10 线程3: 行3, 行7, 行11 模式二：行分块（Block） — ADD、MUL 等使用。每个线程处理连续的一段行，有利于内存局部性： 4线程处理12行: 线程0: 行0, 行1, 行2 线程1: 行3, 行4, 行5 线程2: 行6, 行7, 行8 线程3: 行9, 行10, 行11 模式三：原子抢占（Work Stealing） — MUL_MAT 使用。将输出矩阵切成小块，线程通过原子操作抢占下一个可用块：这种模式的优势：自动负载均衡。快的线程多做几块，慢的线程少做几块，避免木桶效应。 // ggml/src/ggml-cpu/ggml-cpu.c:1235 while (true) { int current_chunk = atomic_fetch_add(&tp->current_chunk, 1); if (current_chunk >= nchunk0 * nchunk1) break; // 没有更多块了 // 计算该块的行列范围，执行矩阵乘法 } 算子分发 ggml_compute_forward() 是从后端到具体算子的最后一跳： // ggml/src/ggml-cpu/ggml-cpu.c:1686 static void ggml_compute_forward(struct ggml_compute_params * params, struct ggml_tensor * tensor) { // ★ 优先尝试硬件加速插件 if (ggml_cpu_extra_compute_forward(params, tensor)) { return; // 插件处理了，直接返回 } // 标准CPU路径：大switch分发 switch (tensor->op) { case GGML_OP_MUL_MAT: ggml_compute_forward_mul_mat(params, tensor); break; case GGML_OP_ADD: ggml_compute_forward_add(params, tensor); break; case GGML_OP_ROPE: ggml_compute_forward_rope(params, tensor); break; case GGML_OP_RMS_NORM: ggml_compute_forward_rms_norm(params, tensor); break; case GGML_OP_SOFT_MAX: ggml_compute_forward_soft_max(params, tensor); break; // ... 100+ 种算子 } } 插件注册机制 ggml_backend_cpu_get_extra_buffer_types() 返回的插件列表： ├→ SpacemiT IME (GGML_USE_CPU_RISCV64_SPACEMIT) │ └→ MUL_MAT + Q4_0/Q4_1/Q4_K → IME内核 │ └→ RMS_NORM/NORM → RVV实现 ├→ Intel AMX (__AMX_INT8__) │ └→ MUL_MAT + 量化类型 → AMX tile指令加速 ├→ ARM KleidiAI (GGML_USE_CPU_KLEIDIAI) │ └→ MUL_MAT → NEON/SVE优化内核 └→ CPU Repack (GGML_USE_CPU_REPACK) └→ MUL_MAT → 重排权重布局以优化缓存访问 上层代码完全不感知硬件差异。无论是 x86 的 AMX、ARM 的 KleidiAI 还是 RISC-V 的 IME，都通过同一个插件接口注入，标准 CPU 路径作为通用 fallback 始终可用。 CPU后端加速算子：以SpacemiT IME（K1）为例 前面介绍了 CPU 后端的通用执行模型和插件机制。本节以 SpacemiT IME 为具体案例，展示一个硬件加速插件如何实现矩阵乘法的全流程优化。 SpacemiT 支持的算子总览 SpacemiT 后端通过 extra_buffer_type 插件机制注册（ggml-cpu/spacemit/ime.cpp），目前支持 3 个算子： 算子 加速方式 支持的权重类型 (src[0]) 输入类型 (src[1]) 额外条件 MUL_MAT IME 硬件矩阵引擎 Q4_0, Q4_1, Q4_K F32 src[0] 为 2D，ne[1] % 16 0 NORM RVV 向量指令 — F32 无 RMS_NORM RVV 向量指令 — F32 无 分发逻辑（ime.cpp 中的 get_tensor_traits()）： // MUL_MAT → 根据量化类型选择对应的 IME tensor_traits static const tensor_traits<block_q4_0, 8, 16> q4_0_16x8_q8_0; // Q4_0 → IME 内核 static const tensor_traits<block_q4_1, 8, 16> q4_1_16x8_q8_0; // Q4_1 → IME 内核 static const tensor_traits<block_q4_K, 8, 16> q4_k_16x8_q8_0; // Q4_K → IME 内核 // NORM / RMS_NORM → 走 RVV 向量实现 static const tensor_traits_common rvv_impl; IME矩阵乘法执行流程 IME（Integer Matrix Extension） 是进迭时空 K1/K1X RISC-V 处理器的硬件加速指令集，专门用于量化整数矩阵乘法。其核心指令是 vmadot（向量乘累加点积）。 整个矩阵乘法分为三个阶段： 阶段一：权重预处理（Repack）:模型加载时，权重从标准 Q4_0 格式重排为 IME 友好的 Q4_0x16 交错格式： // ggml-cpu/spacemit/ime.cpp:230 — Q4_0 重排 // 标准布局: [block0][block1]...[block15] // 交错布局: [scale0..scale15][data0..data15] // 16 个 block 的 scale 连续存放，data 连续存放 // 这样 RISC-V 向量指令可以一次加载 16 个 block 的数据 void repack_q4_0_to_q4_0_16_bl(const block_q4_0 * src, block_q4_0x16 * dst, int64_t n); 阶段二：激活值量化（Runtime）：每次推理时，将 float32 的激活值（A 矩阵）在线量化为 int8： // ggml-cpu/spacemit/ime1_kernels.cpp:79 — 4行同时量化 // 使用 RISC-V 向量指令加速： // vle32.v 加载 float32 数据 // vfabs.v 取绝对值，求 scale // vfdiv.vf 除以 scale // vfcvt.x.f 浮点转整数 // vse8.v 存储 int8 结果 void quantize_a_4row_i8(const float * src, int8_t * dst, float * scales, int64_t K, int64_t lda, int blk_len); 阶段三：IME 内核计算：核心计算由 gemm_kernel_i8i4() 完成（ime1_kernels.cpp:3157），根据 M 的大小选择不同路径： // M >= 4 时使用 4 行内核，M < 4 时使用单行内核 if (CountM >= 4) { SQ4BitGemmM4Kernel_CompInt8(...) // 4行×16列 分块 return 4; } else { SQ4BitGemmM1Kernel_CompInt8(...) // 1行×16列 分块 return 1; } 内核中的关键 RISC-V 汇编指令序列： # ime1_kernels.cpp:1171 — 核心计算宏 vle8.v v4, (s1) # 加载 Q4 量化数据 vand.vi v0, v4, 15 # 提取低 4 位 (值 0-15) vsrl.vi v4, v4, 4 # 提取高 4 位 vmadot v16, v10, v2 # IME 核心指令！int8 × int4 → int32 累加 # 硬件在一个周期内完成多个乘累加 vfcvt.f.x.v v16, v16 # int32 → float32 vfmacc.vv v28, v16, v24 # float 乘累加: result += int_result × scale_A × scale_B vmadot 是 IME 的核心，它在硬件层面实现了 int8 × int4 的点积累加，这就是SpacemiT K1大模型加速的关键。 参考: https://magazine.sebastianraschka.com/p/the-big-llm-architecture-comparison

概述 llama.cpp 的模型加载系统是一个高度优化的、支持多后端、多设备的模型权重加载框架。它通过精心设计的数据结构和加载流程，实现了： 零拷贝加载：通过内存映射（mmap）实现模型文件的零拷贝加载 多设备支持：自动发现和分配GPU、CPU等设备，支持模型分层加载到不同设备 架构无关：通过统一的抽象支持Llama、Qwen、DeepSeek、Mamba、RWKV等多种架构 高性能I/O：支持直接I/O、异步上传、预取等优化策略 本文将从数据结构、加载流程、I/O机制、设备管理等多个维度深入解析模型加载的实现细节。 关键数据结构 llama_model - 核心模型容器 llama_model 是模型加载后的核心容器，它持有模型的所有权重指针、元数据和设备信息。 结构定义 struct llama_model { llm_type type = LLM_TYPE_UNKNOWN; // 模型规模（如3B、7B） llm_arch arch = LLM_ARCH_UNKNOWN; // 模型架构（如Llama、Qwen） std::string name = "n/a"; // 模型名称 llama_hparams hparams = {}; // 超参数 llama_vocab vocab; // 词表 // 全局权重张量 struct ggml_tensor * tok_embd = nullptr; // Token Embeddings struct ggml_tensor * output = nullptr; // Output weights struct ggml_tensor * output_norm = nullptr; // Output normalization // 分类器相关（用于分类模型） struct ggml_tensor * cls = nullptr; // 层权重数组 std::vector<llama_layer> layers; // GGUF元数据 std::unordered_map<std::string, std::string> gguf_kv; // 设备列表 std::vector<ggml_backend_dev_t> devices; // LoRA适配器 std::unordered_set<llama_adapter_lora *> loras; // 性能统计 int64_t t_load_us = 0; int64_t t_start_us = 0; // Pimpl模式隐藏实现细节 private: struct impl; std::unique_ptr<impl> pimpl; }; 核心设计点 结构化抽象 - layers数组 std::vector layers 是核心设计 将重复的层权重封装到 llama_layer 中 支持动态层数，无需修改结构体定义 设备管理 devices 向量记录所有使用的后端设备 dev_layer(int il) 返回第 il 层所在的设备 dev_output() 返回输出层所在的设备 支持多GPU分层加载 Pimpl模式 隐藏底层实现细节（mmap管理、内存映射句柄等） 保持API稳定性，修改实现不影响头文件 权重指针而非数据 llama_model 不存储实际权重数据 只存储指向权重数据的指针（ggml_tensor *） 实际数据通过mmap映射或buffer管理 关键方法 load_arch(): 从loader加载架构信息 load_hparams(): 加载超参数 load_vocab(): 加载词表 load_tensors(): 创建张量并分配设备 build_graph(): 构建计算图 create_memory(): 创建KV Cache内存 llama_layer - 单层权重抽象 llama_layer 封装了Transformer单层的所有权重张量，支持多种架构变体。 结构定义（部分） struct llama_layer { // 归一化层 struct ggml_tensor * attn_norm = nullptr; struct ggml_tensor * attn_norm_b = nullptr; struct ggml_tensor * attn_q_norm = nullptr; // Gemma 2 struct ggml_tensor * attn_k_norm = nullptr; struct ggml_tensor * ssm_norm = nullptr; // Mamba // 注意力机制 struct ggml_tensor * wq = nullptr; // Query投影 struct ggml_tensor * wk = nullptr; // Key投影 struct ggml_tensor * wv = nullptr; // Value投影 struct ggml_tensor * wo = nullptr; // Output投影 struct ggml_tensor * wqkv = nullptr; // 合并QKV（某些架构） // 注意力偏置 struct ggml_tensor * bq = nullptr; struct ggml_tensor * bk = nullptr; struct ggml_tensor * bv = nullptr; struct ggml_tensor * bo = nullptr; // Feed-Forward网络 struct ggml_tensor * ffn_norm = nullptr; struct ggml_tensor * ffn_gate = nullptr; struct ggml_tensor * ffn_up = nullptr; struct ggml_tensor * ffn_down = nullptr; // 混合专家系统（MoE） struct ggml_tensor * ffn_gate_inp = nullptr; // 门控输入 struct ggml_tensor * ffn_gate_exps = nullptr; // 专家门控权重 struct ggml_tensor * ffn_up_exps = nullptr; // 专家上投影 struct ggml_tensor * ffn_down_exps = nullptr; // 专家下投影 // RoPE相关 struct ggml_tensor * rope_freqs = nullptr; struct ggml_tensor * rope_long = nullptr; // LongRoPE struct ggml_tensor * rope_short = nullptr; // ... 更多架构特定张量 }; 设计特点 通用层抽象 合并了对Gemma、Mixtral、Mamba、RWKV、DeepSeek等模型的支持 通过指针是否为nullptr判断架构是否使用该张量 架构变体支持 归一化变体：支持RMSNorm、LayerNorm、Gemma 2的特殊归一化 注意力变体：标准QKV、合并QKV、分组查询注意力（GQA） MoE支持：专家门控、共享专家（Shared Experts） 非Transformer：Mamba的SSM层、RWKV的循环层 加载流程 预分配：model.layers.resize(n_layer) 创建N个空层 点对点赋值：loader解析GGUF，将张量地址赋给对应层的指针 后端绑定：根据设备分配策略，将指针指向CPU或GPU内存 llama_model_loader - 权重加载器 llama_model_loader 是模型加载的"施工队"，负责GGUF文件解析、权重索引、内存映射和数据搬运。 核心成员 struct llama_model_loader { // GGUF元数据 gguf_context_ptr meta; // GGUF解析器上下文 // 文件管理 llama_files files; // 主文件+分片文件句柄 llama_mmaps mappings; // 内存映射句柄 // 权重索引表（核心数据结构） std::map<std::string, llama_tensor_weight, weight_name_comparer> weights_map; // KV覆盖 std::unordered_map<std::string, llama_model_kv_override> kv_overrides; // 张量Buffer类型覆盖 const llama_model_tensor_buft_override * tensor_buft_overrides; // I/O配置 bool use_mmap; bool use_direct_io; bool check_tensors; }; 权重索引表：weights_map 这是loader最核心的数据结构，充当"权重目录"： // Key: tensor名称，如 "tok_embeddings.weight", "blk.0.attn_q.weight" // Value: llama_tensor_weight（定位信息） struct llama_tensor_weight { uint16_t idx; // 文件分片索引（主文件=0，split可能是1、2...） size_t offs; // 在文件中的字节偏移 ggml_tensor * tensor; // 张量元数据（shape/type/name） }; weight_name_comparer：自定义排序规则 解析层索引（如 blk.5 → 5） 按层号排序，而非纯字符串排序 优化磁盘读取顺序，利用预读缓存 关键流程 构造函数：初始化GGUF解析 llama_model_loader::llama_model_loader(...) { // 1. 调用gguf_init_from_file解析文件头 meta.reset(gguf_init_from_file(fname.c_str(), params)); // 2. 识别架构 get_key(LLM_KV_GENERAL_ARCHITECTURE, arch_name, false); llm_kv = LLM_KV(llm_arch_from_string(arch_name)); // 3. 打开文件句柄 files.emplace_back(new llama_file(fname.c_str(), "rb", use_direct_io)); // 4. 构建weights_map索引表 for (ggml_tensor * cur = ggml_get_first_tensor(ctx); cur; ...) { weights_map.emplace(tensor_name, llama_tensor_weight(...)); } // 5. 处理分片文件（如果有） // ... } init_mappings：初始化内存映射 void llama_model_loader::init_mappings(bool prefetch, llama_mlocks * mlock_mmaps) { if (use_mmap) { for (const auto & file : files) { // 创建mmap映射 std::unique_ptr<llama_mmap> mapping = std::make_unique<llama_mmap>(file.get(), prefetch ? -1 : 0, is_numa); mappings.emplace_back(std::move(mapping)); // 可选：内存锁定 if (mlock_mmaps) { std::unique_ptr<llama_mlock> mlock_mmap(new llama_mlock()); mlock_mmap->init(mapping->addr()); mlock_mmaps->emplace_back(std::move(mlock_mmap)); } } } } create_tensor：创建张量对象（空壳） 从weights_map查找元数据 在指定的ggml_context中创建tensor对象 此时tensor->data仍为nullptr load_all_data：加载实际数据 遍历所有tensor 根据use_mmap选择路径： mmap路径：ggml_backend_tensor_alloc(buf_mmap, cur, data) 绑定mmap地址 文件读取路径：file->read_raw(cur->data, n_size) 读取到内存 GPU异步上传：分块读取+异步上传到GPU llama_hparams - 超参数蓝图 llama_hparams 存储模型的数学参数，决定计算图的形状和逻辑。 核心参数分类 基础架构参数 n_embd: 隐藏层维度（如4096） n_layer: 总层数 n_ctx_train: 训练时上下文长度 n_head_arr[]: 每层注意力头数（数组支持每层不同） n_head_kv_arr[]: 每层KV头数（GQA支持） 注意力机制变体 n_embd_head_k_mla: MLA压缩后的KV维度（DeepSeek-V2/V3） n_swa: 滑动窗口大小（Mistral） swa_layers[]: 哪些层启用滑动窗口 位置编码（RoPE & YaRN） rope_freq_base_train: RoPE基础频率 rope_freq_scale_train: RoPE缩放因子 rope_scaling_type_train: 缩放类型（linear/yarn/longrope） yarn_ext_factor: YaRN扩展因子 rope_sections[]: 部分维度旋转（Partial RoPE） 混合专家模型（MoE） n_expert: 总专家数 n_expert_used: 每个token激活的专家数（如8选2） n_ff_exp: 专家FFN维度 n_ff_shexp: 共享专家维度（DeepSeek） expert_gating_func: 门控函数类型 状态空间模型（SSM/Mamba） ssm_d_conv: 卷积维度 ssm_d_inner: 内部维度 ssm_d_state: 状态维度 recurrent_layer_arr[]: 哪些层是循环层（混合架构） 归一化与数值稳定性 f_norm_rms_eps: RMSNorm的epsilon f_attn_logit_softcapping: Gemma 2的注意力分数软上限 f_residual_scale: 残差连接缩放（Granite） 加载流程 void llama_model::load_hparams(llama_model_loader & ml) { const gguf_context * ctx = ml.meta.get(); // 1. 保存所有KV对到gguf_kv（用于元数据查询） for (int i = 0; i < gguf_get_n_kv(ctx); i++) { const char * name = gguf_get_key(ctx, i); const std::string value = gguf_kv_to_str(ctx, i); gguf_kv.emplace(name, value); } // 2. 读取基础参数 ml.get_key(LLM_KV_CONTEXT_LENGTH, hparams.n_ctx_train); ml.get_key(LLM_KV_EMBEDDING_LENGTH, hparams.n_embd); ml.get_key(LLM_KV_BLOCK_COUNT, hparams.n_layer); // 3. 读取数组参数（支持每层不同） ml.get_key_or_arr(LLM_KV_FEED_FORWARD_LENGTH, hparams.n_ff_arr, hparams.n_layer, false); ml.get_key_or_arr(LLM_KV_ATTENTION_HEAD_COUNT, hparams.n_head_arr, hparams.n_layer, false); // 4. 架构特定参数（根据arch分支） switch (arch) { case LLM_ARCH_LLAMA: /* ... */ break; case LLM_ARCH_DEEPSEEK: /* ... */ break; // ... } // 5. 参数校验与推导 // 如果某些参数缺失，根据架构进行推导 } 为什么hparams要独立出来？ 多版本适配：不同架构可能完全没有某些参数（如Mamba没有n_head），但都有n_embd KV Cache管理：hparams决定KV Cache的内存布局和大小 计算图构建：推理时算子需要从hparams读取数学常数 GGML相关结构 gguf_context GGUF文件的内存映射索引，存储所有KV对和张量元信息。 struct gguf_context { uint32_t version; // GGUF协议版本 // KV键值对存储 std::vector<gguf_kv> kv; // 所有元数据KV对 // 张量信息 std::vector<gguf_tensor_info> info; // 所有权重的"档案" // 物理布局 size_t alignment; size_t offset; // 数据区偏移 void * data; // 数据区指针（如果已映射） }; ggml_backend_buffer 张量真实的"物理房产"，解决跨设备内存管理问题。 struct ggml_backend_buffer { struct ggml_backend_buffer_i iface; // 函数指针接口 ggml_backend_buffer_type_t buft; // Buffer类型（CPU/CUDA/Metal等） void * context; // 后端内部状态 size_t size; // 大小 enum ggml_backend_buffer_usage usage; // 用途 }; 关键设计： iface 定义了如何操作这块内存（分配、释放、拷贝等） buft 标识内存由哪个后端分配 支持CPU、CUDA、Metal、Vulkan等多种后端 模型加载完整流程 入口函数：llama_model_load_from_file_impl 这是模型加载的入口点，负责设备发现、模型对象创建和错误处理。 static struct llama_model * llama_model_load_from_file_impl( const std::string & path_model, std::vector<std::string> & splits, struct llama_model_params params) { // 1. 初始化时间统计 ggml_time_init(); // 2. 检查后端是否已加载 if (!params.vocab_only && ggml_backend_reg_count() == 0) { LLAMA_LOG_ERROR("no backends are loaded\n"); return nullptr; } // 3. 设置默认进度回调（如果未提供） if (params.progress_callback == NULL) { params.progress_callback = [](float progress, void * ctx) { // 打印进度点 return true; }; } // 4. 创建模型对象 llama_model * model = new llama_model(params); // 5. 设备发现与选择（详见"设备管理"章节） if (params.devices) { // 使用用户指定的设备 for (ggml_backend_dev_t * dev = params.devices; *dev; ++dev) { model->devices.push_back(*dev); } } else { // 自动发现设备 // - 收集所有GPU设备 // - 收集集成GPU（iGPU） // - 收集RPC服务器 // - 按优先级排序 } // 6. 单GPU模式处理 if (params.split_mode == LLAMA_SPLIT_MODE_NONE) { if (params.main_gpu < 0) { model->devices.clear(); // 只用CPU } else { // 只保留指定的GPU model->devices = {model->devices[params.main_gpu]}; } } // 7. 打印设备信息 for (auto * dev : model->devices) { ggml_backend_dev_props props; ggml_backend_dev_get_props(dev, &props); LLAMA_LOG_INFO("using device %s - %zu MiB free\n", ggml_backend_dev_name(dev), props.memory_free/1024/1024); } // 8. 调用核心加载函数 const int status = llama_model_load(path_model, splits, *model, params); // 9. 错误处理 if (status < 0) { if (status == -1) { LLAMA_LOG_ERROR("failed to load model\n"); } else if (status == -2) { LLAMA_LOG_INFO("cancelled model load\n"); } llama_model_free(model); return nullptr; } return model; } 核心加载函数：llama_model_load 这是模型加载的核心流程，按顺序执行各个加载阶段。 static int llama_model_load( const std::string & fname, std::vector<std::string> & splits, llama_model & model, llama_model_params & params) { // 初始化时间统计 model.t_load_us = 0; time_meas tm(model.t_load_us); model.t_start_us = tm.t_start_us; try { // ========== 阶段0：创建Loader ========== llama_model_loader ml(fname, splits, params.use_mmap, params.use_direct_io, params.check_tensors, params.no_alloc, params.kv_overrides, params.tensor_buft_overrides); ml.print_info(); // 打印模型基本信息 model.hparams.vocab_only = params.vocab_only; model.hparams.no_alloc = params.no_alloc; // ========== 阶段1：架构识别 ========== try { model.load_arch(ml); } catch(const std::exception & e) { throw std::runtime_error("error loading model architecture: " + std::string(e.what())); } // ========== 阶段2：超参数加载 ========== try { model.load_hparams(ml); } catch(const std::exception & e) { throw std::runtime_error("error loading model hyperparameters: " + std::string(e.what())); } // CLIP模型特殊处理 if (model.arch == LLM_ARCH_CLIP) { throw std::runtime_error("CLIP cannot be used as main model"); } // ========== 阶段3：词表加载 ========== try { model.load_vocab(ml); } catch(const std::exception & e) { throw std::runtime_error("error loading model vocabulary: " + std::string(e.what())); } // ========== 阶段4：统计信息 ========== model.load_stats(ml); model.print_info(); // 打印完整模型信息 // ========== 早期退出：仅词表模式 ========== if (params.vocab_only) { LLAMA_LOG_INFO("vocab only - skipping tensors\n"); return 0; } // ========== 阶段5：张量创建与加载 ========== if (!model.load_tensors(ml)) { return -2; // 被进度回调取消 } } catch (const std::exception & err) { LLAMA_LOG_ERROR("error loading model: %s\n", err.what()); return -1; } return 0; } 阶段一：架构识别与元数据解析 load_arch实现 void llama_model::load_arch(llama_model_loader & ml) { arch = ml.get_arch(); // 从loader获取架构枚举 if (arch == LLM_ARCH_UNKNOWN) { throw std::runtime_error("unknown model architecture: '" + ml.get_arch_name() + "'"); } } loader中的架构识别 // llama_model_loader构造函数中 llama_model_loader::llama_model_loader(...) { // 1. 解析GGUF文件头 meta.reset(gguf_init_from_file(fname.c_str(), params)); // 2. 读取架构名称 std::string arch_name; get_key(llm_kv(LLM_KV_GENERAL_ARCHITECTURE), arch_name, false); // 3. 转换为枚举 llm_kv = LLM_KV(llm_arch_from_string(arch_name)); // llm_arch_from_string("llama") -> LLM_ARCH_LLAMA } 架构识别流程： GGUF文件头包含 general.architecture KV对 读取字符串值（如"llama"、"qwen"、"deepseek"） 通过 llm_arch_from_string() 转换为枚举 如果无法识别，抛出异常 阶段二：超参数加载 详见 llama_hparams - 超参数蓝图 章节。 关键点： 从GGUF的KV对中读取所有超参数 支持数组参数（每层不同） 架构特定参数根据arch分支处理 参数校验与默认值推导 阶段三：词表加载 void llama_model::load_vocab(llama_model_loader & ml) { const auto kv = LLM_KV(arch); // 获取架构特定的KV键前缀 vocab.load(ml, kv); // 加载词表 } 词表加载包括： Token ID到字符串的映射 Token类型（普通/BOS/EOS/等） 特殊Token（如<|im_start|>） 子词合并规则（BPE/SPM等） 阶段四：张量创建与设备分配 这是最复杂的阶段，涉及张量创建、设备分配、Buffer类型选择等。 load_tensors流程概览 bool llama_model::load_tensors(llama_model_loader & ml) { const int n_layer = hparams.n_layer; const int n_gpu_layers = this->n_gpu_layers(); // ========== 步骤1：构建Buffer类型列表 ========== pimpl->cpu_buft_list = make_cpu_buft_list(devices, ...); for (auto * dev : devices) { buft_list_t buft_list = make_gpu_buft_list(dev, split_mode, tensor_split); buft_list.insert(buft_list.end(), pimpl->cpu_buft_list.begin(), pimpl->cpu_buft_list.end()); pimpl->gpu_buft_list.emplace(dev, std::move(buft_list)); } // ========== 步骤2：计算设备分割点 ========== std::vector<float> splits(n_devices()); // 根据tensor_split参数或默认策略（按显存）分配 // ========== 步骤3：分配层到设备 ========== auto get_layer_buft_list = [&](int il) -> layer_dev { if (il < i_gpu_start || (il - i_gpu_start) >= act_gpu_layers) { return {cpu_dev, &pimpl->cpu_buft_list}; } const int layer_gpu = /* 根据splits计算 */; return {devices.at(layer_gpu), &pimpl->gpu_buft_list.at(devices.at(layer_gpu))}; }; pimpl->dev_input = {cpu_dev, &pimpl->cpu_buft_list}; pimpl->dev_layer.resize(n_layer); for (int il = 0; il < n_layer; ++il) { pimpl->dev_layer[il] = get_layer_buft_list(il); } pimpl->dev_output = get_layer_buft_list(n_layer); // ========== 步骤4：创建GGML Context ========== // 为每个Buffer类型创建独立的context std::map<ggml_backend_buffer_type_t, ggml_context_ptr> ctx_map; auto ctx_for_buft = [&](ggml_backend_buffer_type_t buft) -> ggml_context * { // 查找或创建context }; // ========== 步骤5：创建张量对象（架构特定） ========== layers.resize(n_layer); auto create_tensor = [&](const LLM_TN_IMPL & tn, const std::initializer_list<int64_t> & ne, int flags) -> ggml_tensor * { // 1. 从weights_map查找元数据 ggml_tensor * t_meta = ml.get_tensor_meta(tn.str().c_str()); // 2. 选择Buffer类型 buft_list_t * buft_list = /* 根据tensor层级选择 */; ggml_backend_buffer_type_t buft = select_weight_buft(hparams, t_meta, op, *buft_list); // 3. 检查覆盖规则 if (ml.tensor_buft_overrides) { // 应用用户指定的覆盖 } // 4. 获取或创建context ggml_context * ctx = ctx_for_buft(buft); // 5. 创建tensor对象（此时data仍为nullptr） return ml.create_tensor(ctx, tn, ne, flags); }; // 架构特定的张量创建 switch (arch) { case LLM_ARCH_LLAMA: tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0); output_norm = create_tensor(tn(LLM_TENSOR_OUTPUT_NORM, "weight"), {n_embd}, 0); output = create_tensor(tn(LLM_TENSOR_OUTPUT, "weight"), {n_embd, n_vocab}, TENSOR_NOT_REQUIRED); for (int i = 0; i < n_layer; ++i) { auto & layer = layers[i]; layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, 0); layer.wq = create_tensor(tn(LLM_TENSOR_ATTN_Q, "weight", i), {n_embd, n_embd_head_k * n_head}, 0); // ... 更多张量 } break; // ... 其他架构 } // ========== 步骤6：初始化内存映射 ========== ml.init_mappings(params.prefetch, &pimpl->mlock_mmaps); // ========== 步骤7：加载实际数据 ========== if (!ml.load_all_data(ctxs_bufs, params.progress_callback, params.progress_callback_user_data)) { return false; // 被取消 } // ========== 步骤8：保存Context和Buffer ========== for (auto & [buft, ctx] : ctx_map) { std::vector<ggml_backend_buffer_ptr> bufs; // 收集该context下的所有buffer pimpl->ctxs_bufs.emplace_back(ctx, std::move(bufs)); } return true; } 设备分配策略 输入层：始终在CPU（dev_input = cpu_dev） 中间层：根据n_gpu_layers和tensor_split分配 输出层：与最后一层相同设备 分割计算： // 默认分割：按显存比例 if (all_zero) { for (size_t i = 0; i < n_devices(); ++i) { ggml_backend_dev_props props; ggml_backend_dev_get_props(devices[i], &props); total_free += props.memory_free; } float acc = 0.0f; for (size_t i = 0; i < n_devices(); ++i) { ggml_backend_dev_props props; ggml_backend_dev_get_props(devices[i], &props); acc += props.memory_free / total_free; splits[i] = acc; } } 阶段五：权重数据加载 load_all_data实现 bool llama_model_loader::load_all_data( std::vector<std::pair<ggml_context_ptr, std::vector<ggml_backend_buffer_ptr>>> & ctxs_bufs, llama_progress_callback progress_callback, void * progress_callback_user_data) { // ========== 准备异步上传（如果支持） ========== ggml_backend_t upload_backend = /* 检查是否支持异步上传 */; std::vector<ggml_backend_buffer_ptr> host_buffers; std::vector<void *> host_ptrs; std::vector<ggml_backend_event_ptr> events; // ========== 遍历所有tensor ========== for (struct ggml_tensor * cur = ggml_get_first_tensor(ctx); cur != NULL; cur = ggml_get_next_tensor(ctx, cur)) { const auto * weight = get_weight(ggml_get_name(cur)); if (weight == nullptr) continue; // 进度回调 if (progress_callback) { if (!progress_callback((float) size_done / size_data, progress_callback_user_data)) { return false; // 用户取消 } } size_t n_size = ggml_nbytes(cur); // ========== 路径A：mmap模式（零拷贝） ========== if (use_mmap) { const auto & mapping = mappings.at(weight->idx); uint8_t * data = (uint8_t *) mapping->addr() + weight->offs; // 验证数据（如果启用） if (check_tensors) { validation_result.emplace_back(std::async(std::launch::async, [cur, data, n_size] { return std::make_pair(cur, ggml_validate_row_data(cur->type, data, n_size)); })); } // 绑定到mmap内存 ggml_backend_buffer_t buf_mmap = bufs.at(weight->idx); if (buf_mmap && cur->data == nullptr) { ggml_backend_tensor_alloc(buf_mmap, cur, data); // tensor->buffer = buf_mmap // tensor->data = data (指向mmap地址) } else { // 拷贝数据（如果tensor已有buffer） ggml_backend_tensor_set(cur, data, 0, n_size); } } // ========== 路径B：文件读取模式 ========== else { const auto & file = files.at(weight->idx); // CPU buffer：直接读取 if (ggml_backend_buffer_is_host(cur->buffer)) { file->seek(weight->offs, SEEK_SET); file->read_raw(cur->data, n_size); if (check_tensors) { validation_result.emplace_back(std::async(std::launch::async, [cur, n_size] { return std::make_pair(cur, ggml_validate_row_data(cur->type, cur->data, n_size)); })); } } // GPU buffer：异步上传 else { if (upload_backend) { // 分块读取+异步上传 size_t offset = weight->offs; size_t aligned_offset = offset & ~(alignment - 1); // ... 对齐读取边界 while (bytes_read < read_end - read_start) { // 等待上一个上传完成 ggml_backend_event_synchronize(events[buffer_idx]); // 读取对齐块 file->read_raw_unsafe(host_ptrs[buffer_idx], read_size); // 异步上传到GPU ggml_backend_tensor_set_async(upload_backend, cur, host_ptrs[buffer_idx], data_read, data_to_copy); ggml_backend_event_record(events[buffer_idx], upload_backend); buffer_idx = (buffer_idx + 1) % n_buffers; // 轮转缓冲区 } } else { // 同步上传 read_buf.resize(n_size); file->seek(weight->offs, SEEK_SET); file->read_raw(read_buf.data(), n_size); ggml_backend_tensor_set(cur, read_buf.data(), 0, n_size); } } } size_done += n_size; } // ========== 验证结果 ========== if (check_tensors) { for (auto & future : validation_result) { auto [tensor, valid] = future.get(); if (!valid) { throw std::runtime_error(format("tensor '%s' has invalid data", ggml_get_name(tensor))); } } } return true; } 三种数据加载路径对比 路径 适用场景 性能 内存占用 mmap绑定 CPU推理，mmap可用 最快（零拷贝） 最小（共享文件页） 文件读取+CPU CPU推理，mmap不可用 中等 中等（需要内存拷贝） 异步上传 GPU推理，支持异步 快（重叠I/O和计算） 中等（需要pinned memory） 同步上传 GPU推理，不支持异步 慢 中等 GGUF协议与I/O抽象 GGUF文件格式解析 GGUF（GPT-Generated Unified Format）是llama.cpp使用的二进制模型格式。 文件结构 ┌─────────────────────────────────────┐ │ Magic Number (4 bytes) │ "GGUF" ├─────────────────────────────────────┤ │ Version (4 bytes) │ 协议版本 ├─────────────────────────────────────┤ │ n_kv (8 bytes) │ KV对数量 ├─────────────────────────────────────┤ │ KV Pairs (变长) │ │ - Key: 字符串长度 + 字符串 │ │ - Type: 类型枚举 │ │ - Value: 根据类型存储 │ ├─────────────────────────────────────┤ │ n_tensors (8 bytes) │ 张量数量 ├─────────────────────────────────────┤ │ Tensor Info Array (变长) │ │ - Name: 字符串长度 + 字符串 │ │ - n_dims: 维度数 │ │ - dims[]: 各维度大小 │ │ - type: 数据类型 │ │ - offset: 数据偏移（相对数据区） │ ├─────────────────────────────────────┤ │ Alignment Padding │ 对齐到16字节 ├─────────────────────────────────────┤ │ Tensor Data (变长) │ 实际权重数据 └─────────────────────────────────────┘ 解析流程 struct gguf_context * gguf_init_from_file_impl(FILE * file, ...) { const struct gguf_reader gr(file); struct gguf_context * ctx = new gguf_context; // 1. 校验Magic Number uint32_t magic = gr.read_u32(); if (magic != GGUF_MAGIC) { throw std::runtime_error("invalid GGUF magic"); } // 2. 读取版本 ctx->version = gr.read_u32(); // 3. 读取KV对数量 int64_t n_kv = gr.read_u64(); // 4. 循环读取所有KV对 for (int64_t i = 0; i < n_kv; ++i) { std::string key = gr.read_string(); gguf_type type = (gguf_type)gr.read_u32(); switch (type) { case GGUF_TYPE_UINT8: ctx->kv.emplace_back(key, gr.read_u8()); break; case GGUF_TYPE_INT32: ctx->kv.emplace_back(key, gr.read_i32()); break; case GGUF_TYPE_FLOAT32: ctx->kv.emplace_back(key, gr.read_f32()); break; case GGUF_TYPE_STRING: ctx->kv.emplace_back(key, gr.read_string()); break; // ... 更多类型 } } // 5. 读取张量数量 int64_t n_tensors = gr.read_u64(); // 6. 循环读取所有张量信息 for (int64_t i = 0; i < n_tensors; ++i) { struct gguf_tensor_info info; info.name = gr.read_string(); info.n_dims = gr.read_u32(); for (uint32_t j = 0; j < info.n_dims; ++j) { info.dims[j] = gr.read_u64(); } info.type = (ggml_type)gr.read_u32(); info.offset = gr.read_u64(); // 相对数据区的偏移 ctx->info.push_back(info); } // 7. 计算数据区偏移 ctx->offset = gr.tell(); // 当前位置就是数据区开始 return ctx; } 内存映射（mmap）机制 mmap是llama.cpp实现零拷贝加载的核心技术。 mmap初始化 llama_mmap::llama_mmap(struct llama_file * file, size_t prefetch, bool numa) { size = file->size(); int fd = file->file_id(); #ifdef _POSIX_MAPPED_FILES int flags = MAP_SHARED; // Linux优化：预读建议 if (posix_fadvise(fd, 0, 0, POSIX_FADV_SEQUENTIAL)) { LLAMA_LOG_WARN("posix_fadvise failed\n"); } // 预填充页面（如果启用） if (prefetch) { flags |= MAP_POPULATE; } // 执行mmap addr = mmap(NULL, file->size(), PROT_READ, flags, fd, 0); if (addr == MAP_FAILED) { throw std::runtime_error(format("mmap failed: %s", strerror(errno))); } // 预取建议（如果启用） if (prefetch > 0) { posix_madvise(addr, std::min(file->size(), prefetch), POSIX_MADV_WILLNEED); } // NUMA优化 if (numa) { posix_madvise(addr, file->size(), POSIX_MADV_RANDOM); } #endif } mmap优势 零拷贝：文件数据直接映射到虚拟地址空间，无需拷贝 延迟加载：页面在首次访问时才从磁盘读取（按需分页） 内存共享：多个进程可以共享同一份映射，节省内存 操作系统优化：OS可以预读、缓存页面 mmap绑定tensor // 在load_all_data中 if (use_mmap) { const auto & mapping = mappings.at(weight->idx); uint8_t * data = (uint8_t *) mapping->addr() + weight->offs; // 关键：将tensor的data指针直接指向mmap地址 ggml_backend_tensor_alloc(buf_mmap, cur, data); // 执行后： // cur->buffer = buf_mmap // cur->data = data (指向mmap虚拟地址) } 注意：mmap地址是虚拟地址，实际数据仍在文件页中，由OS管理。 权重索引表构建 weights_map构建流程 llama_model_loader::llama_model_loader(...) { // 1. 解析GGUF元数据 meta.reset(gguf_init_from_file(fname.c_str(), params)); ggml_context * ctx = contexts[0]; // 2. 遍历GGML context中的所有tensor元数据 for (ggml_tensor * cur = ggml_get_first_tensor(ctx); cur; cur = ggml_get_next_tensor(ctx, cur)) { std::string tensor_name = std::string(cur->name); // 3. 计算文件偏移 size_t offs = gguf_get_data_offset(gguf_ctx) + gguf_get_tensor_offset(gguf_ctx, tensor_idx); // 4. 插入索引表（使用自定义排序） weights_map.emplace(tensor_name, llama_tensor_weight(files[0].get(), 0, meta.get(), cur)); } // 5. 处理分片文件（如果有） for (const auto & split : splits) { // 解析分片GGUF // 合并tensor到weights_map（idx递增） } } weight_name_comparer排序规则 struct weight_name_comparer { bool operator()(const std::string & a, const std::string & b) const { int bid_a = -1, bid_b = -1; // 解析层号（如 "blk.5.attn_q.weight" -> 5） sscanf(a.c_str(), "blk.%d.", &bid_a); sscanf(b.c_str(), "blk.%d.", &bid_b); // 按层号排序 if (bid_a != bid_b) { return bid_a < bid_b; } // 同层按名称排序 return a < b; } }; 排序目的：优化磁盘读取顺序，利用OS预读缓存。 数据加载路径 详见 阶段五：权重数据加载 章节。 设备管理与多GPU支持 设备发现与选择 自动设备发现 // 在llama_model_load_from_file_impl中 std::vector<ggml_backend_dev_t> gpus; std::vector<ggml_backend_dev_t> igpus; // 集成GPU std::vector<ggml_backend_dev_t> rpc_servers; // RPC服务器 for (size_t i = 0; i < ggml_backend_dev_count(); ++i) { ggml_backend_dev_t dev = ggml_backend_dev_get(i); switch (ggml_backend_dev_type(dev)) { case GGML_BACKEND_DEVICE_TYPE_GPU: { // 检查设备ID是否重复 ggml_backend_dev_props props; ggml_backend_dev_get_props(dev, &props); auto it = std::find_if(gpus.begin(), gpus.end(), [&props](ggml_backend_dev_t d) { // 比较device_id return /* 相同ID */; }); if (it == gpus.end()) { gpus.push_back(dev); } break; } case GGML_BACKEND_DEVICE_TYPE_IGPU: igpus.push_back(dev); break; } } // 优先级：RPC服务器 > 独立GPU > 集成GPU（仅当无其他设备时） model->devices.insert(model->devices.begin(), rpc_servers.begin(), rpc_servers.end()); model->devices.insert(model->devices.end(), gpus.begin(), gpus.end()); if (model->devices.empty()) { model->devices.insert(model->devices.end(), igpus.begin(), igpus.end()); } 单GPU模式 if (params.split_mode == LLAMA_SPLIT_MODE_NONE) { if (params.main_gpu < 0) { model->devices.clear(); // 只用CPU } else { // 只保留指定的GPU ggml_backend_dev_t main_gpu = model->devices[params.main_gpu]; model->devices.clear(); model->devices.push_back(main_gpu); } } 层到设备的映射策略 分割点计算 // 默认：按显存比例 std::vector<float> splits(n_devices()); if (all_zero) { // tensor_split全为0 size_t total_free = 0; for (size_t i = 0; i < n_devices(); ++i) { ggml_backend_dev_props props; ggml_backend_dev_get_props(devices[i], &props); total_free += props.memory_free; } float acc = 0.0f; for (size_t i = 0; i < n_devices(); ++i) { ggml_backend_dev_props props; ggml_backend_dev_get_props(devices[i], &props); acc += props.memory_free / total_free; splits[i] = acc; // 累积比例 } } else { // 使用用户指定的tensor_split float acc = 0.0f; for (size_t i = 0; i < n_devices(); ++i) { acc += tensor_split[i]; splits[i] = acc; } } 层分配函数 auto get_layer_buft_list = [&](int il) -> layer_dev { const bool is_swa = il < int(hparams.n_layer) && hparams.is_swa(il); // CPU层 if (il < i_gpu_start || (il - i_gpu_start) >= act_gpu_layers) { return {cpu_dev, &pimpl->cpu_buft_list}; } // GPU层：根据splits查找 float layer_ratio = float(il - i_gpu_start) / act_gpu_layers; const int layer_gpu = std::upper_bound(splits.begin(), splits.begin() + n_devices(), layer_ratio) - splits.begin(); auto * dev = devices.at(layer_gpu); return {dev, &pimpl->gpu_buft_list.at(dev)}; }; 示例：3个GPU，splits = [0.33, 0.67, 1.0] 层0-10 → GPU 0 层11-21 → GPU 1 层22-32 → GPU 2 Buffer类型选择 Buffer类型列表构建 // CPU Buffer类型列表 buft_list_t make_cpu_buft_list(...) { buft_list_t buft_list; // 1. Split buffer（如果启用） if (tensor_split && n_devices() > 1) { // 添加split buffer类型 } // 2. 默认CPU buffer buft_list.emplace_back(cpu_dev, ggml_backend_dev_buffer_type(cpu_dev)); // 3. Extra buffer类型（如果启用） if (use_extra_bufts) { auto * extra_bufts = ggml_backend_dev_get_extra_bufts(cpu_dev); // 添加extra buffer } return buft_list; } // GPU Buffer类型列表 buft_list_t make_gpu_buft_list(ggml_backend_dev_t dev, ...) { buft_list_t buft_list; // 1. Split buffer（如果启用） if (split_mode == LLAMA_SPLIT_MODE_LAYER) { // 添加split buffer } // 2. 默认GPU buffer buft_list.emplace_back(dev, ggml_backend_dev_buffer_type(dev)); // 3. Host buffer（用于CPU-GPU传输） auto * host_buft = ggml_backend_dev_host_buffer_type(dev); if (host_buft) { buft_list.emplace_back(dev, host_buft); } // 4. CPU buffer作为fallback buft_list.insert(buft_list.end(), cpu_buft_list.begin(), cpu_buft_list.end()); return buft_list; } 权重Buffer类型选择 ggml_backend_buffer_type_t select_weight_buft( const llama_hparams & hparams, ggml_tensor * w, ggml_op op, const buft_list_t & buft_list) { // 1. 检查权重是否支持该buffer类型 for (const auto & [dev, buft] : buft_list) { if (weight_buft_supported(hparams, w, op, buft, dev)) { return buft; } } // 2. 回退到CPU return ggml_backend_cpu_buffer_type(); } 选择策略： 优先选择GPU buffer（如果支持该操作） 回退到Host buffer（用于传输） 最后回退到CPU buffer 异步上传机制 异步上传初始化 ggml_backend_t upload_backend = [&]() -> ggml_backend_t { if (use_mmap || check_tensors) { return nullptr; // mmap或验证模式下不使用异步上传 } // 检查backend是否支持异步上传 auto * buf = bufs.count(0) ? bufs.at(0) : nullptr; if (!buf) return nullptr; auto * buft = ggml_backend_buffer_get_type(buf); auto * dev = ggml_backend_buft_get_device(buft); if (!dev) return nullptr; // 检查能力 ggml_backend_dev_props props; ggml_backend_dev_get_props(dev, &props); if (!props.caps.async || !props.caps.host_buffer || !props.caps.events) { return nullptr; } // 创建pinned memory buffers auto * host_buft = ggml_backend_dev_host_buffer_type(dev); for (size_t idx = 0; idx < n_buffers; ++idx) { auto * buf = ggml_backend_buft_alloc_buffer(host_buft, buffer_size); host_buffers.emplace_back(buf); host_ptrs.emplace_back(ggml_backend_buffer_get_base(buf)); auto * event = ggml_backend_event_new(dev); events.emplace_back(event); } return ggml_backend_dev_init(dev, nullptr); }(); 异步上传流程 // 在load_all_data中 if (upload_backend) { size_t offset = weight->offs; size_t aligned_offset = offset & ~(alignment - 1); size_t offset_from_alignment = offset - aligned_offset; file->seek(aligned_offset, SEEK_SET); size_t bytes_read = 0; size_t data_read = 0; while (bytes_read < read_end - read_start) { size_t read_size = std::min<size_t>(buffer_size, read_end - read_start - bytes_read); // 等待上一个上传完成（避免覆盖） ggml_backend_event_synchronize(events[buffer_idx]); // 读取对齐块到pinned memory file->read_raw_unsafe(host_ptrs[buffer_idx], read_size); // 计算实际数据部分（排除对齐填充） uintptr_t ptr_data = /* 对齐后的指针 */; size_t data_to_copy = read_size; if (bytes_read == 0) { ptr_data += offset_from_alignment; data_to_copy -= offset_from_alignment; } // 异步上传到GPU（不阻塞） ggml_backend_tensor_set_async(upload_backend, cur, reinterpret_cast<void *>(ptr_data), data_read, data_to_copy); // 记录事件（用于后续同步） ggml_backend_event_record(events[buffer_idx], upload_backend); data_read += data_to_copy; bytes_read += read_size; buffer_idx = (buffer_idx + 1) % n_buffers; // 轮转缓冲区 } } 优势： 重叠I/O和计算：读取下一块数据时，上一块正在上传 多缓冲区轮转：避免等待上传完成 对齐读取：提高I/O效率 内存管理优化 内存映射（mmap）优化 详见 内存映射（mmap）机制 章节。 补充优化点： 预填充页面：MAP_POPULATE标志立即读取所有页面（适用于小模型） 预读建议：POSIX_MADV_WILLNEED告诉OS预读页面 NUMA感知：POSIX_MADV_RANDOM优化NUMA系统访问模式 内存锁定（mlock）机制 mlock将内存页面锁定在RAM中，防止被swap到磁盘。 struct llama_mlock { void init(void * ptr) { // 初始化mlock } void grow_to(size_t target_size) { // 扩展锁定区域 #ifdef _POSIX_MEMLOCK_RANGE if (mlock(ptr, target_size) != 0) { // 处理错误（可能需要root权限） } #endif } }; 使用场景： 实时推理系统（避免swap延迟） 需要root权限或CAP_IPC_LOCK能力 NUMA感知 NUMA（Non-Uniform Memory Access）系统需要优化内存访问模式。 // 在init_mappings中 bool is_numa = false; auto * dev = ggml_backend_dev_by_type(GGML_BACKEND_DEVICE_TYPE_CPU); if (dev) { auto * reg = ggml_backend_dev_backend_reg(dev); auto * is_numa_fn = (decltype(ggml_is_numa) *) ggml_backend_reg_get_proc_address(reg, "ggml_backend_cpu_is_numa"); if (is_numa_fn) { is_numa = is_numa_fn(); } } // 创建mmap时传递NUMA标志 std::unique_ptr<llama_mmap> mapping = std::make_unique<llama_mmap>(file.get(), prefetch ? -1 : 0, is_numa); NUMA优化： 使用POSIX_MADV_RANDOM提示OS随机访问模式 避免跨NUMA节点访问（如果可能） 内存碎片管理 llama.cpp通过以下策略减少内存碎片： 按Buffer类型分组：相同类型的tensor使用同一个context 对齐分配：tensor数据按16字节对齐 mmap零拷贝：避免额外分配 架构特定加载逻辑 架构识别流程 // 1. 从GGUF读取架构名称 std::string arch_name; ml.get_key(LLM_KV_GENERAL_ARCHITECTURE, arch_name, false); // 2. 转换为枚举 llm_arch arch = llm_arch_from_string(arch_name); // "llama" -> LLM_ARCH_LLAMA // "qwen" -> LLM_ARCH_QWEN // ... // 3. 设置模型架构 model.arch = arch; 张量命名约定 llama.cpp使用统一的张量命名约定： {prefix}.{layer_idx}.{tensor_name}.{suffix} 示例： - tok_embeddings.weight (输入层) - blk.0.attn_norm.weight (第0层，注意力归一化) - blk.0.attn_q.weight (第0层，Query投影) - blk.0.ffn_gate.weight (第0层，FFN门控) - output_norm.weight (输出层) - output.weight (输出权重) 架构特定前缀： Llama: blk.{i} Qwen: blk.{i} DeepSeek: blk.{i} Mamba: blk.{i} (但使用SSM层) 特殊架构处理 MoE架构（Mixtral/DeepSeek） if (n_expert > 0) { // 门控输入 layer.ffn_gate_inp = create_tensor( tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, 0); // 专家权重（3D张量：n_embd x n_ff x n_expert） layer.ffn_gate_exps = create_tensor( tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), {n_embd, n_ff, n_expert}, TENSOR_NOT_REQUIRED); layer.ffn_up_exps = create_tensor( tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), {n_embd, n_ff, n_expert}, 0); layer.ffn_down_exps = create_tensor( tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff, n_embd, n_expert}, 0); // 共享专家（DeepSeek） if (hparams.n_ff_shexp > 0) { layer.ffn_gate_shexp = create_tensor(...); layer.ffn_up_shexp = create_tensor(...); layer.ffn_down_shexp = create_tensor(...); } } Mamba架构 case LLM_ARCH_MAMBA: case LLM_ARCH_MAMBA2: // SSM归一化 layer.ssm_norm = create_tensor( tn(LLM_TENSOR_SSM_NORM, "weight", i), {n_embd}, 0); // SSM参数 layer.ssm_dt_rank = create_tensor(...); layer.ssm_in = create_tensor(...); layer.ssm_out = create_tensor(...); // ... break; RWKV架构 case LLM_ARCH_RWKV6: case LLM_ARCH_RWKV7: // RWKV使用循环层，不使用标准注意力 // 特殊的time_mix参数 layer.time_mix_k = create_tensor(...); layer.time_mix_v = create_tensor(...); // ... break; 错误处理与进度跟踪 异常处理机制 try { llama_model_loader ml(...); model.load_arch(ml); model.load_hparams(ml); model.load_vocab(ml); model.load_tensors(ml); } catch (const std::exception & err) { LLAMA_LOG_ERROR("error loading model: %s\n", err.what()); return -1; } 常见错误： 文件不存在或格式错误 架构不支持 超参数缺失或无效 张量缺失或形状不匹配 内存不足 进度回调系统 typedef bool (*llama_progress_callback)(float progress, void * ctx); // 在load_all_data中 if (progress_callback) { float progress = (float) size_done / size_data; if (!progress_callback(progress, progress_callback_user_data)) { return false; // 用户取消 } } 使用示例： bool my_progress(float progress, void * ctx) { printf("Loading: %.1f%%\r", progress * 100); fflush(stdout); return true; // 返回false可取消加载 } llama_model_params params = llama_model_default_params(); params.progress_callback = my_progress; params.progress_callback_user_data = nullptr; 取消加载机制 进度回调返回false时，加载会立即停止： // 在load_all_data中 if (progress_callback) { if (!progress_callback((float) size_done / size_data, progress_callback_user_data)) { return false; // 取消加载 } } // 在llama_model_load中 if (!model.load_tensors(ml)) { return -2; // 返回-2表示取消 } 数据结构关联关系总结 加载链路关联 ┌─────────────────────────────────────────────────────────────┐ │ 阶段 1: GGUF 读取 │ │ │ │ gguf_init_from_file() │ │ ↓ │ │ 解析所有 KV 对 → 存储在 gguf_context->kv 中 │ │ - "general.architecture" = "llama" │ │ - "llama.context_length" = 4096 │ │ - "llama.embedding_length" = 4096 │ │ - "llama.block_count" = 32 │ │ │ │ 结果：ml.meta (gguf_context_ptr) 持有所有 KV 对 │ └─────────────────────────────────────────────────────────────┘ ↓ ┌─────────────────────────────────────────────────────────────┐ │ 阶段 2: 架构识别 │ │ │ │ model.load_arch(ml) │ │ ↓ │ │ ml.get_arch() → 从gguf_context读取架构名称 │ │ ↓ │ │ model.arch = LLM_ARCH_LLAMA │ └─────────────────────────────────────────────────────────────┘ ↓ ┌─────────────────────────────────────────────────────────────┐ │ 阶段 3: 超参数加载 │ │ │ │ model.load_hparams(ml) │ │ ↓ │ │ ml.get_key(LLM_KV_CONTEXT_LENGTH, hparams.n_ctx_train) │ │ ↓ │ │ 从gguf_context->kv查找并填充hparams │ │ │ │ 结果：model.hparams 包含所有数学参数 │ └─────────────────────────────────────────────────────────────┘ ↓ ┌─────────────────────────────────────────────────────────────┐ │ 阶段 4: 张量创建 │ │ │ │ model.load_tensors(ml) │ │ ↓ │ │ 1. 根据hparams.n_layer创建layers数组 │ │ layers.resize(n_layer) │ │ │ │ 2. 从weights_map查找张量元数据 │ │ weight = ml.get_weight("blk.0.attn_q.weight") │ │ │ │ 3. 创建tensor对象（空壳） │ │ layer.wq = ml.create_tensor(ctx, tn, {n_embd, ...}, 0) │ │ │ │ 4. 分配设备 │ │ pimpl->dev_layer[0] = get_layer_buft_list(0) │ │ │ │ 结果：所有tensor对象已创建，但data仍为nullptr │ └─────────────────────────────────────────────────────────────┘ ↓ ┌─────────────────────────────────────────────────────────────┐ │ 阶段 5: 数据加载 │ │ │ │ ml.load_all_data() │ │ ↓ │ │ 遍历所有tensor： │ │ for (tensor in ctx) { │ │ weight = get_weight(tensor->name) │ │ if (use_mmap) { │ │ data = mapping->addr() + weight->offs │ │ ggml_backend_tensor_alloc(buf, tensor, data) │ │ // tensor->data = data (指向mmap地址) │ │ } else { │ │ file->read_raw(tensor->data, n_size) │ │ } │ │ } │ │ │ │ 结果：tensor->data指向实际权重数据 │ └─────────────────────────────────────────────────────────────┘ 数据结构职责划分 数据结构 职责 生命周期 gguf_context GGUF文件元数据存储 加载期间 llama_model_loader 权重索引、文件管理、数据加载 加载期间 llama_hparams 数学参数存储 模型生命周期 llama_model 权重指针、设备管理、元数据 模型生命周期 llama_layer 单层权重指针 模型生命周期 ggml_tensor 张量元数据+数据指针 模型生命周期 ggml_backend_buffer 物理内存管理 模型生命周期 关键设计模式 Pimpl模式：隐藏实现细节，保持API稳定 索引表模式：weights_map作为"目录"，快速定位权重 策略模式：设备分配、Buffer选择通过策略函数实现 工厂模式：create_tensor根据架构创建不同张量 总结 GGUF读取的调用链 GGUF文件的读取在llama.cpp中遵循以下调用链： llama_model_load_from_file_impl() ↓ llama_model_load() ↓ llama_model_loader::llama_model_loader() // 构造函数 ↓ gguf_init_from_file() // 打开文件并调用解析函数 ↓ gguf_init_from_file_impl() // 实际解析函数 ↓ gguf_reader // 文件读取器 关键调用点： // 在 llama_model_loader 构造函数中（llama-model-loader.cpp:531） struct gguf_init_params params = { /*.no_alloc = */ true, // 只读取元数据，不分配tensor数据 /*.ctx = */ &ctx, // 返回ggml_context用于存储tensor元数据 }; meta.reset(gguf_init_from_file(fname.c_str(), params)); // meta 是 gguf_context_ptr，持有解析后的GGUF元数据 GGUF解析的详细流程 gguf_init_from_file_impl 按以下顺序解析文件： 步骤1：校验Magic Number std::vector<char> magic; gr.read(magic, 4); // 读取4字节 // 验证是否为 "GGUF" 步骤2：读取文件头 gr.read(ctx->version); // 读取版本号（4字节） gr.read(n_kv); // 读取KV对数量（8字节） gr.read(n_tensors); // 读取张量数量（8字节） 步骤3：解析所有KV对 for (int64_t i = 0; i < n_kv; ++i) { std::string key = gr.read_string(); // 读取键名 gguf_type type = (gguf_type)gr.read_u32(); // 读取类型 // 根据类型读取值 switch (type) { case GGUF_TYPE_UINT8: ctx->kv.emplace_back(key, gr.read_u8()); break; case GGUF_TYPE_INT32: ctx->kv.emplace_back(key, gr.read_i32()); break; case GGUF_TYPE_FLOAT32: ctx->kv.emplace_back(key, gr.read_f32()); break; case GGUF_TYPE_STRING: ctx->kv.emplace_back(key, gr.read_string()); break; // ... 更多类型 } } // 结果：所有KV对存储在 ctx->kv 向量中 步骤4：解析所有张量信息 for (int64_t i = 0; i < n_tensors; ++i) { struct gguf_tensor_info info; info.name = gr.read_string(); // 张量名称 info.n_dims = gr.read_u32(); // 维度数 for (uint32_t j = 0; j < info.n_dims; ++j) { info.dims[j] = gr.read_u64(); // 各维度大小 } info.type = (ggml_type)gr.read_u32(); // 数据类型 info.offset = gr.read_u64(); // 数据偏移（相对数据区） ctx->info.push_back(info); } // 结果：所有张量信息存储在 ctx->info 向量中 步骤5：计算数据区偏移 ctx->offset = gr.tell(); // 当前位置就是数据区开始位置 // 此时文件指针已跳过所有元数据，指向实际权重数据 步骤6：创建GGML Context和Tensor元数据（如果no_alloc=false） if (params.ctx && !params.no_alloc) { // 创建ggml_context ggml_context * ctx_data = ggml_init({...}); // 为每个tensor创建ggml_tensor对象 for (size_t i = 0; i < ctx->info.size(); ++i) { struct ggml_tensor * cur = ggml_new_tensor(ctx_data, ...); ggml_set_name(cur, info.t.name); // 如果no_alloc=false，将data指针指向文件中的数据位置 cur->data = (char *) data->data + info.offset; } } 读取后存储的数据结构 GGUF解析后的数据存储在以下结构中： gguf_context（核心存储结构） struct gguf_context { uint32_t version; // GGUF版本 std::vector<gguf_kv> kv; // 所有KV对 // 例如： // kv[0] = {"general.architecture", "llama"} // kv[1] = {"llama.context_length", 4096} // kv[2] = {"llama.embedding_length", 4096} // ... std::vector<gguf_tensor_info> info; // 所有张量信息 // 例如： // info[0] = {name: "tok_embeddings.weight", dims: [4096, 32000], type: F16, offset: 0} // info[1] = {name: "blk.0.attn_q.weight", dims: [4096, 4096], type: F16, offset: 131072000} // ... size_t alignment; // 对齐字节数（通常16或32） size_t offset; // 数据区在文件中的偏移 void * data; // 数据区指针（如果已映射） }; llama_model_loader中的存储 struct llama_model_loader { gguf_context_ptr meta; // 持有gguf_context // meta.get() 返回 gguf_context*，包含所有KV对和张量信息 llama_files files; // 文件句柄列表 // files[0] = 主文件 // files[1..n] = 分片文件（如果有） std::map<std::string, llama_tensor_weight> weights_map; // 权重索引表，Key是tensor名称，Value包含： // - idx: 文件索引（0=主文件，1..n=分片） // - offs: 在文件中的字节偏移 // - tensor: 指向ggml_tensor元数据 }; 在llama.cpp中的调用流程 完整调用链： // 1. 用户调用入口 llama_model * model = llama_model_load_from_file("model.gguf", params); // 2. 内部调用链 llama_model_load_from_file_impl() → llama_model_load() → llama_model_loader ml(fname, ...) // 构造函数开始解析GGUF → gguf_init_from_file(fname, params) → gguf_init_from_file_impl(file, params) // 解析文件头、KV对、张量信息 // 返回 gguf_context* → meta.reset(gguf_context*) // 保存到loader中 → 遍历ctx中的tensor，构建weights_map → 处理分片文件（如果有） // 3. 后续使用 model.load_arch(ml) // 从meta读取架构名称 model.load_hparams(ml) // 从meta读取超参数 model.load_vocab(ml) // 从meta读取词表 model.load_tensors(ml) // 从weights_map定位权重，从文件读取数据 关键访问接口： // 在 llama_model_loader 中访问KV对 ml.get_key(LLM_KV_CONTEXT_LENGTH, hparams.n_ctx_train); // 内部调用： // gguf_find_key(meta.get(), "llama.context_length") // gguf_get_val_u32(meta.get(), key_id) // 访问张量元数据 ggml_tensor * t_meta = ml.get_tensor_meta("blk.0.attn_q.weight"); // 内部从weights_map查找，返回ggml_tensor*（只有元数据，data为nullptr） // 访问权重位置 const auto * weight = ml.get_weight("blk.0.attn_q.weight"); // 返回 llama_tensor_weight，包含： // weight->idx = 0（主文件） // weight->offs = 131072000（文件偏移） // weight->tensor = ggml_tensor*（元数据） 数据流转图 ┌───────────────────────────────────────────────┐ │ GGUF文件（磁盘） │ │ ┌─────────┬─────────┬─────────┬─────────┐ │ │ │ Magic │ Version │ n_kv │ KV Pairs │ │ │ └─────────┴─────────┴─────────┴─────────┘ │ │ ┌─────────┬─────────────────────────────┐ │ │ │n_tensors │ Tensor Info Array │ │ │ └─────────┴─────────────────────────────┘ │ │ ┌───────────────────────────────────────┐ │ │ │ Tensor Data (实际权重) │ │ │ └───────────────────────────────────────┘ │ └──────────────────────────────────────────────┘ ↓ gguf_init_from_file() ┌────────────────────────────────────────────┐ │ gguf_context（内存） │ │ │ │ ctx->kv[]: │ │ [0] = {"general.architecture", "llama"} │ │ [1] = {"llama.context_length", 4096} │ │ [2] = {"llama.embedding_length", 4096} │ │ ... │ │ │ │ ctx->info[]: │ │ [0] = {name: "tok_embeddings.weight", offset: 0, ...} │ │ [1] = {name: "blk.0.attn_q.weight", offset: 131072000,...} │ │ ... │ │ │ │ ctx->offset = 数据区偏移 │ └────────────────────────────────────────────┘ ↓ 存储到 llama_model_loader ┌────────────────────────────────────────────┐ │ llama_model_loader │ │ │ │ meta: gguf_context_ptr → 持有所有KV和tensor信息 │ │ │ │ weights_map: │ │ "tok_embeddings.weight" → {idx:0, offs:0, tensor:...} │ │ "blk.0.attn_q.weight" → {idx:0, offs:131072000, tensor:...} │ │ ... │ │ │ │ files: [主文件句柄, 分片1句柄, ...] │ └────────────────────────── ─────────────────┘ ↓ 使用 ┌───────────────────────────────────────────┐ │ 模型加载各阶段 │ │ │ │ load_arch(): 从 meta->kv 读取架构名称 │ │ load_hparams(): 从 meta->kv 读取超参数 │ │ load_vocab(): 从 meta->kv 读取词表 │ │ load_tensors(): 从 weights_map 定位权重，从文件读取数据 │ └───────────────────────────────────────────┘

后端系统概述 GGML后端系统主要提供如下功能： 统一接口： 不同硬件平台使用相关的API。 自动选择：根据硬件自动选择最优后端。 灵活切换：可以在运行时切换后端。 扩展性：易于添加的新后端。 后端主要涉及如下概念： Backend：执行计算的抽象层（CPU，CUDA， Metal等） Buffer：后端管理的内存缓冲区。 Buffer Type：缓冲区的类型（主机内存、GPU内存等） Graph Plan：计算图的执行。 Devcice：物理设备（CPU，GPU等） 后端接口 核心类型 typedef struct ggml_backend * ggml_backend_t; // 后端句柄 typedef struct ggml_backend_buffer * ggml_backend_buffer_t; // 缓冲区句柄 typedef struct ggml_backend_buffer_type * ggml_backend_buffer_type_t; // 缓冲区类型 主要的API （1）后端管理 // 获取后端名称 const char * ggml_backend_name(ggml_backend_t backend); // 释放后端 void ggml_backend_free(ggml_backend_t backend); // 获取默认缓冲区类型 ggml_backend_buffer_type_t ggml_backend_get_default_buffer_type(ggml_backend_t backend); （2）缓冲区的操作 // 分配缓冲区 ggml_backend_buffer_t ggml_backend_alloc_buffer(ggml_backend_t backend, size_t size); // 释放缓冲区 void ggml_backend_buffer_free(ggml_backend_buffer_t buffer); // 获取缓冲区基址 void * ggml_backend_buffer_get_base(ggml_backend_buffer_t buffer); （3）Tensor操作 // 设置 tensor 数据 void ggml_backend_tensor_set(struct ggml_tensor * tensor, const void * data, size_t offset, size_t size); // 获取 tensor 数据 void ggml_backend_tensor_get(const struct ggml_tensor * tensor, void * data, size_t offset, size_t size); （4）计算图执行 // 同步执行 enum ggml_status ggml_backend_graph_compute(ggml_backend_t backend, struct ggml_cgraph * cgraph); // 异步执行 enum ggml_status ggml_backend_graph_compute_async(ggml_backend_t backend, struct ggml_cgraph * cgraph); 后端注册机制 注册流程 位置: src/ggml-backend-reg.cpp 初始化时注册: struct ggml_backend_registry { std::vector<ggml_backend_reg_entry> backends; ggml_backend_registry() { #ifdef GGML_USE_CUDA register_backend(ggml_backend_cuda_reg()); #endif #ifdef GGML_USE_METAL register_backend(ggml_backend_metal_reg()); #endif #ifdef GGML_USE_CPU register_backend(ggml_backend_cpu_reg()); #endif // ... 其他后端 } }; 后端注册结构 struct ggml_backend_reg { const char * name; // 后端名称 ggml_backend_t (*init_fn)(void); // 初始化函数 size_t (*dev_count_fn)(void); // 设备数量 // ... 其他函数指针 }; 动态加载 支持从动态库加载后端 ggml_backend_reg_t ggml_backend_load(const char * path); 4. demo.c vs demo_backend.c 对比 demo.c (传统模式) （1）特点 使用 ggml_init()直接分配内存 所有 tensor 从 context 分配 使用 ggml_graph_compute_with_ctx()执行 简单直接，适合 CPU 计算 （2）代码流程 // 1. 初始化 context ctx = ggml_init(params); // 2. 创建 tensor (从 context 分配) tensor = ggml_new_tensor_2d(ctx, ...); // 3. 构建计算图 gf = ggml_new_graph(ctx); ggml_build_forward_expand(gf, result); // 4. 执行 (使用 context 的工作内存) ggml_graph_compute_with_ctx(ctx, gf, n_threads); （3）内存管理 Context 管理所有内存 预分配内存池 运行时不再分配 demo_backend.c (后端模式) （1）特点 使用后端系统管理内存 支持 GPU 等后端 使用 ggml_gallocr 分配内存 更灵活，适合生产环境 （2）代码流程 // 1. 初始化后端 backend = ggml_backend_cpu_init(); // 或 CUDA, Metal 等 // 2. 初始化 context (no_alloc = true) ctx = ggml_init(params_no_alloc); // 3. 创建 tensor (不分配内存) tensor = ggml_new_tensor_2d(ctx, ...); // 4. 后端分配内存 buffer = ggml_backend_alloc_ctx_tensors(ctx, backend); // 5. 设置数据 ggml_backend_tensor_set(tensor, data, ...); // 6. 构建计算图 gf = ggml_new_graph(ctx_cgraph); ggml_build_forward_expand(gf, result); // 7. 分配计算图内存 allocr = ggml_gallocr_new(...); ggml_gallocr_alloc_graph(allocr, gf); // 8. 后端执行 ggml_backend_graph_compute(backend, gf); （3）内存管理 后端管理内存分配 支持不同内存类型 (主机内存、GPU 内存) 使用 ggml_gallocr优化分配 关键差异 特性 demo.c demo_backend.c 内存管理 Context 直接管理 后端系统管理 GPU 支持 否 是 内存类型 仅主机内存 支持多种类型 5. 内存分配器 ggml_gallocr (Graph Allocator) 用于 优化计算图的内存分配 减少内存碎片 支持内存复用 Tensor分配 适用场景：你希望某些 tensor 放到特定的内存/设备（比如 Spacemit/特殊 CPU buffer），而另一些 tensor 仍放在默认 CPU buffer。核心 API 是： ggml_backend_buft_get_alloc_size()：算出某个 buft 下该 tensor 需要的实际分配大小 ggml_backend_buft_alloc_buffer()：分配一块后端 buffer ggml_backend_buffer_get_base()：拿到 buffer 的 base 地址 ggml_backend_tensor_alloc(buffer, tensor, base + offset)：把 tensor “绑定”到这块 buffer 的地址上 典型流程（精简版）： // 0) 构建 tensor 时建议 ctx.no_alloc = true（只建元数据，不在 ctx 内部分配） struct ggml_tensor * a = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, cols_A, rows_A); struct ggml_tensor * b = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, cols_B, rows_B); // 1) 为 a 选择特殊 buft（如 Spacemit），为 b 选择默认 CPU buft ggml_backend_buffer_type_t a_buft = ggml_backend_cpu_riscv64_spacemit_buffer_type(); ggml_backend_buffer_type_t cpu_buft = ggml_backend_get_default_buffer_type(backend); // 2) 分配 buffer 并把 tensor 绑定进去 size_t a_size = ggml_backend_buft_get_alloc_size(a_buft, a); ggml_backend_buffer_t a_buf = ggml_backend_buft_alloc_buffer(a_buft, a_size); ggml_backend_tensor_alloc(a_buf, a, ggml_backend_buffer_get_base(a_buf)); size_t b_size = ggml_backend_buft_get_alloc_size(cpu_buft, b); ggml_backend_buffer_t b_buf = ggml_backend_buft_alloc_buffer(cpu_buft, b_size); ggml_backend_tensor_alloc(b_buf, b, ggml_backend_buffer_get_base(b_buf)); // 3) 写入/读取 tensor 数据（后端负责必要的数据搬运） ggml_backend_tensor_set(a, host_ptr_A, 0, ggml_nbytes(a)); ggml_backend_tensor_set(b, host_ptr_B, 0, ggml_nbytes(b)); // ... graph compute ... ggml_backend_tensor_get(out, host_ptr_out, 0, ggml_nbytes(out)); 要点： buffer 的生命周期必须覆盖计算阶段（结束后再 ggml_backend_buffer_free()）。 这种方式本质上是“你负责 placement”，适合做 demo/异构内存验证/特殊平台适配。 Graph分配 适用场景：你希望把“中间 tensor + 输出 tensor”的内存分配交给 allocator 统一规划，减少碎片、提升复用率。前提通常是： 用于建图的 ctx 设置.no_alloc = true（tensor 先不分配） 图 gf 已经 ggml_build_forward_expand()完整展开 典型流程： // 1) 创建分配器：通常用 backend 的默认 buffer type ggml_gallocr_t allocr = ggml_gallocr_new( ggml_backend_get_default_buffer_type(backend) ); // 2) 为计算图分配内存（会给图里需要的 tensor 规划并分配后端 buffer） ggml_gallocr_alloc_graph(allocr, gf); // 3) 执行计算图 ggml_backend_graph_compute(backend, gf); // 4) 释放分配器（以及其内部持有的资源） ggml_gallocr_free(allocr); 可选优化（大模型/多次重复执行同结构图时更常见）： ggml_gallocr_reserve(allocr, gf)：先“规划/预留”一次，再根据 ggml_gallocr_get_buffer_size()观察峰值需求（见一些大 example 的用法）。 Workflow分配 适用场景：你有多个后端（例如主后端 + CPU fallback），希望由 scheduler/allocator 完成： 图的 tensor 分配 跨后端的 placement/拷贝（如果需要） simple/simple-backend.cpp 的核心思路是： 建图时 .no_alloc = true 每次计算前调用 ggml_backend_sched_alloc_graph()触发分配 用 ggml_backend_tensor_set/get在 host 与后端 tensor 之间传数据 典型流程（精简版）： // 1) 初始化多个 backend，并创建 sched ggml_backend_t backend = ggml_backend_init_best(); ggml_backend_t cpu_backend = ggml_backend_init_by_type(GGML_BACKEND_DEVICE_TYPE_CPU, NULL); ggml_backend_t backends[] = { backend, cpu_backend }; ggml_backend_sched_t sched = ggml_backend_sched_new(backends, NULL, 2, GGML_DEFAULT_GRAPH_SIZE, /*pipeline_parallel=*/false, /*graph_copy=*/true); // 2) 建图（ctx.no_alloc=true；图结构固定后可重复执行） struct ggml_cgraph * gf = ...; // ggml_new_graph + ggml_build_forward_expand // 3) 分配 + 执行 ggml_backend_sched_reset(sched); ggml_backend_sched_alloc_graph(sched, gf); ggml_backend_tensor_set(a, host_ptr_A, 0, ggml_nbytes(a)); ggml_backend_tensor_set(b, host_ptr_B, 0, ggml_nbytes(b)); ggml_backend_sched_graph_compute(sched, gf); struct ggml_tensor * out = ggml_graph_node(gf, -1); ggml_backend_tensor_get(out, host_ptr_out, 0, ggml_nbytes(out)); 要点： 这条路线更接近“生产用法”：分配/拷贝/执行由 backend 系统统一处理，你只维护图和输入输出。 如果你需要“像 demo_spacemit.c 那样手动把某个输入 tensor 固定到特定 buft”，就不要完全依赖 scheduler 的自动分配；通常应采用tensor 级手动分配或者在更高层做 placement 策略（取决于后端能力）。

加载后端 void ggml_backend_load_all() { ggml_backend_load_all_from_path(nullptr); } void ggml_backend_load_all_from_path(const char * dir_path) { #ifdef NDEBUG bool silent = true; #else bool silent = false; #endif ggml_backend_load_best("blas", silent, dir_path); ggml_backend_load_best("zendnn", silent, dir_path); ggml_backend_load_best("cann", silent, dir_path); ggml_backend_load_best("cuda", silent, dir_path); ggml_backend_load_best("hip", silent, dir_path); ggml_backend_load_best("metal", silent, dir_path); ggml_backend_load_best("rpc", silent, dir_path); ggml_backend_load_best("sycl", silent, dir_path); ggml_backend_load_best("vulkan", silent, dir_path); ggml_backend_load_best("opencl", silent, dir_path); ggml_backend_load_best("hexagon", silent, dir_path); ggml_backend_load_best("musa", silent, dir_path); ggml_backend_load_best("cpu", silent, dir_path); // check the environment variable GGML_BACKEND_PATH to load an out-of-tree backend const char * backend_path = std::getenv("GGML_BACKEND_PATH"); if (backend_path) { ggml_backend_load(backend_path); } } 动态加载所有可用的计算后端（CPU、CUDA、Metal 等），让 llama.cpp 能在不同硬件上运行。通过评分机制选择最适合的后端进行注册。每个后端会提供一个获取评分的函数，使用该函数进行计算得到分值，比如使用位运算累加看看性能。注册后端注释是加载其对应后端的动态库，将后端添加到backends向量。 加载模型 先看看看GGUF的模型文件。 （1）文件头： 位于文件的最开始位置，用于快速校验文件是否合法。 Magic Number (4 bytes): 0x47 0x47 0x55 0x46，对应的 ASCII 码就是 "GGUF"。加载器首先读取这4个字节，如果不对，直接报错。 Version (4 bytes): 版本号（图中是 3）。这允许加载器处理不同版本的格式变化（向后兼容）。 Tensor Count (8 bytes, UINT64): 模型中包含的张量（权重矩阵）总数。比如一个 7B 模型可能有几百个 Tensor。 Metadata KV Count (8 bytes, UINT64): 元数据键值对的数量。这是 GGUF 最核心的改进点。 （2）元数据键值对：接在头部之后，这部分定义了模型的所有非权重信息 以前的格式（如 GGML）通常把超参数（层数、维度等）写死在代码或特定的结构体里。GGUF 改用 Key-Value 映射。 架构信息: general.architecture = "llama"（告诉加载器用哪套逻辑加载）。 模型参数: llama.context_length = 4096, llama.embedding_length = 4096 等。 Tokenizer (词表): 这一点非常重要。GGUF 将词表（tokenizer.ggml.tokens）直接打包在文件里，不再需要额外的 tokenizer.json 或 vocab.model 文件。这实现了真正的“单文件发布”。 其他: 作者信息、量化方法描述等。 （3）张量信息表：索引目录 在读取完所有元数据后，就是 Tensor 的索引列表。这里不包含实际的权重数据，只是数据的“描述”和“指针”。对于每一个 Tensor（共 tensor_count 个），包含以下字段： Name: 字符串，例如 blk.0.ffn_gate.weight。这对应了模型层中的具体权重名称。 n_dimensions: 维度数量（例如 2 表示矩阵，1 表示向量）。 dimensions: 具体的形状数组，例如 [4096, 32000]。 Type: 数据类型，例如 GGML_TYPE_Q2_K（2-bit 量化）、F16 等。 Offset (8 bytes, UINT64): 这是最关键的字段。它是一个指针（偏移量），告诉程序：“这个张量的实际二进制数据，位于文件数据区的第 X 个字节处”。 （4）数据区：图中右侧箭头指向的 10110... 部分。 这是文件体积最大的部分，包含所有权重矩阵的实际二进制数据。内存映射 (mmap) 的奥秘，因为有了上面的 Offset，llama_model_load_from_file 不需要把整个几 GB 的文件一次性读入 RAM。 它只需要调用系统级 API mmap，将文件映射到虚拟内存。当模型推理需要用到 blk.0 的权重时，CPU/GPU 根据 Offset 直接去磁盘（或文件系统缓存）里拿数据。这就是为什么 GGUF 模型启动速度极快且内存占用低的原因。 模型加载主要是将GGUF加载进来，主要的流程是先检查是否有注册后端，然后选择对应的设备，接着解析GGUF元信息包括KV/张量表，接着按照n_gpu_layers/tensor_split分配层与buffer，再者就是读取/映射/上传权重数据，返回llama_model对象指针。 其中最核心的函数就是llama_mode_load，这个函数主要的作用如下： load_arch：先知道是什么架构，后面的 KV key 解释/张量命名规则才对。 load_hparams：读 GGUF KV 填充超参。 load_vocab：构建 tokenizer/vocab。 load_tensors：决定设备/缓冲区并实际装载权重数据。 load_arch load_arch加载的是模型的架构类型，比如LLaMA、Falcon、Qwen、CLIP等大类，存储到llama_mode::arch字段里面，这个来源是GGUF的metadata key。 get_key(llm_kv(LLM_KV_GENERAL_ARCHITECTURE), arch_name, false); llm_kv = LLM_KV(llm_arch_from_string(arch_name)); load_arch本身只是把loader里解析好的枚举出来。 void llama_model::load_arch(llama_model_loader & ml) { arch = ml.get_arch(); if (arch == LLM_ARCH_UNKNOWN) { throw std::runtime_error("unknown model architecture: '" + ml.get_arch_name() + "'"); } } 之所以要解析加载arch是因为后面所以得操作都依赖着arch，比如怎么解释GGUF、怎么建张量的逻辑依赖arch。l oad_hparams 需要 arch 才知道该读哪些 KV、如何解释超参，而且它还用 arch 做特殊分支（例如 CLIP 直接跳过 hparams）。 load_vocab 会用 arch 构造 LLM_KV(arch)，影响 vocab/tokenizer 的加载规则。 load_tensors 更是直接 switch (arch) 来决定要创建哪些权重张量、张量命名/shape/op 映射等。 load_hparams struct llama_hparams { bool vocab_only; bool no_alloc; bool rope_finetuned; bool use_par_res; bool swin_norm; uint32_t n_ctx_train; // context size the model was trained on uint32_t n_embd; uint32_t n_embd_features = 0; uint32_t n_layer; // ... uint32_t n_rot; uint32_t n_embd_head_k; uint32_t n_embd_head_v; uint32_t n_expert = 0; uint32_t n_expert_used = 0; // ... }; 这里的 hparams 指的是 llama_model 里的 模型超参数/结构参数（hyper-parameters）：决定模型“长什么样”（层数、隐藏维度、头数、RoPE/ALiBi/SWA/MoE 等配置）。它对应的结构体是 struct llama_hparams llama_model::load_hparams(ml) 会从 GGUF 的 KV（metadata）里读取并填充 model.hparams（以及一些辅助状态，比如 gguf_kv 字符串表、type 等） 模型结构： 训练时的context长度、embedding/hidden size、transformer block 数、MoE等。 注意力/FFN：支持 KV 是标量或数组：有些模型每层不同。 RoPE相关：rope_finetuned / rope_freq_base_train / rope_freq_scale_train / rope_scaling_type_train / n_rot / n_embd_head_k/v 等 架构特化项：比如 LLM_ARCH_WAVTOKENIZER_DEC 会额外读 posnet/convnext 参数。 load_vocab load_vocab 加载的是 模型的词表/分词器（tokenizer + vocab）：包括每个 token 的文本、score、属性（normal/control/user_defined/unknown/byte 等），以及各种 special token（BOS/EOS/EOT/PAD/…）和不同 tokenizer（SPM/BPE/WPM/UGM/RWKV/…）所需的附加数据（如 BPE merges、charsmap 等）。 struct llama_vocab { struct token_data { std::string text; float score; llama_token_attr attr; }; void load(llama_model_loader & ml, const LLM_KV & kv); std::string get_tokenizer_model() const; std::string get_tokenizer_pre() const; enum llama_vocab_type get_type() const; enum llama_vocab_pre_type get_pre_type() const; uint32_t n_tokens() const; uint32_t n_token_types() const; // ... }; 模型的推理需要做文本到token id的双向转换，输入的prompt必须先tokenize 才能喂给llama_decode，输出的token id必须detokenize 才能打印成字符串。不同模型（LLaMA / Falcon / ChatGLM / …）的 tokenizer 规则不同（SPM/BPE/WPM/UGM…，以及 pre-tokenization 规则），如果 vocab 没加载或加载错，模型就“无法正确理解输入/输出” load_tensor 这里的tensor是指模型推理所需的所有权重张量（weghts），比如： token embedding：tok_embd 每层 attention 的 Wq/Wk/Wv/Wo（以及可选 bias、norm 权重等） 每层 FFN 的 Wgate/Wup/Wdown（MoE 还会有专家张量、router 等） 输出层/输出 embedding（有些模型与 tok_embd 共享，需要 duplicated 处理） 在load tensors会按arch分支创建这些张量，比如已LLAMA类分支片段。 tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0); // output output_norm = create_tensor(tn(LLM_TENSOR_OUTPUT_NORM, "weight"), {n_embd}, 0); output = create_tensor(tn(LLM_TENSOR_OUTPUT, "weight"), {n_embd, n_vocab}, TENSOR_NOT_REQUIRED); if (output == NULL) { output = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, TENSOR_DUPLICATED); } for (int i = 0; i < n_layer; ++i) { auto & layer = layers[i]; layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, 0); layer.wq = create_tensor(tn(LLM_TENSOR_ATTN_Q, "weight", i), {n_embd, n_embd_head_k * n_head}, 0); layer.wk = create_tensor(tn(LLM_TENSOR_ATTN_K, "weight", i), {n_embd, n_embd_k_gqa}, 0); layer.wv = create_tensor(tn(LLM_TENSOR_ATTN_V, "weight", i), {n_embd, n_embd_v_gqa}, 0); layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", i), {n_embd_head_k * n_head, n_embd}, 0); } （1）tensor对象（ggml_tensor *）存在哪里？ 按语义分组存到 llama_model 的成员里：比如 tok_embd / output_norm / output，以及每层的 layers[i].wq/wk/wv/... 等（这些都是 ggml_tensor * 指针）。 struct llama_model { // ... struct ggml_tensor * tok_embd = nullptr; // ... struct ggml_tensor * output_norm = nullptr; struct ggml_tensor * output = nullptr; // ... std::vector<llama_layer> layers; // ... std::vector<std::pair<std::string, struct ggml_tensor *>> tensors_by_name; // ... }; 按名字查 tensor：load_tensors() 完成后会把每个 ggml_context 里的 tensor 都放进 tensors_by_name（主要用于内部/统计与 get_tensor(name) 这种按名访问）。 // populate tensors_by_name for (auto & [ctx, _] : pimpl->ctxs_bufs) { for (auto * cur = ggml_get_first_tensor(ctx.get()); cur != NULL; cur = ggml_get_next_tensor(ctx.get(), cur)) { tensors_by_name.emplace_back(ggml_get_name(cur), cur); } } 对应按名取 const ggml_tensor * llama_model::get_tensor(const char * name) const { auto it = std::find_if(tensors_by_name.begin(), tensors_by_name.end(), [name](const std::pair<std::string, ggml_tensor *> & it) { return it.first == name; }); if (it == tensors_by_name.end()) { return nullptr; } return it->second; } （2）tensor数据（权重内容）存在哪里？ ggml_tensor 自身只是一个“描述 + 指针”，真正的权重数据会落在三类地方之一： CPU/GPU 后端 buffer 里：load_tensors() 会为不同 buffer type 创建对应的 backend buffer，并把它们与承载 tensor metadata 的 ggml_context 绑在一起。这些 buffer 句柄被保存在：pimpl->ctxs_bufs struct llama_model::impl { // contexts where the model tensors metadata is stored as well as the corresponding buffers: std::vector<std::pair<ggml_context_ptr, std::vector<ggml_backend_buffer_ptr>>> ctxs_bufs; // ... }; mmap 映射区域（文件映射内存）：如果启用 mmap，loader 会把模型文件映射进内存，某些 tensor 的 data 会直接指向映射区域，或者后端 buffer 会通过 buffer_from_host_ptr 包装映射内存。映射对象被保存在loader侧：ml.mappings，model侧：pimpl->mappings。 no_alloc 模式：ml.no_alloc 时只是创建 tensor 元信息与“dummy buffer”，不装载数据，用于某些统计/规划场景。 （3）tensor是怎么分配到那个设备上的？ 在GPU上还是CPU上，决策分为两层： 层级/设备选择（dev）：决定每一层（input/repeating/output）用哪个 ggml_backend_dev_t（CPU / GPU / RPC / IGPU…）。 buffer type 选择（buft）：在选定设备后，再决定该 tensor 用哪个 ggml_backend_buffer_type_t（同一设备也可能有多种 buffer type；并且还会把 CPU buffer type 作为 fallback）。 流程梳理（从“设备列表”到“权重进后端 buffer”）如下： ┌──────────────────────────────────────────────────────────────┐ │ llama_model_load_from_file_impl(path, params) │ │ - 选择/构建 model->devices (GPU/RPC/IGPU...) │ └───────────────┬──────────────────────────────────────────────┘ │ v ┌──────────────────────────────────────────────────────────────┐ │ llama_model_load(...) │ │ - llama_model_loader ml(...) 读取 GGUF 元信息、索引 tensors │ │ - model.load_arch/load_hparams/load_vocab │ │ - model.load_tensors(ml) ← 重点 │ └───────────────┬──────────────────────────────────────────────┘ │ v ┌──────────────────────────────────────────────────────────────┐ │ llama_model::load_tensors(ml) │ │ A) 构建候选 buft 列表 │ │ cpu_buft_list = make_cpu_buft_list(...) │ │ gpu_buft_list[dev] = make_gpu_buft_list(dev,...) + CPU fb │ │ │ │ B) 决定每层放哪台设备(层→dev) │ │ splits = tensor_split 或按 dev free-mem 默认计算 │ │ i_gpu_start = (n_layer+1 - n_gpu_layers) │ │ dev_input=CPU, dev_layer[il]=CPU/GPU..., dev_output=... │ │ │ │ C) 决定每个 tensor 用哪个 buft，并创建 ggml_tensor 元信息 │ │ 对每个权重 tensor: │ │ - 根据层(input/repeating/output)选 buft_list │ │ - buft = select_weight_buft(...) (考虑 op/type/override) │ │ - ctx = ctx_map[buft] (同 buft 归到同 ggml_context) │ │ - ml.create_tensor(ctx, ...) (创建 tensor meta) │ │ ml.done_getting_tensors() 校验数量 │ └───────────────┬──────────────────────────────────────────────┘ │ v ┌──────────────────────────────────────────────────────────────┐ │ D) 分配后端 buffer（真正“存储权重”的地方） │ │ ml.init_mappings(...) (mmap 时建立映射、算 size_data) │ │ 对每个 (buft -> ctx): │ │ - 若满足条件: ggml_backend_dev_buffer_from_host_ptr(...) │ │ （把 mmap 的区间包装成后端 buffer，少拷贝） │ │ - 否则: ggml_backend_alloc_ctx_tensors_from_buft(ctx,buft)│ │ - 结果保存到 model.pimpl->ctxs_bufs 维持生命周期 │ │ - 标记 buffer usage = WEIGHTS (帮助调度) │ └───────────────┬──────────────────────────────────────────────┘ │ v ┌──────────────────────────────────────────────────────────────┐ │ E) 装载权重数据到后端（mmap/读文件/异步上传） │ │ 对每个 ctx: ml.load_all_data(ctx, buf_map, progress_cb) │ │ - progress_cb 返回 false => 取消 => load_tensors 返回 false│ │ 若 use_mmap_buffer: model.pimpl->mappings 接管 ml.mappings │ │ 同时填充 model.tensors_by_name 供按名查询/统计 │ └──────────────────────────────────────────────────────────────┘ │ v ┌──────────────────────────────────────────────────────────────┐ │ 推理阶段 build_graph / llama_decode │ │ - 直接引用 model.tok_embd / model.layers[i].wq... 等 ggml_tensor│ │ - 后端根据 tensor 绑定的 buffer 决定在 CPU/GPU 执行与取数 │ └──────────────────────────────────────────────────────────────┘ buft 是什么？ buft 是 ggml_backend_buffer_type_t（backend buffer type），可以理解为“某个后端设备上，用哪一种内存/分配方式来存放 tensor 数据”的类型句柄。 在 load_tensors() 里它的作用是：给每个权重 tensor 选择一个合适的 buffer type，并据此把 tensor 归到对应的 ggml_context，最终用该 buft 去创建真实的后端 buffer（CPU RAM / GPU VRAM / 远端 RPC buffer / host pinned buffer 等）。 load_tensors() 里按 tensor 所在层选一个 buft_list，再用 select_weight_buft(...) 选出最终 buft： if (!buft) { buft = select_weight_buft(hparams, t_meta, op, *buft_list); if (!buft) { throw std::runtime_error(format("failed to find a compatible buffer type for tensor %s", tn.str().c_str())); } } 随后用 buft 决定该 tensor 属于哪个 ggml_context（ctx_for_buft(buft)），并最终分配后端 buffer（例如 ggml_backend_alloc_ctx_tensors_from_buft(ctx, buft) 或 mmap 的 ggml_backend_dev_buffer_from_host_ptr(...)）。 buft 和 dev（device）的区别 - dev（ggml_backend_dev_t）：是哪块设备（CPU / 某张 GPU / RPC 设备…） - buft（ggml_backend_buffer_type_t）：在这块设备上“用哪种 buffer 类型/内存形式”来存 tensor（默认 device buffer、host buffer、特殊优化 buffer 等） （4）推理阶段如何使用这些权重 后续推理构图时（model.build_graph(...) ——>各 src/models/xx.cpp），会直接用 llama_model 里的这些 ggml_tensor x 作为权重输入，和激活值做 ggml 运算（matmul/add/norm/rope…）。例如 src/models/stablelm.cpp 在构图时直接用： model.tok_embd 做输入 embedding - model.layers[il].wq/wk/wv 做 Q/K/V 投影 - model.layers[il].bq/bk/bv（可选）加 bias - ggml_rope_ext 用 rope 参数对 Q/K 做 RoPE Tokenization const llama_vocab * vocab = llama_model_get_vocab(model); // tokenize the prompt // find the number of tokens in the prompt const int n_prompt = -llama_tokenize(vocab, prompt.c_str(), prompt.size(), NULL, 0, true, true); // allocate space for the tokens and tokenize the prompt std::vector<llama_token> prompt_tokens(n_prompt); if (llama_tokenize(vocab, prompt.c_str(), prompt.size(), prompt_tokens.data(), prompt_tokens.size(), true, true) < 0) { fprintf(stderr, "%s: error: failed to tokenize the prompt\n", __func__); return 1; } 这段代码主要的作用是做文本->token ID的转换。整段的流程可以总结一下。 拿vocab：llama_model_get_vocab(model) 获取 tokenizer/词表句柄； 预估token数：第一次 llama_tokenize(..., NULL, 0, ...)，只算需要多少 token，返回负数取反得到 n_prompt； 分配缓存：std::vector prompt_tokens(n_prompt); 为 prompt 分配精确长度的 token 缓冲区； 真正 tokenize：第二次 llama_tokenize 把 token id 写进 prompt_tokens，如果失败或容量不够返回 < 0 则报错退出。 llama_tokenize() tokenize 的对象是 一段输入文本（UTF-8 字符串），输出是 token id 序列（llama_token）：也就是模型能直接消费的离散符号编号。它依赖 llama_vocab（模型绑定的 tokenizer + 词表 + special token 规则），把文本切分/编码成 token ids。 llama_tokenize会转发到vocab->tokenize，先做 special token 分段，再按 vocab type 分发到具体 tokenizer session。 std::vector<llama_token> llama_vocab::impl::tokenize(const std::string & raw_text, bool add_special, bool parse_special) const { // 1) special token partition（把 raw_text 切成「普通文本片段」和「已识别的特殊 token」） tokenizer_st_partition(fragment_buffer, parse_special); // 2) 按 vocab type 分发到具体 tokenizer session switch (get_type()) { case LLAMA_VOCAB_TYPE_SPM: llm_tokenizer_spm_session(...).tokenize(...); break; case LLAMA_VOCAB_TYPE_BPE: llm_tokenizer_bpe_session(...).tokenize(...); break; case LLAMA_VOCAB_TYPE_WPM: llm_tokenizer_wpm_session(...).tokenize(...); break; case LLAMA_VOCAB_TYPE_UGM: llm_tokenizer_ugm_session(...).tokenize(...); break; case LLAMA_VOCAB_TYPE_RWKV: llm_tokenizer_rwkv_session(...).tokenize(...); break; case LLAMA_VOCAB_TYPE_PLAMO2: llm_tokenizer_plamo2_session(...).tokenize(...); break; default: GGML_ABORT(...); } return output; } SPM SentencePiece BPE-ish 合并，项目里叫 SPM tokenizer，入口类是llm_tokenizer_spm_session，其核心思路是把文本先按 UTF-8 字符拆成链表 symbol，然后不断从优先队列取“得分最高”的可合并 bigram，做合并，最后把每个合并后的片段映射为 token。 struct llm_tokenizer_spm_session { void tokenize(const std::string & text, std::vector<llama_token> & output) { // split into utf8 chars -> symbols // seed work_queue with all 2-char tokens -> try_add_bigram // pop highest score, merge, and push new bigrams // resegment each final symbol into token ids / byte tokens } } 在 impl::tokenize 的 SPM 分支里，会先把空格转义成 “▁”（U+2581），并处理 add_space_prefix、BOS/EOS： if (add_space_prefix && is_prev_special) { text = ' '; } text += fragment... llama_escape_whitespace(text); // ' ' -> ▁ llm_tokenizer_spm_session session(vocab); session.tokenize(text, output); BEP Byte Pair Encoding，入口类：llm_tokenizer_bpe_session，其核心思路是 1) 用“pre-tokenizer regex”（不同模型 pre_type 不同）把文本切成词片段（word_collection）。 2) 每个词片段再按 UTF-8 拆成 symbols。 3) 用 vocab.find_bpe_rank(left,right) 查 merges rank（越小越优先），不断合并。 4) 合并结束后把每个最终 symbol 映射为 token（找不到就回退到逐字节 token）。 WPM WordPiece，入口类：llm_tokenizer_wpm_session，核心思路是： 先做 NFD 归一化 + 按 whitespace/标点切词（preprocess） 每个词前加 “▁” 对每个词做 最长匹配（从当前位置往后尽量取最长、且在 vocab 中存在的 token） 若某个词无法完全分解，回退整个词为unk UGM 入口类：llm_tokenizer_ugm_session，核心思路：典型 unigram LM 的 Viterbi 最优切分： 先 normalize 输入 每个位置沿 trie 找所有可能 token，做动态规划选择 “累计得分最大”的路径 走不通就用 unknown token + penalty 最后从末尾回溯得到 token 序列 初始化上下文 llama_context_params ctx_params = llama_context_default_params(); // n_ctx is the context size ctx_params.n_ctx = n_prompt + n_predict - 1; // n_batch is the maximum number of tokens that can be processed in a single call to llama_decode ctx_params.n_batch = n_prompt; // enable performance counters ctx_params.no_perf = false; llama_context * ctx = llama_init_from_model(model, ctx_params); if (ctx == NULL) { fprintf(stderr , "%s: error: failed to create the llama_context\n" , __func__); return 1; } 这段代码主要的作用是创建llama_context，用于管理推理的状态KV cache、batch 分配器、backend 调度器等）。 配置上下文参数、batch大小、性能统计。创建llama_context包括初始化计算后端，分配KV cache内存，创建batch分配器和调度器，预留计算图内存。llama_context是推理的核心对象，后续的llama_decode调用都会使用它来管理状态和执行计算。 初始化采样器 auto sparams = llama_sampler_chain_default_params(); sparams.no_perf = false; llama_sampler * smpl = llama_sampler_chain_init(sparams); llama_sampler_chain_add(smpl, llama_sampler_init_greedy()); 构造一个“采样器链”（sampler chain），并在其中添加一个贪心采样器（greedy sampler），后面主循环用它从 logits 里选下一个 token。下面按调用顺序拆开。采样器接收一个 llama_token_data_array（包含所有候选 token 的 id、logit、概率），经过处理（过滤、排序、选择等），最终确定一个 token。 处理提示词 // print the prompt token-by-token for (auto id : prompt_tokens) { char buf[128]; int n = llama_token_to_piece(vocab, id, buf, sizeof(buf), 0, true); if (n < 0) { fprintf(stderr, "%s: error: failed to convert token to piece\n", __func__); return 1; } std::string s(buf, n); printf("%s", s.c_str()); } // prepare a batch for the prompt llama_batch batch = llama_batch_get_one(prompt_tokens.data(), prompt_tokens.size()); if (llama_model_has_encoder(model)) { if (llama_encode(ctx, batch)) { fprintf(stderr, "%s : failed to eval\n", __func__); return 1; } llama_token decoder_start_token_id = llama_model_decoder_start_token(model); if (decoder_start_token_id == LLAMA_TOKEN_NULL) { decoder_start_token_id = llama_vocab_bos(vocab); } batch = llama_batch_get_one(&decoder_start_token_id, 1); } llama_token_to_piece是将token ID转换为文本并输出，llama_batch_get_one是为prompt创建batch结构。通过llama_model_has_encoder检查模型是不是编码-解码模型，如果是先进行prompt，在准备decoder的起始batch。 总结一下就是先调用llama_encode编码prompt的token，做完这一步，encoder部分就结束了，接着让decoder开始工作，我们知道大模型是一个自回归模型，需要有第一个token作为起点，类似机器翻译里的 BOS。 最后的llama_batch_get_one这里就是准备decoder起始batch。 llama_encode llama_decode 生成循环 for (int n_pos = 0; n_pos + batch.n_tokens < n_prompt + n_predict; ) { // evaluate the current batch with the transformer model if (llama_decode(ctx, batch)) { fprintf(stderr, "%s : failed to eval, return code %d\n", __func__, 1); return 1; } n_pos += batch.n_tokens; // sample the next token { new_token_id = llama_sampler_sample(smpl, ctx, -1); // is it an end of generation? if (llama_vocab_is_eog(vocab, new_token_id)) { break; } char buf[128]; int n = llama_token_to_piece(vocab, new_token_id, buf, sizeof(buf), 0, true); if (n < 0) { fprintf(stderr, "%s: error: failed to convert token to piece\n", __func__); return 1; } std::string s(buf, n); printf("%s", s.c_str()); fflush(stdout); // prepare the next batch with the sampled token batch = llama_batch_get_one(&new_token_id, 1); n_decode += 1; } } 这段代码是自回归文本生成的核心循环，实现decoder->sample->print的循环。llama_decode进行decoder，然后调用llamap_sampler_sample采样下一个token，获取对应的token ID，通过调用llama_vocab_is_eog来检查结束条件，判断采样到的token是否为结束token，是则退出循环。最后调用llama_token_to_piece将token ID转换为文本并打印。 核心的思想其实跟之前写的：从零实现 Transformer：中英文翻译实例文章是一样的。

准备 硬件信息 硬件信息如下： sudo cat /proc/device-tree/model NVIDIA Jetson Orin NX Enginejetson_releasee Developer Kit(base) nano@nano-desktop:~$ jetson_release Software part of jetson-stats 4.3.2 - (c) 2024, Raffaello Bonghi Model: NVIDIA Jetson Orin NX Engineering Reference Developer Kit - Jetpack 6.2 [L4T 36.4.3] NV Power Mode[0]: 15W Serial Number: [XXX Show with: jetson_release -s XXX] Hardware: - P-Number: p3767-0003 - Module: NVIDIA Jetson Orin Nano (8GB ram) Platform: - Distribution: Ubuntu 22.04 Jammy Jellyfish - Release: 5.15.148-tegra jtop: - Version: 4.3.2 - Service: Active 超级模式 设置电源为超级性能模式，点击 Ubuntu 桌面顶部栏右侧的 NVIDIA 图标 Power mode 0： 15W 1： 25W 2： MAXN SUPER 如果没有MAXN SUPER，执行下面的命令安装。 sudo apt remove nvidia-l4t-bootloader sudo apt install nvidia-l4t-bootloader sudo reboot 如果还是无法配置的话，参考下：super mode 安装docker 本文主要使用https://github.com/dusty-nv/jetson-containers docker来进行部署测试，免去不少的环境搭建过程。 如果设备没有安装docker，执行下面命令进行安装docker。 sudo apt-get update sudo apt-get install -y docker.io nvidia-container-toolkit sudo usermod -aG docker $USER sudo systemctl daemon-reload sudo systemctl restart docker 接着容器启动拉取代码，然后拉取docker镜像： git clone https://github.com/dusty-nv/jetson-containers.git ./jetson-containers/run.sh -v $(pwd)/model_zoo:/opt/tensorrt_llm/models $(autotag tensorrt_llm) 下载模型 先下载一个模型 # Qwen 1.5B modelscope download --model Qwen/Qwen2.5-1.5B-Instruct --local_dir model_zoo/Qwen2.5-1.5B-Instruct # Qwen 0.5B modelscope download --model Qwen/Qwen2.5-0.5B-Instruct --local_dir model_zoo/Qwen2.5-0.5B-Instruct #DeepSeek 1.5B modelscope download --model deepseek-ai/DeepSeek-R1-Distill-Qwen-1.5B --local_dir model_zoo/DeepSeek-R1-Distill-Qwen-1.5B qwen2.5-0.5B INT4 启动容器 ./jetson-containers/run.sh -v $(pwd)/model_zoo:/opt/tensorrt_llm/models $(autotag tensorrt_llm) 转换模型 使用TensorRT-LLM提供的工具进行量化 cd /opt/TensorRT-LLM/examples/qwen python3 convert_checkpoint.py \ --model_dir /opt/tensorrt_llm/models/Qwen2.5-0.5B-Instruct/ \ --output_dir /opt/tensorrt_llm/models/qwen0.5b_ckpt_int4/ \ --dtype float16 \ --use_weight_only \ --weight_only_precision int4 参数,作用说明 model_dir：源模型路径，指向原始 Qwen2.5-1.5B-Instruct的文件夹。 output_dir：输出路径，转换后的 TensorRT-LLM 格式权重将保存在此目录下。 dtype float16：计算精度，设置模型在推理时的基础数据类型为 float16（半精度）。 use_weight_only：启用仅权重处理，告诉脚本不要修改激活值（Activations），只处理权重部分。 weight_only_precision int4：量化等级，将权重从 FP16 压缩到 INT4。这意味着权重占用的空间将缩小为原来的约 1/4。 weight_only_precision和dtype的区别是，前者模型权重以 INT4 格式存放在显存里（非常省空间）。后者了是实际运算用的精度，动态解压 (Dequantization)，当 GPU 准备计算某一层的矩阵乘法时，它会实时地将这一层的 INT4 权重“恢复”成 float16。 构建引擎 构建引擎 trtllm-build --checkpoint_dir /opt/tensorrt_llm/models/qwen0.5b_ckpt_int4/ \ --output_dir /opt/tensorrt_llm/models/qwen0.5b_engine_int4_bs4/ \ --gemm_plugin float16 \ --gpt_attention_plugin float16 \ --max_batch_size 4 \ --max_input_len 1024 \ --max_seq_len 1024 \ --workers 1 测试性能 压力测试 cd /opt/TensorRT-LLM/benchmarks/python # 运行 Batch Size 4 的测试 python3 benchmark.py \ --engine_dir /opt/tensorrt_llm/models/qwen0.5b_engine_int4_bs4/ \ --batch_size 4 \ --input_output_len "512,128" \ --dtype float16 \ --num_runs 10 \ --warm_up 2 engine_dir：指定 Engine 路径,指向已编译好的 TensorRT-LLM Engine。从路径名 qwen0.5b_engine_int4_bs4 可以推断，这个模型是 Qwen-0.5B，采用了 INT4 量化，且构建时设定的最大 Batch Size 为 4。 batch_size： 4,推理并发数,本次测试同时处理 4 个请求。这会直接测试芯片在满载（Full Load）状态下的性能。由于你的 Engine 编译时 max_batch_size 可能是 4，所以这里设置 4 是为了榨干该 Engine 的性能。 input_output_len： "512,128"",输入输出长度,512：输入 Prompt 的长度（影响 PP 阶段）。128：生成 Token 的长度（影响 TG 阶段）。这是一个非常标准的测试组合，模拟了“中等长度提问 + 短回答”的场景。 dtype： float16,计算精度,即使权重是 INT4 存储，实际矩阵运算（Activation）依然使用 FP16。这能保证在利用 INT4 节省显存的同时，推理精度不崩溃。 num_runs： 10,循环次数,连续运行 10 次。取平均值可以消除系统波动（如 CPU 调度、瞬时过热降频）带来的误差，使测试结果更具统计学意义。 warm_up： 2,预热次数,忽略前 2 次的成绩。GPU 在刚开始跑的时候需要进行内存申请、显存对齐等操作，通常第一遍会非常慢。预热可以确保获取的是稳定状态（Steady State)下的真实性能。 qwen2.5-1.5B INT4 启动容器 ./jetson-containers/run.sh -v $(pwd)/model_zoo:/opt/tensorrt_llm/models $(autotag tensorrt_llm) 转换模型 使用TensorRT-LLM提供的工具进行量化，如果内存不够就到PC上去做量化，见下。 cd /opt/TensorRT-LLM/examples/qwen python3 convert_checkpoint.py \ --model_dir /opt/tensorrt_llm/models/Qwen2.5-1.5B-Instruct/ \ --output_dir /opt/tensorrt_llm/models/qwen1.5b_ckpt_int4/ \ --dtype float16 \ --use_weight_only \ --weight_only_precision int4 构建引擎 构建引擎 trtllm-build \ --checkpoint_dir /opt/tensorrt_llm/models/qwen3b_ckpt_int4 \ --output_dir /opt/tensorrt_llm/models/qwen3b_int4_bs1_engine \ --gemm_plugin float16 \ --gpt_attention_plugin float16 \ --max_batch_size 1 \ --max_input_len 1024 \ --max_seq_len 1024 \ --workers 1 如果内存不够的话就把max_batch_size改小一点。 测试对话 python3 ../run.py \ --engine_dir /opt/tensorrt_llm/models/qwen1.5b_int4_bs1_engine/ \ --tokenizer_dir /opt/tensorrt_llm/models/Qwen2.5-3B-Instruct/ \ --max_output_len 128 \ --input_text "你好，请介绍一下你自己，并讲一个关于芯片工程师的冷笑话。" 测试性能 压力测试 cd /opt/TensorRT-LLM/benchmarks/python # 运行 Batch Size 1 的测试 python3 benchmark.py \ --engine_dir /opt/tensorrt_llm/models/qwen1.5b_engine_int4_bs4/ \ --batch_size 1 \ --input_output_len "128,128" \ --dtype float16 \ --num_runs 10 \ --warm_up 2 qwen2.5-3B INT4 启动容器 ./jetson-containers/run.sh -v $(pwd)/model_zoo:/opt/tensorrt_llm/models $(autotag tensorrt_llm) 转换模型 使用TensorRT-LLM提供的工具进行量化，如果内存不够就到PC上去做量化，见下。 python3 convert_checkpoint.py \ --model_dir /data/Qwen2.5-3B-Instruct/ \ --output_dir /data/qwen3b_ckpt_int4 \ --dtype float16 \ --use_weight_only \ --weight_only_precision int4 构建引擎 构建引擎 base) nano@nano-desktop:~$ docker run --runtime nvidia -it --rm \ --network host \ --shm-size=8g \ --volume /home/nano/model_zoo:/opt/tensorrt_llm/models \ --workdir /opt/tensorrt_llm/examples/qwen \ dustynv/tensorrt_llm:0.12-r36.4.0 \ /bin/bash -c "export TRT_MAX_WORKSPACE_SIZE=1073741824 && \ trtllm-build \ --checkpoint_dir /opt/tensorrt_llm/models/qwen3b_ckpt_int4 \ --output_dir /opt/tensorrt_llm/models/qwen3b_int4_bs1_engine \ --gemm_plugin float16 \ --gpt_attention_plugin float16 \ --max_batch_size 1 \ --max_seq_len 1024 \ --workers 1 \ --builder_opt 1" /usr/local/lib/python3.10/dist-packages/transformers/utils/hub.py:128: FutureWarning: Using `TRANSFORMERS_CACHE` is deprecated and will be removed in v5 of Transformers. Use `HF_HOME` instead. warnings.warn( [TensorRT-LLM] TensorRT-LLM version: 0.12.0 [01/24/2026-04:10:02] [TRT-LLM] [I] Set bert_attention_plugin to auto. [01/24/2026-04:10:02] [TRT-LLM] [I] Set gpt_attention_plugin to float16. [01/24/2026-04:10:02] [TRT-LLM] [I] Set gemm_plugin to float16. [01/24/2026-04:10:02] [TRT-LLM] [I] Set gemm_swiglu_plugin to None. [01/24/2026-04:10:02] [TRT-LLM] [I] Set fp8_rowwise_gemm_plugin to None. [01/24/2026-04:10:02] [TRT-LLM] [I] Set nccl_plugin to auto. [01/24/2026-04:10:02] [TRT-LLM] [I] Set lookup_plugin to None. [01/24/2026-04:10:02] [TRT-LLM] [I] Set lora_plugin to None. [01/24/2026-04:10:02] [TRT-LLM] [I] Set moe_plugin to auto. [01/24/2026-04:10:02] [TRT-LLM] [I] Set mamba_conv1d_plugin to auto. [01/24/2026-04:10:02] [TRT-LLM] [I] Set context_fmha to True. [01/24/2026-04:10:02] [TRT-LLM] [I] Set bert_context_fmha_fp32_acc to False. [01/24/2026-04:10:02] [TRT-LLM] [I] Set paged_kv_cache to True. [01/24/2026-04:10:02] [TRT-LLM] [I] Set remove_input_padding to True. [01/24/2026-04:10:02] [TRT-LLM] [I] Set reduce_fusion to False. [01/24/2026-04:10:02] [TRT-LLM] [I] Set enable_xqa to True. [01/24/2026-04:10:02] [TRT-LLM] [I] Set tokens_per_block to 64. [01/24/2026-04:10:02] [TRT-LLM] [I] Set use_paged_context_fmha to False. [01/24/2026-04:10:02] [TRT-LLM] [I] Set use_fp8_context_fmha to False. [01/24/2026-04:10:02] [TRT-LLM] [I] Set multiple_profiles to False. [01/24/2026-04:10:02] [TRT-LLM] [I] Set paged_state to True. [01/24/2026-04:10:02] [TRT-LLM] [I] Set streamingllm to False. [01/24/2026-04:10:02] [TRT-LLM] [W] Implicitly setting QWenConfig.qwen_type = qwen2 [01/24/2026-04:10:02] [TRT-LLM] [W] Implicitly setting QWenConfig.moe_intermediate_size = 0 [01/24/2026-04:10:02] [TRT-LLM] [W] Implicitly setting QWenConfig.moe_shared_expert_intermediate_size = 0 [01/24/2026-04:10:02] [TRT-LLM] [I] Set dtype to float16. [01/24/2026-04:10:02] [TRT-LLM] [W] remove_input_padding is enabled, while opt_num_tokens is not set, setting to max_batch_size*max_beam_width. [01/24/2026-04:10:02] [TRT-LLM] [W] max_num_tokens (1024) shouldn't be greater than max_seq_len * max_batch_size (1024), specifying to max_seq_len * max_batch_size (1024). [01/24/2026-04:10:02] [TRT-LLM] [W] padding removal and fMHA are both enabled, max_input_len is not required and will be ignored [01/24/2026-04:10:02] [TRT] [I] [MemUsageChange] Init CUDA: CPU +12, GPU +0, now: CPU 164, GPU 1447 (MiB) [01/24/2026-04:10:05] [TRT] [I] [MemUsageChange] Init builder kernel library: CPU +947, GPU +1133, now: CPU 1154, GPU 2594 (MiB) [01/24/2026-04:10:05] [TRT] [W] profileSharing0806 is on by default in TensorRT 10.0. This flag is deprecated and has no effect. [01/24/2026-04:10:05] [TRT-LLM] [I] Set weight_only_quant_matmul_plugin to float16. [01/24/2026-04:10:05] [TRT-LLM] [I] Set nccl_plugin to None. [01/24/2026-04:10:06] [TRT-LLM] [I] Total optimization profiles added: 1 [01/24/2026-04:10:06] [TRT-LLM] [I] Build TensorRT engine Unnamed Network 0 [01/24/2026-04:10:06] [TRT] [W] DLA requests all profiles have same min, max, and opt value. All dla layers are falling back to GPU [01/24/2026-04:10:06] [TRT] [W] Unused Input: position_ids [01/24/2026-04:10:06] [TRT] [W] [RemoveDeadLayers] Input Tensor position_ids is unused or used only at compile-time, but is not being removed. [01/24/2026-04:10:06] [TRT] [I] Global timing cache in use. Profiling results in this builder pass will be stored. [01/24/2026-04:10:06] [TRT] [I] Compiler backend is used during engine build. [01/24/2026-04:10:11] [TRT] [I] [GraphReduction] The approximate region cut reduction algorithm is called. [01/24/2026-04:10:11] [TRT] [I] Detected 15 inputs and 1 output network tensors. NvMapMemAllocInternalTagged: 1075072515 error 12 NvMapMemHandleAlloc: error 0 NvMapMemAllocInternalTagged: 1075072515 error 12 NvMapMemHandleAlloc: error 0 [01/24/2026-04:10:20] [TRT] [E] [resizingAllocator.cpp::allocate::74] Error Code 1: Cuda Runtime (out of memory) [01/24/2026-04:10:20] [TRT] [W] Requested amount of GPU memory (1514557910 bytes) could not be allocated. There may not be enough free memory for allocation to succeed. [01/24/2026-04:10:20] [TRT] [E] [globWriter.cpp::makeResizableGpuMemory::433] Error Code 2: OutOfMemory (Requested size was 1514557910 bytes.) Traceback (most recent call last): File "/usr/local/bin/trtllm-build", line 8, in <module> sys.exit(main()) File "/usr/local/lib/python3.10/dist-packages/tensorrt_llm/commands/build.py", line 500, in main parallel_build(model_config, ckpt_dir, build_config, args.output_dir, File "/usr/local/lib/python3.10/dist-packages/tensorrt_llm/commands/build.py", line 377, in parallel_build passed = build_and_save(rank, rank % workers, ckpt_dir, File "/usr/local/lib/python3.10/dist-packages/tensorrt_llm/commands/build.py", line 344, in build_and_save engine = build_model(build_config, File "/usr/local/lib/python3.10/dist-packages/tensorrt_llm/commands/build.py", line 337, in build_model return build(model, build_config) File "/usr/local/lib/python3.10/dist-packages/tensorrt_llm/builder.py", line 1060, in build engine = None if build_config.dry_run else builder.build_engine( File "/usr/local/lib/python3.10/dist-packages/tensorrt_llm/_common.py", line 204, in decorated return f(*args, **kwargs) File "/usr/local/lib/python3.10/dist-packages/tensorrt_llm/builder.py", line 411, in build_engine assert engine is not None, 'Engine building failed, please check the error log.' AssertionError: Engine building failed, please check the error log. 内存不够，测不了了...... DeepSeek R1-1.5B 启动容器 ./jetson-containers/run.sh -v $(pwd)/model_zoo:/opt/tensorrt_llm/models $(autotag tensorrt_llm) 转换模型 使用TensorRT-LLM提供的工具进行量化，如果内存不够就到PC上去做量化，见下。 cd /opt/TensorRT-LLM/examples/qwen python3 convert_checkpoint.py \ --model_dir /data/DeepSeek-R1-Distill-Qwen-1.5B/ \ --output_dir /data/deepseek_r1_1.5b_ckpt_int4 \ --dtype float16 \ --use_weight_only \ --weight_only_precision int4 构建引擎 构建引擎 trtllm-build \ --checkpoint_dir /opt/tensorrt_llm/models/deepseek_r1_1.5b_ckpt_int4 \ --output_dir /opt/tensorrt_llm/models/deepseek_r1_1.5b_int4_bs1_engine \ --gemm_plugin float16 \ --gpt_attention_plugin float16 \ --max_batch_size 1 \ --max_input_len 1024 \ --max_seq_len 1024 \ --workers 1 如果内存不够的话就把max_batch_size改小一点。 测试对话 cd /opt/TensorRT-LLM/examples/qwen python3 ../run.py \ --engine_dir /opt/tensorrt_llm/models/deepseek_r1_1.5b_int4_bs1_engine/ \ --tokenizer_dir /opt/tensorrt_llm/models/DeepSeek-R1-Distill-Qwen-1.5B/ \ --max_output_len 128 \ --input_text "你好，请介绍一下你自己，并讲一个关于芯片工程师的冷笑话。" 测试性能 压力测试 cd /opt/TensorRT-LLM/benchmarks/python # 运行 Batch Size 1 的测试 python3 benchmark.py \ --engine_dir /opt/tensorrt_llm/models/deepseek_r1_1.5b_int4_bs1_engine/ \ --batch_size 1 \ --input_output_len "512,128" \ --dtype float16 \ --num_runs 10 \ --warm_up 2 在PC上做量化 由于在jetson nano上做量化，内存不足，可以考虑在PC上做量化，然后把模型拷贝到设备中去。由于容器内使用的python是3.12，所以我们安装一个PyTorch 24.02的容器。 安装量化环境 （1）启动容器 # 注意：镜像换成了 pytorch:24.02-py3，它是 Python 3.10 环境 docker run --rm -it --ipc=host --gpus all \ -v ~/model_zoo:/data \ -v ~/TensorRT-LLM:/code \ -w /code/examples/qwen \ nvcr.io/nvidia/pytorch:24.02-py3 bash （2）安装TensorRT-LLM 0.12.0 # 这一步会下载约 1GB 的包，请耐心等待 pip install tensorrt_llm==0.12.0 --extra-index-url https://pypi.nvidia.com （3）安装脚本依赖 pip install "transformers==4.38.2" safetensors accelerate （4）验证安装 python3 -c "import tensorrt_llm; print(tensorrt_llm.__version__)" # 输出应该是 0.12.0 过程中如果遇到问题，可以按照下面的办法处理 # 强制重装 TensorRT 相关的核心包 pip install --force-reinstall \ tensorrt==10.3.0 \ tensorrt-cu12==10.3.0 \ tensorrt-cu12-bindings==10.3.0 \ tensorrt-cu12-libs==10.3.0 export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/lib/python3.10/dist-packages/tensorrt_libs #卸载捣乱的 Flash Attention，不需要这个 pip uninstall -y flash-attn # 1. 卸载可能存在的冲突包 pip uninstall -y pynvml nvidia-ml-py # 2. 安装旧版 pynvml (11.4.1 是最稳定的旧版，完美兼容 0.12.0) pip install pynvml==11.4.1 转换模型 python3 convert_checkpoint.py \ --model_dir /data/Qwen2.5-1.5B-Instruct/ \ --output_dir /data/qwen1.5b_ckpt_int4 \ --dtype float16 \ --use_weight_only \ --weight_only_precision int4 然后可以拷贝到设备上。 # 确保 Jetson 上有这个目录 ssh nano@10.0.91.125 "mkdir -p ~/model_zoo" # 开始传输 (PC -> Jetson) rsync -avP ~/model_zoo/qwen1.5b_ckpt_int4 nano@10.0.91.125:~/model_zoo/ 常见问题 找不到库 如果在模型转换的时候出现下面报错，先退出docker容器，安装下面的库 ImportError: libnvdla_compiler.so: cannot open shared object file: No such file or directory wget -O - https://repo.download.nvidia.com/jetson/common/pool/main/n/nvidia-l4t-dla-compiler/nvidia-l4t-dla-compiler_36.4.1-20241119120551_arm64.deb | dpkg-deb --fsys-tarfile - | sudo tar xv --strip-components=5 --directory=/usr/lib/aarch64-linux-gnu/nvidia/ ./usr/lib/aarch64-linux-gnu/nvidia/libnvdla_compiler.so 内存不足 禁用桌面图像界面，暂时禁用 $ sudo init 3 # stop the desktop # log your user back into the console (Ctrl+Alt+F1, F2, ect) $ sudo init 5 # restart the desktop 如果想重启也禁用 #禁用 sudo systemctl set-default multi-user.target #启用 sudo systemctl set-default graphical.target # 强行同步并释放缓存 sudo sh -c 'sync; echo 3 > /proc/sys/vm/drop_caches' 参考资料： jetson SDK jetson-ai-lab jetson官方数据

安装 确认一下是否有/usr/local/cuda/bin/nvcc，有就配置一下环境。 # 临时在当前终端生效 export PATH=/usr/local/cuda/bin:$PATH export LD_LIBRARY_PATH=/usr/local/cuda/lib64:$LD_LIBRARY_PATH # 验证 nvcc 是否能找到 nvcc --version 如果没有就安装 sudo -E apt install cuda-toolkit 下载代码 git clone https://github.com/ggerganov/llama.cpp cd llama.cpp 安装魔塔下载工具 pip install modelscope -i https://mirrors.aliyun.com/pypi/simple/ 编译 mkdir build cmake .. -DGGML_CUDA=ON cmake --build . --config Release -j $(nproc) 跑模型 低功耗模式 下载模型 mkdir model_zoo modelscope download --model Qwen/Qwen2.5-1.5B-Instruct-GGUF qwen2.5-1.5b-instruct-q4_0.gguf --local_dir ./model_zoo 测试模型 ./build/bin/llama-cli -m ../model_zoo/qwen2.5-1.5b-instruct-q4_0.gguf --perf --show-timings -lv 3 换个笑话 |slot get_availabl: id 0 | task -1 | selected slot by LCP similarity, sim_best = 0.952 (> 0.100 thold), f_keep = 1.000 slot launch_slot_: id 0 | task -1 | sampler chain: logits -> ?penalties -> ?dry -> ?top-n-sigma -> top-k -> ?typical -> top-p -> min-p -> ?xtc -> temp-ext -> dist slot launch_slot_: id 0 | task 184 | processing task, is_child = 0 slot update_slots: id 0 | task 184 | new prompt, n_ctx_slot = 4096, n_keep = 0, task.n_tokens = 248 slot update_slots: id 0 | task 184 | n_tokens = 236, memory_seq_rm [236, end) slot update_slots: id 0 | task 184 | prompt processing progress, n_tokens = 248, batch.n_tokens = 12, progress = 1.000000 slot update_slots: id 0 | task 184 | prompt done, n_tokens = 248, batch.n_tokens = 12 slot init_sampler: id 0 | task 184 | init sampler, took 0.11 ms, tokens: text = 248, total = 248 好的，以下是一个不同的笑话： 为什么电脑没有呼吸？ 因为它只有一个输入口，一个输出口。slot print_timing: id 0 | task 184 | prompt eval time = 214.26 ms / 12 tokens ( 17.85 ms per token, 56.01 tokens per second) eval time = 1691.48 ms / 22 tokens ( 76.89 ms per token, 13.01 tokens per second) total time = 1905.74 ms / 34 tokens slot release: id 0 | task 184 | stop processing: n_tokens = 269, truncated = 0 srv update_slots: all slots are idle [ Prompt: 56.0 t/s | Generation: 13.0 t/s ] 看起来速度不是很快。 再跑一个3B的试试看。 ./build/bin/llama-cli -m ../model_zoo/qwen2.5-3b-instruct-q4_0.gguf --perf --show-timings -lv 3 prompt eval time = 1412.32 ms / 15 tokens ( 94.15 ms per token, 10.62 tokens per second) eval time = 48064.10 ms / 244 tokens ( 196.98 ms per token, 5.08 tokens per second) total time = 49476.41 ms / 259 tokens slot release: id 0 | task 148 | stop processing: n_tokens = 469, truncated = 0 srv update_slots: all slots are idle [ Prompt: 10.6 t/s | Generation: 5.1 t/s ] 看起来不太行。 加上个参数重新编译试试，针对CUDA的能力计算数值，加这个参数，llama.cpp 的 cmake 默认可能不会启用针对特定 GPU 架构的深度优化 rm -rf build && mkdir build && cd build cmake .. -DGGML_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=87 #或者cmake .. -DGGML_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=native cmake --build . --config Release -j$(nproc) 再测试执行： ./build/bin/llama-cli -m ../model_zoo/qwen2.5-3b-instruct-q4_0.gguf --perf --show-timings -lv 3 -ngl 99 prompt eval time = 201.92 ms / 13 tokens ( 15.53 ms per token, 64.38 tokens per second) eval time = 39903.36 ms / 386 tokens ( 103.38 ms per token, 9.67 tokens per second) total time = 40105.28 ms / 399 tokens slot release: id 0 | task 159 | stop processing: n_tokens = 598, truncated = 0 srv update_slots: all slots are idle [ Prompt: 64.4 t/s | Generation: 9.7 t/s ] 推理速度上来了，稳定在10 token/s。 再跑跑bench ./build/bin/llama-bench -m ../model_zoo/qwen2.5-3b-instruct-q4_0.gguf -ngl 99 ggml_cuda_init: found 1 CUDA devices: Device 0: Orin, compute capability 8.7, VMM: yes | model | size | params | backend | ngl | test | t/s | | ------------------------------ | ---------: | ---------: | ---------- | --: | --------------: | -------------------: | | qwen2 3B Q4_0 | 1.86 GiB | 3.40 B | CUDA | 99 | pp512 | 259.35 ± 0.46 | | qwen2 3B Q4_0 | 1.86 GiB | 3.40 B | CUDA | 99 | tg128 | 8.98 ± 0.03 | build: 37c35f0e1 (7787) (base) bianbu@ubuntu:~/llama.cpp$ sudo nvpmodel -m 0 # 开启最大功率模式 sudo jetson_clocks # 锁定 CPU/GPU/内存频率到最高 (base) bianbu@ubuntu:~/llama.cpp$ ./build/bin/llama-bench -m ../model_zoo/qwen2.5-3b-instruct-q4_0.gguf -ngl 99 ggml_cuda_init: found 1 CUDA devices: Device 0: Orin, compute capability 8.7, VMM: yes | model | size | params | backend | ngl | test | t/s | | ------------------------------ | ---------: | ---------: | ---------- | --: | --------------: | -------------------: | | qwen2 3B Q4_0 | 1.86 GiB | 3.40 B | CUDA | 99 | pp512 | 264.42 ± 0.27 | | qwen2 3B Q4_0 | 1.86 GiB | 3.40 B | CUDA | 99 | tg128 | 9.17 ± 0.01 | build: 37c35f0e1 (7787) 超级性能模式 上面跑了一下看起来还是很慢，看了一下jetson官网，要跑性能，需要设置电源为超级性能模式，点击 Ubuntu 桌面顶部栏右侧的 NVIDIA 图标 Power mode 0： 15W 1： 25W 2： MAXN SUPER 如果没有MAXN SUPER，执行下面的命令安装。 sudo apt remove nvidia-l4t-bootloader sudo apt install nvidia-l4t-bootloader sudo reboot 如果还是无法配置的话，参考下：super mode 接下来，接着跑。 ./build/bin/llama-cli -m ../model_zoo_gguf/model_zoo/qwen2.5-1.5b-instruct-q4_0.gguf --perf --show-timings -lv 3 -ngl 99 ./build/bin/llama-bench -m ../model_zoo_gguf/model_zoo/qwen2.5-1.5b-instruct-q4_0.gguf -ngl 99 在跑个3B的 ./build/bin/llama-bench -m ../model_zoo_gguf/model_zoo/qwen2.5-3b-instruct-q4_0.gguf -ngl 99 ./build/bin/llama-cli -m ../model_zoo_gguf/model_zoo/qwen2.5-3b-instruct-q4_0.gguf -ngl 99 --perf -lv 3

算子实现 调用流程 主要是ggml_compute_forward_mul_mat函数，该函数把任务拆分，最终计算执行调用ggml_compute_forward_mul_mat_one_chunk实现。 任务拆分 void ggml_compute_forward_mul_mat( const struct ggml_compute_params * params, struct ggml_tensor * dst) { // 1. 获取输入和输出张量的指针 const struct ggml_tensor * src0 = dst->src[0]; const struct ggml_tensor * src1 = dst->src[1]; // 2. 初始化一堆本地变量 (ne0, nb0 等)，方便后面写公式 GGML_TENSOR_BINARY_OP_LOCALS // 3. 获取当前线程的信息 const int ith = params->ith; // 我是第几个线程？ const int nth = params->nth; // 总共有几个线程？ // 4. 获取src0（权重）的数据类型，用于后续查询src1（输入/激活值）必须转换为什么类型才能进行计算。 //在大语言模型中，为了省内存，src0(权重矩阵)通常是经过量化的比如Q4_0，Q5_1，Q8_0等格式。 enum ggml_type const vec_dot_type = type_traits_cpu[src0->type].vec_dot_type; // ... (省略一堆 GGML_ASSERT 检查，确保维度合法，比如 src0 和 src1 能不能相乘) ... #if GGML_USE_LLAMAFILE // 5. 【极速通道】如果系统支持 llamafile 的手写汇编优化，直接丢给它算，算完直接返回。 // 这比下面的 C 代码快得多。 if (src1_cont) { // ... 尝试调用 llamafile_sgemm ... if (!llamafile_sgemm(...)) goto UseGgmlGemm1; // 如果失败，就跳回下面的普通通道 return; } UseGgmlGemm1:; #endif // 6. 如果src1的类型与src0的类型不匹配，需要将src1转换成vec_dot_type类型。 if (src1->type != vec_dot_type) { char * wdata = params->wdata; // ... (循环遍历 src1，把数据解压并复制到 wdata) ... // 这里用了 from_float 函数进行转换 } // 7. 【同步】确保所有线程都知道数据准备好了 //只让线程0做一次初始化，把 current_chunk 设为 nth：因为每个线程先从 current_chunk = ith //开始领一个“私有起始 chunk”，所以下一个“公共可抢”的 chunk 从 nth 开始。 if (ith == 0) { atomic_store_explicit(¶ms->threadpool->current_chunk, nth, memory_order_relaxed); } ggml_barrier(params->threadpool); // barrier 等待：确保所有线程都完成了“src1 预处理/计数器初始化”，再进入真正的 chunk 计算阶段。 // ... (再次尝试 llamafile 优化，省略) ... // 8. 【切蛋糕】决定每个任务块 (chunk) 有多大 // nr0 是结果矩阵的行数，nr1 是其余维度的大小 const int64_t nr0 = ne0; //nr0 代表结果矩阵在“第一维”的长度（可以理解为输出的行数/一个主维度）。 const int64_t nr1 = ne1 * ne2 * ne3; //nr1 把其余维度展平成第二维（把 ne1*ne2*ne3 看成“列块/批次展平后的列方向”）。 int chunk_size = 16; // 默认每块按 16 规模做分块（和 one_chunk 内的 16×16 tiling 呼应）。 if (nr0 == 1 || nr1 == 1) { chunk_size = 64; } // 如果矩阵很细长，就切大一点 // 算出总共有多少块蛋糕 (nchunk0 * nchunk1) int64_t nchunk0 = (nr0 + chunk_size - 1) / chunk_size; int64_t nchunk1 = (nr1 + chunk_size - 1) / chunk_size; // 9. 【NUMA 优化】如果是多路 CPU 服务器，或者切得太碎了，就改用保守策略 if (nchunk0 * nchunk1 < nth * 4 || ggml_is_numa()) { nchunk0 = nr0 > nr1 ? nth : 1; nchunk1 = nr0 > nr1 ? 1 : nth; } // 计算每个块实际包含多少行/列 (dr0, dr1) //把 nr0/nr1 平均分给 nchunk0/nchunk1，得到每块在两维上的跨度 dr0/dr1。 const int64_t dr0 = (nr0 + nchunk0 - 1) / nchunk0; const int64_t dr1 = (nr1 + nchunk1 - 1) / nchunk1; // 10. 【开始干活】线程开始领任务 // 一开始，第 ith 个线程领第 ith 块蛋糕 //每个线程先从 current_chunk = ith 开始，保证一开始每个线程都有活（减少原子争用）。 int current_chunk = ith; while (current_chunk < nchunk0 * nchunk1) { // 算出当前蛋糕块在矩阵的哪个位置 (start 到 end) const int64_t ith0 = current_chunk % nchunk0; const int64_t ith1 = current_chunk / nchunk0; const int64_t ir0_start = dr0 * ith0; const int64_t ir0_end = MIN(ir0_start + dr0, nr0); // ... (同理算出 ir1 的范围) ... // 11. 【核心调用】把这一小块任务交给工人去算 // num_rows_per_vec_dot 是看 CPU 是否支持一次算两行 (MMLA 指令) ggml_compute_forward_mul_mat_one_chunk(..., ir0_start, ir0_end, ir1_start, ir1_end); // 如果线程数比蛋糕块还多，干完这一块就可以下班了 if (nth >= nchunk0 * nchunk1) { break; } // 12. 【抢任务】否则，去原子计数器里“取号”，领下一块没人干的蛋糕 current_chunk = atomic_fetch_add_explicit(¶ms->threadpool->current_chunk, 1, memory_order_relaxed); } } 计算执行 static void ggml_compute_forward_mul_mat_one_chunk( // ... 参数省略，主要是范围 ir0_start 等 ... ) { // 1. 获取数据指针 const struct ggml_tensor * src0 = dst->src[0]; const struct ggml_tensor * src1 = dst->src[1]; // ... 初始化本地变量 ... // 2. 准备点积函数 (vec_dot) // type_traits_cpu 是一张表，根据数据类型查到最优的计算函数 ggml_vec_dot_t const vec_dot = type_traits_cpu[type].vec_dot; // 3. 【广播因子】计算 src1 比 src0 大多少倍 // 比如 src1 有 2 个头，src0 只有 1 个头，那 r2 就是 2。计算时 src0 要重复用。 const int64_t r2 = ne12 / ne02; const int64_t r3 = ne13 / ne03; // 如果任务范围是空的，直接返回 if (ir0_start >= ir0_end || ir1_start >= ir1_end) { return; } // 4. 【Tiling 分块参数】 // 即使是这一小块任务，也要再切碎，为了放进 CPU 缓存 (L1 Cache) const int64_t blck_0 = 16; const int64_t blck_1 = 16; // 5. 定义临时缓冲区，防止多个线程写内存冲突 (False Sharing) float tmp[32]; // 6. 【双重 Tiling 循环】开始遍历任务块 // iir1 每次跳 16 行 for (int64_t iir1 = ir1_start; iir1 < ir1_end; iir1 += blck_1) { // iir0 每次跳 16 列 for (int64_t iir0 = ir0_start; iir0 < ir0_end; iir0 += blck_0) { // 7. 【处理块内的每一行】 for (int64_t ir1 = iir1; ir1 < iir1 + blck_1 && ir1 < ir1_end; ir1 += num_rows_per_vec_dot) { // 8. 【计算索引】这部分数学很绕，简单说就是： // 根据当前的行号 ir1，反推出它在 src1 的第几层、第几页 (i13, i12, i11) const int64_t i13 = (ir1 / (ne12 * ne1)); const int64_t i12 = (ir1 - i13 * ne12 * ne1) / ne1; const int64_t i11 = (ir1 - i13 * ne12 * ne1 - i12 * ne1); // 9. 【应用广播】算出对应的 src0 的位置 // 如果 src0 比较小，就除以倍率 r2, r3 来复用数据 const int64_t i03 = i13 / r3; const int64_t i02 = i12 / r2; // 拿到 src0 这一行的内存地址 const char * src0_row = (const char*)src0->data + (0 + i02 * nb02 + i03 * nb03); // 拿到 src1 这一列的内存地址 (wdata 是之前解压好的数据) const char * src1_col = (const char*)wdata + ...; // (省略复杂的偏移量计算) // 拿到 结果 dst 的写入地址 float * dst_col = (float*)((char*)dst->data + ...); // 10. 【核心计算循环】 // 遍历 src0 的行 (分块内)，调用 SIMD 指令进行计算 for (int64_t ir0 = iir0; ir0 < iir0 + blck_0 && ir0 < ir0_end; ir0 += num_rows_per_vec_dot) { // vec_dot：一次性计算 ne00 个元素的点积 // 结果先存在 tmp 数组里，而不是直接写回 dst vec_dot(ne00, &tmp[ir0 - iir0], ..., src0_row + ir0 * nb01, ..., src1_col, ...); } // 11. 【写回结果】 // 把算好的 tmp 里的数据，一次性拷贝 (memcpy) 到最终的目标 dst_col for (int cn = 0; cn < num_rows_per_vec_dot; ++cn) { memcpy(&dst_col[iir0 + ...], tmp + ..., ...); } } } } } 以一个例子来总结一下。 void ggml_compute_forward_mul_mat( const struct ggml_compute_params * params, // 计算参数（包含线程信息） struct ggml_tensor * dst // 结果矩阵 C = A × B ) 功能是矩阵乘法：结果矩阵C=矩阵A x 矩阵B。其并行的策略如下： 切分对象： 结果矩阵C（不是A和B） 并行方式：2D切分+work_stealing（工作窃取） 写入方式：每个线程直接写入结果矩阵的不同位置。 矩阵A (m×k) × 矩阵B (k×n) = 结果矩阵C (m×n) ↑ ↑ ↑ 只读 只读 切分这个！ 如上，矩阵A和B是只读的，所有线程都需要读取完整的A和B，结果矩阵C需要写入，每个线程写入不同的位置，切分C可以让多个线程并行计算不同的部分，互不干扰。 示例演示 输入数据 矩阵A (4行×3列)： ┌─────────┐ │ 1 2 3 │ │ 4 5 6 │ │ 7 8 9 │ │ 10 11 12│ └─────────┘ 矩阵B (3行×4列)： ┌─────────────┐ │ 1 2 3 4 │ │ 5 6 7 8 │ │ 9 10 11 12 │ └─────────────┘ 要计算：结果矩阵C = A × B (4行×4列) 分配结果矩阵 结果矩阵C（内存已分配，4×4 = 16个float）： ┌─────────────────────┐ │ ? ? ? ? │ ← 行0 │ ? ? ? ? │ ← 行1 │ ? ? ? ? │ ← 行2 │ ? ? ? ? │ ← 行3 └─────────────────────┘ 列0 列1 列2 列3 计算切分参数 假设：2个线程，chunk_size = 2 nr0 = 4, nr1 = 4 nchunk0 = (4 + 2 - 1) / 2 = 3 nchunk1 = (4 + 2 - 1) / 2 = 3 总块数 = 3 × 3 = 9 检查：9 < 2 * 4 = 8？ 不，9 >= 8 -> 不触发优化，保持3×3切分 dr0 = (4 + 3 - 1) / 3 = 2 // 每个块2行 dr1 = (4 + 3 - 1) / 3 = 2 // 每个块2列 切分结果矩阵 结果矩阵C (4×4) 被切分成9个块： ┌──────┬──────┬──────┐ │ 块0 │ 块1 │ 块2 │ ← 行块0（行0-1） │[0,0] │[0,1] │[0,2] │ │[1,0] │[1,1] │[1,2] │ ├──────┼──────┼──────┤ │ 块3 │ 块4 │ 块5 │ ← 行块1（行2-3） │[2,0] │[2,1] │[2,2] │ │[3,0] │[3,1] │[3,2] │ ├──────┼──────┼──────┤ │ 块6 │ 块7 │ 块8 │ ← 行块2（超出边界） │[0,3] │[1,3] │[2,3] │ │[1,3] │[2,3] │[3,3] │ └──────┴──────┴──────┘ 映射表 列0 列1 列2 列3 ┌────┬────┬────┬────┐ │块0 │块0 │块3 │块3 │ ← 行0 │ │ │ │ │ ├────┼────┼────┼────┤ │块0 │块0 │块3 │块3 │ ← 行1 │ │ │ │ │ ├────┼────┼────┼────┤ │块1 │块1 │块4 │块4 │ ← 行2 │ │ │ │ │ ├────┼────┼────┼────┤ │块1 │块1 │块4 │块4 │ ← 行3 └────┴────┴────┴────┘ 块0: ith0=0, ith1=0 → C[0:2, 0:2] → 左上角 2×2 块1: ith0=1, ith1=0 → C[2:4, 0:2] → 左下角 2×2 块3: ith0=0, ith1=1 → C[0:2, 2:4] → 右上角 2×2 块4: ith0=1, ith1=1 → C[2:4, 2:4] → 右下角 2×2 转换的过程 // 步骤1：块编号 → 块坐标 current_chunk = 3 ith0 = 3 % 3 = 0 // 块在第0行（行方向的第0个块） ith1 = 3 / 3 = 1 // 块在第1列（列方向的第1个块） // 步骤2：块坐标 → 实际位置 ir0_start = dr0 * ith0 = 2 * 0 = 0 // 从第0行开始 ir0_end = MIN(0 + 2, 4) = 2 // 到第2行结束（不包括2，实际是0-1行） ir1_start = dr1 * ith1 = 2 * 1 = 2 // 从第2列开始 ir1_end = MIN(2 + 2, 4) = 4 // 到第4列结束（不包括4，实际是2-3列） // 步骤3：实际位置 → 结果矩阵元素 块3 = C[0:2, 2:4] = { C[0][2], C[0][3], C[1][2], C[1][3] } 之所以要做一个映射，是因为work-stealing需要一维的变化，原子操作只能操作一个整数，线程并发是以块为最小运算单位。也是为了方便后面做负载均衡。 并行运算 （1）线程0的工作 current_chunk = 0 // 从块0开始 // 处理块0 ith0 = 0 % 3 = 0, ith1 = 0 / 3 = 0 ir0_start = 0, ir0_end = 2 // 行0-1 ir1_start = 0, ir1_end = 2 // 列0-1 计算： C[0][0] = A[0行] · B[第0列] = [1,2,3]·[1,5,9] = 38 C[0][1] = A[0行] · B[第1列] = [1,2,3]·[2,6,10] = 44 C[1][0] = A[1行] · B[第0列] = [4,5,6]·[1,5,9] = 83 C[1][1] = A[1行] · B[第1列] = [4,5,6]·[2,6,10] = 98 直接写入结果矩阵： dst->data[0*4 + 0] = 38 // C[0][0] dst->data[0*4 + 1] = 44 // C[0][1] dst->data[1*4 + 0] = 83 // C[1][0] dst->data[1*4 + 1] = 98 // C[1][1] // 完成后，抢下一个块 current_chunk = atomic_fetch_add(&current_chunk, 1) = 2 // 处理块2（无效，会被跳过） ir0_start = 4, ir0_end = 4 // 空 → 跳过 // 继续抢下一个块 current_chunk = atomic_fetch_add(&current_chunk, 1) = 3 // 处理块3 （2）线程1的工作 current_chunk = 1 // 从块1开始 // 处理块1 ith0 = 1 % 3 = 1, ith1 = 1 / 3 = 0 ir0_start = 2, ir0_end = 4 // 行2-3 ir1_start = 0, ir1_end = 2 // 列0-1 计算并写入 C[2:4, 0:2]... // 完成后，抢下一个块 current_chunk = atomic_fetch_add(&current_chunk, 1) = 4 // 处理块4 （3）最终的结果 所有线程完成后，结果矩阵C： ┌─────────────────────┐ │ 38 44 50 56 │ ← 所有元素都已计算并写入 │ 83 98 113 128 │ │ 128 152 176 200 │ │ 173 206 239 272 │ └─────────────────────┘

使能方法 在linux 6.12版本之后原生SDK就支持了PREEMPT_RT，使能方式如下： make kernel_menuconfig General setup ---> <*> Fully Preemptible Kernel (Real-Time) 或者直接搜索CONFIG_PREEMPT_RT=y 确认是否已经打开 zcat /proc/config.gz | grep CONFIG_PREEMPT_RT 测试方法 安装软件 sudo apt-get update sudo apt-get install rt-tests stress-ng 空载测试 # -m: 锁定内存，防止交换 # -S: SMP 多核模式 # -p99: 设置最高实时优先级 99 # -i1000: 循环间隔 1000us (1ms) # -D2m: 测试持续 2 分钟 sudo cyclictest -m -S -p99 -i1000 -D2m 压力测试 终端1 # --cpu 8: 占满 8 个核心 # --io 4: 开启 4 个 IO 读写进程（模拟磁盘/存储负载） # --vm 4: 开启 4 个内存压力进程 # --vm-bytes 128M: 每个内存进程反复读写 128M 数据（制造 Cache 抖动） stress-ng --cpu 8 --cpu-method all --io 4 --vm 4 --vm-bytes 128M --fork 4 终端2 # -S: SMP 模式，自动开启 8 个线程并一对一绑核 # -p99: 最高实时优先级 # -m: 锁内存 # -D1: 跑 1 小时 (建议先跑 10 分钟看看情况) sudo cyclictest -c0 -m -S -p99 -i1000 -D1H

矩阵相乘 是神经网络中算力消耗最大的部分，通常占据 LLM 推理计算量的 95% 以上。 矩阵乘法 (Matrix Multiplication / GEMM) 这是最通用的矩阵运算形式，也是 AI 芯片中 Tensor Core 或 MAC 阵列的主要工作内容。 定义： 设矩阵 $A$ 的形状为 $(M \times K)$，矩阵 $B$ 的形状为 $(K \times N)$，则它们的乘积 $C = A \times B$ 的形状为 $(M \times N)$。 计算公式： 目标矩阵中第 $i$ 行第 $j$ 列的元素，等于 $A$ 的第 $i$ 行与 $B$ 的第 $j$ 列的对应元素乘积之和： $$ C_{ij} = \sum_{k=1}^{K} A_{ik} \cdot B_{kj} $$ 工程视角： 维度约束：左矩阵的列数 ($K$) 必须等于右矩阵的行数 ($K$)。这个 $K$ 维度在运算中会被“消掉”（Reduction）。 LLM 应用：全连接层 (Linear Layers)、注意力机制中的 $Q/K/V$ 投影。 硬件特性：典型的计算密集型算子。优化的核心在于提高数据复用率（Data Reuse），减少从 HBM/DRAM 读取数据的次数。 向量点积 (Vector Dot Product) 在数学上，点积是矩阵乘法的一种特例；在物理意义上，它是衡量相似度的工具。 定义：两个同维度向量 $\vec{a}$ 和 $\vec{b}$ 的运算。 计算公式： $$ \vec{a} \cdot \vec{b} = \sum_{i=1}^{n} a_i b_i $$ 结果：结果是一个标量 (Scalar)，即一个单纯的数值。 工程视角： 几何意义：反映两个向量方向的一致性。方向越接近，点积越大。 LLM 应用：Self-Attention 的核心逻辑。虽然代码实现通常是批量矩阵乘法 ($Q \times K^T$)，但其数学本质是计算 Query 向量与 Key 向量的点积来获得注意力分数。 逐元素运算 (Element-wise Operations) 这类运算的特点是不改变矩阵形状，且计算之间互不依赖。它们通常对算力要求不高，但对显存带宽极其敏感。 逐元累积 (Hadamard Product) 常被称为 "Element-wise Product"。 定义： 两个形状完全相同的矩阵 $A$ 和 $B$ 进行运算。 计算公式： $$ (A \odot B){ij} = A{ij} \times B_{ij} $$ 工程视角： LLM 应用： 门控 (Gating)：如 LLaMA 使用的 SwiGLU 激活函数，通过逐元素相乘来控制信息通过量。 掩码 (Masking)：在 Attention 矩阵中，将不需要关注的位置乘以 0（或加负无穷）。 加法与减法 定义： 两个形状相同的矩阵对应位置相加或相减。 计算公式： $$(A \pm B){ij} = A{ij} \pm B_{ij}$$ 工程视角： LLM 应用： 残差连接 (Residual Connection)：$X + \text{Layer}(X)$。这是深层网络能够训练的关键。 偏置 (Bias)：$Y = XW + b$。 硬件特性：典型的 Memory Bound 操作。因为每个数据读进来只做一次简单的加法就写回去了，算术强度（Arithmetic Intensity）极低。 标量乘法 (Scalar Multiplication) 定义： 一个单独的数值 $\lambda$ 乘以矩阵中的每一个元素。 计算公式： $$ (\lambda A){ij} = \lambda \cdot A{ij} $$ 工程视角： LLM 应用： 缩放 (Scaling)：Attention 中的 $\frac{QK^T}{\sqrt{d_k}}$，防止点积数值过大导致 Softmax 梯度消失。 归一化 (Normalization)：LayerNorm 中的 $\gamma$ 参数本质上也是一种特定维度的缩放。 结构变换 (Structural Transformations) 这类操作通常不涉及数值的改变，而是改变数据的排列方式或索引方式。 转置 (Transpose) 定义： 将矩阵的行和列互换。 公式： $$(A^T){ij} = A{ji}$$ 工程视角： LLM 应用：多头注意力机制 (Multi-Head Attention) 中，为了并行计算多个头，需要频繁进行 (Batch, Seq, Head, Dim) 到(Batch, Head, Seq, Dim) 的维度置换（Permute，广义的转置）。 硬件挑战：转置意味着非连续的内存访问。在硬件设计中，通常需要专门的 Transpose Engine 或者利用 SRAM/寄存器文件进行巧妙的数据混洗，否则会严重导致 Cache Miss。 广播 (Broadcasting) 这是工程实现中极为重要的机制，允许不同形状的张量进行算术运算。 定义： 当两个矩阵形状不匹配时，自动将较小的矩阵在特定维度上进行逻辑上的复制，使其与大矩阵形状一致，再进行运算。 规则： 从最后面的维度开始对齐。 如果某个维度大小为 1，则可以扩展到任意大小。 工程视角： 例子：矩阵 $A(100 \times 256)$ 加上向量 $b(1 \times 256)$。系统会将 $b$ 在第 0 维复制 100 次。 硬件优势：优秀的 SDK 或 NPU 设计支持“隐式广播”，即不需要真的在内存中复制数据，只需利用 Stride（步长）为 0 的寻址方式重复读取同一个数据，从而节省大量的显存带宽。

OpenMP是什么 OpenMP是一套用于共享内存并行系统的多线程程序设计标准。通俗的将，它允许通过简单的编译器指令（#pragma）将原本串行执行的C/C++ for循环瞬间变成多线程并发执行，而不需要手动的调用pthread_create的创建、销毁和同步。 核心模型：Fork-Join模型是理解OpenMP的关键。 Fork：主线程运行到并行区，Fork出一组线程。 Join：任务完成后，线程组同步并合并，只剩下主线程继续执行。 核心语法：从串行到并行 并行循环 这是最常用的指令。它告诉编译器：“把下面这个 for 循环的迭代次数拆分，分配给不同的线程同时跑”。 void add_arrays(float* a, float* b, float* c, int n) { for (int i = 0; i < n; i++) { c[i] = a[i] + b[i]; } } 要让他在多核CPU上并行，只需要加上一行： void add_arrays(float* a, float* b, float* c, int n) { // 自动将 n 次循环拆分给当前可用的线程 #pragma omp parallel for for (int i = 0; i < n; i++) { c[i] = a[i] + b[i]; // 每个线程处理一部分 i } } 数据作用域 写并发最怕数据竞争（Data Race）。OpenMP 必须明确变量是“大家共用”还是“每人一份”。 shared(x): 默认属性。所有线程读写同一个内存地址 x。比如输入的大矩阵。 private(i): 每个线程都有自己独立的变量 i 副本，互不干扰。比如循环计数器。 #pragma omp parallel for private(i, j) shared(A, B, C) for (int i = 0; i < rows; i++) { // ... } 规约 是计算“点积”（Dot Product）或“求和”时的神器。如果多个线程同时往同一个变量 sum 里加值，会冲突。 解决：使用 reduction(+:sum)。每个线程并在本地算自己的小 sum，最后 OpenMP 自动把所有线程的小 sum 加在一起。 float dot_product(float* a, float* b, int n) { float sum = 0.0; // 告诉编译器：sum 是归约变量，操作是加法 #pragma omp parallel for reduction(+:sum) for (int i = 0; i < n; i++) { sum += a[i] * b[i]; } return sum; } 循环是如何被切分的？ 初学者最大的疑问通常是：多线程并发执行循环时，i 变量不会冲突吗？大家都执行 i++ 岂不是乱套了？答案是：不会。OpenMP 做的是“分地盘”，而不是“抢地盘”。 （1）空间划分 当编译器看到 #pragma omp parallel for 时，它并没有让多个线程去争夺同一个全局 i。相反，它把 0 到 n 这个迭代范围（Iteration Space）切成了若干块（Chunks），分给了不同的线程。 假设 n = 100，且有 4 个线程，默认调度下： 线程 0 负责 i 属于 [0, 24] 线程 1 负责 i 属于 [25, 49] 线程 2 负责 i 属于 [50, 74] 线程 3 负责 i 属于 [75, 99] （2）变量私有化 在生成的底层代码中，循环变量 i 被自动标记为线程私有（Private）。 线程 1 里的 i 实际上是一个局部变量，它从 25 开始，自己累加到 49。 它完全看不到线程 0 的 i 是多少。 结论：所谓的并发，是大家在各自划定的跑道上同时跑，互不干扰。