NVIDIA AI/GPU Networking 面试两天冲刺作战手册

岗位：System Software Architect, AI and GPU Networking
面试时间：2026-06-15 20:00-22:00 GMT+8
当前目标：用 2 天周末把 JD 技术要求恢复到“面试可深挖、能画图、能 debug、最好有实战证据”的状态。

0. 这份文档怎么用

你要的不是“CV 包装文档”，而是：

1. JD 要求什么技能。
2. 你 CV 里有什么对应证据。
3. 哪些是真匹配，哪些是短板。
4. 面试会问哪些技术问题。
5. 面试官顺着每个技能点会怎么继续深挖。
6. 每个深挖点需要什么实验、代码、benchmark 或架构图作为证据。
7. 到什么程度才算“掌握”，不是只知道概念。

这份文档不要按“任务数量”读。正确读法是：

先看 JD 技能点。
再看你 CV 里能接上的证据。
然后准备这个技能点下的追问链。
最后才做对应的延展任务/实验。

也就是说，任务不是主线。任务只是为了支撑面试回答的证据。

这里的“2 天完全掌握”定义为面试验收标准：

• 能画出核心 data path / control path。
• 能解释关键机制，不只是背定义。
• 能回答二级、三级追问。
• 能给出 debug checklist。
• 至少对 CUDA/NCCL/LLM serving 做出可运行实验或完整实验设计。
• 对 RDMA/UCX/GPUNetIO 这种需要硬件环境的内容，能说明真实搭建方法、可观测指标、失败时怎么排查。

现实边界：

• 2 天内不可能真正成为 UCX/NIXL/GPUNetIO 内核专家。
• 但可以达到这个岗位首轮技术面试需要的“系统架构级掌握”：知道每层解决什么问题、数据怎么走、哪里会慢、怎么验证。

当前准备环境：

• 你现在有一台 Windows GPU 机器，这是本次准备的主环境。
• Windows 原生负责：NVIDIA driver、CUDA Toolkit、Nsight Systems / Nsight Compute GUI、CUDA Samples。
• WSL2 Ubuntu 负责：Linux CUDA CLI、Python/vLLM、UCX/NCCL 工具、接近 NVIDIA AI infra 的 Linux 工作流。
• 没有 RDMA/RoCE/NVIDIA NIC/DOCA 环境时，不硬装 GPUNetIO；这部分用架构图、实验设计、debug checklist 补齐。

推荐目录：

learning-cuda/
  NVIDIA_Interview_Master_Plan_Deep_CN.md
  labs/
    llm_serving/
    cuda_profiling/
    nccl_tests/
    ucx_rdma_notes/
    diagrams/

0.1 最重要：不要按章节号从头读到尾

这份文档现在内容很多，如果你从 0 -> 17 顺序读，会很容易失焦。
正确结构应该是四层：

flowchart TB
    L1[第一层: 面试地图] --> L2[第二层: 技术主线]
    L2 --> L3[第三层: 实战证据]
    L3 --> L4[第四层: 口头回答]

    L1 --> A[JD 要求 / CV 对应 / 追问树]
    L2 --> B[Serving / NIXL / NCCL / UCX / GPUNetIO / CUDA / Architecture]
    L3 --> C[CUDA code / vLLM benchmark / NCCL tests / UCX checklist / data path 图]
    L4 --> D[英文回答 / 深挖回答 / 反问]

你每天打开文档时，应该先问自己：

我现在是在补哪个 JD 技能点？
这个技能点我 CV 里怎么接？
面试官会从哪里开始问？
他继续追问会追到哪一层？
我需要跑什么实验或画什么图来支撑？
我最后能不能用英文 60-90 秒讲清楚？

所以推荐阅读路径不是章节顺序，而是下面这个顺序。

0.2 推荐阅读路径：按面试准备顺序读

第 0 轮：只建立地图，不学细节

目标：知道这个岗位到底考什么，不要被术语淹没。

读这些：

这一轮不要跑代码，也不要深挖 UCX/NIXL/GPUNetIO。
你只需要能回答：

这个岗位的核心不是“会背很多库名”，而是：
AI inference runtime + GPU communication/data movement + CUDA profiling + system architecture。

第 1 轮：先准备最高概率被问的主线

目标：先把面试最可能问到的部分讲顺。

按这个顺序：

这一轮每个技能只做三件事：

1. 读 JD 要求是什么。
2. 读 CV 怎么接。
3. 读面试会问什么和标准回答。

先不要把每个实验都跑完。
先保证你能口头讲出每条技术主线。

第 2 轮：补 NVIDIA 特有技术栈

目标：把 NVIDIA 相关但你 production 经验不强的内容补到“能解释、能设计验证”的程度。

按这个顺序：

这一轮最重要的是“不乱说”。
你要能明确区分：

NCCL: collective communication across ranks。
NIXL: inference state/data movement，比如 KV transfer。
UCX: transport abstraction，可走 TCP/RDMA/shared memory/CUDA-aware path。
GPUDirect RDMA: RNIC 直接访问 GPU memory 的 data path。
GPUNetIO: 更 GPU-centric 的 packet/network processing 编程模型。
GR00T: robotics VLA workload，不是这个 JD 的核心要求。

第 3 轮：做证据实验，不是做学生任务

目标：每个实验都要能服务一个面试追问。

只做这些 P0 实验：

实验完成后不要只保存输出。你要把结果转成面试语言：

我跑了什么？
控制变量是什么？
看了哪些指标？
结果说明什么？
如果结果异常，我会怎么查？
这和 JD 哪个技能点相关？

第 4 轮：只练口头回答

目标：把技术内容压缩成可讲的英文回答。

读这些：

0.3 每个技能点的标准学习结构

后面每个技能点都应该按同一个结构学习。
如果某一章内容太长，你只抓这个结构：

1. 是什么
   这个技术的定义是什么？

2. 为什么需要它
   它解决系统里的什么瓶颈？

3. JD 要求到什么级别
   是会用、会 debug、会设计，还是要能写底层实现？

4. 你 CV 怎么接
   你已有经历能不能自然接上？
   哪些地方不能夸大？

5. 面试第一问
   最可能先问什么？

6. 深挖追问
   如果对方懂这个技术，会继续问什么？

7. Debug playbook
   慢了、错了、P99 高了，你怎么查？

8. 实战证据
   你跑什么实验、画什么图、看什么指标？

9. 英文回答
   60-90 秒能不能讲清楚？

你可以把每个技术点都整理成这张表：

0.4 当前文档的重新分区

为了避免你迷路，可以把全文看成 6 个区：

最推荐的实际学习顺序：

0 -> 1 -> 4 -> 9 -> 6 -> 7 -> 5 -> 8 -> 10 -> 14 -> 12 -> 15 -> 16

不是：

0 -> 1 -> 2 -> 3 -> 4 -> ... -> 17

1. JD 技能要求总表：JD 要什么、你有什么、怎么补

1.0 主索引：按 JD 技能点组织，而不是按任务组织

你现在复习时应该把每个技能点当成一条面试主线：

flowchart LR
    JD[JD 技能要求] --> CV[CV 里的对应证据]
    CV --> Match[匹配点和短板]
    Match --> Q[面试官会问什么]
    Q --> Deep[二级/三级追问]
    Deep --> Evidence[需要什么证据]
    Evidence --> Task[延展任务/实验/图]

下面这张表是全篇文档的主入口。
后面的 CUDA/vLLM/NCCL/UCX 实验都只是这张表最后一列的证据来源。

如果时间不够，只按这个优先级准备：

P0:
  AI inference/model serving
  CUDA/Nsight/GPU performance
  NCCL/collectives
  UCX/RDMA/GPUDirect

P1:
  Dynamo/NIXL
  GPUNetIO
  System architecture/prototype

P2:
  GR00T/Physical AI
  ROCm/Metal/MPS/Vulkan/WebGPU 对照

1.0.1 每个技能点应该怎么复习

每个技能点都按同一个模板准备，不要先陷入实验细节：

1. JD 原话是什么？
2. 这个技能在系统里解决什么问题？
3. 我 CV 里哪段经历可以接上？
4. 我不能夸大的边界是什么？
5. 面试官第一问是什么？
6. 面试官第二层会追到哪里？
7. 如果给一个故障场景，我怎么 debug？
8. 我有什么实验/图/benchmark 能支撑？
9. 我最后怎么用英文 60-90 秒回答？

对应到文档：

1.0.2 面试追问树：问题是主线，任务是延展

下面这些是你真正要准备的“问题链”。
每条问题链最后才对应实验、代码或图。

A. AI inference / model serving

flowchart TB
    A[LLM request path 是什么] --> B[prefill 和 decode 区别]
    B --> C[TTFT / TPOT / P99 分别代表什么]
    C --> D[TTFT 高怎么 debug]
    C --> E[TPOT 高怎么 debug]
    C --> F[P99 高怎么 debug]
    D --> G[queueing / scheduler / prefill / KV allocation]
    E --> H[decode loop / KV cache / memory bandwidth / collectives]
    F --> I[batching / admission control / tail latency]
    G --> J[延展: vLLM benchmark]
    H --> J
    I --> J

你要形成的回答能力：

我不是只知道 vLLM 这个名字。
我能把 serving 请求拆成阶段，并把每个指标映射到具体瓶颈。
如果 TTFT/TPOT/P99 异常，我知道先看什么、后看什么。

B. Dynamo / NIXL

flowchart TB
    A[Dynamo 是什么] --> B[为什么 distributed inference 需要 orchestration]
    B --> C[NIXL 是什么]
    C --> D[为什么 inference state movement 和 collective 不一样]
    D --> E[NIXL vs NCCL]
    D --> F[KV cache movement]
    F --> G[transfer slow 怎么 debug]
    G --> H[延展: KV movement data path / NIXL lifecycle]

你要形成的回答能力：

NIXL 不是 NCCL 的替代品。
NIXL 更像 inference data/state movement layer。
典型场景是 prefill/decode disaggregation 或 KV cache transfer。

C. NCCL / collectives / parallelism

flowchart TB
    A[NCCL 是什么] --> B[collective 是什么]
    B --> C[all-reduce / all-gather / reduce-scatter]
    C --> D[DP / TP / PP / FSDP 分别用什么通信]
    C --> E[small message vs large message]
    E --> F[latency-bound vs bandwidth-bound]
    D --> G[slow all-reduce 怎么定位]
    F --> G
    G --> H[rank mapping / topology / transport / NCCL_DEBUG]
    H --> I[延展: nccl-tests all_reduce_perf]

你要形成的回答能力：

我不会把所有通信慢都说成 NCCL bug。
我会先看 message size、rank、topology、transport、是否能用 nccl-tests 复现。

D. RDMA / RoCE / UCX / GPUDirect

flowchart TB
    A[传统 TCP path 为什么慢] --> B[RDMA 是什么]
    B --> C[为什么需要 registered memory]
    C --> D[QP / CQ / MR 是什么]
    B --> E[RoCE 为什么难]
    B --> F[UCX 是什么]
    F --> G[UCX transport 怎么选]
    B --> H[GPUDirect RDMA 是什么]
    H --> I[GPU-NIC direct path 需要什么条件]
    G --> J[UCX/RDMA 慢怎么 debug]
    I --> J
    J --> K[延展: ucx_info / data path 图 / slow-path checklist]

你要形成的回答能力：

我能画清楚 CPU-staged TCP path、RDMA path、GPUDirect RDMA path。
我也能诚实说明本机没有 RDMA/NIC/DOCA 时不能验证什么。

E. GPUNetIO

flowchart TB
    A[GPUNetIO 是什么] --> B[它和 GPUDirect RDMA 区别是什么]
    B --> C[为什么要 GPU-centric networking]
    C --> D[什么时候值得用]
    D --> E[什么时候不值得用]
    C --> F[代价和风险是什么]
    F --> G[debug 难点是什么]
    G --> H[延展: CPU-centric vs GPU-centric path 图]

你要形成的回答能力：

GPUNetIO 不是所有网络问题的默认答案。
只有当 CPU-mediated networking 成为关键路径，并且 GPU 直接消费网络数据时，它才可能值得考虑。

F. CUDA / Nsight / GPU performance

flowchart TB
    A[CUDA execution model] --> B[thread / block / grid / warp]
    B --> C[memory hierarchy]
    C --> D[global / shared / pinned memory]
    D --> E[coalescing / tiling / copy overlap]
    E --> F[memory-bound vs compute-bound]
    F --> G[Nsight Systems vs Nsight Compute]
    G --> H[GPU utilization 低怎么定位]
    H --> I[延展: CUDA memory / matmul / pinned copy 实验]

你要形成的回答能力：

我先用 Nsight Systems 找 critical path。
再用 Nsight Compute 看具体 kernel 的 memory、occupancy、stall、instruction 指标。

G. System architecture / prototype

flowchart TB
    A[给一个优化方向] --> B[先定义 hypothesis]
    B --> C[找 baseline]
    C --> D[做 microbenchmark]
    D --> E[做最小 prototype]
    E --> F[接入 end-to-end workload]
    F --> G[看 TTFT / TPOT / P99 / throughput / cost]
    G --> H[决定是否进 roadmap]

你要形成的回答能力：

Architect 不是只会说概念。
我会把优化假设、验证指标、风险、fallback、维护成本一起放进判断。

1.1 这个岗位到底按 C++ 还是 Python 准备？

按 JD 原文判断，不是二选一。JD 写的是：

Strong programming background in C++, Python, and ideally CUDA or other GPU programming models.

所以结论是：

语言要求：C++ 和 Python 都要。
面试主战场：C++ / CUDA / Linux systems。
实验和自动化工具：Python / Bash。
加分项：CUDA 或其他 GPU programming models。

为什么不是“只准备 Python”：

• 这是 System Software Architect, AI and GPU Networking，不是普通 Python AI application engineer。
• JD 关注 runtime systems、communication libraries、GPU networking、data movement、prototype、performance optimization。
• UCX、NCCL、GPUNetIO、CUDA、RDMA 这些底层栈主要是 C/C++/CUDA/Linux 语境。
• JD 同时要求 AI frameworks，所以 Python 也重要，但它更偏 PyTorch/vLLM/JAX/TensorFlow 集成、benchmark、实验 orchestration、数据分析和自动化。

为什么也不能说“只要 C++”：

• JD 明确写了 Python。
• AI framework 生态主要靠 Python 暴露接口。
• vLLM、PyTorch、benchmark scripts、实验数据分析会大量用 Python。
• Architect 面试会看你能不能快速 prototype，而 Python 是最快的验证工具。

你面试中要表现出的语言层次：

按面试准备优先级排序：

1. C++/CUDA/Linux systems：最高优先级，因为这是岗位底层技术可信度。
2. Python：必须会，用来跑 AI framework、vLLM benchmark、自动化实验。
3. Bash/Linux tooling：必须会，用来串起 nvidia-smi、nsys、ncu、nccl-tests、ucx_info。

准备优先级：

flowchart TB
    Role[JD Role: System Software Architect AI/GPU Networking]
    Role --> Core[C++ / CUDA / Linux Systems]
    Role --> Support[Python / Bash Tooling]
    Core --> CUDAExp[CUDA Memory / Kernel Profiling]
    Core --> Comm[NCCL / UCX / RDMA Concepts]
    Core --> Runtime[Runtime / Data Movement Design]
    Support --> VLLM[vLLM Serving Benchmark]
    Support --> Automation[Benchmark Automation]
    Support --> Analysis[Result Parsing / Plotting]

你应该这样回答语言问题：

The JD asks for both C++ and Python, ideally with CUDA or another GPU programming model. I would position myself as a systems engineer who can work in C++/CUDA/Linux for performance-critical runtime and communication paths, while using Python for AI framework integration, serving benchmarks, experiment orchestration, and analysis.

如果面试官问你更强的是 C++ 还是 Python，建议回答：

My production background includes both C++ and Python, but for this NVIDIA role I would emphasize systems-level C++ experience, performance-sensitive code paths, and my ability to use Python/Bash to build fast benchmark and diagnostic workflows. I would not present myself as only a Python ML application engineer.

你这两天的代码准备也按这个分工：

• C++/CUDA：至少准备/跑一个 CUDA memory access 或 copy benchmark。
• Python：跑或设计 vLLM benchmark，整理 TTFT/TPOT/P99。
• Bash：记录环境、运行命令、组织 labs/ 输出。

1.2 CUDA 对应的 AMD / macOS 技术怎么理解

JD 主线是 NVIDIA，所以面试准备仍然以 CUDA / NCCL / UCX / GPUDirect / GPUNetIO 为主。
但你有 Vulkan/WebGPU、跨平台 GPU/graphics 经验，可以用 AMD/macOS 技术做类比，证明你理解的是 GPU computing 的共性，而不是只背 CUDA 名词。

核心原则：

面试主语言：CUDA / NVIDIA stack。
迁移类比：AMD ROCm/HIP、Apple Metal/MPS、Vulkan/WebGPU。
不要把 AMD/macOS 讲成这个岗位主技术栈。

技术映射表：

怎么在面试里使用这段经验：

My hands-on GPU experience includes Vulkan/WebGPU-style profiling and GPU workload analysis. I understand that CUDA is the primary stack for this NVIDIA role, so I use those experiences as transferable mental models: GPU execution hierarchy, memory locality, async command execution, profiling timelines, and bottleneck isolation. I would still validate NVIDIA-specific behavior with CUDA, Nsight Systems, Nsight Compute, NCCL, and UCX.

如果面试官问“你不是 CUDA production，很大问题吗？”，可以这样答：

I would be explicit that my production GPU-adjacent work is not CUDA-kernel production ownership. The transferable part is my understanding of GPU workload behavior, profiling, memory/resource trade-offs, and systems debugging. For this role, I am closing the NVIDIA-specific gap by running CUDA/Nsight experiments and mapping my Vulkan/WebGPU mental model to CUDA execution, memory hierarchy, and profiling tools.

两天内的学习优先级：

1. 先用 CUDA/Nsight 跑实验，因为这是 NVIDIA 面试主线。
2. AMD/macOS 只作为对照理解，不花时间搭 ROCm/Metal 环境。
3. 面试里只在被问到“你过去 GPU 经验怎么相关”时讲 AMD/macOS/Vulkan/WebGPU 类比。
4. 讲完类比后立刻落回 NVIDIA：CUDA、Nsight、NCCL、UCX、GPUDirect。

2. 证据实验库：为 JD/CV 追问准备可验证材料

这一章不是新的主线目录，也不是让你机械完成一堆学生作业。
它的作用只有一个：当面试官顺着 JD/CV 技能点追问时，你能拿出“我跑过/我知道怎么跑/我知道怎么看结果”的证据。

所以这一章应该这样使用：

先回到第 1 章，确定当前技能点。
再判断面试官可能追问什么。
最后只做能支撑这个追问的实验。

阅读顺序应该是：

flowchart LR
    A[JD 技能点] --> B[CV 对应证据]
    B --> C[面试追问]
    C --> D[选择实验]
    D --> E[跑代码/看指标]
    E --> F[形成回答证据]
    F --> G[回到 Q&A]

实验优先级：

2.1 你的环境分工：Windows 原生 + WSL2 Ubuntu

flowchart TB
    Win[Windows GPU Machine]
    Native[Windows Native]
    WSL[WSL2 Ubuntu]

    Native --> Driver[NVIDIA Driver / nvidia-smi]
    Native --> CUDAWin[CUDA Toolkit for Windows]
    Native --> NsightGUI[Nsight Systems / Nsight Compute GUI]
    Native --> Samples[CUDA Samples / Visual Studio]

    WSL --> LinuxCUDA[Linux CUDA CLI]
    WSL --> Python[vLLM / Python Benchmarks]
    WSL --> NCCL[NCCL Tests if Available]
    WSL --> UCX[UCX Tools if Available]

推荐优先级：

1. 先确认 Windows 原生 nvidia-smi、nvcc、Nsight 可用。
2. 再确认 WSL2 Ubuntu 里 nvidia-smi 可用。
3. 如果 WSL2 能看到 GPU，后续 vLLM/NCCL/UCX 尽量放 WSL2。
4. 如果 WSL2 暂时不可用，不要卡太久；CUDA/Nsight 先用 Windows 原生跑，Linux 生态内容写实验设计。

2.2 Windows 原生环境检查

在 PowerShell 里跑：

nvidia-smi
nvcc --version
python --version
where nsys
where ncu

你要记录：

Windows environment:
  GPU model:
  Driver version:
  CUDA version:
  Python version:
  Nsight Systems available: yes/no
  Nsight Compute available: yes/no

怎么看结果：

• nvidia-smi 能显示 GPU：driver OK。
• nvcc --version 能显示 CUDA：CUDA Toolkit OK。
• where nsys / where ncu 能找到路径：Nsight CLI 可用。
• 如果 Nsight CLI 找不到，但 GUI 已安装，也可以先用 GUI。

2.3 WSL2 Ubuntu 环境检查

在 WSL2 Ubuntu 里跑：

nvidia-smi
python3 --version
uname -a
which nvcc || true
which nsys || true
which ncu || true

你要记录：

WSL2 environment:
  nvidia-smi works: yes/no
  Python version:
  nvcc available: yes/no
  nsys available: yes/no
  ncu available: yes/no

怎么看结果：

• WSL2 中 nvidia-smi 能看到 GPU：可以跑 Linux CUDA/Python GPU workflow。
• WSL2 中没有 nvcc 但能看到 GPU：仍可跑部分 Python GPU workload，但 CUDA 编译实验需要安装 toolkit。
• WSL2 中看不到 GPU：vLLM/NCCL/UCX 实操会受限，先做 Windows CUDA/Nsight + Linux 实验设计。

2.4 环境检查结果怎么判断

不用先创建任何文档。你现在只需要看命令输出，并判断这台机器能跑什么。

最重要的是这四个结论：

最理想结果：

Windows:
  nvidia-smi works
  nvcc works
  Nsight Systems / Compute available

WSL2:
  nvidia-smi works
  python3 works
  optional: nvcc / nsys / ncu available

如果只有 Windows CUDA 可用，也够你完成最关键的 CUDA/Nsight 主线。
如果 WSL2 也能看到 GPU，再继续做 vLLM/NCCL/UCX。

2.5 任务链 1：CUDA / Nsight，证明你能做 GPU 性能定位

这不是孤立的 CUDA 入门练习。它对应 JD 的这些考察点：

• C++ / Python, ideally CUDA or other GPU programming models
• performance profiling and optimization
• defining and using hardware features
• design and prototype features and optimizations

也对应你 CV 里的这些点：

• Vulkan/WebGPU/GPU workload profiling。
• C++/Python/Linux。
• performance-sensitive path、bottleneck isolation、diagnostic tooling。

你要用这组任务证明三件事：

1. 我知道 GPU 程序慢在哪里要分层看。
2. 我会用 Nsight Systems 看全局 timeline。
3. 我会用 Nsight Compute 看 kernel 内部瓶颈。

任务结构：

flowchart TB
    JD[JD: CUDA or GPU programming + profiling + hardware features]
    CV[CV: Vulkan/WebGPU profiling + C++/Python/Linux + bottleneck isolation]
    Goal[Goal: prove GPU performance debugging ability]

    JD --> Goal
    CV --> Goal

    Goal --> T1[Task 1: Memory Access Pattern]
    Goal --> T2[Task 2: Shared Memory Tiling]
    Goal --> T3[Task 3: Host-Device Transfer]
    Goal --> T4[Task 4: Nsight Workflow]

    T1 --> M1[Metrics: kernel time, bandwidth, load efficiency, stalls]
    T2 --> M2[Metrics: occupancy, shared memory, global load/store]
    T3 --> M3[Metrics: H2D/D2H time, overlap, sync gaps]
    T4 --> M4[Metrics: timeline, CUDA API, kernel detail]

    M1 --> Q1[追问: coalescing / memory-bound / layout]
    M2 --> Q2[追问: shared memory / occupancy / bank conflict]
    M3 --> Q3[追问: pinned memory / DMA / stream overlap]
    M4 --> Q4[追问: Systems vs Compute / critical path]

    Q1 --> Evidence[面试证据: I can locate GPU bottlenecks by measurement]
    Q2 --> Evidence
    Q3 --> Evidence
    Q4 --> Evidence

2.5.0 这组任务里的技术词汇是什么

2.5.1 任务 1：coalesced vs uncoalesced memory access

任务说明：

你要知道是什么：

• GPU global memory 访问不是“每个 thread 随便读都一样快”。
• 相邻 thread 访问连续地址时，硬件可以合并 memory transaction，带宽更高。
• stride/random/uncoalesced access 会造成更多 memory transaction，吞吐下降。

实际怎么跑：

代码文件：memory_access.cu

#include <cuda_runtime.h>
#include <cstdio>
#include <cstdlib>

#define CHECK_CUDA(call) do {                                      \
    cudaError_t err = (call);                                      \
    if (err != cudaSuccess) {                                      \
        std::fprintf(stderr, "CUDA error %s:%d: %s\n",             \
                     __FILE__, __LINE__, cudaGetErrorString(err)); \
        std::exit(1);                                              \
    }                                                             \
} while (0)

__global__ void coalesced_copy(const float* __restrict__ in,
                               float* __restrict__ out,
                               int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        out[i] = in[i];
    }
}

__global__ void strided_copy(const float* __restrict__ in,
                             float* __restrict__ out,
                             int n,
                             int stride) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        int idx = (i * stride) % n;
        out[i] = in[idx];
    }
}

float run_coalesced(const float* d_in,
                    float* d_out,
                    int n,
                    int blocks,
                    int threads,
                    int repeat) {
    cudaEvent_t start, stop;
    CHECK_CUDA(cudaEventCreate(&start));
    CHECK_CUDA(cudaEventCreate(&stop));
    CHECK_CUDA(cudaEventRecord(start));
    for (int r = 0; r < repeat; ++r) {
        coalesced_copy<<<blocks, threads>>>(d_in, d_out, n);
    }
    CHECK_CUDA(cudaEventRecord(stop));
    CHECK_CUDA(cudaEventSynchronize(stop));
    CHECK_CUDA(cudaGetLastError());
    float ms = 0.0f;
    CHECK_CUDA(cudaEventElapsedTime(&ms, start, stop));
    CHECK_CUDA(cudaEventDestroy(start));
    CHECK_CUDA(cudaEventDestroy(stop));
    return ms / repeat;
}

float run_strided(const float* d_in,
                  float* d_out,
                  int n,
                  int blocks,
                  int threads,
                  int repeat,
                  int stride) {
    cudaEvent_t start, stop;
    CHECK_CUDA(cudaEventCreate(&start));
    CHECK_CUDA(cudaEventCreate(&stop));
    CHECK_CUDA(cudaEventRecord(start));
    for (int r = 0; r < repeat; ++r) {
        strided_copy<<<blocks, threads>>>(d_in, d_out, n, stride);
    }
    CHECK_CUDA(cudaEventRecord(stop));
    CHECK_CUDA(cudaEventSynchronize(stop));
    CHECK_CUDA(cudaGetLastError());
    float ms = 0.0f;
    CHECK_CUDA(cudaEventElapsedTime(&ms, start, stop));
    CHECK_CUDA(cudaEventDestroy(start));
    CHECK_CUDA(cudaEventDestroy(stop));
    return ms / repeat;
}

int main() {
    const int n = 1 << 26;          // 67M floats, about 256 MB
    const int repeat = 50;
    const int threads = 256;
    const int blocks = (n + threads - 1) / threads;
    const int stride = 32;
    const size_t bytes = size_t(n) * sizeof(float);

    float* h = static_cast<float*>(std::malloc(bytes));
    for (int i = 0; i < n; ++i) h[i] = float(i % 1024);

    float *d_in = nullptr, *d_out = nullptr;
    CHECK_CUDA(cudaMalloc(&d_in, bytes));
    CHECK_CUDA(cudaMalloc(&d_out, bytes));
    CHECK_CUDA(cudaMemcpy(d_in, h, bytes, cudaMemcpyHostToDevice));

    float coalesced_ms = run_coalesced(d_in, d_out, n, blocks, threads, repeat);
    float strided_ms = run_strided(d_in, d_out, n, blocks, threads, repeat, stride);

    double gb = double(bytes) / 1e9;
    std::printf("coalesced: %.3f ms, %.2f GB/s\n", coalesced_ms, gb / (coalesced_ms / 1000.0));
    std::printf("strided  : %.3f ms, %.2f GB/s, stride=%d\n", strided_ms, gb / (strided_ms / 1000.0), stride);

    CHECK_CUDA(cudaFree(d_in));
    CHECK_CUDA(cudaFree(d_out));
    std::free(h);
    return 0;
}

Windows PowerShell：

nvcc -O3 memory_access.cu -o memory_access.exe
.\memory_access.exe
ncu .\memory_access.exe

WSL2 Ubuntu：

nvcc -O3 memory_access.cu -o memory_access
./memory_access
ncu ./memory_access

程序里写两个 kernel：

kernel A:
  out[i] = in[i]
  连续访问，预期 coalesced。

kernel B:
  out[i] = in[i * stride]
  或通过 index table 做 scattered access，预期 uncoalesced。

固定变量：

数据大小相同。
block size 相同。
重复次数相同。
只改变 memory access pattern。

怎么看结果：

• coalesced 版本 kernel time 更低。
• effective bandwidth 更高。
• uncoalesced 版本可能出现更多 memory dependency / memory throttle。

怎么读 Nsight Compute：

• 看 memory throughput 是否接近硬件能力。
• 看 global load/store efficiency 或 memory transaction 相关指标。
• 看 warp stall reason 是否指向 memory dependency。
• 如果 SM compute utilization 不高但 memory stall 高，说明 memory-bound。

对方可能追问：

• 什么是 memory coalescing？
• 为什么 LLM/GPU networking 也关心 memory layout？
• 如果 GPU utilization 不高但 kernel 很慢，怎么判断是不是 memory-bound？

你可以这样回答：

Memory coalescing means adjacent threads access adjacent memory so the GPU can combine memory transactions. If access is strided or scattered, effective bandwidth drops. I would validate this with Nsight Compute by checking memory throughput, global load efficiency, and warp stall reasons. In inference or networking data paths, layout still matters because KV cache blocks, tensors, and network buffers all eventually become memory access patterns.

2.5.2 任务 2：naive matmul vs tiled shared-memory matmul

任务说明：

你要知道是什么：

• naive matmul 每次从 global memory 读大量重复数据。
• tiled matmul 把一块数据放进 shared memory，让一个 block 内的 threads 复用。
• 这体现了 GPU 优化的核心思想：提高 data reuse，减少 global memory traffic。

实际怎么跑：

代码文件：matmul_tiling.cu

#include <cuda_runtime.h>
#include <cstdio>
#include <cstdlib>
#include <cmath>

#define CHECK_CUDA(call) do {                                      \
    cudaError_t err = (call);                                      \
    if (err != cudaSuccess) {                                      \
        std::fprintf(stderr, "CUDA error %s:%d: %s\n",             \
                     __FILE__, __LINE__, cudaGetErrorString(err)); \
        std::exit(1);                                              \
    }                                                             \
} while (0)

constexpr int TILE = 16;

__global__ void matmul_naive(const float* A,
                             const float* B,
                             float* C,
                             int n) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row >= n || col >= n) return;

    float sum = 0.0f;
    for (int k = 0; k < n; ++k) {
        sum += A[row * n + k] * B[k * n + col];
    }
    C[row * n + col] = sum;
}

__global__ void matmul_tiled(const float* A,
                             const float* B,
                             float* C,
                             int n) {
    __shared__ float As[TILE][TILE];
    __shared__ float Bs[TILE][TILE];

    int row = blockIdx.y * TILE + threadIdx.y;
    int col = blockIdx.x * TILE + threadIdx.x;
    float sum = 0.0f;

    for (int t = 0; t < n; t += TILE) {
        int a_col = t + threadIdx.x;
        int b_row = t + threadIdx.y;

        As[threadIdx.y][threadIdx.x] =
            (row < n && a_col < n) ? A[row * n + a_col] : 0.0f;
        Bs[threadIdx.y][threadIdx.x] =
            (b_row < n && col < n) ? B[b_row * n + col] : 0.0f;

        __syncthreads();

        for (int k = 0; k < TILE; ++k) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }

        __syncthreads();
    }

    if (row < n && col < n) {
        C[row * n + col] = sum;
    }
}

float time_naive(const float* d_A,
                 const float* d_B,
                 float* d_C,
                 int n,
                 dim3 grid,
                 dim3 block,
                 int repeat) {
    cudaEvent_t start, stop;
    CHECK_CUDA(cudaEventCreate(&start));
    CHECK_CUDA(cudaEventCreate(&stop));
    CHECK_CUDA(cudaEventRecord(start));
    for (int r = 0; r < repeat; ++r) {
        matmul_naive<<<grid, block>>>(d_A, d_B, d_C, n);
    }
    CHECK_CUDA(cudaEventRecord(stop));
    CHECK_CUDA(cudaEventSynchronize(stop));
    CHECK_CUDA(cudaGetLastError());
    float ms = 0.0f;
    CHECK_CUDA(cudaEventElapsedTime(&ms, start, stop));
    CHECK_CUDA(cudaEventDestroy(start));
    CHECK_CUDA(cudaEventDestroy(stop));
    return ms / repeat;
}

float time_tiled(const float* d_A,
                 const float* d_B,
                 float* d_C,
                 int n,
                 dim3 grid,
                 dim3 block,
                 int repeat) {
    cudaEvent_t start, stop;
    CHECK_CUDA(cudaEventCreate(&start));
    CHECK_CUDA(cudaEventCreate(&stop));
    CHECK_CUDA(cudaEventRecord(start));
    for (int r = 0; r < repeat; ++r) {
        matmul_tiled<<<grid, block>>>(d_A, d_B, d_C, n);
    }
    CHECK_CUDA(cudaEventRecord(stop));
    CHECK_CUDA(cudaEventSynchronize(stop));
    CHECK_CUDA(cudaGetLastError());
    float ms = 0.0f;
    CHECK_CUDA(cudaEventElapsedTime(&ms, start, stop));
    CHECK_CUDA(cudaEventDestroy(start));
    CHECK_CUDA(cudaEventDestroy(stop));
    return ms / repeat;
}

int main() {
    const int n = 1024;
    const int repeat = 5;
    const size_t bytes = size_t(n) * n * sizeof(float);

    float* h_A = static_cast<float*>(std::malloc(bytes));
    float* h_B = static_cast<float*>(std::malloc(bytes));
    for (int i = 0; i < n * n; ++i) {
        h_A[i] = float((i % 13) - 6) / 13.0f;
        h_B[i] = float((i % 17) - 8) / 17.0f;
    }

    float *d_A = nullptr, *d_B = nullptr, *d_C = nullptr;
    CHECK_CUDA(cudaMalloc(&d_A, bytes));
    CHECK_CUDA(cudaMalloc(&d_B, bytes));
    CHECK_CUDA(cudaMalloc(&d_C, bytes));
    CHECK_CUDA(cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice));
    CHECK_CUDA(cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice));

    dim3 block(TILE, TILE);
    dim3 grid((n + TILE - 1) / TILE, (n + TILE - 1) / TILE);

    float naive_ms = time_naive(d_A, d_B, d_C, n, grid, block, repeat);
    float tiled_ms = time_tiled(d_A, d_B, d_C, n, grid, block, repeat);

    double flops = 2.0 * double(n) * double(n) * double(n);
    std::printf("N=%d\n", n);
    std::printf("naive: %.3f ms, %.2f GFLOP/s\n",
                naive_ms, flops / (naive_ms / 1000.0) / 1e9);
    std::printf("tiled: %.3f ms, %.2f GFLOP/s\n",
                tiled_ms, flops / (tiled_ms / 1000.0) / 1e9);

    CHECK_CUDA(cudaFree(d_A));
    CHECK_CUDA(cudaFree(d_B));
    CHECK_CUDA(cudaFree(d_C));
    std::free(h_A);
    std::free(h_B);
    return 0;
}

Windows PowerShell：

nvcc -O3 matmul_tiling.cu -o matmul_tiling.exe
.\matmul_tiling.exe
ncu .\matmul_tiling.exe

WSL2 Ubuntu：

nvcc -O3 matmul_tiling.cu -o matmul_tiling
./matmul_tiling
ncu ./matmul_tiling

写两个版本：

A:
  naive C = A x B
  直接从 global memory 读。

B:
  tiled C = A x B
  每个 block 把 tile 放入 shared memory，然后复用。

矩阵规模：

512x512 / 1024x1024 / 2048x2048
按你的 GPU 能力选择，先从小规模保证正确，再扩大。

测什么：

kernel time
achieved occupancy
shared memory usage
global memory load/store throughput
SM utilization

怎么看结果：

• tiled 版本通常更快，但 tile size 不一定越大越好。
• shared memory 用太多可能降低 occupancy。
• occupancy 高不代表一定快，要结合 memory throughput 和 stall reason 看。

怎么读 Nsight Compute：

• 如果 global memory load 明显下降，说明 tiling 起作用。
• 如果 shared memory 使用增加但 occupancy 大幅下降，要解释资源 trade-off。
• 如果 stall reason 指向 memory dependency，说明仍可能 memory-bound。
• 如果 compute pipeline utilization 高，才更接近 compute-bound。

对方可能追问：

• shared memory 为什么快？
• occupancy 是什么，为什么不是越高越好？
• bank conflict 是什么？
• 你怎么判断一个 kernel 是 memory-bound 还是 compute-bound？

你可以这样回答：

Shared memory is on-chip and much lower latency than global memory, so tiling can improve data reuse. But shared memory and registers are limited resources; using too much can reduce occupancy. I would not optimize for occupancy alone. I would check whether the kernel is memory-bound