# DRAM 带宽测试

## 测试过程

DRAM 测试通过手写 PTX 指令控制全局内存访问的 cache policy，例如：

```bash
ld.global.cs
ld.global.cg
ld.global.cv

st.global.cs
st.global.cg
```

其中，`.cv` 仅用于 load 指令；由于 PTX 中不存在合法的 `st.global.cv` 形式，因此 `.cv` 模式下的 store 使用 `.cs`。

测试代码使用 grid-stride loop 进行顺序访问，使 warp 内线程访问连续地址，从而保证全局内存访问能够合并。测试分别覆盖两种数据访问粒度：

- `scalar float`
- `vector float4`

并测试以下四类访存模式的有效带宽：

- `Read`
- `Write`
- `Copy_1R1W`
- `Mixed_5R1W`

最终通过比较不同 cache operator 和不同循环模式下的 GB/s，评估对应配置下的 DRAM bandwidth。


## 文件结构

- `gpu_clock_lock.sh`: 锁频脚本


### scalar float

- `1fp/`
  - `makefile`：编译脚本
  - `bench_dram.cu`：源代码
  - `ncu_profile.sh`：Nsight Compute 分析脚本
  - `results/`
    - `result_%1_%2.txt`：cache policy 参数为 `%1`、循环模式参数为 `%2` 时的测量结果
    - `result_%1_%2_%3_%4.csv`：cache policy 参数为 `%1`、循环模式参数为 `%2`、block 数量为 `%3`、每个 block 内线程数量为 `%4` 时的 Nsight Compute 分析结果

### vector float4

- `4fp/`
  - `makefile`：编译脚本
  - `bench_dram.cu`：源代码
  - `ncu_profile.sh`：Nsight Compute 分析脚本
  - `results/`
    - `result_%1_%2.txt`：cache policy 参数为 `%1`、循环模式参数为 `%2` 时的测量结果
    - `result_%1_%2_%3_%4.csv`：cache policy 参数为 `%1`、循环模式参数为 `%2`、block 数量为 `%3`、每个 block 内线程数量为 `%4` 时的 Nsight Compute 分析结果


## 测试选项

### 单线程单次内存访问字节数

- `scalar float`：每次 load/store 访问 `4 B`
- `vector float4`：每次 load/store 访问 `16 B`

### Cache Policy

- `.cv`：仅 load 使用 `.cv`，store 使用 `.cs`
- `.cs`
- `.cg`

### 循环模式

#### `kernel`

迭代发生在 kernel 内部，即 host 端只 launch 一次 kernel，kernel 内部通过 `iters` 控制重复访问次数。

```cpp
// ------------------------------------------------------------
// Mixed 5-read-1-write kernel
// ------------------------------------------------------------
__global__ void dram_mixed_5r1w_kernel(
    const float* __restrict__ a0,
    const float* __restrict__ a1,
    const float* __restrict__ a2,
    const float* __restrict__ a3,
    const float* __restrict__ a4,
    float* __restrict__ dst,
    size_t elems,
    int iters,
    int policy
) {
    ...
    ...

    for (int it = 0; it < iters; ++it) {
        for (size_t i = tid; i < elems; i += stride) {
            float v0 = load_policy_float(a0 + i, policy);
            float v1 = load_policy_float(a1 + i, policy);
            float v2 = load_policy_float(a2 + i, policy);
            float v3 = load_policy_float(a3 + i, policy);
            float v4 = load_policy_float(a4 + i, policy);

            float out = v0 + v1 + v2 + v3 + v4;

            store_policy_float(dst + i, out, policy);
        }
    }

    ...
    ...
}
```

#### `launch`

迭代发生在 kernel 外部，即 host 端多次 launch kernel，每次 kernel 内部只执行一次访问。

```cpp
static float time_benchmark(
    Launcher launcher,
    int warmup_iters,
    int measured_iters,
    LoopMode mode
) {
    ...
    ...

    if (mode == LOOP_LAUNCH) {
        for (int r = 0; r < measured_iters; ++r) {
            launcher(1, r);
        }
    } else {
        launcher(measured_iters, 0);
    }

    ...
    ...
}
```

## 结论和分析

## 结论

选用 cache policy 参数为 `.cv`、循环模式参数为 `launch` 的测量结果作为最终结论。具体结果参见：

- [scalar float](1fp/results/result_cv_launch.txt)
- [vector float4](4fp/results/result_cv_launch.txt)

## 分析

### 结论的正确性

测试结果见：

- [scalar float](1fp/results/result_cv_launch.txt)
- [vector float4](4fp/results/result_cv_launch.txt)

在对应的 `.csv` 文件中，可以观察到平均 L2 缓存命中率 `lts__t_sector_hit_rate.pct` 和平均 L1 缓存命中率 `l1tex__t_sector_hit_rate.pct` 基本接近 0。这说明测试过程中数据主要来自 DRAM，缓存命中对结果的影响较小，因此该配置下测得的带宽可以作为 DRAM bandwidth 的有效参考。



此外，`vector float4` 和 `scalar float` 的测试结果接近，说明当前测试已经主要受底层 DRAM 带宽限制。`float4` 每条指令访问 16B，而 `float` 每条指令访问 4B，因此 `float4` 可以减少访存指令数量并提高单线程单次访问粒度；但在 warp 内访问已经合并、DRAM 带宽已经接近饱和的情况下，`float4` 的有效带宽不一定会达到 `scalar float` 的 4 倍。若两者实测带宽接近，通常说明瓶颈已经从指令数量转移到底层内存带宽。

### 读写资源可能存在共享

单独 `Read` 和单独 `Write` 的实测带宽最高，`Mixed_5R1W` 次之，而 `Copy_1R1W` 相比前三者明显下降。这表明底层读写路径、调度资源或内存系统带宽可能存在共享，读写同时发生时会产生资源竞争，并可能引入读写切换、队列调度等额外开销。