GPU编程的演变

每个十年，GPU都经历了自我革新——从绘制三角形到生成世界，现在，它们能与语言进行推理。我意识到，在我的整个编程旅程中，我一直与GPU紧密合作，并尝试了无数种方法来编程它们。从用GLSL编写像素着色器，到在OpenCL中实现实时3D扫描算法，再到在PyTorch和Tensorflow中优化深度学习模型。有什么比写一篇关于GPU编程演化的博客文章更好的方式呢？这篇文章充满了怀旧和网络迷因？

多年来，GPU编程领域发生了巨大的变化。出现了新的编程模型、新的框架和新的硬件架构。如今研究它们已经没有意义；然而，GPU编程的进化路径非常有趣。如果你是AI专家或其它领域的开发者，它可以帮助你拓宽专业知识，或者提供必要的灵感进入GPU编程的世界。它可以给你新的想法来解决当前的问题，特别是考虑到我们今天在AI中遇到的一些问题，早在25年前图形程序员就已经面临过。

这是一次轻松愉快、充满怀旧感的GPU编程历史之旅，从2000年制作看起来有凹凸感的砖墙到2025年优化大型语言模型中的注意力机制。如果你不感兴趣或已经熟悉这些内容，可以跳过代码片段，享受故事本身。

1、智能像素

在20世纪初，GPU仅用于可视化，渲染管道是完全固定的。这类似于HTML，你会预先定义场景：几何体、纹理、灯光和相机的位置，GPU会负责渲染它。当然，你可以在运行时进行一些自定义，但只能通过改变预定义函数的参数来有限地进行，这种自定义完全是在CPU端完成的。

这是一个使用老式OpenGL渲染三角形的简单示例，摘自这里。

// 将帧缓冲区中的每个像素设置为当前清除颜色。
glClear(GL_COLOR_BUFFER_BIT);  

// 绘制是通过指定一系列顶点来完成的。这些顶点是如何连接的。GL_POLYGON构建一个填充的多边形。
glBegin(GL_POLYGON);  
  glColor3f(1, 0, 0); glVertex3f(-0.6, -0.75, 0.5);  
  glColor3f(0, 1, 0); glVertex3f(0.6, -0.75, 0);  
  glColor3f(0, 0, 1); glVertex3f(0, 0.75, 0);  
glEnd();  

// 刷新绘图命令缓冲区以尽快发生绘图。
glFlush();

在2000年代初，能够实际编程如何在屏幕上渲染像素的想法是非常革命性的。

我第一次接触这些想法是通过一篇文章（2001年）在一家流行的俄罗斯游戏开发网站上关于NV_register_combiners扩展的介绍。令人惊讶的是，这篇文章仍然在线可用。

这个扩展允许你编程如何从各种输入（如纹理颜色和光照）中计算最终像素颜色，从而创建更复杂的视觉效果。这种计算是在GPU上进行的，可以实现实时性能。这相当于在GPU上运行一个小的汇编程序，为每个正在渲染的像素。

图形开发者对这个想法感到着迷，因为它使场景的视觉保真度大幅提升。不久之后，GLSL被构想并正式于2004年引入，允许用类似C的语言编写更复杂的着色器（定义如何操作几何体或像素的小程序）。

你是否感觉GPU贫乏？想象一下那时的情况更糟糕！每一代新的GPU都会引入新功能和能力，显著提高游戏的视觉保真度。拥有新一代GPU是玩最新和最好的游戏的前提条件。对于计算机图形学爱好者来说，等待的挫败感和获得新显卡的兴奋感成倍增加！幸运的是，我可以骗我父母给我买一张支持SHADERS的新显卡！这当然是推进我的计算机科学教育所必需的。能够在高设置下玩《上古卷轴》只是额外的奖励。

这是来自rastertek.com的一个简单的GLSL程序示例，用于执行凹凸映射，这是一种通过扰动纹理的表面法线来模拟物体表面上的小规模凹凸和皱纹的效果。

in vec2 texCoord;  
in vec3 normal;  
in vec3 tangent;  
in vec3 binormal;  

void main(void)  
{  
    // 使用此纹理坐标位置的采样器从纹理中采样像素颜色。  
    vec4 textureColor = texture(shaderTexture1, texCoord);  

    // 从法线贴图中采样像素。  
    vec4 bumpMap = texture(shaderTexture2, texCoord);  

    // 将法线值的范围从 (0, +1) 扩展到 (-1, +1)。  
    vec3 bumpMap = (bumpMap * 2.0f) - 1.0f;  

    // 计算法线贴图中的法线。  
    bumpNormal = (bumpMap.x * tangent) + (bumpMap.y * binormal) + (bumpMap.z * normal);  

    // 归一化结果的凹凸法线。  
    bumpNormal = normalize(bumpNormal);  

    // 根据法线贴图值计算该像素上的光强度。  
    float lightIntensity = clamp(dot(bumpNormal, -lightDirection), 0.0f, 1.0f);  

    // 确定根据漫反射颜色和光强度结合后的最终漫反射颜色。  
    outputColor =  clamp((diffuseLightColor * lightIntensity), 0.0f, 1.0f);  

    // 将最终的光颜色与纹理颜色结合。  
    outputColor = outputColor * textureColor;  
}

所有这些in vec3变量是什么意思？这些是着色器程序的输入。这些是每个顶点指定的，并在渲染的三角形表面上进行插值。插值由GPU硬件完成，并输入到每个正在渲染的像素的着色器程序中。这样，你可以为每个像素有不同的值，从而实现更复杂的效果。这种方式允许在所有正在渲染的像素上并行化计算，因为每个像素可以独立处理。

着色器迅速从简单的像素颜色操作发展到复杂的效应，如阴影、反射和折射。图形程序员特别痴迷于在不增加场景几何复杂性的情况下模拟复杂的表面细节。这一兔子洞的最深处是视差遮蔽映射技术，它在像素着色器中执行一种光线追踪，即遍历空间以确定光线与由高度图纹理定义的表面的交点。这样，一个完全平坦的表面可以呈现出复杂的3D细节。

2、GPU作为通用计算机

此时，你可能会问及LLM、深度学习以及在GPU上执行通用计算的能力。然而，请看看上面的着色器程序。它就像一段C代码。为什么不能用它在GPU上执行任意计算？确实可以，而且人们自2000年代初以来一直在这样做。然而，我们需要首先解决一个问题。我们如何将数据输入和输出到GPU？

输入数据相当直接。我们可以将我们的数据编码为纹理或几何体并上传到GPU。但如何获取数据？为了帮助这一点，我们可以使用渲染到纹理技术。它允许我们将着色器程序的输出渲染到纹理而不是屏幕。然后我们可以将该纹理回传到CPU。

对于不熟悉计算机图形学术语的人。纹理只是一个图像。在计算机图形学中，纹理用于存储可以应用于3D模型表面以赋予其颜色和细节的图像数据。纹理通常是一个二维像素数组，其中每个像素包含颜色信息（例如RGB值）和有时其他信息，如alpha（透明度）或用于凹凸映射的法线向量。

这种技术实际上比着色器本身还要古老，因为在着色器时代之前就用于创建动态反射和阴影等效果。例如，要创建反射效果，你可以从反射摄像机的角度（例如水面上方）渲染场景到纹理，然后使用该纹理来渲染水面。你可以使用像素着色器来扭曲纹理坐标，模拟水面波纹。

一些有创意的人发现，你可以使用这种技术在GPU上执行任意计算，通过将输入数据编码为纹理，编写一个着色器程序来执行计算，将输出渲染到纹理，然后将该纹理回传到CPU。

你能用这种方法实现什么？你今天可以用CUDA实现的一切。早期的GPGPU流行的一种技术是乒乓渲染，其中两个纹理交替用于读取和写入。这样，你可以计算，将输入纹理进行某种函数计算，将结果写入输出纹理，然后使用该输出纹理作为后续计算的输入，依此类推。这样，你可以通过串联多个着色器程序来构建复杂的计算。你不必专门处理图像。你可以将任何数据编码为纹理，例如二维浮点数数组、三维体素体积、图等等。

例如，快速傅里叶变换（FFT）算法可以使用着色器和渲染到纹理技术来实现。这是来自GPU Gems 2的一个基于GPU的FFT实现，以及它的医学图像重建。

下面是一个单个FFT传递的片段着色器示例。它类似于你今天会写的CUDA内核，如下所示。它本质上是一个为输出纹理的每个像素调用的函数。它从输入纹理中读取数据，执行一些计算，并将结果作为像素的颜色写入，然后存储在输出纹理中。

void FragmentProgram(in float4 TexCoordRect  
                     : TEXCOORD0, out float4 sColor0  
                     : COLOR0, out float4 sColor1  
                     : COLOR1, out float4 sColor2  
                     : COLOR2, out float4 sColor3  
                     : COLOR3, uniform samplerRECT Real1,  
                       uniform samplerRECT Imag1, uniform samplerRECT Real2,  
                       uniform samplerRECT Imag2,  
                       uniform samplerRECT ButterflyLookupI,  
                       uniform samplerRECT ButterflyLookupWR,  
                       uniform samplerRECT ButterflyLookupWI)  
{  
  // 读取蝴蝶索引  
  float4 i = texRECT(ButterflyLookupI, TexCoordRect.xy);  
  // 读取打乱坐标  
  float4 WR = texRECT(ButterflyLookupWR, TexCoordRect.xy);  
  // 读取权重  
  float4 WI = texRECT(ButterflyLookupWI, TexCoordRect.xy);  

  // 执行蝴蝶操作，将结果存储在输出颜色中  
  float2 Res;  
  float2 r1 = float2(i.x, TexCoordRect.y);  
  float2 r2 = float2(i.w, TexCoordRect.y);  
  float4 InputX1 = texRECT(Real1, r1);  
  float4 InputY1 = texRECT(Imag1, r1);  
  float4 InputX2 = texRECT(Real1, r2);  
  float4 InputY2 = texRECT(Imag1, r2);  
  Res.x = WR.x * InputX2.x - WI.x * InputY2.x;  
  Res.y = WI.x * InputX2.x + WR.x * InputY2.x;  
  sColor0.x = InputX1.x + Res.x;  
  sColor1.x = InputY1.x + Res.y;  
  float4 InputX1_ = texRECT(Real2, r1);  
  float4 InputY1_ = texRECT(Imag2, r1);  
  float4 InputX2_ = texRECT(Real2, r2);  
  float4 InputY2_ = texRECT(Imag2, r2);  
  Res.x = WR.x * InputX2_.x - WI.x * InputY2_.x;  
  Res.y = WI.x * InputX2_.x + WR.x * InputY2_.x;  
  sColor2.x = InputX1_.x + Res.x;  
  sColor3.x = InputY1_.x + Res.y;  
}

上面的代码是用Cg语言编写的。这是Nvidia早期试图主导图形计算市场的尝试……我是说，简化着色器编程。幸运的是，没人关心它，市场依赖于更广泛支持的GLSL和HLSL语言。

我对这些发展感到着迷！这种技术解锁了计算机图形学、科学和医学等领域中大量新的应用。我个人曾用它来实现高级图形效果。这里有一个使用FFT生成复杂水面的例子。http://www.uraldev.ru/articles/35/page/2。这项技术被用于电影泰坦尼克号和一些先进的游戏中，如刺客信条。

这些文章值得阅读吗？当然不。我想展示的是我如何使用Web-Archive恢复一些不再在线的旧文章，并给帖子添加一个表情包图片。

3、进入CUDA

尽管使用着色器进行通用计算的技术非常强大，但它仍然有些局限。编程模型并不十分友好，因为你必须将数据编码为纹理或其他图形原语。渲染到纹理的方法涉及渲染整个屏幕的矩形区域，确保所有渲染的像素精确对齐输出纹理的纹素。很容易配置错误图形流水线，比如忘记关闭纹理过滤，这会导致错误的结果。

所有这些细节都很分散，使得难以专注于实际计算，尤其是对于非图形程序员来说。因此，NVIDIA在2007年推出了CUDA，它提供了用于在NVIDIA GPU上编写通用计算的C-like编程模型。

编程模型类似于着色器编程模型，因为您仍然编写一个内核函数，该函数由许多线程并行执行。每个线程由其1D、2D或3D索引来标识，您可以使用它来计算想要处理的数据的内存地址。在着色器编程模型中，您会使用纹理坐标或其他变化变量来做到这一点，而在CUDA中，您会使用线程索引。然而，所有设置图形流水线、管理纹理、帧缓冲区等的支撑结构都被消除了。您可以分配GPU上的内存，复制数据到其中，启动内核，然后复制结果回来。

这里是上面的FFT内核在CUDA中的样子。再次提醒，如果您是为了故事而来，可以跳过。

// 帮助函数以执行复数乘法和加法操作。  
__device__ float2 butterfly_op(float2 a, float2 b, float2 twiddle) {  
    // 执行复数乘法和加法  
    float2 temp_result;  
    temp_result.x = b.x * twiddle.x - b.y * twiddle.y;  
    temp_result.y = b.y * twiddle.x + b.x * twiddle.y;  
    return a + temp_result;  
}  

__global__ void fft_stage_kernel(  
    // 输入数据数组（现在使用float2表示复数）  
    float2 *d_input1,  
    float2 *d_input2,  

    // 组合的蝴蝶查找表（现在使用float2表示复数旋转因子）  
    float *d_butterflyLookupI,  
    float2 *d_butterflyTwiddles,  

    // 输出数据数组（现在使用float2）  
    float2 *d_out1,  
    float2 *d_out2,  

    int width,  
    int height  
) {  
    int tx = blockIdx.x * blockDim.x + threadIdx.x;  
    int ty = blockIdx.y * blockDim.y + threadIdx.y;  

    if (tx >= width || ty >= height) {  
        return;  
    }  

    int index = ty * width + tx;  

    // 读取蝴蝶查找索引和复数旋转因子  
    int lookup_i = (int)d_butterflyLookupI[index];  
    float2 twiddle_factor = d_butterflyTwiddles[index];  

    // 使用组合的float2数组读取输入数据  
    int r1_idx = ty * width + tx;  
    int r2_idx = ty * width + lookup_i;  

    float2 input1 = d_input1[r1_idx];  
    float2 input2 = d_input1[r2_idx];  

    // 对第一对输入执行蝴蝶操作  
    d_out1[index] = butterfly_op(input1, input2, twiddle_factor);  

    // 处理第二对数据数组  
    float2 input1_prime = d_input2[r1_idx];  
    float2 input2_prime = d_input2[r2_idx];  

    // 执行第二个蝴蝶操作  
    d_out2[index] = butterfly_op(input1_prime, input2_prime, twiddle_factor);  
}

我一直等着拿到支持CUDA的GPU，再次。我赚了钱，所以不需要再骗我父母了，但高端PC升级仍然是相当大的开支，而且你需要经常做。我第一个支持CUDA的GPU是8800GT，这是历史上最传奇的一代GPU。它采用了全新的架构并引入了CUDA。此外，单张8800 GTX卡能够超越两代之前的7900 GTX卡在SLI中的表现，并且功耗和价格相当（599美元——忍住眼泪吧）。什么时候我们能看到这样的性能和价值飞跃，皮革夹克CEO先生？

4、CUDA护城河？

作为一名真正的开源战士，我没有使用CUDA，而是依靠OpenCL来进行工作。它不像CUDA那样得到很好的支持：调试器和其他工具不够先进，有更多的故障，而且在NVIDIA硬件上CUDA可以获得稍微更好的性能。然而，它的缺点被它是一个开放标准并且适用于AMD和Intel GPU的事实所抵消，因此CUDA当时远非垄断。

在我工作中，我使用OpenCL来实现实时3D扫描算法。Artec Eva是一款用于医疗或工业应用的专业3D扫描仪。实时3D扫描涉及大量的GPU计算来处理输入视频流，识别相对于环境的位置（类似于自动驾驶汽车使用的算法），将所有输入数据融合成一个3D模型，并在屏幕上显示。所有这一切都必须实时发生，以便用户可以立即看到结果并根据需要调整他们的位置。

我选择了OpenCL，这在当时是一个勇敢的选择，可能是一个糟糕的产品决策，因为当你购买价值12000美元的3D扫描仪时，你可以负担得起一台不错的GPU，而不必担心供应商锁定。然而，随着时间的推移，随着GPU变得越来越强大，并且可以在笔记本电脑GPU上运行该管道，特别是Microsoft Surface平板电脑，选择OpenCL变得更加相关。现在，操作员手中有一台轻便的显示器，可以围绕被扫描的对象走动。至少，这就是我告诉自己的，让自己感觉好一点 😅

除了OpenCL之外，还有许多其他硬件无关的GPGPU框架可供选择，包括Halide、ArrayFire和Numba。因此，总的来说，开源和开放标准生态系统当时是CUDA的公平竞争者，而CUDA还没有建立其护城河。

5、深度学习革命

CUDA/OpenCL带来的新的GPU编程能力，使计算机图形学、科学和医学等领域出现了许多新的应用。然而，深度学习（这是在ChatGPT出现之前我们称之为AI的东西）的普及可能是最显著的成果。

许多人认为，由于AI，GPU已成为中心计算平台。事实上，情况正好相反。正是由于GPU，我们才有了AI。深度卷积神经网络早在90年代就已经存在。2012年，一位研究生Alex Krizhevsky在Ilya Sutskever的鼓励下，使用几块GeForce GPU训练了一个深度卷积神经网络，参加了ImageNet挑战赛。该模型被称为AlexNet，数据集包含120万张属于1000个类别的图像。

结果？它们彻底击败了当时的最先进的计算机视觉模型，展示了惊人的9.4%的准确率提升。这是一场变革。它引发了深度学习革命，所有在计算机视觉、自然语言处理和其他领域的突破都是通过在GPU上训练的深度学习模型实现的。

6、数组编程模型

GPU计算在机器学习领域引起了巨大动荡，而后者则通过大幅改变我们编程GPU的方式进行了反击。编程模型从编写操作数组元素的内核转变为编写对整个数组（张量）进行操作的代码。

原因在于，像Tensorflow或PyTorch这样的深度学习框架并不是受到图形编程的启发，而是受到科学计算框架如NumPy和MATLAB的启发。编程模型与CUDA和OpenCL大不相同。你不是编写操作数组元素的内核，而是编写对整个数组（张量）进行操作的代码。框架将操作分解为较小的部分，可以在GPU上并行执行。这种编程模型称为数组编程，可以追溯到60年代，当时APL和Fortran等语言的发展。

我跳过了当时流行的声明式深度学习框架Caffe。它适合定义大量模型，但不适合表达张量上的任意计算。

这种编程模型有一个巨大的优势。它更容易推理代码，因为你不需要考虑如何并行化计算。你编写对整个数组进行操作的代码，框架会处理其余部分。这使得GPU编程对更广泛的受众更加可访问，因为你不一定要成为GPU编程专家才能编写在GPU上运行的代码。它非常方便，许多GPU编程专家，包括我自己，都转向使用这些框架进行工作。它允许你更简洁地表达你的想法，专注于手头的问题，而不是GPU编程的细节。此外，像PyTorch和TensorFlow这样的框架还带有自动微分引擎，可以自动计算函数的梯度。这对于训练神经网络尤其有用，但也可以应用于其他应用。

这是一个简单的numpy程序。即使不知道numpy，你也能猜出它的作用。它创建了一些数组，对它们进行一些基本操作，并打印结果。

import numpy as np  

# 从Python列表创建一维数组  
array1d = np.array([1, 2, 3, 4, 5])  

# 创建二维数组（矩阵）  
array2d = np.array([[10, 20, 30], [40, 50, 60]])  

# 元素相加  
sum_array = array1d + 5  

# 元素相乘  
product_array = array1d * 2  

# 数组中所有元素的总和  
total_sum = np.sum(array1d)  

# 数组元素的平均值  
mean_value = np.mean(array1d)  

# 访问元素  
print("array1d的第一个元素:", array1d[0])  
print("array2d第0行第1列的元素:", array2d[0, 1])

7、为什么编程GPU很难？

数组编程模型的便利性带来了一个重大缺点。优化代码的性能很难。为了理解原因，我们首先需要考虑为什么一开始优化GPU代码就很困难。

有几个原因导致GPU编程是一项复杂的任务。但主要限制是内存带宽严重限制了GPU，因此GPU架构师引入了许多复杂的机制来隐藏内存访问的延迟并最大化可用带宽的利用率。开发人员需要了解这些机制并编写利用它们的代码。这不是一项容易的任务，因为它需要对GPU架构和内存层次结构的深入了解。

想想以下例子。目前最强大的CPU是AMD EPYC 9965。它提供了惊人的192个核心和384个线程。每插槽的内存带宽约为614 GB/s。然而，它的核心数量与目前最强大的GPU相比显得微不足道，即 NVIDIA B200 。它提供了16,896个CUDA核心和高达8TB/s的内存带宽。

现在，你可能会看到问题：每个CPU核心有大约3.2 GB/s的内存带宽，而每个GPU核心只有约0.47 GB/s的内存带宽。这意味着每个GPU核心必须执行更多的工作来隐藏内存访问的延迟并充分利用可用带宽。消费级GPU的情况甚至更糟，例如RTX 5090有21,760个CUDA核心和1,792 GB/s的内存带宽，每个核心只有约0.082 GB/s的带宽。这意味着GPU必须在每次内存访问时执行更多的计算以达到最佳性能。

在GPU计算世界中，计算能力和内存带宽之间的关系被称为ALU到内存比率，它代表GPU核心每内存访问可以执行的操作数量。对于GPU来说，这个比率比CPU高得多。每个内存访问可以执行几十甚至几百次操作。

所有其他并行计算平台，如TPU、神经处理器和FPGA，也存在同样的问题。每个处理单元的内存带宽始终远低于CPU核心。从2017年到2022年，我在苹果公司为他们的定制神经处理器优化神经网络推理。我们已经发布了Animoji、FaceID、Portrait模式以及许多运行在Apple Vision Pro上的模型。对于每个这些模型，我们必须确保在片上内存和DRAM之间没有数据交换，因为内存带宽是主要瓶颈。

为了克服这个限制，GPU采用了几种技术，如使用共享内存——一组线程共享的小量内存。这允许线程协作并共享数据，而无需访问全局内存，这要慢得多。另一种技术是使用内存合并，它允许线程以最小化内存事务的方式访问内存。这是通过确保线程访问连续的内存位置来实现的，这使得GPU可以一次获取多个数据元素。GPU核心还可以访问更多的寄存器，这些寄存器也可以用来存储中间数据。然而，寄存器是共享的（在一个工作组中的线程），所以如果你使用太多，一些核心会被关闭。

足够的复杂术语！如果你从这篇文章中只记住一件事，那就是：优化GPU程序最有效的方法是每次内存访问执行更多的计算。换句话说，确保数据尽可能长时间留在GPU核心中。让我们记住这一点，回到数组编程模型和它引入的性能问题。

8、我爱PyTorch！它有什么问题？

让我们看看一个简单的CUDA内核，用于执行数组操作如A*B + C。在这里，A、B和C是大数组（张量），操作是逐元素的，例如，[1, 2, 3] * [2, 2, 2] + [1, 1, 1] = [3, 5, 7]

__global__ void array_op(const float *A, const float *B, const float *C, float *D, int N) {  
    int idx = blockIdx.x * blockDim.x + threadIdx.x;  
    if (idx < N) {  
        D[idx] = A[idx] * B[idx] + C[idx];  
    }  
}

这个内核很简单。每个线程通过从输入数组A、B和C中读取相应的元素来计算输出数组D的一个元素。

现在，让我们看看在PyTorch中如何执行相同的操作。

import torch  

A = torch.randn(1000000, device='cuda')  
B = torch.randn(1000000, device='cuda')  
C = torch.randn(1000000, device='cuda')  
D = A * B + C

如果你天真地将PyTorch操作（如逐元素乘法和加法）翻译成CUDA，这实际上是实践中所做的，你会得到两个内核：一个用于乘法，另一个用于加法。运行时会启动一个内核来执行逐元素乘法，将结果存储在一个临时数组A1中，然后启动另一个内核来使用A1和B执行逐元素加法，以产生最终的张量C。

__global__ void array_mul(const float *A, const float *B, float *E, int N) {  
    int idx = blockIdx.x * blockDim.x + threadIdx.x;  
    if (idx < N) {  
        E[idx] = A[idx] * B[idx];  
    }  
}  

__global__ void array_add(const float *E, const float *C, float *D, int N) {  
    int idx = blockIdx.x * blockDim.x + threadIdx.x;  
    if (idx < N) {  
        D[idx] = E[idx] + C[idx];  
    }  
}

你现在可以看到问题：

• 在原始示例中，我们从内存中一次性获取数据，并对其进行两次操作（乘法和加法）。
• 在PyTorch示例中，我们从内存中两次获取数据，并只对它进行一次操作（乘法或加法）。

鉴于我们的程序是完全内存受限的，PyTorch版本将比CUDA版本几乎慢两倍，因为每次内存访问的操作次数减少了一半。如果我们再添加一个未融合的逐元素操作，它将慢三倍，依此类推。

你可能会想，难道不能生成一个同时执行这两个操作的融合内核吗？答案是肯定的，我们可以。事实上，PyTorch和TensorFlow都有机制来做到这一点。然而，这并不是一个容易解决的一般性问题。PyTorch官方支持超过1200个可以在张量上执行的操作。这些操作的可能组合数量是天文数字。其中许多操作并不是逐元素的，例如矩阵乘法、卷积、归约等。在一般意义上解决这个问题是一个复杂的问题。对于PyTorch来说，它更具挑战性，因为它是一个动态框架，即计算图是在代码执行时动态构建的。这使得分析整个计算图并确定哪些操作可以融合变得困难。

直到今天，这个问题仍未在一般意义上得到解决，正如我们在讨论Flash Attention时所看到的那样，这涉及到LLM推理。

9、NVIDIA的统治

深度学习革命极大地改变了GPU编程的格局。数组编程模型使GPU编程对更广泛的受众更加可访问，因为你不一定要成为GPU编程专家才能编写在GPU上运行的代码。然而，它也带来了新的挑战，如优化内存访问模式和融合操作以实现最佳性能。

这为NVIDIA创造了一个强大的护城河。虽然CUDA当时只是众多GPGPU框架之一，但CUDA生态系统为社区提供了更多的东西。例如，它有CUDNN，这是一个高度优化的深度学习原语库，如卷积、池化和归一化。这个库被所有主要的深度学习框架如TensorFlow和PyTorch用于在NVIDIA GPU上实现良好的性能。此外，NVIDIA在优化其硬件以适应深度学习工作负载方面投入了大量资金，例如通过引入Tensor Core，这是一种专门设计用于执行矩阵乘法和卷积的硬件单元。

在深度学习时代，NVIDIA GPU已成为深度学习工作的事实标准。所有主要的深度学习框架如PyTorch和Tensorflow都是基于CUDNN构建的，最初甚至没有提供使用其他后端如OpenCL或ROCm的选项。所有的研究都是在NVIDIA硬件上进行的，因为它是唯一支持他们使用的工具的硬件。这创造了一个强大的网络效应，因为每个人都在使用NVIDIA硬件，所以每个人都优化他们的代码以适应NVIDIA硬件，这使得NVIDIA硬件更加吸引人。

从2010年到现在，我只拥有NVIDIA GPU。尽管一些AMD型号提供了更好的性价比，但进行AI相关工作的需求总是让我偏向于Team Green阵营。

讽刺的是，尽管CUDA很有创新性，但护城河并不是由CUDA本身创造的，而是由NVIDIA工程师团队创造的，他们为深度学习工作负载优化了CUDNN和其他库。根本没有好的算法来在一般意义上优化计算图，因此NVIDIA工程师手动优化了深度学习工作负载中最常见的模式。

有许多尝试来提出一种自动优化计算图的方法，或者至少提出一个通用的、硬件无关的AI堆栈，使优化过程更容易，如Google的XLA、Apache基金会的TVM、LLVM的MLIR或Modular AI的MAX。然而，这些尝试都没有能在足够多的真实用例中击败NVIDIA硬件上的手工优化库如CUDNN，并建立足够的网络效应。

10、人工智能时代——越大越好

历史不会重复自己，但往往会重演。GPU的计算能力引发了深度学习的演变。我们使用了自90年代以来已知的算法，但现在我们可以在更大的数据集上训练更大的模型。同样的事情也发生在LLM上。Transformer架构自2017年以来就已为人所知，但直到2020年，我们才看到了第一个大规模的Transformer模型，如GPT-3和BERT。原因是训练这些模型需要大量的计算能力和内存带宽。OpenAI使用10,000块GPU集群训练GPT-3。最大的模型GPT-3有1750亿个参数，并在570GB的文本数据集上进行训练。训练过程持续了几周，花费了数百万美元（可能还提高了全球温度一度）。

AI如何影响GPU编程格局？实际上并没有太多变化。LLM的训练和推理仍然使用相同的数组编程模型。优化内存访问模式和融合操作以实现良好性能的挑战依然存在。然而，模型的规模急剧增加，这带来了新的挑战，如在多个GPU之间分布模型并优化GPU之间的通信。

模型的巨大规模也给推理带来了新的挑战。这些模型太大，无法放入单个GPU的内存中。例如，GPT-3需要大约700GB的内存来存储模型参数，这比当今最强大的GPU的内存大得多。这导致了诸如模型并行性等技术的发展，其中模型被分割到多个GPU上，以及流水线并行性，其中模型的不同部分在不同的GPU上以流水线方式执行。

11、Flash Attention 的案例

Transformer模型中最关键的操作之一是注意力机制。注意力机制允许模型在进行预测时关注输入序列的不同部分。注意力机制是通过一系列矩阵乘法和softmax操作实现的（见下图右侧的图表）。

Softmax操作涉及计算输入矩阵中每个元素的指数，求和，然后将每个组件除以总和。

优化看起来很有挑战性，对吧？我们如何减少这里的内存访问次数呢？原始的实现方式会涉及从内存中读取输入矩阵，将它们相乘，将结果存储在一个临时矩阵中，然后从内存中读取这个临时矩阵，计算每个元素的指数，求和，最后将每个元素除以总和。这将涉及大量的内存访问，并且会非常慢。而且确实很慢！

然而，之前我提到过，GPU配备了一小部分快速的片上内存，称为共享内存（在硬件术语中称为SRAM——静态随机存取存储器）。这是一种比全局内存（GDDR或HBM）快得多的小容量内存，可以用于存储中间结果。最初的Flash Attention实现是在H100上进行的，H100拥有80GB的HBM内存和每SM 192KB的共享内存。SRAM的速度约为19TB/s，而HBM的速度约为1.5–2.0TB/s。

Flash Attention的作者设计了一种方法，将计算划分为一种方式，使得中间结果可以放入共享内存中，从而能够用更少的全局内存访问次数完成整个注意力计算。这是通过将输入矩阵划分为较小的块，对这些块进行计算（包括矩阵乘法和softmax操作），并流式传输结果回全局内存来实现的。结果是，与原始实现相比，速度有了显著提升。

12、结论

过去二十年间，GPU编程领域发生了巨大的变化。CUDA和OpenCL的引入使GPU编程对更广泛的受众变得可行，从而引发了深度学习革命，反过来也改变了我们编写GPU代码的方式。数组编程模型让编写可以在GPU上运行的代码变得更加容易，但也带来了新的挑战，比如优化内存访问模式以及融合操作以达到最佳性能。

现在，当你成为了一名认证的GPU编程专家，享受最后一张表情包，并让你的GPU高速运转吧！

原文链接：Evolution of GPU Programming

汇智网翻译整理，转载请标明出处