C1 答疑：逐行读懂 c1_elementwise.cu

① TORCH_CHECK 在干什么？at::kFloat 从哪来？

你问的是 第 13–16 行：TORCH_CHECK((t).scalar_type() == at::kFloat, #t " must be float32") · drillsc1_elementwise.cu:13–16

一句话：在 kernel 启动之前，确认输入张量满足 kernel 的假设；不满足就抛异常，Python 那边看到一个写清楚原因的 RuntimeError。

这三行在文件里是一个宏的一部分，三项检查一起做：

#define CHECK_F32_CUDA(t)                                                   \
  TORCH_CHECK((t).is_cuda(), #t " must be a CUDA tensor");                  \
  TORCH_CHECK((t).is_contiguous(), #t " must be contiguous");               \
  TORCH_CHECK((t).scalar_type() == at::kFloat, #t " must be float32")

拆开这一行

• TORCH_CHECK(cond, msg...)：PyTorch 的断言宏。展开后大致是 if (!(cond)) torchCheckFail(__func__, __FILE__, __LINE__, msg)，见 pytorchException.h · TORCH_CHECK。
• (t).scalar_type()：返回张量的元素类型，类型是枚举 ScalarType，相当于 Python 的 t.dtype。定义见 pytorchTensorBase.h · scalar_type()。
• #t：预处理器的字符串化。CHECK_F32_CUDA(bias) 里的 #t 变成字面量 "bias"。
• "bias" " must be float32"：相邻的两个字符串字面量会被编译器自动拼成一个，所以最终消息是 "bias must be float32"，直接告诉你是哪个参数出了问题。

CHECK_F32_CUDA(bias);
// 预处理之后，第三行变成：
TORCH_CHECK((bias).scalar_type() == at::kFloat, "bias" " must be float32");
//                                              ^^^^^^ #t     ^^^^^^^^^^^^^^^^ 自动拼接

at::kFloat 的来历

kFloat 不是某个文件里手写的 const，而是用宏批量生成的，再经过两次 using namespace 才能写成 at::kFloat：

对应源码（点开就是那一行）：

1. 枚举本体：pytorchenum class ScalarType；Float 是第 6 号，见 pytorch_(float, Float) /* 6 */
2. 生成 k 前缀常量：pytorchDEFINE_CONSTANT → constexpr kFloat。k##name 是宏的拼接操作，k 和 Float 拼成 kFloat
3. at 导入 c10：pytorchnamespace at { using namespace c10; }
4. torch 导入 at：pytorchnamespace torch { using namespace at; }

C++ 常量和 Python dtype 的对应：

② blockDim.x 从哪来？为什么点不进去？

你问的是 kernel 里的 blockDim.x：它也是 torch 头文件带进来的吗？为什么在编辑器里 click 不到定义？· drillsc1_elementwise.cu:27

不是 torch 带来的。 blockDim、threadIdx、blockIdx、gridDim、warpSize 是 CUDA 语言自带的内置变量，由 nvcc 提供，你不需要 #include 任何东西。

细节：

• nvcc 编译 .cu 时会隐式 include cuda_runtime.h，后者又 include 了 device_launch_parameters.h，里面声明了 extern const dim3 blockDim; 等几行。CUDA Toolkit 的头文件不在 GitHub 上，装好 Toolkit 后在 $CUDA_HOME/include/device_launch_parameters.h。
• clang 也能编 CUDA，它的同等定义是开源的，读起来最直观：llvmstruct __cuda_builtin_blockDim_t，每个分量 x 就是调用 __nvvm_read_ptx_sreg_ntid_x()，也就是读 %ntid.x；变量本身在 llvmblockDim，llvmwarpSize = 32。
• 官方文档：CUDA Programming Guide · Built-in Variables；PTX 寄存器：%ntid、%tid、%ctaid。

为什么在 Mac 上点不进去

解决办法：在仓库根目录运行一次 ./setup_ide.sh（之后在 VS Code 里 Reload Window）。它会：

• 从 NVIDIA 官方的 pip wheel 里取出 CUDA 头文件，放到 .refs/cuda/；
• 生成 .clangd，让 clangd 按 CUDA 模式解析 .cu，并加上 torch、CUDA、FlashInfer 的 include 路径；
• 让 VS Code 用 clangd 代替 C/C++ 扩展的 IntelliSense。

在这台 Mac 上验证过：blockDim 跳到 clang 的 llvm__clang_cuda_builtin_vars.h（本机是 /Library/Developer/CommandLineTools/usr/lib/clang/16/include/ 下的同一个文件），TORCH_CHECK、at::kFloat、getCurrentCUDAStream、C10_CUDA_KERNEL_LAUNCH_CHECK、atomicAdd 都能跳到 torch / CUDA 头文件里的定义。在 Linux GPU 机器上用 Remote-SSH 打开项目也可以，那里有完整的 CUDA toolkit。

③ at 是什么？at / c10 / torch 什么关系？

你问的是 at::kFloat、at::cuda::getCurrentCUDAStream() 里的 at · drillsc1_elementwise.cu:44

at 是 ATen（“A Tensor Library”）的命名空间：PyTorch 底层的 C++ 张量库。PyTorch 的 C++ 部分分三层：

源码：pytorchat 导入 c10、at::cuda 导入 c10::cuda · pytorchtorch 导入 at（附长注释解释为什么） · 你 include 的 pytorchtorch/extension.h 把这一切都带进来。

④ vector_add：threads = 256 和 ceil 除法

你问的是 第 41–42 行：threads = 256 和 blocks = (n + threads - 1) / threads · drillsc1_elementwise.cu:33–49

先看整个 host 函数做了什么。它跑在 CPU 上，自己不做加法，只负责把 GPU 上的活安排好：

核心想法：一个线程负责一个元素

CUDA 不让你直接说“给我 n 个线程”，只能说“给我 blocks 个 block，每个 block threads 个线程”。所以要把 n 个元素按 256 个一组切开，算出要几组。

threads = 256：每个 block 的线程数。

• 必须是 32 的倍数。 GPU 按 warp（32 个线程）调度。写 100 的话，硬件照样分配 4 个 warp = 128 个线程的位置，最后一个 warp 只有 4 个线程在干活。
• 上限 1024。 超过会让启动失败（下面的 LAUNCH_CHECK 就是为了抓这种错误）。
• 为什么是 256： 够大，一个 block 有 8 个 warp，SM 可以在它们之间切换来掩盖访存延迟；又不太大，一个 SM 可以同时驻留好几个 block。对逐元素 kernel，128 到 512 之间差别通常很小。

blocks = (n + threads - 1) / threads：向上取整的整数除法。

整数除法 n / threads 是向下取整，会漏掉尾巴。先加上 threads - 1 再除，就变成 ⌈n / threads⌉。

n = 1000 时，4 个 block 怎么覆盖这 1000 个元素：

block 3 的最后 24 个线程（i = 1000 … 1023）没有对应的元素，被 if (i < n) 挡掉，什么都不做。

所以 blocks 的计算和 kernel 里的边界检查是配套的：前者宁可多发线程也不能漏元素，后者把多发的线程挡掉。

// kernel 里（第 27–29 行）
int64_t i = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;   // 全局编号
if (i < n) c[i] = a[i] + b[i];                                // 多出来的线程直接跳过
// 例：block 3 的第 231 号线程：i = 3 * 256 + 231 = 999 < 1000，干活
//     block 3 的第 232 号线程：i = 3 * 256 + 232 = 1000，不满足，跳过

为什么到处都是 int64_t？ n 可能超过 2³¹（约 21 亿）。blockIdx.x * blockDim.x 两个都是 32 位无符号数，相乘会在 2³² 处回绕，所以第 27 行先把 blockIdx.x 转成 int64_t 再乘。

剩下几行

• if (n == 0) return c;：n = 0 时 blocks 会算出 0，而用 0 个 block 启动 kernel 是 CUDA 错误。
• at::cuda::getCurrentCUDAStream()：拿 PyTorch 当前的 stream，让这个 kernel 和前后的 torch 操作按顺序执行。返回的 CUDAStream 能隐式转换成 cudaStream_t：pytorchoperator cudaStream_t()，函数声明见 pytorchgetCurrentCUDAStream。
• <<<blocks, threads, 0, stream>>>：四个参数依次是 grid 大小、block 大小、动态 shared memory 字节数（这里不用）、stream。见 Programming Guide · Launching Kernels。
• C10_CUDA_KERNEL_LAUNCH_CHECK()：<<<>>> 本身不返回错误。这个宏立刻调用 cudaGetLastError()，把非法的启动配置（例如 threads > 1024）当场报出来：pytorchC10_CUDA_KERNEL_LAUNCH_CHECK。它抓不到 kernel 运行中的错误（比如越界访问），那些要等之后同步时才会暴露。

⑤ __device__ __forceinline__ 是什么意思？

你问的是 第 54 行：__device__ __forceinline__ float gelu_tanh(float x) · drillsc1_elementwise.cu:54–58

两个修饰词各管一件事：__device__ 说“这个函数在 GPU 上运行，给 kernel 调用”；__forceinline__ 说“别生成函数调用，把函数体直接贴到调用处”。

__device__：执行空间

如果把 __device__ 删掉，gelu_tanh 就成了普通 CPU 函数，kernel 里调用它会直接编译报错。

__forceinline__：强制内联

内联之后，第 71 行在编译器眼里相当于：

float v = x[i] + bias[i % cols];
y[i] = 0.5f * v * (1.0f + tanhf(0.7978845608028654f * (v + 0.044715f * v * v * v)));

为什么在 GPU 上这很重要：

• 省掉调用开销。 GPU 上真正的函数调用要按调用约定传参、保存寄存器，可能还要用栈，而栈放在 local memory 里，很慢。对只有几条算术指令的函数，这些开销可能比函数本身还大。
• 让编译器继续优化。 展开以后，编译器能把它和周围代码一起分配寄存器、重排指令、合并成 FMA（乘加）。

nvcc 默认就会积极内联小的 __device__ 函数，__forceinline__ 是把“大概率内联”变成“一定内联”，也写明了作者的意图。反过来的 __noinline__ 用于很大的函数，避免代码膨胀。文档：Execution Space Specifiers · Inline Specifiers。

对照 PyTorch 自己的实现：F.gelu(approximate="tanh") 的 CUDA kernel 用的是同一个公式，kBeta 就是这里的 k0 = sqrt(2/π)，kKappa 就是 k1 = 0.044715：pytorchActivationGeluKernel.cu · tanh GELU。

⑥ grid-stride loop：一个线程处理好几个元素

你问的是 第 68–72 行的 for 循环 · drillsc1_elementwise.cu:68–72

const int64_t stride = (int64_t)blockDim.x * gridDim.x;      // 整个 grid 的线程总数
for (int64_t i = (int64_t)blockIdx.x * blockDim.x + threadIdx.x; i < total; i += stride) {
  y[i] = gelu_tanh(x[i] + bias[i % cols]);
}

和 vector_add 的区别：vector_add 启动“刚好够”的线程，每个线程只做一次；这里 block 数有上限（见第 ⑦ 节），线程可能比元素少，所以每个线程循环：做完一个元素，往后跳 stride 个再做一个。

• stride = blockDim.x * gridDim.x = 每个 block 的线程数 × block 数 = 整个 grid 的线程总数
• 起点 i 和 vector_add 一样，是这个线程的全局编号
• i < total 同时是边界检查

小数字走一遍：2 个 block × 4 个线程，total = 20

stride = 4 × 2 = 8。颜色表示 block，格子里写的是元素下标和处理它的线程（全局编号 g）：

block 0（g0–g3）block 1（g4–g7）i ≥ total，循环结束

• 线程 g0 处理 0 → 8 → 16，下一个是 24，退出
• 线程 g5 处理 5 → 13，下一个是 21，不满足 21 < 20，退出

整个 grid 像一把梳子，每轮整体往前平移 stride 个元素，直到扫完数组。

为什么跨步，而不是每个线程分一段连续的？

另一种直觉是“线程 0 做 0–2，线程 1 做 3–5……”。这样性能很差，原因是合并访存（coalescing）。

关键点：grid-stride 是在同一个线程的不同轮次之间跨步；同一轮里，相邻线程读的永远是相邻地址。延伸阅读：NVIDIA blog · Grid-Stride Loops · Programming Guide · Coalesced Global Memory Access。

⑦ block 数上限 sms × 32 在哪里？为什么要限？

你问的是“为什么要限制 block 数（sms × 32），我没看到”：它在 host 函数 bias_gelu 里，第 84–88 行 · drillsc1_elementwise.cu:84–88

const int threads = 256;
const int sms = at::cuda::getCurrentDeviceProperties()->multiProcessorCount;
const int64_t blocks = std::min<int64_t>((total + threads - 1) / threads, (int64_t)sms * 32);

kernel 里看不到，是因为 kernel 不决定启动多少个 block。block 数是 host 端在 <<<blocks, threads>>> 里传进去的，kernel 只能通过 gridDim.x 读到它。

• getCurrentDeviceProperties() 返回当前 GPU 的 cudaDeviceProp，multiProcessorCount 是 SM（流式多处理器）个数。源码：pytorchgetCurrentDeviceProperties。
• 常见 GPU 的 SM 数：A100 是 108，H100 SXM 是 132，RTX 5090（本仓库实测用的卡）是 170。

用 RTX 5090 算几个例子

sms × 32 = 170 × 32 = 5440 个 block，stride = 5440 × 256 = 1,392,640。

（轮数 = ⌈total / stride⌉，例如 10⁷ / 1,392,640 ≈ 7.2，所以最多 8 轮。）

为什么要封顶

• 数据少时没区别。 min 取的是 ceil(total/256)，每个线程只循环一次。
• 数据多时少调度开销。 与其发几百万个只干一点活就退出的 block，不如让数量合适的 block 常驻在 SM 上循环干活。每个 SM 32 个 block，足够让所有 SM 都忙起来，并有足够多的 warp 轮换来掩盖访存延迟。
• 一个 kernel 适配任意大小。 不管 total 多大，grid 大小都是有限的。

⑧ bias[i % cols]：扁平下标怎么找到列号

你问的是 第 70–71 行：y[i] = gelu_tanh(x[i] + bias[i % cols]); · drillsc1_elementwise.cu:70–71

这一行等价于 PyTorch 的：

y = F.gelu(x + bias, approximate="tanh")   # x: [rows, cols], bias: [cols]

关键是 x + bias 的广播：bias 只有一行，要加到 x 的每一行上。可 kernel 里没有“行”和“列”，只有一个扁平下标 i，所以要用 i % cols 算出当前元素在哪一列。

第一步：二维张量在显存里是一条一维数组

设 rows = 3，cols = 4。颜色表示列：

逻辑上的二维 x（3 行 × 4 列）和 bias（4 个数）：

显存里实际的样子（行优先：先存完第 0 行，再存第 1 行……）：

x[i] 和 bias[i % 4] 的颜色每一格都对得上：每到新的一行，i % cols 就回到 0，bias 从头再读一遍。这就是广播。

第二步：从 i 反推行号和列号

每行有 cols 个元素，所以行号 = i / cols（前面有几个完整的行），列号 = i % cols（在当前行里排第几）。

只需要列号，所以代码里没有算行号。

第三步：这一行按顺序做了什么，以及“融合”省在哪

1. 读 x[i]（显存）
2. 读 bias[i % cols]（几乎总是命中缓存，见下）
3. 相加，结果在寄存器里
4. gelu_tanh，结果还在寄存器里
5. 写 y[i]（显存）

为什么注释说 bias “stays in L1/L2”： x 的每个元素只读一次，实打实地消耗显存带宽；bias 很小（cols = 4096 时只有 16 KB），却被每一行重复读一遍。第一次读之后它就留在缓存里（L2 有几十 MB），之后几乎全部命中，不占显存带宽。所以这个 kernel 的显存流量就是读 x 一次加写 y 一次。

一个小代价： % 是 64 位整数取模，GPU 上要用几十条指令模拟。这个 kernel 受访存限制，取模的时间通常被访存延迟掩盖掉。想进一步优化，可以改成二维 grid，让 block 直接按行列索引，就不需要取模了。

前提：x 必须是连续的。 “列号 = i % cols”只在 x 真的是紧密的行优先布局时成立。传一个 x.t() 进来，逻辑形状没变，内存顺序却变了，算出的列号就错了，而且不会报错。这正是第 ① 节里 is_contiguous() 检查存在的原因。

源码与资料索引

本仓库：drillssolutions/c1_elementwise.cu · drillsmine/c1_elementwise.cu（你写的） · drillscontent/c1.md（题目页）

PyTorch（钉在固定 commit）

• 检查与报错：pytorchTORCH_CHECK · pytorchC10_CUDA_CHECK · pytorchC10_CUDA_KERNEL_LAUNCH_CHECK
• dtype：pytorchenum ScalarType · pytorchkFloat 等常量
• 命名空间：pytorchat ← c10 · pytorchtorch ← at
• Tensor 方法：pytorchis_cuda · pytorchis_contiguous · pytorchnumel · pytorchscalar_type · pytorchdata_ptr<T>
• CUDA 上下文：pytorchgetCurrentCUDAStream · pytorchgetCurrentDeviceProperties
• 对照实现：pytorchPyTorch 的 tanh-GELU kernel

CUDA 内置变量（clang 的开源实现）：llvmblockDim → %ntid · llvmblockDim 变量

文档

• Programming Guide · Built-in Variables · Execution Space Specifiers · Inline Specifiers · __restrict__
• Programming Guide · Launching Kernels · Error Checking · Coalesced Access
• PTX · special registers
• NVIDIA blog · Grid-Stride Loops · How to Access Global Memory Efficiently
• PyTorch · Custom C++ and CUDA Extensions · PyTorch C++ API