GPU 算子与稀疏计算

本课计算同 shape 张量的逐元素表达式：

每个输出元素满足

元素之间互不依赖，因此天然适合数据并行。

视频在 00:07 给出这个目标，在 00:14 用“大而慢/小而快”描述存储层次。

这个抽象可能对应：

不同硬件不必拥有完全相同的层级和显式搬运指令。

视频在 00:34 先说明多个计算单元分担任务。

若有 P 个 worker，可把 N 个元素分成互不重叠的索引集合。

理想均匀负载下，每个 worker 处理约

个元素。

这是并行任务划分。

算子融合回答的是另一个问题：

两种优化可以同时存在，但概念不能混为一谈。

未融合实现分两步：

视频从 01:01 开始描述两个函数，在 01:08 进入乘法 kernel。

若两个 kernel 之间存在全局可见边界，第一步通常需要让 T 可被第二步读取。

理想化地按元素统计：

总计：

融合 kernel 对每个元素直接执行：

t 只需存在于寄存器或编译器管理的局部临时值中。

视频从 01:31 转入融合路径，在 01:44 说明一次完成乘加。

理想化流量是：

即

相对未融合路径，少了 T 的一次完整写入和一次完整读取。

设每个元素占 s bytes。

不考虑 cache line、write allocate 与对齐，未融合流量约为

融合流量约为

节省

bytes，理想化总流量下降

这是流量比例，不等于运行时间必然下降三分之一。

每个元素只有一次乘法和一次加法。

按常见 FLOP 口径，这是 2 FLOPs。

FP32 时，融合路径理想化需要 16 bytes：

算术强度为

未融合路径为

两者算术量几乎相同，fusion 主要提升有效算术强度并减少数据搬运。

除了 temp 流量，融合通常还可减少：

小张量上，launch 开销可能比 HBM 流量更显著。

大张量上，带宽与中间量写回通常更重要。

视频在 01:47 把收益概括为“搬运减少”。

板书把每次算子边界画成外部内存往返，是为了突出最坏的全局中间量成本。

真实执行可能不同：

因此应通过 profiler 检查真实 global memory traffic，而不是把板书箭头直接当成 HBM 事务数。

算子融合是计算图或 kernel 边界优化。

FMA 是单条硬件指令执行

融合 kernel 可能被编译成 FMA，也可能仍使用分离的 multiply 与 add 指令。

反过来，一个 FMA 指令也不自动说明多个高层算子已消除全局 temp。

FMA 只有一次最终舍入，可能与先乘后加的两次舍入产生微小数值差异。

严格数值模式可能限制 contraction。

fusion 会扩大单个 kernel 的工作范围，可能带来：

此外，存在副作用、精确异常语义或必须同步的算子不能任意融合。

板书默认 A、B、C、D 具有相同 shape，并按相同连续索引访问。

若存在 broadcasting，例如

融合仍可行，但 B 的索引、cache 复用和向量化方式会变化。

非连续 stride、转置 view 与动态 shape 也会改变性能模型。

若 T 还有另一个消费者，是否仍能完全删除 T？

不一定。

可选择复制计算、部分融合或保留 T；最优方案取决于 T 的计算成本、消费者数量和内存压力。

本例的数学乘法和加法数量没有减少。

主要减少的是中间量流量与调度边界。

这是跨硬件的教学抽象，实际层级可能完全不同。

是否 contraction 由 ISA、编译器和数值模式决定。

过度融合可能增加寄存器压力、spill、分支和代码尺寸。

运行时间还受 launch、带宽利用率、cache、occupancy 与其他瓶颈影响。

本课使用最简单的逐元素加法：

默认 A、B、C：

视频在 00:04 进入两个张量相加，在 00:07 先讲标量版。

本地字幕在 68 秒后出现重复尾句，课程内容不采用该无效重复。

标量教学伪指令可写为：

``text load_scalar A[i] -> r0 load_scalar B[i] -> r1 add_scalar r0, r1 -> r2 store_scalar r2 -> C[i] i += 1 ``

视频在 00:10 描述循环，在 00:11 进入 A 的标量 load。

FP32 占 4 bytes，因此地址每轮前进

处理 N 个元素，主循环需要约 N 次标量 add 指令。

单条标量 add 只产生一个元素结果。

但现代 CPU 还可能具备：

所以“标量版”只是向量化比较的主循环模型，不表示整个处理器严格串行。

SIMD 是 Single Instruction, Multiple Data。

一条向量算术指令对多个 lane 执行同一运算。

若向量寄存器宽 W bits，元素宽 E bits，理论 lane 数为

视频在 00:29 采用 512-bit 示例。

对 FP32：

这是 AVX-512 风格教学例子。

并非所有 CPU 支持 512-bit 向量：

512-bit、FP32 的概念伪指令：

``text load_vector A[i:i+16] -> v0 load_vector B[i:i+16] -> v1 add_vector v0, v1 -> v2 store_vector v2 -> C[i:i+16] i += 16 ``

视频从 00:35 说明一次读取 16 个数，在 00:38 给出 $A_i$ 到 $A_{i+15}$。

地址每次前进

视频在 00:53 说明 64-byte 步进。

忽略尾部，标量主循环约执行：

向量宽 L 时约执行：

主循环动态指令数近似下降 L 倍。

但 load/store 的总字节数没有下降：

SIMD 改变每条指令覆盖的数据宽度，不会凭空减少必须读取的 A、B 与必须写出的 C。

视频在 00:58 解释“标量和向量指令周期数相近，向量一次处理更多数据”。

这是有用的直觉，但不是跨 CPU 的固定事实。

实际速度受以下上限约束：

对大张量加法，算术强度很低，常见瓶颈是内存带宽。

即使向量 ALU 还能更快，内存也可能喂不满。

视频板书强调“数据连续、周期数相近”。

连续访问有利于：

若张量是转置 view 或大 stride：

向量化可能需要 gather，吞吐通常低于连续 load。

现代 ISA 常支持未对齐向量 load。

“未对齐”不等于一定故障。

但跨 cache line 或页边界时，可能产生额外事务。

编译器可能：

最优策略依赖 ISA 与微架构。

若

主循环处理 q 个完整向量。

剩余 r 个元素可用：

不能直接越界读取或写入。

编译器通常要确认：

语言中的 restrict、alignment hint 或高层 tensor 编译器的 shape/stride 信息可能帮助证明。

错误 hint 会导致未定义行为或错误结果。

CPU SIMD 通常是一条向量指令显式操作一个向量寄存器。

GPU 常以 SIMT 暴露线程模型：多个线程按 warp/wavefront 成组发射。

两者都利用数据并行，但：

不能把 512-bit CPU 向量寄存器直接等同于 GPU warp。

内存带宽、load/store 吞吐、微操作拆分与频率会压低收益。

向量宽度由 ISA 与硬件决定，代码还要处理 dispatch 或 fallback。

许多 ISA 支持未对齐访问，只是成本可能变化。

stride、gather/scatter 与 cache line 利用会显著影响吞吐。

两者都利用数据并行，但执行模型、索引与内存规则不同。

设大张量 A、B 按 tile 分块，逐块计算

每个 tile 经历三个概念阶段：

视频在 00:08 用“大而慢/小而快”建立教学模型，在 00:27 进入单缓冲。

具体硬件可能使用：

板书名称不能直接替代目标平台文档。

设只有一组输入缓冲 $A_0,B_0$。

当前 tile 的输入必须先准备好，计算才能开始。

计算没有读完缓冲前，下一 tile 不能覆盖它。

视频在 00:42 开始搬入，在 00:49 进入计算与写回。

若平台没有其他重叠机制，一个 tile 的时序近似：

视频在 00:53 总结传入、计算、传出三步，并在 00:59 说明它们有依赖。

设每 tile：

K 个 tile 完全串行时：

这个式子忽略 launch、barrier、尾 tile 与 cache 效果，只用于建立上限直觉。

双缓冲为输入准备两套槽位：

计算消费第 0 套时，copy 可以把下一 tile 写入第 1 套。

下一轮交换角色。

视频在 01:08 引入第二个缓冲，在 01:12 明确 A1、A2 两组。

这种角色切换常称 ping-pong：

在计算 tile 0 之前，必须先把 tile 0 放入某个缓冲。

所以流水线有 fill 阶段。

最后一个 tile 计算后，还要完成结果提交，形成 drain 阶段。

双缓冲改善的是中间稳态，不会消除首尾开销。

当 compute 使用 buffer 0 时：

视频从 01:33 进入异步计算与传入，在 01:40 说明下一 tile 写入另一缓冲。

理想情况下，三个资源互相独立，稳态每 tile 周期为

总时间近似

当 K 很大，平均时间趋近最慢阶段，而不是三阶段之和。

视频在 02:00 概括传入、计算、传出同时进行。

这需要额外前提：

若输入与输出共享一个 copy engine，稳态更可能受

约束，而不是三个时间单独取 max。

若它们争用同一外部内存带宽，重叠也可能只改变排程，不增加总可用带宽。

每套缓冲至少有三种状态：

典型顺序：

``text async_copy(next, write_buf) compute(current, read_buf) wait_copy(write_buf) swap(read_buf, write_buf) ``

真实 API 还可能需要：

缺少同步会造成读取未完成数据或覆盖仍在使用的 tile。

K 个 tile 仍要：

双缓冲的主要作用是隐藏 latency，让 copy 与 compute 时间重叠。

它不是压缩，也不是算子融合。

若 kernel 已被内存带宽完全饱和，增加重叠未必提高吞吐。

若单个输入 tile 总大小为 B bytes，单缓冲需约 B，双缓冲需约 2B。

这可能导致：

更高 occupancy 与更深 pipeline 之间存在权衡。

双缓冲不是“无成本加速”。

tile 太小：

tile 太大：

通常要针对 shape、数据类型与硬件 autotune。

在 cache-coherent CPU 上，程序可能不显式把数据搬到 SRAM。

software prefetch、硬件预取和 cache replacement 已在做部分“提前准备”。

显式双缓冲更常见于：

概念相同，实现原语不同。

若 $T_c=1\,\mu s$、$T_{in}=8\,\mu s$，增加更多计算流水级能否突破输入带宽上限？

不能。

稳态仍由输入路径主导，除非减少 bytes 或提高传输带宽。

还需要异步 copy 能力、独立资源与正确依赖。

copy engine 数量、共享带宽和同步可能限制重叠。

它主要隐藏 latency，不减少输入输出的必要 bytes。

更深缓冲消耗片上存储，可能降低 occupancy 并增加同步复杂度。

这是板书抽象；CPU cache 与不同加速器有各自机制。

视频把“两数之和”类比为并行编程中的“两张量之和”。

这个类比只用于引出问题。

LeetCode Two Sum 是在数组中寻找满足目标和的两个索引，通常涉及哈希或搜索。

本课处理的是逐元素算子：

二者算法目标完全不同。

视频在 00:00 提出问题，在 00:08 做上述类比。

最简单模型假设：

把任意维 shape 扁平化为 N 个元素：

标量 baseline：

``text for i in [0, N): C[i] = A[i] + B[i] ``

时间复杂度：

后续优化不会把必须生成 N 个输出的复杂度变成 $o(N)$，主要改善常数与硬件利用率。

视频在 00:20 指出现实系统有多个计算核心，在 00:25 开始任务划分。

对 P 个 worker，可做静态区间划分：

每个 worker 只写自己的 C 区间，因此没有输出数据竞争。

理想计算工作量约为

但整体仍需传输全部 A、B、C 数据。

每个 FP32 输出元素至少需要：

理想化总流量：

每元素只有一次加法，即 1 FLOP。

算术强度：

所以增加 worker 后：

视频从 00:34 进入存储设备快慢差异。

对于没有数据复用的 streaming add，关键通常是：

板书画“搬到小而快存储”是通用抽象。

但简单加法没有 tile 内复用，显式先复制到 scratchpad 再计算未必有利。

若 staging 本身增加一次额外 copy，可能比直接 streaming 更慢。

视频在 00:48 转入向量指令。

向量宽 L 时，一条 vector add 产生 L 个元素结果。

主循环迭代数从 N 降至约

SIMD 可减少动态指令数、分支与地址更新开销。

但必须处理：

视频在 01:02 也提醒向量指令要考虑连续存储。

视频在 01:04 进入异步，在 01:06 给出 DMA 示例。

双缓冲可以：

但逐元素加法计算极少。

如果 copy 时间远大于 add 时间，短计算阶段无法完全隐藏传输。

异步重叠还需要：

视频在 01:12 提到循环展开。

展开因子 U 可让一次 loop body 处理 U 个标量或 U 个向量：

``text for i += U * L: vector_add block 0 vector_add block 1 ... ``

潜在收益：

潜在代价：

现代编译器常会自动展开，手工展开不一定更好。

一个高性能实现可能同时使用：

但瓶颈会迁移。

例如 SIMD 提高算术吞吐后，kernel 更快撞到 DRAM/HBM 带宽上限。

因此不能把“多核倍数 × SIMD 宽度 × 双缓冲倍数”直接相乘。

若

则 B 沿 M 广播。

逻辑元素仍是 MN，但 B 可能被 cache 多次复用。

若 stride=0 表示广播，向量代码可以 broadcast load，而不是读取 MN 个不同 B 元素。

非连续 view、类型转换与 mixed precision 也会改变事务数和 vector width。

一个稳健流程是：

前者生成同 shape 输出，后者寻找索引对。

共享内存带宽会饱和，调度与 NUMA 也有成本。

具体由硬件和编程模型决定；CPU cache 常由硬件管理。

它们可能争用同一瓶颈，必须看新的 roofline 限制。

过度展开会增加寄存器与代码尺寸压力。

上一课从逐元素加法出发，讨论了多 worker、SIMD、异步搬运和循环展开。

本课继续追问更具体的问题：

视频在 00:04 对比两种任务划分：

这里的 worker 可用于讲解 CPU 线程、GPU 线程或抽象计算单元。

但本课公式采用 CUDA 风格的一维 GPU grid。

不能把抽象 worker 与真实硬件核心一一对应。

考虑长度为 N 的一维连续张量：

多维连续张量也可先展平：

若张量非连续、带 broadcasting 或每个元素对应多个坐标，就还需要 stride-aware 的地址计算。

本课先只讨论逻辑索引如何分配，不把 layout 问题混进来。

假设有 G 个 worker，编号为 $g\in[0,G)$。

一种静态范围划分是：

worker g 依次处理：

视频在 00:11 展示这种连续切分，在 00:17 以四个计算单元为例。

范围划分的优点是单个 worker 的局部连续性直观。

若所有元素计算成本相同，每个区间长度也近似相等，它通常已经有不错的静态均衡。

若每个元素的计算成本为 $c_i$，worker g 的工作量是：

整个并行阶段的时间近似由最慢 worker 决定：

视频从 00:49 构造“前半段难、后半段容易”的例子。

如果高代价元素集中在同一连续区域，某些 worker 会很慢，另一些会提前完成。

视频在 01:03 用“同步等待”解释尾部效应。

这里的等待是阶段结束、kernel 结束或后续同步边界上的抽象。

普通 CUDA kernel 内不存在任意 block 之间可直接使用的全局 barrier。

一维 CUDA 风格伪代码如下：

```cpp std::uint64_t i = blockIdx.x blockDim.x + threadIdx.x; std::uint64_t stride = blockDim.x gridDim.x;

for (; i < N; i += stride) { y[i] = f(x[i]); } ```

起始索引是线程的全局线性编号：

整个 grid 的逻辑线程数是：

线程处理的索引序列是：

视频在 01:15 引入网格步进，在 01:22 强调同一 worker 得到的是非连续工作。

任意索引 i 都可唯一写成：

其中

它由起始编号为 r 的线程在第 q 轮处理。

因此，只要边界条件正确，网格步进既不遗漏也不重复逻辑索引。

线程 r 的迭代次数是：

每线程约处理 $\lceil N/G\rceil$ 个元素。

总工作量仍是 $O(N)$，不是因为“跨步”就减少了要计算的元素数。

网格步进等价于按余数类做 round-robin 式静态划分：

若高代价元素在索引空间中形成较宽的连续区域，多个线程可能各取到其中一部分。

视频在 01:40 用“每个计算单元都分到一些难题和简单题”说明这个直觉。

但这不是通用的动态负载均衡算法。

它仍然是静态映射。

以下情况不保证改善：

若负载高度不规则，可能需要 work queue、persistent kernel、任务窃取或其他动态调度。

最容易误解的地方是访问连续性。

单个线程在相邻迭代中访问：

它自己的地址确实跨得很远。

但 GPU 合并访存通常看同一条指令下一个 warp 内线程访问的地址集合。

在第 q 轮，相邻线程 r、r+1 访问：

这些逻辑索引仍相邻。

视频在 02:01 给出“整体仍连续”的直觉。

更准确地说：

所以不能只看一个线程的轨迹判断 coalescing。

即使每元素成本相同，grid-stride loop 仍有实际价值：

它不意味着物理线程会被创建 N 次。

同一个逻辑线程会执行多轮 loop body。

若 $f(x_i)$ 包含数据相关分支，不同线程可能执行不同路径。

在 SIMT 模型中，同一 warp 内的路径分歧会产生 divergence。

网格步进把不同位置交给多个线程，可能改变每轮的成本分布。

但它不会自动消除 warp 内分支发散。

甚至在某些数据布局下，连续范围能把相同分支聚到一起，而 round-robin 会把不同分支混在同一 warp。

负载均衡与分支一致性要分别用 profiler 验证。

视频在 00:26 用大而慢、小而快的存储层级解释 tile。

这是教学模型，不指定某个固定硬件部件。

在不同平台上，它可能对应：

grid-stride 只规定索引循环。

它本身不要求把数据显式搬到 shared memory，也不保证产生数据复用。

对只有一次读写的简单逐元素 kernel，额外 staging 可能没有收益。

当 N 很大时，以下乘法可能超过 32-bit：

实现应根据 N 的取值范围选择足够宽的无符号或有符号索引类型。

还要保证：

对超大 N，可写成在加法前判断剩余范围，或使用经证明不会溢出的 64-bit 逻辑。

不能把两种方案简化为：

如果范围划分把连续索引映射给相邻 GPU 线程，同样可以 coalesce。

如果错误地让一个线程独占超长连续区间、而同一 warp 的其他线程跳到远处，才可能破坏事务模式。

关键是“同一时刻相邻 lane 访问什么”，不是算法名称。

若代价模式为“所有偶数索引很重、奇数索引很轻”，且 G=8，grid-stride 会自动把重任务平均给所有线程吗？

不会。

因为 G 为偶数，偶数起始线程之后始终访问偶数索引，奇数线程始终访问奇数索引。

这是“静态 round-robin 不保证任意负载均衡”的反例。

它只改变索引分配，总体仍处理 N 个元素，复杂度仍是 $O(N)$。

coalescing 主要看同一 warp 同一指令的地址集合。

只有成本与索引位置的结构恰好适合交错时才可能改善。

它是由起始索引和固定 stride 决定的静态映射。

板书表达的是尾部等待；真实同步能力由编程模型和 kernel 边界决定。

索引映射与存储 staging 是两件独立的事。

视频在 00:02 提出问题：两个 shape 不同的张量，在底层线性存储中如何完成广播相加？

示例是：

广播后：

视频用“从右向左读，有变就复制，否则就提取”建立直觉。

实现层更准确的说法是：

本课先假设 row-major 连续输入，再推广到任意 stride。

视频从 00:31 解释输出 shape，在 00:38 给出 A 与 B。

A 的 rank 为 2，B 的 rank 为 3。

给 A 左侧补一个大小 1 的前导轴：

于是对齐为：

左侧补 1 只用于统一索引推导，不需要改变数据。

对轴 k，输入大小分别为 $a_k$ 与 $b_k$。

可广播当且仅当：

视频在 00:57 做兼容检查。

对本例：

三对都兼容。

在普通正尺寸例子中，输出轴大小取两者较大值：

所以视频在 01:09 得到：

如果某一轴既不相等、也没有一方为 1，就应报 shape error，不能靠截断或循环取模补救。

视频从 01:27 说明默认连续存储，并在 01:35 用 shape $(3,2,4)$ 画出一维长条。

对 row-major 连续张量

逻辑坐标 $(i_0,i_1,i_2)$ 的线性 offset 是：

也就是：

范围为：

“线性存储”不意味着张量只有一维语义。

shape 与 strides 保存了从逻辑坐标到物理地址的规则。

给定输出 offset t：

去掉第 0 轴贡献：

再得到：

这就是整数除法与取模形式的 unravel index。

真实 kernel 可能用增量计数、预计算 magic number 或合并循环降低除法成本，但语义相同。

对齐后的 A shape 是：

输出坐标是：

A 的第 0、2 轴大小都是 1，所以输入坐标必须固定为 0：

原始 A shape 为 $(2,1)$，row-major strides 可写为：

因此：

不同 $i_0$、$i_2$ 的输出都可能读取同一个 A 元素。

这就是板书“复制”的真实索引含义。

B shape 是：

它的第 1 轴大小为 1，因此：

row-major strides 为：

所以：

最终每个输出元素执行：

视频从 02:41 开始演示 B 的对应过程。

对任意对齐后的输入轴大小 $d_k$ 与输出坐标 $i_k$：

输入物理 offset 为：

这里 $s_k$ 是输入真实 stride。

这个规则比“真的复制 D 份、E 份”更适合实现。

另一种等价表达是先构造 expanded view。

把广播轴的有效 stride 设为 0：

随后直接用输出坐标：

得到与前面完全相同的地址。

stride 0 表示沿该轴移动时物理地址不变。

视频在 03:38 强调“复制不用真的复制”。

常见实现有两类：

但“广播永远零拷贝”也不严谨。

以下操作可能物化：

输出 C 本身通常仍需分配，除非继续被融合且没有可观察的中间结果。

视频从 03:04 推广到更多维度。

无论 rank 多高，流程都不变：

复杂的是索引 bookkeeping，不是生成一个真正的高维立方体。

若输入来自 transpose、slice 或其他 view，它可能不是 row-major 连续。

此时不能用 shape 自动推导连续 strides。

仍应使用：

但 $s_k$ 必须取 view 的真实 strides。

负 stride、storage offset、aliasing 与设备后端限制也可能影响实现。

广播规则决定逻辑坐标，layout 决定物理地址；两层不能混为一谈。

广播从尾轴对齐，缺失的前导轴视为 1。

输入 shape 与 strides 不同，必须先做坐标映射。

常见实现是索引映射或 stride 0，不物化重复数据。

transpose、slice 与 expand 都可能产生非连续 view。

广播检查逐轴兼容性，不比较总元素数。

视频在 00:02 提问：多个线程执行同一段算子代码时发生了什么？

它选择 CPU 线程和逐元素绝对值作为极简模型：

视频在 00:18 引入算子，在 00:21 设定每线程处理 1000 个元素。

这一模型抓住了核心：

但真实 CPU runtime、线程池、调度器和编译器会增加更多层次。

设固定 block size 为：

线程逻辑编号为 id。

安全版本伪代码：

```cpp start = id * B; end = min(start + B, N);

for (i = start; i < end; ++i) { output[i] = abs(input[i]); } ```

视频在 00:35 获取 id，在 00:39 计算 start，在 00:46 计算 end。

视频板书直接使用 $end=start+1000$。

只有 N 恰好覆盖所有整块、或输入预先 padding 时才不会越界。

通用实现必须 clamp 到 N。

函数体本身没有为线程 0、1、2、3 各写一份。

差异来自运行时状态：

因此：

每个线程把自己的 id 代入同一条指令，得到不同 start/end。

这属于 single program, multiple data 的基本思想。

“像执行了不同代码”只是观察结果不同；机器指令通常来自同一代码段。

视频从 00:58 假设：

区间分别是：

这些半开区间不重叠，合起来恰好覆盖 $[0,4000)$。

视频从 01:18 单独观察线程 2。

它看到：

因此 loop body 处理：

同一个变量名 start 在不同线程中可同时拥有不同值，因为它属于各自的调用与执行上下文。

板书中的 get_thread_id() 是抽象接口。

真实环境可能使用：

OS thread id 未必是从 0 开始的稠密整数，也不应直接乘 1000 做数组分块。

高层 runtime 通常显式提供 worker index 或 range。

同一进程的 CPU 线程通常共享：

每个线程通常拥有自己的：

“共享地址空间”意味着一个线程拿到合法指针时，理论上可以访问另一个线程栈中的对象。

因此“栈私有”主要描述每线程独立分配、调用约定与生命周期，而不是独立进程式的地址隔离。

视频从 01:38 进入寄存器视角。

程序语义上，线程 1 可具有：

线程 2 同时可具有：

但不能据此认为每个软件线程永久占据一整套物理寄存器。

当线程运行时，值驻留在当前核心的 architectural/physical register resources 中。

上下文切换时，必要状态由 OS 保存并恢复。

编译器还可能把变量 spill 到栈，或完全消除某个变量。

视频在 01:55 强调共享代码段。

两个线程可以暂时指向同一条机器指令。

它们不必同步前进：

并发表示执行时间区间重叠。

并行表示某一时刻确实在多个执行资源上同时运行。

线程数大于核心数时，不能假设所有线程物理并行。

视频从 02:09 进入栈视角，并在 02:22 定义局部数组 buf。

概念代码：

``cpp float buf[1000] = {}; buf[0] = id; ``

线程 1 的 buf 与线程 2 的 buf 属于不同调用栈帧。

所以：

视频从 02:34 对比这两个值。

语言语义通常只规定作用域、生命周期与可观察行为。

编译器可把局部变量：

大数组还可能造成 stack overflow。

高性能算子常使用 heap、arena、thread-local scratch buffer 或 runtime workspace，而非在函数栈上放巨型临时张量。

若满足：

则逐元素绝对值没有 output data race。

但仍需正确的 join/barrier 建立 happens-before。

若区间计算错误造成重叠，或共享计数器无同步，就会产生 race。

不同线程写不同元素不等于永远没有共享缓存成本。

若两个线程写的元素落在同一 cache line，可能发生 false sharing。

连续大块只在边界附近可能共享少量 cache line，通常比细粒度交错写更友好。

具体影响取决于：

对浮点数，要考虑：

对有符号整数，最小值的正数可能无法用同 dtype 表示。

例如 32-bit two's complement 的 INT_MIN 没有对应的正 int。

真实算子必须采用框架定义的溢出/饱和/异常语义，不能只凭数学 $|x|$ 推断机器行为。

若希望 P 个线程覆盖任意 N，可定义：

这样每块大小最多相差 1，不要求固定 1000。

真实 runtime 还可能动态切小任务，以改善负载均衡。

固定静态分块适合每元素成本近似相同的情况。

代码可重入时，多个线程可持有各自状态执行同一机器指令。

板书 id 是 runtime 提供的稠密 worker index 抽象。

同一进程线程共享地址空间、heap 与 globals。

编译器可寄存器化、标量替换或消除它。

调度、SMT 与超额订阅都会打破一一对应。

视频在 00:02 提出问题：如何用前缀和理解 CSR？

示例稀疏矩阵是：

它有：

nnz 表示格式中存储的非零条目数。

若格式允许显式存储数值 0 或重复坐标，stored entries 与数学上的非零位置数可能不完全相同。

本课先采用无重复、无显式零的标准示例。

视频从 00:22 说明：一个稀疏条目至少需要知道：

例如：

表示：

这里使用 0-based index。

只有 value 不知道放在哪里；只有坐标也不知道条目取值。

视频在 00:34 引入 COO。

对示例矩阵，一种按行排列的 COO 是：

对任意位置 k，三数组共同组成：

三个数组必须同步重排。

不能只交换 values 而保持坐标不动。

COO 便于增量收集条目，但每个条目都显式重复保存行号。

视频从 01:00 引入每行非零计数。

对示例：

含义是：

视频在 01:11 读出 $[2,1,2]$。

row_count 是构造时的中间数组。

标准 CSR 通常不需要同时保存它，因为 row_ptr 已能通过差分恢复。

对：

定义：

这就是带前导 0 的 exclusive prefix sum。

视频在 01:27 开始求前缀和，在 01:33 命名 row_ptr。

本例：

长度是：

最后一个值满足：

视频在 01:44 解释末项 5。

第 r 行的条目位于：

本例：

这三个区间正好覆盖：

row_ptr 保存的是每一行在扁平 entries 数组中的边界。

视频从 01:54 用差分解释可逆性。

例如第 1 行：

row_ptr 中的 3 表示前两行一共占用 3 个 entries。

它既是第 2 行的起点，也是前两行累计 nnz。

CSR 通常保存：

视频从 02:23 对比 CSR 与 COO，并在 02:31 说明替换关系。

两种格式都有三个主数组，但数组长度不同：

当 $nnz\gg M$ 时，CSR 的行信息通常更紧凑。

对任意 k，只需找到唯一的 r，使：

就知道第 k 个 entry 属于第 r 行。

然后：

因此条目的 value、row、column 信息仍完备。

视频从 02:42 做反向还原，在 03:11 说明行号已经确定。

若某行没有 stored entry：

于是：

也就是 row_ptr 出现重复值。

例如：

表示第 1 行区间为：

即空区间。

重复 row_ptr 不是错误，恰好是空行的合法表示。

视频从 03:32 比较顺序要求。

准确边界是：

很多库定义 sorted/canonical CSR，要求行内列索引升序并合并重复坐标。

但基础 CSR 表示本身不必天然满足这些附加规范。

所以板书“CSR 必须行优先”应理解为按行成段，而不是每行列号必定有序。

给定 M 行、nnz 个 COO entries：

串行复杂度约为：

外加可选排序成本。

并行构造还要处理 histogram 与 scatter 的冲突，可能使用 atomics、分块计数或 segmented prefix sum。

prefix sum 不只是数学定义。

它是 GPU/并行算法中的基础 primitive。典型并行 scan 用树形 upsweep/downsweep 或分层 block scan，把串行依赖重组成 $O(\log M)$ 的并行深度，并保持 $O(M)$ 级总工作量。

具体 work efficiency、同步与内存访问取决于实现。

不能只用并行深度替代实际运行时间。

设 value 占 $s_v$ bytes，index 占 $s_i$ bytes。

CSR 主数组约占：

COO 约占：

CSR 相比 COO 用 row_ptr 替代逐项 row_idx。

是否更省内存取决于 M 与 nnz 的关系、index width 和格式元数据。

CSR 很适合逐行运算，例如 SpMV：

``text for row in [0, M): sum = 0 for k in [row_ptr[row], row_ptr[row+1]): sum += values[k] * x[col_idx[k]] y[row] = sum ``

每行边界可 O(1) 获取。

但性能仍受：

CSR 解决存储与定位，不自动解决所有稀疏计算瓶颈。

它是第 r 行在 entries 数组中的起始 offset。

最后一项是 nnz。

只能把三元组作为整体重排。

是否合并/清理取决于构造与 canonicalization。

性能取决于行长度分布、算子、硬件与访存模式。

上一课用前缀和构造 CSR。

本课回到更直观的 COO，回答视频在 00:00 提出的问题：

SpMV 是：

其中：

是稀疏矩阵，

是稠密向量。

视频在 00:07 定义 SpMV。

本课使用无重复坐标、无显式存储零的 COO 示例，再讨论这些边界。

对每一行 i：

示例：

因此：

数学式中，零项对结果没有贡献。

关键是存储格式能否让程序根本不枚举这些位置。

朴素 dense matvec：

``text for i in [0, M): sum = 0 for j in [0, N): sum += A[i,j] * x[j] y[i] = sum ``

无论 $A_{ij}$ 是否为 0，循环都遍历全部 $M\times N$ 个坐标。

视频从 00:11 说明传统矩阵-向量乘，在 00:16 强调遍历矩阵。

工作量约为：

这里“会计算零”是算法层的说法。

编译器一般不能预先知道任意运行时 dense A 中哪些位置为零。

视频从 00:26 使用 COO 举例。

示例 COO：

三个数组长度都是 $nnz=5$；第 k 个 entry 是 $(row\_idx[k],col\_idx[k],values[k])$。

没有出现在数组中的坐标被解释为隐式零。

视频在 00:47 选择数值 3。

对相应 k：

所以它表示：

三个数组必须按同一个 k 对齐。

如果只重排 values，坐标语义就会损坏。

先把输出初始化为 0：

``text y[0:M] = 0 ``

再遍历 stored entries：

``text for k in [0, nnz): i = row_idx[k] j = col_idx[k] y[i] += values[k] * x[j] ``

视频从 01:05 开始遍历三个数组。