昇腾随机数生成算子库 ops-rand：在昇腾NPU上实现确定性随机数生成的完整技术指南：架构解析、算法原理与实践应用

2401_86049470

16人浏览 · 2026-06-05 10:21:16

2401_86049470 · 2026-06-05 10:21:16 发布

前言

在深度学习训练与推理场景中，随机数生成器的质量直接影响模型初始化的稳定性和实验结果的可复现性。传统 CPU 侧随机数在 NPU 设备端使用时，需要经历一次跨设备拷贝，而 CANN（Compute Architecture for Neural Networks）生态下的 ops-rand 算子库正是为解决这一痛点而生的基础设施。ops-rand 是昇腾NPU随机数生成算子库，完全基于昇腾 AI Core 原生指令集实现，通过指定种子（seed）配合偏移量（offset）的设计，确保每次运行都能产生完全一致的随机数序列。这一特性在分布式训练断点恢复、对抗样本生成、蒙特卡洛模拟等需要结果复现的场景中具有不可替代的价值。

从一台混乱的赌场到一个纪律严明的乐团：为什么 NPU 需要专用随机数生成器

理解 ops-rand 的存在价值，需要先理解一个背景问题：随机数在 AI 硬件上究竟是如何被消费掉的。

在传统 CPU 编程中，随机数生成是简单直接的——调用 rand() 或 std::mt19937，随机字节从 CPU 内存直接抵达你的程序。但到了神经网络训练中，情况变得复杂得多。一张权重张量可能包含数十亿个参数，每个参数都需要用一个随机值初始化；在 Dropout、随机正则化、数据增强等环节，同样需要大量随机数。更关键的是，这些计算必须发生在 NPU 这样的专用加速器上。

这里出现了两种粗犷的解决方案，各有各的问题。第一种方案是完全在 CPU 侧生成随机数，然后通过 PCIe 总线拷贝到 NPU 设备内存中。这个方案看似简单，但当需要生成数十亿个随机数时，跨总线传输的开销会变成明显的性能瓶颈——数据在 CPU 和 NPU 之间来回搬运的过程中，NPU 只能处于空闲等待状态。第二种方案是在 NPU 上直接使用软件模拟的随机数生成器，但没有硬件指令级支持，软件模拟的随机数生成速度远低于 NPU 的计算吞吐量，导致随机数生成成为整条流水线的短板。

ops-rand 的设计哲学恰好对应了这两个问题的反面：它将随机数生成器作为昇腾 AI Core 上的一类一等公民算子实现，直接在 NPU 芯片内部完成全部计算过程，无需跨设备数据搬运。AI Core 内部拥有专门设计的向量计算单元，可以高效执行 Philox 等加密类随机数算法的核心运算，借助tiling 策略将大规模随机数生成任务拆解为适合在片上高速缓存（UB，Unified Buffer）中流转的小块数据处理流程，从而在算子内部就完成随机数的生成与分发。

ops-rand 的三层架构：从 Generator API 到 AI Core 算子

ops-rand 的代码组织可以类比为一栋三层楼的建筑，每层楼都有不同的入口和服务对象，但内部管道相互贯通。

第一层是 Generator API 层，对应头文件 include/cann_ops_rand.h。这是面向应用开发者的最外层接口，模拟了经典 GPU 随机数库的设计风格，提供类似 cuRAND 的编程体验。开发者通过 aclrandCreateGenerator 创建一个生成器句柄，设置种子和偏移量，然后调用 aclrandGenerateUniform 生成均匀分布的随机数。整个交互过程的思维模型非常接近打开一台老虎机：先初始化机器（创建生成器），设定赔率参数（设置种子和偏移量），然后拉动手柄（调用生成函数），硬币就从出币口滚出来了。

#include "cann_ops_rand.h"

aclrandGenerator_t generator;
float output[1024];

aclrandCreateGenerator(&generator, ACLRAND_RNG_PSEUDO_PHILOX4_32_10);
aclrandSetGeneratorSeed(generator, 42);
aclrandSetGeneratorOffset(generator, 0);
aclrandGenerateUniform(generator, output, 1024);
aclrandDestroyGenerator(generator);

为什么需要这套 Generator 接口？原因在于它将随机数生成器的配置状态（种子、偏移量、流）与执行逻辑解耦开来。生成器句柄可以持有内部状态，开发者可以在多次调用之间修改配置而不必每次都重新构建完整的计算图。对于需要分批次生成随机数的训练流程，比如每个 mini-batch 使用不同的 offset 但相同的 seed 来保证同一 epoch 内的数据增强一致性，这种设计模式非常自然。

第二层是算子调用层，对应底层函数 _aclrandStatelessRandomUniformV2。这一层绕过了 Generator 的状态管理开销，直接面向需要集成到计算图中的场景。函数签名直接接受 seed、offset、输出缓冲区指针、数据个数和数据类型，行为完全确定性：给定相同的输入参数组合，永远产生相同的输出序列。这正是"无状态"（stateless）一词的含义——算子本身不维护任何历史状态，序列的唯一性完全由 seed 和 offset 决定。

uint64_t seed = 12345;
uint64_t offset = 0;
int32_t alg = 1;  // 0=Philox算法，1表示特定配置
float output[100];

aclError ret = _aclrandStatelessRandomUniformV2(
    seed,           // 随机种子
    offset,         // 偏移量，用于产生不同序列片段
    alg,            // 算法类型标识
    output,         // 输出缓冲区（Host内存）
    100,            // 生成数量
    ACL_FLOAT,      // 数据类型：FP32
    stream          // ACL运行时流
);

为什么选择 seed 和 offset 两个参数的组合而不是单一参数？这是出于对复现粒度的精细控制。seed 决定整条随机数序列的初始状态，改变 seed 会产生完全不同的随机序列；offset 则允许在同一条序列上"快进"到任意位置。训练中断后恢复时，使用相同的 seed 配合相同的 global_step 乘以 batch_size 计算出的 offset，可以精确保留训练中断时的随机状态。而在同一个 epoch 的不同 step 之间使用相同的 seed 但递增的 offset，则保证同一批数据增强参数在不同 step 之间互不重复。

第三层是 AI Core 内核层，即 stateless_random_uniform_v2.cpp 中定义的 stateless_random_uniform_v2 核函数。这一层是整个技术栈中距离硬件最近的部分，完全使用昇腾 C 语言（AscendC）编写，直接操作 GM（Global Memory）、UB（Unified Buffer）和 AI Core 的向量化执行单元。

核函数的入口签名遵循 AscendC 的约定：使用 __global__ 和 __aicore__ 标记修饰函数，GM_ADDR 类型参数用于接收 Global Memory 地址，TilingData 结构体通过寄存器传递分块参数。进入核函数后，程序首先根据数据类型选择对应的模板实例化类：如果是 float，实例化 StatelessRandomUniformV2<float>；如果是 half，则实例化 StatelessRandomUniformV2<half>；如果是 bfloat16_t，则实例化 StatelessRandomUniformV2<bfloat16_t>。这一模板机制保证了三种数据类型在算法逻辑层面的代码复用。

extern "C" __global__ __aicore__ void stateless_random_uniform_v2(
    GM_ADDR y, StatelessRandomUniformV2TilingData tilingData)
{
    KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
    TPipe pipe;
    if (tilingData.tilingKey == FLOAT_TILING_KEY) {
        StatelessRandomUniformV2<float> op;
        op.Init(y, &tilingData, &pipe);
        op.Process();
    } else if (tilingData.tilingKey == FLOAT16_TILING_KEY) {
        StatelessRandomUniformV2<half> op;
        op.Init(y, &tilingData, &pipe);
        op.Process();
    } else if (tilingData.tilingKey == BFLOAT16_TILING_KEY) {
        StatelessRandomUniformV2<bfloat16_t> op;
        op.Init(y, &tilingData, &pipe);
        op.Process();
    }
}

为什么用模板而不是普通的分支判断来处理数据类型？因为 AI Core 的向量化指令一次处理多个数据元素，使用模板可以在编译期确定数据类型对应的 SIMD（单指令多数据）操作宽度，float 类型每条向量指令处理 16 个元素，half 处理 32 个元素，bfloat16 也是 32 个元素。编译器据此生成最优的向量指令序列，而运行时只需要一个简单的 tilingKey 判断来决定加载哪个模板实例即可。

Philox 算法：随机数的密码学品质与硬件友好性

ops-rand 当前唯一实现的随机数生成算法是 Philox 4x32-10。理解这个选择，有助于理解为什么算子库选择了这个特定的算法而不是常见的 Mersenne Twister 或 LCG。

Philox 算法是一种专门为硬件实现而设计的加密类伪随机数生成器，由 Sandra S. Hotz 和 James E. G. Wilcox 于 2012 年提出。它的核心结构由两部分组成：一组 4 个 32 位计数器（counter）和一个 2 元素的密钥（key），通过 10 轮循 MXQ（Multiply-XOR-Quxstep）结构产生 4 个 32 位随机数输出。每轮循 MXQ 操作的数学表达为：取两个 32 位无符号整数，将它们相乘得到 64 位结果，低位 32 位与高位 32 位分别与不同的轮常数做异或操作后交换位置，形成新的两个 32 位值。在 Philox 的语境中，这个操作被描述为一对数字"互咬"——两个数不断相互乘以对方并做异或变换，这就是"Philox"名字的含义（源自希腊语"噬咬"）。

template <uint32_t ROUNDS>
__aicore__ inline void PhiloxRandom(LocalTensor<uint32_t>& output,
    const std::array<uint32_t, 2>& key,
    const std::array<uint32_t, 4>& counter,
    const uint32_t count)
{
    // 每4个counter值产生4个随机数，结果数为count
    // 内部通过向量化方式批量处理多组counter
    // 10轮MXQ操作在循环中展开执行
}

为什么选择 Philox 而非其他算法？第一个原因在于可逆性验证。密码学类随机数算法的一个核心优势是其输出可以通过种子和偏移量精确反向推导。对于科学计算和深度学习场景，这意味着给定相同的参数组合，结果的每一 bit 都可以被精确复现和验证。Mersenne Twister 虽然周期极长，但状态转移函数的复杂性使得无法从中间结果推断出之前的状态，而 Philox 的 counter 模式天然支持任意位置的"快进"操作。第二个原因在于 AI Core 的指令集亲和性。Philox 算法的核心运算是 32 位乘法和异或，这两类操作在 AI Core 上都有高度优化的硬件支持。乘法的延迟在 AI Core 中经过特殊调优，而异或操作可以零延迟映射到向量执行单元。第三个原因在于向量化效率。Philox 的 4x32 结构天然对应 SIMD 的 4 通道处理模型，向量化指令一次可以处理 4 组计数器的运算，充分利用 AI Core 的 512 位向量宽度。

Tiling 策略：如何将亿级随机数拆分给多个 AI Core 并行处理

在大规模张量初始化场景中，一次性生成数百万个随机数是常态。如果只有一个 AI Core 来承担这个任务，即使每个元素只需要几次乘法和异或操作，总耗时也会变得难以接受。ops-rand 通过 tiling 策略将任务分配给芯片上所有的 AI Core 并行处理。

Tiling 的思想可以类比为一个印刷厂接到了一份百万张海报的大订单：如果只有一台印刷机，即使它速度极快，印刷百万张也需要数天时间；但如果有 100 台印刷机同时工作，每台只承担一万张，总时间就缩短到原来的百分之一。tilling 策略的核心就是将输出缓冲区拆分成多个"块"（block），每个 AI Core 独立处理其中一个块。

stateless_random_uniform_v2.cpp 中的 CalcTilingData 函数负责计算 tiling 参数。计算过程遵循以下逻辑：首先获取当前平台的有效 AI Core 数量（GetCoreNumAiv），然后将总输出大小按 Core 数量等分，得到每个 Core 负责的基础数据量；接着将这个数据量向上对齐到 AI Core 内存访问的最优边界（CORE_ALIGN_SIZE = 512 字节），确保每个 Core 访问的数据起始地址和大小都是 512 字节的倍数；最后还需要考虑 UB 的大小限制，如果每个 Core 的数据量超过了 UB 容量的四分之一，就需要将块进一步拆分为 UB 可以容纳的 sub-tiling 单位。

void CalcTilingData(
    size_t outputSize,
    uint32_t dtypeSize,
    platform_ascendc::PlatformAscendC* platform,
    StatelessRandomUniformV2TilingData& tilingData)
{
    auto coreNum = platform->GetCoreNumAiv();
    auto coreAlignFactor = CORE_ALIGN_SIZE / dtypeSize;
    auto blockFactor = CeilDiv(outputSize, coreNum);
    auto blockAlignFactor = CeilDiv(blockFactor, coreAlignFactor) * coreAlignFactor;

    // 计算主块大小，确保字节对齐
    tilingData.blockTilingSize = std::max(static_cast<uint32_t>(blockAlignFactor), MIN_TILING_SIZE);
    tilingData.blockNum = CeilDiv(outputSize, tilingData.blockTilingSize);

    // 最后一个块的大小可能小于标准块大小
    tilingData.tailBlockTilingSize = outputSize - tilingData.blockTilingSize * (tilingData.blockNum - 1);

    // 计算UB分块大小（UB总容量的1/4）
    uint64_t ubSize = 0;
    platform->GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
    auto quarterUbSize = ubSize / UB_DIVISOR;
    auto alignFactor = BLOCK_SIZE_BYTES / dtypeSize;
    tilingData.ubTilingSize = CeilDiv(quarterUbSize / dtypeSize, alignFactor) * alignFactor;
}

为什么要对最后一个块单独处理？因为总数据量通常不能被 Core 数量整除。以 1,000,000 个 float 元素和 1,000 个 AI Core 为例，每个 Core 理论分配 1,000 个元素，但 1,000,000 除以 1,000 得出的 blockNum 实际上是按照 tiling 策略重新计算的块数——实际部署时可能根据 UB 大小重新划分，最终一个 Core 可能负责 512 个元素而另一个 Core 负责 488 个元素。tailBlockTilingSize 记录的就是最后一批元素的数量，确保计算不会出现 off-by-one 错误。

数据类型转换：从无符号整数到浮点随机数的技术路径

随机数生成器产生的原始输出是 32 位无符号整数序列，这些整数的取值范围是 [0, 2^32-1]。但深度学习模型初始化和计算需要的是 [0, 1) 区间的浮点数。ops-rand 实现了三种数据类型转换路径，分别对应 float、half（FP16）和 bfloat16。

float 类型转换是三者中最直接的 IEEE 754 浮点数构造过程。Philox 输出的 32 位无符号整数被看作一个 IEEE 754 单精度浮点数的尾数部分，在此基础上人工设置指数段为 127（二进制 01111111），符号位为 0。按照 IEEE 754 的内存布局，构造出来的 float 值范围恰好是 [1.0, 2.0)。最后通过一个简单的减法操作（加一个 -1.0 的偏置），将整个区间平移到 [0.0, 1.0)。

这个转换过程的妙处在于：无需任何除法或乘法指令，只需要按位与（And）、按位或（Or）和一次浮点加法。AI Core 的向量执行单元对这几类指令都有专用硬件支持，可以在一个时钟周期内完成一个向量寄存器（16 个 float）的全部转换操作。

// IEEE 754 单精度格式: |1位符号|8位指数|23位尾数|
// 原始 32 位无符号整数 -> 取低 23 位尾数
// 构造 float: 尾数保持不变，将指数段设为 127（即 (127 << 23) = 0x3f800000）
// float 结果 = 构造值 - 1.0f
// 等价于: ((raw_uint32 & 0x007fffff) | 0x3f800000) - 1.0f

half 类型转换的逻辑类似，但尾数宽度只有 10 位而非 23 位，因此先通过位截断取低 10 位，再设置指数段 15（5 位指数的偏移量）。由于 half 的动态范围较窄，这种转换在大多数实际场景中已经足够精确。bfloat16 则是一个特殊的变体，它保留了 8 位指数（与 float 相同），但尾数只有 7 位。这种格式的动态范围与 float 一致，但精度略低，适合需要大幅动态范围的场景如某些正则化操作。

效率对比：ops-rand 与 CPU 生成方案的实际差异

在实际部署中，使用 ops-rand 相比传统的 CPU 生成方案，在大规模随机数场景中展现出明显的效率优势。以下数据基于典型的权重初始化场景进行估算。

对于一次完整的模型权重初始化任务，假设模型包含 1 亿个参数（FP32），需要生成 100,000,000 个随机数。CPU 生成方案的实施路径是：首先在 CPU 侧通过 std::mt19937 生成 100,000,000 个 32 位随机整数（约 400MB 数据），然后通过 PCIe Gen4 x16 总线拷贝到 NPU 设备内存，总耗时包括 CPU 生成时间（约 50-100ms）、PCIe 拷贝时间（约 5-10ms，400MB 在 PCIe Gen4 上的理论带宽约为 16GB/s，实际利用率约 50%）以及 NPU 侧的内存写入时间（约 2-3ms）。综合来看，CPU 方案的总耗时约为 60-110ms。

ops-rand 方案的执行路径则完全位于 NPU 芯片内部。首先，Host 侧通过 ACL API 启动核函数，参数（seed、offset、tilingData）通过 PCIe 发送到 NPU，总数据量不足 100 字节，拷贝开销可忽略不计。然后，1,000 个 AI Core 并行执行，每个 Core 负责约 100,000 个元素，按照 tiling 参数逐块处理。基于 AI Core 每核每秒可处理数百万次 Philox 迭代的能力估算，单核处理 100,000 个元素的时间约为 0.5-1ms，1,000 个核并行总时间约为 0.5-1ms（加上核间同步开销）。最后，结果直接从 NPU 内存被写入权重张量，无需额外的跨设备拷贝。

对比结果如下：ops-rand 方案去掉了 CPU 生成和跨 PCIe 拷贝两个高延迟环节，在相同硬件条件下将总耗时从 60-110ms 降低到 1-3ms，缩短了约 30-50 倍。对于需要周期性重新初始化的训练流程（如课程学习、课程调度等），这种效率提升的累积效果非常显著。

更重要的效率维度在于流水线的连贯性。在 CPU 方案中，NPU 必须等待随机数生成完成并完成拷贝后才能开始权重初始化计算，数据在 CPU 和 NPU 之间往返一次意味着 NPU 在这段时间内完全处于空闲状态。使用 ops-rand 后，随机数生成和后续计算可以组成更紧密的流水线：在上一批数据完成计算的同时，下一批的随机数可以在 NPU 上同时生成。固有的跨设备延迟消失后，NPU 的实际有效计算时间占比大幅提升。

编译与部署：从源码到可执行包的完整流程

ops-rand 的编译体系基于 CMake 构建，提供了高度自动化的端到端打包流程。项目的根目录包含 build.sh 脚本，这是所有构建操作的统一入口。

编译过程分为两个阶段。第一阶段是依赖准备：编译脚本会自动检测系统中是否已安装 makeself（用于生成自解压安装包的开源工具），如果未安装且处于联网环境，脚本会自动从 CANN 社区的第三方仓库下载。如果处于离线环境，开发者需要提前准备好 makeself 并放置在 third_party 目录或通过 --cann_3rd_lib_path 参数指定路径。第二阶段是算子源码编译：CMake 读取 CMakeLists.txt 中的算子模块定义，为每个算子生成独立的编译目标，然后调用 AscendC 的编译器 ascendc 将核函数源码编译为适配 AI Core 指令集的可执行二进制。

# 基本编译（使用默认 8 线程并行编译）
bash build.sh

# 编译指定算子（仅编译 stateless_random_uniform_v2）
bash build.sh --ops=stateless_random_uniform_v2

# 编译并运行单元测试
bash build.sh --run

# 编译生成安装包（.run 格式）
bash build.sh --pkg --soc=ascend950

# 多线程编译加速
bash build.sh -j16 --pkg --soc=ascend950

编译完成后，--pkg 参数会生成一个自解压的 .run 安装包，文件命名遵循 cann-${soc_name}-ops-rand_${cann_version}_linux-${arch}.run 的格式。以 Ascend 950 平台为例，生成的安装包名为 cann-950-ops-rand_9.0.0_linux-x86_64.run（或 _aarch64.run）。这个安装包可以直接分发给最终用户，无需提供源码。

# 安装到默认路径（需要 root 权限）
sudo ./cann-950-ops-rand_9.0.0_linux-x86_64.run

# 安装到自定义路径
sudo ./cann-950-ops-rand_9.0.0_linux-x86_64.run --install-path=/opt/ascend

# 查看安装包信息
./cann-950-ops-rand_9.0.0_linux-x86_64.run --help

# 卸载
sudo ./cann-950-ops-rand_9.0.0_linux-x86_64.run --uninstall

# 升级（保留配置和历史版本信息）
sudo ./cann-950-ops-rand_9.0.0_linux-x86_64.run --upgrade

安装完成后，ops-rand 的文件被部署在 ${install_path}/cann 目录中，具体来说，set_env.bash 脚本位于 $install_path/cann/set_env.bash，算子二进制位于 $install_path/cann/ops/rand_lib/ 目录下。使用前需要先加载环境变量：source ${install_path}/cann/set_env.bash。这条命令会将 CANN 运行时库路径、AscendC 工具链路径以及 ops-rand 自身的路径添加到 shell 的环境变量中，确保后续的编译和运行时能找到所需的库和头文件。

单元测试体系：如何验证随机数算子的正确性

ops-rand 自带了一套轻量级的测试框架，框架代码位于 tests/ 目录下。这套框架的设计目标是在不引入第三方测试库（如 GoogleTest）的前提下，提供基本的测试用例注册、断言判断、超时控制和结果统计功能。

测试框架的核心宏包括 TEST_CASE_BEGIN/TEST_CASE_PASS 用作测试用例的分界标记，TEST_ASSERT 用于条件判断，TEST_ASSERT_ARRAY_EQ 用于数组比较。以 stateless_random_uniform_v2_test.cpp 为例，测试用例覆盖了以下维度：

边界条件测试验证了空指针和零长度的处理逻辑。这两个测试用例确保当 API 接收到异常参数时，能够返回明确的错误码而非触发未定义行为。边界条件的测试覆盖率是算子正确性的基础保障，因为异常参数路径在实际调用中虽然罕见但并非不可能发生。

可复现性测试是整个测试体系中最核心的用例。测试逻辑非常简洁：使用相同的 seed 和 offset 调用两次算子，将两次输出数组逐元素比较。任何一处不匹配都会导致测试失败。这个测试用例验证了 ops-rand 对确定性承诺的兑现能力——在昇腾 NPU 的硬件级别上给定相同参数产生相同输出，这不是一个理所当然的特性，硬件指令流水线的分支预测、数据通路的微架构状态变化都可能引入不确定因素。

种子差异测试验证了不同种子产生不同序列的基本属性。在测试中分别使用 11111 和 22222 作为种子生成两组随机数，然后验证两组的对应元素不完全相同。这个测试的隐含前提是：如果两个不同的种子恰好产生了相同的随机序列，那意味着生成器本身的周期长度或状态空间存在严重缺陷。

偏移量测试验证了 offset 参数的"快进"功能。使用相同的 seed 但不同的 offset（0 和 100）生成两次数组，确认两组结果存在差异。这个测试确保 offset 参数确实在控制序列的位置而非被忽略或错误处理。

运行测试的方式非常直接：bash build.sh --run。执行后测试框架会打印每个用例的执行状态和最终的统计结果。成功的输出格式为 100% tests passed, 0 tests failed。

未来路线图：从均匀分布到完整的随机数工具箱

ops-rand 当前处于 Phase 1 阶段，仅提供了最基础的均匀分布随机数生成能力。根据 implementation.md 中记录的路线图，后续发展分为四个阶段。

Phase 2 将扩展数据类型和分布类型。数据类型方面，将新增 FP16（half）和 FP64（double）的原生支持，其中 FP64 的支持需要确认目标 SoC 是否具备双精度浮点运算能力。分布类型方面，将实现正态分布（高斯分布）和对数正态分布，这两者在神经网络权重初始化（如 Xavier/Glorot 初始化）和变分推断等场景中有广泛需求。正态分布的实现预计将使用 Box-Muller 变换算法，将均匀分布的随机数转换为正态分布的随机数。

Phase 3 将扩展随机数生成器引擎种类。除了当前的 Philox 4x32-10，还将引入 XORWOW（一种基于 XOR-shift 的生成器，周期较短但计算速度极快）、MRG32K3A（多重递归生成器，具有极长的周期）、MTGP32（并行 Mersenne Twister 的硬件优化版本）以及 Sobol 准随机序列生成器。Sobol 序列的优势在于其低差异性（low discrepancy），在蒙特卡洛积分和金融衍生品定价等场景中，使用 Sobol 序列比使用伪随机序列能以更少的采样点达到同等的数值精度。

Phase 4 将开放 Device API 级别的接口，允许 AscendC 算子开发者在自己的核函数内部直接调用 ops-rand 的随机数生成能力。这意味着未来的自定义算子也可以使用确定性随机数，而无需在算子外部预先生成再传入。这对于需要在内核中进行随机 dropout、随机量化等操作的场景特别有价值。

应用场景实践：从权重初始化到蒙特卡洛积分

ops-rand 的应用场景远不止权重初始化一个。以几个典型的深度学习场景为例说明其使用方式。

分布式训练断点恢复。在多卡分布式训练中，每个设备的随机状态需要被精确记录以便从检查点恢复。使用 ops-rand，可以通过一个全局的 seed 加上设备 ID 和当前 step 序号构造出一个唯一的 (seed, offset) 组合来保证每个设备和每个 step 的随机状态都是确定且可恢复的。

对抗样本生成。生成对抗样本需要对输入图像施加微小的扰动，随机扰动的方向和幅度通常需要精确控制。使用固定的 seed 配合不同的 offset 生成随机扰动向量，可以确保对抗样本生成过程在相同的输入下产生相同的结果，这对于攻击可复现性和防御方测试都有重要意义。

蒙特卡洛物理模拟。在科学计算场景中，蒙特卡洛方法需要大量随机数进行积分估计。使用 ops-rand 在 NPU 上直接生成 Sobol 序列并就地完成积分计算，可以省去跨设备数据传输的开销。结合 Phase 3 计划中的 Sobol 序列支持，这种应用将获得更高效的实现。

金融风险模拟。期权定价和风险价值（VaR）计算需要在大量路径上进行蒙特卡洛模拟，每条路径依赖的随机数序列要求可审计可复现。使用 ops-rand 给定相同的风险因子 seed，可以精确复现某一次风险评估的结果，满足监管合规和内部审计的要求。

这些场景的共同特征是：对随机数有"确定性"的要求——不是要求随机性越强越好，而是要求随机性在可控可复现的前提下尽量高效。ops-rand 的设计正是围绕这一核心需求展开的。

项目地址为 https://atomgit.com/cann/ops-rand