# L1 和 L2 缓存带宽测试

## 文件结构

- `gpu_clock_lock.sh`: 锁频脚本
- `l1_cache/`
    - `l1_cache_1fp/`: scalar float 下 L1 缓存带宽测量
    - `l1_cache_4fp/`: vector float4 下 L1 缓存带宽测量
- `l2_cache/`
    - `l2_cache_1fp/`: scalar float 下 L2 缓存带宽测量
    - `l2_cache_4fp/`: vector float4 下 L2 缓存带宽测量

## L1 缓存带宽测试

测量 GPU L1 缓存的读带宽，分别使用 `ld.global.ca.f32` 和 `ld.global.ca.v4.f32` 指令（ca = cache at all levels，经过 L1 和 L2）。
每个 SM 上的工作集大小可控，且每个 block 访问自己独立的数据块，使 L1 命中。

## 文件结构

以 `l1_cache_1fp/` 举例:
    - `makefile`：编译脚本
    - `bench_l1_cache.cu`: 源代码
    - `ncu_profile.sh`：Nsight Compute 分析脚本
    - `results/`
        - `result_l1_cache.txt`：l1 cache 带宽测量结果
        - `result_l1_cache_%1_%2.csv`：block 数量为 `%1`、每个 block 内线程数量为 `%2` 时的 Nsight Compute 分析结果

注意，`l1_cache_4fp/`，`l2_cache_1fp/` 和 `l2_cache_4fp/` 遵循类似的文件结构。


## L2 缓存带宽测试

测量 GPU L2 缓存的读带宽，使用 `ld.global.cg.f32` 和 `ld.global.cg.v4.f32` 指令（cg = cache globally，通常绕过 L1，只经过 L2/DRAM 路径）。
