题目1: 求解二维数组每行的累加值 20%

基于 DCU 编写完整 并行程序 ：

1. 实现初始化二维数组 float A[10000][10000] ；
2. 求该数组 每行的累加值 ，并将该值保存到相应行的第一个元素中。注意优化程序性能，并说明优化方法 。

环境要求

注意

该项目需在支持 ROCm 的 Linux 机器（如 海光DCU ） 上编译运行（详见请至：环境配置）

目录结构

.
── q1
   ├─ hip.cpp           海光DCU性能优化版
   ├─ raw.cpp           纯CPU版
   └─ run.sh            编译脚本

源代码拉取

$ git clone https://static.ksuser.cn/dcu/q1.git

编译运行

请转到实际部署章节

实现概述

矩阵维度为 N = 10000，共 N×N = 1e8 个 float，约 400 MB 显存。

• 先在 GPU 上生成 10000×10000 随机矩阵，再对每一行执行块级并行归约，最终将行和写回该行的首元素 (A[i][0]) 并同步保存在 rowSums[i]
• sum_rows_kernel 通过每行一个线程块 + 共享内存树形归约显著提升性能；线程 0 同时更新 rowSums[row] 与 A[row * N]，使该值被保存到相应行的第一个元素中
• 主程序使用 HIP 事件测量两个核函数的运行时间，并在主机侧同时打印行首元素与 rowSums，验证写回结果一致。

程序包含两个 GPU 内核：

• 初始化内核 init_kernel：
  • 逐元素生成伪随机数（线性同余法 LCG），写入 A[idx]。
• 行求和内核 sum_rows_kernel：
  • 每“行”对应一个线程块（block），块内线程并行遍历该行的不同列并累加，随后用共享内存做块内归约，得到该行总和
  • 为减少无谓拷贝，行和写入独立的 rowSums 数组，主机侧仅拷贝需要的若干个结果进行打印验证

性能优化方法

1. 行级并行归约

   • 每行一个线程块，块内线程以 stride 方式遍历列元素并局部累加，随后在共享内存中做树形归约（多轮 __syncthreads()）
   • 相比“每行一个线程串行相加 10000 次”，吞吐通常提升一个数量级以上

2. 独立行和输出

   • 将行和写入连续的 rowSums，主机侧仅拷贝所需的前几个元素进行校验，避免拷贝整行数据，减少 PCIe/D2H 开销与缓存污染

3. 使用设备事件计时

   • 采用 hipEventRecord/hipEventElapsedTime 统计 kernel 的纯 GPU 时间，避免主机计时误差，便于对比不同并行度与块大小

4. 其他细节

   • 维度常量化 constexpr int N = 10000; 便于编译器优化
   • 指针加 __restrict__，帮助别名分析与自动向量化
   • 线程块大小默认 256（2 的幂便于归约），可按设备实际情况调为 512/1024 做对比

实际部署

在 海光DCU 设备上：

chmod +x ./run.sh
./run.sh
./raw # 运行纯CPU版本
./hip # 运行海光DCU性能优化版本

run.sh 中的编译参数详解：

参数	含义	说明
hipcc	HIP 编译器	自动调用 clang++（或 ROCm 版的 clang）进行编译，负责将 HIP 源码编译为可在 HCU 上运行的代码。
-O3	优化等级 3（最高级别）	启用激进的优化策略（循环展开、内联、常量传播、消除冗余计算等），让代码运行得更快，但编译时间更长。
-ffast-math	启用快速数学优化	允许编译器假设浮点运算满足结合律、分配律等数学规律（即使在 IEEE754 下不完全正确），牺牲精度换取性能。
-DNDEBUG	定义宏 NDEBUG	表示关闭断言（assert()），避免运行时断言检查造成性能损失。
raw/hip.cpp	输入源文件	需要编译的 C++ 源文件。
-o raw/hip	输出文件名	指定生成的可执行文件名称为 raw/hip（如果省略则默认输出为 a.out）。

运行输出示例（数值视设备不同会变化）：

Hip

初始化耗时: 0.0024677 秒
累加计算耗时: 0.000630364 秒
第 0 行的累加值 = 4997.21
第 1 行的累加值 = 5000.17
第 2 行的累加值 = 5000.12
第 3 行的累加值 = 5000.08
第 4 行的累加值 = 5001.03
...

Raw

初始化耗时: 1.11158 秒
累加计算耗时: 0.0321607 秒
第 0 行的累加值 = 4965.18
第 1 行的累加值 = 5005.03
第 2 行的累加值 = 5008.86
第 3 行的累加值 = 4975.35
第 4 行的累加值 = 5016.61
...

复杂度与资源占用

• 计算：合计约 1e8 次加法；并行实现将每行工作分给 256/512 个线程，大幅缩短 wall-time。
• 显存：矩阵数据约 400 MB，另有 rowSums 约 40 KB。
• 拷贝：仅拷贝需要的行和（例如前 5 行约 20 字节），远小于拷贝整行数据。

验证与注意事项

• 保证 hipDeviceSynchronize()/事件同步后再读取与计时，避免异步影响。
• 确保线程配置覆盖所有元素/行，并做好越界判断。
• 确保 hipcc 编译可用；编辑器可能提示找不到 hip/hip_runtime.h，那是本地未配置头文件路径所致，不影响在目标环境中的编译。

许可与引用

• 本项目遵行 Apache License 2.0协议