Zhangzhe's Blog

Toolformer: Language Models Can Teach Themselves to Use Tools

Posted on 2025-07-04 Edited on 2025-07-07 In Agent Valine:

URL

paper: https://arxiv.org/pdf/2302.04761

TL;DR

大模型并不是全能的，大模型 + 工具可以更好地解决问题，本文提出一种新的训练方法，让大模型可以自我学习使用如何工具。
本文的主要贡献是提出了一种新的训练方法，包括通过大模型构造工具调用数据集、清洗数据集、使用工具调用数据集微调大模型。

Algorithm

ToolFormer 的意义

众所周知，大模型在很多任务上表现出色，但它们并不是全能的，比如去数 “strawberry” 这个单词中有多少个 “r”。
工具可以帮助大模型更好地解决问题，比如计算器、日历、知识库等。
一种常见的大模型和工具结合的方式是：通过 Agent 多角色（user / llm / function）多轮对话 形式调用，简单来说就是：
1. 大模型在需要调用工具的时候，输出一段特定格式的文本
2. 外部程序解析这段文本，调用相应的工具，调用得到结果
3. 新的结果作为 Function 角色的输入，继续和大模型对话
而 ToolFormer 采用的方式和 Agent 有相似之处，也有不同的地方：
- 相似点：
  1. 都需要大模型输出一段特定格式的文本来调用工具
  2. 都需要一段 endless loop 程序来解析大模型的输出，调用工具
- 不同点：
  1. ToolFormer 不是通过 多角色多轮对话 的方式调用工具，而是通过 单角色单轮对话 的方式调用工具
  2. ToolFormer 需要对大模型进行微调，而 Agent 不需要
ToolFormer 可以将存在确定答案的专用任务转化为工具调用任务（例如：计算、翻译、问答等），让大模型可以更专注在通用任务上（例如：上下文理解、常识知识运用等）。

ToolFormer 的工作方式

ToolFormer 是经过工具调用微调的大模型，知道有哪些工具可以调用，也知道如何调用这些工具。
假设模型输入的问题是：
1
Pittsburgh is also known as
这个时候，模型会意识到这个问题可以通过调用 Question Answering 工具来解决，于是模型会续写：
1
Pittsburgh is also known as <API>QA(Pittsburgh is also known as)</API>
输出 </API> 之后，推理进程会暂停推理模型，等待外部监听程序的调用结束。
外部监听程序会解析模型输出文本中的工具调用指令（通过 <API> </API> 格式），然后调用 Question Answering 工具，得到结果：
1
the Steel City
推理程序将工具调用结果和模型的历史输出（去掉调用相关信息）拼接起来，继续推理模型：
1
Pittsburgh is also known as the Steel City.

应该怎么得到 ToolFormer

假设上面提到的 ToolFormer 工作方式是一个愿景，那么接下来就是考虑应该如何得到一个这样的模型。
显然，直接拿一个预训练或经过微调的大模型来使用是行不通的，因为它并不知道哪些工具可以用，以及如何调用。
那么需要做的事就是：通过微调，让大模型自己学习如何使用工具
最困难的部分就是：如何构造工具调用数据集，因为不存在现成的工具调用数据集，需要自己构造。
这篇论文花了很大的篇幅就在讲一件事：如何用大模型来构造工具调用数据集

如何构造工具调用数据集

上面这张图展示了如何用大模型来构造工具调用数据集的流程。

首先，还是用预训练数据做为基础，假设一条预训练数据是：
1
Pittsburgh is also known as the Steel City.

然后，使用大模型（例如：GPT-3）来在数据中找到可以调用工具的位置和工具类型，并给出调用工具的参数，例如：

1	Pittsburgh is also known as <API>QA(What other name is Pittsburgh known by?) -> Steel City</API> the Steel City.

1	Pittsburgh is also known as <API>QA(Which country is Pittsburgh in?) -> United States</API> the Steel City.

计算调用工具带来的损失收益，剔除负收益的数据，保留正收益的数据。
- 损失的计算方式（离调用工具位置越远，损失权重 $w_{j-i}$ 越小）：
$L_i(z) = -\sum_{j=i}^n w_{j-i}\cdot \log p_M(x_j|z,x_{i;j-1})$
- 调用工具的损失：
$L_i^+=L_i(e(c_i,r_i))$
- 不调用工具 / 调用工具没有返回的损失：
$L_i^- = \min(L_i(\epsilon),L_i(e(c_i,\epsilon)))$
- 只保留调用工具损失收益大于阈值的样本：
$L_i^- - L_i^+ > \tau_f$
用筛选后的数据来微调大模型，得到 ToolFormer。

上图展示了 ToolFormer 支持的五种工具类型：
1. Question Answering：问答工具
2. Wikipedia Search：维基百科搜索工具
3. Calculator：计算器工具
4. Calendar：日历工具
5. Machine Translation：机器翻译工具

Thoughts

ToolFormer 这种通过大模型自我学习使用工具的方式感觉挺好的，但似乎在实际使用中没有得到大范围推广，目前主流外挂工具的方式基本还是 Agent 的多角色多轮对话方式。
可能是因为 ToolFormer 的方式需要对大模型进行微调，而 Agent 的方式不需要微调，直接使用预训练模型就可以。

ReAct: Synergizing Reasoning and Acting in Language Models

Posted on 2025-07-04 In Agent Valine:

URL

paper: https://arxiv.org/pdf/2210.03629
code: https://github.com/ysymyth/ReAct

TL;DR

人类通过语言推理（分解目标、调整计划）与行动（获取外部信息）的协同机制高效完成任务（如烹饪时动态调整步骤）。
受到人类智能启发，本文提出一种 Reasoning 和 Acting 相结合的框架，称为 ReAct，这种新的推理范式可以使语言模型在复杂任务中表现更好。
主要优势包括：
- 事实性：减少模型幻觉（虚构信息）。
- 决策鲁棒性：增强模型在复杂环境中的泛化能力。
- 可解释性：人类可以追踪模型的推理轨迹。

Algorithm

ReAct 框架

纯推理模型（如 Chain-of-Thought）：易产生事实幻觉（如虚构信息）和错误传播（如算术推理错误），缺乏实时环境交互能力（图 1b）。
纯行动模型（如 WebGPT）：缺乏高层规划能力，难以处理多步决策（图1c）。
ReAct 模型：交替进行推理和行动，能够在复杂任务中表现更好（图 1d）。

核心贡献

首提协同框架：统一推理与行动，解决静态推理与无规划行动的缺陷。
实践价值：
- 提升模型事实性（减少幻觉）与决策鲁棒性（复杂环境泛化）。
- 增强可解释性：人类可追踪推理轨迹。

Thoughts

论文提出的算法一眼开门，非常符合直觉，从哲学角度讲，一个事物的超集一般都优于其本身，因为最差的情况就是超集退化为本身（因此常见的一个情况是：新的论文称一篇老的论文是其某个参数设置下的一个特例）。
这些所有的 Reasoning 和 Acting 等 Agent 功能，都建立在大模型超长上下文的基础上，因此模型支持超长上下文长度是模型是否具有高级智能潜质的先决条件（至少现有范式是这样）。

Chain-of-Thought Prompting Elicits Reasoning in Large Language Models

Posted on 2025-07-01 In LLM Valine:

URL

paper: https://arxiv.org/pdf/2201.11903

TL;DR

本文是 Google Research 团队发表的一篇论文，这篇论文是在已开源的 非推理 模型上，不做微调，而是通过 Chain-of-Thought Prompting 的方式来引导模型进行推理，取得了很好的效果。
具体来说，就是在输入 prompt 中加入少量的 输入-思维链-输出 三元组示例，引导模型生成中间推理步骤，最终给出答案。
在算术、常识和符号推理等任务上，Chain-of-Thought Prompting 的效果都非常好，比直接给出答案的效果好很多。
Chain-of-Thought Prompting 是在模型规模达到一定程度（>= 100B）后，才涌现出的能力，小模型没有这个能力。

论文详情

1. 核心方法：思维链提示（Chain-of-Thought Prompting）

论文提出一种简单方法：在提示（prompt）中提供 输入-思维链-输出 的三元组示例，引导大型语言模型生成一系列中间推理步骤（称为 “思维链”），再得出最终答案。
思维链类似于人类逐步推理的过程（例如，解决数学题时先分解步骤：“先计算A，再计算B，最后得出答案”）。

2. 关键优势

提升复杂推理能力：在算术（如数学题 GSM8K）、常识（如 CSQA）和符号推理（如字母拼接游戏）任务上，思维链提示显著优于标准提示（standard prompting）。
在 GSM8K 数学题基准上，PaLM 540B 模型使用思维链提示后，准确率从 17.9% 提升至 56.9%，甚至超过微调的 GPT-3 模型。
在常识推理任务（如 StrategyQA）上，准确率从 68.6% 提升至 77.8%。
模型规模涌现特性：思维链推理是大型模型（约 100B 参数以上）的 “涌现能力” —— 小模型无法生成逻辑思维链，但足够大的模型（如 GPT-3 175B、PaLM 540B）能自然学习此模式。
无需微调：仅需在提示中添加少量示例（如 8 个），即可激发模型能力，无需额外训练或数据标注。
可解释性与泛化性：思维链提供透明推理路径，便于调试；且适用于多种任务（数学、常识、符号等），甚至能泛化到更长序列。

3. 实验验证

任务覆盖：
- 算术推理：在 GSM8K、SVAMP 等数据集上，思维链提示将性能提升高达 39%（PaLM 540B）。
- 常识推理：在 StrategyQA、Date Understanding 等任务上，模型表现接近或超越人类水平。
- 符号推理：在硬币翻转（coin flip）和字母拼接（last letter concatenation）任务中，模型能处理未见过的长序列。
鲁棒性：不同注释者编写的思维链示例均有效，且对示例顺序、数量变化不敏感。

4. 局限性与启示

模型规模依赖：思维链仅在大型模型（≥100B 参数）中有效，小模型生成逻辑混乱。
潜在错误：生成的推理路径可能不准确（如算术计算错误或语义误解），需外部验证（如添加计算器）。
应用意义：该方法拓展了提示技术的边界，证明大型模型能通过自然语言示例学习复杂推理，减少对标注数据的依赖。

论文核心贡献

思维链提示是一种低成本、高效的方法，通过模拟人类逐步推理过程，释放大型语言模型在复杂任务上的潜力。论文强调，这是 “模型规模涌现” 的典型例子——推理能力随模型增大而自然出现，为未来 AI 推理研究提供了新方向。

Thoughts

一定要注意，这篇论文讨论的对象 不是 Reasoning 模型（这个论文出来的时候还没有 Reasoning 模型的概念），而是普通的 LLM 模型。
本质是一种通过 prompt 引导模型通过增加推理计算预算的方式，来提升模型的推理能力的方法。
依托于 LLM 恐怖的指令遵循和上下文学习能力。

Large Language Diffusion Models

Posted on 2025-05-28 Edited on 2025-05-29 In Diffusion LLM Valine:

URL

paper: https://arxiv.org/pdf/2502.09992
code: https://ml-gsai.github.io/LLaDA-demo/

TL;DR

LLaDA 提出了一个新概念，叫 “扩散语言模型”，和主流的自回归语言模型 predict next token 的方式不同，LLaDA 使用类似 Diffusion 去噪的方法，一次性生成多个 token，通过多次生成，得到一个完整的生成文本。
但细看就会发现，Diffusion 就是一个彻头彻尾的噱头，和经典的热力学扩散过程没有鸡毛关系，LLaDA 本质就是一个大 BERT 模型，用完形填空的方式来生成文本（一次可以做多个完形填空），只是下图所示的每轮迭代的过程看起来有点像 Diffusion 的去噪（没关系硬蹭）。

上图来自官方 repo 的 README

Algorithm

总体流程

虽然多少有点标题党，但这篇论文本身是值得一读的，将文本生成任务做了重新定义，确实可大幅提高生成速度。

模型架构

纯纯 Transformer encoder 架构，和 BERT 类似，双向注意力，模型参数规模达 8B

训练过程

预训练
- 使用随机 mask 一定比例的 token，然后使用 Transformer 预测被 mask 的 token（完形填空）
- 损失函数：mask 部分的 cross-entropy 损失
- 数据规模：2.3 万亿 token，包含通用文本、代码和多语言数据
SFT
- 目标：使模型具备指令跟随能力
- 数据格式：成对数据 $(p_\theta,r_\theta)$ ，其中 $p_\theta$ 是指令， $r_\theta$ 是响应
- 掩码策略：仅对响应部分掩码，保持指令完整
- 损失函数：仅对响应部分计算 cross-entropy 损失
- 数据规模：450 万对指令响应对，涵盖代码、数学和多轮对话

推理与生成

过程：从全掩码的响应开始，逐步预测并更新掩码 token，直到生成完整响应
重掩码策略（预测之后 mask 一部分生成结果做二次生成）：
- 随机重掩码：基础策略，与扩散过程对齐
- 低置信度重掩码：优先掩码预测置信度低的 token
- 半自回归策略（SFT后）：分块生成，块内并行预测以提高效率
生成效果：支持多轮对话、多语言翻译和复杂推理任务

和自回归模型对比

特性	自回归模型（如`GPT`）	`LLaDA`
生成顺序	严格从左到右逐 `token` 生成	并行预测 + 动态调整
计算效率	需串行预测多次	仅需少量迭代（块级并行）
错误修正能力	无法修改已生成 `token`	通过重掩码可修正低置信度位置
逆向推理支持	受限于单向建模	双向注意力机制支持逆向推理

Thought

预测下一个词的大模型范式是否一定是最优的？可能未必。这篇论文就提出了一个不错的思路
make bert great again ~~手动滑稽~~

High-Resolution Image Synthesis with Latent Diffusion Models

Posted on 2025-05-11 Edited on 2025-06-30 In stable diffusion Valine:

URL

paper: https://arxiv.org/pdf/2112.10752
code: https://github.com/CompVis/latent-diffusion

TL;DR

LDM 是 stable diffusion 系列的开山之作，让 Diffusion Model 在 Image Synthesis 领域大放异彩
传统的 Diffusion Model 有两大问题：
1. 没办法控制生成内容，只能确保生成的内容和训练数据集风格比较类似
2. 在像素尺度上做去噪，计算量大，导致只能生成较小分辨率的图，且很慢
LDM 解决了上述的两个问题：
1. 通过编码器将文本 / mask / bbox 等条件信息转成 conditioning embedding，再通过 cross attention 机制将条件信息和 latent space 中的噪声结合起来做去噪，让条件信息可引导图片生成
2. 通过 VAE 将图片压缩到 latent space 中，再进行去噪，计算量小，速度快

Algorithm

总体流程

生成隐空间下的随机噪声
将条件信息通过各自类型的编码器编码成 conditioning embedding，例如文本编码使用 CLIP text encoder
将 latent noise 和 conditioning embedding 和 timestep embedding 输入到 UNet 中进行多轮迭代去噪（50 step）
去噪后的 latent 通过 VAE decoder 解码成图片

Conditioning UNet

通过 交叉注意力机制 将 文本条件 和时序通过 UNet 嵌入到 Noised Latent 中
具体来说：TimeEmbedding 直接和图像特征相加，TextEmbedding 和图像特征做 CrossAttention，（TextEmbedding 作为 KV，图像特征作为 Q）
Conditioning UNet 示意图，两次下采样 + 中间块 + 两次上采样

graph TD
    Input[Noised Latent: 32x32x4] --> DownBlock1[CrossAttnDownBlock2D]
    DownBlock1 --> DownBlock2[CrossAttnDownBlock2D]
    DownBlock2 --> MidBlock[UNetMidBlock2DCrossAttn]
    MidBlock --> UpBlock1[CrossAttnUpBlock2D]
    UpBlock1 --> UpBlock2[CrossAttnUpBlock2D]
    UpBlock2 --> Output[Denoised Latent: 32x32x4]
  
    TextEncoder[Text Encoder] -->|Text Embedding| DownBlock1
    TextEncoder -->|Text Embedding| DownBlock2
    TextEncoder -->|Text Embedding| MidBlock
    TextEncoder -->|Text Embedding| UpBlock1
    TextEncoder -->|Text Embedding| UpBlock2
  
    Time[Timestep] -->|Time Embedding| DownBlock1
    Time -->|Time Embedding| DownBlock2
    Time -->|Time Embedding| MidBlock
    Time -->|Time Embedding| UpBlock1
    Time -->|Time Embedding| UpBlock2

CrossAttnBlock2D 结构示意

graph TD
    %% 输入节点
    Input[输入特征图 h_in] --> ResNet
    TimeEmb[时间嵌入 t_emb] --> MLP
    TextEmb[文本条件 y_text] --> ProjText
  
    %% 主干计算路径
    ResNet[ResNet块] --> Add
    MLP[MLP时间投影] --> Add
    Add[逐元素相加] --> GroupNorm
    GroupNorm[GroupNorm] --> Conv1
    Conv1[Conv2D 1x1] --> CrossAttn
  
    %% 交叉注意力分支
    ProjText[文本投影 W_k/W_v] --> CrossAttn
    Conv2[Conv2D 1x1] --> Merge
    CrossAttn[交叉注意力层] --> Merge
  
    %% 残差连接
    Input --> Conv2
    Merge[特征合并] --> LayerNorm
    LayerNorm[LayerNorm] --> Output[输出特征图 h_out]

DecoderAttentionBlock2D 结构示意

graph TD
    X[("Input x
Shape: 1,512,32,32")] --> Norm["Normalize (GroupNorm)
Output: 1,512,32,32"]
  
    Norm --> Q["Q Conv2d(1x1)
Output: 1,512,32,32"]
    Norm --> K["K Conv2d(1x1)
Output: 1,512,32,32"]
    Norm --> V["V Conv2d(1x1)
Output: 1,512,32,32"]
  
    Q --> ReshapeQ["Reshape & Permute
1,512,32,32 → 1,1024,512"]
    K --> ReshapeK["Reshape
1,512,32,32 → 1,512,1024"]
  
    ReshapeQ --> MatmulQK["Matmul(Q,K)
1,1024,512 × 1,512,1024 → 1,1024,1024"]
    ReshapeK --> MatmulQK
  
    MatmulQK --> Scale["Scale (×1/√512)
1,1024,1024"]
    Scale --> Softmax["Softmax
1,1024,1024"]
  
    V --> ReshapeV["Reshape
1,512,32,32 → 1,512,1024"]
    Softmax --> PermuteSoftmax["Permute
1,1024,1024 → 1,1024,1024"]
  
    ReshapeV --> MatmulVW["Matmul(V, Softmax)
1,512,1024 × 1,1024,1024 → 1,512,1024"]
    PermuteSoftmax --> MatmulVW
  
    MatmulVW --> ReshapeOut["Reshape
1,512,1024 → 1,512,32,32"]
  
    ReshapeOut --> ProjOut["Proj_out Conv2d(1x1)
1,512,32,32"]
    ProjOut --> Add["Add (x + h_)
1,512,32,32"]
  
    X --> Add
    Add --> Output[("Final Output
1,512,32,32")]

下面附上 Conditioning UNet block 的实现代码，可以看出非常优雅：

import torch
import torch.nn as nn
import torch.nn.functional as F


class Attention(nn.Module):
    def __init__(self, in_dim, context_dim=None):
        super().__init__()
        self.to_q = nn.Linear(in_dim, in_dim, bias=False)
        self.to_k = nn.Linear(
            context_dim if context_dim else in_dim, in_dim, bias=False
        )
        self.to_v = nn.Linear(
            context_dim if context_dim else in_dim, in_dim, bias=False
        )
        self.to_out = nn.Sequential(nn.Linear(in_dim, in_dim), nn.Dropout(0.0))

    def forward(self, x, context=None):
        q = self.to_q(x)
        k = self.to_k(context if context is not None else x)
        v = self.to_v(context if context is not None else x)

        attn = torch.einsum("b i d, b j d -> b i j", q, k) * (x.shape[-1] ** -0.5)
        attn = F.softmax(attn, dim=-1)
        out = torch.einsum("b i j, b j d -> b i d", attn, v)
        return self.to_out(out)


class GEGLU(nn.Module):
    def __init__(self, in_dim, hidden_dim):
        super().__init__()
        self.proj = nn.Linear(in_dim, hidden_dim * 2)

    def forward(self, x):
        x_proj = self.proj(x)
        x1, x2 = x_proj.chunk(2, dim=-1)
        return x1 * F.gelu(x2)


class FeedForward(nn.Module):
    def __init__(self, in_dim, hidden_dim):
        super().__init__()
        self.net = nn.Sequential(
            GEGLU(in_dim, hidden_dim), nn.Dropout(0.0), nn.Linear(hidden_dim, in_dim)
        )

    def forward(self, x):
        return self.net(x)


class BasicTransformerBlock(nn.Module):
    def __init__(self, dim):
        super().__init__()
        self.norm1 = nn.LayerNorm(dim, eps=1e-5)
        self.attn1 = Attention(dim)
        self.norm2 = nn.LayerNorm(dim, eps=1e-5)
        self.attn2 = Attention(dim, context_dim=768)
        self.norm3 = nn.LayerNorm(dim, eps=1e-5)
        self.ff = FeedForward(dim, 1280)

    def forward(self, x, context=None):
        # Self attention
        x = self.attn1(self.norm1(x)) + x
        # Cross attention
        x = self.attn2(self.norm2(x), context=context) + x
        # Feed forward
        x = self.ff(self.norm3(x)) + x
        return x


class Transformer2DModel(nn.Module):
    def __init__(self, in_channels):
        super().__init__()
        self.norm = nn.GroupNorm(32, in_channels, eps=1e-6, affine=True)
        self.proj_in = nn.Conv2d(in_channels, in_channels, kernel_size=1)
        self.transformer_blocks = nn.ModuleList([BasicTransformerBlock(in_channels)])
        self.proj_out = nn.Conv2d(in_channels, in_channels, kernel_size=1)

    def forward(self, x, context=None):
        b, c, h, w = x.shape
        x_in = x
        x = self.norm(x)
        x = self.proj_in(x)
        x = x.permute(0, 2, 3, 1).reshape(b, h * w, c)

        for block in self.transformer_blocks:
            x = block(x, context)

        x = x.reshape(b, h, w, c).permute(0, 3, 1, 2)
        x = self.proj_out(x)
        return x + x_in


class ResnetBlock2D(nn.Module):
    def __init__(self, in_channels):
        super().__init__()
        self.norm1 = nn.GroupNorm(32, in_channels, eps=1e-5, affine=True)
        self.conv1 = nn.Conv2d(in_channels, in_channels, kernel_size=3, padding=1)
        self.time_emb_proj = nn.Linear(1280, in_channels)
        self.norm2 = nn.GroupNorm(32, in_channels, eps=1e-5, affine=True)
        self.dropout = nn.Dropout(0.0)
        self.conv2 = nn.Conv2d(in_channels, in_channels, kernel_size=3, padding=1)
        self.nonlinearity = nn.SiLU()

    def forward(self, x, time_emb=None):
        h = x
        h = self.norm1(h)
        h = self.nonlinearity(h)
        h = self.conv1(h)

        if time_emb is not None:
            time_emb = self.nonlinearity(time_emb)
            time_emb = self.time_emb_proj(time_emb)[:, :, None, None]
            h = h + time_emb

        h = self.norm2(h)
        h = self.nonlinearity(h)
        h = self.dropout(h)
        h = self.conv2(h)
        return h + x


class Downsample2D(nn.Module):
    def __init__(self, channels):
        super().__init__()
        self.conv = nn.Conv2d(channels, channels, kernel_size=3, stride=2, padding=1)

    def forward(self, x):
        return self.conv(x)


class CrossAttnDownBlock2D(nn.Module):
    def __init__(self, in_channels=320):
        super().__init__()
        self.attentions = nn.ModuleList(
            [Transformer2DModel(in_channels) for _ in range(2)]
        )
        self.resnets = nn.ModuleList([ResnetBlock2D(in_channels) for _ in range(2)])
        self.downsamplers = nn.ModuleList([Downsample2D(in_channels)])

    def forward(self, x, context=None, time_emb=None):
        for attn, resnet in zip(self.attentions, self.resnets):
            x = attn(x, context)
            x = resnet(x, time_emb)

        for downsampler in self.downsamplers:
            x = downsampler(x)

        return x


# 测试代码
if __name__ == "__main__":
    block = CrossAttnDownBlock2D(in_channels=320)
    x = torch.randn(1, 320, 64, 64)
    context = torch.randn(1, 77, 768)  # 文本条件，77 个 token 组成的文本序列经过 CLIP 编码成向量 
    time_emb = torch.randn(1, 1280)  # 时间嵌入，一个时间步（例如：961）变成一个 enbedding

    output = block(x, context, time_emb)
    print(f"输入形状: {x.shape} -> 输出形状: {output.shape}")
    # 预期输出: torch.Size([1, 320, 32, 32])

Thoughts

论文思路无比清晰，且说服力很强，把很多领域的知识结合起来，真正把图像生成在实用性方面推到了一个新的高度

Denoising Diffusion Implicit Models

Posted on 2025-05-10 Edited on 2025-06-30 In Diffusion Valine:

URL

paper: https://arxiv.org/pdf/2010.02502

TL;DR

这篇论文提出了去噪扩散隐式模型 Denoising Diffusion Implicit Models (DDIM) 模型，可以看作是对 Denoising Diffusion Probabilistic Models (DDPM) 模型的改进。
DDPM 的采样过程是一个 Markov 过程，Markov 过程只有知道 t 时刻的状态才能计算第 t-1 时刻，而 DDIM 的采样过程是一个非 Markov 过程。
DDIM 的优势是：
1. 去噪速度更快。可以在采样（去噪）时使用更少的时间步数，从而加快采样速度，并且用 DDPM 训练的模型可以直接用于 DDIM 的采样，二者可以无缝衔接。
2. 确定性。在 DDIM 中，给定一个模型和一个噪声图像和时间步数，可以确定性地生成一个图像（运行再多次也是同一个图）。在 DDPM 中，给定一个模型和一个噪声图像和时间步数，生成的图像是随机的（每次跑都不一样）。

Algorithm

DDIM 的公式是从 DDPM 公式通过复杂推导得到的，推导过程比较复杂，这里不做详细介绍，重点讲二者逆向过程（去噪）公式区别和使用上的区别。

DDPM 逆向过程公式

x_{t-1}=\frac{1}{\sqrt{\alpha_t}}(x_t-\frac{\beta_t}{\sqrt{1-\bar\alpha_t}}\cdot \epsilon_\theta(x_t,t))

其中：
- $\epsilon_\theta(x_t,t)$ 是模型预测的噪声，即 model(x_t, t)
- $\alpha_t=1-\beta_t$
- $\bar\alpha_t=\prod_{s=1}^t\alpha_s$
- t 取值是 999, 998,..., 1, 0，即从 T-1 到 0 逐步去噪

DDIM 逆向过程公式

x_{t-1}=\sqrt {\bar\alpha_t-1}(\frac{x_t-\sqrt {\bar\alpha_t-1}\cdot\epsilon_\theta(x_t,t)}{\sqrt {\bar\alpha_t}})+\sqrt{1-\bar\alpha_{t-1}}\cdot\epsilon_\theta(x_t,t)

其中：
- 大多数符号含义都和 DDPM 一样
- 只有 t 取值是 980, 960,..., 20, 0 这种非连续的取值

`DDIM` 公式拆解：

1. 预测原输入

pred_{x0}=\frac{x_t-\sqrt {\bar\alpha_t-1}\cdot\epsilon_\theta(x_t,t)}{\sqrt{\bar\alpha_t}}

通过当前噪声隐变量 $x_t$ 和模型预测的噪声 $\epsilon_\theta(x_t,t)$ 估计原始输入 $x_0$ ，即去噪

2. 计算调整方向

direction\_point=\sqrt{1-\bar\alpha_{t-1}}\cdot\epsilon_\theta(x_t,t)

根据噪声预测结果，计算从当前时间步 t 到前一个时间步 t-1 的调整方向，这一方向结合了噪声预测和噪声调度参数，用于引导隐变量的更新。

3. 更新隐变量

x_{t-1}=\sqrt {\bar\alpha_t-1}\cdot pred_{x0}+direction\_point

将预测的原始输入 $pred_{x0}$ 与调整方向结合，生成前一时刻的隐变量 $x_{t-1}$ 。此步骤通过线性组合逐步去噪，最终逼近目标数据 $x_0$ 。

DDPM 和 DDIM 使用上的区别

DDPM 去噪过程

model = UNet2DModel.from_pretrained("google/ddpm-celebahq-256").to(device)
scheduler = DDPMScheduler.from_pretrained("google/ddpm-celebahq-256")
# Get precalculated alphas and alpha bars from the scheduler
alphas = scheduler.alphas
alphas_cumprod = scheduler.alphas_cumprod
# Initialize sample with static random noise
sample = torch.load("random_noise.pt").to(device)
# DDPM denoising loop
# scheduler.timesteps = [999, 998, ..., 1, 0]
for t in tqdm.tqdm(scheduler.timesteps):
    with torch.no_grad():
        # Model prediction (noise residual)
        residual = model(sample, t).sample
        # DDPM denoising formula
        sample = (
            sample - (1 - alphas[t]) / torch.sqrt(1 - alphas_cumprod[t]) * residual
        ) / torch.sqrt(alphas[t])
        # Add random noise only for t > 1
        if t > 1:
            noise = torch.randn_like(sample).to(device)
            sample += torch.sqrt(1 - alphas[t]) * noise

DDIM 去噪过程

model = UNet2DModel.from_pretrained("google/ddpm-celebahq-256").to(device)
scheduler = DDIMScheduler.from_pretrained("google/ddpm-celebahq-256")
# set inference steps
scheduler.set_timesteps(num_inference_steps=50)
# Initialize sample with static random noise
sample = torch.load("random_noise.pt").to(device)
# DDIM denoising loop
# scheduler.timesteps = [980, 960,..., 20, 0]
for i, t in enumerate(tqdm.tqdm(scheduler.timesteps)):
    # 将时间步转换为LongTensor并确保在正确设备上
    t = t.to(device).long()
    # 获取当前和上一步的alpha累积乘积
    alpha_cumprod_t = scheduler.alphas_cumprod[t]
    alpha_cumprod_prev = (
        scheduler.alphas_cumprod[scheduler.timesteps[i + 1]]
        if i + 1 < len(scheduler.timesteps)
        else torch.tensor(1.0)
    )
    # 将alpha值转换到相同设备
    alpha_cumprod_t = alpha_cumprod_t.to(device)
    alpha_cumprod_prev = alpha_cumprod_prev.to(device)
    with torch.no_grad():
        # 1. 预测噪声残差
        residual = model(sample, t).sample
        # 2. 计算预测的原始图像x0（去噪后的图像）
        pred_x0 = (sample - torch.sqrt(1.0 - alpha_cumprod_t) * residual) / torch.sqrt(
            alpha_cumprod_t
        )
        # 3. 计算下一步的样本方向
        direction_xt = torch.sqrt(1.0 - alpha_cumprod_prev) * residual
        # 4. 组合得到新的样本
        sample = torch.sqrt(alpha_cumprod_prev) * pred_x0 + direction_xt

二者对比分析

DDIM 只需要 50 次迭代就能生成高质量的图像，而 DDPM 需要 1000 次迭代。
生成的图像质量相似，DDIM 生成的图像质量略高。
上面的代码中加载的噪声图是静态的，DDIM 跑多次生成的图像是一样的，而 DDPM 跑多次生成的图像是不一样的。
二者去噪结果对比，左侧是 DDIM，右侧是 DDPM：

Thoughts

DDIM 解决了 DDPM 的两大痛点，算是一个很好的改进。
为后续的 LDM 等模型打下了基础。

Qwen3 技术报告先导篇

Posted on 2025-05-05 Edited on 2025-05-07 In LLM Valine:

URL

blog: https://qwenlm.github.io/zh/blog/qwen3/

TL;DR

Qwen3 系列模型四月二十九日正式发布，但到目前为止，还没发布技术报告，只有一篇官方博客，介绍了 Qwen3 的一些基本信息
本文围绕着官方博客介绍，结合实际使用情况，给出一些个人的理解

Qwen3 系列模型

本次 Qwen3 主要发布了八个版本的模型，其中包含两个 MOE 模型和六个 dense 模型：
- Qwen3-235B-A22B：MOE 模型，235B 参数量，22B 激活参数量
- Qwen3-30B-A3B：MOE 模型，30B 参数量，3B 激活参数量
- Qwen3-32B：dense 模型，32B 参数量
- Qwen3-14B：dense 模型，14B 参数量
- Qwen3-8B：dense 模型，8B 参数量
- Qwen3-4B：dense 模型，4B 参数量
- Qwen3-1.7B：dense 模型，1.7B 参数量
- Qwen3-0.6B：dense 模型，0.6B 参数量
还有对应的 Base 模型（只经过预训练）和 fp8 Quantized 模型（量化模型）
Qwen3 相较于上一代 Qwen2.5，一个较大的技术进步是：统一了 Reasoning 和非 Reasoning 模式，即不再区分推理模型和非推理模型（或者叫思考模型和非思考模型），而是一个模型可以通过 prompt 来选择推理模式和非推理模式

Qwen3 模型的主要特性

预训练

相较于上一代 Qwen2.5 使用了 18 万亿个 token 做预训练，Qwen3 使用了 36 万亿个 token 做预训练（整整翻了一倍，这是多大的数据团队才能搞出来的，~~太壕了~~）
包含了 119 种语言和方言，有大量数据是合成数据和在 PDF 上识别的文本数据
预训练分成了三个阶段：
- 第一阶段：模型在超过 30 万亿个 token 上进行了预训练，上下文长度为 4K token
- 第二阶段：通过增加知识密集型数据（如 STEM、编程和推理任务）的比例来改进数据集，在 5 万亿个 token 上进行了预训练
- 第三阶段：使用高质量的长上下文数据将上下文长度扩展到 32K token
预训练实际效果：官方的说法是 Qwen3-1.7B/4B/8B/14B/32B-Base 分别与 Qwen2.5-3B/7B/14B/32B/72B-Base 表现相当

后训练

简单概括：用 SFT → RL → SFT+RL 混合 → RL 四个阶段后训练 Qwen3-235B-A22B 和 Qwen3-32B 模型，然后蒸馏得到其他小尺寸模型（大 MOE 蒸小 MOE，大 dense 蒸小 dense）
四个阶段分别是：
- 长思维链冷启动：使用监督微调（SFT）训练模型生成初步的长思维链推理能力，作为初始阶段的基础。
- 长思维链强化学习：通过强化学习（RL）进一步提升模型在复杂推理任务中的表现，优化生成思维链的质量和连贯性。
- 思维模式融合：结合 SFT 和 RL 的混合策略，将不同思维模式（如逻辑推理、知识检索等）整合到统一框架中，增强模型的灵活性。
- 通用强化学习：以 RL 为主导，对模型进行全局优化，强化其在多样化任务中的通用性和鲁棒性。

Qwen3 使用

使用上和 Qwen2.5 没有太大区别，主要是增加了 Reasoning 模式的选择开关
这个开关也非常简便，大概有两种方式可以选择：
1. 在最新版 Transformer 库中 tokenizer.apply_chat_template 中可以设置 enable_thinking=True # Switch between thinking and non-thinking modes. Default is True.
2. 另外一种方式就是始终保持上面的开关常开，然后在 prompt 中添加 /no_think 来关闭推理模式。/no_think 可加到 system content 或 user content 结尾。
在一些下游任务上做了 SFT 微调测试，发现 Qwen3-1.7B 经过 SFT 之后和 Qwen2.5-3B-Instruct 经过 SFT 之后的效果差不多，这一点和官方对预训练的说法一致

Thoughts

在当今这个大模型结构同质化时代，预训练数据越多，模型的能力越强（同样效果下数据翻一倍模型尺寸可缩减一半！！），所以小厂没有这么多数据工程师，搞不到大量的高质量的数据，是很难在大模型上追赶大厂
到今天为止，Llama 系列官方模型都不能做到原生支持中文，和 Qwen 系列模型原生支持 119 种语言和方言相比，感觉非常小家子气，注定会被扫进历史的垃圾堆里
DeepSeek 对 Qwen 的影响很大，比如：
- 从 DeepSeek v3 之后，Qwen 系列模型预训练和最终版的命名从：Qwen2.5-7B/Qwen2.5-7B-Instruct 变成了 Qwen3-8B-Base/Qwen3-8B
- 广泛的使用了模型蒸馏，确实极大的提高了小尺寸模型的能力（之前 Qwen2.5-0.5B-Instruct 基本就是个答非所问的傻子，现如今的 Qwen3-0.6B 在不开推理模式的情况下也可以解决很多数学问题，非常强）
- 四段式后训练也和 DeepSeek 使用的后训练非常相似
合并推理和非推理模型的做法，是今后大模型的趋势，有不少模型都在朝着这个方向发展
官方表示 Qwen 系列后续会向着 Agent 方向发展，铺垫了很久的 MCP 可能会产生新的 Agent 应用变革

CUDA 学习笔记 01 —— CUDA 基础

Posted on 2025-03-24 Edited on 2025-05-22 In CUDA Valine:

0. 学习 CUDA 的目的

作为一个算法工程师，平时接触 HPC (High Performance Computing) 的机会并不多，那为什么还要学习 CUDA 呢？
学习 CUDA 的目的不是为了用 CUDA 做模型加速，而是从 CUDA 角度理解目前较新的大模型设计理念，这些高性能模型是如何从原理上做到又快又好的。
例如火出圈的 DeepSeek 系列模型，在模型设计角度做了较多创新，并开源了部分 CUDA 代码，对于不了解 CUDA 的工程师，很难 get 到算法设计的优雅之处。
反观某家大模型基座公司，曾开源某个模型结构，论文中一通自夸，分析理论计算量有多低。但很多人实测发现速度并没有很快，究其原因，实际上是这家公司还用的小模型时代的旧思维，即：一个模型理论计算量低，那就是快。
大模型时代不了解硬件，不尊重硬件，在算法创新上不太可能走的远。

1. Hello World

cuda 代码

#include <cuda_runtime_api.h>
#include <stdio.h>
// cuda 中 host 表示 cpu 端，device 表示 gpu 端
// __device__ 是设备函数的声明符号，表明该函数在 device 执行，且只能在 device
// 中调用
__device__ const char *device_hello_world(void) {
  return "GPU: Hello world!\n";
}
// __host__ 是主机函数的声明符号，表明该函数在 host 执行，且只能在 host 中调用
__host__ const char *host_hello_world(void) { return "CPU: Hello world!\n"; }
// __global__ 是核函数的声明符号，表明该函数在 device 执行，且只能在 host 中调用
__global__ void hello_world(void) {
  const char *str = device_hello_world();
  printf("%s", str);
}
int main(int argc, char **argv) {
  printf("%s", host_hello_world());
  // <<<grid_dim, block_dim>>> 是核函数的调用符号，表示启动 grid_dim 个 block，
  // 每个 block 有 block_dim 个线程
  hello_world<<<1, 10>>>();
  cudaDeviceReset();
  return 0;
}

cuda 的三个函数声明符号：
- __host__：主机函数，表示该函数在 CPU 上执行，且只能在 CPU 中调用
- __device__：设备函数，表示该函数在 GPU 上执行，且只能在 GPU 中调用
- __global__：核函数，表示该函数在 GPU 上执行，且只能在 CPU 中调用
其中 __global__ 声明的函数类型被称为 核函数，是 CUDA 中最重要的函数类型
- 核函数通过 <<<grid_dim, block_dim>>> 的方式调用，其中 <<<>>> 是 cuda 扩展关键字
- grid_dim 表示启动的 block 数量，block_dim 表示每个 block 中的线程数量
- grid_dim 和 block_dim 都是 dim3 类型的变量，表示三维数组，如果使用整形则模型 y 和 z 维度都为 1

编译

1	nvcc hello_world.cu -g -o hello_world

运行

./hello_world
# CPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!

2. dimension 测试

#include <cuda_runtime.h>
#include <stdio.h>
__global__ void checkIndex(void) {
  printf("threadIdx:(%d,%d,%d) blockIdx:(%d,%d,%d) blockDim:(%d,%d,%d)\
  gridDim(%d,%d,%d)\n",
         threadIdx.x, threadIdx.y, threadIdx.z, blockIdx.x, blockIdx.y,
         blockIdx.z, blockDim.x, blockDim.y, blockDim.z, gridDim.x, gridDim.y,
         gridDim.z);
}
int main(int argc, char **argv) {
  int nElem = 6;  // number of elements
  dim3 block(3);  // block size
  int nBlock = (nElem + block.x - 1) / block.x; // number of blocks
  dim3 grid(nBlock);  // grid size
  printf("grid.x %d grid.y %d grid.z %d\n", grid.x, grid.y, grid.z);  
  printf("block.x %d block.y %d block.z %d\n", block.x, block.y, block.z);
  checkIndex<<<grid, block>>>();
  cudaDeviceReset();
  return 0;
}

执行结果：

./check_dimension
# grid.x 2 grid.y 1 grid.z 1
# block.x 3 block.y 1 block.z 1
# threadIdx:(0,0,0) blockIdx:(1,0,0) blockDim:(3,1,1)  gridDim(2,1,1)
# threadIdx:(1,0,0) blockIdx:(1,0,0) blockDim:(3,1,1)  gridDim(2,1,1)
# threadIdx:(2,0,0) blockIdx:(1,0,0) blockDim:(3,1,1)  gridDim(2,1,1)
# threadIdx:(0,0,0) blockIdx:(0,0,0) blockDim:(3,1,1)  gridDim(2,1,1)
# threadIdx:(1,0,0) blockIdx:(0,0,0) blockDim:(3,1,1)  gridDim(2,1,1)
# threadIdx:(2,0,0) blockIdx:(0,0,0) blockDim:(3,1,1)  gridDim(2,1,1)

3. CUDA 向量加法

代码

#include "freshman.h"
#include <cuda_runtime.h>
#include <stdio.h>
__host__ void sumArrays(float *a, float *b, float *res, const int size) {
  for (int i = 0; i < size; i++) {
    res[i] = a[i] + b[i];
  }
}
__global__ void sumArraysGPU(float *a, float *b, float *res, const int size) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < size) // 线程索引越界检查
    res[i] = a[i] + b[i];
}
int main(int argc, char **argv) {
  // set up device
  initDevice(0);
  // allocate host memory
  int nElem = 1 << 24;
  printf("Vector size:%d\n", nElem);
  int nByte = sizeof(float) * nElem;
  float *a_h = (float *)malloc(nByte);
  float *b_h = (float *)malloc(nByte);
  float *res_h = (float *)malloc(nByte);
  float *res_from_gpu_h = (float *)malloc(nByte);
  memset(res_h, 0, nByte);
  memset(res_from_gpu_h, 0, nByte);
  // allocate device memory
  float *a_d, *b_d, *res_d;
  CHECK(cudaMalloc((float **)&a_d, nByte));
  CHECK(cudaMalloc((float **)&b_d, nByte));
  CHECK(cudaMalloc((float **)&res_d, nByte));
  // randomly initialize the input data
  initialData(a_h, nElem);
  initialData(b_h, nElem);
  CHECK(cudaMemcpy(a_d, a_h, nByte, cudaMemcpyHostToDevice));
  CHECK(cudaMemcpy(b_d, b_h, nByte, cudaMemcpyHostToDevice));
  // set up execution configuration
  // 1. 计算最佳块大小
  int minGridSize, bestBlockSize;
  cudaOccupancyMaxPotentialBlockSize(&minGridSize, &bestBlockSize,
                                     (void *)sumArraysGPU,
                                     0, // 动态共享内存大小
                                     0  // 无块大小限制
  );
  // 2. 设置网格和块维度
  dim3 block(bestBlockSize);
  dim3 grid((nElem + bestBlockSize - 1) / bestBlockSize);
  // 3. 设备执行并统计耗时
  double iStart, iElaps;
  iStart = cpuSecond();
  sumArraysGPU<<<grid, block>>>(a_d, b_d, res_d, nElem);
  iElaps = cpuSecond() - iStart;
  CHECK(cudaGetLastError());
  CHECK(cudaDeviceSynchronize());
  CHECK(cudaMemcpy(res_from_gpu_h, res_d, nByte, cudaMemcpyDeviceToHost));
  printf("Execution configuration<<<%d,%d>>> Time elapsed %f sec\n", grid.x,
         block.x, iElaps);
  // 4. CPU执行并统计耗时
  iStart = cpuSecond();
  sumArrays(a_h, b_h, res_h, nElem);
  iElaps = cpuSecond() - iStart;
  printf("CPU Time elapsed %f sec\n", iElaps);
  // 5. 检查结果
  checkResult(res_h, res_from_gpu_h, nElem);
  // 6. 释放内存
  cudaFree(a_d);
  cudaFree(b_d);
  cudaFree(res_d);
  free(a_h);
  free(b_h);
  free(res_h);
  free(res_from_gpu_h);
  return 0;
}

cudaOccupancyMaxPotentialBlockSize 函数用于计算最佳块大小
max thread per block 是 cuda 中的一个限制，表示每个块中最多可以有多少个线程，一般为 1024，当超过这个限制时，CHECK(cudaGetLastError()); 会报错

运行结果

./sum_arrays
# Using device 0: NVIDIA GeForce RTX 4060 Laptop GPU
# Vector size:16777216
# Execution configuration<<<21846,768>>> Time elapsed 0.000030 sec
# CPU Time elapsed 0.076604 sec
# Check result success!

可以看出，cuda 的执行时间远远小于 cpu 的执行时间，相差了 2553 倍

总体流程

graph TD
    A[Host 端申请内存] --> B[Host 端初始化输入数据]
    B --> C[Device 端申请内存]
    C --> D[拷贝 Host 输入数据到 Device 端]
    D --> E[Device 端执行核函数]
    E --> F[拷贝 Device 端输出数据到 Host 端]
    F --> G[Host 端检查结果]
    G --> H[释放 Host 端和 Device 端全部内存]
    B --> J[Host 端执行普通函数] --> G

细节分析

1. `cuda` 内存分配是怎么做的？

1 2	float a_d; // 空指针 cudaMalloc((float *)&a_d, nByte); // 将指针的地址转成二级指针（指针的指针）传入内存分配函数

这里的二级指针应用很巧妙，由于 c++ 中的指针是值传递，所以如果是一级指针传入 cudaMalloc 函数时，指针 a_d 的值不会改变，因此只能将指针的地址转成二级指针传入内存分配函数

2. `cuda` 内存拷贝是怎么做的？

1	cudaMemcpy(a_d, a_h, nByte, cudaMemcpyHostToDevice);

cudaMemcpy 函数的四个参数分别是：
- dst：目标地址
- src：源地址
- size：拷贝的字节数
- kind：拷贝的类型，属于 cudaMemcpyKind 枚举类型
  - cudaMemcpyHostToHost = 0, /**< Host -> Host */
  - cudaMemcpyHostToDevice = 1, /**< Host -> Device */
  - cudaMemcpyDeviceToHost = 2, /**< Device -> Host */
  - cudaMemcpyDeviceToDevice = 3, /**< Device -> Device */
  - cudaMemcpyDefault = 4 /**< Direction of the transfer is inferred from the pointer values. Requires unified virtual addressing */

3. `cuda` 核函数如何解析线程索引？

1
2
3

int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < size) // 线程索引越界检查
  res[i] = a[i] + b[i];

blockIdx：表示当前块的索引
blockDim：表示当前块的维度（每个块中的线程数）
threadIdx：表示当前线程的索引
每个线程中计算两个标量的和
由于 gridDim * blockDim 可能大于 size，所以需要判断线程索引是否越界

4. 如何计算最佳块大小？

int minGridSize, bestBlockSize;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &bestBlockSize,
                                     (void *)sumArraysGPU,
                                     0, // 动态共享内存大小
                                     0  // 无块大小限制
);

cudaOccupancyMaxPotentialBlockSize 函数用于计算最佳块大小，五个参数分别是：
- minGridSize：最小网格大小变量地址
- bestBlockSize：最佳块大小变量地址
- kernel：核函数指针
- dynamicSMemSize：动态共享内存大小
- blockSizeLimit：块大小限制
函数名就是函数地址，可强转为 void * 函数指针（也可以写成：(void *)&sumArraysGPU）

4. CUDA 编程模型

线程块

线程块 block 是 CUDA 中的逻辑执行单元，是一个三维逻辑结构：
- block.x：表示块的 x 维度大小
- block.y：表示块的 y 维度大小
- block.z：表示块的 z 维度大小
- 其中 block.x 是最内层的循环，block.y 是第二层循环，block.z 是最外层的循环
- 用三维数组可以表示为：tread[z][y][x]，即 tid = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y

线程束

线程束 warp 是 CUDA 基本调度执行单元，一个 warp 由 32 个线程组成
- 一个 warp 中的线程在一个时钟周期内执行同一条指令（单指令多线程，SIMT）
- 一个 warp 中的线程可以共享指令指针和执行资源（如寄存器、缓存等）
- Warp 调度器（warp scheduler）负责将 warp 分配到物理执行单元上执行
线程块会被划分为多个 warp， $WarpsPerBlock=ceil(\frac{ThreadsPerBlock}{32})$

CUDA core

CUDA core 是 CUDA 物理执行单元，负责实际的计算任务
一个 CUDA core 一个时钟周期只能计算一个线程的指令

StreamMultiprocessor

StreamMultiprocessor 流式多处理器（简称 SM），负责执行 CUDA 线程块中的并行计算任务
每个 GPU 包含多个 SM，每个 SM 包含多个 CUDA core，例如：RTX 4060 有 24 个 SM，每个 SM 有 128 个 CUDA core

5. Reduce

规约（Reduce）是 CUDA 编程中常见的操作，主要用于将多个数据元素规约为一个数据元素
规约操作通常是一个二元操作，例如：sum、mul、max、min 等，简单的规约可以合并成强大的算子，甚至可以说规约算子是神经网络的基础

规约求和

// CPU 规约求和
int recursiveReduce(int *data, int const size) {
  // terminate check
  if (size == 1)
    return data[0];
  // renew the stride
  int const stride = size / 2;
  if (size % 2 == 1) {
    for (int i = 0; i < stride; i++) {
      data[i] += data[i + stride];
    }
    data[0] += data[size - 1];
  } else {
    for (int i = 0; i < stride; i++) {
      data[i] += data[i + stride];
    }
  }
  // call
  return recursiveReduce(data, stride);
}
// GPU 规约相邻求和
__global__ void reduceNeighbored(int *g_idata, int *g_odata, unsigned int n) {
  // set thread ID
  unsigned int tid = threadIdx.x;
  // boundary check
  if (tid >= n)
    return;
  // convert global data pointer to local point of this block
  int *idata = g_idata + blockIdx.x * blockDim.x;
  // in-place reduction in global memory
  for (int stride = 1; stride < blockDim.x; stride *= 2) {
    if ((tid % (2 * stride)) == 0) {
      idata[tid] += idata[tid + stride];
    }
    // synchronize within block
    __syncthreads();
  }
  // write result for this block to global mem
  if (tid == 0)
    g_odata[blockIdx.x] = idata[0];
}
// GPU 规约相邻求和（简化版）
__global__ void reduceNeighboredLess(int *g_idata, int *g_odata,
                                     unsigned int n) {
  unsigned int tid = threadIdx.x;
  unsigned idx = blockIdx.x * blockDim.x + threadIdx.x;
  // convert global data pointer to the local point of this block
  int *idata = g_idata + blockIdx.x * blockDim.x;
  if (idx > n)
    return;
  // in-place reduction in global memory
  for (int stride = 1; stride < blockDim.x; stride *= 2) {
    // convert tid into local array index
    int index = 2 * stride * tid;
    if (index < blockDim.x) {
      idata[index] += idata[index + stride];
    }
    __syncthreads();
  }
  // write result for this block to global men
  if (tid == 0)
    g_odata[blockIdx.x] = idata[0];
}
// GPU 规约交错求和，主要是 stride 的计算方式不同
__global__ void reduceInterleaved(int *g_idata, int *g_odata, unsigned int n) {
  unsigned int tid = threadIdx.x;
  unsigned idx = blockIdx.x * blockDim.x + threadIdx.x;
  // convert global data pointer to the local point of this block
  int *idata = g_idata + blockIdx.x * blockDim.x;
  if (idx >= n)
    return;
  // in-place reduction in global memory
  for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
    if (tid < stride) {
      idata[tid] += idata[tid + stride];
    }
    __syncthreads();
  }
  // write result for this block to global men
  if (tid == 0)
    g_odata[blockIdx.x] = idata[0];
}

效率分析

本代码中使用了三种核函数实现方式做同一个规约操作，分别是：
- reduceNeighbored：相邻线程规约
- reduceNeighboredLess：相邻线程规约（简化版）
- reduceInterleaved：交错线程规约
三者效率从高到低依次是：
- reduceInterleaved > reduceNeighboredLess > reduceNeighbored
三者的示意图分别如下：
- reduceNeighbored：相邻线程规约的实现
- reduceNeighboredLess：相邻线程规约的简化版实现（注意圆圈中的符号已和上图不一致）
- reduceInterleaved：交错线程规约的实现
三者效率差异主要来自于 线程分支分化，后续会详细介绍

6. 循环展开

循环展开（Loop Unrolling）是 CUDA 中常用的优化手段，主要用于减少循环控制开销和提高指令级并行度
简单说就是一个线程不再只计算一个数据，而是计算多个数据，而且是直接在代码中展开，而不是在编译器中展开
可以简单理解成：启动线程是需要花时间的，启动一个线程只算一个数据，太浪费了，所以我们可以让一个线程计算多个数据，这样就能减少启动线程的时间开销，所以就省时间了

代码

// 数据总长是 8 * blockDim.x * gridDim.x
// 线程数是 blockDim.x * gridDim.x
// 每个线程计算 8 个数据
__global__ void reduceUnroll8(int *g_idata, int *g_odata, unsigned int n) {
  unsigned int tid = threadIdx.x;
  unsigned int idx = blockDim.x * blockIdx.x * 8 + threadIdx.x;
  if (tid >= n)
    return;
  int *idata = g_idata + blockIdx.x * blockDim.x * 8;
  // 循环展开，每个线程计算 8 个数据
  // 直接把 8 * blockDim.x * gridDim.x 的数据总长
  // 聚合到了 blockDim.x * gridDim.x 的线程数上
  if (idx + 7 * blockDim.x < n) {
    g_idata[idx] += g_idata[idx + blockDim.x];
    g_idata[idx] += g_idata[idx + blockDim.x * 2];
    g_idata[idx] += g_idata[idx + blockDim.x * 3];
    g_idata[idx] += g_idata[idx + blockDim.x * 4];
    g_idata[idx] += g_idata[idx + blockDim.x * 5];
    g_idata[idx] += g_idata[idx + blockDim.x * 6];
    g_idata[idx] += g_idata[idx + blockDim.x * 7];
  }
  // 这里需要同步，也就是线程阻塞直到所有线程都执行完
  __syncthreads();
  // 然后就是一个最简单的规约操作了，和上面一样
  for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
    if (tid < stride) {
      idata[tid] += idata[tid + stride];
    }
    __syncthreads();
  }
  // write result for this block to global mem
  if (tid == 0)
    g_odata[blockIdx.x] = idata[0];
}

7. 核函数递归调用

和 CPU 一样，CUDA 也支持核函数递归调用，调用方式和普通递归函数一样
需要注意的是在编译的时候需要加上 -rdc=true 选项

代码

__global__ void nesthelloworld(int iSize, int iDepth) {
  unsigned int tid = threadIdx.x;
  printf("depth : %d blockIdx: %d,threadIdx: %d\n", iDepth, blockIdx.x,
         threadIdx.x);
  if (iSize == 1)
    return;
  int nthread = (iSize >> 1);
  if (tid == 0 && nthread > 0) {
    // 递归调用核函数
    nesthelloworld<<<1, nthread>>>(nthread, ++iDepth);
    printf("-----------> nested execution depth: %d\n", iDepth);
  }
}

8. 固定内存

Pinned Memory 是 CUDA 中的一种特殊内存类型（不是显存，是内存），主要用于提高数据传输效率
普通内存是分页管理，分页管理存在两个问题：
1. 一页内存逻辑上连续，但物理上不连续
2. 操作系统可能会将内存页交换到磁盘上，导致数据不在物理内存中
Pinned Memory 就是解决了这两个问题，分配了一块连续物理地址且固定的主机内存（host 内存），方便整块拷贝数据到显存（DMA）

代码

#include "freshman.h"
#include <cuda_runtime.h>
int main(int argc, char **argv) {
  int dev = 0;
  cudaSetDevice(dev);
  int nElem = 1 << 14;
  int nByte = sizeof(float) * nElem;
  float *a_h, *b_h, *res_h, *res_from_gpu_h;
  // 注意这里的 cudaMallocHost 和 cudaMalloc 是不同的
  // 前者申请的是 host 固定内存，后者申请的是 device 显存
  // cudaMallocHost 是 malloc 的一个平替
  CHECK(cudaMallocHost((float **)&a_h, nByte));
  CHECK(cudaMallocHost((float **)&b_h, nByte));
  CHECK(cudaMallocHost((float **)&res_h, nByte));
  CHECK(cudaMallocHost((float **)&res_from_gpu_h, nByte));
  // 初始化数据
  memset(res_h, 0, nByte);
  memset(res_from_gpu_h, 0, nByte);
  initialData(a_h, nElem);
  initialData(b_h, nElem);
  // 申请设备显存
  float *a_d, *b_d, *res_d;
  CHECK(cudaMalloc((float **)&a_d, nByte));
  CHECK(cudaMalloc((float **)&b_d, nByte));
  CHECK(cudaMalloc((float **)&res_d, nByte));
  // 拷贝数据到设备显存
  CHECK(cudaMemcpy(a_d, a_h, nByte, cudaMemcpyHostToDevice));
  CHECK(cudaMemcpy(b_d, b_h, nByte, cudaMemcpyHostToDevice));
  // 跑核函数
  dim3 block(1024);
  dim3 grid(nElem / block.x);
  sumArraysGPU<<<grid, block>>>(a_d, b_d, res_d);
  printf("Execution configuration<<<%d,%d>>>\n", grid.x, block.x);
  // 结果拷贝回主机并检查两个设备计算结果是否一致
  CHECK(cudaMemcpy(res_from_gpu_h, res_d, nByte, cudaMemcpyDeviceToHost));
  sumArrays(a_h, b_h, res_h, nElem);
  checkResult(res_h, res_from_gpu_h, nElem);
  // 释放内存
  // 注意这里的 cudaFreeHost 和 cudaFree 是不同的
  // 前者释放的是 host 固定内存，后者释放的是 device 显存
  cudaFree(a_d);
  cudaFree(b_d);
  cudaFree(res_d);
  cudaFreeHost(a_h);
  cudaFreeHost(b_h);
  cudaFreeHost(res_h);
  cudaFreeHost(res_from_gpu_h);
  return 0;
}

关键点

主存 普通内存 的分配和释放函数是 malloc 和 free
主存 固定内存 的分配和释放函数是 cudaMallocHost 和 cudaFreeHost
显存的分配和释放函数是 cudaMalloc 和 cudaFree

9. 零拷贝内存和统一虚拟地址

Zero-Copy Memory 是 CUDA 中一种允许 GPU 直接访问主机内存的技术，避免了显式的数据拷贝操作（不需要用 cudaMemcpy 函数）
实际上，Zero-Copy Memory 在很多时候并不快，因为 GPU 访问主机内存的速度远远低于访问显存的速度，因此，Zero-Copy Memory 只适用于一些特殊的场景，例如：
- 主机内存中的数据只需要被 GPU 使用一次
- 数据量太大，显存放不下
- 调试用途
Zero-Copy Memory 的实现方式是将主机内存映射到 GPU 的地址空间中，GPU 通过访问这个地址空间来访问主机内存，实际上走的是 PCIe 总线
由于不需要先完成所有数据的拷贝再开始执行核函数，因此 Zero-Copy Memory 使用异步拷贝的方式来实现，可将部分拷贝数据的时间和核函数执行的时间重叠，但并不多
Unified Virtual Addressing (UVA) 是 CUDA 中的一种内存管理机制，允许 CPU 和 GPU 共享同一虚拟地址空间

代码

float *a_host, *b_host, *res_d;
// 申请主机固定内存，添加特殊 flag cudaHostAllocMapped
CHECK(cudaMallocHost((float **)&a_host, nByte, cudaHostAllocMapped));
CHECK(cudaMallocHost((float **)&b_host, nByte, cudaHostAllocMapped));
// a_host 和 b_host 是可直接作为核函数的输入参数
// 也可以转成 Device 地址空间，如下：
float *a_dev, *b_dev;
// 映射主机内存到设备地址空间
CHECK(cudaHostGetDevicePointer((void **)&a_dev, (void *)a_host, 0));
CHECK(cudaHostGetDevicePointer((void **)&b_dev, (void *)b_host, 0));
// 用 a_dev 和 b_dev 作为核函数的输入参数计算

和 Zero-Copy Memory 一样，UVA 甚至不需要将 a_host 转成 a_dev，直接用 a_host 就可以调用核函数

10. Aos 和 SoA

CUDA 不仅支持最简单的原生数据类型，还支持自定义数据类型（struct），例如 Aos 和 SoA 等
Aos（Array of Structures）和 SoA（Structure of Arrays）是两种不同的数据存储方式，这两种方式由于变量的排布方式不同，导致了访问内存的效率差异

Aos

Aos 是将多个结构体存储在一个数组中，每个结构体的成员变量是连续存储的
例如：

struct AoSStruct {
  float a;
  float b;
};
__global__ void sumArraysGPU(float *a, float *b, struct naiveStruct *res, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n)
    res[i].a = a[i] + b[i];
}

这里的 res 是一个 结构体数组，a 和 b 是结构体的成员变量，每个变量都是一个 float 类型的标量
每个结构体的成员变量是连续存储的，即：a1 b1 a2 b2 a3 b3 ...

SoA

SoA 是将多个结构体的成员变量存储在一个数组中，每个成员变量是连续存储的
例如：

struct SoAStruct {
  float a[SIZE];
  float b[SIZE];
};
__global__ void sumArraysGPU(float *a, float *b, struct SoAStruct *res, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n)
    res->a[i] = a[i] + b[i];
}

这里的 res 是一个 结构体，a 和 b 是结构体的成员变量，每个变量都是一个 float 类型的数组
每个成员变量的数组是连续存储的，即：a1 a2 a3 ... b1 b2 b3 ...

11. 行主序和列主序

行主序（Row Major Order）和 列主序（Column Major Order）是两种不同的数组存储方式
行主序 是将数组的每一行存储在连续的内存中，列主序 是将数组的每一列存储在连续的内存中
例如：

int a[3][4] = {
    {1, 2, 3, 4},
    {5, 6, 7, 8},
    {9, 10, 11, 12}};

行主序 存储方式是：1 2 3 4 5 6 7 8 9 10 11 12
列主序 存储方式是：1 5 9 2 6 10 3 7 11 4 8 12
默认情况下，C / C++ / CUDA 语言是 行主序 存储方式
在行主序存储下，如果按行序访问数组元素，访问效率会更高，因为连续的内存访问会提高缓存命中率，反之如果按列序访问数组元素，访问效率会更低

速度对比

WIP

Transformers without Normalization

Posted on 2025-03-17 Edited on 2025-03-21 In Transformer Valine:

URL

paper: https://arxiv.org/pdf/2503.10622
code: https://github.com/jiachenzhu/DyT

TL;DR

这是由恺明和杨立昆提出的一篇关于 transformer 算子优化的论文，主要观点是去掉 transformer 结构中的 normalization 层，改成 tanh 层
改用 tanh 算子的 transformer 模型，在大多数任务上可达到使用归一化层的模型相同的性能，甚至更好

Algorithm

简单来说，这篇论文的核心思想是将 transformer 中的 normalization 层（可以是 LayerNorm 或 RMSNorm）替换成 dynamic tanh 层（简称 DyT）
normalization 计算公式：

\text{normalization}(x) = \gamma \times \frac{x - \mu}{\sqrt{\sigma^2+\epsilon}} + \beta

其中 $\mu$ 和 $\sigma$ 分别是 mean 和 std， $\gamma$ 和 $\beta$ 是 scale 和 shift 参数

DyT 计算公式：

\text{DyT}(x) = \gamma \times \tanh(\alpha x) + \beta

其中 $\alpha$ 是个可学习参数， $\gamma$ 和 $\beta$ 是 scale 和 shift 参数（和 normalization 一样）

DyT 实现伪代码：

# input x has the shape of [B, T, C]
# B: batch size, T: tokens, C: dimension
class DyT(Module):
    def __init__(self, C, init_α):
        super().__init__()
        self.α = Parameter(ones(1) * init_α)
        self.γ = Parameter(ones(C))
        self.β = Parameter(zeros(C))
    def forward(self, x):
        x = tanh(self.alpha * x)
        return self.γ * x + self.β

$\alpha$ 默认初始化值为 0.5

Results

作者在多个领域的知名模型上都对比了修改前后训练精度，DyT 的性能和 normalization 的性能基本一致，打的有来有回
作者还对比了 DyT 和 normalization 的训练/推理速度，DyT 的训练/推理速度要快很多
作者同时做 tanh 和 $\alpha$ 做了消融实验，发现 tanh 和 $\alpha$ 都是必要的

Thoughts

属于是恺明和立昆的梦幻联动了…，这种对最火的结构的优化，非大佬不能为也，想象下如果这篇论文是大学实验室发表的，大家第一反应恐怕是：Who think you are? 😂
之前算是稍微接触过硬件，DyT 这种 element-wise op 比 normalization 这种 reduce op 一定快多了，想怎么 tiling 都行…

2025.02 DeepSeek 开源周第五弹 —— 3FS

Posted on 2025-03-14 Edited on 2025-06-04 In file-system Valine:

URL

code: https://github.com/deepseek-ai/3FS

TL;DR

3FS (Fire-Flyer File System) 是一个高性能的分布式文件系统，旨在提供低延迟和高吞吐量的存储解决方案，利用现代 SSD 和 RDMA 网络带全宽的并行文件系统，解决 AI 训练和推理存储问题

Algorithm

效果

集群高吞吐：在 180 节点集群中，3FS 实现了高达 6.6 TiB/s 的聚合读取吞吐量
基准测试优异：在 25 节点集群的 GraySort 基准测试中，3FS 达到了 3.66 TiB /min 的吞吐量
单节点高性能：每个客户端节点的 KVCache 查找峰值吞吐量超过 40 GiB/s
架构先进： 3FS 采用去中心化架构，并具备强一致性语义

系统介绍

还是看大佬讲吧，传送门

Thoughts

这种底层架构一般只有大厂可以做，deepseek 有点东西

URL

TL;DR

Algorithm

ToolFormer 的意义

ToolFormer 的工作方式

应该怎么得到 ToolFormer

如何构造工具调用数据集

Thoughts

URL

TL;DR

Algorithm

ReAct 框架

核心贡献​​

Thoughts

URL

TL;DR

论文详情

1. 核心方法：思维链提示（Chain-of-Thought Prompting）

2. 关键优势

3. 实验验证

4. 局限性与启示

论文核心贡献

Thoughts

URL

TL;DR

Algorithm

总体流程

模型架构

训练过程

推理与生成

和自回归模型对比

Thought

URL

TL;DR

Algorithm

总体流程

Conditioning UNet

Thoughts

URL

TL;DR

Algorithm

DDPM 逆向过程公式

DDIM 逆向过程公式

DDIM 公式拆解：

1. 预测原输入

2. 计算调整方向

3. 更新隐变量

DDPM 和 DDIM 使用上的区别

二者对比分析

Thoughts

URL

TL;DR

Qwen3 系列模型

Qwen3 模型的主要特性

预训练

后训练

Qwen3 使用

Thoughts

0. 学习 CUDA 的目的

1. Hello World

cuda 代码

编译

运行

2. dimension 测试

3. CUDA 向量加法

代码

运行结果

总体流程

细节分析

1. cuda 内存分配是怎么做的？

2. cuda 内存拷贝是怎么做的？

3. cuda 核函数如何解析线程索引？

4. 如何计算最佳块大小？

4. CUDA 编程模型

线程块

线程束

CUDA core

StreamMultiprocessor

5. Reduce

规约求和

核心贡献

`DDIM` 公式拆解：

1. `cuda` 内存分配是怎么做的？

2. `cuda` 内存拷贝是怎么做的？

3. `cuda` 核函数如何解析线程索引？

9. 零拷贝内存和统一虚拟地址