(23) Philox可复现随机数与确定性训练

——“一个人用AI如何写出比PyTorch更快的自研深度学习框架”系列文章之二十三

训练一个深度学习模型时，随机数看起来只是背景噪音：权重初始化、数据增强、Dropout 掩码、数据集洗牌、BatchNorm 的某些实现……它们无处不在，却很少被放在聚光灯下。然而，当你真正需要一个可复现的训练结果时——比如调试一个诡异的收敛问题、做消融实验、或者向审稿人证明“这个改动确实有效”——随机数立刻从幕后走到台前，成为决定成败的关键。

可复现性这件事，说起来简单：固定种子，不就能得到固定结果吗？但真正在生产级框架里把它做稳，远比想象中复杂。多线程数据加载会打乱采样顺序，多卡并行的原子加法会改变浮点求和顺序，CUDA Graph 里的随机状态不能每次从 CPU 重新喂，而传统基于状态递推的伪随机数生成器（如 Mersenne Twister）一旦遇到并行执行顺序变化，整个序列就会错位。

这篇文章就来聊聊 Tech-Renaissance 如何把Philox 计数器型随机数生成器贯穿到训练全流程，并以此为基础构建端到端的确定性训练能力。

一、为什么深度学习训练的可复现性这么难？

随机性在深度学习里大致有三类用途：

1. 初始化：权重、偏置的随机初始化决定优化的起点；
2. 数据侧：数据集洗牌、随机裁剪、随机翻转、颜色抖动、Random Erasing 等增强；
3. 模型侧：Dropout、某些采样或扰动操作。

要让这些环节全部可复现，核心不是“有没有随机数”，而是随机数的消费顺序必须完全一致。传统 PRNG（如 C++ std::mt19937、NumPy 默认的 PCG64）都是“状态机”：每生成一个数就推进一次内部状态。如果两个线程同时请求随机数，谁先谁后会改变后续所有人拿到的序列；如果第 3 个 epoch 的 worker 数量变了，全局状态也会不同。

1.1 状态机式 RNG 的级联故障

状态机式 RNG 最大的麻烦在于“级联偏移”。假设有两个预处理 worker，共享同一个 Mersenne Twister 生成器。worker A 在某次迭代里比 worker B 多消费了一个随机数——可能是因为 JPEG 解码速度不同，也可能是因为 CPU 调度抖动——那么从这一刻开始，两个 worker 拿到的所有后续随机数都会发生错位。数据增强的随机裁剪位置变了，Random Erasing 的区域变了，如果它们还共享同一 RNG 做其他事，权重初始化也会跟着漂移。最终你看到的结果不是“某个样本增强不同”，而是整条训练轨迹的分叉。

更棘手的是，这种问题通常不会崩溃，而是“悄悄地错”。模型照样收敛，loss 曲线看起来正常，但两次运行的权重、精度、甚至超参数敏感性都不同。调试时你根本无法判断这是算法本身的特性，还是随机数顺序被扰动导致的噪声。

主流框架的做法是“分而治之”：PyTorch 为 CPU 和 CUDA 分别维护 RNG 状态，CPU 用 Mersenne Twister，CUDA 用 Philox；用户需要手动设置 torch.manual_seed、torch.cuda.manual_seed_all，为每个 DataLoader worker 写 worker_init_fn，再关掉 cudnn.benchmark、打开 cudnn.deterministic，甚至启用 torch.use_deterministic_algorithms(True)。即便如此，官方文档也坦承：跨 PyTorch 版本、跨平台、CPU 与 GPU 之间仍无法保证完全一致。

根本原因在于，状态机式 RNG 的“可复现”是建立在执行顺序严格不变的前提上的，而现代训练管线天然就是乱序并行的。要跳出这个困境，最好是让随机数生成本身不再依赖执行顺序——这就是计数器型 RNG（Counter-Based RNG）的出发点。

二、Philox：把“第 n 个随机数”变成纯函数

Philox 由 Salmon 等人于 2011 年提出，是 Random123 库的代表性算法之一，也是 PyTorch CUDA、TensorFlow、NVIDIA cuRAND 等广泛采用的并行随机数方案。它的核心思想非常优雅：

给定一个种子 seed 和一个计数器 offset，输出 f(seed, offset) 就是确定性的随机数。

换句话说，第 n 个随机数不再由第 n-1 个状态递推而来，而是可以直接计算。这个性质带来了几个直接好处：

• 并行友好：每个线程用不同的 offset 独立生成，无需竞争共享状态；
• 可跳转：想要第 100 亿个随机数，不需要先跑 100 亿步；
• 顺序无关：只要 (seed, offset) 映射不变，无论哪个线程先算、后算，结果集合都不变。

Philox4x32-10 的命名已经说明了结构：4 个 32 位计数器、2 个 32 位密钥，经过 10 轮“乘高-乘低-XOR”的 Feistel 结构变换。每一轮都会用 Weyl 序列常量 0x9E3779B9、0xBB67AE85 更新密钥，保证充分混淆。它通过 TestU01 等统计检验，周期长达 2^128，对任何实际训练任务都绰绰有余。

在单轮变换里，Philox 把两个 64 位乘积拆成高 32 位和低 32 位，再与密钥和另一半计数器做 XOR。这种“乘高-乘低”操作是现代 CPU 和 GPU 都原生支持的位运算与整数乘法，因此 Philox 既快又容易向量化。10 轮迭代是一个经验上足够安全的参数：再多几轮只会增加计算量，而 10 轮已经能让每个输出位都充分依赖每个输入位。

Tech-Renaissance 没有依赖外部库，而是在 include/renaissance/core/philox.h 中实现了一份CPU/GPU 通用的 Philox4x32-10 内核：

constexpr uint32_t PHILOX_M4x32_0 = 0xD2511F53u;
constexpr uint32_t PHILOX_M4x32_1 = 0xCD9E8D57u;
constexpr uint32_t PHILOX_W32_0   = 0x9E3779B9u;  // golden ratio
constexpr uint32_t PHILOX_W32_1   = 0xBB67AE85u;  // sqrt(3)-1

TR_HOST_DEVICE TR_FORCEINLINE
void philox4x32_round(
    uint32_t* ctr0, uint32_t* ctr1, uint32_t* ctr2, uint32_t* ctr3,
    uint32_t key0, uint32_t key1)
{
    uint32_t lo0, lo1;
    uint32_t hi0 = mulhilo32(PHILOX_M4x32_0, *ctr0, &lo0);
    uint32_t hi1 = mulhilo32(PHILOX_M4x32_1, *ctr2, &lo1);

    uint32_t new_ctr0 = hi1 ^ *ctr1 ^ key0;
    uint32_t new_ctr1 = lo1;
    uint32_t new_ctr2 = hi0 ^ *ctr3 ^ key1;
    uint32_t new_ctr3 = lo0;

    *ctr0 = new_ctr0;
    *ctr1 = new_ctr1;
    *ctr2 = new_ctr2;
    *ctr3 = new_ctr3;
}

TR_HOST_DEVICE TR_FORCEINLINE
void philox4x32_10(
    uint32_t* ctr0, uint32_t* ctr1, uint32_t* ctr2, uint32_t* ctr3,
    uint32_t key0, uint32_t key1)
{
    for (int round = 0; round < 10; ++round) {
        philox4x32_round(ctr0, ctr1, ctr2, ctr3, key0, key1);
        key0 += PHILOX_W32_0;
        key1 += PHILOX_W32_1;
    }
}

变量说明：

• ctr0..ctr3：4 个 32 位计数器，输入的低 64 位由 offset 填充，高 64 位补 0；
• key0, key1：2 个 32 位密钥，由 64 位 seed 拆分得到；
• PHILOX_M4x32_*：两个固定乘数，负责把计数器信息扩散到高 32 位；
• PHILOX_W32_*：Weyl 序列常量，每轮加到密钥上，让每轮使用的密钥不同。

这里的 TR_HOST_DEVICE 宏在 CUDA/MUSA 下展开为 __host__ __device__，意味着同一份算法代码既能在 CPU 上运行，也能被编译进 GPU kernel。这是保证跨设备数值一致的第一步。如果 CPU 和 GPU 使用不同算法生成同一组“正态分布”随机数，即便种子相同，初始化出来的权重也会不同，后续训练就无从谈起。

2.1 为什么 mulhilo 是 Philox 的核心

Philox 的安全性并不来自复杂的查表或 S-box，而是来自乘法的高位扩散。32 位 × 32 位 = 64 位，高 32 位包含了低 32 位经过乘法后的“混合信息”。把高 32 位和低 32 位拆开，分别赋给不同的计数器分量，再用 XOR 与密钥结合，就完成了单轮的混淆。

TR_HOST_DEVICE TR_FORCEINLINE
uint32_t mulhilo32(uint32_t a, uint32_t b, uint32_t* lo) {
#ifdef __CUDA_ARCH__
    uint32_t hi;
    asm("mul.hi.u32 %0, %1, %2;" : "=r"(hi) : "r"(a), "r"(b));
    *lo = a * b;
    return hi;
#else
    uint64_t product = static_cast<uint64_t>(a) * static_cast<uint64_t>(b);
    *lo = static_cast<uint32_t>(product);
    return static_cast<uint32_t>(product >> 32);
#endif
}

在 GPU 上，我们直接用 PTX 指令 mul.hi.u32 取高 32 位，避免 64 位乘法的额外寄存器消耗；在 CPU 上，直接用 64 位整数乘法一次拿到高、低两半。两种路径的输出必须逐位一致，否则 CPU 和 GPU 生成的随机数就会分叉。

2.2 从整数到浮点：均匀分布与正态分布

Philox 一轮输出 4 个 32 位无符号整数。要得到训练常用的浮点分布，需要做一次转换。

均匀分布 [0, 1) 采用经典的高位映射法：取第一个 32 位整数的高 24 位，乘以 2^-24，避免整数除法带来的不均匀性：

TR_HOST_DEVICE TR_FORCEINLINE
float philox_uniform_float(uint64_t seed, uint64_t offset) {
    uint32_t r[4];
    philox_generate_4x32(seed, offset, r);

    constexpr float scale = 1.0f / 16777216.0f;  // 2^-24
    return static_cast<float>(r[0] >> 8) * scale;
}

正态分布则通过 Box-Muller 变换，把两个均匀随机数转换成一对标准正态随机数：

TR_HOST_DEVICE TR_FORCEINLINE
void philox_normal_pair(uint64_t seed, uint64_t offset, float* out0, float* out1) {
    uint32_t r[4];
    philox_generate_4x32(seed, offset, r);

    constexpr float scale = 1.0f / 16777216.0f;
    float u1 = (static_cast<float>((r[0] >> 8) | 1)) * scale;  // (0, 1]
    float u2 = static_cast<float>(r[1] >> 8) * scale;          // [0, 1)

    constexpr float two_pi = 6.283185307179586f;

#ifdef __CUDA_ARCH__
    float radius = sqrtf(-2.0f * logf(u1));
    float theta = two_pi * u2;
    float sin_theta, cos_theta;
    sincosf(theta, &sin_theta, &cos_theta);
    *out0 = radius * cos_theta;
    *out1 = radius * sin_theta;
#else
    float radius = std::sqrt(-2.0f * std::log(u1));
    float theta = two_pi * u2;
    *out0 = radius * std::cos(theta);
    *out1 = radius * std::sin(theta);
#endif
}

因为 philox_normal_pair 一次消耗一个 offset 却产出两个正态随机数，所以 CPU 正态分布生成器在调用 gen.next_offset() 时，预留的是 (count + 1) / 2 个 offset，而不是 count 个。这是一个实现细节，但对可复现性至关重要：如果生成器消费 offset 的方式发生变化，即使 seed 相同，输出序列也会变。

2.3 周期与统计质量

Philox4x32-10 的周期是 2^128（4 个 32 位计数器总共 128 位），对于任何实际深度学习任务来说都绰绰有余。即使每天生成 10^12 个随机数，也要超过 10^19 年才能耗尽周期。

在统计检验方面，Philox 通过了 TestU01 的 SmallCrush / Crush / BigCrush 等主流随机数统计检验，其均匀分布、正态分布、位级独立性都与 MT19937 处于同一水平。对于深度学习训练而言，我们更关心的是：

• 给定相同 seed，不同运行是否产生完全相同的序列；
• 给定相同 seed，CPU 和 GPU 是否产生完全相同的序列；
• 多线程并发时是否不引入额外的不确定性。

Philox 在这三点上都表现得非常干净。

三、Generator：一个种子 + 一个原子计数器

在 Philox 之上，Tech-Renaissance 封装了 include/renaissance/core/rng.h 中的 Generator 类。它的状态极简：

class Generator {
public:
    explicit Generator(uint64_t seed = 0) noexcept;
    ~Generator();
    Generator(const Generator&) = delete;
    Generator& operator=(const Generator&) = delete;
    Generator(Generator&& other) noexcept;
    Generator& operator=(Generator&& other) noexcept;

    void set_seed(uint64_t seed);
    uint64_t seed() const noexcept;

    std::pair<uint64_t, uint64_t> get_state() const;
    void set_state(uint64_t seed, uint64_t offset);

    uint64_t next_offset(uint64_t count);
    uint64_t current_offset() const noexcept;

    int random_int(int low, int high);

private:
    class Impl;
    std::unique_ptr<Impl> impl_;
};

Impl 里只有三个成员：uint64_t seed_、std::atomic<uint64_t> offset_ 和一个保护 seed/offset 读写的 std::mutex。所有并发请求都通过 next_offset(count) 原子地预留一段连续的 offset 区间：

uint64_t Generator::next_offset(uint64_t count) {
    return impl_->offset_.fetch_add(count, std::memory_order_relaxed);
}

返回的是本次调用的起始 offset，调用者可以安全地占用 [offset, offset + count) 这段空间。即便多个线程同时申请，由于 fetch_add 的原子性，每个人拿到的区间互不重叠；而只要申请顺序相同，最终生成的随机数序列就完全相同。

get_state 和 set_state 让训练中断后可以从检查点精确恢复随机状态，对长周期训练尤为重要。seed 与 offset 的二元组足够小，可以直接序列化到 checkpoint 文件里，不需要像 Mersenne Twister 那样保存 2.5 KB 的内部状态数组。这意味着 checkpoint 更小，恢复逻辑也更清晰。

这里有一个值得注意的设计细节：头文件使用 Pimpl 模式 隐藏 std::atomic。原因是 MUSA SDK 的 musa/std/atomic 与标准库 <atomic> 存在命名空间冲突，如果在公共头文件里直接包含 <atomic>，会在某些国产 GPU 编译环境下触发大量错误。把原子成员放到 .cpp 的实现类中，既解决了兼容性问题，也没有牺牲 next_offset 的性能。

3.1 保存与恢复 RNG 状态

对于长周期训练，checkpoint 不仅要保存权重，还要保存随机状态，否则恢复训练后随机序列会从头开始，导致结果与不间断训练不同。Generator 提供了非常轻量的状态接口：

// 训练前保存
auto [seed, offset] = gen.get_state();
save_checkpoint({seed, offset, model_weights});

// 恢复训练时加载
gen.set_state(loaded_seed, loaded_offset);

整个状态只有两个 64 位整数，相比 Mersenne Twister 需要保存 624 个 32 位整数的状态数组，Philox 的 checkpoint 负担几乎为零。这也是计数器型 RNG 在工程上的另一个优势。

3.2 用户侧种子设置

用户侧的种子设置通过 GlobalRegistry::manual_seed(seed) 完成：

GlobalRegistry& GlobalRegistry::manual_seed(uint64_t seed) {
    // 只要设定种子，就必定是想要可复现
    // 因此会自动使用可复现模式，以微小的性能代价换来随机可复现性
    reproducible();
    rng_set_seed(seed);
    return *this;
}

只要调用 manual_seed，框架就会自动开启“可复现性保险”，把 TransferStation、数据洗牌、随机增强等全部切换到确定性模式。这是 Tech-Renaissance 的一个高层约定：用户只需要说一次“我要固定种子”，剩下的事情框架负责。相比之下，PyTorch 用户通常需要为一堆独立的 RNG 分别设种子，并手动管理各种后端开关。

3.3 一个最小的可复现示例

在实际代码里，用户不需要关心 Philox 的细节。以 tests/example/mlp_mnist.cpp 为例，开启可复现训练只需要在配置链里加一行：

GLOBAL_SETTING
    .use_gpu("0")
    .amp(true)
    .manual_seed(123)      // 固定随机种子，保证结果可复现
    .global_batch_size(200)
    .input_resolution(28);

manual_seed(123) 背后发生的事情包括：

• 设置全局 Generator 的 seed 为 123，offset 归零；
• 开启 reproducibility_insurance_，让 TransferStation 进入可复现模式；
• 每个 PreprocessWorker 用 global_seed ^ (worker_id << 32) 派生自己的初始种子；
• 每个 rank 的 Dropout seed 通过 SplitMix64 哈希从 123 派生；
• 权重初始化、数据增强、数据集洗牌全部接入同一套 Philox 体系。

也就是说，一行调用就替代了 PyTorch 里常见的“seed everything”模板函数，把种子管理从用户侧转移到框架侧。

3.4 为什么不用 Mersenne Twister

Mersenne Twister（MT19937）是 C++ 标准库和许多科学计算库的默认选择，它周期极长（2^19937 - 1）、统计质量好，但有一个根本问题：它是一个状态机。要生成第 n 个随机数，必须先生成前 n-1 个。这个特性导致它在并发场景下非常别扭：

• 多线程共享一个生成器必须加锁，锁竞争和调度顺序都会破坏可复现性；
• 每个线程一个生成器，线程创建顺序变化会导致种子分配顺序变化；
• 状态数组 2.5 KB，保存到 checkpoint 或传递给 GPU 都不方便。

Philox 用计数器替代状态机后，这些问题迎刃而解。它的状态只有两个 64 位整数，每个线程可以独立计算自己那一份，不需要锁，也不需要维护历史状态。在 GPU 上，这意味着成千上万个线程可以同时生成随机数，而不会有任何竞争。

3.5 统计特性：Philox 与 MT19937 的对比

我们在 tests/correction/rng_comparison_test.cpp 中实现了一个简单的统计对比程序，用同样的种子分别生成 Philox 和 MT19937 的均匀分布、正态分布样本，然后比较均值、标准差、最小值、最大值。Philox 与 MT19937 的统计特性在同一量级，完全满足深度学习训练的需求。需要注意的是，这个对比程序目前尚未接入 tests/correction/CMakeLists.txt，它更像是一份随用随取的验证脚本，而不是每日持续集成的一部分。

四、从权重初始化到数据增强：Philox 贯穿全流程

4.1 参数初始化

src/tensor/tensor.cpp 中的 Tensor::normal、Tensor::uniform 等方法直接调用 cpu_rand_* 系列函数：

void Tensor::normal(float mean, float stddev) {
    const int64_t total_elems = numel();
    float* data = static_cast<float*>(ptr_);
    cpu_rand_normal_float(data, total_elems, mean, stddev);
}

这些函数的实现（src/core/rng.cpp）统一遵循同一个模式：先 next_offset 预留整个张量所需的 offset，再用 OpenMP 并行生成，每个线程用 base_offset + i 作为 Philox 计数器：

void cpu_rand_normal_float(float* ptr, size_t count, float mean, float std, Generator& gen) {
    // Box-Muller 每次生成 2 个数，所以 offset 消耗量为 (count + 1) / 2
    uint64_t pairs_needed = (count + 1) / 2;
    uint64_t base_offset = gen.next_offset(pairs_needed);
    uint64_t seed = gen.seed();

    int num_threads = get_num_threads(count);
    int64_t pair_count = static_cast<int64_t>(count / 2);

#if TR_USE_OPENMP
    #pragma omp parallel for num_threads(num_threads) schedule(static)
#endif
    for (int64_t i = 0; i < pair_count; ++i) {
        float n0, n1;
        detail::philox_normal_pair(seed, base_offset + i, &n0, &n1);
        ptr[i * 2]     = mean + std * n0;
        ptr[i * 2 + 1] = mean + std * n1;
    }

    // 处理奇数尾部（如果有）
    if (count % 2 == 1) {
        float n0, n1;
        detail::philox_normal_pair(seed, base_offset + pair_count, &n0, &n1);
        ptr[count - 1] = mean + std * n0;
    }
}

注意 schedule(static) 的重要性：它确保线程 i 永远处理同一段 offset，即便线程调度有抖动，只要 base_offset 和 count 相同，输出就相同。如果换成 schedule(dynamic)，随机数与线程的绑定关系会随运行时变化，可复现性就会被破坏。

对于 DTensor 在 GPU 上的初始化，src/backend/infra_kernels.cu 提供了一个 Philox 正态分布 kernel，被 TaskBase::randn 调用：

__global__ void tr_philox_normal_float_kernel(
    int n, uint64_t seed, uint64_t base_offset,
    float mean, float std, float* out)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= n) return;

    float n0, n1;
    philox_normal_pair(seed, base_offset + idx / 2, &n0, &n1);

    if (idx % 2 == 0) {
        out[idx] = mean + std * n0;
    } else {
        out[idx] = mean + std * n1;
    }
}

CPU 与 GPU 用的是同一个 philox_normal_pair，只是承载并行粒度的载体不同：一个是 OpenMP 线程，一个是 CUDA 线程。这避免了“CPU 初始化一套、GPU 初始化一套”带来的数值分叉。

在框架初始化流程中，TaskBase::init_all() 会遍历 MemoryPlan 中的每个 DTensor，根据 InitConfig 决定调用哪种初始化方式。对于 Kaiming/Xavier 等常见初始化，最终都会落到 Tensor::normal 或 Tensor::uniform（CPU 路径），或 TaskBase::randn（GPU 路径）。这些入口统一使用全局 Generator，所以只要 manual_seed 相同，所有层的初始化结果都完全确定。

4.2 数据增强与洗牌

数据预处理是随机性的另一个重灾区。Tech-Renaissance 的每个 PreprocessWorker 都持有自己独立的 Generator rng_，其种子在 ensure_rng_initialized() 中按 worker 和 phase 衍生：

void PreprocessWorker::ensure_rng_initialized() {
    if (!rng_initialized_) {
        uint64_t base_seed = initial_seed_;
        uint64_t worker_seed = base_seed ^ (static_cast<uint64_t>(param_.phase_id) << 16);
        rng_.set_seed(worker_seed);
        rng_initialized_ = true;
    }
}

其中 initial_seed_ = global_seed ^ (worker_id << 32)，在构造时就固定下来。这意味着不同 worker 的种子不同，同一 worker 不同 epoch 的种子也不同，但全部由全局种子确定性派生。worker 之间不会因为共享全局 RNG 而互相“吃掉”对方的随机数。

对于每个 epoch 的数据集洗牌，src/data/preprocess_worker.cpp 中的 shuffle_s_indices() 没有使用 Generator::random_int()——因为那会修改 Generator 状态，可能被其他随机操作污染——而是直接调用 Philox 原语：

for (int i = n - 1; i > 0; --i) {
    uint32_t r[4];
    detail::philox_generate_4x32(shuffle_seed, static_cast<uint64_t>(i), r);
    uint32_t j = r[0] % (i + 1);
    std::swap(s_shuffled_indices_[i], s_shuffled_indices_[j]);
}

shuffle_seed 由 initial_seed_ ^ (phase_id << 32) 得到，offset 直接取循环变量 i。这是一个完全独立的随机流，不与其他任何随机操作共享计数器，因此无论训练过程中参数初始化、Dropout、Random Erasing 消耗了多少随机数，洗牌结果都不会受影响。

在 src/data/fused_normalization.cpp 里，Random Erasing、随机水平翻转等增强同样通过 Generator* 完成。FusedNormalization::uniform 和 randint 都是基于 Philox 原语的薄封装：

float FusedNormalization::uniform(float min_val, float max_val, Generator* rng) const {
    uint64_t offset = rng->next_offset(1);
    float rand_val = detail::philox_uniform_float(rng->seed(), offset);
    return min_val + rand_val * (max_val - min_val);
}

int FusedNormalization::randint(int min_val, int max_val, Generator* rng) const {
    return rng->random_int(min_val, max_val);
}

这确保这些增强与参数初始化、Dropout 使用同一套随机数体系，而不是从某个独立库里再拉一条 RNG 进来。对于 Random Erasing，它会根据 erase_p_ 决定是否生成擦除矩形，再依次用 uniform 采样面积比例、长宽比、以及矩形左上角位置；所有这些随机消费都通过当前 worker 的 Generator 完成，并且消费顺序固定。

4.3 可复现的数据搬运：TransferStation 双模式

即使随机数本身确定了，如果多个 worker 往 GPU 搬运数据时的写入顺序不确定，最终 batch 内的样本排列仍可能变化。为此，src/data/transfer_station.cpp 设计了两种模式，由 require_reproducibility_ 标志控制：

• 可复现模式：每个 worker 按固定公式计算自己在 batch 中的 position 和 batch_id，快 worker 必须等慢 worker 完成当前 batch 才能继续，传输时机完全确定；
• 非可复现模式：用原子计数器动态分配 slot，性能更高，但每次运行顺序可能不同。

if (require_reproducibility_) {
    std::unique_lock<std::mutex> lock(mutex_);
    cv_batch_ready_.wait(lock, [this, batch_id]() {
        return current_batch_id_.load() >= batch_id || finished_.load();
    });

    int buf_id = current_buffer_.load(std::memory_order_acquire);
    buffer_labels_[buf_id][position] = label;
    size_t offset = position * snapshot_sample_bytes;
    return buffer_data_[buf_id] + offset;
}

这是 Tech-Renaissance 对“可复现”做出的一个明确 trade-off：开启可复现模式后，数据管线的吞吐量会略有下降，因为快 worker 会被慢 worker 拖住；但换来的是字节级一致的样本顺序。对于调试和科学实验，这是值得的。对于追求极限性能的生产跑分，可以关闭可复现模式，让 TransferStation 走无锁快速路径。

4.4 Dropout：在 CUDA Graph 内自助旋转种子

Dropout 的难点在于它发生在 GPU kernel 里，而 Tech-Renaissance 的训练循环被完整捕获为 CUDA Graph，CPU 不能在每次 forward 时插进去改种子。

解决方案是在 MemoryPlan 里为每个 rank 分配一个 dropout_seed 张量（形状 {1,1,1,2}，INT32），放在 S_SCALAR_INT32 区域：

Shape seed_shape{1, 1, 1, 2};
auto seed_dt = alloc_impl(seed_shape, DType::INT32, Region::S_SCALAR_INT32);
set_init_config(seed_dt.id, InitConfig{0.0f, InitKind::NONE, FanMode::FAN_IN});
baseline_.dropout_seed = seed_dt.id;

初始化时，TaskBase::init_all() 用全局种子通过一个 SplitMix64 风格的哈希为每个 rank 生成独立种子，并 H2D 拷贝到该张量：

if (active_memory_plan_->baseline().dropout_seed >= 0) {
    uint64_t global_seed = get_default_generator().seed();
    for (int rank = 0; rank < num_gpus_; ++rank) {
        uint64_t z = global_seed + static_cast<uint64_t>(rank) + 0x9e3779b97f4a7c15ULL;
        z = (z ^ (z >> 30)) * 0xbf58476d1ce4e5b9ULL;
        z = (z ^ (z >> 27)) * 0x94d049bb133111ebULL;
        uint64_t rank_seed = z ^ (z >> 31);

        int32_t seed_data[2] = {
            static_cast<int32_t>(rank_seed & 0xFFFFFFFFULL),
            static_cast<int32_t>(rank_seed >> 32)
        };

        Tensor host_seed(Shape{1, 1, 1, 2}, DType::INT32);
        memcpy(host_seed.data(), seed_data, sizeof(seed_data));
        transfer_to_rank(host_seed, active_memory_plan_->get_dtensor(
            active_memory_plan_->baseline().dropout_seed), rank);
    }
}

在每次 Dropout forward 的 CUDA Graph 节点里，kernel 先调用 rotate_dropout_seed_kernel 用 Xorshift64* 原地旋转种子，再从这个设备上的种子生成 Philox 样本：

__global__ void rotate_dropout_seed_kernel(int32_t* seed_ptr) {
    uint64_t s = (static_cast<uint64_t>(static_cast<uint32_t>(seed_ptr[1])) << 32)
               | static_cast<uint32_t>(seed_ptr[0]);
    if (s == 0) s = 0x9e3779b97f4a7c15ULL;  // 防御 stuck
    s ^= s >> 12;
    s ^= s << 25;
    s ^= s >> 27;
    s *= 0x2545F4914F6CDD1DULL;
    seed_ptr[0] = static_cast<int32_t>(s & 0xFFFFFFFFULL);
    seed_ptr[1] = static_cast<int32_t>(s >> 32);
}

CPU 路径在 dropout_op.cpp 中做了逐位一致的实现。这样无论走 GPU CUDA Graph 还是 CPU 回退，同一个全局种子都会得到相同的 Dropout 掩码序列。关键在于，种子的旋转完全在设备端完成，CUDA Graph 捕获一次后就可以反复重放，不需要 CPU 介入。

值得一提的是，Dropout kernel 内部使用的是 Philox 2×64 简化版，而不是 philox.h 中的 4×32-10：

__device__ float philox_sample(uint64_t seed_lo, uint64_t seed_hi, uint64_t offset) {
    uint64_t counter = seed_lo + offset;
    uint64_t key0 = seed_hi;
    uint64_t key1 = seed_lo ^ 0x9e3779b97f4a7c15ULL;

    for (int i = 0; i < 10; ++i) {
        // mulhilo：用 64 位乘法把计数器扩散到高 64 位
        uint64_t lo = 0xD2B74407B1CE6E93ULL * counter;
        uint64_t hi = __umul64hi(0xD2B74407B1CE6E93ULL, counter);
        uint64_t new_counter = hi ^ key0 ^ key1;
        uint64_t new_key = lo;
        counter = new_counter;
        key0 = new_key;
        key1 = new_key ^ key1;
    }

    uint32_t upper = static_cast<uint32_t>(counter >> 32);
    return __uint2float_rn(upper) * (1.0f / 4294967296.0f);
}

变量说明：

• seed_lo / seed_hi：由 per-rank dropout seed 拆分出的两个 64 位量；
• offset：每个输出元素的全局扁平索引；
• counter / key0 / key1：Philox 2×64 的核心状态，每轮通过 64 位乘高-乘低扩散；
• 最终取 counter 的高 32 位映射到 [0, 1)。

这种 2×64 变体比 4×32-10 更适合 Dropout：每个线程只需要一个 float，无需像 4×32 那样一次生成 4 个整数再浪费 3 个。

需要强调的是，Dropout 的 CPU 路径和 GPU 路径虽然在不同文件中实现，但算法是逐位对齐的：

• 两者都用 SplitMix64 从全局种子派生 per-rank seed；
• 两者都用 Xorshift64* 对 seed 做同构旋转；
• 两者都用 Philox 2×64 采样 offset = 元素索引；
• 两者都把高 32 位映射到 [0, 1)。

因此，即使某次运行因为环境没有 GPU 而回退到 CPU，只要全局种子相同，Dropout 掩码序列仍然一致。

五、随机数之外：确定性训练还需要什么？

把随机数换成 Philox 只是必要条件，不是充分条件。Tech-Renaissance 为了端到端可复现，还做了几件事：

1. 静态图 + CUDA Graph 全捕获：算子序列、内存地址、调度顺序在编译期固定，消除了运行时调度和地址变化带来的不确定性。动态图框架里常见的 graph break、动态形状、运行时分配，都会成为不可复现的来源。Tech-Renaissance 的 MemoryPlan 在编译期就为每个 DTensor 计算出固定的 offset，运行期所有张量都从这个静态池里按偏移访问。这意味着同一次训练、同一个模型、同一个 batch，每次 forward/backward 访问的显存地址完全相同。CUDA Graph 捕获一次后反复重放，CPU 不再参与每次迭代的 kernel 启动，自然也不会因为 CPU 调度抖动而改变 GPU 上的执行时序。
2. 避免求和顺序变化的归约 kernel：浮点加法不满足结合律，分散的 atomicAdd 会导致结果随线程调度波动。例如，多个 CUDA block 同时往同一个地址累加梯度时，block 的到达顺序不同，最终舍入结果就会不同。框架在关键路径上尽量使用稳定的归约策略，减少多卡之间的数值漂移。举个具体例子：假设三个梯度值分别是 1.0f、1e-8f、1e-8f，如果先用大数加小数：float a = (1.0f + 1e-8f) + 1e-8f;  // 结果仍是 1.0f（第一次加已被舍入）
float b = 1e-8f + 1e-8f + 1.0f;    // 结果可能是 1.0000002f在单精度下，1.0f + 1e-8f == 1.0f，因为小数被大数的尾数“吃掉”了。不同的累加顺序会导致最终结果不同。如果 AllReduce 的 block 到达顺序每次运行都不一样，那么即使梯度本身完全相同，聚合后的权重更新也会出现微小差异，并在后续优化步骤中被放大。
3. 确定性的 MaxPool 反向实现：CUDA 路径中，src/backend/ops/dtensor/maxpool_op.cu 实现了 maxpool_bwd_deterministic_kernel，通过”输入遍历 + 收集累加”的方式避免传统 atomicAdd 带来的非确定性；如果走 cuDNN pooling 路径，则显式选择 CUDNN_POOLING_MAX_DETERMINISTIC。CPU 路径则使用固定遍历顺序的实现，保证同一输入下结果稳定。
4. 禁用 TF32：GlobalRegistry::use_tf32(false) 通过设置环境变量 NVIDIA_TF32_OVERRIDE=0，强制 Ampere 及更新架构的 GPU 使用纯 FP32 计算，避免 TF32 低精度带来的运行间差异。
5. 一致的 Region 布局：所有 rank 共享同一份 MemoryPlan，DTensor 偏移跨卡一致，保证多卡随机数消费顺序一致。

即便如此，我们也必须诚实地说：绝对的跨平台、跨硬件可复现性在工程上很难承诺。GPU 驱动版本、cuDNN/cuBLAS 版本、不同 GPU 的 Tensor Core 行为都可能引入微小差异。Tech-Renaissance 能保证的是：在相同代码、相同硬件、相同种子、开启可复现模式的前提下，多次运行得到一致结果。

5.1 一个值得注意的细节：LARS 信任系数的确定性

在 B20《融合优化器》中我们提到过，LARS 需要对每个可训练层分别计算 ||w||^2 和 ||g||^2，这是典型的归约操作。为了保证可复现，LARS 的信任系数计算采用两阶段固定顺序归约：第一阶段每个 block 计算部分和并写入固定的部分和缓冲区，第二阶段按固定顺序累加这些部分和。只要 block 网格大小不变，求和顺序就不变，结果也就完全一致。

六、与 PyTorch 路径的对比

PyTorch 的 RNG 设计是“分层的”：

• CPU 用 Mersenne Twister，必须顺序消费；
• CUDA 用 Philox，但每个张量操作有自己的 offset 计数；
• DataLoader 需要单独为每个 worker 设种子，典型写法如下：

def seed_worker(worker_id):
    worker_seed = torch.initial_seed() % 2**32
    numpy.random.seed(worker_seed)
    random.seed(worker_seed)

g = torch.Generator()
g.manual_seed(0)
DataLoader(train_dataset, batch_size=..., num_workers=...,
           worker_init_fn=seed_worker, generator=g)

• 用户还需要手动处理 cudnn.deterministic、cudnn.benchmark、use_deterministic_algorithms、CUBLAS_WORKSPACE_CONFIG 等开关。

Tech-Renaissance 的做法是“统一的”：从权重初始化、数据增强、数据集洗牌到 Dropout，全部基于同一份 Philox 原语；通过 manual_seed 一个入口，把 Generator、worker 衍生种子、Dropout per-rank seed、TransferStation 可复现模式全部串起来。它牺牲了 DataLoader 在可复现模式下的极限吞吐，但换来的是“一键可复现”的体验。

这种差异也反映在设计哲学上。PyTorch 倾向于把控制权交给用户：你可以为 CPU、CUDA、每个 DataLoader worker、每个第三方库分别设种子，灵活但繁琐。Tech-Renaissance 倾向于把控制权收回到框架：只要用户表达“我要可复现”，框架就自动把所有随机消费点串到同一条确定性轨道上。这不是说 PyTorch 做不到可复现，而是说在 Tech-Renaissance 里，做对的成本更低。

七、性能与 trade-off

Philox 算法本身的计算开销很低：10 轮整数乘法和 XOR，现代 CPU 和 GPU 都能高效执行。在我们的实现里，Dropout kernel 中随机数生成只占整个 kernel 很小一部分时间，主要开销仍然是显存读写；参数初始化几百万个 float 也只需要微秒级。因此，随机数生成本身并不是可复现性的性能瓶颈。

相比 Mersenne Twister，Philox 在并行场景下甚至可能更快：MT19937 需要维护一个 624 字的状态数组，频繁访问会带来缓存压力；而 Philox 每次生成只需要少量寄存器，没有状态数组，缓存友好。在 CPU 多线程场景下，Philox 不需要锁，多个线程可以无竞争地同时生成随机数。

真正的性能代价来自数据预处理阶段的同步。在可复现模式下，TransferStation 必须“快 worker 等慢 worker”，保证每个样本写入预定位置。如果 worker 之间处理速度差异很大，快 worker 会被慢 worker 拖慢，整体吞吐量会有所下降。这个开销与数据集、worker 数量、预处理复杂度都相关，无法给出一个放之四海而皆准的数字；用户可以根据场景在“可复现”和“极限吞吐”之间切换。

另一个潜在的代价是某些算子必须选择确定性实现。例如 MaxPool 的反向传播，如果使用传统的“输出遍历 + atomicAdd”策略，不同 block 的到达顺序会导致结果波动；我们的确定性实现采用“输入遍历 + 收集累加”，避免了原子操作，但可能牺牲一点点峰值性能。这是一种可预测、可量化的取舍。

场景	是否建议开启可复现	理由
调试 bug	✅	每次运行代码路径相同，方便定位
论文实验	✅	结果可复现是科学研究的基本要求
超参数搜索	✅	排除随机性干扰，准确比较超参数效果
生产环境大规模跑分	❌	性能优先，微小差异不影响最终结果
ablation study	✅	准确比较不同组件的贡献

如果不调用 manual_seed，框架会默认使用 auto_seed() 取时间戳作为种子，并关闭可复现性保险，此时 TransferStation 走无锁快速路径，数据管线可以达到最高吞吐。这种“要可复现就一行开启，要性能就不额外同步”的设计，让同一个框架既能服务科研调试，也能服务生产跑分。

7.1 可复现性的边界

虽然我们把确定性训练做到了框架层面的“一键开启”，但仍有一些边界需要说明：

• 跨 GPU 架构：不同架构的 Tensor Core、FP32 单元实现可能有微小差异，不能保证 A100 和 RTX 4090 跑出完全一致的 loss 曲线；
• 跨驱动 / 库版本：cuDNN、cuBLAS 的小版本更新有时会改变某些 kernel 的取舍或舍入行为；
• 跨编译器 / 优化级别：CPU 路径的 OpenMP 划分、向量化可能影响浮点累加顺序；
• 分布式异步通信：虽然我们尽量让通信与计算重叠并保持确定性，但极端情况下网络延迟波动仍可能影响某些时序敏感的操作。

因此，我们不说“跨平台字节级一致”，而是说“在相同代码、相同硬件、相同软件栈、相同种子下可复现”。这是一个更诚实、也更有工程价值的承诺。

八、小结

随机数不是深度学习框架的装饰品，而是决定训练能否被信任的基础设施之一。Philox 的计数器型设计把“第 n 个随机数”变成纯函数，从根本上解除了并行执行顺序对随机序列的绑定。Tech-Renaissance 围绕这一思想，打造了覆盖初始化、数据增强、洗牌、Dropout 的确定性管线，并通过 GlobalRegistry::manual_seed 和可复现性保险把它暴露为一行调用。

可复现性有代价：它要求你放弃一些因乱序执行而获得的性能，要求所有随机消费点都被框架精确计数，也要求底层算子不能引入非确定性。但在调试和科研场景里，这种代价通常是值得的。毕竟，一个跑得快却不可解释的训练过程，很难被称为真正可靠的训练。

回顾一下本文的要点：

1. 传统状态机式 RNG 在并发场景下天生难以保证可复现，线程调度顺序会影响状态更新顺序；
2. Philox 是基于计数器的无状态 RNG，相同 (seed, offset) 永远得到相同随机数，天然支持无锁并行；
3. Tech-Renaissance 的 Generator 用原子 offset_ 预留区间，配合 Pimpl 模式解决 MUSA 兼容性；
4. manual_seed() 一行调用即可开启端到端可复现模式；
5. Dropout 的 per-rank seed 用 SplitMix64 派生、Xorshift64* 自旋转，完全在 CUDA Graph 内完成；
6. 可复现性是全链路问题，还需要静态图、确定性归约、TF32 关闭等配合；
7. 在相同代码、相同硬件、相同种子下，Tech-Renaissance 可以保证多次运行结果一致。

下一篇，我们将把视角从单卡拉向多卡，聊聊 NCCL 分布式通信与数据并行训练。