这是一个简单的 CUDA 程序。它添加了两个向量。__global__ void vadd ( const float* a, const float* b, float* c, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) c[i] = a[i] + b[i]; } int main () { int n = 1 << 20 ; // 一百万个浮点数（1,048,576） size_t bytes = n * sizeof ( float ); float * a = ( float* ) malloc (bytes), * b = ( float* ) malloc (bytes), * c = ( float* ) malloc (bytes); for ( int i = 0 ; i < n; i ++ ) a[i] = b[i] = 1.0 f ; float * da, * db, * dc; cudaMalloc ( & da, bytes); cudaMalloc ( & db, bytes); cudaMalloc ( & dc, bytes); cudaMemcpy (da, a, bytes, cudaMemcpyHostToDevice); cudaMemcpy (db, b, bytes, cudaMemcpyHostToDevice); vadd <<< 4096 , 256 >>> (da, db, dc, n); // 4096 * 256 = n 线程，每个浮点数一个 cudaMemcpy (c, dc, bytes, cudaMemcpyDeviceToHost); printf ( "c[0]= %f c[n-1]= %f\n " , c[ 0 ], c[n - 1 ]); } 编译针对 RTX 4090，并运行，它确实正确地计算了 1+1=21+1=2，一百万次我没有检查所有。 . $ nvcc -arch=sm_89 -o vadd vadd.cu && ./vadd c[0]=2.000000 c[n-1]=2.000000 告诉你这涉及到数千万条 CPU 指令，几个设备文件，九百个 ioctl，以及一个内存映射的门铃寄存器。在这篇文章中，我们将跟随这个内核从代码到处理单元，再回到答案。一提，这篇文章是代理人所产生的“可读性转变”的一个实例。确实很少有关于计算机的内容是你不能通过好奇和（机器增强的）坚持去发现的。有关于可读性对 AI 能帮助我们了解的含义的有趣讨论在这里。。。用 nvcc 编译我们的程序§我们应该从如何将这个 CUDA 程序转换为设备可以实际读取的内容开始。为此，我们需要一个编译器。实际上，我们需要许多编译器。nvcc 是一个驱动程序，它运行几个其他编译器并组合它们的输出。如果你传递 --keep 它会将整个管道保留在磁盘上供你查看：$ nvcc --keep -arch=sm_89 -o vadd vadd.cu && ls ... vadd.ptx # 设备代码作为 PTX（来自 cicc） vadd.sm_89.cubin # 设备代码作为 SASS（来自 ptxas） vadd.fatbin # cubin + PTX，捆绑在一起（来自 fatbinary） vadd.cudafe1.stub.c # 主机启动存根 + 核心注册 vadd.o # 最终主机对象，fatbin 嵌入 ... 主机代码传递给你的主机编译器。设备代码（vadd）需要更多步骤：cicc，一个基于 LLVM 的编译器，将其转换为 PTX，然后 ptxas 将 PTX 转换为 SASS。PTX 是一种虚拟 ISA。它有无限多的类型寄存器，并且没有关于硬件实际拥有多少个的概念。以下是 vadd 的（省略）主体：$ cat vadd.ptx ... mad.lo.s32 %r1, %r3, %r4, %r5 ; // 将寄存器 r1 设置为 ctaid*ntid + tid setp .ge.s32 %p1, %r1, %r2 ; // 如果 i >= n 设置谓词 p1 @%p1 bra $L__BB0_2 ; // 如果越界，跳到退出 cvta.to.global.u64 %rd4, %rd1 ; // 将通用指针 %rd1 转换为全局地址，储存在 %rd4 中 mul .wide.s32 %rd5, %r1, 4 ; // 将 r1 乘以 4，结果存储在 %rd5 中 add .s64 %rd6, %rd4, %rd5 ; // 将 %rd4、%rd5 相加，结果在 %rd6 中 ld.global.f32 %f2, [%rd6] ; // 将 a[i] 加载到 %f2 中 ... add .f32 %f3, %f2, %f1 ; // 将 %f1 和 %f2 相加，结果在 %f3 中 st .global.f32 [%rd10], %f3 ; // 将 c[i] = ... 存储在全局内存中 虚拟寄存器看起来像 %rd1 – %rd10，%f1 – %f3 前缀是类型：%r 是 32 位整数，%rd 是 64 位，%f 是 32 位浮点数，%p 是一个位谓词。 . PTX 比你期望的“长篇幅”要多。例如，在 %rd6 中形成一个地址需要三个 PTX 指令。这是因为 PTX 是设备无关的。为什么要三个？CUDA 指针默认是“通用的”，意味着它们可以指向全局、共享或局部内存。cvta.to.global 断言指针在全局窗口中，因此稍后可以使用更便宜的 ld.global。mul.wide.s32把索引 i 乘以 4（sizeof(float)）并在一步中将其宽度从 32 转为 64 位。add.s64 将结果加到基指针中。接下来，ptxas 将我们的 PTX（设备无关的）转换为你的架构的 SASS（设备相关的），它看起来不同：$ cuobjdump -sass vadd /*0000*/ MOV R1 , c [0x0][0x28] ; // 设置堆栈指针（ABI；在这里未使用） /*0010*/ S2R R6 , SR_CTAID.X ; // R6 = blockIdx.x /*0020*/ S2R R3 , SR_TID.X ; // R3 = threadIdx.x /*0030*/ IMAD R6 , R6 , c [0x0][0x0] , R3 ; // i = ctaid*ntid + tid /*0040*/ ISETP.GE.AND P0 , PT , R6 , c [0x0][0x178] , PT ; // P0 = (i >= n) /*0050*/ @P0 EXIT ; // 如果是，退出 /*0060*/ MOV R7 , 0x4 ; // 将字面量 4（sizeof(float)）加载到 R7 作为乘数 /*0070*/ ULDC.64 UR4 , c [0x0][0x118] ; // 从驱动程序提供的系统值统一加载 /*0080*/ IMAD.WIDE R4 , R6 , R7 , c [0x0][0x168] ; // &b[i] /*0090*/ IMAD.WIDE R2 , R6 , R7 , c [0x0][0x160] ; // &a[i] /*00a0*/ LDG.E R4 , [R4.64] ; // b[i] /*00b0*/ LDG.E R3 , [R2.64] ; // a[i] /*00c0*/ IMAD.WIDE R6 , R6 , R7 , c [0x0][0x170] ; // &c[i] /*00d0*/ FADD R9 , R4 , R3 ; // a[i] + b[i] /*00e0*/ STG.E [R6.64] , R9 ;