📚 课程目录第3章 CUDA编程模型在IGX上的映射

3. CUDA编程模型在IGX上的映射

CUDA 编程模型本身是跨 GPU 产品线统一的抽象：线程按 grid → block → warp → thread 组织，内存按 register / shared / L1 / L2 / global 分层，执行以 warp 为调度单位、以 SM 为资源容器。但 IGX 与桌面 GPU 的差别不在抽象层，而在抽象层落到硬件时的“参数取值”和“约束条件”。理解 IGX 上的 CUDA 映射，本质上是理解同一套编程模型在嵌入式、长生命周期、功能安全场景下的具体落地形态。

线程层次：从 grid 到 warp 的映射

在 IGX 上，一个 kernel 启动时，grid 被切分为若干 block，block 被分配到各个 SM 上执行。每个 block 内部再按 32 个线程一组划分为 warp，warp 是 SM 上真正被调度和发射的最小单位。IGX Orin 系列（如 IGX Orin 700）的 GPU 部分基于 Ampere 架构，每个 SM 包含 4 个 warp 调度器，每个调度器每周期可发射一条指令。这意味着一个 SM 上同时驻留的 warp 数量直接决定了延迟隐藏能力。

与桌面 GPU 不同，IGX 的 SM 数量通常更少（例如 IGX Orin 700 的 GPU 约 16 个 SM，而桌面 RTX 4090 有 128 个 SM）。因此，在 IGX 上“用 block 数量堆并行度”的策略收益有限，更关键的是让每个 SM 上的 warp 占用率（occupancy）足够高。一个常见的适配要点是：block 大小优先选择 128 或 256 线程，而不是桌面 GPU 上常见的 512 或 1024。原因是 IGX 上 SM 的寄存器文件和 shared memory 总量与桌面同架构产品接近，但 SM 数量少，过大的 block 会导致单个 SM 上可驻留的 block 数下降，反而降低 warp 并发度。

内存层次：容量、带宽与一致性

IGX 的内存层次与桌面 GPU 结构一致，但容量和带宽的绝对值不同。以 IGX Orin 700 为例，GPU 与 CPU 共享 64 GB LPDDR5 内存，GPU 可访问的 global memory 是其中的一部分（通常通过 carveout 配置，例如 32 GB 给 GPU）。这与桌面 GPU 拥有独立显存（如 24 GB GDDR6X）的模式不同：IGX 是统一内存架构，CPU 和 GPU 共享物理内存，但访问路径和带宽不同。

在 IGX 上，global memory 的带宽约为 204.8 GB/s（LPDDR5 256-bit），而桌面 RTX 4090 的显存带宽超过 1 TB/s。这意味着在 IGX 上，内存带宽更容易成为瓶颈。适配要点是：尽量让数据在 shared memory 或寄存器中复用，减少对 global memory 的重复访问。例如，在矩阵乘法中，将 tile 加载到 shared memory 后，每个线程从 shared memory 读取多次，而不是每次都从 global memory 读取。

另一个关键差异是 L2 cache 的容量。IGX Orin 的 L2 约为 4 MB，而桌面同架构产品可能为 6 MB 或更大。对于需要跨 block 复用数据的 kernel，L2 命中率会直接影响性能。在 IGX 上，如果数据量超过 L2 容量，性能下降会比桌面 GPU 更明显。

执行模型：功耗、持续性与功能安全

桌面 GPU 的执行模型假设“短时爆发、散热充足、供电稳定”。IGX 则假设“持续运行、功耗受限、环境温度可能较高”。这带来两个直接后果：第一，IGX 的 GPU 频率会在持续负载下动态调整，kernel 的执行时间可能比桌面 GPU 更不稳定；第二，IGX 上不适合长时间满占用运行，需要留出功耗余量给 CPU 和其他加速器。

在 IGX 上做 CUDA 优化时，一个实用的做法是：先用 nvidia-smi 或 tegrastats 观察 GPU 利用率和功耗，再决定是否要降低 block 数量或引入 __syncthreads() 之外的同步点。如果 GPU 利用率长期接近 100% 而功耗接近上限，说明 kernel 可能触发了降频，此时应优先减少计算量或提高数据复用率，而不是继续增加并行度。

功能安全方面，IGX 支持 ECC 内存和部分 ASIL 相关机制。在 CUDA 编程中，这意味着需要关注内存访问的确定性：避免依赖未初始化内存、避免越界访问、避免在 kernel 中使用不确定的浮点归约顺序。这些在桌面 GPU 上可能只是“性能问题”，在 IGX 上可能升级为“安全问题”。

与桌面 GPU 的差异与适配要点

综合来看，IGX 与桌面 GPU 在 CUDA 映射上的差异可以归纳为三点：SM 数量少、内存带宽低、功耗约束强。对应的适配策略是：

线程层次：block 大小优先 128 或 256，grid 维度按实际数据量计算，不要盲目放大。使用 cudaOccupancyMaxActiveBlocksPerMultiprocessor 检查每个 SM 上可驻留的 block 数。
内存层次：把 shared memory 当作“第一公民”，在 kernel 内部显式做 tile 加载和复用。对于跨 block 复用的数据，评估 L2 容量是否足够，必要时用 cudaMemAdvise 或 cudaMemPrefetchAsync 管理统一内存的迁移。
执行模型：用 tegrastats 或 nvidia-smi dmon 监控功耗和频率，避免长时间满占用。在功能安全相关场景中，使用 compute-sanitizer 检查内存访问的确定性。

在 IGX 上移植桌面 CUDA 代码时，建议先做一次“参数审计”：把 block 大小、grid 大小、shared memory 用量、寄存器用量四个参数列出来，对照 IGX 的 SM 数量和内存容量重新计算。很多在桌面 GPU 上跑得通的配置，在 IGX 上会因为 SM 数量不足或 L2 容量不够而性能骤降。

最后需要强调的是，IGX 上的 CUDA 优化不是“把桌面 GPU 的代码改小”，而是“重新理解资源约束下的并行度分配”。桌面 GPU 的优化目标是“用更多 SM 和更大带宽换时间”，IGX 的优化目标是“在有限 SM 和带宽下，让每个 warp 做更多有效工作”。这个思维转变，是 IGX CUDA 编程的核心。