欢迎来到并行计算的真实世界。如果说传统的串行程序是一辆追求极致单步响应速度的跑车那么我们即将学习的异构数据并行程序就是一支能够同时处理海量任务的超级舰队。本章是引导你从“串行思维”跨越到“并行思维”的破冰之旅。在这里我们将彻底打破传统围绕for循环构建的编程习惯带你从底层硬件架构出发去理解 GPU 天生为“吞吐量”而生的设计哲学。通过本章的学习你将掌握以下核心内容架构认知深刻理解 CPU低延迟导向与 GPU高吞吐量导向在底层硅片设计和调度策略上的根本差异。思维转换区分任务并行与数据并行理解为何图像处理、科学模拟与矩阵运算是 GPU 的天然主场。核心抽象掌握 CUDA 编程模型中“网格Grid-线程块Block-线程Thread”的三级组织架构并熟练运用黄金公式计算全局唯一的线程索引。代码实战掌握异构内存管理 API从零开始编写并运行你的第一个真正的 GPU 并行程序——向量加法。工程规范了解nvcc编译管线掌握应对 GPU 异步执行特性的标准错误检查Error Checking宏定义与调试机制。引言异构计算架构简介在进入具体的并行代码编写之前我们需要先建立一个核心概念异构计算。传统的计算机程序主要在中央处理器CPU上单线程或多线程运行这是一种同构的计算模式。然而随着深度学习、科学模拟和图像处理等应用对算力需求的呈指数级增长仅仅依靠提升单核 CPU 的时钟频率早已触及物理极限即“功耗墙”。异构计算由此成为主流它指的是在一个系统中协同使用不同类型的处理器或计算单元如 CPU、GPU、FPGA 或专用的 AI 加速器遵循“让最合适的硬件做最合适的工作”的原则。在现代异构数据并行计算模型例如 CUDA 或 OpenCL中程序架构通常围绕两个基本实体展开主机Host通常指 CPU 及其所在的系统主存RAM。它是整个程序的“指挥官”擅长处理复杂的控制逻辑、操作系统交互、网络 I/O 以及任务的调度。设备Device通常指 GPU图形处理器及其独立的显存VRAM或其他加速卡。它是程序的“主力计算军团”专门负责处理计算密集型、高度并行化的任务。一个典型的异构计算程序的生命周期通常包含以下几个步骤主机在系统内存中初始化数据。主机将数据通过 PCIe 总线拷贝到设备的显存中。主机向设备发送指令启动设备上的成千上万个独立线程进行并行计算。计算完成后主机将结果从设备显存拷贝回系统内存。CPU 与 GPU 的设计哲学对比延迟导向 vs. 吞吐量导向要深刻理解为什么 GPU 在数据并行任务上拥有远超 CPU 的性能我们需要剖析它们底层设计哲学的根本分歧。它们天生是为了解决完全不同类型的计算问题而诞生的。CPU致力于降低延迟Latency-OrientedCPU 被设计为尽可能快地完成单一的、复杂的指令流。现实中的通用程序充满了复杂的逻辑跳转如海量的if-else分支和不可预测的数据内存访问。为了让串行代码跑得飞快CPU 芯片内部的大部分晶体管并没有用于计算而是被用来构建巨大的多级缓存L1/L2/L3 Cache以及极其复杂的控制逻辑部件如分支预测器、乱序执行单元。因此CPU 只有少量的算术逻辑单元ALU但它的每个核心都极其强大能够以极低的延迟响应复杂的串行任务。GPU致力于最大化吞吐量Throughput-OrientedGPU 的设计初衷是处理图形渲染——这要求对屏幕上数以百万计的独立像素进行极其相似的矩阵或向量数学运算。这是一种极其纯粹的“数据并行”任务。基于这一需求GPU 的设计者无情地砍掉了复杂的控制逻辑和庞大的单线程缓存将绝大部分宝贵的硅片面积投入到了计算核心ALU上。一个现代 GPU 可能拥有数以千计甚至上万个简单的流处理器。当 GPU 中的某一组线程在读取显存遇到延迟时GPU 的控制策略并非像 CPU 那样利用巨大的缓存去弥补而是零开销地瞬间切换到另一组已经准备好数据的线程上继续执行计算。通过这种用海量线程来“隐藏”内存访问延迟的策略GPU 实现了极为惊人的整体浮点运算吞吐量。这种硬件架构的根本差异决定了我们的编程范式必须发生巨大的转变在编写 CPU 程序时我们致力于优化单一线程的执行效率并减少内存延迟而在编写 GPU 异构程序时我们的首要任务是寻找并最大程度地暴露程序中的数据并行性以提供足够庞大的线程规模去喂饱那成千上万个嗷嗷待哺的计算核心。第一章数据并行性基础当我们谈论“并行计算”时我们通常指的是同时执行多个计算操作。然而根据我们拆分工作的方式并行计算主要可以划分为两大基本范式任务并行和数据并行。理解这两者的区别是编写高效异构程序的关键。任务并行任务并行是指将一个复杂的应用程序分解为多个不同的功能模块或任务并将它们分配给不同的处理单元同时执行。生活中的类比想象一个厨房在准备一场晚宴。主厨负责烤牛排副手负责切沙拉糕点师负责烤甜点另一个人负责布置餐桌。每个人都在做完全不同的事情使用不同的工具但他们共同推进了“准备晚宴”这个大目标。计算机中的体现现代多核 CPU 非常擅长任务并行。在一个现代视频游戏中CPU 可能用一个核心处理物理引擎碰撞Physics另一个核心处理人工智能逻辑AI第三个核心负责网络通信Network I/O第四个核心负责音频解码Audio。这些任务逻辑复杂、分支众多且彼此差异巨大。数据并行数据并行则是指针对一个庞大的数据集将相同的计算操作或指令同时应用于数据集合的不同部分。生活中的类比想象有一万份单项选择题考卷需要批改。如果你有一百位老师最快的方法不是让每个老师负责不同的改卷步骤而是给每位老师分配一百份考卷并且所有老师都拿着同一份标准答案相同的指令同时进行批改不同的数据。计算机中的体现这正是 GPU 的统治领域。GPU 拥有成千上万个轻量级核心它们被设计为在同一时刻对内存中不同位置的数据执行一模一样的数学运算。这种模式被称为SIMD单指令多数据流Single Instruction, Multiple Data或在 CUDA 架构中更准确地称为SIMT单指令多线程Single Instruction, Multiple Threads。为什么图像处理和矩阵运算天然适合数据并行让我们看两个经典的例子这也是我们后续编写代码时最常遇到的场景1. 图像处理如彩色转灰度一张 4K 分辨率的图片包含超过 800 万个像素。要将这张彩色图片转换为黑白灰度图我们需要对每一个像素的红R、绿G、蓝B通道执行相同的数学公式在这 800 万次计算中像素 A 的计算结果绝对不会影响像素 B 的计算结果。它们是完全独立的。这就是完美的数据并行场景。如果在一个拥有 8000 个核心的 GPU 上运行GPU 可以瞬间抓取 8000 个像素在一个时钟周期或极短的时间内同时完成这 8000 次相同的公式计算。2. 矩阵加法 ()假设我们有两个的大型矩阵和相加。我们需要执行一百万次独立的加法操作计算。同样每一个元素的相加操作与其它元素完全无关非常适合分配给成千上万个 GPU 线程去并发执行。核心总结异构编程的本质就是剥离程序中复杂的控制逻辑交给 CPU 做任务并行剥离出那些“数据量巨大、计算规则统一且相互独立”的循环体将它们转化为成千上万个线程交给 GPU 做数据并行。左边就是CPU的任务并行右边就是GPU的数据并行。第二章CUDA 编程模型初探网格与线程块当我们说 GPU 可以同时运行数以万计的线程时你可能会问程序员该如何管理这上万个线程如果这上万个线程都在执行同一段代码它们怎么知道自己该去处理内存中的哪一部分数据为了解决这个问题NVIDIA 的 CUDA 编程模型提出了一种高度结构化、可扩展的线程组织层次。这种层次结构不仅方便程序员管理也完美契合了底层 GPU 硬件的物理架构。1. 核函数Kernel Function在 CUDA 中我们在 CPU主机上调用而在 GPU设备上执行的函数被称为核函数Kernel。 在 C/C 代码中我们通过在函数定义前加上__global__限定符来声明它。当主机启动一个核函数时GPU 会生成成千上万个线程来并发执行这段代码。2. 三级线程层级结构Grid - Block - Thread为了组织这些海量的线程CUDA 采用了“网格-线程块-线程”的三级结构线程Thread最基本的执行单元。每个线程执行核函数的一个实例拥有自己的指令地址寄存器和局部变量。线程块Block由一组并发执行的线程组成。这是一个极其重要的概念因为同一个 Block 内的线程可以通过一块极速的“共享内存Shared Memory”进行数据交换并且可以通过内置函数进行同步Synchronization。网格Grid由一组执行相同核函数的 Block 组成。一个核函数的启动对应一个 Grid。不同 Block 之间的线程通常是相互独立的无法进行轻量级的同步。通俗的比喻假设我们要建一栋摩天大楼执行一个 Kernel/Grid。我们将工程分包给多个施工队Block。每个施工队里有许多工人Thread。同一个施工队的工人可以轻松地互相递砖头、喊话协调共享内存与同步但不同施工队之间距离太远各自独立干活互不干扰。3. 我是谁我在哪内置坐标变量因为所有的线程都在执行同一段核函数代码SPMD单程序多数据线程必须依靠自己独特的“坐标”来计算出自己应该读取哪里的数据。CUDA 为每个线程提供了几个极其关键的内置变量Built-in Variables它们由硬件直接填充threadIdx当前线程在其所属 Block 中的坐标索引。blockIdx当前线程所属的 Block 在整个 Grid 中的坐标索引。blockDim一个 Block 中包含的线程维度大小例如一个 Block 里有多少个线程。gridDim一个 Grid 中包含的 Block 维度大小。注这些变量都是内置的dim3向量类型包含.x,.y,.z三个分量这意味着你可以将线程组织成一维、二维或三维的结构以完美适配向量、矩阵或空间体素的数据计算。4. 黄金公式计算全局唯一索引对于最常见的一维数据处理例如我们马上要在下一节写的向量加法我们如何将两级坐标Block 索引和 Thread 索引展平计算出该线程在所有线程中的全局唯一 ID 呢这就引出了并行编程中最著名的一维索引公式计算出在当前 Block 之前已经有多少个线程存在了。加上当前线程在自己 Block 内的偏移量。最终得到的就是这个线程用来访问一维数组如Array[i]的全局索引。第三章第一个并行程序向量加法假设我们有两个包含 100 万个浮点数的一维数组向量和我们需要将它们对应位置的元素相加并将结果存储在数组中。即。如果是在传统的 CPU C 程序中我们会写一个简单的for循环void vecAdd_CPU(float *A, float *B, float *C, int N) { for (int i 0; i N; i) { C[i] A[i] B[i]; } }但在 GPU 编程中我们要彻底抛弃这种串行循环的思维。我们要让成千上万个线程同时去执行加法操作。一个完整的 CUDA 程序分为两部分设备代码Device Code即核函数和主机代码Host Code负责统筹管理。1. 编写设备代码核函数 (Kernel)这是真正在 GPU 上成百上千个核心中并发执行的代码。我们用__global__关键字来修饰它。// CUDA 核函数每个线程计算向量中的一个元素 __global__ void vecAdd_GPU(float *A, float *B, float *C, int N) { // 黄金公式计算当前线程的全局唯一索引 i int i blockDim.x * blockIdx.x threadIdx.x; // 边界检查防止线程索引超出数组实际长度 if (i N) { C[i] A[i] B[i]; // 剥离了 for 循环每个线程只做一次加法 } }核心思想转变注意看这里没有for循环因为循环被转化为空间上的并行了。每一个启动的线程都会通过计算自己的索引i精确地找到自己负责的那个元素完成一次加法后就退出。2. 编写主机代码统筹帷幄 (Host)GPU 是个纯粹的运算机器它自己不会去内存里抓数据。因此我们需要 CPU主机来为它准备好一切。一个典型的 CUDA 主机代码包含四个标准步骤这套“分配显存 - 拷贝进去 - 执行计算 - 拷贝出来 - 释放显存”的流程是几乎所有异构计算程序的不变基石。#include stdio.h int main() { int N 1000000; // 100万个元素 size_t size N * sizeof(float); // [主机端] 分配内存并初始化 A 和 B (代码略) float *h_A ...; float *h_B ...; float *h_C ...; // --------------------------------------------------------- // 步骤 1: [设备端] 在 GPU 显存中分配空间 (cudaMalloc) // --------------------------------------------------------- float *d_A, *d_B, *d_C; cudaMalloc((void**)d_A, size); cudaMalloc((void**)d_B, size); cudaMalloc((void**)d_C, size); // --------------------------------------------------------- // 步骤 2: 数据传输 - 将数据从 CPU 拷贝到 GPU (cudaMemcpy) // --------------------------------------------------------- cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice); // --------------------------------------------------------- // 步骤 3: 启动核函数 - 定义网格和线程块大小 // --------------------------------------------------------- int threadsPerBlock 256; // 每个 Block 包含 256 个线程 // 计算需要多少个 Block向上取整确保覆盖所有元素 int blocksPerGrid (N threadsPerBlock - 1) / threadsPerBlock; // 核心语法 Grid大小, Block大小 vecAdd_GPUblocksPerGrid, threadsPerBlock(d_A, d_B, d_C, N); // --------------------------------------------------------- // 步骤 4: 数据回传 - 将计算结果从 GPU 拷贝回 CPU // --------------------------------------------------------- cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost); // [设备端] 释放显存 cudaFree(d_A); cudaFree(d_B); cudaFree(d_C); return 0; }第四章程序的编译与错误检查在传统的 C/C 开发中我们习惯使用 GCC、Clang 或 MSVC 等编译器。但当我们编写了包含主机代码CPU和设备代码GPU的混合源文件通常以.cu为扩展名时传统的编译器就“傻眼”了因为它根本不认识__global__或者...这些奇怪的语法。这时候我们就需要引入专门的异构编译工具链。1. NVCC 编译器的工作原理在 CUDA 生态中标准编译器是nvcc(NVIDIA CUDA Compiler)。 准确地说nvcc更像是一个“编译器驱动程序Driver”。它的核心工作逻辑是代码分离Separation分离nvcc首先解析.cu文件将其中的主机代码标准的 C/C和设备代码CUDA 扩展部分剥离开来。分发主机代码将纯主机代码移交给系统自带的宿主编译器比如 Linux 下的 GCCWindows 下的 MSVC进行编译。编译设备代码调用 NVIDIA 自己的底层编译器组件将设备代码编译为 PTX一种中间汇编指令或特定 GPU 架构的机器码SASS。链接最后将编译好的主机目标文件和设备目标文件链接在一起生成最终的可执行文件。2. 异步执行带来的“沉默地雷”在异构计算中有一个极易让初学者崩溃的特性核函数的启动是完全异步的Asynchronous。当你执行vecAdd_GPUgrid, block(...)时CPU 只是向 GPU 发送了一个“开始干活”的指令然后 CPU 会瞬间返回继续执行下一行代码根本不会等待 GPU 把活干完。这种设计是为了让 CPU 和 GPU 能够并行工作重叠计算最大化系统吞吐量。但这带来了一个致命问题如果 GPU 在执行计算时崩溃了比如数组越界访问CPU 在调用核函数的那一瞬间是毫无察觉的。程序不会在...这一行报错而是可能在后续尝试拷贝数据回 CPU 时突然崩溃或者直接静默失败输出一堆乱码。3. 编写标准的错误检查宏 (Error Checking Macro)为了捕获这些“沉默的地雷”我们必须养成在每一个 CUDA API 调用和核函数执行后进行严格错误检查的习惯。在工程实践中我们通常会定义一个统一的宏Macro来包装这些调用。#include stdio.h // 宏定义捕获并打印 CUDA 错误信息 #define CUDA_CHECK(call) \ { \ const cudaError_t error call; \ if (error ! cudaSuccess) \ { \ fprintf(stderr, Error: %s:%d, , __FILE__, __LINE__); \ fprintf(stderr, code: %d, reason: %s\n, error, \ cudaGetErrorString(error)); \ exit(1); \ } \ } // 检查核函数异步执行的错误 #define CHECK_KERNEL_ERRORS() \ { \ /* 捕获核函数启动时的同步错误如无效的 block 参数*/ \ CUDA_CHECK(cudaGetLastError()); \ /* 强制 CPU 等待 GPU 执行完毕捕获执行过程中的异步错误如越界*/ \ CUDA_CHECK(cudaDeviceSynchronize()); \ }如何使用它我们将 2.4 节中的代码用这个宏包装起来它就变成了符合工程标准的安全代码// 1. 包装内存分配和拷贝 CUDA_CHECK(cudaMalloc((void**)d_A, size)); CUDA_CHECK(cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice)); // 2. 启动核函数注意核函数启动本身不能直接被宏包装 vecAdd_GPUblocksPerGrid, threadsPerBlock(d_A, d_B, d_C, N); // 3. 立即捕获核函数可能产生的错误 CHECK_KERNEL_ERRORS(); // 4. 包装数据回传 CUDA_CHECK(cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost));通过引入cudaDeviceSynchronize()我们强制 CPU 停下来等待 GPU 汇报工作结果。如果在计算过程中发生了内存越界系统会精准地打印出错误发生的文件名、行号以及具体的错误原因比如an illegal memory access was encountered关于API的额外说明API的全称是Application Programming Interface应用程序编程接口。在 CUDA 中的意义NVIDIA 的工程师写好了底层与 GPU 硬件沟通的极其复杂的驱动代码然后将其包装成了一个个简单的函数如cudaMalloc、cudaMemcpy。这些函数就是 CUDA API。我们调用这些 API就是在给 GPU 下达指令。在 CUDA C/C 中几乎所有的 CUDA API 调用除了核函数...本身都会返回一个状态码用来告诉我们“后厨的菜做成功了没有”。这个返回值的专门类型叫做cudaError_t。它本质上是一个枚举enum类型如果一切顺利它的值会是cudaSuccess也就是数字 0。做个全景示范#include stdio.h #include cuda_runtime.h int main() { int V 1000; // 假设我们要分配 1000 个整数 // 1. 计算需要的总字节数 // 整数个数 乘以 每个整数占用的字节数通常是4字节 size_t size V * sizeof(int); // 2. 声明设备端指针这只是一张暂时空白的房卡 int *d_array; // 3. 声明用于接收 API 状态的 ERR 变量 cudaError_t ERR; // 4. 调用 API 分配内存并接收返回值 // 注意一定要传入指针的地址 d_array并强制转换为 (void**) ERR cudaMalloc((void**)d_array, size); // 5. 检查 API 是否执行成功 if (ERR ! cudaSuccess) { // 如果失败打印出具体的错误原因 printf(GPU内存分配失败原因: %s\n, cudaGetErrorString(ERR)); return -1; // 退出程序 } else { printf(成功在 GPU 上分配了 %d 个整数的内存\n, V); } // ... 后续计算 ... // 释放内存 cudaFree(d_array); return 0; }概率习题习题 1基础线程到数据的映射题目如果我们要让网格Grid中的每个线程计算向量加法的一个输出元素将线程/块索引映射到数据索引i的 C 语言表达式是什么习题 2步长与多元素处理进阶映射题目假设我们要修改程序让每个线程计算向量加法中的两个相邻元素。将线程/块索引映射到该线程处理的第一个数据元素索引i的表达式是什么习题 3网格与线程边界计算题目对于向量加法假设向量的长度为 8000且每个线程计算一个输出元素。如果你设定线程块大小Block Size为 1024 个线程那么整个网格Grid中总共会被启动多少个线程习题 4CUDA 内存分配的 C 语言陷阱题目假设我们有一个指向浮点数的指针float *A_d;。以下哪一行代码是正确地在设备GPU上分配内存的语法A.cudaMalloc((void*)A_d, size);B.cudaMalloc(A_d, size);C.cudaMalloc((void**)A_d, size);D.cudaMalloc((void**)A_d, size);习题 5数据传输方向题目假设A_h是主机端CPU数组A_d是设备端GPU数组两者都已经分配了 3000 字节的空间。现在需要将数据从 CPU 传输到 GPU正确的cudaMemcpy语句应该怎么写解答:1.答案i blockIdx.x * blockDim.x threadIdx.x;思路点拨这是贯穿整个 CUDA 编程的最核心公式。blockIdx.x * blockDim.x计算的是在当前线程块Block之前所有排在前面的线程块所包含的线程总数。加上当前线程在自己所属 Block 内部的偏移量threadIdx.x就能得到该线程在整个网格中的全局绝对位置。2.答案i (blockIdx.x * blockDim.x threadIdx.x) * 2;思路点拨既然每个线程要处理 2 个元素这就意味着数据数组在逻辑上被按每 2 个元素分成了多个“组”。线程的全局唯一 ID 依然是blockIdx.x * blockDim.x threadIdx.x这代表它是第几个“组”的负责人。为了找到它负责的这组数据的起始内存地址只需将它的全局线程 ID 乘以 2 即可。3.答案8192 个线程。思路点拨这是一道极易踩坑的经典题。很多初学者会回答 8000。但硬件启动线程块时只能启动完整的块。 计算所需 Block 数量的公式是向上取整(8000 / 1024) 8 个 Block。 既然启动了 8 个完整的 Block且每个 Block 有 1024 个线程那么总启动线程数就是 8 * 1024 8192 个。这也是为什么我们在 2.4 节的核函数中必须要写if (i N)这个边界检查的原因因为会有 192 个多余的线程被启动如果不加限制它们就会导致内存越界访问4.答案D思路点拨这一题考察的是 C 语言指针的基础以及 API 的设计。cudaMalloc函数的任务是去 GPU 上开辟一块空间并且将这块空间的起始地址写入到你提供的指针变量中。 因为函数内部要修改你传入的指针A_d的值它就必须知道这个指针本身在内存中的地址因此必须传入指针的地址A_d这是一个二级指针。同时为了兼容各种数据类型API 定义为void**所以需要进行强制类型转换。5.答案cudaMemcpy(A_d, A_h, 3000, cudaMemcpyHostToDevice);思路点拨重点在于参数的顺序(目标地址, 源地址, 字节数, 传输方向常量)。记住它和 C 标准库里的memcpy是一脉相承的永远是目标在前源在后就像赋值语句A_d A_h一样。同时必须显式地指定cudaMemcpyHostToDevice告诉系统底层的 DMA 控制器数据流向。实战习题考虑以下 CUDA 核函数以及调用它的相应主机函数__global__ void foo_kernel(float* a, float* b, unsigned int N) { unsigned int i blockIdx.x * blockDim.x threadIdx.x; if (i N) { b[i] 2.7f * a[i] - 4.3f; } } void foo(float* a_d, float* b_d) { unsigned int N 200000; // 执行配置(N 128 - 1)/128, 128 foo_kernel(N 128 - 1) / 128, 128(a_d, b_d, N); }问题a.每个线程块Block中有多少个线程b.整个网格Grid中总共有多少个线程c.整个网格中总共有多少个线程块d.有多少个线程执行了计算if条件语句即判断i N的代码e.有多少个线程执行了if语句内部的计算代码a. 每个线程块Block中有多少个线程答案128解析核函数的执行配置语法是网格大小, 线程块大小。在主机代码foo中第二个参数显式指定了 128这意味着每一个 Block 被硬编码为包含 128 个线程。b. 整个网格Grid中总共有多少个线程答案200,064解析网格中的总线程数 线程块总数 × 每个线程块的线程数。根据 c 题的计算网格中启动了 1563 个线程块因此总线程数为 1563 × 128 200,064 个线程。c. 整个网格中总共有多少个线程块答案1,563解析线程块的数量由公式(N 128 - 1) / 128决定。这是一种在 C 语言整数除法中实现“向上取整”的经典编程技巧。 代入 N 200000 进行计算(200000 127) / 128 200127 / 128。 在 C/C 的整数除法中小数部分会被直接截断舍弃因此 200127 / 128 1563。这意味着 GPU 实际启动了 1563 个完整的线程块。d. 有多少个线程执行了if条件判断的代码答案200,064解析这是一个极易踩坑的核心概念硬件在分配和启动线程时只能以完整的 Block 为最小单位进行分配。既然系统最终启动了 1563 个 Block那么包含在其中的全部 200,064 个线程都会被真实地分配到流处理器上并开始运行。因此所有被启动的线程都会首先计算出自己的全局索引i并执行if (i N)这行条件评估代码。e. 有多少个线程执行了if语句内部的计算代码答案200,000解析只有满足条件即全局索引i的范围在 0 到 199,999 之间的真实有效线程才会进入if代码块内部执行实际的内存读写和浮点数运算。超出这个范围的 64 个多余线程200,064 - 200,000在条件判断为假后直接越过了计算代码并结束生命周期。这就是我们在核函数中必须编写边界保护代码的根本原因。
大规模并发处理器程序设计(PMPP)讲解(CUDA架构)：第二期：异构数据并行计算
欢迎来到并行计算的真实世界。如果说传统的串行程序是一辆追求极致单步响应速度的跑车那么我们即将学习的异构数据并行程序就是一支能够同时处理海量任务的超级舰队。本章是引导你从“串行思维”跨越到“并行思维”的破冰之旅。在这里我们将彻底打破传统围绕for循环构建的编程习惯带你从底层硬件架构出发去理解 GPU 天生为“吞吐量”而生的设计哲学。通过本章的学习你将掌握以下核心内容架构认知深刻理解 CPU低延迟导向与 GPU高吞吐量导向在底层硅片设计和调度策略上的根本差异。思维转换区分任务并行与数据并行理解为何图像处理、科学模拟与矩阵运算是 GPU 的天然主场。核心抽象掌握 CUDA 编程模型中“网格Grid-线程块Block-线程Thread”的三级组织架构并熟练运用黄金公式计算全局唯一的线程索引。代码实战掌握异构内存管理 API从零开始编写并运行你的第一个真正的 GPU 并行程序——向量加法。工程规范了解nvcc编译管线掌握应对 GPU 异步执行特性的标准错误检查Error Checking宏定义与调试机制。引言异构计算架构简介在进入具体的并行代码编写之前我们需要先建立一个核心概念异构计算。传统的计算机程序主要在中央处理器CPU上单线程或多线程运行这是一种同构的计算模式。然而随着深度学习、科学模拟和图像处理等应用对算力需求的呈指数级增长仅仅依靠提升单核 CPU 的时钟频率早已触及物理极限即“功耗墙”。异构计算由此成为主流它指的是在一个系统中协同使用不同类型的处理器或计算单元如 CPU、GPU、FPGA 或专用的 AI 加速器遵循“让最合适的硬件做最合适的工作”的原则。在现代异构数据并行计算模型例如 CUDA 或 OpenCL中程序架构通常围绕两个基本实体展开主机Host通常指 CPU 及其所在的系统主存RAM。它是整个程序的“指挥官”擅长处理复杂的控制逻辑、操作系统交互、网络 I/O 以及任务的调度。设备Device通常指 GPU图形处理器及其独立的显存VRAM或其他加速卡。它是程序的“主力计算军团”专门负责处理计算密集型、高度并行化的任务。一个典型的异构计算程序的生命周期通常包含以下几个步骤主机在系统内存中初始化数据。主机将数据通过 PCIe 总线拷贝到设备的显存中。主机向设备发送指令启动设备上的成千上万个独立线程进行并行计算。计算完成后主机将结果从设备显存拷贝回系统内存。CPU 与 GPU 的设计哲学对比延迟导向 vs. 吞吐量导向要深刻理解为什么 GPU 在数据并行任务上拥有远超 CPU 的性能我们需要剖析它们底层设计哲学的根本分歧。它们天生是为了解决完全不同类型的计算问题而诞生的。CPU致力于降低延迟Latency-OrientedCPU 被设计为尽可能快地完成单一的、复杂的指令流。现实中的通用程序充满了复杂的逻辑跳转如海量的if-else分支和不可预测的数据内存访问。为了让串行代码跑得飞快CPU 芯片内部的大部分晶体管并没有用于计算而是被用来构建巨大的多级缓存L1/L2/L3 Cache以及极其复杂的控制逻辑部件如分支预测器、乱序执行单元。因此CPU 只有少量的算术逻辑单元ALU但它的每个核心都极其强大能够以极低的延迟响应复杂的串行任务。GPU致力于最大化吞吐量Throughput-OrientedGPU 的设计初衷是处理图形渲染——这要求对屏幕上数以百万计的独立像素进行极其相似的矩阵或向量数学运算。这是一种极其纯粹的“数据并行”任务。基于这一需求GPU 的设计者无情地砍掉了复杂的控制逻辑和庞大的单线程缓存将绝大部分宝贵的硅片面积投入到了计算核心ALU上。一个现代 GPU 可能拥有数以千计甚至上万个简单的流处理器。当 GPU 中的某一组线程在读取显存遇到延迟时GPU 的控制策略并非像 CPU 那样利用巨大的缓存去弥补而是零开销地瞬间切换到另一组已经准备好数据的线程上继续执行计算。通过这种用海量线程来“隐藏”内存访问延迟的策略GPU 实现了极为惊人的整体浮点运算吞吐量。这种硬件架构的根本差异决定了我们的编程范式必须发生巨大的转变在编写 CPU 程序时我们致力于优化单一线程的执行效率并减少内存延迟而在编写 GPU 异构程序时我们的首要任务是寻找并最大程度地暴露程序中的数据并行性以提供足够庞大的线程规模去喂饱那成千上万个嗷嗷待哺的计算核心。第一章数据并行性基础当我们谈论“并行计算”时我们通常指的是同时执行多个计算操作。然而根据我们拆分工作的方式并行计算主要可以划分为两大基本范式任务并行和数据并行。理解这两者的区别是编写高效异构程序的关键。任务并行任务并行是指将一个复杂的应用程序分解为多个不同的功能模块或任务并将它们分配给不同的处理单元同时执行。生活中的类比想象一个厨房在准备一场晚宴。主厨负责烤牛排副手负责切沙拉糕点师负责烤甜点另一个人负责布置餐桌。每个人都在做完全不同的事情使用不同的工具但他们共同推进了“准备晚宴”这个大目标。计算机中的体现现代多核 CPU 非常擅长任务并行。在一个现代视频游戏中CPU 可能用一个核心处理物理引擎碰撞Physics另一个核心处理人工智能逻辑AI第三个核心负责网络通信Network I/O第四个核心负责音频解码Audio。这些任务逻辑复杂、分支众多且彼此差异巨大。数据并行数据并行则是指针对一个庞大的数据集将相同的计算操作或指令同时应用于数据集合的不同部分。生活中的类比想象有一万份单项选择题考卷需要批改。如果你有一百位老师最快的方法不是让每个老师负责不同的改卷步骤而是给每位老师分配一百份考卷并且所有老师都拿着同一份标准答案相同的指令同时进行批改不同的数据。计算机中的体现这正是 GPU 的统治领域。GPU 拥有成千上万个轻量级核心它们被设计为在同一时刻对内存中不同位置的数据执行一模一样的数学运算。这种模式被称为SIMD单指令多数据流Single Instruction, Multiple Data或在 CUDA 架构中更准确地称为SIMT单指令多线程Single Instruction, Multiple Threads。为什么图像处理和矩阵运算天然适合数据并行让我们看两个经典的例子这也是我们后续编写代码时最常遇到的场景1. 图像处理如彩色转灰度一张 4K 分辨率的图片包含超过 800 万个像素。要将这张彩色图片转换为黑白灰度图我们需要对每一个像素的红R、绿G、蓝B通道执行相同的数学公式在这 800 万次计算中像素 A 的计算结果绝对不会影响像素 B 的计算结果。它们是完全独立的。这就是完美的数据并行场景。如果在一个拥有 8000 个核心的 GPU 上运行GPU 可以瞬间抓取 8000 个像素在一个时钟周期或极短的时间内同时完成这 8000 次相同的公式计算。2. 矩阵加法 ()假设我们有两个的大型矩阵和相加。我们需要执行一百万次独立的加法操作计算。同样每一个元素的相加操作与其它元素完全无关非常适合分配给成千上万个 GPU 线程去并发执行。核心总结异构编程的本质就是剥离程序中复杂的控制逻辑交给 CPU 做任务并行剥离出那些“数据量巨大、计算规则统一且相互独立”的循环体将它们转化为成千上万个线程交给 GPU 做数据并行。左边就是CPU的任务并行右边就是GPU的数据并行。第二章CUDA 编程模型初探网格与线程块当我们说 GPU 可以同时运行数以万计的线程时你可能会问程序员该如何管理这上万个线程如果这上万个线程都在执行同一段代码它们怎么知道自己该去处理内存中的哪一部分数据为了解决这个问题NVIDIA 的 CUDA 编程模型提出了一种高度结构化、可扩展的线程组织层次。这种层次结构不仅方便程序员管理也完美契合了底层 GPU 硬件的物理架构。1. 核函数Kernel Function在 CUDA 中我们在 CPU主机上调用而在 GPU设备上执行的函数被称为核函数Kernel。 在 C/C 代码中我们通过在函数定义前加上__global__限定符来声明它。当主机启动一个核函数时GPU 会生成成千上万个线程来并发执行这段代码。2. 三级线程层级结构Grid - Block - Thread为了组织这些海量的线程CUDA 采用了“网格-线程块-线程”的三级结构线程Thread最基本的执行单元。每个线程执行核函数的一个实例拥有自己的指令地址寄存器和局部变量。线程块Block由一组并发执行的线程组成。这是一个极其重要的概念因为同一个 Block 内的线程可以通过一块极速的“共享内存Shared Memory”进行数据交换并且可以通过内置函数进行同步Synchronization。网格Grid由一组执行相同核函数的 Block 组成。一个核函数的启动对应一个 Grid。不同 Block 之间的线程通常是相互独立的无法进行轻量级的同步。通俗的比喻假设我们要建一栋摩天大楼执行一个 Kernel/Grid。我们将工程分包给多个施工队Block。每个施工队里有许多工人Thread。同一个施工队的工人可以轻松地互相递砖头、喊话协调共享内存与同步但不同施工队之间距离太远各自独立干活互不干扰。3. 我是谁我在哪内置坐标变量因为所有的线程都在执行同一段核函数代码SPMD单程序多数据线程必须依靠自己独特的“坐标”来计算出自己应该读取哪里的数据。CUDA 为每个线程提供了几个极其关键的内置变量Built-in Variables它们由硬件直接填充threadIdx当前线程在其所属 Block 中的坐标索引。blockIdx当前线程所属的 Block 在整个 Grid 中的坐标索引。blockDim一个 Block 中包含的线程维度大小例如一个 Block 里有多少个线程。gridDim一个 Grid 中包含的 Block 维度大小。注这些变量都是内置的dim3向量类型包含.x,.y,.z三个分量这意味着你可以将线程组织成一维、二维或三维的结构以完美适配向量、矩阵或空间体素的数据计算。4. 黄金公式计算全局唯一索引对于最常见的一维数据处理例如我们马上要在下一节写的向量加法我们如何将两级坐标Block 索引和 Thread 索引展平计算出该线程在所有线程中的全局唯一 ID 呢这就引出了并行编程中最著名的一维索引公式计算出在当前 Block 之前已经有多少个线程存在了。加上当前线程在自己 Block 内的偏移量。最终得到的就是这个线程用来访问一维数组如Array[i]的全局索引。第三章第一个并行程序向量加法假设我们有两个包含 100 万个浮点数的一维数组向量和我们需要将它们对应位置的元素相加并将结果存储在数组中。即。如果是在传统的 CPU C 程序中我们会写一个简单的for循环void vecAdd_CPU(float *A, float *B, float *C, int N) { for (int i 0; i N; i) { C[i] A[i] B[i]; } }但在 GPU 编程中我们要彻底抛弃这种串行循环的思维。我们要让成千上万个线程同时去执行加法操作。一个完整的 CUDA 程序分为两部分设备代码Device Code即核函数和主机代码Host Code负责统筹管理。1. 编写设备代码核函数 (Kernel)这是真正在 GPU 上成百上千个核心中并发执行的代码。我们用__global__关键字来修饰它。// CUDA 核函数每个线程计算向量中的一个元素 __global__ void vecAdd_GPU(float *A, float *B, float *C, int N) { // 黄金公式计算当前线程的全局唯一索引 i int i blockDim.x * blockIdx.x threadIdx.x; // 边界检查防止线程索引超出数组实际长度 if (i N) { C[i] A[i] B[i]; // 剥离了 for 循环每个线程只做一次加法 } }核心思想转变注意看这里没有for循环因为循环被转化为空间上的并行了。每一个启动的线程都会通过计算自己的索引i精确地找到自己负责的那个元素完成一次加法后就退出。2. 编写主机代码统筹帷幄 (Host)GPU 是个纯粹的运算机器它自己不会去内存里抓数据。因此我们需要 CPU主机来为它准备好一切。一个典型的 CUDA 主机代码包含四个标准步骤这套“分配显存 - 拷贝进去 - 执行计算 - 拷贝出来 - 释放显存”的流程是几乎所有异构计算程序的不变基石。#include stdio.h int main() { int N 1000000; // 100万个元素 size_t size N * sizeof(float); // [主机端] 分配内存并初始化 A 和 B (代码略) float *h_A ...; float *h_B ...; float *h_C ...; // --------------------------------------------------------- // 步骤 1: [设备端] 在 GPU 显存中分配空间 (cudaMalloc) // --------------------------------------------------------- float *d_A, *d_B, *d_C; cudaMalloc((void**)d_A, size); cudaMalloc((void**)d_B, size); cudaMalloc((void**)d_C, size); // --------------------------------------------------------- // 步骤 2: 数据传输 - 将数据从 CPU 拷贝到 GPU (cudaMemcpy) // --------------------------------------------------------- cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice); // --------------------------------------------------------- // 步骤 3: 启动核函数 - 定义网格和线程块大小 // --------------------------------------------------------- int threadsPerBlock 256; // 每个 Block 包含 256 个线程 // 计算需要多少个 Block向上取整确保覆盖所有元素 int blocksPerGrid (N threadsPerBlock - 1) / threadsPerBlock; // 核心语法 Grid大小, Block大小 vecAdd_GPUblocksPerGrid, threadsPerBlock(d_A, d_B, d_C, N); // --------------------------------------------------------- // 步骤 4: 数据回传 - 将计算结果从 GPU 拷贝回 CPU // --------------------------------------------------------- cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost); // [设备端] 释放显存 cudaFree(d_A); cudaFree(d_B); cudaFree(d_C); return 0; }第四章程序的编译与错误检查在传统的 C/C 开发中我们习惯使用 GCC、Clang 或 MSVC 等编译器。但当我们编写了包含主机代码CPU和设备代码GPU的混合源文件通常以.cu为扩展名时传统的编译器就“傻眼”了因为它根本不认识__global__或者...这些奇怪的语法。这时候我们就需要引入专门的异构编译工具链。1. NVCC 编译器的工作原理在 CUDA 生态中标准编译器是nvcc(NVIDIA CUDA Compiler)。 准确地说nvcc更像是一个“编译器驱动程序Driver”。它的核心工作逻辑是代码分离Separation分离nvcc首先解析.cu文件将其中的主机代码标准的 C/C和设备代码CUDA 扩展部分剥离开来。分发主机代码将纯主机代码移交给系统自带的宿主编译器比如 Linux 下的 GCCWindows 下的 MSVC进行编译。编译设备代码调用 NVIDIA 自己的底层编译器组件将设备代码编译为 PTX一种中间汇编指令或特定 GPU 架构的机器码SASS。链接最后将编译好的主机目标文件和设备目标文件链接在一起生成最终的可执行文件。2. 异步执行带来的“沉默地雷”在异构计算中有一个极易让初学者崩溃的特性核函数的启动是完全异步的Asynchronous。当你执行vecAdd_GPUgrid, block(...)时CPU 只是向 GPU 发送了一个“开始干活”的指令然后 CPU 会瞬间返回继续执行下一行代码根本不会等待 GPU 把活干完。这种设计是为了让 CPU 和 GPU 能够并行工作重叠计算最大化系统吞吐量。但这带来了一个致命问题如果 GPU 在执行计算时崩溃了比如数组越界访问CPU 在调用核函数的那一瞬间是毫无察觉的。程序不会在...这一行报错而是可能在后续尝试拷贝数据回 CPU 时突然崩溃或者直接静默失败输出一堆乱码。3. 编写标准的错误检查宏 (Error Checking Macro)为了捕获这些“沉默的地雷”我们必须养成在每一个 CUDA API 调用和核函数执行后进行严格错误检查的习惯。在工程实践中我们通常会定义一个统一的宏Macro来包装这些调用。#include stdio.h // 宏定义捕获并打印 CUDA 错误信息 #define CUDA_CHECK(call) \ { \ const cudaError_t error call; \ if (error ! cudaSuccess) \ { \ fprintf(stderr, Error: %s:%d, , __FILE__, __LINE__); \ fprintf(stderr, code: %d, reason: %s\n, error, \ cudaGetErrorString(error)); \ exit(1); \ } \ } // 检查核函数异步执行的错误 #define CHECK_KERNEL_ERRORS() \ { \ /* 捕获核函数启动时的同步错误如无效的 block 参数*/ \ CUDA_CHECK(cudaGetLastError()); \ /* 强制 CPU 等待 GPU 执行完毕捕获执行过程中的异步错误如越界*/ \ CUDA_CHECK(cudaDeviceSynchronize()); \ }如何使用它我们将 2.4 节中的代码用这个宏包装起来它就变成了符合工程标准的安全代码// 1. 包装内存分配和拷贝 CUDA_CHECK(cudaMalloc((void**)d_A, size)); CUDA_CHECK(cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice)); // 2. 启动核函数注意核函数启动本身不能直接被宏包装 vecAdd_GPUblocksPerGrid, threadsPerBlock(d_A, d_B, d_C, N); // 3. 立即捕获核函数可能产生的错误 CHECK_KERNEL_ERRORS(); // 4. 包装数据回传 CUDA_CHECK(cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost));通过引入cudaDeviceSynchronize()我们强制 CPU 停下来等待 GPU 汇报工作结果。如果在计算过程中发生了内存越界系统会精准地打印出错误发生的文件名、行号以及具体的错误原因比如an illegal memory access was encountered关于API的额外说明API的全称是Application Programming Interface应用程序编程接口。在 CUDA 中的意义NVIDIA 的工程师写好了底层与 GPU 硬件沟通的极其复杂的驱动代码然后将其包装成了一个个简单的函数如cudaMalloc、cudaMemcpy。这些函数就是 CUDA API。我们调用这些 API就是在给 GPU 下达指令。在 CUDA C/C 中几乎所有的 CUDA API 调用除了核函数...本身都会返回一个状态码用来告诉我们“后厨的菜做成功了没有”。这个返回值的专门类型叫做cudaError_t。它本质上是一个枚举enum类型如果一切顺利它的值会是cudaSuccess也就是数字 0。做个全景示范#include stdio.h #include cuda_runtime.h int main() { int V 1000; // 假设我们要分配 1000 个整数 // 1. 计算需要的总字节数 // 整数个数 乘以 每个整数占用的字节数通常是4字节 size_t size V * sizeof(int); // 2. 声明设备端指针这只是一张暂时空白的房卡 int *d_array; // 3. 声明用于接收 API 状态的 ERR 变量 cudaError_t ERR; // 4. 调用 API 分配内存并接收返回值 // 注意一定要传入指针的地址 d_array并强制转换为 (void**) ERR cudaMalloc((void**)d_array, size); // 5. 检查 API 是否执行成功 if (ERR ! cudaSuccess) { // 如果失败打印出具体的错误原因 printf(GPU内存分配失败原因: %s\n, cudaGetErrorString(ERR)); return -1; // 退出程序 } else { printf(成功在 GPU 上分配了 %d 个整数的内存\n, V); } // ... 后续计算 ... // 释放内存 cudaFree(d_array); return 0; }概率习题习题 1基础线程到数据的映射题目如果我们要让网格Grid中的每个线程计算向量加法的一个输出元素将线程/块索引映射到数据索引i的 C 语言表达式是什么习题 2步长与多元素处理进阶映射题目假设我们要修改程序让每个线程计算向量加法中的两个相邻元素。将线程/块索引映射到该线程处理的第一个数据元素索引i的表达式是什么习题 3网格与线程边界计算题目对于向量加法假设向量的长度为 8000且每个线程计算一个输出元素。如果你设定线程块大小Block Size为 1024 个线程那么整个网格Grid中总共会被启动多少个线程习题 4CUDA 内存分配的 C 语言陷阱题目假设我们有一个指向浮点数的指针float *A_d;。以下哪一行代码是正确地在设备GPU上分配内存的语法A.cudaMalloc((void*)A_d, size);B.cudaMalloc(A_d, size);C.cudaMalloc((void**)A_d, size);D.cudaMalloc((void**)A_d, size);习题 5数据传输方向题目假设A_h是主机端CPU数组A_d是设备端GPU数组两者都已经分配了 3000 字节的空间。现在需要将数据从 CPU 传输到 GPU正确的cudaMemcpy语句应该怎么写解答:1.答案i blockIdx.x * blockDim.x threadIdx.x;思路点拨这是贯穿整个 CUDA 编程的最核心公式。blockIdx.x * blockDim.x计算的是在当前线程块Block之前所有排在前面的线程块所包含的线程总数。加上当前线程在自己所属 Block 内部的偏移量threadIdx.x就能得到该线程在整个网格中的全局绝对位置。2.答案i (blockIdx.x * blockDim.x threadIdx.x) * 2;思路点拨既然每个线程要处理 2 个元素这就意味着数据数组在逻辑上被按每 2 个元素分成了多个“组”。线程的全局唯一 ID 依然是blockIdx.x * blockDim.x threadIdx.x这代表它是第几个“组”的负责人。为了找到它负责的这组数据的起始内存地址只需将它的全局线程 ID 乘以 2 即可。3.答案8192 个线程。思路点拨这是一道极易踩坑的经典题。很多初学者会回答 8000。但硬件启动线程块时只能启动完整的块。 计算所需 Block 数量的公式是向上取整(8000 / 1024) 8 个 Block。 既然启动了 8 个完整的 Block且每个 Block 有 1024 个线程那么总启动线程数就是 8 * 1024 8192 个。这也是为什么我们在 2.4 节的核函数中必须要写if (i N)这个边界检查的原因因为会有 192 个多余的线程被启动如果不加限制它们就会导致内存越界访问4.答案D思路点拨这一题考察的是 C 语言指针的基础以及 API 的设计。cudaMalloc函数的任务是去 GPU 上开辟一块空间并且将这块空间的起始地址写入到你提供的指针变量中。 因为函数内部要修改你传入的指针A_d的值它就必须知道这个指针本身在内存中的地址因此必须传入指针的地址A_d这是一个二级指针。同时为了兼容各种数据类型API 定义为void**所以需要进行强制类型转换。5.答案cudaMemcpy(A_d, A_h, 3000, cudaMemcpyHostToDevice);思路点拨重点在于参数的顺序(目标地址, 源地址, 字节数, 传输方向常量)。记住它和 C 标准库里的memcpy是一脉相承的永远是目标在前源在后就像赋值语句A_d A_h一样。同时必须显式地指定cudaMemcpyHostToDevice告诉系统底层的 DMA 控制器数据流向。实战习题考虑以下 CUDA 核函数以及调用它的相应主机函数__global__ void foo_kernel(float* a, float* b, unsigned int N) { unsigned int i blockIdx.x * blockDim.x threadIdx.x; if (i N) { b[i] 2.7f * a[i] - 4.3f; } } void foo(float* a_d, float* b_d) { unsigned int N 200000; // 执行配置(N 128 - 1)/128, 128 foo_kernel(N 128 - 1) / 128, 128(a_d, b_d, N); }问题a.每个线程块Block中有多少个线程b.整个网格Grid中总共有多少个线程c.整个网格中总共有多少个线程块d.有多少个线程执行了计算if条件语句即判断i N的代码e.有多少个线程执行了if语句内部的计算代码a. 每个线程块Block中有多少个线程答案128解析核函数的执行配置语法是网格大小, 线程块大小。在主机代码foo中第二个参数显式指定了 128这意味着每一个 Block 被硬编码为包含 128 个线程。b. 整个网格Grid中总共有多少个线程答案200,064解析网格中的总线程数 线程块总数 × 每个线程块的线程数。根据 c 题的计算网格中启动了 1563 个线程块因此总线程数为 1563 × 128 200,064 个线程。c. 整个网格中总共有多少个线程块答案1,563解析线程块的数量由公式(N 128 - 1) / 128决定。这是一种在 C 语言整数除法中实现“向上取整”的经典编程技巧。 代入 N 200000 进行计算(200000 127) / 128 200127 / 128。 在 C/C 的整数除法中小数部分会被直接截断舍弃因此 200127 / 128 1563。这意味着 GPU 实际启动了 1563 个完整的线程块。d. 有多少个线程执行了if条件判断的代码答案200,064解析这是一个极易踩坑的核心概念硬件在分配和启动线程时只能以完整的 Block 为最小单位进行分配。既然系统最终启动了 1563 个 Block那么包含在其中的全部 200,064 个线程都会被真实地分配到流处理器上并开始运行。因此所有被启动的线程都会首先计算出自己的全局索引i并执行if (i N)这行条件评估代码。e. 有多少个线程执行了if语句内部的计算代码答案200,000解析只有满足条件即全局索引i的范围在 0 到 199,999 之间的真实有效线程才会进入if代码块内部执行实际的内存读写和浮点数运算。超出这个范围的 64 个多余线程200,064 - 200,000在条件判断为假后直接越过了计算代码并结束生命周期。这就是我们在核函数中必须编写边界保护代码的根本原因。
