GELU（Gaussian Error Linear Unit）是 BERT 的灵魂激活函数，后来被 GPT-2/3 沿用。两种实现：精确版（调用 erf，慢但数学精确）和 tanh 近似版（快但误差 ~0.1%）。BERT 的训练耗时分析：GELU 占用了 11% 的前向时间——如果换成 tanh 近似版，降到 4%。

差距在哪？精确版要算 Φ(x)（标准正态累积分布函数），Φ 内部是 erf(x/√2)——erf 在 NPU 上没有硬件指令，靠多项式展开。

精确版：多项式展开 erf

GELU(x) = x × Φ(x) = x × ½ × (1 + erf(x/√2))

erf(z) = 2/√π × ∫₀ᶻ exp(-t²) dt

erf 没有硬件支持——靠 7 次多项式展开：

erf(z) ≈ 1 - (a₁×t + a₂×t² + ... + a₅×t⁵) × exp(-z²)
其中 t = 1/(1 + p×|z|),  p=0.3275911
a₁=0.254829592, a₂=-0.284496736, a₃=1.421413741,
a₄=-1.453152027, a₅=1.061405429

// ops-nn/kernels/gelu/gelu_exact.cpp

__aicore__ void GELUExactKernel(
    GlobalTensor<float16>& x,      // [N] 输入
    GlobalTensor<float16>& y,      // [N] 输出
    int N
) {
    const float INV_SQRT2 = 0.7071067811865475f;

    // 多项式系数
    const float P = 0.3275911f;
    const float A1 = 0.254829592f;
    const float A2 = -0.284496736f;
    const float A3 = 1.421413741f;
    const float A4 = -1.453152027f;
    const float A5 = 1.061405429f;

    for (int i = threadIdx.x; i < N; i += 256) {
        float val = float(x[i]);

        // z = x / sqrt(2)
        float z = val * INV_SQRT2;
        float abs_z = fabsf(z);

        // t = 1 / (1 + p * |z|)
        float t = 1.0f / (1.0f + P * abs_z);

        // Horner 方法计算多项式（5 次，6 次 FMA）
        // a₁t + a₂t² + a₃t³ + a₄t⁴ + a₅t⁵
        float poly = A5;
        poly = poly * t + A4;   // a₅t + a₄
        poly = poly * t + A3;   // a₅t² + a₄t + a₃
        poly = poly * t + A2;   // a₅t³ + a₄t² + a₃t + a₂
        poly = poly * t + A1;   // a₅t⁴ + a₄t³ + a₃t² + a₂t + a₁

        // erf(|z|) ≈ 1 - poly * exp(-z²)
        float erf_abs = 1.0f - poly * expf(-abs_z * abs_z);

        // erf(z) = sign(z) × erf(|z|)
        float erf_val = (z >= 0.0f) ? erf_abs : -erf_abs;

        // Φ(x) = ½ × (1 + erf(x/√2))
        float phi = 0.5f * (1.0f + erf_val);

        // GELU(x) = x × Φ(x)
        y[i] = float16(val * phi);
    }
}

Horner 展开用了 6 次 FMA，加上 expf 和 2 次乘法 → 总共 ~14 次浮点操作。不算慢，但相比 tanh 近似版还是多了不少。

tanh 近似版：4 次 FMA

GELU(x) ≈ 0.5 × x × (1 + tanh(√(2/π) × (x + 0.044715 × x³)))

// ops-nn/kernels/gelu/gelu_tanh_approx.cpp

__aicore__ void GELUTanhKernel(
    GlobalTensor<float16>& x,
    GlobalTensor<float16>& y,
    int N
) {
    const float SQRT_2_OVER_PI = 0.7978845608028654f;
    const float COEFF = 0.044715f;

    for (int i = threadIdx.x; i < N; i += 256) {
        float val = float(x[i]);

        // inner = √(2/π) × (x + 0.044715 × x³)
        float x2 = val * val;
        float x3 = x2 * val;
        float inner = SQRT_2_OVER_PI * (val + COEFF * x3);

        // tanh(inner)
        float tanh_val = tanhf(inner);

        // GELU(x) = 0.5 × x × (1 + tanh(inner))
        y[i] = float16(0.5f * val * (1.0f + tanh_val));
    }
}

4 次 FMA + 1 次 tanhf → ~8 次浮点操作。tanhf 在 Ascend NPU 上有硬件支持（Vector 单元内置），一个周期完成。比 exact 版快 ~3×。

性能对比

Ascend 910 NPU，FP16，N=4096×4096

| 实现           | BF16 延迟 | 最大误差 | LLaMA 7B 训练 loss |
|---------------|-----------|---------|-------------------|
| 精确版（erf）  | 47.2 μs   | 0       | 1.8543 (基线)     |
| tanh 近似版    | 15.8 μs   | 1.2e-3  | 1.8544 (+0.0001)  |
| 加速比         | 2.99×     | —       | —                 |

BERT-base（12 层 × hidden=768）：
  精确版：12 × 47.2 = 566 μs/layer
  tanh版： 12 × 15.8 = 190 μs/layer
  省 376 μs/layer → 1M steps × 12 layers = 4.5 秒省（单卡）

loss 差异 0.0001——在训练误差范围内，对收敛无影响。

tanh 近似版的误差分布

x ∈ [-3, 3]：误差 < 2e-4（主要使用范围，误差极小）
x ∈ [-5, -3] ∪ [3, 5]：误差 ~5e-4（激活值的边缘区域）
x < -5 或 x > 5：误差 ~1.2e-3（饱和区，GELU ≈ 0 或 x）

训练中 99.7% 的激活值在 [-3, 3] 内 → 实际误差 < 2e-4

反向传播

精确版的反向：

GELU'(x) = Φ(x) + x × φ(x)
其中 φ(x) = (1/√(2π)) × exp(-x²/2)  # 标准正态密度函数

tanh 近似版的反向（同样用近似）：

GELU'(x) ≈ 0.5 × (1 + tanh(T)) + 0.5 × x × (1 - tanh²(T)) × √(2/π) × (1 + 3×0.044715×x²)
其中 T = √(2/π) × (x + 0.044715 × x³)

// ops-nn/kernels/gelu/gelu_tanh_backward.cpp

__aicore__ void GELUTanhBackwardKernel(
    GlobalTensor<float16>& x,        // [N] 前向输入
    GlobalTensor<float16>& dy,       // [N] 上游梯度
    GlobalTensor<float16>& dx,       // [N] 输出梯度
    int N
) {
    const float SQRT_2_OVER_PI = 0.7978845608028654f;
    const float COEFF = 0.044715f;

    for (int i = threadIdx.x; i < N; i += 256) {
        float val = float(x[i]);
        float grad_in = float(dy[i]);

        float x2 = val * val;
        float x3 = x2 * val;
        float T = SQRT_2_OVER_PI * (val + COEFF * x3);

        float tanh_T = tanhf(T);
        float sech2_T = 1.0f - tanh_T * tanh_T;  // sech² = 1 - tanh²

        float dT_dx = SQRT_2_OVER_PI * (1.0f + 3.0f * COEFF * x2);

        // GELU'(x) = 0.5 × (1 + tanh(T)) + 0.5 × x × sech²(T) × dT_dx
        float gelu_grad = 0.5f * (1.0f + tanh_T) +
                          0.5f * val * sech2_T * dT_dx;

        dx[i] = float16(grad_in * gelu_grad);
    }
}

踩坑一：x³ 在 FP16 下溢出

tanh 近似版中 x³ = x × x × x——如果 x = 10（FP16 下合法的 logit 值），x³ = 1000——FP16 最大值 65504，不溢出。但如果模型训练不稳定，某个 step 的 logit 飙到 40→ x³ = 64000 > 65504→溢出。

// ❌ FP16 下直接算 x³ → 溢出风险
float16 x3 = x * x * x;  // x=40.0 → x³=64000 → 溢出 = inf

// ✅ 内部先用 FP32 算，最后转 FP16
float x3 = float(val) * float(val) * float(val);  // FP32 范围 3.4e38 → 安全
float inner = SQRT_2_OVER_PI * (float(val) + COEFF * x3);  // 全 FP32
y[i] = float16(0.5f * val * (1.0f + tanhf(inner)));  // 最后才转 FP16

踩坑二：tanhf 的 FP16 实现在负半轴不精确

Ascend NPU 的 tanhf 硬件实现是为 FP32 设计的。FP16 输入下，负半轴 (x < -3) 的 tanhf 误差 ~5e-4。GELU 的负半轴 (x < -2) 的激活值接近 0，一个小误差会被 x 放大。

x = -4.0: 精确 GELU(−4) = −4 × Φ(−4) ≈ −4 × 3.2e-5 ≈ -1.3e-4
tanh 近似: 0.5 × (−4) × (1 + tanh(−3.21)) = −2 × (1 − 0.9969) = −2 × 0.0031 = −6.2e-3
偏差: 6.2e-3 vs 1.3e-4 → 48×

x = -6.0: 精确 GELU(−6) ≈ −6 × 9.9e-10 ≈ −5.9e-9
tanh 近似: −3 × (1 − 0.999998) = −3 × 2e-6 = −6e-6
偏差: 1000×（但绝对值都接近 0）

实际情况：训练中 x < -3 的激活值占比 < 0.3%，对总 loss 影响 < 1e-6。不是实用层面的问题，但数值特性值得了解。

踩坑三：推理时还在用精确版

训练中 tanh 近似版 loss 和精确版一样，但很多推理代码还是用精确版——没改预处理管线。

# ❌ 推理时用了精确版（从训练 checkpoint 改过来的）
model.layers[i].activation = GELUExact()

# ✅ 用 tanh 近似版——loss 相同，延迟降 3×
model.layers[i].activation = GELUTanhApprox()

BERT-base 推理：精确版 activation 占 11% 延迟 → tanh 版占 4%。batch=1 下从 12ms 降到 11.2ms（省 0.8ms）。1M 次推理省 800 秒。

---

GELU 的 tanh 近似版误差 < 0.1%，对训练 loss 无影响，但延迟省 3×。精确版的 erf 多项式展开（14 次浮点操作）vs tanh 版（4 次 FMA + 硬件 tanhf）。决定很简单：训练和推理都用 tanh 近似版。唯一的坑：内部计算全用 FP32（防 x³ 溢出），最后才转 FP16。