系列第二篇 · 第一部分：跑通第一个 DCU 程序，并彻底理解线程是如何被组织和执行的

系列第二篇 · 第一部分本系列共三篇①硬件架构 →②编程与优化基础→ ③Qwen3.5 推理优化实战第二篇篇幅较长分为若干部分。本文是第一部分跑通第一个 DCU 程序并彻底理解线程是如何被组织和执行的。阅读本文不需要 GPU 编程经验但建议先读过第一篇《DCU 架构详解》——本文会多次用到其中的 CU、SIMD、线程束等概念。引言第一篇我们把 DCU 的硬件拆开看了一遍数千个计算核心、CU 与 SIMD、寄存器与共享内存、以及延迟靠并发来掩藏的核心机制。这一篇开始写代码。本文要回答两个问题一个 DCU 程序长什么样它和 CPU 程序的区别在哪里怎么编译、怎么确认它真的跑在 DCU 上。threadIdx.x blockIdx.x * blockDim.x到底是什么这行代码几乎出现在每一个 DCU 程序里。看懂它就看懂了 DCU 的执行模型。顺序上我们先建立最关键的那个观念转变再看代码。一、思维方式的转变从一个人做一万次到一万个人各做一次这是学习 DCU 编程要跨过的第一道、也是最高的一道坎。它不是语法问题是思考方式的问题。先看一个最普通的需求两个数组相加C[i] A[i] B[i]数组长度是 10000。在 CPU 上你会很自然地写一个循环for (int i 0; i N; i) { C[i] A[i] B[i]; }这段代码的含义是一个执行者依次做一万次加法。而在 DCU 上思路要反过来。你不再描述整件事怎么做完而是描述假设我是其中一个线程我该做什么答案是我只负责一个位置。我先搞清楚自己是第几号线程然后把 A、B 中对应位置的两个数加起来写进 C 的对应位置。做完就结束。然后你一次性启动一万个这样的线程让它们同时各做一次加法。用官方的表述从 CPU 编程到 DCU 编程你最需要转换的是你的思考方式。在设计核函数时你需要站在一个线程的角度去思考你的问题如何被分解为一个个相似的小问题而每一个小问题需要一个线程去做什么样的事情来解决它。这个转变可以概括成一句话CPUDCU你写的代码描述的是整个任务循环一万次一个线程的职责做一次谁来重复你写的循环硬件启动一万个线程形象比喻一个人做一万次一万个人各做一次这个观念一旦建立后面的一切——线程编号、线程块、线程束——都只是在回答同一个问题的细节我是谁我该处理哪个数据二、第一个 DCU 程序数组相加学习 CPU 平台上的编程语言第一个程序一般是 Hello world!。但鉴于 DCU 多任务数据处理架构的特点我们的第一个 DCU 程序换一个更贴合的例子——两数组相加数组 A 与数组 B 相加结果赋值给数组 C。这个程序足够简单但它能让你掌握 DCU 编程的基本方法、了解 DCU 编程和 CPU 编程之间的重要差别以及编译运行的方法。2.1 CPU 版本先看 C 语言版本就是一个 for 循环的简单程序#include stdlib.h #define N 10000 int main() { // 申请数据空间 float *A (float *) malloc(N * sizeof(float)); float *B (float *) malloc(N * sizeof(float)); float *C (float *) malloc(N * sizeof(float)); // 数据初始化 for (int i 0; i N; i) { A[i] 1; B[i] 1; C[i] 0; } // 进行数组相加 for (int i 0; i N; i) { C[i] A[i] B[i]; } free(A); free(B); free(C); return 0; }CPU 程序采用 gcc 编译器编译。假设文件名为vector_add_cpu.cpp在 Linux 平台下用下面的命令编译gcc vector_add_cpu.cpp -o vector_add_cpu运行后就可以看到 C 数组内的值都被置为了 2数组相加完成。2.2 DCU 版本接下来看 DCU 版程序。先不用急着看明白每一行的意思后面会详细拆解#include iostream #include hip/hip_runtime.h #include hip/hip_runtime.h #define N 10000 __global__ void add(float *d_A, float *d_B, float *d_C) { int tid threadIdx.x blockIdx.x * blockDim.x; if (tid N) { d_C[tid] d_A[tid] d_B[tid]; } } int main() { // 申请数据空间 float *A (float *) malloc(N * sizeof(float)); float *B (float *) malloc(N * sizeof(float)); float *C (float *) malloc(N * sizeof(float)); float *d_A NULL; float *d_B NULL; float *d_C NULL; hipMalloc((void **) d_A, N * sizeof(float)); hipMalloc((void **) d_B, N * sizeof(float)); hipMalloc((void **) d_C, N * sizeof(float)); // 数据初始化 for (int i 0; i N; i) { A[i] 1; B[i] 1; C[i] 0; } hipMemcpy(d_A, A, sizeof(float) * N, hipMemcpyHostToDevice); hipMemcpy(d_B, B, sizeof(float) * N, hipMemcpyHostToDevice); hipMemcpy(d_C, C, sizeof(float) * N, hipMemcpyHostToDevice); dim3 blocksize(256, 1); dim3 gridsize(N / 256 1, 1); // 进行数组相加 addgridsize, blocksize(d_A, d_B, d_C); // 结果验证 hipMemcpy(C, d_C, sizeof(float) * N, hipMemcpyDeviceToHost); for (int i 0; i N; i) { std::cout C[i] std::endl; } // 释放申请空间 free(A); free(B); free(C); hipFree(d_A); hipFree(d_B); hipFree(d_C); }注意那个for循环消失了。它被addgridsize, blocksize(...)这一行取代——这正是第一节说的转变循环不再由你写而是由硬件启动的成千上万个线程并行地完成。2.3 编译DCU 程序需要采用专用的编译器hipcc进行编译编译出的二进制代码可以被设备端加载运行。编译命令与 gcc 相似hipcc vector_add_dcu.cpp -o vector_add_dcu编译后运行如果成功应该会在屏幕上看到 2 这个数字不断被打印。到这里你已经运行了第一个 DCU 程序。抛开代码来看编译方法上二者似乎并没有什么区别仅仅是更换了编译器。2.4 怎么确认它真的跑在 DCU 上答案是使用rocm-smi命令。它的作用与 Linux 上的top命令类似——top看 CPU 负载rocm-smi看 DCU 负载当然它还有更多用途。在执行上面程序的同时使用rocm-smi可以实时监控 DCU 负载情况。输出中第一列是 DCU 的序号如果你的机器里有不止一块 DCU 就会有多行%DCU列代表 DCU 的实时负载。运行程序时可以看到第一块 DCU 处于满载状态——这就证明计算确实发生在 DCU 上。三、拆解一个 HIP 程序的四个部分运行完程序回头对比 CPU 端和 DCU 端的代码会发现二者还有不小的区别。这些区别可以归纳为四个部分。3.1 空间管理显存是独立的首先要明确一点数组相加最开始的输入数据和输出数据都是在 CPU 端初始化的这些数据一开始都在主机内存当中。而 DCU 要参与计算就必须拿到这些数据。要实现这一目的需要在DCU 的显存上开辟一块空间大小与 CPU 端内存相对应经过PCI-E 总线将数据传输到 DCU 当中。这正对应第一篇讲过的硬件事实DCU 有独立于主机内存的全局内存两者通过 PCI-E 相连。CPU 的指针和 DCU 的指针指向的是两个完全不同的地址空间——这就是为什么代码里既有A主机又有d_A设备d_前缀是一个约定俗成的命名习惯用来提醒你这是设备端的指针。与 CPU 内存管理类似调用hipMalloc就可以完成 DCU 内存分配hipMalloc((void **) d_A, N * sizeof(float));hipMalloc是最基础的内存分配接口此外还有一些其他的内存分配方式例如页锁定内存这部分将在后续部分介绍。HIP 程序其实是 C 程序的一种扩展所以同样需要注意管理空间的申请和释放以防止内存泄露。程序结束前要用hipFree释放显存hipFree(d_A);3.2 数据拷贝注意方向完成空间申请后下一步就是把 CPU 端准备好的数据传输给 DCU 显存使用hipMemcpy完成hipMemcpy(d_A, A, sizeof(float) * N, hipMemcpyHostToDevice);需要注意的是hipMemcpy具有方向性。最后一个参数指示了方向参数含义hipMemcpyHostToDevice主机端 → 设备端hipMemcpyDeviceToHost设备端 → 主机端计算完成后结果还在显存里必须再拷回主机才能打印hipMemcpy(C, d_C, sizeof(float) * N, hipMemcpyDeviceToHost);3.3 函数定义核函数__global__ void add(float *d_A, float *d_B, float *d_C)__global__是一个函数标识符它代表这个函数在主机端被调用但执行是在设备端。这个标识符会被 hipcc 编译器识别到然后按照 DCU 端的指令去翻译下面的代码段。被__global__标识的函数又被称为核函数或kernel。核函数代表了设备端的入口——一旦进入核函数函数的执行和控制就交给了设备端。除了__global__HIP 还有其他标识符提供给开发者使用通过标识符可以管理设备端的存储并让开发者有更灵活的程序设计方式。这些内容会在后续部分详细介绍。进入核函数体中int tid threadIdx.x blockIdx.x * blockDim.x; // ① 我是谁 if (tid N) { // ② 越界了吗 d_C[tid] d_A[tid] d_B[tid]; // ③ 干活 }三步首先使用threadIdx.x、blockIdx.x、blockDim.x定位当前的线程编号根据线程编号算出当前线程读取数据的偏移地址接着判断线程索引的数据是否超越数据边界如果没有超过运行核函数真正的计算逻辑进行数据相加。第 ① 步具体是什么意思这正是第四节要回答的问题。3.4 函数执行三个尖括号dim3 blocksize(256, 1); dim3 gridsize(N / 256 1, 1); addgridsize, blocksize(d_A, d_B, d_C);核函数需要采用三个尖括号的方式调起。一旦调用起来设备端就会自动创建gridsize × blocksize数量的线程。这些线程如何被组织第四节会讲。这里只需要知道一共有gridsize × blocksize个线程都会运行add函数的代码括号里的参数就是执行该函数时需要的数据。这些启动的线程将会被设备端的调度器分配到众多的计算核心中运行从而实现并行化。这一点就是 DCU 能实现计算加速的关键。一个必须知道的陷阱核函数调用是异步的这里有一个初学者极易踩的坑核函数被调用后执行权限交给设备端但主机端【不会】等待核函数执行完毕再执行下面的代码而是立刻往下执行。也就是说主机端和设备端是异步的。这意味着如果你想给核函数计时在add...(...)这一行的前后各取一次时间你会测出一个接近于 0 的数值——因为主机根本没等它算完就往下跑了你测的只是发出命令的时间不是执行完成的时间。那为什么上面的程序还能打印出正确的结果因为紧接着的那句hipMemcpy是隐式同步的一旦主机端遇到hipMemcpy就会等待设备端的所有操作包括核函数执行、数据传输执行完毕以后才会往下执行。这样会效率低下吗当然会。不过 HIP 也具有异步拷贝功能这需要配合流来使用——通过这个技术可以同时启动多个核函数、多个拷贝剩下的就交给设备端调度执行。这部分内容会在本篇的并发执行部分详细展开。四、执行模型threadIdx.x到底是什么现在回答那个悬而未决的问题。4.1 三层组织网格、线程块、线程当核函数被调用起来后会自动启动数千个线程来执行核函数的代码。这些线程不是平铺的一大堆而是被组织成三个层次网格 Grid一次核函数调用启动的全部线程 └── 线程块 Block线程的分组 └── 线程 Thread最小执行单位回到三个尖括号addgridsize, blocksize(d_A, d_B, d_C); // ↑ 第一个参数 ↑ 第二个参数 // 启动多少个线程块 每个线程块里有多少线程第二个参数blocksize配置线程块的大小每块多少线程第一个参数gridsize配置一共启动多少个线程块。例如gridsize为 1、blocksize为 256翻译过来就是一共启动1 × 256个线程来执行add的代码。图 3-3 多维线程组织模型图 3-3 展示了这种层次关系一个 GRID 中包含多个 BLOCK每个 BLOCK 中又包含多个 Thread。为什么要分成两层初学者常有的疑问是既然最终就是一堆线程为什么不直接说启动 10000 个线程非要分成40 个块 × 256 个线程原因在于线程块是硬件调度和资源分配的单位。回忆第一篇的内容一个线程块只能在一个 CU 上被调度一旦被调度到某个 CU 上就会保存在该 CU 上直到执行完成CU 上的共享内存被分配给该 CU 上的常驻线程块寄存器被分配给线程同一个块里的线程因为共处一个 CU可以通过共享内存相互合作和通信。所以块不是一个多余的抽象它对应着实实在在的硬件边界块内的线程能协作块之间基本各干各的。你把多少线程放进一个块直接决定了这个块要吃掉 CU 上多少资源——这一点会在第六节占用率变得非常重要。4.2 多维组织与dim3线程块在逻辑上可以被配置为一维、二维或三维。如果需要使用多维线程组织模型在设置gridsize和blocksize时就需要配置为dim3类型dim3 dimGrid(90, 90); // 网格x 轴 90 个块y 轴 90 个块 dim3 dimBlock(64, 1); // 每个块64 个线程 hipLaunchKernelGGL(myGpuMatrix, dimGrid, dimBlock, 0, 0, dev_A, dev_B, dev_C_t);dim3是基于uint3指定维度的整数向量类型分别用x、y、z来表示网格三个维度上线程块的个数或者线程块三个维度上线程的个数。x、y、z中未指定的任何部分默认初始化为 1结构体定义如下typedef struct dim3 { uint32_t x; uint32_t y; uint32_t z; dim3(uint32_t _x 1, uint32_t _y 1, uint32_t _z 1) : x(_x), y(_y), z(_z) {}; };注意上面出现了两种启动核函数的写法——三尖括号addgrid, block(...)和hipLaunchKernelGGL(kernel, grid, block, 0, 0, ...)。两者等价后者是 HIP 提供的宏形式多出的两个参数分别是动态共享内存大小和流stream暂时填 0 即可后续部分会讲。维度只是给你看的这里有一个重要的观念所谓的维度其实是人为的一种划分。在执行时线程都被组织成一维形式。这个概念就好像我们用 C 语言申请了一个二维数组a[N][M]虽然使用时有维度这个概念但在物理空间中这些数据其实是连续存储的。我们当然可以用一维数组去替代二维的只要计算好偏移就可以。线程也是一样。多维只是为了让你的代码更贴合问题的形状比如处理图像时用二维更自然硬件并不在乎。在一个块中每个线程都有唯一 ID。如果将x 作为第一个维度最内层、y 作为第二个维度、z 作为第三个维度最外层那么不管是二维还是三维线程块都可以转化为一维线程布局。线程的增长方向是沿着 x、y、z 依次递增——这个规则就像在访问一个二维或三维数组一样。记住这一点后面讲共享内存的 bank 冲突时还会用到。4.3 索引计算回答我是谁现在可以看懂那行代码了int tid threadIdx.x blockIdx.x * blockDim.x;三个内置变量的含义变量含义threadIdx.x当前线程在它所属的块内的编号blockIdx.x当前块在网格内的编号blockDim.x每个块的大小块里有多少线程所以这个公式的逻辑是全局编号 我前面有多少个完整的块 × 每块的线程数 我在本块内排第几 blockIdx.x × blockDim.x threadIdx.x举个具体的例子blockDim.x 256当前线程在2 号块的5 号位上那么它的全局编号就是5 2 × 256 517。它负责处理A[517] B[517]。由于 GRID 和 BLOCK 的大小和维度可以自由组合就会形成不同的线程索引方式组合线程全局 ID 计算式1D grid, 1D blockblockIdx.x * blockDim.x threadIdx.x1D grid, 2D blockblockIdx.x * (blockDim.x * blockDim.y) blockDim.x * threadIdx.y threadIdx.x3D grid, 1D block(gridDim.x * gridDim.y * blockIdx.z gridDim.x * blockIdx.y blockIdx.x) * blockDim.x threadIdx.x3D grid, 3D block(gridDim.x * gridDim.y * blockIdx.z gridDim.x * blockIdx.y blockIdx.x) * (blockDim.x * blockDim.y * blockDim.z) blockDim.x * blockDim.y * threadIdx.z blockDim.x * threadIdx.y threadIdx.x公式看着吓人但本质都一样先算出我所在的块前面有多少个线程再加上我在块内的位置。而块内位置的计算又遵循 x 最内、z 最外的规则。4.4 为什么必须有边界检查if (tid N) { ... }这句看似多余的判断其实是必需的。原因藏在启动配置里dim3 gridsize(N / 256 1, 1); // N 1000010000 / 256 39.06C 语言的整数除法会截断为39加 1 得到40个块。于是实际启动的线程数是40 × 256 10240 个线程比 N 10000多出了 240 个。这多出来的 240 个线程它们的tid从 10000 一直排到 10239。如果不加边界检查它们就会去读写d_A[10000]这种根本不存在的位置——轻则算出垃圾数据重则程序崩溃。这是 DCU 编程的一个常规操作因为线程数通常要凑成块的整数倍几乎总会启动比实际需要更多的线程所以核函数里几乎总要写一句边界检查。五、硬件的真相线程束wavefront到这里你已经能写出正确的 DCU 程序了。但要写出快的程序还得再往下看一层——硬件其实并不按线程块来执行。5.1 逻辑视图 vs 硬件视图图 3-5 线程块逻辑视图与硬件视图关系图图 3-5 展示了这层关系逻辑视图你写代码时的视角线程块是一堆线程的集合可以是一维、二维或三维硬件视图硬件执行时的视角线程块是一维线程束的集合——每 64 个连续的线程组成一个线程束执行线程束被送到多处理器CU上执行。用一句话概括从逻辑角度看线程块是线程集合可以被组织为一维、二维或三维布局但从硬件角度看线程块是一维线程束的集合在线程块中线程被组织成一维布局每 64 个连续线程组成一个线程束。这就是第一篇讲过的wavefront一个线程束中的所有线程按单指令多线程SIMT方式执行——它们步调一致同一时刻执行同一条指令。核函数启动后的完整链条是这样的核函数启动 → 产生大量线程按你的配置分成若干线程块 → 线程块被分发到各个 CU 上一个块只去一个 CU → 到了 CU 上块内线程被切成一个个 64 线程的线程束 → 线程束被 SIMD 调度执行5.2 凑不满 64 会浪费既然硬件按 64 个线程一束来切那么如果线程块大小不是线程束大小的整数倍会发生什么答案是那些除以 64 余下的线程依然会组成一个线程束。图 3-4 线程束分配示意图图 3-4 给了一个具体例子一个 x 轴为 80、y 轴为 2 的二维线程块整个块中有160 个线程。160 ÷ 64 2 余 32所以它会被切成3 个线程束前两个各装满 64 个线程最后 32 个线程单独组成第三个线程束。问题在于这最后一个只装了 32 个线程的线程束却要占用 64 个线程的调度资源和存储资源。图中灰色的那半截就是被白白占掉、却什么也不干的部分。你以为启动了160 个线程 硬件实际占用64 × 3 192 个线程的资源 浪费32 个槽位约 17%所以这种现象的发生会降低性能。这也就直接推出了那条最基本的配置原则保持每个块中的线程数量是 64 的倍数。5.3 线程束分化if / else 的代价线程束步调一致的特性带来了另一个后果。在同一线程束中的线程执行不同的指令被称为线程束分化divergence这种情况最常见于程序中的if和else。因为DCU 没有类似 CPU 那样复杂的分支预测机制所以在一个指令周期中只有满足条件的线程才会执行不满足条件的线程什么都不做、等待其他线程执行完毕然后同一个线程束内的线程继续同步往下执行。看这段代码A; if (threadIdx.x 4) { B; } else { C; } A;图 3-6 线程束分化图 3-6 以 32 个线程为例展示了执行过程时间自上而下A 语句线程束中所有线程都执行没有分歧。分支 B 和 C条件threadIdx.x 4将线程束分成了两派。前 4 个线程走B后 28 个线程走C。硬件先执行B前 4 个线程干活后 28 个等待再执行C后 28 个线程干活前 4 个等待。两条路径串行执行总耗时翻倍。A 语句分歧结束所有线程再次同步执行。这就是线程束分化的代价本该并行执行的线程因为分支而被迫串行化。如果分支条件在同一个线程束内差异很大性能会显著下降。避免线程束分化的常见策略包括让分支条件与线程束边界对齐例如用threadIdx.x / 64代替threadIdx.x 4使整个线程束走同一条路径用数据预处理代替分支将条件判断提前到数据准备阶段核函数内只做统一计算减少分支深度尽量用三元运算符或简单算术代替深层if/else嵌套。线程束分化是 DCU 编程中一个重要的性能考量因素。理解它就能写出既正确又高效的核函数。六、总结本文从零开始带你跑通了第一个 DCU 程序并深入理解了线程的组织和执行模型。回顾一下关键点思维转变从 CPU 的“一个人做一万次”到 DCU 的“一万个人各做一次”是学习 DCU 编程的第一道坎。HIP 程序的四个部分空间管理hipMalloc、数据拷贝hipMemcpy、核函数定义__global__、核函数执行grid, block。执行模型线程被组织为网格→线程块→线程三层结构。threadIdx.x blockIdx.x * blockDim.x是计算全局线程编号的核心公式。边界检查因为线程数要凑成块的整数倍几乎总会启动多余线程所以核函数里必须写边界检查。线程束硬件按 64 个线程为一束执行。块大小要是 64 的倍数避免浪费分支要尽量对齐线程束边界避免分化。掌握了这些基础下一篇我们将进入更深入的优化话题共享内存、同步原语、以及如何通过调整线程块大小来提升占用率。敬请期待。参考文献DCU 架构详解 - CSDN博客编写第一个 DCU 程序 - 开发者源文档