想象你是一家高档餐厅的老板。你的餐厅里有一位全能主厨,他手艺精湛,能做分子料理、烤全羊、法式甜点,但整个后厨只有他一个人。晚市高峰时,二十桌客人同时下单,主厨再快也得一道菜一道菜地做。你意识到瓶颈不在厨艺,而在“只靠一个人”。 于是你请来了一支一千人的帮厨团队。每位帮厨只懂几个基础动作——切丝、焯水、摆盘,但他们能在同一秒内把一千份食材全部备好,再把一千份餐盘刷完。这支团队没有主厨的创造力,却有主厨望尘莫及的“吞吐量”。 这就是 CPU 与 GPU 的本质差异:CPU 是那位全能主厨,用少量核心把复杂逻辑执行到极致;GPU 是那支庞大帮厨团,用海量简单核心做重复劳动。GPGPU(General-Purpose GPU,通用 GPU 计算)要做的事情,就是让这支帮厨团不只服务于图形渲染,还能去帮 AI 训练、科学计算、数据并行这些“通用菜肴”。从今天起,你会发现“显卡”远不止打游戏那么简单——它是现代 AI 世界的数字帮厨团。

岗位关联

学完本节,你将理解 GPU 并行架构的底层逻辑,能看懂 CUDA、OpenCL 等主流 GPGPU 框架的核心概念,为后续掌握并行计算、性能优化与 AI 加速奠定基础。这对 AI 工程实践与面试都是高频考点。

目录
  1. 可视化演示
  2. 并行架构概述
  3. 线程海量并发
  4. CUDA编程模型
  5. OpenCL跨平台
  6. 着色器通用化
  7. 协作与数据流
  8. 性能优化方向
  9. AI 应用场景
  10. Python 演示
  11. 常见误区
  12. 面试题

可视化演示

展示GPGPU从CPU任务分发到GPU并行执行、结果回传的全流程,通过模式切换理解GPU通用计算的核心架构与编程模型。

本节配有交互动效(进入视口后自动加载,约 3KB)。

CPU-GPU协作模式 CPU分发计算任务至GPU GPU并行执行后回传结果 PCIe总线实现高带宽数据传输

并行架构概述

想象那家餐厅后厨。主厨单人独臂,纵然手艺再精,一份惠灵顿牛排从备料到上桌也需要四十分钟;而晚市同时有几十桌客人在等。直到你雇来一千位帮厨,每个人只负责切菜、焯水、摆盘这三个基础动作,瞬间就能把整条出菜流水线铺满。这就是 CPU 与 GPU 最本质的差异:CPU 用少量复杂核心追求“单点极致”,GPU 用海量简单核心追求“整体吞吐”。

GPGPU 的全称是 General-Purpose GPU,意为“通用目的 GPU”。过去 GPU 只服务图形渲染,把三角形变成像素;后来人们发现这台机器做重复数学运算特别快,于是把 AI 矩阵乘法、物理仿真、大数据规约这些“通用任务”也搬了上来。这里的“通用”不是指什么都能干,而是指走出图形领域、应用到各行各业的并行计算。

但要小心:不是所有任务都能靠并行受益。Amdahl 定律给出了一个冷静的边界:

回到厨房:如果一道招牌菜从腌制到装盘有 30% 的步骤只能由主厨亲自完成,那么即使帮厨团从 100 人增加到 10000 人,出菜速度最多也只能提升约 3.3 倍。GPGPU 的第一步,永远是甄别任务中“可切分的菜”所占的比例。

S=1(1−p)+pN\displaystyle S = \frac{1}{(1 - p) + \dfrac{p}{N}}
公式通俗说明:$S$ 是加速比,$p$ 是任务中可并行的比例,$N$ 是并行核心数。它告诉我们:无论你雇多少帮厨,任务里“必须串行”的部分都会成为天花板上限。
  1. 可并行比例 $p$整个任务中能拆成多份同时处理的比例,比如切一千份土豆丝,$p$ 接近 1。
    p=可并行部分工作量总工作量p = \frac{\text{可并行部分工作量}}{\text{总工作量}}
  2. 串行比例 $1 - p$必须由单人顺序完成的部分,比如最后装盘需要主厨亲自把关。
    (1−p)+pN(1 - p) + \frac{p}{N}
  3. 极限加速比当 $N$ 趋向无穷大时,加速比逼近 $\dfrac{1}{1 - p}$。
    lim⁡N→∞1(1−p)+pN=11−p\displaystyle \lim_{N \to \infty} \frac{1}{(1 - p) + \frac{p}{N}} = \frac{1}{1 - p}

线程海量并发

帮厨团虽然人多,却不能一拥而上。好的后厨会先把 1000 人分成若干小组,每组负责几张桌子,组里还有组长统一协调。GPU 内部采用了类似的层级结构:线程(thread)→ 线程块(block)→ 网格(grid)。线程是执行计算的最小个体,线程块是可协作、可共享内存的小团体,网格则是整个被调度的大任务集合。

GPU 在硬件层面执行并行时,并不一个线程一个线程地调度,而是以 warp 为单位——一个 warp 固定包含 32 个线程。当 GPU 执行一条指令时,这 32 个线程同时执行同一条指令,只是各自处理不同的数据。这种模型叫 SIMT(单指令多线程),相当于组长喊一声“切丝!”,整组帮厨同时挥刀,每人切自己面前那根萝卜。

并发线程数的上限由硬件资源决定:

以 NVIDIA A100 为例:108 个 SM,每个 SM 可驻留 64 个 warp,那么理论并发线程上限就是 $108 \times 64 \times 32 \approx 22$ 万个线程同时活跃。帮厨团顷刻间涌动起来——这正是 GPU 能碾压 CPU 做浮点运算的秘密武器。

并发线程数上限=SM 数×每 SM Warp 槽位×32\text{并发线程数上限} = \text{SM 数} \times \text{每 SM Warp 槽位} \times 32
公式通俗说明:SM(流式多处理器)是 GPU 上真正干活的“小厨房”;每个小厨房能同时容纳若干个 warp;每个 warp 有 32 位帮厨。三者相乘就是 GPU 同时能容纳的最大线程数。
  1. SM 数(小厨房个数)GPU 芯片上的计算单元数量,类比后厨里能同时开火的灶台组数。
    SM=GPU 全部流处理器每 SM 流处理器数量\text{SM} = \frac{\text{GPU 全部流处理器}}{\text{每 SM 流处理器数量}}
  2. 每 SM Warp 槽位每个小厨房最多能同时驻留多少个 32 人小组。
    Warp 槽位=每 SM 可驻留线程数32\text{Warp 槽位} = \frac{\text{每 SM 可驻留线程数}}{32}
  3. 32(warp 大小)一次并行调度包含的线程数,是硬件设计上固定的组队规模。
    warp 内线程数=32\text{warp 内线程数} = 32

CUDA编程模型

主厨不可能手把手教一千个帮厨做菜,他只需要把“麻婆豆腐标准作业指引”写成一份文档,复印一千份发下去。在 CUDA 中,这份“文档”就是内核函数(kernel):由 CPU 启动,却在 GPU 的成千上万个线程中并行执行。

CUDA 是 NVIDIA 推出的 GPGPU 编程框架,它把程序世界清晰地分成两半:主机(Host)是 CPU,负责串行逻辑、数据准备、任务调度;设备(Device)是 GPU,负责执行海量并行的内核函数。一个典型流程是:CPU 分配内存 → 数据拷贝到显存 → 启动 kernel → GPU 并行计算 → 结果拷回主存。

为了让每个线程知道自己该处理哪个数据,CUDA 在内核函数中提供了三个内置变量:

这个公式是 CUDA 编程里最基础也最常用的“定位咒语”。掌握它,你就懂得如何让一千位帮厨各自走向自己该切的萝卜,而不是三百人抢一根萝卜、七百人闲着。

全局线程 ID=blockIdx.x×blockDim.x+threadIdx.x\text{全局线程 ID} = \text{blockIdx.x} \times \text{blockDim.x} + \text{threadIdx.x}
公式通俗说明:想象所有线程排成一个大方阵,先找到自己属于哪个块(blockIdx),再乘以每块的人数(blockDim),加上自己在块内的序号(threadIdx),就得到了在全军中的唯一编号。
  1. blockIdx.x(块编号)当前线程块在整个网格中的索引,类比第几组帮厨。
    blockIdx.x∈[0,gridDim.x−1]\text{blockIdx.x} \in [0, \text{gridDim.x} - 1]
  2. blockDim.x(块内线程数)每个线程块包含多少个线程,类比每组的人数。
    blockDim.x=每块线程数量\text{blockDim.x} = \text{每块线程数量}
  3. threadIdx.x(块内序号)当前线程在块中的位置,类比组内第几号帮厨。
    threadIdx.x∈[0,blockDim.x−1]\text{threadIdx.x} \in [0, \text{blockDim.x} - 1]

OpenCL跨平台

不是所有餐厅都愿意买同一品牌厨具。有的后厨用进口灶台,有的用国产集成灶,菜谱却要通用。OpenCL(Open Computing Language)正是这样一本“世界通用菜谱”——由 Khronos 组织维护的开放异构计算标准,CPU、GPU、DSP、FPGA 都可以作为执行设备。

OpenCL 与 CUDA 在思路上高度相似。它定义了 kernel(核函数) 作为设备端执行单元,用 C 风格语言编写;用 NDRange(N 维索引空间) 描述工作项的调度方式。开发者把任务组织成 N 维数组,在每一维上指定全局规模和局部规模,硬件会自动映射到具体设备上并行执行。

OpenCL 中工作项(work-item)的全局下标计算方式如下:

CUDA 与 OpenCL 的选择就像品牌厨具与世界通用菜谱的选择:CUDA 在 NVIDIA 生态中性能挖掘更深、工具链更成熟;OpenCL 则胜在跨平台普适性。前者适合深度绑定 NVIDIA 的场景,后者适合多厂商异构环境。

global_id=group_id×workgroup_size+local_id\text{global\_id} = \text{group\_id} \times \text{workgroup\_size} + \text{local\_id}
公式通俗说明:这与 CUDA 的线程定位公式如出一辙:先看自己属于哪个工作组,乘以组的大小,再加组内偏移,就得到全局唯一的身份编号。
  1. group_id(工作组编号)当前工作组在全局索引空间中的位置。
    group_id=⌊global_idworkgroup_size⌋\text{group\_id} = \lfloor \frac{\text{global\_id}}{\text{workgroup\_size}} \rfloor
  2. workgroup_size(工作组大小)每个工作组包含的 work-item 个数,类比 CUDA 的 blockDim。
    workgroup_size=每小组工作项数量\text{workgroup\_size} = \text{每小组工作项数量}
  3. local_id(组内编号)当前 work-item 在自己工作组内的偏移量。
    local_id=global_id mod workgroup_size\text{local\_id} = \text{global\_id} \bmod \text{workgroup\_size}

着色器通用化

传统 GPU 编程都依附于图形渲染管线——顶点着色器、片元着色器像宴席上的固定流程:洗菜、切菜、热锅、摆盘,每一道工序都有专门环节。但计算着色器(Compute Shader)的出现改变了游戏规则:它相当于在宴会厅旁边单独开了一个小厨房,不需要走宴席流程,也能随时开火炒菜。

计算着色器是 DirectX 11(2010 年)引入的概念,随后被 OpenGL、Vulkan、Metal 广泛支持。它不参与顶点变换、不输出像素颜色,而是作为一个独立阶段,让开发者直接用 GPU 的并行能力完成任意计算任务——粒子物理、布料模拟、后处理特效都能塞进去。

这种模式的本质是提高 GPU 的“任务吞吐量”:

如今 GPGPU 已不再是图形学的专属领地:深度学习中卷积算子、推荐系统的嵌入向量检索、科学计算的粒子模拟,全都借用了这套“着色器通用化”的思想。AI 与图形学在同一块芯片上再次交汇。

计算吞吐量=任务总数执行时间×并行单元数\text{计算吞吐量} = \frac{\text{任务总数}}{\text{执行时间} \times \text{并行单元数}}
公式通俗说明:每秒完成的任务量,取决于总任务量、耗时以及参与计算的并行单元数。计算着色器越多地利用并行单元,单位时间完成的通用计算就越多。
  1. 任务总数需要被并行处理的工作总量,比如 10000 个粒子的位置更新。
    任务总数=∑i=1N任务i\text{任务总数} = \sum_{i=1}^{N} \text{任务}_i
  2. 执行时间从内核启动到全部线程完成所经过的时间。
    执行时间=完成时刻−启动时刻\text{执行时间} = \text{完成时刻} - \text{启动时刻}
  3. 并行单元数GPU 上能同时工作的计算核心数量。
    并行单元数=着色器单元×并发线程数\text{并行单元数} = \text{着色器单元} \times \text{并发线程数}

协作与数据流

后厨里最忙的往往不是锅铲,而是那扇传菜窗口。主厨与帮厨之间隔着一堵墙,食材、调料、半成品都要从窗口递进递出;如果窗口大小只有巴掌宽,后厨再多人也只能干等。在 GPGPU 世界中,这扇窗口就是 PCIe 总线——连接 CPU 主存与 GPU 显存的数据通道。

GPU 不能直接读取 CPU 的主存,因为两者物理隔离。一次完整的 GPGPU 任务通常分四步:CPU 把数据从主存拷贝到显存 → CPU 启动 kernel → GPU 并行计算 → 结果从显存拷回主存。如果你每调用一次并行计算就搬运一次数据,那么 PCIe 传输时间很可能远超 GPU 算力时间。

传输耗时由数据量和总线带宽共同决定:

实践中,CUDA 的异步拷贝和双缓冲技术能把“搬运”与“计算”重叠起来:主厨一边把下一批食材递进窗口,帮厨一边处理当前这批。这让传菜窗口不再闲着,也是 GPGPU 性能优化中必须掌握的一课。

传输时间=数据量总线带宽\text{传输时间} = \frac{\text{数据量}}{\text{总线带宽}}
公式通俗说明:传菜窗口的宽度(带宽)固定时,要传递的食材越多(数据量越大),耗费在窗口的时间就越长。数据搬运是 GPGPU 中最典型的隐形瓶颈。
  1. 数据量需要在主存与显存之间搬运的字节总数,比如 1 个 128MB 的张量。
    数据量=元素个数×每元素字节数\text{数据量} = \text{元素个数} \times \text{每元素字节数}
  2. 总线带宽PCIe 每秒钟能传输的最大字节数,类比窗口每秒能递多少盘菜。
    带宽=通道数×每通道速率\text{带宽} = \text{通道数} \times \text{每通道速率}
  3. 传输时间搬运耗时与数据量成正比,与带宽成反比。
    ttransfer=DBt_{\text{transfer}} = \frac{D}{B}

性能优化方向

帮厨团队效率高不高,不看人多人少,而看三点:备料是否顺手、站位是否合理、动作是否统一。GPU 性能优化同样围绕这三个维度展开。

第一是内存合并访问(coalesced access)。一个 warp 的 32 个线程如果同时访问连续的内存地址,硬件能把 32 个请求合并成一次事务传输;如果线程访问的是随机跳跃的地址,就要拆成多次事务,带宽利用率骤降。好比备料时,帮厨们按照顺序从相邻位置拿取食材,比到处乱跑高效得多。

第二是占用率(occupancy)平衡。占用率指活跃 warp 数与硬件理论上限的比值。占用率越高,越能掩盖内存延迟;但过高的占用率会让每个线程可用的寄存器变少,反而增加内存往返。找到甜点区是调优的关键。

第三是向量化指令。用 float4 一次性处理四个浮点数,指令数降为原来的四分之一,等价于一名帮厨一次端四盘菜。

衡量内存访问效率的核心指标:

回到厨房,真正的效率高手不是让每个帮厨都忙到飞起,而是让食材、动作、走位三条链路全部顺畅。GPU 优化也是如此:合并访问让数据流顺畅,占用率让调度不空转,向量化让指令更精炼。三者齐头并进,才配得上“通用计算”这顶帽子。

内存带宽利用率=理论最小请求数实际内存请求数\text{内存带宽利用率} = \frac{\text{理论最小请求数}}{\text{实际内存请求数}}
公式通俗说明:硬件为完成数据访问“必须”发起的最小请求数,除以实际发起的请求数。合并访问时利用率接近 100%,随机访问时可能跌到 10% 以下。
  1. 理论最小请求数在理想合并情况下,32 个线程所需数据仅需一次事务请求即可完成。
    最小请求数=warp 访问字节数总线事务宽度\text{最小请求数} = \frac{\text{warp 访问字节数}}{\text{总线事务宽度}}
  2. 实际内存请求数随机访问时硬件被迫拆分的实际事务数量。
    实际请求数=非合并访问时的事务拆分次数\text{实际请求数} = \text{非合并访问时的事务拆分次数}
  3. 利用率两者的比值越接近 1,内存带宽越被充分利用。
    利用率=最小请求数实际请求数≤1\text{利用率} = \frac{\text{最小请求数}}{\text{实际请求数}} \le 1

AI 应用场景

深度学习训练加速

阶段:模型训练与推理优化

解决的问题:深度学习训练本质是海量矩阵乘法与卷积运算。传统 CPU 串行执行时,训练一个 ResNet-50 在 ImageNet 上跑一个 epoch 可能耗时数小时;把这些矩阵运算并行地拆分到 GPU 的成千上万个核心上,训练时间可缩短一到两个数量级。

具体案例:某算法团队用 PyTorch 训练图像分类模型,CPU 版本每个 epoch 需要 40 分钟,切换为 CUDA 版本的 GPU 后只需 3 分钟。底层正是 GPGPU 让每个线程负责一个卷积窗口的部分乘加运算。

药物分子动力学模拟

阶段:科学计算与仿真

解决的问题:分子动力学需要反复计算大量原子间的相互作用力,每个原子对之间相互独立,天然适合并行。OpenCL 让同一套模拟程序能在不同厂商的 GPU 集群上运行,解决跨平台适配难题。

具体案例:药企在筛选候选分子时,需要在微秒级时间尺度上模拟数万个原子的运动轨迹。使用 OpenCL 内核在 GPU 上并行计算非键相互作用,模拟速度提升 20 倍以上,使虚拟筛选从“不可能”变为日常工作流。

实时渲染中的物理特效

阶段:渲染后处理与游戏物理

解决的问题:游戏中上万颗粒子(火焰、烟雾、水流)需要每帧更新位置与速度,若放在 CPU 上计算,渲染线程会被拖垮。使用计算着色器在 GPU 内部直接完成粒子更新,无需 CPU 参与,也无需数据回传。

具体案例:某 3A 游戏项目在 DirectX 12 中用计算着色器处理 50000 粒子的流体模拟,每帧粒子更新压缩到 0.8ms 以内,同时算力可继续用于下一帧的渲染。

Python 演示

python33 行
import numpy as npimport time def vector_add_serial(a, b):    """模拟 CPU 串行执行:逐个元素相加,速度慢"""    c = np.zeros_like(a)    for i in range(len(a)):        c[i] = a[i] + b[i]    return c def vector_add_parallel(a, b):    """模拟 GPU 并行执行:numpy 一次完成百万级元素相加"""    return a + b def main():    N = 1_000_000    a = np.random.rand(N).astype(np.float32)    b = np.random.rand(N).astype(np.float32)     t0 = time.perf_counter()    c1 = vector_add_serial(a, b)    t1 = time.perf_counter()     c2 = vector_add_parallel(a, b)    t2 = time.perf_counter()     print(f"串行耗时: {t1 - t0:.3f} 秒")    print(f"并行耗时: {t2 - t1:.3f} 秒")    print(f"加速比: {(t1 - t0) / (t2 - t1):.2f} 倍")    print("结果一致:", np.allclose(c1, c2)) if __name__ == "__main__":    main()
目的:用 Python 模拟 GPGPU 的核心思想——同样一个向量相加任务,串行逐元素执行与“并行”一次性执行的时间差异,帮助理解 GPU 为何在并行任务上胜出。
  1. for i in range(len(a))模拟 CPU 逐个元素执行计算,对应串行瓶颈。
  2. return a + bnumpy 底层用 SIMD 并行一次处理大量元素,对应 GPU 并行内核。
  3. time.perf_counter()精确计时,对比串行与并行的真实耗时差异。
  4. np.allclose(c1, c2)验证并行计算结果与串行结果完全一致,体现“并行不改逻辑,只改速度”。
  5. N = 1_000_000百万级元素规模,让并行优势充分体现。

常见误区

误区我听说 GPU 核心多、算力高,那我把所有程序放 GPU 上跑肯定更快,干脆连操作系统也让 GPU 管算了。

正解GPU 只在“大量可并行简单任务”上才有优势。如果任务是强串行的(比如每一步都依赖上一步结果),GPU 那上万核心大部分只能闲置;此时 CPU 凭借高主频和强单核反而更胜一筹。Amdahl 定律明确说明:串行比例决定了加速上限,雇一万帮厨也救不了一道只能由主厨亲手完成的菜。

为何常见:GPU 厂商宣传的 TFLOPS 算力指标让人误以为“算力高=什么都快”,忽略了并行适配度这一前提。

误区每次听人聊 GPGPU 都是在讲 CUDA,那是不是学会了 CUDA 就等于搞懂了全部 GPGPU?别的都不需要看了。

正解GPGPU 是“用 GPU 做通用计算”这一技术方向的总称,CUDA 只是 NVIDIA 提供的专有实现。OpenCL、DirectCompute、Vulkan、Metal 都是 GPGPU 的落地框架。CUDA 生态最成熟,但 OpenCL 跨平台更广。学 CUDA 是很好的入口,但把它等同于 GPGPU 就一叶障目了。

为何常见:NVIDIA 在深度学习时代市场份额太高,CUDA 曝光率远超其他框架,让新手产生了路径依赖的错觉。

误区GPU 显存那么大,我直接把数据丢给 GPU 就行,反正显存比内存还快,不用管中间过程。

正解主存与显存是两块独立物理空间,中间隔着 PCIe 总线。每一轮 GPGPU 任务都要经历“主存→显存→计算→回拷”的循环,PCIe 带宽远低于显存内部带宽。如果任务只运行几毫秒却搬运几 GB 数据,耗时全耗在传菜窗口上了。异步复制与双缓冲正是为此而生。

为何常见:新手通常只盯着 kernel 执行时间,忽略整体流程;GPU 本身拥有更好的内存带宽,让人误以为整个链路都快。

误区GPU 能同时跑几万个线程,那我直接开一亿线程,并行度拉满,性能一定最大化。

正解GPU 以 warp 为单位调度,硬件能驻留的线程数有上限。线程开得远超硬件槽位时,调度器反复切换、寄存器被过度瓜分、内存请求排队,反而拖慢速度。正确做法是依据 SM 数和每 SM 槽位计算合理线程规模,并通过占用率曲线找到甜点区。

为何常见:初学并行编程时容易从“分工越细越好”的直觉出发,低估了硬件资源的物理限制。

面试题

Q1. 关于 CPU 与 GPU 的核心差异,下列哪项正确?

单选题 · easy

  • A.CPU 拥有大规模简单运算核心
  • B.GPU 擅长海量并行简单运算← 正确答案
  • C.GPU 核心数量通常少于 CPU
  • D.CPU 一定慢于任何 GPU 任务

答案:B。CPU 用少量复杂核心追求单核极致,GPU 用海量简单核心追求并行吞吐。GPU 核心数量远多于 CPU(B选项正确),但“一定慢于”是绝对化错误。GPU 在串行任务上反而不如 CPU。

Q2. GPGPU 中的“通用”最准确的理解是什么?

单选题 · easy

  • A.GPU 能运行所有 Python 程序
  • B.GPU 可执行图形以外通用计算← 正确答案
  • C.GPU 可以取代 CPU 管理内存
  • D.GPU 只服务科学计算超算

答案:B。GPGPU 的“通用”指 GPU 不只做渲染,还承接 AI、仿真等通用计算。它不是 Python 解释器,也不能自主管理内存,更不局限于超算场景,因此 B 正确。

Q3. 某任务 60% 可并行,其余必须串行。按 Amdahl 定律,无论增加到多少核,理论加速比上限约为多少?

单选题 · mid

  • A.1.6 倍
  • B.2.5 倍← 正确答案
  • C.3.3 倍
  • D.5.0 倍

答案:B。加速比上限 $S = \frac{1}{1-p} = \frac{1}{0.4} = 2.5$ 倍。串行部分 40% 成为天花板,即使核数无穷多,也无法突破 2.5 倍。这正说明并行化前要先提高可并行比例 $p$。

Q4. 关于 SIMT 执行模型,下列哪些说法正确?

多选题 · mid

  • A.同一 warp 线程执行相同指令← 正确答案
  • B.线程可以处理不同的数据← 正确答案
  • C.分支分歧不会产生性能损耗
  • D.线程粒度小于 CPU 线程← 正确答案

答案:ABD。SIMT 是“单指令多线程”:一个 warp 内所有线程执行同一条指令(A),但各自处理不同数据(B),线程粒度远小于 CPU 线程(D)。分支分歧会导致 warp 内部分线程等待,必然有损耗(C 错)。

Q5. 在 CUDA kernel 中,已知 blockIdx.x 等于 2,blockDim.x 等于 256,threadIdx.x 等于 10,全局线程 ID 是多少?

单选题 · mid

  • A.512
  • B.522← 正确答案
  • C.768
  • D.256

答案:B。全局 ID = blockIdx.x × blockDim.x + threadIdx.x = 2 × 256 + 10 = 522。A 选项漏加 threadIdx,C 选项误把 3 当块数,D 选项只取了 blockDim。定位公式是 CUDA 编程的基石。

Q6. OpenCL 中用于描述多维工作调度空间的概念是?

单选题 · mid

  • A.NDRange 索引空间← 正确答案
  • B.顶点缓冲区域
  • C.纹理映射单元
  • D.像素着色管线

答案:A。NDRange(N-Dimensional Range)是 OpenCL 的 N 维索引空间,定义了 work-item 在多维网格中的排列方式。顶点缓冲、纹理映射、像素着色都属于图形渲染概念,与工作调度无关。

Q7. 下列哪些做法能有效提升 GPU 内存带宽利用率?

多选题 · hard

  • A.让同一 warp 访问连续内存地址← 正确答案
  • B.让线程随机跳跃访问显存
  • C.提高占用率但严格控制寄存器数← 正确答案
  • D.尽量让每个线程单独复制大块数据

答案:AC。合并访问让 32 个线程的内存请求合并成一次事务(A),提高占用率能隐藏延迟但寄存器不能超限(C)。随机跳跃访问会拆散事务(B 错),每线程复制大块数据反而加重内存压力(D 错)。

Q8. CPU 在 GPGPU 任务中的主要职责是什么?

单选题 · easy

  • A.逐元素执行所有并行计算
  • B.初始化数据并启动内核调度← 正确答案
  • C.直接改写 GPU 内部寄存器
  • D.替代 GPU 完成矩阵乘法

答案:B。CPU 是主机端,负责分配内存、拷贝数据、启动 kernel、回收资源;GPU 是设备端,执行海量并行计算。CPU 直接改写寄存器或替代 GPU 做矩阵乘法都不符合协作模型。

Q9. 关于主机与设备间的数据搬运,下列说法正确的是?

多选题 · hard

  • A.搬运耗时受 PCIe 带宽限制← 正确答案
  • B.异步拷贝可使传输与计算重叠← 正确答案
  • C.小批量高频率搬运通常更高效
  • D.显存足够大时搬运不算瓶颈

答案:AB。传输时间 = 数据量 / 总线带宽,带宽是硬限制(A);异步拷贝让数据传输与 GPU 计算并行(B)。小批量高频率传输会放大启动开销(C 错),显存再大也解决不了 PCIe 带宽瓶颈(D 错)。

Q10. 计算着色器与传统渲染着色器最大的区别是什么?

单选题 · mid

  • A.只能在顶点阶段运行
  • B.不依赖渲染管线独立计算← 正确答案
  • C.只输出固定颜色像素
  • D.必须配合光栅化器执行

答案:B。计算着色器独立于渲染管线运行,不参与顶点变换或像素着色,可执行任意通用计算(B 正确)。A、C、D 都是传统渲染着色器的特征。它是 GPGPU 走向图形学融合的标志性产物。