CUDA编程实战02：认识 GPU、SM、线程束与计算能力

CUDA编程实战02：认识 GPU、SM、线程束与计算能力

系列定位：从零开始，用能够运行、能够观察、能够验证的程序学习 CUDA。

本篇适合：已经完成第一篇环境搭建，能运行一个简单 Kernel，但仍然分不清 GPU、SM、Block、Thread、Warp 和计算能力的读者。

阅读后的实际成果：你将得到一份自己电脑的 GPU 能力报告，能够算出线程的全局编号与 Warp 编号，能够亲眼看到 32 位活动掩码，并能判断常见 CUDA 配置是否合理。

配套环境：Windows 或 Linux、NVIDIA GPU、CUDA Toolkit、CMake；文中代码使用 CUDA C++。

本文资料核对日期：2026 年 7 月。CUDA 会持续演进，具体上限应以程序查询结果与当前 NVIDIA 官方文档为准。

上一篇：CUDA编程实战01：开发环境搭建与第一个 GPU 程序

一、这篇文章要解决的，不是“背几个硬件名词”

第一篇文章完成后，你应该已经见过这样的 Kernel 启动：

hello_from_gpu<<<2, 4>>>();

你知道它代表 2 个 Block，每个 Block 有 4 个 Thread，总共会创建 8 个 CUDA 线程。但是，从这一行继续往下追问，很快就会遇到一连串问题：

• 这些线程是不是一个线程对应一颗“CUDA Core”？
• Block 是不是一块固定硬件？
• SM 到底是什么，为什么性能文章总在谈 SM？
• 为什么很多示例喜欢用 128、256 个线程，而不是 100、200 个？
• Warp 为什么总是 32 个线程？
• threadIdx.x、blockIdx.x 如何组合成全局编号？
• 同一个 Block 的线程是否一定在一起执行？
• 不同 Block 谁先执行？
• 显卡型号、CUDA Core 数量、计算能力、CUDA Toolkit 版本是什么关系？
• sm_75 是显卡型号，还是编译器的优化等级？
• 为什么程序逻辑完全正确，换一个 Block 大小性能却可能不同？

如果这些问题没有理清，后面学习向量加法、共享内存、归约、矩阵乘法时，很容易陷入一种状态：代码能够照着写，数字也能照着改，但不知道每一个数字为什么存在，更不知道错误发生在哪一层。

因此，本篇不会把 GPU 硬件知识写成产品参数百科。我们只保留真正会影响 CUDA 编程判断的部分，并用三个程序完成验证：

读完后，你不应该只会复述“Warp 有 32 个线程”，而应该能回答：

当一个 Block 有 40 个线程时，硬件会形成几个 Warp？第二个 Warp 有多少有效 Lane？下一个 Block 的第 0 个线程属于前一个 Block 的第二个 Warp 吗？

答案会在程序输出中出现，而不是只写在结论里。

二、先建立一张不会混乱的地图

CUDA 初学者最常见的困难，是把“软件抽象”和“硬件实体”混在了一起。先看两个层次。

2.1 软件层：你在代码里创建什么

当 CPU 启动一个 CUDA Kernel 时，程序员描述的是：

Grid
└── Block
└── Thread

• Grid 表示本次 Kernel 启动产生的全部线程；
• Grid 由多个 Block 组成；
• 每个 Block 由多个 Thread 组成；
• Grid、Block 都可以是一维、二维或三维；
• Thread 是程序员视角下最基本的逻辑执行者。

这些概念直接出现在 CUDA 语法中：

kernel<<<grid_size, block_size>>>();

以及 Kernel 内的内建变量中：

gridDim
blockDim
blockIdx
threadIdx

2.2 硬件层：GPU 用什么执行

为了建立第一版正确心智模型，可以把硬件层简化为：

GPU
├── 多个 SM
│ ├── Warp 调度相关资源
│ ├── 运算功能单元
│ ├── 寄存器文件
│ ├── Shared Memory / L1 相关片上资源
│ └── 可同时驻留的若干 Block 和 Warp
├── L2 Cache
└── Global Memory

SM 的英文是 Streaming Multiprocessor，通常译为流式多处理器。CUDA 中大量线程真正被组织和推进执行的核心场所，就是 SM。

这里暂时不要把 SM 简单等同为“CPU 核心”，因为两者的设计目标和执行方式不同。CPU 核心擅长少量复杂控制流、低延迟和强单线程能力；GPU 的 SM 擅长让大量线程保持在执行现场，并在某些 Warp 等待数据时快速切换到其他可运行 Warp，以提高整体吞吐量。

2.3 软件层和硬件层如何连接

最关键的映射规则是：

• 一个 Grid 可以包含远多于 SM 数量的 Block；
• 每个 Block 在执行期间归属于一个 SM，不会拆成两半跨越两个 SM；
• 一个 SM 在资源允许时可以同时驻留多个 Block；
• Block 被分配到哪个 SM、谁先执行，通常不由程序员指定；
• Block 内的线程会被划分为 Warp；
• Warp 是硬件调度和执行时必须理解的重要分组。

在这里插入图片描述

这张图必须读出三个结论。

第一，Block 是软件组织单位，不是芯片上焊死的一块区域。Block A 这次可能在 SM 0，下一次运行不应该假定还在 SM 0。

第二，一个 SM 不一定一次只执行一个 Block。只要寄存器、Shared Memory、线程数、Warp 数等资源允许，多个 Block 可以同时驻留在同一个 SM。

第三，不同 Block 不能依赖普通 Kernel 中的执行先后顺序。你不能写出“Block 1 等待 Block 0 设置一个标志”的天真实现，然后期待所有显卡都正确运行。Block 0 可能还没有得到执行机会，而 Block 1 却一直占着资源等待它，最终产生错误或死锁风险。

NVIDIA 当前的 CUDA Programming Guide 明确描述了 Grid、Block 与 SM 的关系，也强调 Block 之间不应依赖调度顺序。学习时可以把官方 Programming Model 作为长期参考。

三、GPU 与 CPU 的差异：重点不是“谁更快”

初学者常问：“GPU 比 CPU 快多少倍？”这个问题没有脱离任务类型的统一答案。

3.1 CPU 更像少量能力全面的负责人

CPU 的典型优势包括：

• 复杂分支与控制逻辑；
• 操作系统、文件、网络、用户交互；
• 低延迟响应；
• 强单线程性能；
• 大缓存与复杂预测机制；
• 处理任务数量不大、相互依赖明显的工作。

3.2 GPU 更像能够容纳大量在途任务的并行工厂

GPU 的典型优势包括：

• 大量结构相似、彼此相对独立的运算；
• 高吞吐量数值计算；
• 大规模向量、矩阵、图像、科学计算；
• 用大量线程隐藏部分内存访问等待；
• 对规则数据并行任务进行批量处理。

这张概念图不是 GPU 物理版图。它表达的是：

• SM 内可以驻留多个待执行 Warp；
• Warp 调度器会从可运行的 Warp 中选择工作推进；
• 寄存器和 Shared Memory 等片上资源离执行单元更近；
• Global Memory 容量大，但访问代价与片上资源不同；
• 当某个 Warp 因数据依赖等原因暂时不能推进时，SM 可以让其他就绪 Warp 工作。

因此，GPU 的强项不是“单个线程特别快”，而是“同时管理并推进大量线程”。一个 CUDA 线程很轻量，创建上百万个逻辑线程也很常见。它与操作系统线程不是同一个成本模型。

3.3 不要把“CUDA 线程”理解成“一个固定物理核心”

这是整篇文章最需要消除的误区：

CUDA Thread 是逻辑执行实例，不是一颗被长期独占的 CUDA Core。

Grid 中可能有几百万个 Thread，而 GPU 不可能有几百万个独立物理计算核心。硬件会分批调度 Block 和 Warp，复用 SM 内的执行资源。

类似地，Warp 的 32 个 Lane 也不能简单画成 32 颗固定物理核心并建立永久一一对应。不同架构的功能单元数量、指令吞吐和调度方式会变化。CUDA 编程模型保证的是可观察的软件语义，而不是要求每一代 GPU 使用完全相同的微架构实现。

这就是为什么本文不会根据计算能力和 SM 数量“猜 CUDA Core 总数”。产品规格表可以提供对应型号的硬件信息，但 cudaDeviceProp 没有一个跨架构稳定的“CUDA Core 数量”字段。对于实际编程，SM 数量、资源限制、内存特征和经过测量的 Kernel 表现通常更有意义。

四、SM 到底是什么

4.1 一句话定义

SM 是 GPU 中负责驻留、调度并执行大量线程的核心多线程处理单元。

注意这句话中的三个动词：

• 驻留：线程状态和所需资源需要在 SM 上占据位置；
• 调度：并不是所有 Warp 每个时刻都能发射指令；
• 执行：SM 中不同类型的功能单元完成整数、浮点、加载存储等工作。

4.2 为什么一个 SM 能同时保留很多线程

GPU 依靠大量在途线程提高吞吐。当一个 Warp 等待 Global Memory 数据或等待前一条指令结果时，调度器可以选择另一个已经准备好的 Warp。

这不是说切换完全没有任何硬件成本，也不是说等待总能被完全隐藏，而是说明 GPU 的设计思路与依靠少量重量级线程的直觉不同。

想象一个车间：

• 每个订单是一个 Block；
• 每个订单被拆成若干组，每组是一个 Warp；
• 每组里有 32 个 Lane；
• 车间拥有有限的工作台、工具柜和临时存放空间；
• 一个订单需要的工具和空间越多，车间能同时接纳的订单就越少。

映射到 CUDA：

• 工作台与工具资源可以类比寄存器、Shared Memory 和功能单元；
• Block 线程数越多，占用的 Warp 名额越多；
• 每个线程使用的寄存器越多，整个 Block 需要的寄存器越多；
• 每个 Block 申请的 Shared Memory 越多，SM 能同时容纳的 Block 可能越少。

4.3 “驻留”不等于“同一时刻每个线程都在算”

一个 Warp 驻留在 SM 上，表示它的执行上下文和资源已经就位。它可能处于：

• 可以发射下一条指令；
• 等待数据；
• 等待同步；
• 等待某项依赖；
• 暂时没有被调度器选中。

因此，“SM 上驻留了很多 Warp”与“所有 Warp 每个时钟周期都同时执行”不是同一个意思。

4.4 一个 SM 能放多少 Block

没有只看 Block 数量就能得到的固定答案。至少会受到以下因素限制：

• 每个 Block 的线程数量；
• 每个 Block 形成的 Warp 数量；
• Kernel 每个线程使用的寄存器数量；
• 每个 Block 使用的静态与动态 Shared Memory；
• 架构允许的每 SM 最大线程数、Warp 数和 Block 数；
• 某些架构特定限制。

可以把“同时能驻留多少 Block”理解为多个资源约束共同求最小值。这里只建立形式，不代入具体架构数字：

B

r

=

min

⁡

(

B

t

,

B

w

,

B

g

,

B

s

,

B

h

)

B_r = \\min(B_t, B_w, B_g, B_s, B_h)

Br​=min(Bt​,Bw​,Bg​,Bs​,Bh​)

其中：

• $B_r$
• $B_t$
• $B_w$
• $B_g$
• $B_s$
• $B_h$

这个公式不是让你现在手算 Occupancy，而是提醒你：只看 Block 大小，无法完整推断驻留情况。

后续做性能优化时，可以使用 NVIDIA Nsight Compute、Occupancy Calculator 或 Runtime Occupancy API。现在先知道性能不是由一个数字决定的。

4.5 Block 为什么不能跨 SM

同一个 Block 的线程可以：

• 访问该 Block 的 Shared Memory；
• 使用 __syncthreads() 进行 Block 级同步；
• 以较低代价交换某些局部数据。

如果 Block 跨越多个 SM，实现这些语义会复杂得多。CUDA 的基本执行模型让一个 Block 完整驻留在一个 SM，进而建立清晰的共享和同步边界。

这也解释了为什么 Block 线程数存在上限：所有线程需要在同一个 SM 上获得资源。程序一会儿会查询你设备的 maxThreadsPerBlock，不要把网上某个数字永远写成所有 GPU 的绝对真理。

五、Grid、Block、Thread：从坐标到唯一编号

5.1 四个必须记住的内建变量

它们都是 dim3 风格的三维值，可以访问 .x、.y、.z。

5.2 一维全局编号

假设：

• 当前 Block 的一维编号是 blockIdx.x；
• 每个 Block 有 blockDim.x 个线程；
• 当前线程在 Block 中的编号是 threadIdx.x。

最常见的一维全局索引是：

int i = blockIdx.x * blockDim.x + threadIdx.x;

为了让 Typora 稳定预览，本文公式都使用单行公式体，并把 CUDA 标识符放在公式后的文字中：

i

=

b

x

B

x

+

t

x

i = b_x B_x + t_x

i=bx​Bx​+tx​

其中，

b

x

b_x

bx​ 对应 blockIdx.x，

B

x

B_x

Bx​ 对应 blockDim.x，

t

x

t_x

tx​ 对应 threadIdx.x。

例如 Grid 有 3 个 Block，每个 Block 有 4 个 Thread：

你会发现，每进入下一个 Block，全局编号就向后移动一个 blockDim.x。

5.3 为什么几乎总要做边界判断

假设需要处理

N

=

1000

N=1000

N=1000 个元素，希望每个 Block 使用 256 个线程。需要的 Block 数可以用整数运算写成：

int blocks = (N + threads – 1) / threads;

这里会得到 4 个 Block，总共创建 1024 个线程。最后 24 个线程没有对应数据，因此 Kernel 内必须写：

int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) {
// 只处理有效元素
}

对应关系是：

N

l

=

G

x

B

x

N_l = G_x B_x

Nl​=Gx​Bx​

其中

N

l

N_l

Nl​ 是创建的逻辑线程总数，

G

x

G_x

Gx​ 是一维 Block 数，

B

x

B_x

Bx​ 是每个 Block 的线程数。通常

N

l

N_l

Nl​ 可以大于真实数据量

N

N

N，所以要判断

i

<

N

i<N

i<N。

5.4 二维坐标

图像和矩阵天然是二维数据，常见写法是：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;

数学形式：

x

=

b

x

B

x

+

t

x

x = b_x B_x + t_x

x=bx​Bx​+tx​

y

=

b

y

B

y

+

t

y

y = b_y B_y + t_y

y=by​By​+ty​

如果二维数组按行连续存储，宽度为 width，线性下标通常是：

int index = y * width + x;

对应：

i

=

y

W

+

x

i = yW + x

i=yW+x

其中

W

W

W 是每行元素数量。

5.5 三维 Block 内线程如何线性化

Warp 的划分需要把三维 threadIdx 线性化。CUDA 中 x 维变化最快，然后是 y，最后是 z。Block 内线性线程编号可以写成：

t

l

=

t

x

+

B

x

t

y

+

B

x

B

y

t

z

t_l = t_x + B_x t_y + B_x B_y t_z

tl​=tx​+Bx​ty​+Bx​By​tz​

其中

t

l

t_l

tl​ 是 Block 内线性编号。

这条规则非常重要。Warp 不是“先按 y 方向凑 32 个线程”，而是根据线性线程编号连续分组。设计二维 Block 时，例如 dim3 block(16, 16)：

• 总线程数是 256；
• 线性编号 0～31 组成第一个 Warp；
• 因为 x 维长度为 16，第一个 Warp 覆盖 y=0 的 16 个线程和 y=1 的 16 个线程。

所以，二维形状不仅影响代码可读性，也会影响连续 Lane 映射到数据的方式，进而与内存访问合并等性能问题相关。

六、Warp：为什么总有人建议 Block 大小是 32 的倍数

6.1 Warp 是什么

在当前 CUDA 编程模型中，同一 Block 的线程会被组织成每组 32 个线程的 Warp。Warp 中线程拥有各自的寄存器状态和线程编号，但以 SIMT 方式推进同一份 Kernel 代码。

你可以把 Warp 理解为：

硬件调度时一起考虑的一组 32 个 Lane。

“Lane”表示 Warp 内的位置，通常是 0～31。

6.2 Warp 编号从每个 Block 重新开始

一维 Block 内：

int warp_in_block = threadIdx.x / warpSize;
int lane = threadIdx.x % warpSize;

当 blockDim.x == 40 时：

• threadIdx.x 为 0～31 的线程属于 Warp 0；
• threadIdx.x 为 32～39 的线程属于 Warp 1；
• Warp 1 只有 8 个 Lane 承载这个 Block 的线程。

下一个 Block 的 threadIdx.x == 0 会重新成为该 Block 的 Warp 0、Lane 0。不能拿全局线程编号直接除以 32，认为 Warp 会跨 Block 延续。

6.3 一个 Block 有多少个 Warp

设每个 Block 的线程数为

T

b

T_b

Tb​，Warp 大小为

W

s

W_s

Ws​，则 Warp 数量需要向上取整：

W

b

=

⌈

T

b

/

W

s

⌉

W_b = \\lceil T_b / W_s \\rceil

Wb​=⌈Tb​/Ws​⌉

当前设备通常查询到

W

s

=

32

W_s=32

Ws​=32。

对于 40 个线程：

W

b

=

⌈

40

/

32

⌉

=

2

W_b = \\lceil 40 / 32 \\rceil = 2

Wb​=⌈40/32⌉=2

图中第一行是完整的 Warp 0，第二行是 Warp 1。Warp 1 的前 8 个 Lane 对应线程 32～39，剩余 24 个位置没有承载这个 Block 的线程。

6.4 不足 32 个线程，程序会不会错

不会因为“不足 32”而自动出错。以下配置在语法上都可能合法：

kernel<<<1, 1>>>();
kernel<<<1, 17>>>();
kernel<<<1, 40>>>();
kernel<<<1, 100>>>();

但最后一个 Warp 可能不完整。未承载有效线程的 Lane 无法为你的任务贡献工作，因此可能降低资源利用率。

这就是“Block 线程数通常选择 32 的倍数”的来源。但它不是“任何 Kernel 都必须使用 256”的死规则。Block 大小还要考虑：

• 数据形状；
• 寄存器使用；
• Shared Memory；
• 占用率；
• 内存访问模式；
• 指令类型；
• 实际测量结果。

对于刚开始的一维逐元素 Kernel，128 或 256 经常是合理起点；最终选择需要测量，而不是把经验值当定律。

6.5 为什么本文代码使用 warpSize

Kernel 中存在内建常量 warpSize。示例写：

int warp_in_block = threadIdx.x / warpSize;
int lane = threadIdx.x % warpSize;

而不是把 32 到处散落在逻辑中。这能清楚表达“计算依赖 Warp 大小”。主机端也可以通过设备属性查询 prop.warpSize。

NVIDIA 官方文档当前明确说明 Warp 为 32 个线程，但工程代码仍应尽量让硬件属性的来源清晰。

七、SIMT：看起来像一起执行，但每个线程仍有自己的身份

7.1 SIMT 不是简单的 SIMD 改名

SIMT 是 Single Instruction, Multiple Threads。可以翻译为单指令、多线程。

在 CUDA Kernel 中，每个线程：

• 有自己的 threadIdx；
• 有自己的寄存器状态；
• 可以计算不同地址；
• 可以读取不同数据；
• 功能上可以进入不同分支。

这使程序员能够按照“一个线程处理一个元素”的自然方式写代码：

int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) {
output[i] = input[i] * 2.0f;
}

所有线程运行同一份 Kernel 函数，但 i 不同，所以处理的数据不同。

7.2 什么是分支分歧

看下面的代码：

if ((threadIdx.x % 2) == 0) {
// 路径 A
} else {
// 路径 B
}

一个 Warp 中：

• Lane 0、2、4……走 A；
• Lane 1、3、5……走 B。

同一 Warp 内出现不同控制路径，这就是 Warp Divergence，通常译为线程束分歧或分支分歧。

为了保持编程模型语义，硬件需要在不同活动掩码下推进相关路径。图中先让偶数 Lane 对应的活动位置执行 A，再让奇数 Lane 对应的位置执行 B，最后汇合。

这是一幅编程模型示意图，不应被理解为所有架构都必须采用完全相同的固定微架构时序。尤其从 Volta 开始存在 Independent Thread Scheduling，需要更加谨慎地编写依赖 Warp 内同步的代码。但对初学者而言，下面的性能直觉仍然成立：

同一 Warp 中控制流越一致，通常越容易充分利用执行资源。

7.3 分支不是敌人

不要看到 if 就害怕。分支是否造成 Warp 分歧，取决于同一 Warp 内线程的条件结果是否不同。

例如：

if (blockIdx.x == 0) {
// …
}

同一个 Block 内所有线程看到相同的 blockIdx.x。若每个 Warp 完全位于一个 Block 内，那么 Warp 内条件一致，不会因为这个条件产生 Lane 之间的分歧。

再例如边界判断：

if (i < N) {
output[i] = input[i] * 2.0f;
}

通常只有最后一个 Block 的最后部分 Warp 可能出现少量无效 Lane。这种边界判断是保证正确性所必需的，不能为了“消除分支”而越界访问内存。

优化的第一原则是正确，第二原则是测量。不要在没有性能证据时破坏清晰性。

7.4 为什么不能再依赖隐式 Warp 同步

一些老代码假设：

同一 Warp 的线程天然在每一条指令上严格同步，因此共享数据后不需要任何同步操作。

这种写法可能在新架构上出现问题。需要 Warp 内同步时，应使用有明确语义的同步与 Warp 级原语，例如 __syncwarp()，并正确处理参与掩码。

本篇的 __ballot_sync 示例让所有活动 Lane 在进入分支前调用相同投票函数和相同参与掩码。这是有意设计的安全模式。

八、计算能力：不是 CUDA Toolkit 版本，也不是性能分数

8.1 计算能力表示什么

Compute Capability 通常写为 major.minor，例如：

7.5
8.6
8.9
9.0

它描述 GPU 硬件架构向 CUDA 暴露的功能集合和部分资源特征。不同计算能力可能影响：

• 支持哪些指令和硬件特性；
• 每个 Block、SM 的部分资源上限；
• 某些数据类型和原子操作是否可用；
• 编译器应生成哪类目标代码；
• 某些优化策略与吞吐特征。

8.2 计算能力不会因为升级驱动而改变

假设一张 GPU 的计算能力是 7.5：

• 更新 NVIDIA 驱动后，它仍然是 7.5；
• 安装 CUDA Toolkit 11、12 或其他兼容版本后，它仍然是 7.5；
• 更换 CMake 版本后，它仍然是 7.5。

因为计算能力是 GPU 硬件架构属性，不是软件包版本号。

8.3 CUDA Toolkit 版本是什么

CUDA Toolkit 是开发工具集合，包含：

• NVCC 编译器驱动；
• CUDA Runtime 与开发库；
• 头文件；
• 调试与分析工具；
• 各类配套组件。

CUDA Toolkit 12.0 和 Compute Capability 12.0 只是数字形式看起来相似，含义完全不同。不要把二者建立一一对应。

8.4 sm_75 是什么

当计算能力是 7.5 时，常见目标架构写法是：

sm_75

使用 NVCC 时可能看到：

nvcc –arch=sm_75 program.cu –o program.exe

CMake 中通常写不带小数点的数值：

cmake –S . –B build –DCMAKE_CUDA_ARCHITECTURES=75

它表示为相应 SM 版本生成目标代码，不表示“优化等级 75%”，也不是显卡显存大小。

8.5 compute_75 与 sm_75 的初步区别

在传统 NVCC 目标语义中，可以先建立以下直觉：

• compute_75 通常表示虚拟架构与 PTX 级目标；
• sm_75 通常表示对应真实架构的机器代码目标。

Fat Binary 可以同时包含多种目标，使程序覆盖多个 GPU 架构。更完整的前向兼容、PTX JIT 与 -gencode 会在编译专题展开。本篇只要求你能根据设备查询结果配置当前实验。

8.6 如何查自己的计算能力

最可靠的方式有三种：

可尝试：

nvidia-smi —query-gpu=name,compute_cap —format=csv

如果本机驱动版本不支持这个字段，不代表 GPU 不支持 CUDA，直接运行本文查询程序即可。

8.7 同一计算能力是否代表性能一样

不代表。两张 GPU 即使计算能力相同，也可能有不同的：

• SM 数量；
• 显存容量；
• 显存带宽；
• 功耗与频率；
• 散热条件；
• 产品定位；
• 某些硬件资源规模。

计算能力回答的是“架构功能与目标兼容性”问题，不是完整性能评分。评价实际任务必须结合硬件规模、数据特征和真实基准测试。

九、实战一：生成自己的 GPU 能力报告

9.1 这个程序解决什么需求

网上的 CUDA 教程经常写：

• Warp 是 32；
• 每个 Block 最大 1024 个线程；
• 某卡有多少 SM；
• Shared Memory 有多少。

但学习真正有价值的动作，是让程序查询你正在使用的设备。这样你可以把抽象名词变成真实数据，并避免把别人的显卡参数当成自己的。

程序文件位于：

blogs/code/02/device_query_lite.cu

9.2 完整代码

#include <cuda_runtime.h>

#include <cstdio>
#include <cstdlib>

#define CUDA_CHECK(call) \\
do { \\
const cudaError_t error__ = (call); \\
if (error__ != cudaSuccess) { \\
std::fprintf(stderr, "CUDA error at %s:%d: %s\\n", \\
__FILE__, __LINE__, cudaGetErrorString(error__)); \\
return EXIT_FAILURE; \\
} \\
} while (0)

int main() {
int device_count = 0;
CUDA_CHECK(cudaGetDeviceCount(&device_count));

if (device_count == 0) {
std::fprintf(stderr, "No CUDA-capable GPU was found.\\n");
return EXIT_FAILURE;
}

std::printf("CUDA device count: %d\\n", device_count);

for (int device = 0; device < device_count; ++device) {
cudaDeviceProp prop{};
CUDA_CHECK(cudaGetDeviceProperties(&prop, device));

const double global_memory_gib =
static_cast<double>(prop.totalGlobalMem) /
(1024.0 * 1024.0 * 1024.0);

const double peak_bandwidth_gb_s =
2.0 * static_cast<double>(prop.memoryClockRate) * 1000.0 *
(static_cast<double>(prop.memoryBusWidth) / 8.0) / 1.0e9;

std::printf("\\n=== Device %d ===\\n", device);
std::printf("Name : %s\\n", prop.name);
std::printf("Compute capability : %d.%d\\n",
prop.major, prop.minor);
std::printf("Streaming multiprocessors : %d\\n",
prop.multiProcessorCount);
std::printf("Warp size : %d threads\\n", prop.warpSize);
std::printf("Maximum threads per block : %d\\n",
prop.maxThreadsPerBlock);
std::printf("Maximum block dimensions : (%d, %d, %d)\\n",
prop.maxThreadsDim[0], prop.maxThreadsDim[1],
prop.maxThreadsDim[2]);
std::printf("Maximum grid dimensions : (%d, %d, %d)\\n",
prop.maxGridSize[0], prop.maxGridSize[1],
prop.maxGridSize[2]);
std::printf("Global memory : %.2f GiB\\n",
global_memory_gib);
std::printf("Shared memory per block : %zu bytes\\n",
static_cast<size_t>(prop.sharedMemPerBlock));
std::printf("Registers per block : %d\\n", prop.regsPerBlock);
std::printf("Core clock (reported) : %.0f MHz\\n",
prop.clockRate / 1000.0);
std::printf("Memory bus width : %d bits\\n",
prop.memoryBusWidth);
std::printf("Peak memory bandwidth (est.) : %.1f GB/s\\n",
peak_bandwidth_gb_s);
std::printf("Concurrent kernels : %s\\n",
prop.concurrentKernels ? "yes" : "no");
std::printf("Unified virtual addressing : %s\\n",
prop.unifiedAddressing ? "yes" : "no");
}

std::printf(
"\\nNote: CUDA core count is intentionally not estimated here. "
"It is architecture-specific and is not a stable cudaDeviceProp field.\\n");
return EXIT_SUCCESS;
}

9.3 Windows 编译

在“x64 Native Tools Command Prompt for VS 2022”或已经正确加载 MSVC 环境的终端中：

cd cuda-notes\\blogs\\code\\02
nvcc –std=c++14 device_query_lite.cu –o device_query_lite.exe
.\\device_query_lite.exe

如果使用 CMake：

cd cuda-notes\\blogs\\code\\02
cmake –S . –B build –G "Visual Studio 17 2022" –A x64 –DCMAKE_CUDA_ARCHITECTURES=75
cmake —build build —config Release
.\\build\\Release\\device_query_lite.exe

请把 75 改成自己 GPU 的计算能力数字。如果暂时不知道，可以第一次配置时省略 CMAKE_CUDA_ARCHITECTURES，先运行查询程序，再重新明确目标。

9.4 Linux 编译

cd cuda-notes/blogs/code/02
nvcc -std=c++14 device_query_lite.cu -o device_query_lite
./device_query_lite

使用 CMake：

cmake -S . -B build -DCMAKE_BUILD_TYPE=Release -DCMAKE_CUDA_ARCHITECTURES=75
cmake –build build -j
./build/device_query_lite

9.5 本机实测输出

本文在 RTX 2060 上得到：

CUDA device count: 1

=== Device 0 ===
Name : NVIDIA GeForce RTX 2060
Compute capability : 7.5
Streaming multiprocessors : 30
Warp size : 32 threads
Maximum threads per block : 1024
Maximum block dimensions : (1024, 1024, 64)
Maximum grid dimensions : (2147483647, 65535, 65535)
Global memory : 6.00 GiB
Shared memory per block : 49152 bytes
Registers per block : 65536
Core clock (reported) : 1200 MHz
Memory bus width : 192 bits
Peak memory bandwidth (est.) : 264.0 GB/s
Concurrent kernels : yes
Unified virtual addressing : yes

你的输出不同是正常的。重要的不是得到和本文一样的数字，而是能解释每一行。

9.6 每个字段到底有什么用

name

设备名称。多 GPU 机器可能列出多个设备，编号从 0 开始。代码不应只凭名称硬编码能力，因为产品名称不是 Runtime 的功能检测机制。

major 与 minor

组成计算能力。例如 7 和 5 组合为 7.5。编译目标通常对应 sm_75。

multiProcessorCount

SM 数量。这是理解 GPU 并行规模的重要入口，但不能直接乘一个固定常数就跨架构比较所有性能。

warpSize

Warp 的线程数。本文实测为 32。后面的映射代码会使用设备端内建 warpSize。

maxThreadsPerBlock

单个 Block 允许的最大线程总数。即使三个维度分别有较大的上限，三个维度乘积仍不能超过这个总数。

例如下面的 Block：

dim3 block(32, 32);

总线程数是 1024，可能处在设备上限。下面的配置：

dim3 block(32, 32, 2);

总线程数是 2048，即使每个单独维度没有越过对应维度上限，乘积也可能超过 maxThreadsPerBlock，导致 Kernel 启动失败。

三维 Block 总线程数为：

T

b

=

B

x

B

y

B

z

T_b = B_x B_y B_z

Tb​=Bx​By​Bz​

必须满足：

T

b

≤

T

m

T_b \\le T_m

Tb​≤Tm​

其中

T

m

T_m

Tm​ 是设备允许的每 Block 最大线程数。

maxThreadsDim

分别表示 Block 的 x、y、z 维度上限。除了检查总线程数，还要检查每个维度。

maxGridSize

Grid 在各维度的 Block 数上限。对于大多数初学示例，一维 x 上限非常大，但工程代码仍应基于查询属性和问题规模。

totalGlobalMem

设备 Global Memory 总量。本文转换为 GiB：

M

g

=

M

b

/

2

30

M_g = M_b / 2^{30}

Mg​=Mb​/230

其中

M

b

M_b

Mb​ 是字节数，

M

g

M_g

Mg​ 是 GiB。

系统中显示的可用显存可能小于总显存，因为驱动、图形界面、其他进程和 CUDA 上下文会占用资源。

sharedMemPerBlock

单个 Block 可使用的 Shared Memory 基础上限之一。不同架构还可能提供可配置 Shared Memory、Opt-in 上限等属性，不能仅靠这一行覆盖所有高级情况。本系列会在共享内存专题中查询更完整属性。

regsPerBlock

一个 Block 相关的寄存器资源上限。每个线程用多少寄存器由编译结果和代码决定。线程数乘以每线程寄存器需求，会影响 Block 能否驻留和同时驻留数量。

clockRate

Runtime 报告的核心时钟属性。真实运行频率会受到功耗、温度、负载、动态加速和系统状态影响，所以不要拿这一行当作固定性能承诺。

memoryBusWidth 与估算带宽

示例使用报告的显存时钟和总线位宽给出一个理论估算。简化公式：

B

p

=

2

f

m

w

/

8

B_p = 2 f_m w / 8

Bp​=2fm​w/8

其中

f

m

f_m

fm​ 是显存时钟，

w

w

w 是总线位宽，系数 2 表示双沿数据传输。代码还完成了 kHz 与 GB/s 的单位换算。

这是理论峰值估算，不是应用实际带宽。真实 Kernel 会受到访问模式、缓存、协议开销、并发度和硬件状态影响。后续应该通过基准测试或分析工具测量。

9.7 为什么程序故意不打印 CUDA Core 数量

网上有些 deviceQuery 简化代码根据计算能力维护一张“每 SM CUDA Core 数”表，然后计算总数。这在覆盖的架构范围内可以得到产品规格近似，但需要持续维护架构映射，也容易让初学者误以为：

CUDA Thread 数量 = CUDA Core 数量

这是错误的。

因此，本程序明确不猜。需要产品规格时查官方规格页；需要程序兼容性时查计算能力和设备属性；需要性能判断时做真实测量。三个问题使用三种证据，不混在一起。

十、实战二：观察 Thread、Warp、Lane 和 SM

10.1 为什么选择 2 个 Block、每个 40 个线程

如果使用 32 或 64 个线程，Warp 都是完整的，不容易看见边界。

40 是一个故意选择的实验数字：

• 一个 Block 形成 2 个 Warp；
• Warp 0 有 32 个线程；
• Warp 1 只有 8 个线程；
• 两个 Block 各自重新从 Warp 0、Lane 0 开始；
• 总共只有 80 行，仍然容易阅读。

10.2 为什么不直接在 GPU 中 printf

设备端 printf 的输出顺序不保证，很适合观察并发乱序，但不适合制作清晰的映射表。

本程序让每个 GPU 线程把自己的信息写入数组，然后复制回 CPU，由 CPU 按全局编号顺序打印。这样：

• GPU 仍然负责生成每个线程的记录；
• 输出顺序稳定；
• 更容易比较两个 Block；
• 顺便预习 cudaMalloc、cudaMemcpy 和 cudaFree。

这些内存 API 会在第三篇系统讲解，本篇只需要理解“GPU 写，CPU 取回并打印”。

10.3 完整代码

#include <cuda_runtime.h>

#include <cstdio>
#include <cstdlib>

#define CUDA_CHECK(call) \\
do { \\
const cudaError_t error__ = (call); \\
if (error__ != cudaSuccess) { \\
std::fprintf(stderr, "CUDA error at %s:%d: %s\\n", \\
__FILE__, __LINE__, cudaGetErrorString(error__)); \\
return EXIT_FAILURE; \\
} \\
} while (0)

struct ThreadInfo {
int block;
int thread;
int global;
int warp_in_block;
int lane;
unsigned int sm;
};

__device__ __forceinline__ unsigned int read_smid() {
unsigned int sm;
asm("mov.u32 %0, %smid;" : "=r"(sm));
return sm;
}

__global__ void collect_thread_info(ThreadInfo* info) {
const int global =
static_cast<int>(blockIdx.x * blockDim.x + threadIdx.x);

info[global].block = static_cast<int>(blockIdx.x);
info[global].thread = static_cast<int>(threadIdx.x);
info[global].global = global;
info[global].warp_in_block =
static_cast<int>(threadIdx.x / warpSize);
info[global].lane =
static_cast<int>(threadIdx.x % warpSize);
info[global].sm = read_smid();
}

int main() {
constexpr int grid_size = 2;
constexpr int block_size = 40;
constexpr int total_threads = grid_size * block_size;

ThreadInfo* device_info = nullptr;
CUDA_CHECK(cudaMalloc(
&device_info, total_threads * sizeof(ThreadInfo)));

collect_thread_info<<<grid_size, block_size>>>(device_info);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());

ThreadInfo host_info[total_threads]{};
CUDA_CHECK(cudaMemcpy(
host_info,
device_info,
total_threads * sizeof(ThreadInfo),
cudaMemcpyDeviceToHost));

std::printf("Configuration: grid=%d, block=%d, total=%d\\n",
grid_size, block_size, total_threads);
std::printf("Each block has ceil(40 / 32) = 2 warps.\\n\\n");
std::printf(" block thread global warp_in_block lane sm\\n");
std::printf(" —– —— —— ————- —- –\\n");

for (const ThreadInfo& item : host_info) {
std::printf(" %5d %6d %6d %13d %4d %2u\\n",
item.block,
item.thread,
item.global,
item.warp_in_block,
item.lane,
item.sm);
}

CUDA_CHECK(cudaFree(device_info));

std::printf(
"\\nImportant: warp numbering restarts in every block. "
"The SM id is diagnostic only; never rely on it "
"for program correctness.\\n");
return EXIT_SUCCESS;
}

10.4 编译运行

直接使用 NVCC：

nvcc –std=c++14 thread_warp_mapping.cu –o thread_warp_mapping.exe
.\\thread_warp_mapping.exe

Linux：

nvcc -std=c++14 thread_warp_mapping.cu -o thread_warp_mapping
./thread_warp_mapping

CMake 在一次构建中会生成三个程序：

cmake –S . –B build –G "Visual Studio 17 2022" –A x64 –DCMAKE_CUDA_ARCHITECTURES=75
cmake —build build —config Release
.\\build\\Release\\thread_warp_mapping.exe

10.5 先看 Block 0 的关键边界

输出前半部分会类似：

block thread global warp_in_block lane sm
—– —— —— ————- —- —
0 0 0 0 0 0
0 1 1 0 1 0
…
0 30 30 0 30 0
0 31 31 0 31 0
0 32 32 1 0 0
0 33 33 1 1 0
…
0 39 39 1 7 0

观察：

• Thread 31 是 Warp 0 的 Lane 31；
• Thread 32 进入 Warp 1，并从 Lane 0 开始；
• Thread 39 是 Warp 1 的 Lane 7；
• 这个 Block 的 Warp 1 没有 Lane 8～31 对应的线程。

10.6 再看 Block 1

本机输出的末尾是：

1 30 70 0 30 2
1 31 71 0 31 2
1 32 72 1 0 2
1 33 73 1 1 2
1 34 74 1 2 2
1 35 75 1 3 2
1 36 76 1 4 2
1 37 77 1 5 2
1 38 78 1 6 2
1 39 79 1 7 2

Block 1 的 Thread 0 虽然全局编号是 40，但它仍然是本 Block 的 Warp 0、Lane 0。

这揭示了一个常见错误。下面的 Warp 计算不正确：

int wrong_warp = global_index / warpSize;

它把 Warp 当成可以跨 Block 连续编号。CUDA 中 Block 才是 Warp 划分的边界。对于一维 Block，应使用 Block 内线程编号：

int warp_in_block = threadIdx.x / warpSize;

10.7 为什么两个 Block 的 SM 编号可能不同

本机一次运行中：

• Block 0 记录到 SM 0；
• Block 1 记录到 SM 2。

你的机器可能出现：

• 两个 Block 在同一个 SM；
• 两个 Block 在不同 SM；
• 多次运行看到不同 SM 编号。

这些都可能正常。%smid 是 PTX 特殊寄存器，适合诊断与采样，不应进入程序正确性逻辑。官方 PTX ISA Special Registers 对 %smid 的定位也是诊断和分析用途，并提醒该值具有易变性。

代码中：

asm("mov.u32 %0, %smid;" : "=r"(sm));

使用内联 PTX 读取它。这一行不是初学者日常 Kernel 必需写法，只是为了让“Block 被分配到 SM”变成可观察事实。

10.8 从输出能证明什么，不能证明什么

能够证明：

• 线程全局编号如何计算；
• Warp 编号在 Block 内如何划分；
• Lane 如何在 Warp 内从 0 开始；
• 一个 Block 的线程采样到同一个 SM；
• 不同 Block 可能被调度到不同 SM。

不能证明：

• Block 永远固定在某个 SM；
• SM 编号连续等价于物理空间顺序；
• 两个 Block 同时执行；
• 每个 Lane 对应一颗固定物理核心；
• 当前输出就是所有架构的底层时序。

一个实验的价值不仅是看到什么，还要知道证据边界在哪里。

十一、实战三：把 Warp 的 32 个 Lane 变成一个掩码

11.1 为什么需要 Warp 投票

仅仅打印 lane 仍然容易把 Warp 当作一个抽象名词。CUDA 提供 Warp Vote Functions，让 Warp 中线程共同对条件进行投票。

本篇使用：

unsigned int __ballot_sync(unsigned int mask, int predicate);

直观理解：

• 每个参与 Lane 提供一个真假条件；
• 返回 32 位无符号整数；
• 第 n 位对应 Lane n；
• 条件为真且 Lane 参与时，对应位为 1；
• 条件为假时，对应位为 0。

11.2 完整代码

#include <cuda_runtime.h>

#include <cstdio>
#include <cstdlib>

#define CUDA_CHECK(call) \\
do { \\
const cudaError_t error__ = (call); \\
if (error__ != cudaSuccess) { \\
std::fprintf(stderr, "CUDA error at %s:%d: %s\\n", \\
__FILE__, __LINE__, cudaGetErrorString(error__)); \\
return EXIT_FAILURE; \\
} \\
} while (0)

struct WarpResult {
unsigned int active_mask;
unsigned int even_mask;
unsigned int odd_mask;
int even_count;
int odd_count;
int values[32];
};

__global__ void inspect_warp(WarpResult* result) {
const int lane = static_cast<int>(threadIdx.x);
const unsigned int active = __activemask();
const bool is_even = (lane % 2) == 0;

const unsigned int even_mask =
__ballot_sync(active, is_even);
const unsigned int odd_mask =
__ballot_sync(active, !is_even);

if (is_even) {
result->values[lane] = 100 + lane;
} else {
result->values[lane] = 200 + lane;
}

if (lane == 0) {
result->active_mask = active;
result->even_mask = even_mask;
result->odd_mask = odd_mask;
result->even_count = __popc(even_mask);
result->odd_count = __popc(odd_mask);
}
}

int main() {
WarpResult* device_result = nullptr;
CUDA_CHECK(cudaMalloc(&device_result, sizeof(WarpResult)));
CUDA_CHECK(cudaMemset(device_result, 0, sizeof(WarpResult)));

inspect_warp<<<1, 32>>>(device_result);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());

WarpResult host_result{};
CUDA_CHECK(cudaMemcpy(
&host_result,
device_result,
sizeof(WarpResult),
cudaMemcpyDeviceToHost));

std::printf("active mask : 0x%08x\\n", host_result.active_mask);
std::printf("even mask : 0x%08x, active lanes = %d\\n",
host_result.even_mask,
host_result.even_count);
std::printf("odd mask : 0x%08x, active lanes = %d\\n",
host_result.odd_mask,
host_result.odd_count);

std::printf("\\nValues after the branch:\\n");
for (int lane = 0; lane < 32; ++lane) {
std::printf(
"lane %2d -> %d%s",
lane,
host_result.values[lane],
(lane % 4 == 3) ? "\\n" : " ");
}

CUDA_CHECK(cudaFree(device_result));

const bool ok =
host_result.active_mask == 0xffffffffu &&
host_result.even_mask == 0x55555555u &&
host_result.odd_mask == 0xaaaaaaaau &&
host_result.even_count == 16 &&
host_result.odd_count == 16;

std::printf(
"\\nVerification: %s\\n",
ok ? "PASSED" : "FAILED");
return ok ? EXIT_SUCCESS : EXIT_FAILURE;
}

11.3 编译运行

nvcc –std=c++14 warp_vote_demo.cu –o warp_vote_demo.exe
.\\warp_vote_demo.exe

Linux：

nvcc -std=c++14 warp_vote_demo.cu -o warp_vote_demo
./warp_vote_demo

11.4 本机实测结果

active mask : 0xffffffff
even mask : 0x55555555, active lanes = 16
odd mask : 0xaaaaaaaa, active lanes = 16

Values after the branch:
lane 0 -> 100 lane 1 -> 201 lane 2 -> 102 lane 3 -> 203
lane 4 -> 104 lane 5 -> 205 lane 6 -> 106 lane 7 -> 207
lane 8 -> 108 lane 9 -> 209 lane 10 -> 110 lane 11 -> 211
lane 12 -> 112 lane 13 -> 213 lane 14 -> 114 lane 15 -> 215
lane 16 -> 116 lane 17 -> 217 lane 18 -> 118 lane 19 -> 219
lane 20 -> 120 lane 21 -> 221 lane 22 -> 122 lane 23 -> 223
lane 24 -> 124 lane 25 -> 225 lane 26 -> 126 lane 27 -> 227
lane 28 -> 128 lane 29 -> 229 lane 30 -> 130 lane 31 -> 231

Verification: PASSED

11.5 如何读 0xffffffff

一个十六进制数字表示 4 个二进制位。8 个十六进制数字表示 32 位。

0xffffffff
= 11111111111111111111111111111111

32 位全部为 1，表示 32 个 Lane 在调用 __activemask() 时都处于活动状态。

11.6 如何读 0x55555555

十六进制 5 对应二进制：

5 = 0101

从最低位 Lane 0 开始看，0x55555555 的第 0、2、4……30 位是 1，恰好对应偶数 Lane。

Lane: … 7 6 5 4 3 2 1 0
Mask: … 0 1 0 1 0 1 0 1

因为打印二进制时高位在左、低位在右，第一次阅读掩码很容易把方向看反。只要记住最低有效位对应 Lane 0。

11.7 如何读 0xaaaaaaaa

十六进制 A 对应二进制：

A = 1010

因此 0xaaaaaaaa 的第 1、3、5……31 位为 1，对应奇数 Lane。

偶数掩码和奇数掩码满足：

M

e

∣

M

o

=

M

a

M_e \\mathbin{|} M_o = M_a

Me​∣Mo​=Ma​

其中

M

e

M_e

Me​ 是偶数掩码，

M

o

M_o

Mo​ 是奇数掩码，

M

a

M_a

Ma​ 是全部活动 Lane 掩码，竖线表示按位或。

并且：

M

e

&

M

o

=

0

M_e \\mathbin{\\&} M_o = 0

Me​&Mo​=0

表示两组没有重叠 Lane。

11.8 __popc 做什么

__popc 统计 32 位整数中 1 的数量：

result->even_count = __popc(even_mask);

偶数 Lane 和奇数 Lane 各 16 个，所以结果都是 16。

这类 Warp 原语在后续可以用于：

• Warp 内条件统计；
• 快速判断是否有线程满足条件；
• 构建紧凑索引；
• Warp 级归约与协作；
• 复杂并行算法中的活动线程管理。

11.9 为什么投票放在分支前

示例先让所有活动 Lane 调用：

__ballot_sync(active, is_even);
__ballot_sync(active, !is_even);

然后才进入偶数/奇数分支。

参与 _sync Warp 原语的线程与掩码必须满足接口要求。不能随意让掩码声明的某些线程不调用对应原语，否则结果可能未定义。

NVIDIA 官方 Warp Vote Functions 对 __all_sync、__any_sync、__ballot_sync 和 __activemask 有完整语义说明。现阶段请优先使用带 _sync 的现代接口，不要照抄早期教程中的旧 __ballot。

十二、把三个实验串起来：CUDA 程序究竟怎样执行

现在把从 CPU 启动到 GPU 执行的过程串成一条链。

#mermaid-svg-TT1bsLgFBmnLcruT{font-family:\”trebuchet ms\”,verdana,arial,sans-serif;font-size:16px;fill:#333;}@keyframes edge-animation-frame{from{stroke-dashoffset:0;}}@keyframes dash{to{stroke-dashoffset:0;}}#mermaid-svg-TT1bsLgFBmnLcruT .edge-animation-slow{stroke-dasharray:9,5!important;stroke-dashoffset:900;animation:dash 50s linear infinite;stroke-linecap:round;}#mermaid-svg-TT1bsLgFBmnLcruT .edge-animation-fast{stroke-dasharray:9,5!important;stroke-dashoffset:900;animation:dash 20s linear infinite;stroke-linecap:round;}#mermaid-svg-TT1bsLgFBmnLcruT .error-icon{fill:#552222;}#mermaid-svg-TT1bsLgFBmnLcruT .error-text{fill:#552222;stroke:#552222;}#mermaid-svg-TT1bsLgFBmnLcruT .edge-thickness-normal{stroke-width:1px;}#mermaid-svg-TT1bsLgFBmnLcruT .edge-thickness-thick{stroke-width:3.5px;}#mermaid-svg-TT1bsLgFBmnLcruT .edge-pattern-solid{stroke-dasharray:0;}#mermaid-svg-TT1bsLgFBmnLcruT .edge-thickness-invisible{stroke-width:0;fill:none;}#mermaid-svg-TT1bsLgFBmnLcruT .edge-pattern-dashed{stroke-dasharray:3;}#mermaid-svg-TT1bsLgFBmnLcruT .edge-pattern-dotted{stroke-dasharray:2;}#mermaid-svg-TT1bsLgFBmnLcruT .marker{fill:#333333;stroke:#333333;}#mermaid-svg-TT1bsLgFBmnLcruT .marker.cross{stroke:#333333;}#mermaid-svg-TT1bsLgFBmnLcruT svg{font-family:\”trebuchet ms\”,verdana,arial,sans-serif;font-size:16px;}#mermaid-svg-TT1bsLgFBmnLcruT p{margin:0;}#mermaid-svg-TT1bsLgFBmnLcruT .label{font-family:\”trebuchet ms\”,verdana,arial,sans-serif;color:#333;}#mermaid-svg-TT1bsLgFBmnLcruT .cluster-label text{fill:#333;}#mermaid-svg-TT1bsLgFBmnLcruT .cluster-label span{color:#333;}#mermaid-svg-TT1bsLgFBmnLcruT .cluster-label span p{background-color:transparent;}#mermaid-svg-TT1bsLgFBmnLcruT .label text,#mermaid-svg-TT1bsLgFBmnLcruT span{fill:#333;color:#333;}#mermaid-svg-TT1bsLgFBmnLcruT .node rect,#mermaid-svg-TT1bsLgFBmnLcruT .node circle,#mermaid-svg-TT1bsLgFBmnLcruT .node ellipse,#mermaid-svg-TT1bsLgFBmnLcruT .node polygon,#mermaid-svg-TT1bsLgFBmnLcruT .node path{fill:#ECECFF;stroke:#9370DB;stroke-width:1px;}#mermaid-svg-TT1bsLgFBmnLcruT .rough-node .label text,#mermaid-svg-TT1bsLgFBmnLcruT .node .label text,#mermaid-svg-TT1bsLgFBmnLcruT .image-shape .label,#mermaid-svg-TT1bsLgFBmnLcruT .icon-shape .label{text-anchor:middle;}#mermaid-svg-TT1bsLgFBmnLcruT .node .katex path{fill:#000;stroke:#000;stroke-width:1px;}#mermaid-svg-TT1bsLgFBmnLcruT .rough-node .label,#mermaid-svg-TT1bsLgFBmnLcruT .node .label,#mermaid-svg-TT1bsLgFBmnLcruT .image-shape .label,#mermaid-svg-TT1bsLgFBmnLcruT .icon-shape .label{text-align:center;}#mermaid-svg-TT1bsLgFBmnLcruT .node.clickable{cursor:pointer;}#mermaid-svg-TT1bsLgFBmnLcruT .root .anchor path{fill:#333333!important;stroke-width:0;stroke:#333333;}#mermaid-svg-TT1bsLgFBmnLcruT .arrowheadPath{fill:#333333;}#mermaid-svg-TT1bsLgFBmnLcruT .edgePath .path{stroke:#333333;stroke-width:2.0px;}#mermaid-svg-TT1bsLgFBmnLcruT .flowchart-link{stroke:#333333;fill:none;}#mermaid-svg-TT1bsLgFBmnLcruT .edgeLabel{background-color:rgba(232,232,232, 0.8);text-align:center;}#mermaid-svg-TT1bsLgFBmnLcruT .edgeLabel p{background-color:rgba(232,232,232, 0.8);}#mermaid-svg-TT1bsLgFBmnLcruT .edgeLabel rect{opacity:0.5;background-color:rgba(232,232,232, 0.8);fill:rgba(232,232,232, 0.8);}#mermaid-svg-TT1bsLgFBmnLcruT .labelBkg{background-color:rgba(232, 232, 232, 0.5);}#mermaid-svg-TT1bsLgFBmnLcruT .cluster rect{fill:#ffffde;stroke:#aaaa33;stroke-width:1px;}#mermaid-svg-TT1bsLgFBmnLcruT .cluster text{fill:#333;}#mermaid-svg-TT1bsLgFBmnLcruT .cluster span{color:#333;}#mermaid-svg-TT1bsLgFBmnLcruT div.mermaidTooltip{position:absolute;text-align:center;max-width:200px;padding:2px;font-family:\”trebuchet ms\”,verdana,arial,sans-serif;font-size:12px;background:hsl(80, 100%, 96.2745098039%);border:1px solid #aaaa33;border-radius:2px;pointer-events:none;z-index:100;}#mermaid-svg-TT1bsLgFBmnLcruT .flowchartTitleText{text-anchor:middle;font-size:18px;fill:#333;}#mermaid-svg-TT1bsLgFBmnLcruT rect.text{fill:none;stroke-width:0;}#mermaid-svg-TT1bsLgFBmnLcruT .icon-shape,#mermaid-svg-TT1bsLgFBmnLcruT .image-shape{background-color:rgba(232,232,232, 0.8);text-align:center;}#mermaid-svg-TT1bsLgFBmnLcruT .icon-shape p,#mermaid-svg-TT1bsLgFBmnLcruT .image-shape p{background-color:rgba(232,232,232, 0.8);padding:2px;}#mermaid-svg-TT1bsLgFBmnLcruT .icon-shape .label rect,#mermaid-svg-TT1bsLgFBmnLcruT .image-shape .label rect{opacity:0.5;background-color:rgba(232,232,232, 0.8);fill:rgba(232,232,232, 0.8);}#mermaid-svg-TT1bsLgFBmnLcruT .label-icon{display:inline-block;height:1em;overflow:visible;vertical-align:-0.125em;}#mermaid-svg-TT1bsLgFBmnLcruT .node .label-icon path{fill:currentColor;stroke:revert;stroke-width:revert;}#mermaid-svg-TT1bsLgFBmnLcruT :root{–mermaid-font-family:\”trebuchet ms\”,verdana,arial,sans-serif;}

CPU 准备数据与启动配置

发起 Kernel Grid

Grid 拆分为多个 Block

Block 动态分配到可用 SM

每个 Block 划分为多个 Warp

Warp 中 32 个 Lane 以 SIMT 模型推进

线程按全局索引处理各自数据

CPU 同步并取得结果

12.1 程序员决定的内容

• Grid 维度；
• Block 维度；
• Kernel 代码；
• 每个线程如何根据索引选择数据；
• 使用多少 Shared Memory；
• 是否同步；
• 数据如何传输；
• 编译目标架构。

12.2 Runtime 与硬件负责的内容

• Block 具体调度到哪个 SM；
• Block 执行的先后顺序；
• 什么时候选择哪个就绪 Warp；
• 实际指令如何映射到架构资源；
• 缓存命中和动态硬件状态；
• 在符合 CUDA 语义的前提下完成底层实现。

12.3 正确的关注边界

程序员应该理解 Warp，因为它影响：

• 分支分歧；
• 连续线程如何访问内存；
• Warp 级原语；
• Block 大小选择；
• 资源利用。

但不应该依赖没有被编程模型保证的偶然时序，例如：

• 认为 Block 0 必然最先完成；
• 认为相邻 Block 必然在相邻 SM；
• 认为同一个 Block 每次都在同一 SM；
• 认为一次观察到的输出顺序永远不变；
• 认为 Lane 0 永远对应某颗固定物理核心。

理解硬件，是为了写出符合模型、容易优化的代码，不是为了把偶然实现细节变成脆弱依赖。

十三、Block 大小到底怎么选

这是读者完成本篇后最希望解决的实际问题之一。

13.1 入门阶段的可执行建议

对于简单的一维逐元素 Kernel：

示例：

constexpr int threads = 256;
const int blocks = (N + threads – 1) / threads;
kernel<<<blocks, threads>>>(data, N);

13.2 为什么不是“线程越多越快”

如果每个 Block 线程过多，可能：

• 单个 Block 占据太多寄存器；
• 单个 Block 需要太多 Shared Memory；
• SM 同时可驻留 Block 数下降；
• 可调度 Warp 数受限；
• Kernel 因资源超限无法启动。

如果每个 Block 线程过少，可能：

• 形成大量不完整 Warp；
• 每个 Block 的固定管理成本占比提高；
• 无法提供足够并行度；
• 难以覆盖内存等待；
• 对后续 Shared Memory 协作不方便。

所以 Block 大小是资源与问题结构的折中。

13.3 Occupancy 越高越好吗

Occupancy 常用于描述 SM 上活动 Warp 数相对于硬件上限的比例。简化表达：

O

=

W

a

/

W

m

O = W_a / W_m

O=Wa​/Wm​

其中

W

a

W_a

Wa​ 是活动 Warp 数，

W

m

W_m

Wm​ 是设备允许的最大活动 Warp 数。

高 Occupancy 有助于提供更多可切换 Warp，但不是性能的唯一目标。某些 Kernel 即使 Occupancy 没有达到最高，也可能因为：

• 每个线程拥有更多寄存器；
• 数据复用更好；
• 指令级并行更高；
• 内存访问更有效；
• 计算吞吐已经饱和；

而获得更好性能。

正确说法是：

Occupancy 太低可能提示并行资源不足，但 Occupancy 最高不保证 Kernel 最快。

13.4 用测量代替猜测

一个可靠实验至少应该：

• 使用足够大的数据；
• 预热 GPU；
• 多次运行；
• 使用 CUDA Event 计时；
• 分开数据传输与 Kernel 时间；
• 检查结果正确；
• 固定或记录环境；
• 比较多个 Block 候选值。

这些内容会在性能测量专题中写成完整程序。本篇只建立选择 Block 大小的第一层依据。

十四、十五个常见误区，一次纠正

误区 1：一个 CUDA Thread 对应一颗 CUDA Core

错误。Thread 是逻辑执行实例，硬件分批调度大量线程并复用执行资源。

误区 2：Block 就是 SM

错误。Block 是软件工作组，SM 是硬件执行资源。Block 会被调度到 SM。

误区 3：一个 SM 一次只能执行一个 Block

错误。资源允许时，一个 SM 可以同时驻留多个 Block。

误区 4：一个 Block 可以拆到多个 SM 上

普通 CUDA 执行模型中，一个 Block 的线程在一个 SM 上执行，不跨多个 SM。

误区 5：Block 0 必须先于 Block 1 执行

错误。普通 Kernel 不保证 Block 调度顺序。

误区 6：全局线程编号除以 32 就是正确 Warp 编号

如果想求 Block 内 Warp 编号，这种写法在 Block 大小不是 32 整数倍时会跨 Block 混淆。Warp 划分从每个 Block 重新开始。

误区 7：Block 大小不是 32 的倍数，Kernel 就不能运行

通常仍可运行，但最后一个 Warp 可能不完整，存在利用率损失。

误区 8：所有分支都会造成 Warp 分歧

只有同一 Warp 内 Lane 选择不同路径时才形成该意义上的分歧。Warp 内条件一致的分支不产生 Lane 路径分裂。

误区 9：看到 if 就应该删除

错误。边界判断保证内存安全。优化不能以越界和错误结果为代价。

误区 10：计算能力 8.6 表示 CUDA 8.6

错误。前者是 GPU 架构能力，后者若出现则是软件平台版本概念，二者不同。

误区 11：更新驱动能提高计算能力

错误。驱动不会改变硬件架构版本。

误区 12：sm_90 比 sm_75 优化程度高，所以都写 90

错误。它们表示不同架构目标。目标必须与 Toolkit 支持范围和运行设备兼容。

误区 13：SM 多的 GPU 在任何程序中都一定更快

错误。还取决于显存带宽、频率、架构、数据规模、Kernel 特征、功耗状态和软件实现。

误区 14：Occupancy 达到 100% 就完成优化

错误。Occupancy 是资源和并发状态指标，不是最终性能分数。

误区 15：设备端打印顺序可以表示真实调度顺序

错误。printf 有缓冲和序列化影响，输出先后不能完整还原硬件执行时间线。

十五、遇到问题时如何定位

15.1 invalid configuration argument

常见原因：

• 每个 Block 总线程数超过 maxThreadsPerBlock；
• 某一维度超过 maxThreadsDim；
• Grid 某一维超过 maxGridSize；
• 使用了不支持的启动配置。

定位方法：

15.2 no kernel image is available for execution on the device

常见原因是编译产物没有包含当前 GPU 可执行的目标代码。检查：

• GPU 计算能力；
• CMAKE_CUDA_ARCHITECTURES；
• NVCC -arch 或 -gencode；
• Toolkit 是否认识目标架构；
• 二进制是否由其他机器以过窄架构集合构建。

15.3 invalid device function

同样可能与目标架构不匹配有关，也可能是函数编译和链接配置问题。先缩小为本文最小工程，明确生成目标。

15.4 程序输出的 SM 编号每次不一样

这是允许的。%smid 只用于诊断，不要把某个编号写进正确性判断。

15.5 40 线程实验没有出现 Lane 8～31

这是预期结果。第二个 Warp 对这个 Block 只有 8 个实际线程。程序只为实际 CUDA Thread 创建记录，不会为未承载线程的 Lane 创建虚构记录。

15.6 Warp 掩码不是预期值

先确认启动配置仍然是：

inspect_warp<<<1, 32>>>(device_result);

再确认：

• 所有 32 个线程都到达投票位置；
• __ballot_sync 在分支前调用；
• 参与掩码来自同一位置的 __activemask()；
• 没有提前 return；
• 条件仍然是 Lane 奇偶。

如果你把 Block 改成少于 32 个线程，active_mask 本来就不会是 0xffffffff，自动验证也需要相应修改。

15.7 CMake 配置后修改架构没有生效

可以显式重新配置：

cmake –S . –B build –DCMAKE_CUDA_ARCHITECTURES=86
cmake —build build —config Release

如果生成器或缓存状态混乱，创建一个新的、明确命名的构建目录通常更安全：

cmake –S . –B build-sm86 –DCMAKE_CUDA_ARCHITECTURES=86

不要在不确认路径的情况下递归删除目录。

十六、把知识连接到真实需求

16.1 我只想用 PyTorch，为什么还要懂 SM 和 Warp

即使不手写大量 CUDA Kernel，这些概念仍能帮助你理解：

• 为什么 Batch 太小可能无法充分利用 GPU；
• 为什么不同张量形状性能差异明显；
• 为什么某些算子受显存带宽限制；
• 为什么分支和不规则数据可能不适合 GPU；
• 为什么编译扩展需要指定目标架构；
• 为什么同一模型换卡后需要重新测量。

你不需要为了使用框架而成为微架构专家，但正确的执行模型会让调优和排错更有依据。

16.2 我想比较两张显卡，应该看什么

不要只比较一个指标。至少要区分：

• 计算能力：功能与编译目标；
• SM 数量与架构：并行执行资源；
• 显存容量：能否容纳数据和模型；
• 显存带宽：数据供给能力；
• 对目标数据类型的吞吐：例如 FP32、FP16、BF16、INT8；
• 功耗与散热：持续性能；
• 软件栈与驱动支持；
• 你的真实工作负载测试。

对于深度学习训练，显存容量可能先决定“能不能跑”；对于带宽受限的逐元素操作，显存带宽可能比理论算力更关键；对于复杂自定义 Kernel，代码质量可能掩盖硬件差距。

16.3 我只想让程序先跑起来，Block 应该设多少

一维逐元素任务可以从 256 开始：

int threads = 256;
int blocks = (N + threads – 1) / threads;
kernel<<<blocks, threads>>>(...);

同时保留边界判断。确认正确后，比较 128、256、512。不要在没有测量时宣称 256 永远最佳。

二维图像可以从：

dim3 block(16, 16);

开始。它有 256 个线程和 8 个 Warp，二维坐标自然。但某些内存访问或算法可能更适合 32×8 等形状，仍要结合连续内存方向和测量。

16.4 为什么我的 GPU 利用率不高

低利用率可能来自多种原因：

• 数据规模太小；
• CPU 数据准备成为瓶颈；
• 频繁同步；
• Host 与 Device 传输过多；
• Kernel 很短，启动开销占比高；
• Block 配置不足以提供并行度；
• 寄存器或 Shared Memory 限制驻留；
• 内存访问不连续；
• 分支和工作量不均；
• 算法本身串行依赖强；
• 监控采样方式不合适。

本篇只能帮你排除线程层级与基础配置问题。真正定位需要计时和性能分析工具，后续会逐步进入。

16.5 为什么我有很多 Block，却只看到少量 SM

Grid 的 Block 数不要求等于 SM 数。大量 Block 会被分批调度到有限 SM 上。这正是 CUDA 可扩展编程模型的重要价值：同一个 Grid 可以在 SM 数量不同的 GPU 上运行，只是完成速度和调度批次不同。

十七、动手实验：不要只阅读

下面的实验按难度递增。每个实验都要求先写预测，再运行验证。

实验 1：记录自己的 GPU 身份卡

运行：

device_query_lite

在学习笔记中填写：

GPU 名称：
计算能力：
SM 数量：
Warp 大小：
每 Block 最大线程数：
Global Memory：
每 Block Shared Memory：
编译架构：

如果有多张 GPU，为每张设备各记录一份。

实验 2：把 Block 大小改成 32

修改：

constexpr int block_size = 32;

预测：

• 每个 Block 有几个 Warp；
• Block 1 的全局编号从多少开始；
• Block 1 的 Thread 0 是哪个 Warp、哪个 Lane。

再运行确认。

实验 3：把 Block 大小改成 33

预测：

• 每个 Block 仍然需要几个 Warp；
• 第二个 Warp 有几个有效 Lane；
• 与 32 相比，多创建一个线程为什么需要多一个 Warp 分组。

这个实验能非常直观地解释“向上取整”。

实验 4：把 Block 大小改成 64

此时每个 Block 应有两个完整 Warp。检查边界：

thread 31 -> warp 0, lane 31
thread 32 -> warp 1, lane 0
thread 63 -> warp 1, lane 31

实验 5：把 Grid 改成 100 个 Block

输出会很多，可以修改程序，只打印每个 Block 的 Thread 0：

if (item.thread == 0) {
std::printf(...);
}

观察多少个不同 SM 被采样到。不要期待每次完全一致，也不要把采样数当作所有 SM 一定同时工作的证明。

实验 6：修改 Warp 投票条件

把偶数条件改为：

const bool selected = lane < 8;

预测掩码。Lane 0～7 为 1，十六进制结果应该是：

0x000000ff

再尝试：

const bool selected = lane >= 16;

预测：

0xffff0000

实验 7：创建半个 Warp

把启动改成：

inspect_warp<<<1, 16>>>(device_result);

先不要运行，预测：

• active_mask；
• 偶数掩码；
• 奇数掩码；
• 各自的 __popc。

注意还要调整结果校验，否则程序会故意报告 FAILED。这个失败不是 CUDA 出错，而是测试期望仍然写着 32 个活动线程。

实验 8：二维 Block 的 Warp 边界

将映射程序改成二维：

dim3 block(16, 4);
dim3 grid(1, 1);

计算线性线程编号：

int local = threadIdx.x + blockDim.x * threadIdx.y;
int warp = local / warpSize;
int lane = local % warpSize;

预测：

• y=0 的 16 个线程属于哪些 Lane；
• y=1 的 16 个线程是否与 y=0 组成同一 Warp；
• y=2 从哪个 Warp、哪个 Lane 开始。

完成这个实验，你就真正理解了“x 维变化最快”。

十八、自测题：检验是否真正理解

问题 1

一个 Grid 有 100 个 Block，GPU 有 20 个 SM。是否意味着只有 20 个 Block 能运行，剩余 80 个会丢失？

问题 2

一个 Block 有 96 个线程，会形成多少个 Warp？

问题 3

一个 Block 有 100 个线程，会形成多少个 Warp？最后一个 Warp 有多少个有效线程？

问题 4

Block 1 的 Thread 0 能否与 Block 0 的最后几个线程组成同一个 Warp？

问题 5

blockDim.x == 40 时，threadIdx.x == 37 的 Warp 编号与 Lane 编号是多少？

问题 6

更新显卡驱动后，Compute Capability 7.5 会不会变成 8.0？

问题 7

为什么不能根据一次 %smid 输出，把 Block 永久绑定关系写进算法？

问题 8

同一 Warp 中所有线程都执行 if (blockIdx.x == 0) 的同一分支，这一定产生 Warp 分歧吗？

问题 9

为什么 dim3 block(32, 32, 2) 可能非法，即使 x、y、z 单独看都不大？

问题 10

Occupancy 达到 100%，是否证明 Kernel 已经是最快实现？

参考答案

如果你能不看正文回答 8 道以上，就已经建立了进入数据处理 Kernel 的基础模型。

十九、术语表：以后看到这些词不再慌

Host

通常指 CPU 端程序与主机内存环境。负责准备、启动、同步和结果处理。

Device

CUDA 语境下通常指 GPU 设备及其执行和内存环境。

Kernel

由 Host 启动、在 Device 上被大量线程执行的函数。

Grid

一次 Kernel 启动产生的全部 Block 集合。

Thread Block

一组可在同一 SM 上协作的线程。Block 内可使用 Shared Memory 和 Block 级同步。

Thread

CUDA 编程模型中的逻辑执行实例，拥有自己的索引和状态。

SM

Streaming Multiprocessor，负责驻留、调度和执行大量 CUDA 线程的 GPU 硬件单元。

Warp

同一 Block 中按线性线程编号划分的 32 线程组，是理解执行与性能的重要单位。

Lane

线程在 Warp 内的位置，通常为 0～31。

SIMT

Single Instruction, Multiple Threads。线程运行同一 Kernel，但保留自己的身份和控制流状态。

Warp Divergence

同一 Warp 的 Lane 进入不同控制路径，导致硬件需要在不同活动掩码下推进相关路径。

Active Mask

用位集合表示 Warp 中当前参与或活动的 Lane。

Compute Capability

GPU 硬件架构能力版本，写成 major.minor，例如 7.5。

sm_xy

常见真实架构目标标记，例如计算能力 7.5 对应 sm_75。

PTX

CUDA 工具链中的虚拟指令集与中间表示层之一，可由驱动在合适条件下进一步编译。

Register

SM 上的重要片上资源，通常保存线程局部变量。每线程寄存器使用会影响资源驻留。

Shared Memory

Block 内线程可共同访问的片上内存空间，容量有限但延迟和带宽特征与 Global Memory 不同。

Global Memory

GPU 上容量较大的设备内存。多个 Block 和 Kernel 可以访问，但访问模式对性能影响明显。

Occupancy

活动 Warp 相对硬件上限的比例或相关资源驻留指标。它有参考价值，但不是最终性能分数。

Coalescing

相邻线程的内存访问能够以硬件友好的方式合并。将在内存专题详细学习。

Runtime API

本文使用的 cudaGetDeviceProperties、cudaMalloc、cudaMemcpy 等接口所属的 CUDA Runtime 编程接口。

Device Property

通过 Runtime 查询到的设备属性，例如计算能力、SM 数量和线程上限。

Intrinsic

编译器和硬件提供的内建函数，例如 __ballot_sync、__popc。

Synchronization

协调线程或 Host/Device 执行进度的机制。同步范围和内存可见性必须明确。

Scheduling

硬件和 Runtime 对 Block、Warp 等工作的安排过程。程序不能依赖未被接口保证的具体顺序。

二十、建议的 120 分钟学习安排

如果你希望这篇文章真正转化为能力，可以按下面节奏完成。

第 0～15 分钟：恢复第一篇环境

• 执行 nvidia-smi；
• 执行 nvcc –version；
• 重新运行第一篇 hello_cuda；
• 确认终端与构建环境没有变化。

第 15～35 分钟：建立设备身份卡

• 编译 device_query_lite；
• 保存完整输出；
• 查官方计算能力表；
• 将自己的 Compute Capability 转为 sm_xy 和 CMake 数字。

第 35～60 分钟：手算线程映射

• 在纸上画 2 个 Block；
• 每个 Block 写 40 个线程；
• 标记 0～31 与 32～39；
• 算出 Warp 和 Lane；
• 暂时不要运行程序。

第 60～80 分钟：用程序验证

• 运行 thread_warp_mapping；
• 对照手算结果；
• 记录两个 Block 的 SM 编号；
• 再运行三次，观察编号是否变化；
• 明确哪些现象是保证，哪些只是一次采样。

第 80～100 分钟：理解活动掩码

• 运行 warp_vote_demo；
• 把 0x55555555 手写成二进制；
• 标记 Lane 0 在最低位；
• 修改条件为 lane < 8；
• 预测并验证 0x000000ff。

第 100～120 分钟：输出自己的总结

不要复制本文，尝试用自己的话解释：

能讲清楚，才是真正掌握。

二十一、完整构建工程

目录结构：

blogs/
├── CUDA编程实战02_认识GPU_SM_线程束与计算能力.md
├── assets/
│ └── 02/
│ ├── 01-hero-gpu-sm-warp.png
│ ├── 02-sm-parallel-workshop.png
│ ├── 03-software-hardware-mapping.svg
│ ├── 04-forty-threads-two-warps.svg
│ └── 05-warp-divergence.svg
└── code/
└── 02/
├── CMakeLists.txt
├── device_query_lite.cu
├── thread_warp_mapping.cu
└── warp_vote_demo.cu

CMakeLists.txt：

cmake_minimum_required(VERSION 3.18)

project(cuda_blog_02 LANGUAGES CXX CUDA)

set(CMAKE_CXX_STANDARD 14)
set(CMAKE_CXX_STANDARD_REQUIRED ON)
set(CMAKE_CUDA_STANDARD 14)
set(CMAKE_CUDA_STANDARD_REQUIRED ON)

add_executable(device_query_lite device_query_lite.cu)
add_executable(thread_warp_mapping thread_warp_mapping.cu)
add_executable(warp_vote_demo warp_vote_demo.cu)

foreach(target IN ITEMS
device_query_lite
thread_warp_mapping
warp_vote_demo)
if(MSVC)
target_compile_options(${target} PRIVATE
$<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=/utf-8>
$<$<COMPILE_LANGUAGE:CXX>:/utf-8>)
endif()
endforeach()

Windows 一次运行全部程序：

cd cuda-notes\\blogs\\code\\02
cmake –S . –B build –G "Visual Studio 17 2022" –A x64 –DCMAKE_CUDA_ARCHITECTURES=75
cmake —build build —config Release
.\\build\\Release\\device_query_lite.exe
.\\build\\Release\\thread_warp_mapping.exe
.\\build\\Release\\warp_vote_demo.exe

Linux：

cd cuda-notes/blogs/code/02
cmake -S . -B build -DCMAKE_BUILD_TYPE=Release -DCMAKE_CUDA_ARCHITECTURES=75
cmake –build build -j
./build/device_query_lite
./build/thread_warp_mapping
./build/warp_vote_demo

成功标准

不要要求 SM 编号和本文相同。只检查：

• 三个程序退出码为 0；
• 设备程序能显示正确 GPU 名称和计算能力；
• 映射程序中每个 Block 的 Thread 0 都是 Warp 0、Lane 0；
• Thread 31 是 Warp 0、Lane 31；
• Thread 32 是 Warp 1、Lane 0；
• Warp 投票程序显示 32 个活动 Lane；
• 偶数与奇数掩码各有 16 位；
• 最终显示 Verification: PASSED。

二十二、本篇真正应该带走的十条结论

如果只想记一句话，请记：

程序员创建 Grid 和 Block，GPU 把 Block 调度到 SM，再把 Block 内线程组织成 Warp；性能优化的第一步，是尊重这个层级而不是猜测底层顺序。

二十三、你已经为第三篇准备好了什么

第一篇解决的是：

如何让 CPU 成功启动第一个 GPU Kernel。

第二篇解决的是：

大量 CUDA 线程如何组织，它们与 SM、Warp 和硬件能力有什么关系。

但目前的程序还没有完成一个真正有用的数据计算闭环。下一篇将进入：

CUDA编程实战03：向量加法——显存分配、数据传输与第一个有效计算

下一篇会解决这些现实问题：

• CPU 数组如何复制到 GPU；
• cudaMalloc 到底分配在哪里；
• cudaMemcpyHostToDevice 和 cudaMemcpyDeviceToHost 有什么区别；
• 一个线程如何负责一个数组元素；
• 数据数量不能整除 Block 大小时如何处理；
• 如何验证 GPU 结果不是“看起来能跑”；
• 为什么要检查最大误差；
• 如何区分计算错误、越界和浮点误差；
• 如何用 CUDA Event 进行第一版正确计时；
• 为什么小数组用 GPU 反而可能更慢。

今天三个程序中已经短暂使用了 cudaMalloc、cudaMemcpy 和 cudaFree。下一篇会把它们拆开讲清楚，并完成第一条真正可复用的 CUDA 数据处理流水线。

如果你已经运行了本文程序，请保留自己的设备身份卡和 40 线程映射结果。进入第三篇时，你会直接用这些信息选择 Block 大小和编译目标，不再盲目复制数字。

二十四、八个真实场景：把概念变成判断能力

下面不再介绍新名词，而是用真实开发场景检验你是否能够利用已有知识作出判断。每个场景都按照“现象—推理—行动”展开。

场景 1：一百万个元素应该启动多少线程

假设数组有 1,000,000 个元素，每个 Block 使用 256 个线程。Block 数使用向上取整：

const int threads = 256;
const int blocks = (N + threads – 1) / threads;

计算结果为 3907 个 Block，总共创建：

N

l

=

3907

×

256

=

1000192

N_l = 3907 \\times 256 = 1000192

Nl​=3907×256=1000192

比真实数据多 192 个逻辑线程。最后一个 Block 只有 64 个线程对应有效元素，恰好是两个完整 Warp，其他线程会被边界判断挡住。

行动建议：

• 不要为了让线程总数刚好等于数据量而使用奇怪 Block 大小；
• 选择便于硬件组织和调优的 Block 大小；
• 用向上取整覆盖全部数据；
• 用 if (i < N) 保证最后一批线程不会越界；
• 把“创建了额外逻辑线程”和“访问了额外数组元素”区分开，前者正常，后者是错误。

如果数据量改为 1,000,001，最后一个 Block 有 65 个有效线程，需要三个 Warp，其中第三个 Warp 只有一个 Lane 承担有效工作。这可能稍有利用率损失，但边界处只占整个 Grid 的极小部分，通常不值得为此破坏简单、通用的索引方案。

场景 2：GPU 有 30 个 SM，程序却只启动 4 个 Block

假设设备报告 30 个 SM，而 Kernel 只有：

kernel<<<4, 256>>>();

每个 Block 只能落到一个 SM。即使一个 Block 内有 256 个线程，也最多只有 4 个 SM 得到这批 Block，其他 SM 没有来自该 Grid 的 Block 可执行。此时首先应检查的是问题规模与 Grid 数量，而不是立刻怀疑显卡损坏。

行动建议：

• 查看数据量是否本来就很小；
• 确认 Block 数计算没有使用错误的整数除法；
• 检查是不是把 blocks 和 threads 写反；
• 减少不必要的频繁小 Kernel；
• 在算法允许时合并批次或提高一次处理的数据量；
• 使用分析工具确认真正的空闲原因。

但不要反向得出“Block 数必须等于 SM 数”。实际应用通常会创建远多于 SM 数量的 Block，让硬件有足够工作分批调度。

场景 3：把每个 Block 从 256 改成 1024，为什么没有更快

1024 可能仍在设备允许的每 Block 线程上限内，但合法不等于最佳。

256 线程对应 8 个 Warp，1024 线程对应 32 个 Warp。更大的 Block 可能导致：

• 每个 Block 占用更多线程和 Warp 名额；
• 每个线程的寄存器需求累积成更大的 Block 资源需求；
• 一个 SM 同时能驻留的 Block 数下降；
• Shared Memory 的组织方案受到限制；
• 某些 Block 完成较慢，尾部调度粒度变粗；
• 实际瓶颈根本不在并行线程数量。

行动建议：

经验值只能提供实验起点，不能替代测量。

场景 4：两种分支写法，分歧程度为什么不同

比较两个条件。

条件 A：

if ((i % 2) == 0) {
// A
} else {
// B
}

连续 Lane 的全局编号通常连续，所以奇偶交替。同一 Warp 几乎每次都会一半走 A、一半走 B，形成明显分歧。

条件 B：

if (i < N / 2) {
// A
} else {
// B
}

如果线程按连续全局编号排列，大部分 Warp 会整体处于前半区或后半区，Warp 内条件一致。只有跨越 N/2 边界的少数 Warp 可能分歧。

两段代码都有 if/else，但 Warp 行为不同。真正应该分析的是“相邻 Lane 的条件模式”，而不是统计源代码中有几个 if。

行动建议：

• 观察条件与线程索引的关系；
• 将相似工作尽可能映射给同一 Warp；
• 不要为了消除分支做高成本的数据重排，除非测量证明值得；
• 先保证语义和边界正确；
• 使用 Warp Execution Efficiency 等分析指标验证推测。

场景 5：二维图像为什么有人使用 32×8，有人使用 16×16

两种 Block 都有 256 个线程和 8 个 Warp：

dim3 block_a(16, 16);
dim3 block_b(32, 8);

但是 Warp 内 Lane 映射不同。

对于 16×16：

• x 维只有 16；
• 第一个 Warp 会覆盖两行，每行 16 个线程；
• 如果图像宽度远大于 16，两段地址之间存在跨行间隔。

对于 32×8：

• x 维正好 32；
• 一个 Warp 可以对应同一行中连续的 32 个 x 坐标；
• 对行优先连续数组，这经常有利于建立直观的连续访问模式。

这不表示 32×8 永远快于 16×16。二维邻域、Shared Memory Halo、边界比例、缓存和寄存器压力都可能改变结果。选择二维 Block 时，要同时考虑：

• 数据在内存中的连续方向；
• Warp 线性化规则；
• 算法需要的局部邻域；
• Shared Memory 布局；
• 图像尺寸与边界；
• 实际性能测量。

场景 6：程序在开发机运行，复制到另一张 GPU 却失败

开发机是计算能力 8.6，构建时只生成 sm_86 机器代码；目标机器可能是另一架构。如果二进制没有合适目标代码或可 JIT 的兼容 PTX，就可能出现没有可用 Kernel Image 等错误。

正确的发布思路不是随便把数字改成“最大的架构”，而是：

CMake 可以配置多个目标，例如具体写法需结合当前 CMake 和 Toolkit：

cmake –S . –B build –DCMAKE_CUDA_ARCHITECTURES="75;86"

发布兼容性是一个明确的产品需求，不应到部署失败时才临时猜测。

场景 7：为什么一次实验没有采样到所有 SM

即使 GPU 有 30 个 SM，启动 100 个 Block，也不保证一次 %smid 输出必然覆盖 0～29 的每个编号。原因包括：

• SM 标识编号不一定适合作为连续物理数量解释；
• Block 执行时间太短，调度采样具有偶然性；
• 系统同时有其他 GPU 工作；
• 某些 SM 在采样窗口中没有接到该 Grid 的 Block；
• 打印或记录方式会影响观察；
• %smid 本来就是诊断值，不是分布证明接口。

如果要研究负载分布，应设计持续时间足够的基准、减少观察扰动、多轮采样，并使用 Nsight 工具。本文程序只为证明“Block 最终在某个 SM 上执行，且分配由硬件动态完成”。

场景 8：边界分支有分歧，要不要删除

典型边界判断：

if (i < N) {
output[i] = input[i] * scale;
}

最后一个 Warp 可能部分 Lane 有效、部分 Lane 无效，确实可能形成分歧。但删除判断会导致越界访问，后果可能是：

• 破坏其他数据；
• 得到偶发错误结果；
• 触发非法内存访问；
• 错误在后续同步位置才报告；
• 不同运行环境表现不同，难以定位。

边界 Warp 通常只占整个 Grid 很小比例。除非有充分测量和安全填充策略，否则不应为了这点局部分歧牺牲正确性。

成熟优化顺序是：

结果正确 → 边界安全 → 计时可靠 → 找到瓶颈 → 修改 → 再次验证

这比“看到分支就删”更接近真实 CUDA 工程。

二十五、二十个高频问答

Q1：CUDA Thread 和 Windows/Linux 线程是同一种东西吗？

不是。操作系统线程由 CPU 与操作系统调度，创建和切换成本较高；CUDA Thread 是 GPU 编程模型中的轻量逻辑执行实例，通常成千上万甚至上百万地创建。二者都叫 Thread，但成本模型、调度主体和使用方式不同。

Q2：Kernel 只启动一个线程有没有意义？

可以运行，适合验证工具链或执行一次设备端操作，但无法体现 GPU 的大规模并行优势。第一篇 Hello CUDA 是连通性测试，不是性能基准。

Q3：Grid 可以有零个 Block 吗？

不要依赖零尺寸启动。实际代码应在 Host 端处理 N==0 等空输入，避免构造无意义或非法启动配置。

Q4：一个 Block 的三个维度都没超过上限，就一定合法吗？

不一定。还要检查三个维度乘积是否超过 maxThreadsPerBlock，并考虑 Kernel 资源是否允许启动。

Q5：为什么 threadIdx 和 blockIdx 看起来不是普通整数？

它们是具有 x、y、z 分量的内建三维值。参与算术时要明确取 .x、.y 或 .z，并注意其无符号类型与目标索引类型的转换。

Q6：Warp 大小可以在 Host 端读取吗？

可以，通过 cudaGetDeviceProperties 得到 prop.warpSize，也可以通过设备属性查询接口获得。Kernel 内可使用 warpSize。

Q7：为什么不直接把 Warp 大小永远写成 32？

当前 CUDA 文档明确为 32，很多 Warp 级位掩码也直接体现 32 位。但使用 warpSize 能让意图更清楚，设备报告也为诊断提供依据。某些算法和位掩码仍会基于当前 32-Lane 语义设计，应明确写出假设。

Q8：多个 Block 能共享普通 Shared Memory 吗？

普通 Block 级 Shared Memory 属于 Block 的协作范围，不同 Block 不能把它当作共同数组直接共享。更新架构存在 Cluster 与 Distributed Shared Memory 等高级能力，但那是有额外启动和同步语义的功能，不能反向改变普通 Kernel 的基本规则。

Q9：同一 Block 的线程一定同时开始吗？

不要建立这种要求。它们被划分成多个 Warp，具体推进时机由硬件调度。Block 内需要阶段同步时使用正确的同步原语，而不是依赖“大家应该差不多同时到达”。

Q10：__syncthreads() 能同步整个 Grid 吗？

不能。它是 Block 范围的屏障。不同 Block 的普通 Grid 级同步通常需要拆成多个 Kernel、使用特定 Cooperative Groups 能力或重新设计算法。

Q11：__activemask() 可以在任意分支里调用再拿到完整 Warp 吗？

不能这样假设。它返回调用位置当前活动 Lane 的掩码。如果调用已经处在分歧路径中，得到的可能只是当前路径的活动子集。设计 _sync Warp 原语时必须严格遵守参与掩码要求。

Q12：分支分歧会导致计算结果错误吗？

正常分支分歧不会自动导致错误，硬件会维护各线程语义。它主要可能影响执行效率。错误通常来自错误同步、错误参与掩码、数据竞争或越界，而不是“Lane 选择了不同合法分支”本身。

Q13：为什么 Kernel 启动后要同时调用 cudaGetLastError() 和 cudaDeviceSynchronize()？

前者可以尽早发现启动配置等即时错误；Kernel 启动对 Host 通常是异步的，执行阶段错误往往要在同步或后续 Runtime 调用时才暴露。教学程序同时检查两处，能让错误更靠近源头。

Q14：为什么设备属性中的频率与监控软件不一样？

Runtime 属性是设备报告值，真实运行频率受动态加速、温度、功耗、负载和系统策略影响。需要研究频率时，应在受控负载下使用合适监控工具，而不是把一个静态字段当作实时传感器。

Q15：SM 数量能直接决定应该启动多少 Block 吗？

不能直接决定，但可以帮助判断 Grid 是否明显太小。实际通常启动足够多的 Block，让每个 SM 获得多批工作；具体数量还与数据规模、Block 资源和负载均衡有关。

Q16：为什么一个 SM 会同时驻留多个 Warp，而不是只运行当前 Warp 到结束？

大量驻留 Warp 为隐藏等待提供机会。当某 Warp 因数据依赖或访存等原因暂时不能推进时，调度器可以选择其他就绪 Warp，从而提高吞吐。

Q17：Debug 构建和 Release 构建的性能能直接比较吗？

不应该。调试选项和优化级别会改变代码、寄存器使用和执行行为。正确性调试可以使用 Debug，性能比较应使用明确且一致的优化配置，并记录编译参数。

Q18：计算能力更高的 GPU 一定更快吗？

不一定。计算能力表示架构功能版本，不是完整性能分数。低一级架构的大型数据中心 GPU 可能在某些任务上远快于更新架构的入门产品，反之也可能因特定新功能而改变。

Q19：可以在一份程序中支持多种计算能力吗？

可以。CUDA 工具链能够在二进制中包含多个架构目标和 PTX。具体组合要根据目标用户、二进制大小、Toolkit 支持和兼容策略设计。

Q20：完成本篇后，最值得保留的文件是什么？

保留设备身份卡、三个可执行示例、构建命令和你修改实验后的输出。它们是后续排查“架构不匹配、Block 超限、Warp 计算错误”的最小工具箱，比只收藏网页更有用。

二十六、官方参考资料

本文概念与接口优先依据 NVIDIA 官方文档：

• CUDA Programming Guide
• Programming Model：Grid、Block、Warp 与 SM
• Writing SIMT Kernels
• CUDA Runtime API Device Attributes
• CUDA GPU Compute Capability 查询表
• PTX ISA Special Registers

官方文档内容会随 CUDA 版本更新。文章中的数字示例用于建立方法，实际工程应优先查询当前设备属性，并查看与你使用的 Toolkit 对应的官方文档。

结语：从“知道 GPU 很并行”走到“能解释并行如何发生”

很多人学 CUDA 的第一道坎是环境，第二道坎就是本篇这些看似相似的名词。只要把软件层和硬件层分开，关系就会清楚：

程序员：Grid → Block → Thread
硬件执行：GPU → SM → Warp / Lane
连接规则：Block 调度到 SM，Block 内线程划分为 Warp
兼容依据：Compute Capability 与编译目标

你不需要在第二篇就背下每一代 GPU 的功能单元数量，也不需要马上手算复杂 Occupancy。更重要的是获得一种可持续的学习方法：

当你能看到 0x55555555 并解释每一个 1 为什么对应偶数 Lane 时，Warp 就不再是书上的一句定义；当你能看到两个 Block 采样到不同 SM，又知道程序不能依赖该编号时，你已经开始以 CUDA 工程师的方式思考。

完成本文的三个程序、八个实验和十道自测题后，再进入第三篇。那时，向量加法不会只是“复制一段经典代码”，而会成为你第一次独立设计线程数量、索引、边界、数据流和验证方法的 GPU 程序。