Learn Kokkos Module 1 - 简介、构建与并行分发Learn Kokkos Module 1 - 简介

Learn Kokkos Module 1 - 简介、构建与并行分发

1. 核心定位与工程目标

1.1 性能可移植性

Kokkos 的核心目标是实现 性能可移植性（Performance Portability） ：

用一套 C++ 源码
面向多种硬件后端编译运行
包括：
- 多核 CPU
- NVIDIA GPU
- AMD GPU
- Intel GPU

这并不只是“代码能运行”，更强调 在不同架构上都尽量获得高性能。

1.2 Kokkos 解决的核心痛点

Kokkos 不只是一个并行 API 库，它更重要的价值在于解决：

并行执行问题
数据布局问题（Data Layout Problem）

在异构计算中，单纯把循环并行化远远不够。
如果数据布局与底层硬件访问模式不匹配，就会产生严重性能问题，尤其在 GPU 上会出现：

非合并访存（Non-coalesced Access）
带宽利用率大幅下降
整体性能显著劣化

因此，Kokkos 的设计从一开始就把 执行模型 和 数据布局 一起考虑。

1.3 技术路线

Kokkos 采用 C++ 模板库 的方式实现，具有以下特点：

零运行时抽象开销
编译期分发能力强
易于与现代 C++ 工程集成

同时它深度对齐 ISO C++ 标准演进，例如：

std::atomic_ref
std::mdspan

这些方向都与 Kokkos 的设计思想高度一致。

2. 并行执行的工程抽象：组件解耦

传统基于 Pragma 的并行模型，例如：

#pragma omp parallel for

虽然适合简单 CPU 并行，但很难优雅处理：

异构设备执行
复杂数据布局
显式内存空间映射

Kokkos 将并行控制流拆解为三个正交概念。

2.1 Pattern：模式

定义“做什么类型的并行计算”。

常见模式：

parallel_for
- 独立迭代
parallel_reduce
- 并行归约，例如求和、最值等

2.2 Policy：策略

定义“工作如何映射到硬件”。

通常包括：

迭代范围
执行空间
调度粒度
线程层级映射规则

也就是说，Policy 决定的是 运行方式。

2.3 Body：计算体

定义“单个工作项具体做什么”。

它是真正的业务逻辑代码，可以通过以下形式提供：

Functor
Lambda

2.4 工程契约：不保证顺序与并发数

这是使用 Kokkos 时非常重要的一条工程原则：

不保证执行顺序
不保证具体并发数
不保证某轮迭代由哪个线程执行

这意味着底层系统可以自由优化，例如：

向量化
重排
分块
乱序执行

因此，计算体必须满足：

各次迭代之间相互独立
不能依赖“前一次”或“后一次”迭代结果
不能隐式共享不安全状态

3. 计算体分发与内存隔离限制

在异构架构中，Host（CPU）与 Device（GPU）通常处于不同内存空间，因此状态传递必须非常谨慎。

3.1 计算体的传递方式

Kokkos 支持两种主要形式：

Functor
C++11 Lambda

但无论哪种方式，都必须满足设备端可见性要求。

3.2 必须按值捕获

在 Lambda 中应使用：

[=]

或使用：

KOKKOS_LAMBDA

原因是：

CPU 栈变量位于 Host 内存
GPU 无法直接访问 Host 栈地址
如果使用 [&] 引用捕获，Device 会尝试访问无效地址

结果通常是：

非法内存访问
Kernel 失败
程序崩溃

3.3 禁止捕获标准容器

例如：

std::vector
其他深拷贝语义容器

不应直接按值捕获进 Device Lambda。

原因

第一，深拷贝代价高

会在 Host 端发生昂贵拷贝

第二，设备端通常不支持

GPU 代码不能可靠使用 std::vector

第三，语义错误

即使捕获成功，也只是副本
Device 对副本的修改不会自动反映回原对象

3.4 正确做法

应使用：

具备浅拷贝语义
支持设备访问
生命周期与空间可控

的抽象结构，例如：

Kokkos::View

这一部分通常在后续模块详细展开。

4. 并行规约机制与竞态防御

并行循环里最典型的错误之一，就是对共享状态进行直接更新。

例如：

total += local_val;

如果多个线程同时执行这类操作，就会产生：

读-改-写冲突
数据竞争（Race Condition）
结果不确定甚至错误

4.1 为什么普通写法不安全

共享变量的更新通常包含三个步骤：

读取旧值
计算新值
写回结果

多个线程并发执行时，这三个步骤可能交错，导致部分更新丢失。

4.2 `parallel_reduce` 的防御设计

Kokkos 不要求用户自己加锁，而是通过接口设计规避竞态。

机制一：线程私有累加器

Lambda 会被注入一个局部更新变量，例如：

double& update

这个变量是当前执行上下文的私有状态，不与其他线程共享。

机制二：局部无锁更新

每个线程只更新自己的局部副本，因此不会发生全局冲突。

机制三：最终硬件归并

所有局部结果由 Kokkos 运行时统一合并。
底层会使用目标平台最优实现，例如：

CPU 树形归约
GPU 共享内存归约
Warp 级优化

4.3 工程结论

凡是出现以下需求时：

求和
求最大值 / 最小值
计数
聚合统计

都应优先考虑：

Kokkos::parallel_reduce(...)

而不是在 parallel_for 中直接修改共享变量。

5. 性能模型与调度开销优化

并行不是免费的。
一次并行执行的时间可以粗略写成：

[
Time = \alpha + \frac{\beta \times N}{P}
]

其中：

(\alpha)：启动开销
(\beta)：单次迭代成本
(N)：任务总量
(P)：并发度

5.1 启动开销不可忽视

在 HPC 和 GPU 场景中，(\alpha) 往往并不小，例如：

CUDA Kernel Launch 往往是微秒级

如果任务很小，可能出现：

启动成本远高于真正计算成本
GPU Offload 反而比 CPU 更慢

也就是说，不是所有循环都值得并行，更不是所有循环都值得搬到 GPU。

5.2 优化手段一：内核融合

当存在多个连续操作并且针对同一批数据时，应优先考虑融合为一个内核。

好处

减少 Kernel Launch 次数
降低总启动成本
提升缓存局部性
减少显存往返与中间写回
降低内存带宽压力

思路

将原本多个独立的：

parallel_for
parallel_reduce

整合为一个更大的执行体

前提是逻辑允许合并。

5.3 优化手段二：内核命名

建议为每个并行分发加标签：

Kokkos::parallel_for("MyKernel", N, ...);

作用

这几乎没有运行时成本，但价值很高：

便于 Nsight 定位瓶颈
便于 Tau 分析调用轨迹
便于大型工程中的性能回溯
便于识别异步流中的具体内核

在复杂系统里，没有名字的 Kernel 几乎无法高效分析。

Learn Kokkos Module 1 - 简介、构建与并行分发