TensorPlay 1.1.0 Release Notes
- 亮点 Highlights
- 不兼容变更 Backwards Incompatible Changes
- 弃用 Deprecations
- 新特性 New Features
- 改进 Improvements
- 问题修复 Bug fixes
- 性能 Performance
- 文档 Documentation
- 开发者 Developers
- 安装 Installation
亮点 Highlights
- omni_attention 可编程注意力:
score_mod/mask_mod编程接口 +BlockMask块稀疏编译;CPU 融合 flash 内核(分块 + 线程池)与 CUDA fp32 3xTF32 融合前向/反向 原生实现,eager 未融合回退与编译降级路径齐备。 - 图编译管线全面收口:注意力调用自动降级为生成内核;AOT 编译产出 C++ wrapper 并在运行时构建加载;CuteDSL 后端自研化(移除上游依赖);单次编译耗时从 11.3s 降至 1.3s。
- CUDA graphs 编译器后端:
tensorplay.compile新增mode="reduce-overhead", 编译产物经 CUDA graphs 回放;GEMM 走 TunableOp 调优上下文,结果持久化。 - 前向模式 AD 原生引擎:dual-tensor 引擎单趟传播切向,
autograd.functional.jvp默认改走前向模式,jacfwd随之可用。 - API 面大幅扩展:新增
audio、vision、hub、onnx、package、ao量化、masked、signal、nested、monitor、func控制流算子、meta设备、稀疏分发键等 一大批公开面。 - 随机数工厂接入 Generator:
randint、poisson、multinomial与多种就地采样 把调用方 generator 从 Python 层一路传到内核。 - CUDA 内核性能批量刷新:
unique/一维排序/三角掩码/外层维归约/直方图/嵌入梯度 等重写,多组负载从落后数十倍收敛到持平或领先。 - 构建与发布增强:cutlass / flash-attention 收为 git 子模块,wheel 携带 JIT 运行时头文件,CPU 轮子恢复 PyPI 发布。
- 白皮书新增图编译管线章节:补上 08b 编译前端、08c stax 降级、08d 原生图运行时
三章,以
path:line援引具体机制。 - 现代解码栈进入模块面:
RotaryEmbedding/apply_rotary_emb、GroupedQueryAttention、MultiheadLatentAttention、GatedDeltaNet/KimiDeltaAttention(含 chunked 与 recurrent 两套规则)、CompressedSparseAttention、CrossAttention/CausalEncoderDecoder、SwiGLUMLP/TransformerBlock;KV 缓存补齐StaticCache/SlidingWindowCache/PagedCache/QuantizedCache,配套tensorplay.nn.generation.generate自回归循环 (top-k / top-p 过滤 + 重复惩罚)。 - 高阶梯度大面积补齐:
config/derivatives.yaml新增 172 条可微条目(净增 161), 覆盖卷积 / 池化 / 分数池化 / ROI / 填充 / 上采样 / 嵌入 / FFT / STFT / 线性代数 / 特殊函数 / 损失 / 激活 / 稀疏 softmax / CTC 的反向再可微;_segment_reduce、 稀疏sum与稀疏 softmax 的 COO 分支同样可再求导。 - 每个可微算子都有前向模式切向:原先只声明反向公式的算子补上
result:切向;gradcheck新增前向模式、批处理与 fast 模式校验(双引擎 Jacobian 对中心差分, 批模式经 vmap 校验 Jacobian-向量积与向量-Jacobian 积)。 - 频域与线性代数补全:原生多维 FFT 原语(
_fft_c2c/_fft_r2c/_fft_c2r)与 Hermitian 族 API,fft/stft全部可再求导;linalg的qr/lstsq(CUDA 侧gels走 QR 驱动,高、瘦、宽三种形态)、lu_factor/cholesky/solve/pinv及其伴随均已可微。 - 编译器原生下沉成规模:
op_lowerings.py新增约 5.5k 行,norm / softmax / 激活 / 视图 / 池化 / 上采样 / 散射 /searchsorted/bucketize/ldexp/unfold/prelu从"库调用"改写为"区域内生成",新增tensorplay._decomp分解表(618 条分解, 每条同时注册函数式与就地两种拼写)供后端沿分解路径下沉。 - CPU 指令集分层向量化:AVX2 与 AVX512 双路径并行铺开(广播算术、dot 归约、
clamp/lerp/power、三角行拆分、标量算术旁路),channels-last 池化 / 上采样 / grid sampling / bicubic 走 tier-compiled 专用通道。 - CUDA 归约 / softmax / 排序重写:非对齐输入头尾分离 + 连续输出 vec4/vec2/标量
分派、
hash_tensorXOR 归约、min/max 合并为单趟打包键归约、Welford 分块、空间维 softmax 与宽行尾部、half→float 前向、大张量中位数走基数排序、超大排序分块、cummax网格按设备上限封顶。 - 导出(export)按符号尺寸推导:捕获期在克隆的样例值上跑一遍每个记录调用,按逐算子 形状规则经符号输入维度推导计算值的维度范围(与样例具体形状交叉校验),规则在符号 范围上做的决定变成读输入形状的运行时守卫;声明为 dynamic 的维度在捕获期是符号本身, 读它得到图值而不是被样例尺寸定死。导出从此知道程序内部算出来的形状(两项不兼容变更)。
- ONNX 导出的动态轴打通:
dynamic_axes进入捕获(成为Dim.DYNAMIC捕获规格),dynamic_shapes按可调用对象签名映射回 ONNX 输入;计算值上的.shape/.size()/.dim()/.numel()/.ndim/.T/.dtype读取降到Shape/Gather运行时 读数,b, n, c = x.shape这类解包不再把样例 batch 烤进图里。 tensorplay._assert:把 Pythonassert作为可调用算子暴露(随包的vit_b_16/maxvit一直在调它)。符号跟踪会记录它,编译捕获把张量值上的断言转成程序实际执行的_assert_async.msg检查,导出按"尺寸断言"与"数值检查"分别处理;检查类算子登记为 side-effectful,死代码消除不再把它们删掉(此前编译后的_assert_async从不执行)。- profile 表格补齐捕获列:
profile.table()逐 span 渲染,record_shapes、with_stack、输出显存、输入形状、FLOPs、源码位置等此前采集后即丢弃的列都能出表;header/time_unit/top_level_events_only生效,嵌套 span 的设备时间与显存不再 汇总成 0。
不兼容变更 Backwards Incompatible Changes
注意力
注意力入口统一为 omni_attention
注意力编程入口的公开命名全部收口为 omni_attention:模块路径、配置项与钩子统一到该
命名;上下文并行集合算子注册为 cplib::omni_cp_allgather /
cplib::omni_cp_allgather_backward,命名旧注册名的已存图需重新编译。
from tensorplay.nn.attention import omni_attention
自动微分
layer_norm 反向改为保存统计量
layer_norm 现组合自 native_layer_norm(输出 + 每行 mean / rstd),autograd 节点保存
这两份统计量并由 native_layer_norm_backward 消费,反向不再从输入重算矩。六参
layer_norm_backward 算子保留但不再位于 autograd 路径上。
functional.jvp 默认前向模式
jvp(func, inputs, v) 的 mode 默认值由 "reversed" 改为 "forward":默认经
dual-tensor 引擎逐算子传播切向,不再依赖 double-backward 技巧;mode="reversed"
仍可用,create_graph=True 仍走反向路径。
导出
声明为 dynamic 的维度在捕获期保持符号
捕获过去把每个输入维度都按样例定死,于是 dynamic_shapes={'x': {0: Dim('batch')}}
下读 x.shape[0] 的程序被录成 x + 2,在任何非样例 batch 上都是错的;对该维度的分支
也只判一次就忘了条件。现在声明为 dynamic 的维度代表一个符号(ShapeEnv 区间,SymInt
携带样例作为提示),读它得到图值 getitem(getattr(x, 'shape'), axis),其上的算术保持
符号;对符号做 Python 决定(分支、int()、len()、range())按提示作答并留下守卫。
捕获后逐条校验守卫:区间蕴含的丢弃、只排除小于 2 的(如 x.shape[0] != 1)转成运行时
断言、落在具名 Dim 上的拒绝并给出 refine_dynamic_shapes_from_suggested_shapes 能
应用的修复、把 Dim.DYNAMIC 固定成一个尺寸的拒绝、其余(dim 提示)收窄并断言。
dynamic_shapes 现在先按可调用对象自己的参数归一化——模块带参数时,序列规格过去会
绑到第一个被提升的权重而不是第一个入参。
行为变化:export(..., dynamic_shapes=...) 在程序收窄或固定一个具名 Dim、或固定一个
Dim.DYNAMIC 维度时抛 ConstraintsExceededError,不再静默特化到样例尺寸。照提示改用
refine_dynamic_shapes_from_suggested_fixes、把该维度声明为 static,或用 Dim.AUTO
让导出自行特化。
计算值的形状按逐算子规则推导并落守卫
导出过去只知输入的维度范围,程序内部算出来的每个值都回读固定的样例形状,于是在动态
程序里对 h.shape[0] 分支、或据此 reshape 的地方把样例尺寸静默烤死。捕获现在在克隆
的样例值(克隆输入、缓冲区与常量;共享参数在就地写前克隆;随机生成器 fork)上跑每个
记录调用,经逐算子形状规则跨符号输入维度推导每个结果的范围,并与样例具体形状交叉校验;
规则在符号范围上做的决定成为守卫,导出为读输入形状的运行时断言;广播与 slice 钳位对
尺寸无关,不额外加守卫。没有规则、或尺寸由张量数据决定的算子给出程序可读但不可据以
分支的范围,并报出它来自哪里。
溯源记录程序从哪些输入读了固定范围;ONNX 导出器对这些轴上的 dynamic_axes 直接拒绝,
而不是导出一张在任何其他尺寸上都错的图,并在降级期间跳过运行时检查。
行为变化:
- 形状规则开始记录守卫:合并具名维度不同的输入(例如相加)抛
ConstraintsExceededError并给出建议修复;Dim.AUTO程序为其假设的关系带上运行时 断言。 - 以固定尺寸捕获后再做 ONNX 导出,若
dynamic_axes标记的轴正是程序读过的轴,抛ValueError;请传模块本身(或声明dynamic_shapes)而不是预导出的静态程序。 - 在前导维度固定的张量上迭代会展开;由动态维度定尺寸的则抛错。
from tensorplay.export import export
from tensorplay.export.dynamic_shapes import Dim, refine_dynamic_shapes_from_suggested_fixes
内核契约
索引类算子契约与 ABI 对齐
index_put/index_put_的索引参数改为Tensor?[](可含 None);C++ 调用方需传std::vector<std::optional<Tensor>>。rrelu_with_noise的 noise 缓冲改为可变入参并接受 generator;rrelu/rrelu_接受 generator;rrelu_with_noise_增加导数。- 默认
slice重载删除,统一走slice.Tensor。 - stub 签名中拥有型参数(
Scalar、Tensor?、str、列表)改为 const 引用传递; 自定义算子内核声明随之更新。 out=变体统一经write_out(resize + dtype/device 校验)写目标,不再重绑定目标。
CUDA
cuDNN benchmark 默认关闭
cuDNN benchmark 开关的默认值改为 False:默认取启发式首选配置,
避免每个新形状先实测多份候选计划(实测落到的引擎可能比启发式首选慢近两倍)。张量
描述符对齐度改为按存储地址计算(1..32 字节),不再写死 16。
unique / 一维排序 / 三角掩码重写
unique按需计算逆映射与段长(此前 Python 层写死请求后丢弃);不需要逆映射时只排键, 不需要段长时走无游程表路径。- 一维排序改走全局基数排序;
unique改为一趟UniqueByKey,段表只随输出大小增长。 tril/triu三角掩码按列铺开到块内,窄行回退一行一线程。- 数值保持逐位一致(76 组 unique 用例、119 组 sort 用例核验)。
跨维 softmax 竞态修复
dim 非最后一维的分块核在读出行最大值后缺少屏障,慢线程会在部分和写入后误读
tile[0],导致整行被错误缩放(bf16 (512, 8192) dim=0 log_softmax 曾观测到单列误差
2.0)。现已在读最大值后补屏障;同时修复寄存器驻留核空 lane 的 exp(-inf - -inf)
产生 NaN 的问题,并把波形核行长上限调低。
前端
卷积 padding 改用模式词
conv1d / conv2d / conv3d 的 padding 参数接受 'valid'(不填充)与 'same'
(保留每个输出位置)。'same' 只在步长为 1 时存在,请求覆盖卷积核的触及范围,
奇数格留在尾部一侧;因逐轴计数对两侧对称,尾部多出的量预先补到输入前部,卷积
保留前部计数。其余取值仍按边界词解释并被拒绝。
block_diag 落到原生内核
block_diag 的 Python 组合实现被原生内核取代(分块放置与周边零的分配都在操作数
所在设备上完成);Python 结果 dtype 辅助函数改按 C++ 侧的次序排序操作数
(带维张量 > 0 维张量 > 普通数字,数字按其包装数类型计);meshgrid 导出到包根。
弃用 Deprecations
无。
新特性 New Features
注意力(omni_attention)
omni_attention可编程注意力入口:score_mod(五元坐标(score, b, h, q_idx, kv_idx)任意表达式)与mask_mod谓词编程,BlockMask块稀疏编译形态。- 掩码工具:
create_block_mask/create_mask/and_masks/or_masks/noop_mask。 - 输出选项:
return_lse(对数求和指数)与return_aux(反向复用辅助输出AuxOutput)。 - eager 未融合回退实现(显式物化完整分数矩阵,运行时给出融合指引),
tp.compile捕获后降级为融合内核。 - CPU 融合 flash 内核:查询轴分块 + BLAS gemm 分片、向量化行 softmax、跨 intra-op 线程池并行,前向/反向均支持;query head 数可为 key head 的倍数,mask 可广播。
- CUDA 融合内核:fp32 flash 前向(3xTF32 tensor core,head 64)、混合精度 SDPA 反向、宽 tile 双缓冲、按 query 长度选 tile、因果对角直接声明(非方形因果不再走 组合路径)。
- 上下文并行(context parallel)集合算子与反向梯度规则,配套
cplib::omni_cp_*注册名。
现代解码栈(tensorplay.nn)
- 位置编码:
RotaryEmbedding模块与函数式apply_rotary_emb。 - 注意力模块:
GroupedQueryAttention(GQA)、MultiheadLatentAttention(MLA)。 - 线性注意力混合器:
GatedDeltaNet/GatedRMSNorm(gated delta rule)与KimiDeltaAttention,两者都提供 chunked 与 recurrent 两套规则函数 (chunk_gated_delta_rule/recurrent_gated_delta_rule、chunk_kimi_delta_rule/recurrent_kimi_delta_rule)。 - 压缩稀疏注意力:
CompressedSparseAttention(窗口压缩器 + 闪电索引器)。 - 编解码器与块:
CrossAttention/CausalEncoderDecoder(CED 脚手架)、SwiGLUMLP/TransformerBlock。 - KV 缓存:
StaticCache(预留定长)、SlidingWindowCache(滑窗 + 窗口长度查询)、PagedCache(分块分配、按需扩容)、QuantizedCache(int8 分组量化存取), 与既有的DynamicCache共用Cache协议。 - 自回归生成:
tensorplay.nn.generation.generate,含 top-k / top-p 过滤 (top_k_top_p_filtering)与重复惩罚。
频域与线性代数
- 原生多维 FFT 原语:
_fft_c2c/_fft_r2c/_fft_c2r(不带归一化的多轴变换),fft/ifft/rfft/irfft及其n维变体(fftn/rfftn/irfftn)与 Hermitian 族(hfft/ihfft及其 2 维 / n 维变体)直接建立在这三个原语上, 不再由 Python 层递归分解。 fft/ifft/rfft/irfft/fft2系列与stft的反向全部可再求导。linalg因子分解与求解器可微:qr、lu_factor/lu_factor_ex/lu_solve、cholesky/cholesky_ex/cholesky_solve/cholesky_inverse、solve/solve_ex、solve_triangular、det、inv_ex、pinv(含 atol/rtol 张量变体)。- CUDA
lstsq走gels驱动(QR 分解,高 / 瘦 / 宽三种系统形态分别取路),残差从 列主序工作缓冲读取。
编译器(tensorplay.compile / Stax)
- 注意力图→核降级:从捕获的注意力调用中识别 score/mask 语义,选择 flash 或逐位置
生成的 CPU/CUDA 内核;
var走单趟 Welford 归约。 - AOT 收口:
partition_default/partition_min_cut分割后生成 C++ wrapper 并在 运行时构建加载;build产物可命名导出二进制位置与 launch 自述信息。 - CuteDSL 后端自研化:模板依赖替换为自有后端;
cutedsl内核编写、lane 分析、注意力 模板覆盖。 - 调优体系:候选内核跨进程测量、决策持久化、
TUNING_VERSION决策盐、bench_candidates交错轮询抗噪;tunable 上下文经tensorplay.cuda.tunable暴露。 - 融合与布局:channels-last 布局沿卷积链传播、group-norm backward 模板核、默认融合 内存阈值提升至 10、IO 缓冲上限提升至 32。
- 控制流算子下沉:
cond/while_loop/map/scan前端支持。 - 原生下沉扩面:norm / softmax / 激活 / 视图 / 池化 / 上采样 / 散射族 /
searchsorted/bucketize/ldexp/unfold/prelu/split/ 排名分位数 / 按样本归一化由"库调用"改写为"区域内生成";op_lowerings.py新增约 5.5k 行。 - 后端分解追踪:新增
tensorplay._decomp分解表(618 条分解),编译器可沿 后端分解路径下沉算子,分解本身带导数。 - 区域与状态:
cond/while_loop区域按其实际执行的片段求导;编译模块状态经 其表格读取,不再每次调用重建;allclose/equal从各自调用的操作数求值。 - 采样与池化:自适应采样 / 重采样 / 池化进区域,抽样与重采样按源映射读取。
- 卷积路径:fp32 channels-last 卷积可走 TF32;降精度卷积的操作数按 channels-last
交给内核;
einsum在捕获期按公开形态记录。
CUDA
tensorplay.compile新增cudagraphs后端:mode="reduce-overhead"将 stax 编译 产物包装为 CUDA graphs 回放。- TunableOp:GEMM 调优上下文、结果持久化(文件读写)、环境变量种子与失效候选自愈; cuBLASLt GEMM 选择接入调优上下文。
- NVRTC 运行时编译 elementwise 内核(
jiterator风格),并扩展到广播布局。 Tensor.record_stream:分配器流跟踪。- cudart profiler 绑定与
tensorplay.cuda.profiler面。 - 新内核:embedding renorm、fused RMSNorm、batch norm 专用内核。
- 注意力内核:fp32 flash 前向(3xTF32)、head 64 全宽度、宽 tile 双缓冲、channels-last 梯度直读(卷积反向不再先拷贝成 contiguous)。
- 归约引擎修订 14:非对齐输入的头部与向量主体分离、连续输出按 vec4 / vec2 / 标量 分派(归约长度 < 32768 走标量输出以控合并成本)、输出分组只要求输入布局连续读取、 超出 64 个 iterator 维度时的动态元数据。
hash_tensor归约:CUDA XOR 归约内核实现(此前无内核可解析)。- 稀疏 softmax:
SparseCUDA后端新增softmax/log_softmax(稠密路径复用同一 出口壳层)。 - softmax 专项:空间维(spatial)softmax、宽行尾部处理、
half_to_float前向、 掩码 softmax 融合、宽 softmax 反向向量化;波次分层上限降至 1024。 - 排序与扫描:大张量平坦中位数走基数排序、超大排序分块并去掉小排序的暂存、
扩宽 cub scan 覆盖并并行化复数与反向路径、
topk多块路径在 tile 超出 uint32 时 回退、cummax最内层网格按设备maxGridSize封顶。 - 其他:min/max 合并为单趟打包键归约、
median之外的小规模全归约免迭代器、 foreach 归约资格单趟判定、foreach 启动批量在 CUDA 13 工具链上放宽、like分配按张量实际 stride 排布、转置左操作数的 GEMM 就地读取、 float32 卷积在普通引擎上不走 TF32。
前端
tensorplay._assert:Pythonassert的算子化入口。随包的vit_b_16/maxvit一直调用它,此前在 eager 下就会失败。符号跟踪记录断言;编译捕获把张量值上的 断言转成程序执行的_assert_async.msg检查(样例为假时行为与 Pythonassert一致); 导出对尺寸断言按其他守卫处理、对数值断言按普通检查处理。_assert、两个 assert 算子 与sym_constrain_range登记为 side-effectful,死代码消除不再删掉它们——此前编译后的_assert_async从不执行。ONNX 降级丢弃检查以及只为其存在的计算。profile.table()补齐捕获列:逐 span 一行,渲染record_shapes、with_stack、 输出显存、输入形状、FLOPs 与源码位置等此前采集后即丢弃的列;header/time_unit/top_level_events_only生效,长名字按给定宽度截断并保持列对齐,FLOPs 整表统一单位,首帧以下的调用栈单独成行。key_averages()仍是聚合视图,两者共用同一 套布局。onnx.export的dynamic_axes进入捕获:输入上的dynamic_axes成为Dim.DYNAMIC捕获规格(经input_names或实参命名);反向映射按可调用对象签名把每个 实参对到它的 ONNX 输入,带提示的轴命名为<argument>_dim_<axis>。- 多媒体:
audio子包(WAV 原生编解码、FFmpeg 多格式加载、SIMD 音频转换)、vision批量 PNG 编解码、ROI / NMS / deformable conv 算子族。 - hub:
list/help/load完整实现,缓存固定到 commit,下载文件校验与 MEGA 工件来源验证。 - 导出与打包:
onnx导出公开 API 与自定义算子支持;package导出/导入器。 - 量化:
ao量化包(eager 工作流、剪枝、ao.ns数值比较套件)。 - 兼容层:
tensorplay._numpyNumPy 兼容层、numpy 关键字别名与dtype=None绑定接受。 - 分布与信号:
distributions包移植;signal窗函数模块与 STFT 往返测试。 - 张量扩展:
masked张量模块、nestedjagged 张量(构造/元数据/padding)、meta设备(仅形状)、sparse专用稀疏分发键、foreach批量算子面。 - 监控:
monitor事件、窗口统计与等待计数器。 - JIT/AOT:
tensorplay.utils.cpp_jit(load_inline/build_inline/load/build/load_module/system_lib),FFI 引擎原生加载,Windows 同样可用。 - 随机:
randint/randint_like/rand_like/poisson/multinomial/exponential_/geometric_/log_normal_与 normal 系列接受调用方Generator。 - 线程池:
get_num_interop_threads/set_num_interop_threads入口。 - 设备归因:
tensorplay.accelerator设备接口;device_count/current_stream等按当前加速器路由。 - 形状与数值辅助:
meshgrid导出到包根;包装数结果类型统一(见不兼容变更)。
自动微分
- 原生前向模式 AD 引擎:dual-tensor 表示、
jvp/jacfwd单趟传播。 autograd.functional.jvp默认mode="forward"。- 全算子前向模式切向:每个可微算子都带
result:切向公式(原先只声明反向的算子 补齐前向),gradcheck用 dual 张量算出的前向模式 Jacobian 对中心差分校验。 - gradcheck 覆盖面扩大:
check_forward_ad/check_batched_grad/check_batched_forward_grad与fast_mode(把两个 Jacobian 投影到随机向量而不完整 构造);gradgradcheck接受check_fwd_over_rev与批处理 / fast 模式;复数之外的 实数输入一并校验,稀疏输入经其稠密等价物校验。 - 高阶梯度补齐(新增 172 条可微条目,净增 161):
- 卷积族(
conv1d/2d/3d、conv_transpose1d/2d/3d的 grad_input / grad_weight / grad_bias)、im2col/col2im反向; - 池化族(
max_pool2d/3d_with_indices、avg_pool2d/3d、adaptive_avg_pool2d/3d、adaptive_max_pool2d/3d、fractional_max_pool2d/3d、max_unpool2d/3d)反向,adaptive_max_pool3d反向按保存的 indices 做散射; - 分数池化
roi_align/roi_pool/ps_roi_align/ps_roi_pool反向; - 填充族(
constant/reflection/replication/circular的 nd 与 1d/2d/3d)反向; - 上采样族(
nearest/exact/linear/bilinear/bicubic/trilinear及_aa变体)反向; - FFT / STFT 反向可再求导,多维原语自带梯度;
- 线性代数因子分解与求解器的伴随(见"频域与线性代数");
- 特殊函数(Bessel j/y 与 i0/i1/i1e、k0/k1、
zeta、igamma/igammac/gammainc/gammaincc)、ldexp、_trapezoid/_cumulative_trapezoid、median.dim/nanmedian.dim、cdist/pdist; - 损失族(
smooth_l1/huber/binary_cross_entropy/l1/kl_div/hinge_embedding/soft_margin/poisson_nll/nll_loss/nll_loss2d/multi_margin_loss/multilabel_margin_loss/mse_loss)反向可再求导; - 激活族反向(
gelu/elu/leaky_relu/softplus/hardswish/hardsigmoid/log_sigmoid/hardshrink/softshrink/glu/sigmoid/tanh/threshold/hardtanh/rrelu_with_noise/softmax/log_softmax)可再求导; - 嵌入族(
embedding_dense/_embedding_bag_dense/ 带per-sample-weights 变体) 反向; - 视图族(
unfold/diagonal/permute/squeeze/inner/as_strided)与fmax/fmin/clamp/eq_等就地与比较拼写; - 稀疏
_sparse_softmax_backward_data/_sparse_log_softmax_backward_data/_sparse_sum_backward/sparse_add反向可再求导(COO 紧凑池和形式); _segment_reduce_backward反向可再求导;ctc_loss动态规划与已保存反向输入可再求导;- 就地拼写与其函数式孪生同记(
_alpha_dropout_等)。
- 卷积族(
- backward 节点可保存符号尺寸(SymInt)与缺省输出;
save_for_backward上下文存活期 收紧为数据所需。 - 自定义算子调用折叠为一次原生穿越(native crossing),并跳过无需梯度的图节点。
- layer norm 统计量经原生算子保存与消费(见不兼容变更)。
- 评估与求值纪律:反向公式只为本次 pass 需要的输入求值;缺失的输出梯度按节点槽位 元数据物化(落在该输出的设备上);导数式遍历与设备转换在 dispatch 变换下整块记录; 保留的梯度只向 backward 的输入累加;变换与设备状态在反向 worker 上保留; 释放后仍保留边、已释放的图被再次遍历时报错。
内核
- 新算子族:
roi_align/roi_pool/ps_roi_*/nms/deform_conv2d、_masked_softmax(前向/反向、拒绝项置零)、_segment_reduce前向 + CUDA stub、quantized_linear(含动态变体)、量化激活/池化/卷积 CPU 内核、_embedding_renorm。 - 嵌套 jagged 访问器、构造器与反向内核;
_C层Tensor内存重绑定。 - 融合注意力契约(
scaled_dot_product_attention族)原生前向/反向实现,覆盖所有 服务 head 宽度。 - 复类型
sgn求值。 - 原生多维 FFT:
_fft_c2c/_fft_r2c/_fft_c2r在 CPU 与 CUDA 两侧直接实现, Python 层的递归分解退场。 - 线性代数:
qr给出有效反射打包与完整三角 R;CUDAlstsq的gelsQR 驱动。 - 生成内核暴露:
bit_cast与特殊数学函数对生成的 CPU 内核可见;quantile/nanquantile注册为组合内核。 - 批量算子:一元逐元素算子、共轭与填充在
vmap下批处理;复数 CUDA 张量与实数 标量比较;打印覆盖每种 dtype。
改进 Improvements
- 前端热路径:算子句柄查找去注册表互斥锁(开放寻址槽数组 + release/acquire)、
tensor 视图 / 属性 / 下标经 fastcall descriptor 与 getset 表解析、
t/permute/expand/unfold/eye等工厂下沉到 C trampoline、按实参个数裁剪重载候选、 变参折叠缓冲携带 kwargs、模式栈与设备键原生探测、工厂 trampoline 绑定到模块自身、 类型无钩子答案跨调用复用、norm 只在必要时跨设备搬运、OpenMP worker 自旋在算子 突发之间封顶。 - autograd:自定义 Function apply 融合为一次 C 穿越;
expand_as反向把展开维求和; dtype/device 转换与其他算子一样被记录;梯度跨流排序;软/对数软最大化反向复用已存 输出(不再重读输入)。 - 编译器:图构建从 O(n²) 计数改为单节点维护(编译 11.3s → 1.3s 的主因);外部内核
输入拷贝按签名共享;只读请求共享同一份排列拷贝;
evalf支持IntInfinity; 进区域前的主机侧工作削减、区域在降级图作用域内运行、持久化捕获缓存默认关闭、 训练用 batch norm 由生成而非回调、编译模块状态经其表格读取、外部状态每次调用读取 且区域内的就地更新保留。 - 编译器资源生命周期:编译反向按需消费已保存输入而不是替整段保存值兜底;捕获结束即释放 跟踪输入并把 fake 值的追踪改为弱引用(区域导出期间不再钉住样例);AOT 降级后释放已 记录的激活元数据,长驻的编译模块不再随编译次数增长占用。
- CPU 内核:OpenMP 团队按性能核命名封顶、免每循环 resize;channels-last 卷积走 NHWC 原语;1x1 卷积梯度合一为单次 gemm;oneDNN 引擎上限尊重显式线程数;降精度无分组 卷积按通道主序执行。
- 调度:捕获守卫仅在活动 tracer 时开启;
if constexpr化散射分发。 - CPU 向量化(AVX2 / AVX512 双层):广播算术与 dot 归约按 ISA 层级向量化、
clamp走 AVX2 并去掉masked_fill的多余拷贝、lerp/power单趟向量化、 标量算术在 AVX512 旁走 AVX2、三角行拆成零段与拷贝段、末维极值用向量 lane、 gather 行遍历免坐标除法、匹配布局的稠密二元运算按存储序流式、dot / vdot 的部分和 按线程合并、内层行和在寄存器内折叠后标量存储、leading-dim 归约的 lane 直写新输出、 方差归约拆并行 Welford 块、整张量矩用向量 Welford 链、leading-dim 求和按任务行块 拆分(小行数走流式)、逐元素 pass 在保留的稠密布局上平铺、strided 一元不物化中间 结果、连续 CPU fill 在粒度阈值以上并行、bitwise/distance/arange按 CPU 层级编译、窄浮点均值在单精度下计算。 - channels-last 专用路径:池化 / 上采样 / grid sampling / bicubic / 自适应 3d 池化
走 tier-compiled 专用通道;max-pool 扫描致密化并补齐 2d / 3d 散射;
avg_pool2d反向与自适应平均池化反向同样分层编译。 - CUDA 归约:向量 Welford 分块归约、小规模全归约 max/min 不经迭代器、
var_mean/std_mean的整数修正走融合 pass、生成 CPU store 补向量原子加、 cub scan 覆盖扩宽。
问题修复 Bug fixes
- 导出与捕获:导出任何
nn层都得到调不起来的程序(图把占位符改名却保留原名作 target,KeyError 'input',module()(input=x)同样失败);捕获期在计算值上读dim()/numel()把绑定方法交给程序(Tensor * builtin_function_or_method);符号尺寸上的size != 1经bool()就地定死(SymNode缺__ne__,所有保留守卫的不等式全部 丢失);导出 LSTM / GRU / RNN 报 "cannot create weak reference to 'Proxy' object" (Proxy.__slots__缺__weakref__,模块对已跟踪权重持弱引用);带proj_size的 LSTM 首步即失败(门按hx宽度切分,而投影后门是细胞状态宽度);被禁用的模块递归触发 自身编译;ONNX 导出计算值上的形状读取此前无降级(改为Shape/Gather运行时读数、reshape接受运行时尺寸、二元运算按节点产出的 dtype 计算,整数张量乘0.5不再被 截断)。 - 自动微分:autocast 上下文退出时未清理转置后的权重(按缓存生命周期补测试);值级引擎
trace(
TP_ENGINE_TRACE=3)用toDouble()打印每个被跟踪元素,携带复数梯度的反向在 追踪期直接抛 "Cannot convert a complex Scalar to double"(改按re+imj打印);CPU 构建 仍编译 CUDA 事件簿记。 - CUDA(图内存池 / 设备端断言 / 算法工作区):
_assert_async在 release 构建里形同虚设 (打印后调assert(false),而NDEBUG把它编掉,检查失败后程序继续跑;改为像索引区间 守卫那样 trap,故障在下一次同步处浮现);图内存池取出块后只从尺寸桶移除、仍留在地址索引 里,释放物理相邻块时会与仍存活的块合并再停放,导致下一次捕获期分配别名到活跃张量 (UNet 经make_graphed_callables报 "unsupported memory overlap between output and input",较小网络静默产生 NaN 梯度;graphed-callables 的 eager 引用此前调用被替换的 forward,拿图跟自己比);GEMM 工作区按所选算法设上限;V8 conv planner 的 cuDNN 前端头 文件加包含守卫;空 welford NaN 标量改名以避开 Windows cu130 工具链的 nvcc 解析错误。 - 内核与跨平台可移植性:
var_mean的REGISTER_DISPATCH写在匿名命名空间里,类模板静态成员的 特化不能在那里定义,各 tier 全部编译失败;lerp_tensor_fast住在 x86 段却被 tensor-lerp 无条件调用(arm64 与 MSVC 断链)、TP_NO_FP_CONTRACT为 MSVC 无法解析的-ffp-contract=off属性补上开关;BIN_*枚举与 AVX-512/AVX2 dot 累加器定义在 x86 段内 却无限定引用(arm64 未声明);MSVC 的/openmp不接受#pragma omp simd;arm64 设备 绑定里 SVE 探针带了tensorplay::cpu::限定。 - 前端:numpy 标量无法进入 number 类型槽、
like工厂选项未记录为关键字、随机工厂 忽略 generator、expand_as反向展开维未求和、已知符号尺寸未作为普通数字返回、 变参折叠缓冲丢失关键字值、layer_norm未路由到原生算子、MultiheadAttention取batch_first、自注意力保持单次投影、卷积操作数与 groups 不符时报错、layer_norm仿射项形状不符时报错、符号位运算标量搬到张量所在设备、var_mean/std_mean接受修正参数与无维调用、undef张量在绑定处被拒绝、 带号码取重载优先于无尺寸孪生、生成代码的存储从目标张量 stride 排布。 - 编译器:stax 辅助导入指向缺失模块、整数无穷大进入共享符号注册表、常量比较与
无限边界求值、channels-last 卷积区域以 NameError 失败、捕获常量命名冲突、
容器参数内的张量未获得独立输入、
IntInfinity代码生成支持、库算子的函数体在 记录它的 trace 之外运行、FFT 不再递归 eager 分解、自适应池化的参差行对齐行轴、 余数用本语言运算符打印、对角视图保持其所视的值、未声明填充类型读为未定义、 不透明常量的改写换成其非就地等价形式、高阶算子在 dispatch trace 中整块记录。 - 内核:
index_put契约与 ABI 对齐(见不兼容变更)、负尺寸在稠密张量分配处拒绝、 mapped tensor 在 vmap 下转换、CUDA 复数广播内核越界、平均池化向上取整计数、 strided 池化操作数经 row-major 拷贝读取、copy_源广播与置换目标拷贝、index_fill/index_copy触达每个前导位置、gappedempty_strided按完整跨度分配、logsumexp可在所有后端上归约任意维度集合、点积读 strided 向量、addmv族服务 每种 dtype、pow.Scalar标量自变量、掩码通道保持全一 bit 模式、poisson的 Stirling 项按真值掩码选取、低精度 max / min 传播 NaN、conv / dropout / masked softmax 拼写拿到各自契约、half_to_float 与频率 API 路由修正、 FFT 批处理与执行路径补全、quantile/nanquantile注册、CPU 激活在双精度张量上 走双精度、位运算按两个操作数提升到的类型应答、_trapezoid/_cumulative_trapezoid缺省 x 时求值、ps_roi_pool反向分箱对齐前向、越界 batch 下标的 ROI 池化输出置零、 自适应最大池化 3d 的 channels-last 分派恢复、cummax/cummin索引在相等时取 最新极值、布尔与复数比较与打印覆盖每种 dtype、CPU batch / group norm 保双精度、conv_transpose1dbias 梯度、稀疏半结构化 gather 梯度、stride 的 leading-dim 求和 链深度用进位级联封顶。 - CUDA:跨维 softmax 竞态(见不兼容变更)、bool 二元运算跳过尾部元素并忽略广播、
loss 反向线性读取扩展梯度、nll 反向仅在 mean 时除以 total_weight、
unique多算 输出、一维排序最坏情形、嵌入梯度原子堆积、ldexp的整数对在单精度下求值、 GEMM 计划路径上转置左操作数改拷贝、softmax 入口路径加保护、 归约大张量的输出地址拆分与 packed 极值索引定位、hash_tensor索引在 NaN 与相等 情形下取首次出现、FFT 信号长度、超过 64 个 iterator 维度的额外元数据上传与寻址。 - 构建:Python 绑定头文件名大小写冲突、JIT 运行时头文件未进 wheel、sdist 排除项、 wheel 根目录 include/share 锚定、tprun 弹性启动器 console script 注册。
性能 Performance
测量方法均为静默机器 min-of-window;硬件标注于各条目。
- CPU flash 注意力(64×4×49×32 float32,20 线程):前向 0.24 ms,前向+反向约 1.0 ms(分块 + 线程池路径)。
- 编译耗时:单次编译从 11.3s 降至 1.3s(图构建 O(n²) → O(n))。
- CUDA unique:1M 元素
0.60ms → 0.31ms(按需计算逆映射/段长);配合一维基数 排序,unique 1M 由落后 49.5x 收敛到 2.4x;一维 sort 0.43ms。 - 三角掩码(tril/triu):2048×2048 由 11.1x → 1.21x,100000×8 由 12.8x → 2.3x。
- 外层维归约:100000×256 沿 dim0 标准差 4.1ms → 0.33ms(多 CTA 拆分放开)。
- 直方图:1M 元素 64 桶负载全局原子串行化 12.8x → 共享内存暂存后块尾发布。
- cummax:2048×2048 由 1.10x → 0.97x(比较方向改编译期常量)。
- 嵌入梯度:524288×64 落后 1.27x → 排序分段累加并双路径择优后领先。
- 融合优化器检查:Adam/Adagrad 融合路径的 Python 指纹重建(约 0.4 ms/步)改为 单次原生遍历检查布局。
- fp32 注意力分数驻留片上:32 batch × 32 head × 4096 token 调用省去分数矩阵三次 片外往返(约 20 ms)。
- 非方形因果注意力:此前落后 17x–183x 的调用改走融合路径。
- cuDNN benchmark 模式:同候选列表内择优比启发式首选快约 11%。
- CUDA 归约引擎修订 14(RTX 3090,CUDA event 测 7 组 × 100 次 C++ 归约启动的中位
时间,单位微秒;完整测量记录见
docs/cuda_reduction_optimizations.md): 非对齐输入 float32 64M 元素318.657 → 309.954(-2.7%)、half→float32 64M 元素172.993 → 157.727(-8.8%);连续输出分组 float32 1048576×1693.060 → 85.699(-7.9%)、half→float3266.787 → 48.946(-26.7%)。短归约 half→float32 16384×128 从10.568 → 11.817(+11.8%),非对齐路径的 1M 元素 float32 测试无收益(6.646 → 7.158)——该轮数字只代表这些 dtype 与布局。 - CPU 向量化与 channels-last 铺开:AVX2 / AVX512 双层路径、tier-compiled 的 channels-last 池化 / 上采样 / grid sampling 通道、max-pool 扫描致密化, 以逐算子基准为准(未在本轮重测全表)。
- 排序与扫描:超大排序分块并去掉小排序暂存、大张量平坦中位数走基数排序、
topk多块路径 tile 超 uint32 时回退、cummax网格按设备maxGridSize封顶。
文档 Documentation
- 白皮书:新增 08b 图编译管线(捕获、六 pass、AOT 联合图与最小割、守卫、区域缓存)、 08c stax 降级(12 级降级查找链、循环 IR、物化阈值、三条代码生成路线)、08d 原生图 运行时(ValueNode/OpNode/Graph、最后使用即释放、fused_mul_add、custom_op 桥)。
- 指南:新增 getting-started 系列(tensors / autograd / models / training / datasets / saveload / transforms)与扩展指南(自定义算子、C++ 内核 JIT/AOT、dispatch 协议)。
- 笔记:新增
notes/系列——cuda semantics、ddp internals(双进程 gloo 实测)、 multiprocessing best-practices、amp、序列化、FAQ、modules。 - 参考页:补全
func(whirlwind tour / API / batch_norm / 局限)、distributed(fsdp / tensor parallel / device mesh / pipelining / RPC / elastic / checkpoint)、ao量化、masked、signal、nested、monitor、accelerator、cuda.tunable、omni_attention、varlen等页面。 - 概念页:
size、storage、tensor_view、tensor_attributes、type_info、complex_numbers、deterministic、dlpack、meta等。 - 内核编写指南:JIT 区新增 kernel authoring guide,覆盖内核侧经环境分配器的分配。
- API 参考:多个参考页的 autodoc 修复(docstring 附着时机、重复摘要、限定名引用)。
- README 重写:quick-start 改为围绕现代 decoder-only LLM 组织(中英双份同步), mixer 段落改写,API 路径修正,修正表格与目录数字;README 内嵌反向传播演示视频 (含 B 站入口)。
- 交叉引用清理:全部跨引用指向确实存在的 API;dispatch-key 与 scale 章节按当前 树重建;移除遗留的参照实现指涉与虚构引用;AMP 笔记指向确实存在的设置。
- 测量记录:新增
docs/cuda_reduction_optimizations.md(归约引擎修订 14 的测量 数据、可重复验证命令与构建记录)。 - 测试规模改为实测:README 徽章与 benchmark 段落此前手写"5,680 tests",树一长大就
过期。新增
tools/update_readme_stats.py,解析测试树统计测试函数与测试类并改写tp-stats标记之间的区域,--check在已提交文本与实测不一致时失败;用扫描而非 collection 是因为数字必须是树的属性——collect 会让公布值随编译扩展、GPU 与各模块importorskip的可选包而漂移。 - 子模块 bot 的指引改指真正的开关:两处"如何解封"此前都让读者去配 token,而 bot 撞到
的拒绝(
GitHub Actions is not permitted to create or approve pull requests)点名的是 工作流凭证,且在到达 API 之前就被拒。改为指向仓库设置 "Allow GitHub Actions to create and approve pull requests",并把它与"SUBMODULE_UPDATE_PAT 才能让 bump PR 触发 lint / smoke-build"分成两件事说清(一个是复选框,一个是 secret)。 - 删除无人引用的视频 poster:embed 早已改为从仓库附件 URL 取片,本地 poster 不再有读者。
开发者 Developers
- 子模块:cutlass 与 flash-attention 收为 git 子模块,CUDA lane 拉取。删掉
docs/ebook这个悬空 gitlink——它在索引里指向一个提交,但本仓库历史里从未有任何提交 为它写.gitmodules条目,谁也拉不到、也没人声明;每次检出都要为它付出代价,post-checkout 清理遍历子模块时撞上这个无 url 的路径并以 128 退出,该步把失败吞掉,于是主干每个 job 的 检出都留一条警告注解。 - wheel 内容:Python 桥接头文件、JIT 运行时头文件随 wheel 发布(C++ 扩展可用); include/share 排除项锚定到 wheel 根。
- 发布:CPU 轮子恢复 PyPI trusted publishing;nightly 版本号折叠 build number。
- CI:CodeQL 扫描告警修复、PR 标题/提交规范校验、Roadmap 项目板接线与状态同步、 移除 CircleCI / AppVeyor 集成。
- 构建:
MAX_JOBS伞形并行度、环境变量直通、MKL 前缀探测、生成文件内容不变时 跳过重写。 - 命令行入口:
tprun弹性启动器 console script 注册(tensorplay.distributed.run)。 - 测试基建:新增 52 个测试文件,覆盖二阶梯度(激活 / 卷积归一化 / 形状算子 /
网格采样 / SDPA / 核级 / 特殊函数 / 损失 / 线性代数 / 保留梯度 / 就地拼写)、
现代注意力与生成栈、stax 原生下沉与算子图区域、FFT / STFT、CUDA 归约布局 /
hash_tensor/ softmax 与 cummax 修复 / 批量归一化 / GEMM 转置操作数 / 卷积精确 float32 / 大张量归约。 - Windows CUDA 轮子:
_C因图像 IO 绑定链接 nvjpeg,而 Windows CUDA 的 DLL 清单从未 列入nvjpeg64_*.dll——构建照链接导入库通过,每次导入轮子却报 "DLL load failed while importing _C"(Python 只搜tensorplay/lib与 System32)。改为随其他 toolkit DLL 一起 打包。新增tools/check_wheel_dlls.py:读 Windows 轮子里每个二进制的加载期导入表, 点名轮子既未携带、System32 也找不到的 DLL(在出问题的 cu130 轮子上正好报出nvjpeg64_13.dll),构建 lane 直接跑它,缺 DLL 在构建 job 里以名字失败而不是在测试 job 里变成 import 错误。CrashHandler 的 Windows 分支补齐无条件栈捕获 (capture_stacktrace_unconditional只在 POSIX 与 fallback 分支有定义,_C.pyd以 LNK2001 链接失败)、getpid(改用process.h)、strsignal(换成返回常量信号名的 可移植帮助函数)。 - Linux nvjpeg 预载:
_C链接 libnvjpeg,而 Linux 预载只对缺失的 cudart / cublas / curand / nvrtc 反应。从 NVIDIA wheel 而非LD_LIBRARY_PATH上的 toolkit 拿 CUDA 库的安装会在 libnvjpeg 上失败且不重试;现在按 nvrtc 一样预载它(可选, 缺失时仍抛出原始错误)。令牌扫描改调builtins.any——包随后把any()重绑到张量归约, 在导入之后被调用时预载会坏掉。 - arm64 轮子:修复 trunk CI 的 linux-arm64 与 macos-arm64 构建——
FindMKLDNN.cmake把MKL_INCLUDE_DIR的 NOTFOUND 哨兵转发进dnnl_cpu的 include 目录(在走 OpenBLAS 的平台上直接中止 CMake generate),其无条件USE_MKL定义本会让 oneDNN 去拉不存在的 MKL 头;ReduceMomentsKernels.cpp无条件包含<immintrin.h>;ReductionKernelsImpl.cpp在无守卫的 dispatch 处调用 x86-only 的 leading-rows 帮助函数。 - ISA 分层链接顺序:AVX2/AVX512(VSX、ZVECTOR、SVE、RVV)的分发内核副本此前是对象库,
CMake 会把它们排在 p10 自身对象之前(链接行与静态归档都是)。这类副本发射到线外的每个
inline 函数或模板实例化都是弱定义、用该 tier 的 ISA 标志编译,而链接器保留它先遇到的那
一个——在没有 AVX512 的 CPU 上,
dot()因此死在归约内核 AVX512 tier 里带出的std::vector<float>fill 构造的 AVX512 副本上(SIGILL)。改为把 ISA 副本作为 p10 自身 的源(单行包装器)编译,并追加在所有其他 p10 源之后(含 tpx),从要求最低的 tier 起排; DEFAULT 副本是基线代码,保留其对象库。 - 轮子泄漏检查:新增
tools/check_isa_tier_leaks.py,列出一个库里执行 AVX / AVX512 指令却不属于CPU_CAPABILITYtier 命名空间、也不在名字里标出 ISA 的弱函数——这种函数 是赢下链接、绕过分发器直接跑的 tier 副本。在b460137e构建的轮子主干上它正好报出 拖垮 nn-utils lane(SIGILL)的那个构造;现在每个 x86-64 Linux 轮子的libp10都跑它。 - CI 稳定性:托管 Linux CUDA 构建加 24 GB swap(16 GB 内存 / 3 GB swap 的 runner 上,
仅
ReductionKernels.cu的 ptxas 就越过 14 GB,第二个编译 job 旁边把 cu130 runner 饿到 被 GitHub 丢弃);sccache导出SCCACHE_BASEDIRS(此前报 "Base directories (none)", cache key 带着绝对检出路径、runner 临时目录与解释器前缀,跨 runner 与 Python 版本必然 失配);空 tpx 探针不再生成= {}这种 MSVC 以 C2466 拒绝的零长度数组;readme-stats 在主干拒绝工作流 token 时把过期数字报成警告并给出落地命令,不再因仓库设置而红(其他 push 错误仍然失败);子模块 bot 在凭证被禁时把该事实与逐子模块故障区分开——第一次拒绝 置标志,其余过期项照报但不再推分支,run 以绿色收尾并给出补救说明,漂移检测的语义不变。
安装 Installation
# CPU 轮子(PyPI)
pip install tensorplay==1.1.0
# CUDA 13.0 轮子
pip install tensorplay==1.1.0 \
--index-url https://download.tensorplay.cn/whl/cu130/ \
--extra-index-url https://pypi.org/simple
1.2.0.dev20261027
pre-releaseRolling preview builds of TensorPlay, published automatically from main.
Expect the newest capabilities — and the occasional rough edge.
Install
Wheels cover Python 3.10–3.13. --pre is required to select dev builds.
CPU (Linux x86_64 / aarch64, macOS arm64, Windows x86_64)
pip install --pre tensorplay \
--index-url https://download.tensorplay.cn/whl/nightly/cpu/ \
--extra-index-url https://pypi.org/simple
CUDA 13.0 (Linux x86_64, Windows x86_64)
pip install --pre tensorplay \
--index-url https://download.tensorplay.cn/whl/nightly/cu130/ \
--extra-index-url https://pypi.org/simple
Pin a specific build
pip install --pre tensorplay==1.0.0.dev20260907+cu130 \
--index-url https://download.tensorplay.cn/whl/nightly/cu130/ \
--extra-index-url https://pypi.org/simple
About this channel
- Built and smoke-tested on every code-bearing
mainpush; a scheduled run at 02:30 UTC fills any gaps. - Only the newest build of each variant is kept here — older dev wheels are removed as new ones land.
- The same wheels are served from the plain indexes at
https://download.tensorplay.cn/whl/nightly/<variant>/. - For production work prefer the stable releases or PyPI.
- Hit a problem? Open a bug report and include the version printed by
python -c "import tensorplay; print(tensorplay.__version__)".
TensorPlay 1.0.0 Release Notes
- 亮点 Highlights
- 不兼容变更 Backwards Incompatible Changes
- 弃用 Deprecations
- 新特性 New Features
- 改进 Improvements
- 问题修复 Bug fixes
- 性能 Performance
- 文档 Documentation
- 开发者 Developers
- 安装 Installation
亮点 Highlights
- autograd.Function 完整实现:ctx-first backward 契约、grad_fn/ctx 双侧 hooks 与引擎协同逐项核验,并新增原生 forward-mode AD。
tensorplay.library自定义算子体系:覆盖 custom_op/triton_op/ register_fake/register_vmap/opcheck 等全量公开面,调度层净开销 2.3µs; tile-lang 一等支持、tvm-ffi 零拷贝直通。tensorplay.compile(Stax 编译器)收口:CUDA 减归约/逐元素链完成 11 组负载验证(RTX 4090 D)。- CUDA graphs 完整实现:条件节点(if/while 图)、多设备并发捕获、
make_graphed_callables、共享内存池,外加单调用回放快路径。 - RNG 逐位稳定:CPU 原生随机引擎 + CUDA 设备端 Philox4_32_10, 同种子逐值一致,state 字节布局固定。
- 构建/CI/发布全链路收口:scikit-build-core PEP 517、统一版本链路 与 CI 形态、CPU 走 PyPI、CUDA 走自建 PEP 503 索引,并新增 nightly 预览通道。
不兼容变更 Backwards Incompatible Changes
CUDA graphs
旧 _C.cuda_graph_* 自由函数绑定全部移除
原生层重写为 tensorplay::cuda::graph::CUDAGraph 类,
Python 层改用 tensorplay/cuda/graphs.py 的原生类封装:capture_end() 即完成
instantiate,replay() 直接持 exec 指针发射,支持 capture_error_mode、
graph_pool_handle() 共享池与 debug_dump()。依赖旧整型句柄接口的代码需迁移。
图 API
Graph.output() 改为单例替换语义
先擦旧 output 再建新。此前 _interpret 取首个 output 而 recompile 取末个,
双 output 时可能返回不同值;现在全链路一致。
tensorplay/fx.py 移除
公共图 API 迁移为 tensorplay/graph.py 门面(实现仍在 tensorplay/_stax/graph.py)。
构建
自建构建脚本删除,统一 PEP 517
rebuild.py(Windows 本地构建编排)与 release.py(cibuildwheel 发布编排)移除。
构建一律走 pip install . / python -m build --wheel(scikit-build-core),
python setup.py install 不再支持。USE_CUDA=OFF pip install .、
CMAKE_CUDA_ARCHITECTURES="70;75;86" pip install . 等环境变量直通生效。
版本号改为构建期生成
tensorplay/version.py 不再硬编码,由 tools/generate_tensorplay_version.py 在构建时
从 TENSORPLAY_BUILD_VERSION/BUILD_NUMBER → PKG-INFO → version.txt + git SHA
三级生成。main 分支版本形如 1.0.0a0+git<sha7>。
弃用 Deprecations
无。
新特性 New Features
Autograd
autograd.Function完整实现:ctx-first backward、grad_fn/ctx 双侧 hooks、to_save/metadata/next_functions/mark_dirty/once_differentiable/InplaceFunction/NestedIOFunction、梯度数校验含"多余 None 截断"规则。- 新增
custom_function_apply单入口融合:节点创建/unpack_input/forward/ setup_context/wrap_outputs 一次 C 穿越完成。 - 引擎 InputMetadata 原生化:Edge 记录 shape/dtype/device,Engine 在 C++ 层 零填充缺失梯度槽。
saved_tensors_hooks(pack 在 save 时、unpack 随 ctx 快照)。- forward-mode AD 原生核心:
tensorplay/autograd/forward_ad.py(level 栈/DualTensor/种子算子 JVP)。
自定义算子
tensorplay/library.py:custom_op/triton_op装饰器、register_kernel/register_fake/register_autograd/register_vmap/register_autocast、Library(DEF/IMPL/FRAGMENT)、infer_schema、opcheck四项检测、set_kernel_enabled/get_kernel、顶层define/impl/impl_abstract。- tile-lang 一等支持:
tile_lang_op+wrap_tilelang(同契约超集扩展)。 - tvm-ffi 直通:TP Tensor 实现 DLPack 协议,
tvm_ffiJIT/AOT 两路tvm::ffi::TensorView参数零拷贝直通。 - Stax 原生图新增
custom_op节点:编译图中原生执行(非解释器), autograd 可穿透编译图求梯度。 - Composite 分发键:形状/视图组合算子批
在 CUDA 上不再
Kernel not found。
编译器(tensorplay.compile / Stax)
- 控制流静态特化:
if x.shape[0] > 2、range(x.ndim)、len(x)切片等元数据 驱动的 Python 控制流可 fullgraph 编译(数据依赖分支仍拒绝)。 - 守卫版符号形状:被读形状的参数按尺寸精确重特化,未触碰参数保持 dynamic 通配。
- AOT 双 partitioner:
partition_default(标签切分)与partition_min_cut(自研 Edmonds-Karp 最大流/最小割,支持 memory_budget)。 - Pass 体系:
PassManager至不动点执行, 内置 DeadCodeElimination/ConstFold/ShapeProp。 tensorplay.graph门面:语义化节点命名、erase_node/replace_all_uses_with、is_leaf_module钩子、concrete_args、create_feature_extractor、to_dot()/draw()可视化。- Triton 代码生成:四类 launcher(single/dims/split/pw)+ 静态 fast-launch 路径。
CUDA
- CUDA graphs 条件节点(if/while 图,CUDA ≥ 12.4)、多设备并发捕获、
make_graphed_callables/make_graphed_autograd_function。 - 共享内存池:
graph_pool_handle(),多 graph 共享私有池,引用计数释放, 池内相邻空闲块合并防碎片。 stage_and_launch回放快路径:全部输入单次 Python→C++ 调用内暂存并发射。- 原生
memory_stats(device):segments/free_blocks/graph_pools/capturing 等, 接入tp.cuda.memory_stats*。 - OOM 防御梯队:分配失败先排空跨流 pending 块,再 flush 缓存重试, 报文附 reserved/allocated/free 统计。
内核
- LayerNorm CUDA 前向/反向自定义 kernel(不再依赖 cuDNN):Welford warp 归约、 4 宽向量化、半精度 fp32 累加、确定性反向归约。
- argmax packed warp-shuffle:(value, index) 打包 u64 单次整数 max, NaN/±0/首现 tie 位等价。
- Loss 半精度扩展(nll_loss/mse_loss_backward → Float64/Half/BFloat16)、 二元算子与半精度一元浮点 4 宽向量化快路径。
RNG
- CPU:原生随机引擎 + 分布变换公式(Box-Muller 缓存、Fisher-Yates randperm、poisson 等),normal_fill 走 AVX2 向量化路径,大张量 randn 逐位一致。
- CUDA:curand 设备端 Philox4_32_10 + host (seed, offset) 预留。
- Python:
get_rng_state/set_rng_state/fork_rng、Generator.get_state/set_state(5056 字节 POD 布局)。
改进 Improvements
- CUDAStream 热路径:设备数进程级缓存、每线程当前流改扁平数组(消除哈希查找)。
- 当前设备线程级缓存:
currentDevice()/getCurrentCUDAStream(-1)降为一次 TLS 读。 - 自定义算子热路径:捕获守卫无 trace 时跳过 proxy 扫描、内核选择无锁化。
- 分布算子 dtype 补齐:Half/BFloat16、random_/geometric_ 全整型谱 + Bool、randperm Int32。
- Tracer:
call_module节点 + 限定模块路径消歧。
问题修复 Bug fixes
- DLPack:CPU/-1 设备归一为 device_id=0(此前 -1 使 tvm_ffi 走 workspace 复制路径段错误)。
- CUDA scatter 分发:
if (Add)改if constexpr,assign 模式全 dtype 实例化。 - PyNode backward 输入槽语义修正为前向输出数。
- eager 自定义算子调用丢弃 kwargs 的 bug。
is_autocast_key恒假 /is_autograd_key上界吞新键两个区间谓词。- persistent split 代码生成修复(误发 classic tail、死 preamble、stride 整除越界等)。
性能 Performance
测量方法均为静默机器 min-of-window;硬件标注于各条目。
- 自定义 Function 端到端 fwd+bwd 40-45µs(backward 腿 13.5µs)。
- 自定义算子调度层净开销 2.3µs,端到端调用 ~2.9× 于 Python 装饰器路径。
tensorplay.compile(RTX 4090 D,iters=200):11 组负载基准, dims/epilogue 链 2.25-2.51x, pw 链 0.98x(内核已验证,残差为 Python 启动路径)。- kernel_launch CPU 开销 45 → 9.7 µs/call;compiled 全链 51.3 → 20.2 µs/call。
- sum full 16M:62.5 → 41.0 µs。
文档 Documentation
docs/source/library.md增补 tvm-ffi 用法章节。CONTRIBUTING.md构建说明更新为 PEP 517 命令。- README(中英)新增 CUDA 索引与 nightly 通道安装说明。
开发者 Developers
- 构建统一 scikit-build-core(
MAX_JOBS伞形旋钮、editable redirect 模式、 sdist 裁剪),cmake/EnvVarForwarding.cmake实现环境变量直通。 - 版本链路 provider 化(PEP 440 校验 + sdist 一致性断言)。
- CI 重排:可复用
_binary-*工作流 +pull/trunk/publish/lint编排,BUILD_ENVIRONMENT/SHA1环境变量与显式 permissions。 - 发布:CPU 轮子 PyPI trusted publishing;CUDA 轮子 GitHub Release +
pypi-pages分支静态 PEP 503 索引(Cloudflare Pages);新增手动 nightly 预览通道(X.Y.0.dev<日期>+cuXXX版本规则)。 - Release notes 机制:PR 按路径自动打
release notes: *标签 (.github/labeler.yml),发布文案存于docs/release-notes/。
安装 Installation
# CPU 轮子(PyPI)
pip install tensorplay==1.0.0
# CUDA 13.0 轮子
pip install tensorplay==1.0.0 \
--index-url https://download.tensorplay.cn/whl/cu130/ \
--extra-index-url https://pypi.org/simple

