CUDA 与昇腾数据流对照：存储层级、分块与异步流水
问题链与阅读目标
来源：2026-09-21—23的AI Core/CUDA Core、UB/L1/L0、TBuf/TQue、SIMD/SIMT、不规则访存和Tensor Core追问。本篇是整理时补充的教学解释，不是已在你机器上验证的性能结果。
核心问题是：数据放在哪里，由谁搬运，谁消费，何时允许复用？把这四项画清，平台名词就有共同结构。
对应关系不能当成硬件等号
| 层次 | CUDA侧 | 昇腾侧 | 类比边界 |
|---|---|---|---|
| 计算簇 | SM | AI Core及具体AIC/AIV组织 | 层级接近，内部组成、数量口径不同 |
| 矩阵计算 | Tensor Core | Cube | 数据类型、操作数布局、指令不同 |
| 设备主存 | Global memory；物理介质依设备可为GDDR/HBM | GM；物理存储依设备规格 | 地址空间与物理存储不是同一概念 |
| 软件管理片上空间 | Shared memory | UB及矩阵侧缓冲 | UB不等于shared；访问路径和消费者不同 |
| 矩阵操作数/累加 | 寄存器fragment或代际特定空间 | L0A/L0B/L0C等 | 不能把所有L0统一叫CUDA寄存器 |
CUDA不只有shared memory：还有寄存器、L1/L2、local地址空间等。local memory名称也不保证物理片上。昇腾暴露更多专用缓冲和搬运路径，是执行单元与编程抽象的结果，不能直接推为先进或落后。
此前“910C确定只有32个AI Core”的回答缺乏统一可靠规格，不采纳。也不能把一个AI Core与一个CUDA Core直接比数目。
为什么CUDA的一个if不能消除不规则访存成本
Gather教学例：y[i]=x[index[i]]。CUDA线程可自然写条件和索引，但仍需处理越界、非合并访问、分支分歧及负载不均；语法短不等于硬件代价小。昇腾tile/vector路径往往更显式地处理索引、掩码和搬运。
SIMD/SIMT需要按具体芯片和编程模式判断，不能把950相关文档套到910A/B/C。910A到B某条搬运路径的变化，不足以判定整体架构退步；迁移成本应分别看语义、API、二进制、性能四层。
从GEMM理解分块
C[M,N]=A[M,K]×B[K,N]。每个计算块负责输出的BM×BN部分，沿K以BK累加。片上保留A/B tile以提高复用；块越大不总越快，因为片上容量、寄存器、尾块和并行数有约束。
最简教学账本：A/B缓冲约为(BM×BK+BK×BN)×输入字节；双缓冲约乘2；再加累加器、转换和临时空间。实际不同空间分别核算，不把UB和L0容量直接相加当作统一池。
TBuf、TQue和双缓冲
TBuf表达临时局部缓冲；TQue表达阶段间张量的生产、排队、消费和释放。TQue不是跨卡网络队列，也不是Python线程队列。TPipe和相关API协助管理资源和同步。
双缓冲只是有两套存储。正确流水还需：先预取首块；发起下一块搬运；计算当前块；确认搬运完成再消费；确认旧数据已不用再覆盖；最后排空。普通compute(); load();交换索引的代码并不能证明异步。
教学流水上界：串行约n×(搬运+计算)，理想重叠约启动成本+(n−1)×max(搬运,计算)+排空成本。实际要减去争用、屏障及流水气泡；本式不是设备实测。
Tensor Core调用层级
框架或cuBLAS/CUTLASS隐藏较多细节；WMMA暴露warp协同矩阵操作；PTX更靠近虚拟指令层。以下仅演示调用关系，省略分配、对齐、加载和边界，不是可直接运行的完整kernel：
// 同一warp协同执行；a、b、acc是匹配类型和shape的fragment
wmma::fill_fragment(acc, 0.0f);
wmma::load_matrix_sync(a, A, lda);
wmma::load_matrix_sync(b, B, ldb);
wmma::mma_sync(acc, a, b, acc);
wmma::store_matrix_sync(C, acc, ldc, wmma::mem_row_major);
WMMA的逻辑tile不必一对一对应一条同形状mma PTX。旧回答将16×16×16例子直接映射为固定PTX字符串，不能当作已核实汇编；应编译到目标架构后检查生成代码。
最小实验与来源
用同一GEMM或向量操作比较简单版本、显式tiling、双缓冲；固定输入dtype/shape，先验数值，再看搬运等待和关键路径。性能结论需记录设备、编译器、版本、预热和计时范围。
核查2026-09-23：NVIDIA PTX ISA、Ascend搬运路径API、Ascend InitBuffer。各来源仅支持对应版本，未核实你机器每项特性。
2026-10-07 深化：CUDA执行层级与寄存器的边界
来源：近期warp、lane、SM、寄存器、PTX及CC追问摘录；以下是核查后的教学解释，不是实机规格报告。
grid包含blocks，block包含threads；block可协作共享数据和同步，普通blocks不能依赖隐含执行先后。一个warp有32个lane；三维block的线性thread编号为x+blockDim.x×(y+blockDim.y×z)，lane编号是其对32取模。threadIdx.x%32只有一维或布局满足相应条件时才可直接代用。
驻留warp、可发射warp和真正同时进行的指令工作不同。“SM能驻留64个warp，所以每周期执行64个warp”不成立；也不能把某代的4个分区、CUDA/Tensor Core数量套到全部SM。Ada官方上限是每SM驻留48个warp、64K个32bit寄存器；这是Ada架构限制，不是每个kernel的实际占用。Ada Tuning Guide
寄存器保存线程所需的操作数/状态，由编译器分配；L1缓存来自内存访问的数据。二者都可片上，但管理方式、可寻址语义不同，因此寄存器不是“GPU的CPU L1”。动态索引数组、过多局部数据和spill可能进入local地址空间，后者可位于设备内存，不能由local这个名字推断极快。SIMT与存储说明
Occupancy是驻留活跃warp相对支持上限的比例，受线程、寄存器、shared和block上限制约；提高它可能改善延迟隐藏，也可能因压缩寄存器导致spill。指标不是最终优化目标。
CC是设备能力版本，sm_89是相应机器代码目标；GPU型号还决定实际SM数量。cluster是特定架构功能，不是所有CUDA卡都有；Ada 8.9不能因CUDA版本新便取得Hopper的cluster功能。后续架构是否支持应查官方能力表，而非“版本数字更大就全部拥有”。Compute Capabilities
PTX是虚拟ISA，cubin包含目标设备代码，SASS用于查看具体机器指令；不要把PTX长度当运行时指令数。global、__device__表达调用/执行位置，设备端动态启动又有单独规则；原提问里的__manager__未取得完整代码，不擅自改写为__managed__后宣称已解释。
Triton也不是永远慢于CUDA：性能取决于算法、布局、编译器、目标设备和基线。较新的布局/warp specialization能力会随版本变化，早期限制不能写成永久事实。
索引、尾块、tile成本和Roofline的完整推导见GEMM、Triton 与 TileLang：从分块代码到 Roofline 性能账本。复习实验：固定相同数学操作，改变block/tile并读取寄存器、spill和时间，说明为什么高occupancy与低耗时可以不一致。