# 向量加法

## 实现目标

`vector_addition.py` 使用 Triton 实现两个等长向量的逐元素加法，并与 PyTorch
参考结果和吞吐量进行比较。该示例展示了 Triton 内核最基本的启动、寻址和边界
保护方式。

## 核心数据流

1. `add()` 分配与输入同形状的输出张量，并根据元素总数构造一维启动网格。
2. 启动网格中的每个程序通过 `tl.program_id(0)` 获得 PID。
3. `add_kernel()` 用 `program_id * block_size` 计算程序负责的起始位置，再生成一组
   连续元素偏移量。
4. 内核使用 `offsets < n_elements` 屏蔽末尾越界位置，从 DRAM 加载两个输入块，
   在 SRAM 中相加并将结果写回输出。

## 关键细节与限制

- `block_size=1024` 是编译期元参数；改变它会生成不同的内核实例。
- 块大小不必整除向量长度，因为加载和存储都使用同一个边界掩码。
- 输入和输出必须位于当前 CUDA 设备，且实现不支持广播。
- 包装函数没有自定义反向传播，但 PyTorch 可在需要时使用普通加法作为参考。

## 测试与基准

直接运行会先用 PyTorch 结果执行数值测试：

```bash
python3 vector_addition.py
```

传入 `--benchmark` 会额外比较 Triton 与 PyTorch 的有效内存带宽：

```bash
python3 vector_addition.py --benchmark
```

## 配套视频

[bilibili](https://www.bilibili.com/video/BV1fMyWBgERM?spm_id_from=333.788.videopod.episodes&vd_source=e98b669ccbafff4b5aa59dd6303b722f&p=9)
