想象你是一家 AI 大饭店的厨师长，PyTorch 是后厨里的标准中央厨房：蒸箱、烤箱、炒锅一应俱全，绝大多数菜谱都能直接完成。可有一天，客人点了一道“三秒锁鲜融合菜”——需要把洗菜、切菜、下锅、调味一次完成；用普通灶台一道道传菜，等端上桌已经老了。这时，你只能自己搭一个“私房小灶间”。C++/CUDA 扩展开发，就是搭建这个私房小灶间的全套手艺：C++ 是菜谱框架，CUDA Kernel 是那块能精确控制火候的专业灶台，编译工具链是装修管道，显存管理是食材仓库，性能分析与调试则是后厨的计时器和试菜流程。读完这篇文章，你会掌握一门 PyTorch 高阶技能：把 Python 层的慢节奏，换成 GPU 上贴着硬件执行的快节奏，并让它从实验室灶台平稳走向生产连锁店。

岗位关联

掌握 C++/CUDA 扩展开发，意味着你能突破现成算子的边界，把模型训练与推理的高频瓶颈用最贴合硬件的方式消灭，这是 AI 工程师从原型走向生产的高阶通行证。

目录
  1. 可视化演示
  2. 扩展开发概述
  3. 编译工具链
  4. C++算子接口
  5. CUDA核函数
  6. 显存管理
  7. 性能陷阱
  8. 调试分析
  9. 生产落地
  10. AI 应用场景
  11. Python 演示
  12. 常见误区
  13. 面试题

可视化演示

演示Extension模块编译、CUDA Kernel编写与访存优化，理解GPU编程中自定义高性能算子的实现路径。

本节配有交互动效（进入视口后自动加载，约 3KB）。

编译流程：扩展构建 C++/CUDA 编译链接 Ext 高性能算子

扩展开发概述

中央厨房里不缺菜，但缺“一锅出”的菜。PyTorch 作为中央厨房，已经准备好了卷积、矩阵乘法、归一化等常见算子；但当你写 z = torch.relu(x * w).sum() 这样的表达式时，Python 需要多次创建中间张量、检查类型、调用底层调度器。小算子越多，传菜损耗越高：真正在下锅炒菜的时间，远小于在厨房过道里跑来跑去的时间。

C++/CUDA 扩展开发解决的就是这个痛点。它适合几类场景：多个算子反复读写中间结果、PyTorch 尚未提供的高性能算子、以及访存密集且需要手工融合的算子。开发流程也可以想象成搭小灶间：先写 C++ 接口把菜谱框架搭好，再写 CUDA Kernel 控制火候，然后用编译工具把灶台安装好，最后调试、测速、打包。

回到比喻里，原生的 PyTorch 算子像一道道需要中途传递的菜；C++/CUDA 扩展则把配菜、调味、爆炒放进同一口锅。私房小灶间不是每道菜都需要，但遇到“三秒锁鲜融合菜”时，它就是唯一能保住锅气的方案。

R=TdispatchTcompute\displaystyle R = \frac{T_{\text{dispatch}}}{T_{\text{compute}}}
这个公式表示“调度开销”与“实际计算开销”的比值。当 R 远大于 1 时，说明 Python 调度已经喧宾夺主，正适合用扩展合并算子。
  1. 分子 T_dispatchPython 解释器、张量类型检查、算子路由、中间变量创建等不产生实际数学计算的时间。
    TdispatchT_{\text{dispatch}}
  2. 分母 T_computeCUDA Kernel 或底层 BLAS 库真正完成加、乘、归一化等计算的时间。
    TcomputeT_{\text{compute}}
  3. 比值 R 远大于 1传菜耗时比做菜耗时还高，把多个步骤合并进一个自定义扩展能显著降低总延迟。
    R≫1R \gg 1

编译工具链

小灶间不是凭空出现的，它需要接通燃气、水管和排风管道。编译工具链就是扩展开发中的“装修队”。setuptools 是最常用的大门，torch.utils.cpp_extension 则是一张标好了“哪些墙能拆、哪些管线不能碰”的施工图。load_inline 甚至可以在 Python 脚本里直接传入源码并现场装修，适合快速原型验证。

cpp_extension 会自动替你找到 PyTorch 的头文件、库文件和 CUDA 安装路径；配合 Ninja 构建系统，增量编译就像装修队只补刷一间房，而不是整栋楼重装。CUDA_HOME 则是施工队寻找燃气总阀的地图：如果这个路径没配置好，编译时就会抱怨找不到 cuda_runtime.h。

装修有没有捷径？有，用 load_inline 就好比找一家“整装公司”：你把菜谱源码给它，它替你跑完预处理、编译、链接，最后把一把钥匙交给你。但在生产环境中，我们仍然要理解这条流水线，才能在报错时知道是燃气管道问题，还是墙拆错了。

源码→预处理中间码→编译目标文件→链接扩展库\displaystyle \text{源码} \xrightarrow{\text{预处理}} \text{中间码} \xrightarrow{\text{编译}} \text{目标文件} \xrightarrow{\text{链接}} \text{扩展库}
这是一条标准编译流水线：C++ 和 CUDA 源码经过预处理、编译、链接，最终变成 Python 可导入的 .so 扩展库。
  1. 预处理展开头文件、处理宏定义，让编译器看到完整的原料清单。
    预处理\text{预处理}
  2. 编译C++ 源码交给 C++ 编译器，.cu 源码交给 NVCC 编译器，分别生成目标文件。
    编译\text{编译}
  3. 链接把所有目标文件和 PyTorch 的库黏合在一起，生成最终扩展库。
    链接\text{链接}

C++算子接口

小灶间装好后，需要有服务员把客人的需求翻译成厨房能懂的指令。C++ 算子接口就是这层翻译。torch::Tensor 是 PyTorch 张量在 C++ 世界里的“餐盘”，它装满了数据、形状和梯度信息；pybind11 则是菜单，把 C++ 函数暴露给 Python 调用。

代码入口通常是一个普通 C++ 函数，直接在扩展模块中注册；如果你希望它支持自动微分，就需要继承 torch::autograd::Function，实现前向和反向逻辑。前向方法像服务员记单：把计算中需要保存的张量用 save_for_backward 放进后厨预留的抽屉；反向方法则像结账：根据顾客退回的“菜量误差”，算出该找补给每个输入的“梯度余额”。

在 C++ 接口里，前向是“做菜”，反向是“算清楚每份食材的功劳和锅”。只有把 forward 和 backward 都写对，私房小灶间才能融入 PyTorch 的自动微分流水线；否则就像菜做得再香，客人结账时账却对不上。

∂L∂x=∂L∂y⋅∂y∂x\displaystyle \frac{\partial L}{\partial x} = \frac{\partial L}{\partial y} \cdot \frac{\partial y}{\partial x}
这是反向传播的链式法则：最终损失 L 对输入 x 的梯度，等于 L 对中间输出 y 的梯度，乘以 y 对 x 的局部梯度。自定义扩展的反向函数，就是在手动计算这条链。
  1. ∂L/∂y上游传下来的梯度，即“客人对这道菜的不满程度”。
    ∂L∂y\frac{\partial L}{\partial y}
  2. ∂y/∂x自定义算子本身的前向导数，即“食材变化如何影响成品”。
    ∂y∂x\frac{\partial y}{\partial x}
  3. 两者相乘把上游误差一路传回输入，让参数更新知道该往哪个方向走。
    ∂L∂y⋅∂y∂x\frac{\partial L}{\partial y} \cdot \frac{\partial y}{\partial x}

CUDA核函数

真正点火爆炒的动作，发生在 CUDA Kernel 里。GPU 像一座拥有成百上千个灶眼的厨房，线程就是每一个灶眼上的厨师。CUDA 把线程组织成三层：网格（Grid）、线程块（Block）、线程（Thread）。你决定每个线程块有多少人、整个网格有多少个块，就像决定每个包间有几张桌、大厅一共摆多少桌。

Kernel 里的代码会同时被成千上万个线程执行。为了让每个线程知道自己负责哪一份食材，就必须计算全局唯一编号。这个编号像餐厅叫号屏幕上的号码：先看自己是第几号桌（blockIdx.x），每桌有多少人（blockDim.x），再看自己是桌上第几位（threadIdx.x）。

小灶间的火候也讲究“边界兜底”。如果数据总量不能被线程块大小整除，就会有余下的空位，Kernel 中必须有 if (i < n) 这样的检查，防止厨师去取不存在的食材。网格配置得当、边界判断扎实，CUDA 核函数才算真正能上灶。

global_id=blockIdx.x×blockDim.x+threadIdx.x\displaystyle \text{global\_id} = \text{blockIdx.x} \times \text{blockDim.x} + \text{threadIdx.x}
这是 CUDA 编程中最基础的“叫号公式”：线程的全局索引由线程块编号、块大小和块内编号共同决定。
  1. blockIdx.x当前线程所在的线程块编号，从 0 开始数。
    blockIdx.x\text{blockIdx.x}
  2. blockDim.x每个线程块包含多少个线程，相当于一桌的人数。
    blockDim.x\text{blockDim.x}
  3. threadIdx.x当前线程在块内的编号，从 0 开始。
    threadIdx.x\text{threadIdx.x}
  4. 三者相乘再相加先跳过前面所有桌的人数，再加上本桌内的位置，得到全局号。
    blockIdx.x×blockDim.x+threadIdx.x\text{blockIdx.x} \times \text{blockDim.x} + \text{threadIdx.x}

显存管理

食材要在砧板上有固定位置，显存就是 GPU 厨房里的食材货架。每个张量的数据指针 data_ptr() 就是一个货架编号，Kernel 通过这个编号直接读写显存。PyTorch 的缓存分配器像一个聪明的仓库管理员，用完的货架不会立刻退回系统，而是留在缓存池里供下一个张量复用，从而减少 cudaMalloc 这种昂贵的进货操作。

自定义扩展中一个常见动作是拿到 data_ptr<float>() 并传给 Kernel。但在访问之前，必须保证张量是连续的；contiguous() 就像把散落在各处货架上的食材重新码成一排。显存地址的计算也很朴素：只要知道起始货架编号、元素序数和每个元素占用的格子数，就能精确找到目标。

显存管理最怕“张量已经被回收，指针还在锅里炒”。PyTorch 的缓存池虽然会复用显存，但只要原始张量对象还活着，指针就是安全的；一旦它被销毁，同一块显存可能被分配给其他张量。因此自定义扩展必须明确数据所有权，必要时用 torch::from_blob 包裹外部指针时要格外小心生命周期。

address=base_ptr+offset×elem_size\displaystyle \text{address} = \text{base\_ptr} + \text{offset} \times \text{elem\_size}
这是指针寻址公式：从起始地址出发，跳过 offset 个元素，每个元素占 elem_size 字节，就得到目标地址。
  1. base_ptr张量首元素的显存地址，也就是货架第一格的位置。
    base_ptr\text{base\_ptr}
  2. offset目标元素是第几个元素，从 0 开始。
    offset\text{offset}
  3. elem_size每个元素占用的字节数，例如 float 是 4 字节。
    elem_size\text{elem\_size}
  4. 相乘再相加跳过前面所有元素，得到目标元素的实际地址。
    offset×elem_size\text{offset} \times \text{elem\_size}

性能陷阱

后厨最怕的不是做菜慢，而是所有人挤在同一个过道上。GPU 性能优化同样如此：线程虽然多，但如果访存不合并，相邻线程东一下西一下地取数据，硬件每次都要多搬几趟货。理想的模式是相邻线程访问连续地址，这就是“访存合并”，像一条流水线上的人手递手传菜。

另一个经典陷阱是共享内存的 Bank 冲突。共享内存被分成若干等宽“银行”，如果多个线程同时访问同一银行不同地址，硬件就会被迫拆分交易，拖慢速度。解决办法通常是在数组里填充一些空白单元，把访问错开，就像在挤满人的过道里加入几个隔板。除此之外，频繁同步会让所有灶眼停下等最慢的那位，而过低的 Occupancy 又会让 GPU 无菜可做。

性能调优不是把菜谱里的调料加倍，而是要看清厨房的物理限制。访存合并、Bank 冲突、同步开销、Occupancy，四个问题就像排烟、过道、灶眼和备菜速度：只有整体平衡，私房小灶间才能在高峰时段稳定出菜。

Occupancy=活跃 Warp 数每个 SM 最大可容纳 Warp 数\displaystyle \text{Occupancy} = \frac{\text{活跃 Warp 数}}{\text{每个 SM 最大可容纳 Warp 数}}
Occupancy 表示 GPU 计算单元里“上座率”有多高。上座率太低，访存延迟就无法被隐藏；但也不是越高越好，还要看每位厨师手头是否有足够独立的任务。
  1. 活跃 Warp 数当前在一个 SM 上真正被调度执行的 Warp 数量。
    Active Warps\text{Active Warps}
  2. 最大可容纳 Warp 数硬件架构决定的每个 SM 最多能同时容纳多少个 Warp。
    Max Warps per SM\text{Max Warps per SM}
  3. 两者相除得到硬件资源被利用的程度，通常在 0 到 1 之间。
    Active WarpsMax Warps per SM\frac{\text{Active Warps}}{\text{Max Warps per SM}}

调试分析

菜做得再快，如果味道不对，一切都是白费。CUDA 调试工具就是后厨里的尝菜勺。torch.cuda.synchronize() 是强制“后厨暂停，等所有灶眼熄火”的命令；没有它，Kernel 是异步返回的，时间测量就像客人还没动筷就宣布光盘。

NVIDIA Nsight 系列工具能分析每个 Kernel 的执行时间、访存带宽和占用率，相当于在每道菜出餐口装一个秒表。CUDA 调试器则允许你进入 Kernel 内部打断点，像厨师长巡视每一道工序。真正定位瓶颈时，要分清楚时间是耗在数据传输、Kernel 本身还是同步等待上。

调试分析的意义在于：小灶间不能只凭“感觉变快了”来验收。用秒表锁定最耗时的 Kernel，用同步检查排除异步陷阱，再用数据对比验证收益，私房菜才能心服口服地端上桌。

Speedup=Tbefore−TafterTbefore×100%\displaystyle \text{Speedup} = \frac{T_{\text{before}} - T_{\text{after}}}{T_{\text{before}}} \times 100\%
这是最常用的优化收益公式：优化前耗时减去优化后耗时，再除以优化前耗时，得到提升百分比。
  1. T_before优化前算子或整个模型的耗时。
    TbeforeT_{\text{before}}
  2. T_after优化后耗时。
    TafterT_{\text{after}}
  3. T_before - T_after节省下来的时间。
    Tbefore−TafterT_{\text{before}} - T_{\text{after}}
  4. 除以 T_before 并乘 100%把节省量换算成可汇报的提升比例。
    Tbefore−TafterTbefore×100%\frac{T_{\text{before}} - T_{\text{after}}}{T_{\text{before}}} \times 100\%

生产落地

小灶间在实验室里做得再好，如果不能复制到连锁店，价值就有限。C++/CUDA 扩展的生产落地，首先要考虑环境兼容。PyTorch 版本、Python 版本、CUDA 版本、编译器版本，都会影响扩展库是否能够被正常加载。不同编译器还可能产生 ABI 不兼容，就像两家连锁店用的水电标准不一样，同一套厨具无法直接装。

打包发布时，可以通过 wheel 将扩展封装成独立安装包，但要针对不同环境分别构建。CI 流水线则可以自动在多个 CUDA/Python 版本上完成编译、测试和性能回放，确保每一次修改不会让某家分店悄悄变慢。持续优化也不是一次性行为，而是每次算法迭代后都重新跑一遍性能基准。

从私房小灶间到连锁店，最重要的不是菜谱本身，而是“标准化复制”的能力。用清晰的构建脚本、自动化的 CI 矩阵和可复现的性能基准，C++/CUDA 扩展才能真正成为生产系统中的稳定资产。

发布矩阵=Python 版本×CUDA 版本×编译器 ABI\displaystyle \text{发布矩阵} = \text{Python 版本} \times \text{CUDA 版本} \times \text{编译器 ABI}
发布适配不是只测“自己这台电脑”，而是把 Python、CUDA、编译器 ABI 这三个维度组合成矩阵，逐一验证。
  1. Python 版本CPython 的 ABI 随主版本变化，扩展库需要为 3.8、3.9、3.10 等分别构建。
    Python 版本\text{Python 版本}
  2. CUDA 版本运行时库和依赖的算子库版本会影响可执行环境。
    CUDA 版本\text{CUDA 版本}
  3. 编译器 ABIGCC 版本的细微差异可能导致二进制接口不一致。
    编译器 ABI\text{编译器 ABI}
  4. 三者相乘组合成多组“客户环境”，每组都要经过打包、安装、运行验证。
    Python 版本×CUDA 版本×编译器 ABI\text{Python 版本} \times \text{CUDA 版本} \times \text{编译器 ABI}

AI 应用场景

融合 Transformer 注意力算子

阶段：模型推理优化与部署

解决的问题：标准 PyTorch 的注意力计算包含 QK^T、Softmax、PV 等多个 Kernel，需要反复写回显存；用 C++/CUDA 扩展把三步融合成一个 Kernel，可以把中间结果留在片上，大幅降低访存开销和延迟。

具体案例：某大模型推理服务中，自注意力部分每 Token 延迟约 0.8ms；通过自定义 Fused Attention Kernel 降到 0.35ms，整体推理吞吐量提升约 20%。

推荐系统 Embedding 多表聚合

阶段：特征工程与训练加速

解决的问题：推荐模型中往往有几十张 Embedding 表，循环查表再求和会引入大量 Python 调度和 GPU 小 Kernel 开销；自定义 CUDA Kernel 可以批量完成多表查询和求和。

具体案例：CTR 预估模型的 Embedding 聚合步骤占训练步长的 35%，使用扩展后聚合部分耗时下降 60%，训练步长时间缩短近 20%。

自定义信号处理算子

阶段：数据预处理与科学计算加速

解决的问题：语音、雷达等场景中的窗函数、STFT、滤波器设计常需要定制数学变换；用 C++/CUDA 扩展把 Python 循环改成 GPU Kernel，即可在保持算法灵活性的同时获得数量级加速。

具体案例：一段 10 秒音频的短时傅里叶变换原实现耗时 120ms，使用自定义扩展后降到 5ms，满足实时流式处理要求。

Python 演示

python45 行
import torchfrom torch.utils.cpp_extension import load_inline cpp_source = """#include <torch/extension.h>torch::Tensor cuda_add_one(torch::Tensor x);""" cuda_source = """#include <torch/extension.h> __global__ void add_one_kernel(const float* x, float* y, int64_t n) {  int64_t i = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;  if (i < n) {    y[i] = x[i] + 1.0f;  }} torch::Tensor cuda_add_one(torch::Tensor x) {  auto x_contig = x.contiguous();  auto y = torch::empty_like(x_contig);  int64_t n = x_contig.numel();  int64_t block = 256;  if (n > 0) {    int64_t grid = (n + block - 1) / block;    add_one_kernel<<<grid, block>>>(x_contig.data_ptr<float>(), y.data_ptr<float>(), n);  }  return y;}""" ext = load_inline(    name="add_one_ext",    cpp_sources=cpp_source,    cuda_sources=cuda_source,    functions=["cuda_add_one"],    with_cuda=True,    verbose=False,) x = torch.randn(1 << 20, device="cuda")y = ext.cuda_add_one(x)torch.cuda.synchronize()err = (y - (x + 1.0)).abs().max().item()print("max error:", err)
目的：演示如何用 load_inline 把一个 CUDA Kernel 编译成 PyTorch 扩展，并在 GPU 上执行“逐元素加一”的自定义算子。
  1. add_one_kernel<<<grid, block>>>(...)这是 CUDA Kernel 启动语法，后面三括号分别指定网格大小、线程块大小和可选参数。
  2. int64_t i = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;用全局索引公式让每个线程找到自己负责的元素。
  3. if (i < n)边界检查，防止线程越界访问显存。
  4. auto x_contig = x.contiguous();确保输入张量内存连续，Kernel 才能用线性地址安全遍历。
  5. torch.empty_like(x_contig)为输出张量分配一块连续显存。
  6. torch.cuda.synchronize()等待所有异步 Kernel 执行完毕，再检查结果。

常见误区

误区C++/CUDA 扩展不就是一个 `.cpp` 文件吗？我直接用 `g++` 编译成 `.so` 再导入不就行了？

正解CUDA 源码必须由 NVCC 编译，并且要链接 PyTorch 的头文件和库。torch.utils.cpp_extension 会帮你设置好所有编译参数；如果手动 g++，通常连 torch/extension.h 都找不到，更不用说处理 CUDA 运行时依赖了。

为何常见：初学者容易觉得编译就是“把一个文件变成另一个文件”，忽略了 CUDA 的编译器和 PyTorch 的复杂依赖体系。

误区既然 GPU 有几千个核心，那我就把 block 设得很大、grid 设得很大，让线程数远远超过数据量，速度肯定最快。

正解线程数量只是影响性能的维度之一。Kernel 还受访存合并、Bank 冲突、Occupancy、同步开销和启动开销共同影响。盲目增加线程会让边界判断变重、资源占用过高，甚至因为每个线程做的工作太少而得不偿失。

为何常见：人们习惯用“并行核心数”理解 GPU，以为堆线程就等于堆性能，却忽略了硬件资源分配和访存瓶颈。

误区我已经把 `data_ptr()` 传进 Kernel 并启动了线程，后面 Python 里的张量释放掉也没关系，反正显存还在那块地址上。

正解PyTorch 的缓存分配器会复用显存。原始张量一旦销毁，同一块地址可能被新张量占据；此时 Kernel 读到的已经不是原来的数据。自定义扩展需要保证传入张量在 Kernel 执行期间保持存活。

为何常见：因为 PyTorch 在 Python 层做了很多自动化管理，初学者容易忘记 C++ 层收到的是裸指针，而不是带引用计数的安全对象。

误区我在自己电脑上编译成功，也跑出了正确结果，那放到服务器的 Docker 里应该同样能跑。

正解扩展库对 Python 版本、PyTorch 版本、CUDA 版本、编译器 ABI 都非常敏感。服务器上任何一项不同，都可能出现 undefined symbol 或加载失败。生产环境必须用与目标一致的镜像重新编译，并通过 CI 矩阵验证。

为何常见：本地开发环境通常只有一套固定的 CUDA/Python 组合，容易让人误以为全世界只有这一种“燃气管道标准”。

面试题

Q1. 在 PyTorch 中快速编译并加载一个简单的 C++/CUDA 扩展，最常用的现成工具是什么？

单选题 · easy

  • A.torch.load 加载权重
  • B.load_inline 内联源码← 正确答案
  • C.load_state_dict 加载字典
  • D.torch.compile 编译模型

答案：B。load_inline 接收 C++/CUDA 源码并编译成扩展模块，是快速原型的首选。torch.load 加载权重，load_state_dict 恢复模型参数，torch.compile 是 PyTorch 的图编译能力，都不直接构建自定义扩展。

Q2. CUDA Kernel 中每个线程的全局唯一索引应如何计算？

单选题 · easy

  • A.blockIdx.x * blockDim.x + threadIdx.x← 正确答案
  • B.threadIdx.x * blockDim.x + blockIdx.x
  • C.blockIdx.x + blockDim.x * threadIdx.x
  • D.gridIdx.x * blockIdx.x + threadIdx.x

答案：A。全局索引应先跳过前面所有线程块中的线程：块编号乘以块大小，再加上块内线程编号。B 和 C 交换了编号顺序，D 使用了不存在的 gridIdx 维度，都会造成线程定位错误。

Q3. 自定义扩展中为什么常先调用 x.contiguous()？

单选题 · easy

  • A.把非连续张量变成连续内存布局← 正确答案
  • B.把张量复制到 CPU 内存
  • C.提高浮点数计算精度
  • D.给张量增加引用计数

答案：A。Kernel 用线性指针遍历张量时，要求内存是连续的。contiguous() 会把转置、切片等操作产生的非连续张量重新排布成连续内存；它不会把张量移到 CPU，也不改变数值精度或引用计数。

Q4. 自定义 torch.autograd.Function 的 backward 需要返回梯度对象，数量应与什么一致？

单选题 · mid

  • A.forward 中全部参数的数量
  • B.forward 中需要梯度的张量入参数量← 正确答案
  • C.forward 输出的所有元素总数
  • D.自定义类中成员变量的数量

答案：B。backward 的返回梯度数量必须与 forward 中接受梯度的张量入参一一对应。如果 forward 有两个需要梯度的输入，backward 就要返回两个梯度对象；非张量参数不需要返回梯度。

Q5. GPU 访存合并指的是哪种理想模式？

单选题 · mid

  • A.相邻线程访问连续内存地址← 正确答案
  • B.所有线程同时访问同一地址
  • C.每个线程随机访问全局内存
  • D.把全部数据放入共享内存

答案：A。当线程块内相邻的线程访问连续地址时，硬件可以把这些访问合并成少数几笔大事务，从而提升带宽利用率。所有线程访问同一地址是广播，随机访问是访存冲突，共享内存也不能替代全局内存访存优化。

Q6. load_inline 编译报找不到 cuda_runtime.h，最可能是什么原因？

单选题 · mid

  • A.CUDA_HOME 或 CUDA 路径未正确配置← 正确答案
  • B.Python 解释器版本太低
  • C.没有安装 setuptools
  • D.PyTorch 的 CPU 版本不匹配

答案：A。cuda_runtime.h 是 CUDA 工具链的头文件，编译器必须通过 CUDA_HOME 等路径找到它。Python 版本和 setuptools 主要影响 Python 扩展构建，CPU 版 PyTorch 虽然也不能编译 CUDA 算子，但直接报这个头文件缺失更常见的原因是 CUDA 路径没配好。

Q7. 减少共享内存 Bank 冲突的常用手段是什么？

单选题 · mid

  • A.给数组填充一个或多个空白单元← 正确答案
  • B.增大 blockDim.x 到 1024
  • C.将 shared memory 换成 global memory
  • D.减少线程块中的线程数量

答案：A。Bank 冲突发生在多个线程同时访问共享内存中不同地址但同一 Bank 时。在数组行之间填充空白单元，可以让相邻行错开 Bank，从而避免冲突。调整块大小并不直接解决 Bank 编号碰撞，换成全局内存通常更慢，减少线程数也只是降低并行度。

Q8. 下面哪些场景适合用 C++/CUDA 扩展来提速？

多选题 · mid

  • A.多个小算子反复读写中间结果← 正确答案
  • B.超大规模 Embedding 查表聚合← 正确答案
  • C.需要自动微分跟踪动态控制流
  • D.逐元素向量加法和 ReLU 融合← 正确答案

答案：ABD。融合多个小算子可以减少中间显存读写，Embedding 多表聚合适合定制 Kernel，逐元素加法和 ReLU 融合也能消除一次中间访存。动态控制流通常更适合在 Python 中表达，自定义扩展虽然能做但成本较高，不是首选提速场景。

Q9. CUDA Kernel 性能优化需要重点考虑哪些因素？

多选题 · hard

  • A.相邻线程访存是否合并← 正确答案
  • B.Shared memory 是否发生 bank conflict← 正确答案
  • C.Python 的 GIL 是否释放
  • D.Kernel 占用率是否过低← 正确答案

答案：ABD。访存合并、Bank 冲突和 Occupancy 直接决定 Kernel 能否高效利用 GPU 硬件。Python GIL 影响的是 Python 线程并发，而 CUDA Kernel 一旦启动就是异步执行，和 GIL 没有直接关系。

Q10. 关于 PyTorch 的 Caching Allocator，下列说法正确的是？

多选题 · extreme

  • A.通过缓存复用减少 cudaMalloc 开销← 正确答案
  • B.释放张量后显存立即归还操作系统
  • C.可通过环境变量调整分配策略← 正确答案
  • D.使用显存指针时要留意张量生命周期← 正确答案

答案：ACD。Caching Allocator 会把释放的显存留在缓存池中以便复用，因此不是立即归还操作系统。PyTorch 也提供环境变量控制策略。正因为缓存池会复用地址，裸指针的使用必须保证原张量存活，否则数据可能被覆盖。