Elliot-CUDA-编程笔记-全-

Elliot CUDA 编程笔记（全）

0：课程介绍 🚀

在本节课中，我们将一起了解这门CUDA编程课程的全貌。课程将引导你学习如何利用GPU进行高性能计算，涵盖从基础概念到实际应用的全过程。

课程概述

欢迎来到这门CUDA编程课程。你将学习如何利用GPU进行高性能计算。课程从深度学习生态系统的概述开始，指导你完成CUDA环境的设置，并回顾必要的C和C++概念。你将探索GPU架构，并编写你的第一个CUDA内核。高级主题包括优化矩阵乘法，以及通过实际应用（如为MNIST数据集实现一个多层感知机）来扩展PyTorch。

本课程由Elliot Olridge创建。

什么是CUDA？

那么，什么是CUDA？CUDA是英伟达推出的“统一计算设备架构”。我的名字是Elliot。我是freeCodeCamp平台的一名讲师，同时也是一名攻读计算机科学学位的学生。在这门课程中，我将为你带来面向深度学习的CUDA知识。但如果你不从事深度学习，请不要因此却步，因为我们仍将涵盖大量通用内容。

以及并行编程的许多其他领域。因此，本课程虽然更偏向于深度学习，但并非专门针对它。这里将涵盖很多内容。

课程最终项目

首先，我将展示最终项目是什么，让你能提前预览我们课程结束时将要构建的内容。然后，我们再从头开始。

在开始任何复杂内容之前，我需要包含一个免责声明。当你观看本课程时，它可能不是完全最新的。如果你在我发布课程十年后才观看，情况可能已大不相同。可能会有更新，新的计算能力可能强大得多，可能发生了许多变化。我不太确定十年后生态系统会变成什么样。但截至2024年，这几乎是你能找到的最好的内容。我尽量让所有内容不完全围绕特定时间点，这样你可以回到这个版本或特定的代码版本，重现所有相同的内容。只是如果你在很久以后观看，操作上可能会有些许不同。

创建课程的初衷

那么，我究竟为什么要创建这门课程呢？首先，许多性能和内核工程岗位需要大量知识和行业经验。要达到能够与顶尖性能工程师竞争的水平非常困难。这些人负责编写像GPT-4、GPT-5等大型模型的训练流程。要优化在大型数据中心或计算集群上进行的大规模神经网络训练和推理，需要大量技能。

本课程旨在减少你手动摸索的部分，虽然仍鼓励你独立探索，但避免了你从零开始、独自深入研究的那种高强度劳动。这是我创建课程的原因之一。

另一个原因是，一般来说，编写GPU内核或在GPU上进行任何编程的目的都是为了让某些东西运行得更快。如果你有一个嵌套循环，例如 for i in range(4): for j in range(4): for k in range(4): ...，并行编程和CUDA允许我们展开这些循环。举例来说，对于 for i in range(4)，你可以将循环中的每个小任务分配到不同的CUDA核心上执行。因此，如果你有10000个CUDA核心，并且你的循环有10000次迭代，那么你可以在GPU上通过单个指令或单个线程有效地完成每次迭代。这就是它允许我们做的事情之一。

你将运用在本课程中学到的GPU架构知识、内核启动配置以及其他很酷的东西，来使代码尽可能快地运行。

最后一个原因是，如今数据量非常庞大。人们常说我们有太多数据，但经过清洗的数据却很少。我汇集了所有其他视频课程、互联网和YouTube上的内容，并将它们整合到一门课程中。我过滤掉了很多无用的、过时的内容，以及一些可能未被充分涵盖的新内容，并将精华提炼到这门杰作中。这包括付费课程涵盖的主题。我虽然没有实际付费，但我查看了它们涵盖的章节，并将其中一些重要概念纳入本课程。

我提供了YouTube视频和所有这些资源的链接，我只筛选了高质量的内容。我浏览了很多这些视频和资源，它们都将放在描述中的GitHub链接里。因此，你需要的所有内容都会在那里，我把所有链接都集中放在了那个链接里。

CUDA的应用场景

那么，CUDA和并行GPU编程有哪些应用场景呢？首先，是图形和光线追踪。你在电子游戏中看到的计算机图形、用户界面等都属于此类。其次，是流体模拟，用于物理和建模，例如引擎动力学。第三，是视频编辑。我现在编辑这个视频就在使用并行计算进行渲染。第四，是加密货币挖矿，很多人可能已经在做了，这会利用你的GPU硬件及其优势来解决挖矿问题。然后是像Blender这样的3D建模软件。当你有很多不同的点需要处理并渲染物体时，本质上与视频编辑类似，只是从2D变成了3D。

最后一个，你可能已经猜到了，就是深度学习。目前CUDA最主要的应用场景，也是本课程将主要涵盖的内容，就是深度学习。我们不会像讲解卷积那样深入，但为了理解如何优化像矩阵乘法这样的算法，我们会进行相当深入的探讨。

课程要求与前提

现在你可能会问，Elliot，这门课程有什么要求或先决条件？有些是学术知识上的，有些则不是。

首先，本课程严格针对英伟达GPU。如果你没有，可以考虑在云上租用最便宜的型号。我建议你在明确拒绝某些云GPU的定价前，先了解一下价格。起初，我对一些云实例（尤其是那些对计算要求不高的实例）的低成本感到惊讶。因此，如果你只有CPU或内存密集型机器，其成本可能远低于带有GPU的机器。GPU实例仍然非常便宜。你可以使用像vast.ai这样的服务，我稍后会详细介绍。你可以用它来获得非常便宜的消费级硬件，通过SSH连接到云端，然后在那上面进行所有实验并完成课程。

你可以继续使用任何英伟达GTX、RTX或数据中心级GPU。基本上所有英伟达显卡都支持，也许那些15年前的低端型号不行，但一般来说，如果你有像GTX 1660这样的显卡，那就没问题。

关于课程先决条件：Python编程知识会有所帮助，因为我们会在底层语言中实现，所以理解整个编程概念才是真正需要的。重申一下，所有这些不同的语言只是语法上的变化。我们将使用基本的微分和向量微积分，如果你已经了解，会使学习更容易。这主要是为了理解反向传播背后的直觉，以及我们从头构建神经网络将用到的一些东西。线性代数肯定会让你的生活更轻松，因为你不需要从头学习基本算法。如果你对矩阵乘法还没有直观理解，或者没有深入接触过，可能跟上进度会有点困难。但矩阵乘法其实很简单，回想起来非常容易理解，只是其中的直觉和优化技巧，如果你之前没有大量实践，可能会有点困难。

如果你真的很在意，我建议你复习一下矩阵转置、矩阵乘法、微积分中的链式法则，以及梯度和导数之间的区别。可能还有我遗漏的一些点，但这些是你入门需要掌握的大致概念。

另外请注意，如果你使用的是Windows机器，可能会稍微困难一些。我确实有一个关于Windows硬件的简短设置指南，但我在这里的所有操作都是在Ubuntu Linux上完成的。这就是我在本地机器上运行的系统，也是我们将用来完成课程的系统。你可以在Windows上使用WSL来模拟Linux系统，或者使用Docker。Docker是一个很棒的工具，可以让你在Windows的终端中启动一个模拟的Linux小机器，并通过它完成所有操作。我认为它通过Windows直接支持英伟达GPU，我不完全确定，还没有测试过。但如果你使用的是Windows机器，我推荐WSL或Docker。

遇到问题怎么办？

如果你在整个过程中遇到错误或问题，我建议你查看GitHub、Stack Overflow、英伟达开发者论坛、PyTorch文档。如果你的问题与任何课程材料相关。因此，如果你需要解决课程材料中未出现的错误，你手头有很多资源。你还可以使用非常强大的语言模型。最近发布了很多语言模型，它们非常擅长解决和处理编码问题。因此，如果其他方法都失败了，我建议你尝试一下。

代码与资源

本课程的所有代码和笔记都保存在描述中的Git仓库里。生态系统会不断变化。因此，如果这个视频不是最新的，GitHub仓库将会是，因为我可以推送并实际进行更改。所以，如果这里有些内容过时了，你可能需要去仓库查看实际的样子，以便正确编写代码，也许那里有更优化的版本。情况会变化，但你明白我的意思。

我建议跟随该仓库，以保持结构化的学习方法。我使用Excalidraw绘制图表，这将有助于说明高层次的概念、我们处理问题的方法，以及如何进行内核级别的优化——从上到下，涵盖所有内容。Excalidraw对于图示说明非常棒，而且完全免费。所有图表都将包含在GitHub仓库和课程中。

你可以随时通过我的Discord服务器联系我，链接也会在GitHub仓库中。你可以通过那里联系我并与社区交流。将会有很多其他学生在学习，会有专门为此设置的频道。因此，如果你遇到困难、想讨论某事，或者只是想在该服务器中进行很酷的聊天，你完全可以加入。

课程范围说明

我想提前说明，本课程不仅仅关于CUDA。我会涵盖一些CUDA之外的内容，包括PyTorch扩展、深入Triton，以及不包含CUDA的C和C++内容，以帮助说明算法原始版本的工作原理。所以，一方面是代码方面。另一方面，我也会提供一些先决条件，或者更确切地说，是对整个深度学习生态系统的良好理解。实际上，下一章的内容之一就是关于整个生态系统如何运作，以及我可以在哪里应用CUDA。如果我只告诉你如何优化内核使其在你的硬件上运行得非常快，却没有给你一些坚实的应用案例，那会有点愚蠢。你可能已经知道应用场景，但如果你只是想学习CUDA，并且可能想了解一些可以应用它的方式，我也会提供这些资源。

课程核心收获

剧透警告，但从本课程中你可能获得的一些收获是：通过实验和研究，你会发现GPU性能的主要瓶颈是内存带宽。在深度学习中，我们有这些巨大、难以理解的矩阵，无法一次性放入芯片上的内存。想象一下，如果你有一个巨大的GPU集群，每个GPU都有非常快的张量核心（这些核心针对深度学习中的张量操作进行了超级优化）。但如果你在多个GPU上进行这些操作，你必须在它们之间交换、混合和管理信息。因此，你最终需要将电子从一个节点发送到另一个节点，再到另一个节点，这中间会有大量的通信发生。所以，你从芯片内部的计算中获得了巨大的速度提升，但当涉及到通信时，实际上存在一个相当大的瓶颈。这是你可能从中学到的一点。

此外，还有芯片上的限制。例如，你有GPU显存，它实际上比芯片上的东西要慢。显存位于实际的核心等组件之外，它必须与核心、芯片上所有共享内存以及所有寄存器进行通信。这最终也会成为一个瓶颈。因此，瓶颈不仅出现在跨多个GPU通信的巨大矩阵上，实际上很多芯片内的通信也是瓶颈。会出现多个瓶颈，但这些只是你稍后通过优化会遇到并能够解决的问题。

另一个关键收获是，能够获取现有的实现并使其更快。很多时候，你会看到一篇新的研究论文发布，里面有一个很酷的算法，但你可能不完全知道它是如何工作的。或者，也许你知道它如何工作，只是想让它更快，并想将其集成到PyTorch中。例如，这实际上是我们将在本课程中做的事情。我们将构建一个算法，优化它，然后将其移植到PyTorch扩展中，以便你可以在Python中调用它，这非常酷。学习如何将自己的研究成果集成到事物中，使其更快，并使其在生产规模下运行，这些都是你将必须做的一些非常重要的事情。

当你开始非常深入地使用CUDA时，另一件事是：Karpathy的llm.c。你们很多人可能听说过这个。如果你在谷歌（不是YouTube）上搜索llm.c，你会找到一个名叫Andrej Karpathy的人，他几乎用C语言从头构建了一个巨大的GPT-2训练流程。它使用了C和CUDA等。内容非常明确。起初我觉得很难理解，对于一个不是超级资深、没有像做了20年CUDA那样经验丰富的人来说，一开始理解起来有点困难。因此，拥有像本课程这样非常好的基础，让你真正理解如何使用CUDA、它的真正优势在哪里以及如何使用它，将使你能够更好地阅读和理解Karpathy的llm.c。这实际上是我创建这门课程的原因之一，为了让人们更容易进入llm.c并理解其中发生的事情。

额外资源

在我的GitHub仓库的Notion文档和GitHub链接中，你会在介绍部分看到这些内容。里面有很多关于CUDA如何工作、Transformer如何工作的酷炫视频，以及一些非常有趣、能让你对所有这些内容保持动力和兴奋的视频。所以，我们有一些技术性的东西，也有一些由Fireship制作的趣味视频。

一般来说，这些是一些很酷的资源。CUDA编程、CUDA模式是一个非常好的服务器，我强烈建议你加入。这只是一个Discord社区，里面有很多真正对CUDA感兴趣的人。我相信Andrej Karpathy也在里面，还有很多很酷的知名程序员和工程师，他们在这里讨论如何让某些内核工作，以及一般的CUDA内容，所以它才叫“CUDA模式”。非常酷的服务器，我强烈建议你也加入那个服务器，以及我的服务器（链接也在GitHub仓库中）。

本节课中，我们一起了解了CUDA编程课程的总体框架、学习目标、应用场景以及所需的基础知识。我们明确了课程将引导你从环境搭建开始，逐步深入到GPU内核编写与优化，最终完成一个实际的深度学习项目。准备好开始你的高性能计算之旅了吗？让我们在下一节中正式启程。

1：深度学习生态系统概述 🧠

在本节课中，我们将概览当前的深度学习生态系统。了解这个全景图有助于你将后续学到的具体技术（如CUDA）与实际应用场景连接起来，明确学习目标，并知道如何将所学技能付诸实践。

上一节我们介绍了课程的整体方向，本节中我们来看看构成现代深度学习世界的各个关键部分。请注意，生态系统的具体工具和技术会快速演变，以下内容旨在为你提供一个理解框架和起点。

研究与应用框架 🧪

我们从最上层、最易用的框架开始。这些是进行深度学习研究和开发的主要工具。

以下是几个核心框架：

PyTorch：由Meta（原Facebook）开发，以其动态计算图和Pythonic的接口而广受欢迎，是当前学术研究和许多工业应用的首选。
TensorFlow：由Google开发，拥有强大的生产部署工具链和广泛的社区支持。
JAX：同样由Google开发，专注于函数式编程和自动微分，在高性能计算和研究中越来越受关注。
MLX：由Apple开发，专为Apple Silicon设备优化的开源框架。

此外，还有一些旨在简化开发的工具：

PyTorch Lightning：一个基于PyTorch的轻量级封装，旨在减少样板代码。例如，设置TF32精度以启用张量核心计算这类优化操作，在PyTorch Lightning中会被自动处理。

生产与推理优化 🚀

当模型需要部署到实际环境中提供服务时，就进入了生产阶段。这主要涉及训练和推理两个环节，有些工具同时支持两者，有些则专门优化其一。

以下是生产与推理环节的关键工具：

vLLM：一个专为大型语言模型（LLM）推理设计的高吞吐量、内存高效的服务系统。它在性能基准测试中常与TensorRT-LLM进行比较。
TensorRT / TensorRT-LLM： NVIDIA推出的高性能深度学习推理SDK和运行时。TensorRT-LLM专门针对LLM推理进行了大量底层GPU硬件优化。
Triton：由OpenAI开发，后来开源。它是一个用于编写高效GPU内核的编程语言和编译器。其核心思想源于论文《Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations》中提出的分块计算。公式上，这类似于将大矩阵乘法 C = A @ B 分解为多个小块的并行计算，以充分利用GPU的并行架构。
torch.compile： PyTorch 2.0引入的特性。只需一行代码 model = torch.compile(model)，即可将PyTorch的动态图转换为静态图，并应用内核融合等优化，通常能带来显著的性能提升。
TorchScript： PyTorch的早期静态图导出工具，用于将模型转换为可在非Python环境中运行的格式。
ONNX Runtime：建立在ONNX格式之上的高性能推理引擎。ONNX（Open Neural Network Exchange）是一个开放的模型格式标准，允许你在不同框架（如PyTorch, TensorFlow）间转换和运行模型。导出ONNX模型的代码示例如下：
```
# PyTorch 导出 ONNX
torch.onnx.export(model, dummy_input, "model.onnx")
```
Detectron2：由Meta（原Facebook）开发的计算机视觉库，提供了目标检测、图像分割等任务的先进算法实现。

底层硬件编程 ⚙️

这是本课程的核心领域，涉及直接与GPU硬件交互的编程。

以下是主要的底层编程平台：

CUDA：全称Compute Unified Device Architecture，是NVIDIA推出的并行计算平台和编程模型。它是我们本课程的重点。
ROCm： AMD推出的开源软件平台，为AMD GPU提供类似CUDA的生态支持。
OpenCL：一个开放的、跨厂商的通用并行计算框架，支持CPU、GPU、DSP等多种硬件。

边缘计算与嵌入式系统 📱

边缘计算指的是在数据产生的源头（如物联网设备、自动驾驶汽车）进行本地计算，而非全部上传到云端。

以下是适用于边缘设备的工具：

TensorFlow Lite： TensorFlow的轻量级版本，专为移动和嵌入式设备设计。
PyTorch Mobile： PyTorch针对移动端的优化版本。
Core ML： Apple的机器学习框架，用于在Apple设备（iOS, macOS, watchOS等）上集成模型。

易用性与工具链 🛠️

这部分包含了一些能极大提升开发效率的库和服务。

以下是一些提高易用性的工具：

Fast.ai：一个旨在让深度学习更易用的库，建立在PyTorch之上，提供了高级API和最佳实践。
Weights & Biases：一个实验跟踪工具，可以方便地记录和可视化训练过程中的指标（如损失、准确率），帮助管理机器学习项目。
云服务提供商：提供强大的计算资源，是开发生态的重要组成部分。
- AWS：提供EC2实例（通用虚拟机）和SageMaker（集成的ML平台，支持从数据标注、训练到部署的全流程）。
- Google Cloud：提供Vertex AI和Compute Engine VM。
- Microsoft Azure：提供Azure Machine Learning等服务。
- RunPod / Lambda Labs：提供按小时租用GPU的服务器，通常比大型云厂商更具价格优势，适合需要特定GPU（如RTX 4090, H100）的用户或初创公司。

编译器基础设施 🏗️

编译器在将高级代码转换为高效机器指令的过程中扮演着关键角色。

以下是深度学习领域重要的编译器项目：

LLVM：一个模块化、可重用的编译器及工具链集合。许多现代编译器（包括CUDA的NVCC）都使用或借鉴了LLVM的技术。
MLIR：多级中间表示，是LLVM项目的一部分，旨在解决构建领域特定编译器（如AI编译器）的复杂性。由Chris Lattner（也是Swift和LLVM的创建者）等人推动。
NVCC： NVIDIA CUDA Compiler，是编译CUDA C/C++代码的工具。
XLA：加速线性代数，是TensorFlow和JAX使用的编译器，用于优化线性代数计算。

模型与数据集社区 🤗

最后，但绝非最不重要的，是模型和数据的集散地。

以下是Hugging Face平台的核心组成部分：

Models：托管了数以万计的开源预训练模型，涵盖自然语言处理、计算机视觉、音频等多个领域。
Datasets：提供了大量公开可用的数据集，用于训练和评估模型。
Spaces：允许用户轻松部署和分享机器学习演示应用。

本节课中我们一起学习了深度学习生态系统的全貌，从高层的易用框架（如PyTorch），到生产部署的优化工具（如TensorRT, vLLM），再到本课程的核心——底层CUDA编程。我们还了解了边缘计算、云服务、编译器以及Hugging Face这样的核心社区。理解这个生态系统将帮助你在后续深入学习具体技术时，始终保持清晰的视野和明确的目标。下一节，我们将开始深入CUDA编程的具体细节。

2：CUDA环境搭建 🛠️

在本节课中，我们将学习如何在Windows和Ubuntu系统上搭建CUDA编程环境。我们将分步完成系统配置、必要组件的安装以及CUDA工具包的设置，并最终运行一个简单的测试程序来验证环境是否配置成功。

Windows系统配置

上一节我们介绍了课程概述，本节中我们来看看如何在Windows系统上开始配置CUDA环境。

首先，我们需要在Windows上启用几个关键功能。以下是具体步骤：

以管理员身份打开终端。
导航到“启用或关闭Windows功能”设置。
在列表中，找到并勾选“Hyper-V”。
继续找到并勾选“虚拟机平台”。
确保“适用于Linux的Windows子系统”也已勾选启用。

为了确保这些功能正常工作，您可能需要在计算机的BIOS/UEFI设置中启用虚拟化技术。

安装WSL与Ubuntu

完成Windows功能启用后，下一步是安装WSL（Windows Subsystem for Linux）和一个Linux发行版。

在管理员终端中，输入以下命令来安装WSL：

wsl --install

系统会列出可用的Linux发行版。我们可以选择安装Ubuntu：

wsl --install -d Ubuntu

安装过程可能需要一些时间。安装完成后，系统可能会提示您重启计算机以使更改生效。

重启后，您可能会看到一个命令行界面。按照提示设置您的用户名和密码。成功登录后，您就进入了一个模拟的Linux环境。

配置Ubuntu环境

现在我们已经进入了WSL下的Ubuntu环境，接下来需要更新系统并安装一些必要的软件包。

首先，运行以下命令来更新软件包列表并升级所有已安装的软件：

sudo apt update && sudo apt upgrade -y

这个过程可能需要一些时间。更新完成后，安装一些后续会用到的工具，如 wget、curl 和 git：

sudo apt install wget curl git -y

接着，安装Python 3和pip，这对于许多开发任务都是必需的：

sudo apt install python3 python3-pip -y

下载并安装CUDA工具包

系统基础环境准备就绪后，我们就可以安装核心的CUDA工具包了。

打开浏览器，搜索“CUDA Toolkit download”，进入NVIDIA官方网站。选择与您的系统匹配的最新版本（例如12.5或12.6）。对于WSL环境，请选择：

操作系统：Linux
架构：x86_64
发行版：WSL-Ubuntu
安装类型：runfile（本地）

页面上会提供下载和安装命令。通常包括两个步骤。首先，使用 wget 下载安装文件：

wget https://developer.download.nvidia.com/compute/cuda/12.5.0/local_installers/cuda_12.5.0_555.42.02_linux.run

然后，运行下载的 .run 文件进行安装：

sudo sh cuda_12.5.0_555.42.02_linux.run

在安装向导中，您只需选择安装“CUDA Toolkit”即可，驱动部分通常不需要在WSL内单独安装。

配置环境变量

安装完成后，安装程序会提示需要将CUDA路径添加到环境变量中。我们需要手动配置。

使用文本编辑器（如vim或nano）打开用户主目录下的 .bashrc 文件：

vim ~/.bashrc

在文件末尾添加以下几行，请根据您安装的CUDA版本（例如12.5）修改路径：

export CUDA_HOME=/usr/local/cuda-12.5
export PATH=$CUDA_HOME/bin:$PATH
export LD_LIBRARY_PATH=$CUDA_HOME/lib64:$LD_LIBRARY_PATH

保存并退出编辑器。然后，让配置立即生效：

source ~/.bashrc

验证安装

环境变量配置好后，我们可以验证CUDA是否安装成功。

首先，检查NVIDIA CUDA编译器 nvcc 的版本：

nvcc --version

如果成功，会显示版本信息。接着，运行 nvidia-smi 命令来查看GPU状态：

nvidia-smi

这个命令会显示显卡驱动版本、CUDA版本以及GPU的详细信息。如果这两个命令都能正常运行，说明CUDA环境基本配置成功。

编写并运行测试程序

最后，我们通过一个简单的“Hello World”程序来测试整个开发链路是否通畅。

创建一个测试目录并进入：

mkdir cuda_setup_test && cd cuda_setup_test

创建一个名为 main.cu 的CUDA源文件：

vim main.cu

在文件中输入以下测试代码：

#include <iostream>
#include <cuda_runtime.h>

int main() {
    std::cout << "Hello World from CUDA!" << std::endl;
    return 0;
}

使用 nvcc 编译器编译这个文件，生成可执行文件：

nvcc -o main main.cu

运行生成的可执行文件：

./main

如果终端成功输出“Hello World from CUDA!”，则证明您的CUDA开发环境已经完全配置成功，可以开始进行CUDA编程了。

Ubuntu原生系统安装（补充说明）

上一节我们详细介绍了在Windows WSL下的安装过程，本节中我们简要看看在原生Ubuntu系统上的安装有何不同。

对于原生Ubuntu系统，步骤更为直接。您同样需要访问NVIDIA CUDA Toolkit下载页面。根据您的Ubuntu版本（如22.04）和架构选择对应的deb（网络）安装包。

页面上会提供类似以下的安装命令，通常包含下载、安装密钥、更新仓库和安装工具包几个步骤：

wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2204/x86_64/cuda-keyring_1.1-1_all.deb
sudo dpkg -i cuda-keyring_1.1-1_all.deb
sudo apt update
sudo apt install cuda-toolkit-12-5

安装完成后，同样需要配置环境变量（步骤与WSL中相同），并运行 nvcc --version 和 nvidia-smi 进行验证。如果遇到问题，重启系统通常是有效的解决方法。

本节课中我们一起学习了在Windows（通过WSL）和Ubuntu系统上搭建CUDA编程环境的完整流程。我们从系统基础配置开始，逐步完成了WSL启用、Ubuntu安装、系统更新、CUDA工具包下载安装、环境变量配置，并最终通过编译和运行一个测试程序验证了整个开发环境。现在，您已经拥有了一个可以开始高性能GPU编程的工作环境。

3：C/C++ 语言回顾

在本章节中，我们将回顾C和C++编程语言的核心概念，为后续的CUDA编程打下坚实基础。理解这些基础概念，特别是内存管理和指针，对于高效地进行GPU编程至关重要。

概述

为了真正理解如何使用CUDA，我们首先需要掌握C和C++。本节课程并非专门教授C/C++，而是提供关键知识点的回顾和资源指引。我们将从指针开始，逐步深入到自定义类型、类型转换、宏、编译器和调试器等主题。

学习资源

对于C和C++的初学者，以下是一些推荐的学习资源：

Learn C++：网站 learncpp.com 提供了结构化的C++学习路径。
《C Programming: A Modern Approach》：这是一本广受好评的、适合初学者的C语言教材。
freeCodeCamp：该平台提供了大量关于C和C++编程的免费博客和教程。
W3Schools：提供了易于阅读和理解的C和C++语法介绍及示例。

建议新手逐一学习这些资源中的主题，并完成相关的练习。即使某些内容在本课程中不会直接使用，掌握这些底层知识也能帮助你更好地分析和解决未来可能遇到的复杂问题。

上一节我们介绍了学习资源，本节中我们来看看C/C++中一个核心且强大的概念：指针。

指针

指针是存储变量内存地址的变量。理解指针对于管理内存和构建高效的数据结构至关重要。

示例 1：基础指针

以下是一个简单的指针示例：

int x = 10; // 变量 x 存储数据 10
int *pointer = &x; // 指针变量 pointer 存储 x 的内存地址
printf("内存地址: %p\n", (void*)pointer); // 打印指针（地址）
printf("存储的值: %d\n", *pointer); // 解引用指针，获取地址处的值（10）

& 运算符用于获取变量的地址。
* 在声明中表示指针类型，在表达式中表示解引用操作（获取指针指向地址的值）。

示例 2：多级指针

指针可以指向另一个指针，形成多级间接引用。

int value = 42;
int *pointer1 = &value; // pointer1 指向 value
int **pointer2 = &pointer1; // pointer2 指向 pointer1
int ***pointer3 = &pointer2; // pointer3 指向 pointer2
printf("值: %d\n", ***pointer3); // 三级解引用，最终得到 42

每一级指针都存储着下一级指针（或最终数据）的地址。解引用操作（*）用于逐级向下访问。

示例 3：Void 指针

Void 指针 (void*) 是一种通用指针类型，可以指向任何数据类型的数据，但在使用前必须进行类型转换。

int num = 10;
float f_num = 3.14;
void *void_ptr;

void_ptr = &num; // void_ptr 指向整数
printf("整数值: %d\n", *((int*)void_ptr)); // 转换为 int* 后解引用

void_ptr = &f_num; // void_ptr 现在指向浮点数
printf("浮点数值: %f\n", *((float*)void_ptr)); // 转换为 float* 后解引用

malloc 函数就返回一个 void* 指针，调用者需要将其转换为具体的指针类型。

示例 4：NULL 指针

NULL 指针不指向任何有效的内存地址。检查指针是否为 NULL 可以避免程序崩溃（如段错误）。

int *ptr = NULL; // 初始化为 NULL

if (ptr == NULL) {
    printf("指针为 NULL，无法解引用。\n");
}

ptr = (int*)malloc(sizeof(int)); // 动态分配内存
if (ptr != NULL) {
    *ptr = 42; // 安全地使用指针
    printf("值: %d\n", *ptr);
    free(ptr); // 释放内存
    ptr = NULL; // 再次设为 NULL，避免“释放后使用”错误
}

示例 5：指针与数组

数组名在大多数情况下可以视为指向其第一个元素的指针。

int array[] = {12, 23, 34, 45, 56};
int *ptr = array; // ptr 指向数组首元素

printf("第一个元素: %d\n", *ptr); // 解引用得到 12

for (int i = 0; i < 5; i++) {
    printf("值[%d]: %d, 地址: %p\n", i, *ptr, (void*)ptr);
    ptr++; // 指针算术：移动到下一个整数（地址增加 4 字节）
}

指针递增 (ptr++) 会根据指针类型的大小移动。对于 int*，每次增加 4 字节（假设 int 为 32 位）。

示例 6：指针数组（模拟矩阵）

可以使用指针数组来模拟二维数据结构。

int array1[] = {1, 2, 3, 4};
int array2[] = {5, 6, 7, 8};
int *matrix[] = {array1, array2}; // matrix 是指针数组

for (int i = 0; i < 2; i++) {
    for (int j = 0; j < 4; j++) {
        printf("%d ", *(matrix[i] + j)); // 等价于 matrix[i][j]
    }
    printf("\n");
}

matrix 存储了两个一维数组的起始地址，通过它可以访问所有元素。

上一节我们深入探讨了指针，本节中我们将了解如何创建和使用自定义数据类型。

自定义类型

C语言允许使用 typedef 关键字创建自定义的数据类型别名，这能提高代码的可读性和可维护性。

示例：自定义 `size_t` 和结构体

标准库定义了 size_t 类型，通常用于表示对象大小或数组索引。它是一个无符号长整型，确保能容纳大尺寸对象。

#include <stdio.h>
#include <stddef.h> // 包含 size_t 的定义

int main() {
    int array[] = {10, 20, 30, 40, 50};
    size_t array_length = sizeof(array) / sizeof(array[0]); // 计算数组长度
    printf("数组长度: %zu\n", array_length); // 使用 %zu 格式化 size_t
    printf("size_t 的大小: %zu 字节\n", sizeof(size_t));
    return 0;
}

我们也可以定义自己的结构体类型：

typedef struct {
    float x;
    float y;
} Point; // 定义了一个名为 Point 的新类型

int main() {
    Point p = {3.5, 2.8};
    printf("点坐标: (%.1f, %.1f)\n", p.x, p.y);
    printf("Point 类型大小: %zu 字节\n", sizeof(Point)); // 通常是 8 字节 (两个 float)
    return 0;
}

上一节我们创建了自定义类型，本节中我们来看看如何在不同类型之间进行转换。

类型转换

类型转换允许你将一种数据类型的值转换为另一种类型。在C++中，有几种不同的类型转换运算符，static_cast 是最常用且最安全的一种。

#include <iostream>
int main() {
    float f = 69.69f;
    int i = static_cast<int>(f); // 将 float 转换为 int，小数部分被截断
    std::cout << "浮点数: " << f << std::endl;
    std::cout << "转换为整数: " << i << std::endl; // 输出 69

    int num = 69;
    char c = static_cast<char>(num); // 将 int 转换为 char (ASCII 码)
    std::cout << "整数: " << num << std::endl;
    std::cout << "转换为字符: " << c << std::endl; // 输出 'E' (ASCII 69)
    return 0;
}

上一节我们进行了类型转换，本节中我们来看看如何使用宏和全局变量来简化代码。

宏与全局变量

宏 (#define) 是预处理器指令，用于定义常量或简单的函数式代码片段，在编译前进行文本替换。

#include <stdio.h>

#define PI 3.14159
#define AREA(r) (PI * (r) * (r)) // 带参数的宏，计算圆面积

#ifndef RADIUS // 如果 RADIUS 未定义，则定义它
    #define RADIUS 7
#endif

int main() {
    int radius = RADIUS;
    printf("半径为 %d 的圆面积: %.2f\n", radius, AREA(radius));
    return 0;
}

宏有助于减少魔法数字，并使全局常量的修改变得容易。但要注意，带参数的宏可能因运算符优先级导致意外结果，所以参数通常用括号括起来。

上一节我们使用了宏，本节中我们简要了解一下代码是如何从文本变成可执行文件的。

编译器简介

编译器（如 GCC 和 G++）负责将人类可读的C/C++源代码翻译成机器可执行的二进制指令。这个过程大致包括：预处理（处理宏等）、编译（生成汇编代码）、汇编（生成目标文件）、链接（合并库文件和目标文件，生成最终可执行文件）。

理解编译过程有助于调试，但编写功能代码通常不需要深入了解其内部数学原理。我们主要通过命令行调用编译器。

上一节我们提到了编译器，本节中我们来看看如何利用 make 工具自动化编译过程。

Makefile

Makefile 是一个包含构建规则的文件，用于自动化编译过程。它定义了目标、依赖项和生成命令。

一个简单的 Makefile 示例：

CC = gcc
CFLAGS = -Wall
TARGET = myprogram
SRC = main.c utils.c

$(TARGET): $(SRC)
	$(CC) $(CFLAGS) -o $(TARGET) $(SRC)

clean:
	rm -f $(TARGET)

CC 和 CFLAGS 是变量。
$(TARGET): $(SRC) 表示目标 myprogram 依赖于 main.c 和 utils.c。
下一行缩进的是生成目标的命令。
clean 是一个伪目标（通常用 .PHONY: clean 声明），用于清理生成的文件。

使用命令 make 会默认构建第一个目标，make clean 会执行清理操作。Makefile 能极大提升多文件项目的管理效率。

上一节我们自动化了编译，本节中我们介绍最后一个重要工具：调试器。

调试器 (GDB)

调试器（如 GDB）允许你逐步执行程序、检查变量状态、设置断点，是比插入大量 printf 语句更强大的调试工具。

一些常用的 GDB 命令包括：

gdb ./program：启动 GDB 并加载可执行文件。
break main 或 b main：在 main 函数开头设置断点。
run 或 r：开始运行程序。
next 或 n：执行下一行代码（不进入函数内部）。
step 或 s：执行下一行代码（会进入函数内部）。
print variable 或 p variable：打印变量的值。
continue 或 c：继续运行直到下一个断点或程序结束。
backtrace 或 bt：显示函数调用栈。
quit 或 q：退出 GDB。

掌握 GDB 的基本用法能显著提高定位和修复代码逻辑错误或内存问题的效率。

总结

本节课中我们一起回顾了CUDA编程所需的C/C++核心知识。我们从指针的基础概念和多级应用开始，学习了如何管理内存地址。接着，我们探讨了自定义数据类型的创建，以及如何进行安全的类型转换。我们还了解了如何使用宏来定义常量和小型代码块，并简要介绍了编译器的工作流程。最后，我们掌握了使用 Makefile 自动化编译过程以及使用 GDB 调试器进行代码调试的基本方法。这些概念和工具是进行高效、可靠的CUDA编程的重要基石。

4：GPU硬件入门 🚀

在本章中，我们将了解不同类型的计算硬件，特别是CPU与GPU的区别，以及GPU为何在并行计算任务中如此高效。我们还将介绍一些CUDA编程的基本术语和概念。

硬件类型概述

首先，我们来比较几种主要的计算硬件：CPU、GPU、TPU和FPGA。

以下是它们的主要特点：

CPU (中央处理器)
- 用途：通用计算。
- 核心：数量少，但每个核心的时钟频率很高。
- 片上内存：缓存容量大，用于预加载数据以减少访问主内存的延迟。
- 特点：延迟低，旨在以最快速度完成单个任务并返回结果。
- 吞吐量：相对较低，每秒能处理的操作数（OPS）有限，尤其是在处理简单数学运算（如矩阵乘法）时。
GPU (图形处理器)
- 用途：专用并行计算。
- 核心：数量极多，但每个核心的时钟频率较低。
- 片上内存：缓存较小，但拥有专用的显存 (VRAM)，访问带宽极高（可达每秒数百GB）。
- 特点：延迟较高，但吞吐量巨大，专为同时处理海量简单任务而优化。
- 速度优势：核心数量远超CPU，因此在并行任务上速度极快。例如，若有12000个任务，6000个GPU核心只需2轮操作，而4个CPU核心则需要3000轮。
TPU (张量处理器)
- 用途：专为现代深度学习设计，高效执行张量运算（线性代数、矩阵乘法）。
- 特点：速度极快，但价格昂贵且高度专业化，通常不是消费级硬件。
FPGA (现场可编程门阵列)
- 用途：可通过编程定制硬件逻辑，实现特定任务的极致优化。
- 特点：提供精细的控制，延迟极低，吞吐量极高，但价格昂贵，具备模块化特性。

GPU发展简史

上一节我们介绍了各类硬件，本节中我们来看看GPU的发展历程。了解其演进有助于理解当前架构的设计思路。

GPU性能的提升主要依赖于核心数量的不断增加以及架构的持续优化。从早期的GeForce系列，到Tesla、Fermi、Kepler架构，再到后来的Maxwell、Pascal，性能逐步提升。从Volta架构开始，GPU在深度学习领域的计算能力（如浮点运算性能）实现了飞跃。随后的Turing、Ampere（当前主流架构，如RTX 30/40系列）、Hopper（如H100）以及最新的Blackwell架构，性能更是达到了新的高度。

例如，Volta架构已能提供约6 TFLOPs的双精度计算能力，而现代的Ampere架构显卡在CUDA核心上进行单精度矩阵乘法（通过cuBLAS库）时，可达到约20 TFLOPs以上的性能。

GPU为何适合深度学习？

我们已经看到GPU拥有海量核心，但为何这种架构特别适合深度学习等任务呢？关键在于其设计哲学。

CPU的设计目标是快速完成复杂、串行的任务。其芯片上大部分面积被大型控制单元和缓存占据，留给计算核心的空间有限。

相比之下，GPU的设计目标是高吞吐量。其芯片上大部分面积是大量简单的计算核心，以及为这些核心服务的高速缓存和显存控制器。控制单元相对简单小巧。

这就像一个拼图游戏：

CPU 如同几个高手，一次只能拼几块，但能处理复杂的拼图策略（复杂指令）。
GPU 如同成千上万的工人，每人一次只拼简单的一小块（简单指令），但可以同时进行，最终快速完成整幅拼图。

深度学习中的许多运算（如矩阵乘法和卷积）正是这种可以高度并行化的“拼图”任务，因此GPU能发挥巨大优势。

CUDA编程核心概念

了解了硬件背景后，我们现在进入CUDA编程的核心概念。这些术语将贯穿整个学习过程。

以下是CUDA编程中的基本术语：

主机 (Host) 与设备 (Device)
- 主机 (Host)：指CPU及其内存。负责运行常规的C/C++函数。
- 设备 (Device)：指GPU及其显存。负责运行并行计算函数（内核）。
- 性能关注点：主机关注延迟（完成任务的速度），设备关注吞吐量（单位时间完成的任务量，如每秒渲染的像素数）。
内核 (Kernel)
- 这是在GPU上执行的并行函数。在代码中，通过 __global__ 关键字来定义。
- 一个内核看起来像一个串行程序，但它会被成千上万个线程同时执行。
- 注意：此处的“内核”不同于操作系统内核、卷积核或玉米粒，特指GPU上的并行函数。

线程层次结构：Thread, Block, Grid
- 这是CUDA并行编程模型的核心，我们将在下一章详细展开。简单来说：
  - 线程 (Thread)：最基本的执行单元。
  - 线程块 (Block)：一组线程的集合，块内的线程可以协作。
  - 网格 (Grid)：由多个线程块组成。
- 当启动一个内核时，你需要指定执行这个内核的网格和线程块的维度。
通用矩阵乘法 (GEMM)
- 这是深度学习等领域的核心运算。其公式比简单的矩阵乘法更通用：
  C = α * (A * B) + β * C
  其中 α 和 β 是标量，A, B, C 是矩阵。
- SGEMM 特指单精度浮点数 (float) 的GEMM运算。此外还有半精度 (FP16)、双精度 (FP64) 等版本。

典型的CUDA程序流程

最后，我们来看一个典型的CUDA程序是如何工作的。这将把前面所有的概念串联起来。

一个典型的CUDA程序遵循以下流程：

主机端内存分配：在CPU内存中使用 malloc 或 new 分配空间。
数据拷贝至设备：将数据从主机内存复制到GPU显存。
启动内核：在GPU上调用内核函数，指定网格和线程块的配置，执行并行计算。
结果拷贝回主机：将计算结果从GPU显存复制回CPU内存。
后续处理：在CPU上对结果进行进一步处理或输出。

这个流程可以反复迭代，形成复杂的工作流：CPU准备数据 -> GPU高速计算 -> CPU处理结果。

本章总结

在本节课中，我们一起学习了：

硬件对比：了解了CPU（低延迟、通用）、GPU（高吞吐量、并行）、TPU（专用张量计算）和FPGA（可编程硬件）的特点。
GPU优势：明白了GPU通过海量简单核心实现高并行度，特别适合矩阵运算等任务，因此成为深度学习的首选硬件。
核心术语：掌握了CUDA编程的基础概念，包括主机/设备、内核、线程/块/网格层次结构以及GEMM运算。
编程流程：熟悉了典型的CUDA程序从数据准备、传输、并行计算到结果回收的基本步骤。

接下来，我们将进入实践环节，开始编写第一个简单的CUDA内核（例如向量加法），亲身体验这些概念是如何在代码中实现的。

5：编写你的第一个内核 🚀

在本节课中，我们将学习CUDA编程的核心概念，包括如何编写和启动你的第一个CUDA内核。我们将从了解GPU硬件规格开始，逐步深入到CUDA的软件抽象层次、内存管理、内核编写以及性能分析工具的使用。

概述：从硬件规格到软件抽象

上一节我们介绍了GPU的基本概念。本节中，我们来看看如何获取和理解你GPU的具体规格，这是编写高效CUDA代码的第一步。

你可以通过查阅维基百科或使用CUDA工具来获取GPU的详细信息。例如，Pascal架构用于GTX 1080/1070，Ampere架构用于RTX 30系列，而最新的Blackwell架构则代表了最先进的技术。了解你的GPU的计算能力（Compute Capability，如8.6）至关重要，因为它决定了支持哪些CUDA功能。

在终端中，你可以通过编译并运行CUDA样例中的deviceQuery程序来获取本地GPU的详细信息。

cd /usr/local/cuda/samples/1_Utilities/deviceQuery
sudo make
./deviceQuery

输出会显示GPU型号、驱动版本、计算能力等关键信息。计算能力是一个核心指标，你可以在CUDA官方文档的“Compute Capability”章节查询特定功能（如线程块簇）所需的最低版本。

CUDA基础：运行时与内存模型

了解了硬件后，我们来看看CUDA编程的基本运行时流程。

典型的CUDA程序流程如下：

在主机（CPU）内存中定义和初始化数据。
在设备（GPU）内存中分配空间。
将数据从主机内存复制到设备内存。
在GPU上启动内核（Kernel）执行计算。
将结果从设备内存复制回主机内存。
释放设备内存。

以下是相关的命名约定和关键函数：

变量命名：通常使用h_前缀表示主机变量，d_前缀表示设备变量。例如：h_matrixA, d_matrixA。
函数限定符：
- __global__：定义内核函数，由CPU调用，在GPU上执行。
- __device__：定义设备函数，由内核或其他设备函数调用。
- __host__：定义主机函数（通常省略）。
设备内存管理：
- cudaMalloc(&d_ptr, size)：在GPU上分配内存。
- cudaMemcpy(dst, src, size, kind)：在主机与设备间复制数据。kind可以是cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost或cudaMemcpyDeviceToDevice。
- cudaFree(d_ptr)：释放GPU内存。

CUDA代码通过NVCC编译器编译。主机代码被编译为CPU指令，而设备代码（内核）则被编译为PTX（并行线程执行）指令，最终在GPU上转换为特定硬件的着色器汇编指令。

CUDA层次结构：网格、块与线程 🧱

CUDA使用一个分层的软件抽象来实现大规模并行。理解这个层次结构是编写内核的关键。

你可以将CUDA的执行模型想象成一个三维的网格（Grid）。网格中包含许多线程块（Block），每个线程块又包含许多线程（Thread）。线程是执行计算的基本单位。

这个层次结构通过以下内置变量在核函数中访问：

gridDim：网格的维度（每个维度上有多少个块）。
blockIdx：当前线程块在网格中的索引（坐标）。
blockDim：线程块的维度（每个维度上有多少个线程）。
threadIdx：当前线程在线程块中的索引（坐标）。

一个线程在整个网格中的全局ID可以通过这些变量计算得出。例如，在一维情况下：

int global_id = blockIdx.x * blockDim.x + threadIdx.x;

线程被分组为线程束（Warp），通常是32个线程一组。线程束是GPU调度和执行的基本单位。一个块内的线程可以通过共享内存（Shared Memory/L1缓存） 进行快速通信，其带宽远高于全局内存（GPU显存）。整个网格中的所有线程都可以访问全局内存。

实战：向量加法内核 ✨

现在，我们将理论付诸实践，编写一个简单的向量加法内核，并比较CPU和GPU的实现。

CPU版本的向量加法使用循环顺序处理每个元素：

for (int i = 0; i < n; i++) {
    c[i] = a[i] + b[i];
}

GPU内核版本则“展开”这个循环，让每个线程处理一个独立的加法操作：

__global__ void vectorAdd(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 blockSize = 256;
int numBlocks = (n + blockSize - 1) / blockSize; // 确保覆盖所有元素
vectorAdd<<<numBlocks, blockSize>>>(d_a, d_b, d_c, n);

if (i < n) 是必要的边界检查，因为线程总数可能略大于向量长度。

在性能对比中，即使是这个简单的内核，GPU也能带来上百倍的加速。需要注意的是，使用三维索引的内核可能会因为更复杂的索引计算而比一维内核稍慢，因此应仅在算法需要时使用多维索引。

深入：矩阵乘法内核 ⚡

矩阵乘法是许多科学计算和机器学习应用的核心。我们首先实现一个基础（朴素）版本。

一个 M x K 的矩阵 A 与一个 K x N 的矩阵 B 相乘，得到一个 M x N 的矩阵 C。CPU的朴素实现使用三层嵌套循环：

for (int i = 0; i < M; i++) {
    for (int j = 0; j < N; j++) {
        float sum = 0;
        for (int k = 0; k < K; k++) {
            sum += a[i * K + k] * b[k * N + j];
        }
        c[i * N + j] = sum;
    }
}

对应的GPU内核让每个线程计算输出矩阵C中的一个元素：

__global__ void matrixMultiplyNaive(float* a, float* b, float* c, int M, int N, int K) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < M && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < K; k++) {
            sum += a[row * K + k] * b[k * N + col];
        }
        c[row * N + col] = sum;
    }
}

内核启动时，我们使用二维的网格和块来映射输出矩阵的行和列：

dim3 threadsPerBlock(16, 16);
dim3 numBlocks((N + threadsPerBlock.x - 1) / threadsPerBlock.x,
               (M + threadsPerBlock.y - 1) / threadsPerBlock.y);
matrixMultiplyNaive<<<numBlocks, threadsPerBlock>>>(d_a, d_b, d_c, M, N, K);

这个朴素内核虽然直观，但效率不高。后续优化的关键思想是分块（Tiling），即将矩阵分成小块，利用共享内存来减少对全局内存的访问，从而极大提升性能。

性能剖析：使用Nsight Compute 🔍

编写内核后，我们需要工具来分析和优化其性能。NVIDIA Nsight Compute 是一个强大的GPU内核性能分析器。

首先，我们可以使用NVTX（NVIDIA Tools Extension）在代码中标记范围，以便在时间线上直观看到不同阶段（如内存分配、拷贝、内核执行）的耗时。

#include <nvtx3/nvtx3.hpp>
...
nvtxRangePushA("Memory Allocation");
// ... 内存操作
nvtxRangePop(); // Memory Allocation
nvtxRangePushA("Kernel Execution");
myKernel<<<...>>>(...);
cudaDeviceSynchronize();
nvtxRangePop(); // Kernel Execution

使用nsys命令行工具生成性能分析报告：

nsys profile -o my_report ./my_cuda_program

然后使用Nsight Compute GUI打开生成的.nsys-rep文件。在“CUDA HW”视图中，你可以看到内核和内存拷贝的时间线。点击某个内核，选择“Profile Kernel”并进行“PM Sampling”（性能指标采样），可以获取详细的硬件性能指标，如：

计算吞吐量与内存吞吐量的百分比（接近100%为佳）。
内存带宽利用率（GB/s）。
L1/L2缓存命中率。

通过比较朴素矩阵乘法和优化后的分块矩阵乘法，可以观察到后者具有显著更高的内存吞吐量，这是性能提升的直接体现。

高级概念：原子操作与流 ⚛️

最后，我们介绍两个高级概念：原子操作和CUDA流。

原子操作确保对同一内存地址的读-修改-写操作作为一个不可分割的整体执行，防止多线程同时访问导致的数据竞争（Race Condition）。例如，多个线程同时递增一个计数器时，使用原子操作能保证结果的正确性。

// 非原子操作，结果可能小于实际线程数
__global__ void nonAtomicIncrement(int* counter) {
    int old_val = *counter;
    *counter = old_val + 1;
}
// 原子操作，结果准确
__global__ void atomicIncrement(int* counter) {
    atomicAdd(counter, 1); // 原子加
}

原子操作会序列化对内存的访问，可能降低性能，但保证了正确性。

CUDA流用于实现操作间的并发，隐藏数据传输延迟。默认情况下，CUDA操作在默认流（0流）中顺序执行。创建多个流可以实现：

并发数据传输与内核执行：当一个流在执行内核时，另一个流可以同时进行数据拷贝。
并发内核执行：多个内核在不同流中可能同时执行（如果硬件支持）。

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 在流1中异步拷贝数据A并执行内核A
cudaMemcpyAsync(d_a, h_a, size, cudaMemcpyHostToDevice, stream1);
kernelA<<<..., stream1>>>(d_a, ...);
// 在流2中异步拷贝数据B并执行内核B（可能与流1的操作并发）
cudaMemcpyAsync(d_b, h_b, size, cudaMemcpyHostToDevice, stream2);
kernelB<<<..., stream2>>>(d_b, ...);
// 等待所有流完成
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);

结合事件（Event） 和回调（Callback），可以更精细地控制流间的依赖关系和执行时机，这对于构建高效、复杂的数据处理流水线至关重要。

总结

本节课中我们一起学习了CUDA编程的核心实践。我们从查询和理解GPU硬件规格出发，深入探讨了CUDA的网格-块-线程层次结构。我们动手编写了第一个向量加法和矩阵乘法内核，并学习了如何使用Nsight Compute工具来剖析内核性能。最后，我们介绍了用于保证数据安全访问的原子操作和用于提升并发性能的CUDA流。掌握这些基础知识，是进行高性能GPU编程和后续更深入优化的关键。

6：CUDA API 🚀

在本章中，我们将学习CUDA API的核心组成部分，特别是cuBLAS和cuDNN库。这些库提供了高度优化的函数，是GPU加速计算，尤其是深度学习领域的关键。我们将了解它们的基本概念、使用方法，并通过代码示例进行对比。

概述 📋

上一章我们深入探讨了CUDA的内存模型和内核执行。本节中，我们将转向更高级别的编程接口——CUDA API。这些API封装了复杂的底层操作，使我们能够轻松调用经过极致优化的GPU计算例程，从而显著提升应用程序性能。

我们将重点介绍两个核心库：

cuBLAS：用于执行基础线性代数运算，如矩阵乘法。
cuDNN：专为深度神经网络设计的库，提供卷积、池化、激活函数等操作。

理解这些API是构建高效GPU应用，特别是深度学习框架（如PyTorch）底层实现的关键。

访问官方文档 🔍

学习CUDA API的最佳起点是官方文档。你可以访问 docs.nvidia.com/cuda。这个网站提供了丰富的资源，包括：

安装指南：适用于Windows和Linux系统。
编程指南与最佳实践：涵盖从Maxwell到Hopper的各种架构。
PTX参考：CUDA编译后的汇编指令。
API参考：包括运行时API、驱动API和数学API。
工具：如用于调试CUDA程序的cuda-gdb，以及我们之前用过的nsight-compute。

对于我们本章的内容，主要关注的是 cuBLAS 和 cuDNN 的API参考。

理解“不透明结构类型” ⚫

cuBLAS和cuDNN中的函数并非我们手动编写内核。它们更像是“黑盒”函数，你调用它们，它们则在硬件上执行预编译好的、高度优化的代码。这些函数使用所谓的 不透明结构类型。

这意味着你看不到其内部实现代码（因为它们已被编译成人类难以阅读的二进制格式），只能通过文档中描述的结构和函数接口来调用它们。通常，这些由NVIDIA提供的API函数在大多数情况下都是最快的选择。

错误检查宏 🛡️

在调用CUDA API时，进行错误检查至关重要。我们通常会定义一些宏来包装函数调用。例如，对于cuBLAS函数：

#define CHECK_CUBLAS(err) do { \
    cublasStatus_t err_ = (err); \
    if (err_ != CUBLAS_STATUS_SUCCESS) { \
        printf("cuBLAS error at %s:%d code=%d\n", __FILE__, __LINE__, err_); \
        exit(EXIT_FAILURE); \
    } \
} while (0)

当调用一个cuBLAS函数（如 cublasSgemm）后，使用 CHECK_CUBLAS(...) 来检查是否返回错误。如果出错，它会打印错误信息和行号。对于cuDNN函数，我们也需要类似的检查宏。这能确保程序在出现意外时能清晰地报告问题。

cuBLAS：CUDA基础线性代数子程序 🧮

cuBLAS是 CUDA Basic Linear Algebra Subprograms 的缩写。顾名思义，它用于线性代数运算，其中最核心的就是矩阵乘法。例如，Sgemm 代表 单精度通用矩阵乘法。

在深度学习（如Transformer或MLP）中，矩阵乘法是关键算法。为了获得最快的推理速度，我们需要消除瓶颈。使用cuBLAS中的子程序通常能获得接近硬件极限的性能。

cuBLAS的不同版本

cuBLAS有几个变体，针对不同场景进行了优化：

以下是主要版本及其特点：

标准 cuBLAS：最易用、最通用的起点，支持基本的FP32和FP16矩阵乘法。
cuBLASLt：cuBLAS的轻量级扩展，提供更灵活的API，主要针对特定工作负载提升性能。它尤其擅长处理大矩阵和低精度计算（如FP16, FP8），在这些情况下可能比标准cuBLAS更快。
cuBLASXt：此版本支持在多个GPU和CPU之间互联以解决问题。它适用于那些因矩阵过大而无法完全放入单个GPU显存的大规模计算。但需要注意，由于CPU和GPU之间的内存带宽限制，跨设备计算可能会带来显著的性能开销。
cuBLASDx：此版本在主机端运行，其文档和优化可能不如其他版本完善。在Transformer等需要融合操作（如矩阵乘后接激活函数）的场景中，我们可能更倾向于使用在设备端执行完整内核的库。
CUTLASS：这是一个模板库，不属于官方cuBLAS，但值得了解。它允许开发者组合和定制高度优化的线性代数内核，实现操作融合（例如，将矩阵乘法、偏置加和激活函数融合为一个内核）。著名的Flash Attention论文实现就是手工编写融合内核的典范，能带来5-10倍的性能提升。CUTLASS为这类优化提供了工具。

代码示例：cuBLAS Sgemm 和 Hgemm

让我们通过一个具体例子来理解如何使用cuBLAS。以下代码演示了单精度(Sgemm)和半精度(Hgemm)矩阵乘法。

首先，包含必要的头文件和定义宏：

#include <cublas_v2.h>
#include <cuda_fp16.h>
#define M 3
#define K 4
#define N 2
// ... 错误检查宏 CHECK_CUBLAS ...

初始化矩阵（在CPU上）：

float h_A[M * K] = {1,2,3,4,5,6,7,8,9,10,11,12};
float h_B[K * N] = {1,2,3,4,5,6,7,8};
float h_C_cpu[M * N] = {0}; // CPU结果
float h_C_cublas[M * N] = {0}; // cuBLAS单精度结果
__half h_C_cublas_half[M * N]; // cuBLAS半精度结果

创建cuBLAS句柄（用于管理上下文）：

cublasHandle_t handle;
CHECK_CUBLAS(cublasCreate(&handle));

执行单精度矩阵乘法。关键点在于处理列主序：cuBLAS默认使用列主序存储矩阵，而我们通常使用行主序。一种技巧是通过交换参数来“欺骗”cuBLAS，使其按我们的意图计算：

float alpha = 1.0f, beta = 0.0f;
CHECK_CUBLAS(cublasSgemm(handle,
    CUBLAS_OP_N, CUBLAS_OP_N, // 操作：不转置
    N, M, K, // 注意维度顺序：N, M, K 而非 M, K, N
    &alpha,
    d_B, N, // 设备B矩阵，领先维度N
    d_A, K, // 设备A矩阵，领先维度K
    &beta,
    d_C, N // 设备C矩阵，领先维度N
));

执行半精度矩阵乘法，需要先将数据转换为half类型：

__half alpha_half = __float2half(1.0f);
__half beta_half = __float2half(0.0f);
CHECK_CUBLAS(cublasHgemm(handle,
    CUBLAS_OP_N, CUBLAS_OP_N,
    N, M, K,
    &alpha_half,
    d_B_half, N,
    d_A_half, K,
    &beta_half,
    d_C_half, N
));

最后，将结果拷贝回主机并打印验证。运行代码后，可以看到CPU计算结果、cuBLAS单精度结果和转换回浮点数的半精度结果都是一致的。

cuBLASLt 注意事项

使用cuBLASLt时，有一个重要限制：矩阵的维度（M, N, K）以及领先维度必须是4的倍数。例如，3x4的矩阵将无法成功运行，而4x4或12x16的矩阵则可以。因此，在处理大矩阵时，将维度设置为4096这类值是很好的实践。

性能对比实验

我们编写了一个基准测试脚本，比较不同方法处理大矩阵（4096x1024 * 1024x4096）的性能：

Naive内核：28毫秒
cuBLAS FP32：2.5毫秒
cuBLASLt FP32：2.8毫秒
cuBLASLt FP16：0.63毫秒

可以看到，cuBLASLt FP16的速度极快，相比Naive内核有巨大提升。而cuBLASXt由于涉及CPU-GPU通信，在相同规模计算下耗时约3.5秒，远慢于纯GPU计算，这凸显了注意数据位置的重要性。

cuDNN：CUDA深度神经网络库 🧠

如果说cuBLAS负责基础的矩阵运算，那么cuDNN则专注于深度学习中的其他核心操作。当你执行 pip install torch 时，cuDNN就是其底层依赖之一。

cuDNN覆盖了除矩阵乘法外的大部分常用深度学习操作：

卷积（Convolutions）
池化（Pooling）
激活函数（如Softmax, ReLU, Tanh）
丢弃层（Dropout）
批归一化（Batch Normalization）
张量变换（如重塑、连接）
层归一化（Layer Norm）

cuDNN的架构：图API与融合引擎

cuDNN的强大之处在于其图API和融合引擎的概念。在深度网络中，我们经常按顺序执行多个操作，例如“卷积 -> 加偏置 -> ReLU激活 -> 最大池化”。传统上，每个操作都需要单独的函数调用和内存读写。

cuDNN允许我们将这些操作融合成一个计算图。在这个图中，节点代表操作（如卷积、偏置加法），边代表流动的数据（张量）。然后，cuDNN可以对这个图进行整体优化，并选择或编译一个高效的“融合引擎”来执行，从而减少内核启动开销和内存访问，极大提升性能。

cuDNN提供几种类型的引擎：

预编译单操作引擎：针对单一操作（如特定卷积）高度优化，速度快但不灵活。
通用运行时融合引擎：能在运行时动态融合多个操作，灵活性高，但性能可能不如特化版本。
特化运行时融合引擎：针对特定操作模式（如“卷积+激活函数”）优化，兼顾灵活性和性能。
预编译序列引擎：针对固定操作序列预编译，能获得与单操作引擎相近的高性能。

代码示例：cuDNN Tanh激活函数

让我们看一个使用cuDNN执行Tanh激活函数的例子。虽然这是一个简单操作，但能展示基本流程。

首先，创建cuDNN句柄和描述符：

cudnnHandle_t cudnn_handle;
CHECK_CUDNN(cudnnCreate(&cudnn_handle));

cudnnTensorDescriptor_t tensor_desc;
CHECK_CUDNN(cudnnCreateTensorDescriptor(&tensor_desc));
// 设置张量描述符：格式为NCHW，数据类型为单精度浮点数
CHECK_CUDNN(cudnnSetTensor4dDescriptor(tensor_desc,
                                        CUDNN_TENSOR_NCHW,
                                        CUDNN_DATA_FLOAT,
                                        batch_size, channels, height, width));

cudnnActivationDescriptor_t activation_desc;
CHECK_CUDNN(cudnnCreateActivationDescriptor(&activation_desc));
// 设置激活描述符：模式为Tanh
CHECK_CUDNN(cudnnSetActivationDescriptor(activation_desc,
                                          CUDNN_ACTIVATION_TANH,
                                          CUDNN_NOT_PROPAGATE_NAN,
                                          0.0));

然后，调用前向传播函数：

float alpha = 1.0f, beta = 0.0f;
CHECK_CUDNN(cudnnActivationForward(cudnn_handle,
                                    activation_desc,
                                    &alpha,
                                    tensor_desc, d_input,
                                    &beta,
                                    tensor_desc, d_output));

在这个例子中，我们创建了一个很大的张量（约1.6GB）并进行Tanh计算。有趣的是，基准测试发现，一个手写的简单Naive CUDA内核（仅逐元素计算Tanh）比调用cuDNN的 cudnnActivationForward 略快一点（约1.3%）。

这可能是因为：

cuDNN函数是不透明的黑盒，内部可能包含我们未知的额外开销。
cuDNN API支持alpha和beta参数，即使我们设为1和0，也可能引入微小开销。
不过，这点性能差异在实际生产环境中几乎可以忽略。但对于更复杂的操作（如卷积），cuDNN的优势将非常明显。

代码示例：cuDNN卷积

卷积是cuDNN的强项。以下示例展示了如何设置并执行一个2D卷积，并与PyTorch的结果进行对比。

设置卷积描述符和过滤器描述符：

cudnnConvolutionDescriptor_t conv_desc;
CHECK_CUDNN(cudnnCreateConvolutionDescriptor(&conv_desc));
// 设置2D卷积描述符：填充、步长、膨胀等
CHECK_CUDNN(cudnnSetConvolution2dDescriptor(conv_desc,
    pad_h, pad_w, stride_h, stride_w, dilation_h, dilation_w,
    CUDNN_CROSS_CORRELATION, CUDNN_DATA_FLOAT));

cudnnFilterDescriptor_t filter_desc;
CHECK_CUDNN(cudnnCreateFilterDescriptor(&filter_desc));
// 设置过滤器描述符：格式为输出通道x输入通道x高度x宽度
CHECK_CUDNN(cudnnSetFilter4dDescriptor(filter_desc,
                                        CUDNN_DATA_FLOAT,
                                        CUDNN_TENSOR_NCHW,
                                        out_channels, in_channels, kernel_h, kernel_w));

为卷积操作选择算法。cuDNN提供了多种算法（如IMPLICIT_GEMM, FFT_TILING等），我们可以让cuDNN自动寻找最佳算法，也可以手动遍历选择：

cudnnConvolutionFwdAlgo_t algo;
CHECK_CUDNN(cudnnGetConvolutionForwardAlgorithm(cudnn_handle,
                                                 tensor_desc_input,
                                                 filter_desc,
                                                 conv_desc,
                                                 tensor_desc_output,
                                                 CUDNN_CONVOLUTION_FWD_PREFER_FASTEST,
                                                 0, // 无内存限制
                                                 &algo));

执行卷积前向传播：

CHECK_CUDNN(cudnnConvolutionForward(cudnn_handle,
                                     &alpha,
                                     tensor_desc_input, d_input,
                                     filter_desc, d_kernel,
                                     conv_desc,
                                     algo,
                                     d_workspace, workspace_size,
                                     &beta,
                                     tensor_desc_output, d_output));

在这个例子中，我们使用一个较小的输入（4x4图像）和卷积核（3x3）进行测试，以确保逻辑正确并与PyTorch结果匹配。对于小问题，cuDNN可能由于设置开销而比Naive内核慢。

但是，当我们进行大规模卷积的性能对比时（例如，批大小4，通道32/64，图像224x224，卷积核11x11），结果截然不同：

Naive卷积内核：82毫秒
cuDNN卷积：14.8毫秒

cuDNN带来了超过5.5倍的性能提升！这充分展示了为何PyTorch等框架要依赖cuDNN——它能够以惊人的速度执行极其复杂的计算。

大规模计算：cuBLASMg， NCCL 与 MIG 🖥️

当你的工作负载扩展到多GPU甚至GPU集群时，还需要了解其他一些工具。

以下是用于大规模计算的关键技术：

MIG：多实例GPU技术。它允许将一块强大的数据中心GPU（如A100）分割成多个独立的、更小的GPU实例。每个实例可以安全地运行不同的工作负载，提高大卡的利用率和隔离性。
NCCL：NVIDIA集体通信库。它是用于多GPU和多节点间高性能通信的库。它不进行计算，而是负责在集群中高效地分发、收集和同步数据，操作包括All-Reduce、Broadcast、Gather、Scatter等。在PyTorch中，DistributedDataParallel就建立在NCCL之上。
cuBLASMg：专为分布式密集线性代数设计的库。如果你需要在多个GPU上进行张量或矩阵运算，cuBLASMg可以处理这种分布式的计算任务。

对于大多数使用PyTorch等框架的用户，DistributedDataParallel已经封装了底层的分布式通信和计算。但如果你需要深入构建或优化数据中心级的基础设施，理解cuBLASMg和NCCL将非常有用。

总结 🎯

本节课中，我们一起深入探讨了CUDA API的核心部分。

我们首先了解了如何利用官方文档来学习和查找API。然后，重点学习了两个最重要的库：

cuBLAS：用于加速线性代数运算，特别是矩阵乘法。我们比较了其不同版本（标准版、Lt、Xt）的适用场景，并通过代码示例掌握了如何使用Sgemm和Hgemm函数，其中特别需要注意列主序存储格式的处理技巧。
cuDNN：专为深度学习设计的库。我们理解了其图API和融合引擎如何通过将多个操作合并执行来大幅提升性能。通过Tanh激活函数和卷积的例子，我们学习了cuDNN的基本调用方法，并亲眼见证了其在复杂操作（如卷积）上带来的巨大性能优势。

最后，我们简要介绍了用于大规模集群计算的NCCL和cuBLASMg，以及用于GPU虚拟化的MIG技术，为你未来的扩展学习指明了方向。

掌握这些CUDA API，意味着你能够更直接地驾驭GPU的强大算力，并为理解和优化高级深度学习框架打下坚实基础。

7：更快的矩阵乘法 🚀

概述

在本节课中，我们将深入学习如何优化CUDA中的矩阵乘法（MatMul）核心。矩阵乘法是深度学习等高性能计算领域的基石算法。我们将从最基础的实现开始，逐步引入一系列优化技术，最终达到接近NVIDIA cuBLAS库的性能水平。本教程将遵循Simon Boehm的“SGEMM CUDA”仓库和博客文章的思路，通过多个优化步骤，让初学者也能理解如何编写高效的CUDA内核。

7.1 基础实现与性能基准

如果你已经坚持学习到这里，值得为自己鼓掌。到目前为止，我们已经涵盖了大量内容。现在，我们将进入课程中技术性更强的部分：矩阵乘法及其优化。

这将是技术性最强的部分之一，因为我们将关注底层优化，探究如何在硬件上真正加速计算。这不再仅仅是理解其工作原理的通用概念，我们将运用已有的知识以及即将分享的新知识，使基础的矩阵乘法算法变得非常、非常快。这个算法在深度学习中无处不在，因此我认为教授如何优化内核的最佳方式就是以此为例。

幸运的是，我们有一个由Simon Boehm创建的仓库。他目前在Anthropic担任性能或内核工程师，经验丰富。他创建了一个名为“SGEMM CUDA”的酷炫仓库以及配套的博客文章。我将跟随他的思路进行讲解。本教程的目标是逐步优化，最终达到接近甚至超越cuBLAS性能的水平（具体取决于你的硬件）。

如果你已经看过那篇文章但觉得太难，别担心，我将详细讲解，我们会深入到非常底层的细节。完成这部分学习后，你将清楚地理解如何优化CUDA内核。

我们不会在主要的CUDA课程仓库中进行这部分操作。我会在README文件中提供链接，方便你跟随学习。现在，我将克隆Simon的仓库到主目录。

首先，删除旧版本（如果有的话），然后克隆新仓库。

git clone https://github.com/simonboehm/sgemm-cuda.git

克隆完成后，在VS Code中打开这个新仓库。

首先，查看README文件中的构建说明。通常需要创建一个build目录，然后使用cmake和make进行构建。

mkdir build
cd build
cmake ..
make

构建过程需要一点时间。构建完成后，我们将运行一系列基准测试，依次展示每种优化后的性能，从最基础的实现开始，逐步到最复杂的优化，并与cuBLAS的性能进行对比。

7.2 朴素（Naive）实现

让我们从朴素实现开始。实际上，我们在之前的课程中已经实现过朴素的矩阵乘法内核。

回顾一下我们之前编写的内核：它接收矩阵A、B、C以及参数alpha和beta。它执行的操作类似于cuBLAS中的SGEMM：计算alpha * (A * B) + beta * C。我们主要关注计算A * B的矩阵乘法核心部分。

在朴素实现中，我们定义了矩阵的维度：A的形状是M x K，B的形状是K x N，输出C的形状是M x N。

在内核中，我们通过threadIdx和blockIdx计算出当前线程负责的C矩阵中的元素坐标(row, col)。然后，我们使用一个循环遍历K维度，累加A的一行和B的一列的点积结果。

核心计算循环如下：

float sum = 0.0f;
for (int i = 0; i < K; ++i) {
    sum += A[row * K + i] * B[i * N + col];
}
C[row * N + col] = alpha * sum + beta * C[row * N + col];

这就是朴素的矩阵乘法。现在，让我们运行这个内核的基准测试。

进入构建目录，运行对应的可执行文件（例如sgemm_naive）。

./sgemm_naive

输出会显示不同矩阵大小（如128, 256, 512, 1024, 2048, 4096）下的性能，单位是GFLOPS（每秒十亿次浮点运算）。

例如，在4096x4096的矩阵上，朴素内核可能达到约166 GFLOPS。这听起来很高，但实际上相对于硬件潜力来说非常低。随着优化，这个数字会大幅提升。

请注意，在4096大小上运行50次迭代，朴素内核可能耗时约0.83秒。我们将看到优化如何显著减少这个时间。

7.3 全局内存合并访问优化

上一节我们介绍了朴素的实现，其性能有巨大的提升空间。本节中，我们来看看第一个关键优化：全局内存合并访问。

首先，理解内存布局至关重要。在内存中，一个M x N的矩阵是按行优先顺序连续存储的，即第一行之后紧跟着第二行，以此类推。

在朴素内核中，当线程计算点积时，对矩阵A的访问是连续的（沿着行），但对矩阵B的访问是不连续的（沿着列）。这导致了非合并的内存访问，严重降低了带宽利用率。

CUDA架构中，一个Warp（32个线程）内的内存访问如果满足特定模式（例如访问连续的内存地址），硬件可以将这些访问合并为一个或少数几个内存事务，从而极大提高效率。

优化思路：改变线程到数据映射的索引方式，确保每个Warp内的线程访问连续的内存地址。

在优化后的内核中，我们调整了row和col的计算方式：

int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
// 但为了合并访问，我们可能交换或修改索引计算
// 一种常见技巧是：
int row = (blockIdx.y * blockDim.y + threadIdx.y);
int col = (blockIdx.x * blockDim.x) * ELEMENTS_PER_THREAD + threadIdx.x; // 假设每个线程处理多个元素

更具体的实现可能涉及将线程ID的X维度用于列索引，并确保相邻线程访问相邻的列元素。

Simon的代码中，通过以下计算实现了合并访问：

int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + (threadIdx.x / ELEMENTS_PER_THREAD); // 示例

其核心是让threadIdx.x对应连续的内存地址。这样，当Warp中的线程threadIdx.x = 0, 1, 2, ...执行时，它们访问的全局内存地址是连续的。

性能提升：运行这个优化后的内核（例如sgemm_coalesced），在4096矩阵上，性能可能从166 GFLOPS跃升至约1183 GFLOPS，提升了约7倍。每次运行的耗时也从0.83秒降至约0.12秒。

这个优化清晰地展示了内存访问模式对CUDA内核性能的巨大影响。

7.4 共享内存缓存分块（Tiling）

上一节我们通过合并内存访问获得了显著的性能提升。本节中，我们引入一个更强大的概念：使用共享内存进行缓存分块。

目前，我们的内核直接从全局内存（VRAM）读取数据，其带宽虽然高（例如约700 GB/s），但延迟也高。GPU上每个流多处理器（SM）都有自己的一小块共享内存（SRAM），其带宽极高（例如超过1 TB/s），且延迟极低。

优化思路：与其让每个线程在每次计算时都从全局内存读取数据，不如让一个线程块（Block）协作，将一小块（Tile）数据从全局内存加载到共享内存中。然后，线程块内的所有线程可以高速、重复地访问共享内存中的这块数据进行计算。处理完当前数据块后，再加载下一块。

这个过程称为分块（Tiling）。我们将大的矩阵乘法分解为许多小的矩阵乘法。每个线程块负责计算输出矩阵C中的一个子块。为了计算这个子块，它需要从A加载若干行块，从B加载若干列块。

内核结构变化：

在共享内存中声明两个数组：__shared__ float As[TILE_SIZE][TILE_SIZE]; 和 __shared__ float Bs[TILE_SIZE][TILE_SIZE];。
计算线程块对应的输出子块在C中的起始位置。
在一个外层循环中，迭代K维度（A的列/B的行）。每次迭代：
a. 线程块协作，将A和B的相应数据块从全局内存加载到共享内存As和Bs中。
b. 调用__syncthreads()确保所有线程都完成加载。
c. 每个线程使用共享内存中的As和Bs计算其负责的部分点积，并累加到本地寄存器变量中。
d. 调用__syncthreads()确保所有线程完成计算，然后再加载下一个数据块。
循环结束后，将寄存器中累加的结果写回全局内存的C矩阵。

索引计算：这是最复杂的部分。需要仔细计算每个线程应该加载As和Bs中的哪个元素，以及如何将共享内存中的索引映射到全局内存。

以下是一个简化的加载步骤示例：

// 假设 TILE_SIZE = BLOCK_SIZE
int tx = threadIdx.x;
int ty = threadIdx.y;

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/elliot-cuda-prog/img/0f9ea5e8bacc4cc914094ea7b918190a_17.png)

// 加载 A 的 tile
int aRow = blockIdx.y * TILE_SIZE + ty;
int aCol = k * TILE_SIZE + tx; // k 是外层循环索引
if (aRow < M && aCol < K) {
    As[ty][tx] = A[aRow * K + aCol];
} else {
    As[ty][tx] = 0.0f;
}

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/elliot-cuda-prog/img/0f9ea5e8bacc4cc914094ea7b918190a_19.png)

// 加载 B 的 tile
int bRow = k * TILE_SIZE + ty;
int bCol = blockIdx.x * TILE_SIZE + tx;
if (bRow < K && bCol < N) {
    Bs[ty][tx] = B[bRow * N + bCol];
} else {
    Bs[ty][tx] = 0.0f;
}
__syncthreads();

性能提升：运行共享内存分块内核（例如sgemm_tiled）。在4096矩阵上，性能可能从1183 GFLOPS进一步提升至约1600 GFLOPS。这证明了利用更快的内存层次结构的好处。

注意：这个实现是“分块”的一种形式，但每个线程仍然只计算输出C中的一个元素。接下来，我们将让每个线程计算多个元素，以进一步提高效率。

7.5 一维块分块（1D Block Tiling）

上一节我们利用共享内存减少了全局内存访问。本节中，我们通过一维块分块来增加每个线程的计算量，从而提升算术强度。

在之前的共享内存内核中，每个线程只负责输出矩阵C中的一个元素。这意味着，如果输出矩阵很大，就需要启动大量的线程。虽然GPU线程很轻量，但启动和调度它们仍有开销。更重要的是，每个线程只执行很少的计算（一个点积），相对于内存操作（加载数据到共享内存）的比例较低，这被称为算术强度低。

优化思路：让一个线程计算输出子块中的多个元素（例如一列）。这样，每个线程需要加载的数据量（特别是从共享内存）与其执行的计算量之比增加，能更好地隐藏内存延迟，提高硬件利用率。

内核变化：

线程块结构：线程块的大小（blockDim）可能变小，但每个线程的任务变多。
输出子块：每个线程块现在计算C中一个更大的矩形区域（例如 BM x BN）。
每个线程的任务：每个线程计算该矩形区域中的一小列（大小为 TM）。TM 是一个调优参数（例如8）。
寄存器使用：每个线程使用一组寄存器（例如 float reg_M[TM]）来累加这一列中每个元素的部分和。
计算循环：在外层循环遍历K维度的数据块时，内层循环遍历 TM，计算每个元素的部分点积。

索引计算变得更加复杂：需要计算：

thread_row：线程在块内负责的输出区域的起始行。
thread_col：线程在块内负责的输出区域的起始列（在一维分块中，通常线程负责一列，所以thread_col是固定的）。
在共享内存中加载数据时，需要根据thread_row和循环索引k来计算偏移。

核心计算部分伪代码示意：

float reg_C[TM] = {0}; // 每个线程计算TM个结果
for (int k_offset = 0; k_offset < K; k_offset += TILE_SIZE_K) {
    // 1. 协作加载 As 和 Bs 的 tile 到共享内存 (As[TILE_SIZE_K][BM], Bs[TILE_SIZE_K][BN])
    // 2. __syncthreads()
    
    // 3. 每个线程从共享内存加载数据到寄存器并计算
    for (int i = 0; i < TILE_SIZE_K; ++i) {
        float reg_A = As[i][thread_local_row]; // 加载A的一个元素
        for (int j = 0; j < TM; ++j) { // TM是线程计算的行数
            float reg_B = Bs[i][thread_local_col + j]; // 加载B的一列中的TM个元素（需要仔细索引）
            reg_C[j] += reg_A * reg_B;
        }
    }
    // 4. __syncthreads()
}
// 5. 将 reg_C[TM] 写回全局内存

性能提升：运行一维块分块内核（例如sgemm_1d_block）。在4096矩阵上，性能可能从1600 GFLOPS提升至约4800 GFLOPS。这是一个巨大的飞跃，主要归功于每个线程计算量的增加，更好地利用了核心的计算单元。

7.6 二维块分块（2D Block Tiling）

上一节中，每个线程计算一列（1D）元素。本节中，我们进一步扩展，让每个线程计算一个小的二维块（2D Tile），这被称为二维块分块。

一维分块中，线程计算一列，对共享内存中B矩阵的访问是连续的（合并的），但对A矩阵的访问可能不是最理想的。二维分块旨在更平衡地利用内存带宽和计算资源。

优化思路：每个线程计算输出子块中一个更小的二维矩形区域（例如 TM x TN，如 8x8）。这样，线程需要从共享内存中加载A的一小列（TM个元素）和B的一小行（TN个元素），然后计算它们的并积（outer product），更新线程本地存储的 TM x TN 个累加器。

优势：

计算与内存访问比更高：加载 TM + TN 个数据，进行 TM * TN 次乘加运算。
更好的指令级并行：循环展开和软件流水线更容易优化。

内核变化：

线程块结构：线程块大小根据 (BM/TM) * (BN/TN) 计算，确保每个输出子块（BM x BN）被合理分配。
寄存器使用：每个线程声明一个寄存器数组 float thread_results[TM][TN] 或展开成一维数组 float reg_C[TM*TN]。

核心计算循环：在外层循环加载数据块到共享内存后，内层是一个双重循环（或展开）：

for (int k_inner = 0; k_inner < TILE_SIZE_K; ++k_inner) {
    // 从共享内存 As 加载 TM 个元素到寄存器 reg_A[TM]
    // 从共享内存 Bs 加载 TN 个元素到寄存器 reg_B[TN]
    for (int i = 0; i < TM; ++i) {
        for (int j = 0; j < TN; ++j) {
            thread_results[i][j] += reg_A[i] * reg_B[j];
        }
    }
}

索引计算：需要仔细计算每个线程在加载As和Bs、以及写入最终结果时的位置。

性能提升：运行二维块分块内核（例如sgemm_2d_block）。在4096矩阵上，性能可能从4800 GFLOPS进一步提升至约9100 GFLOPS。这使我们离cuBLAS的性能（约11500 GFLOPS）更近了一步。

7.7 向量化内存访问

上一节我们通过二维分块大幅提升了性能。本节中，我们进行最后一项关键优化：向量化内存访问，以最大化内存带宽利用率。

即使在使用共享内存之后，从全局内存加载数据到共享内存，以及从共享内存加载数据到寄存器，仍然是性能瓶颈。CUDA支持向量数据类型，如float4（包含4个float）。使用这些类型可以进行向量化加载/存储。

优化思路：将多个标量内存访问合并为一个向量内存访问指令。例如，一次加载一个float4（128位）而不是四次单独的float（32位）加载。这减少了指令数量，提高了内存事务的效率，并有助于实现更宽的内存合并。

实施步骤：

数据类型：在全局内存和共享内存的加载/存储操作中使用float4。
指针转换：使用reinterpret_cast<float4*>将普通的float*指针转换为float4*指针。这向编译器表明数据是128位对齐的，允许生成向量加载指令（如LDG.E.128）。
索引调整：因为一次操作处理4个元素，所以循环步长和索引计算需要相应调整。例如，原来循环步长为1，现在可能步长为4。
转置共享内存布局：为了确保从共享内存到寄存器的加载也是合并的，有时需要改变共享内存中数据的布局（例如，将A tile存储为转置形式）。这样，当线程读取一列数据（在转置后是连续存储的）时，访问是连续的。

代码示例（全局内存加载到共享内存）：

// 假设每个线程负责加载一个 float4（4个float）
int load_idx = ... // 计算全局内存索引
float4* A_vec = reinterpret_cast<float4*>(&A[load_idx]);
float4 loaded_val = *A_vec; // 一次向量加载

// 将 loaded_val 的四个分量 (x, y, z, w) 存储到共享内存的适当位置
__shared__ float As_tile[TILE_SIZE][TILE_SIZE];
int tile_row = ...;
int tile_col = ...;
As_tile[tile_row][tile_col] = loaded_val.x;
As_tile[tile_row][tile_col+1] = loaded_val.y; // 注意列索引+1
// ... 存储 z 和 w

性能提升：运行向量化内核（例如sgemm_vectorized）。在4096矩阵上，性能可能从9100 GFLOPS提升至约10800 GFLOPS。这已经达到了cuBLAS性能（~11500 GFLOPS）的94%左右，是一个巨大的成功！

验证：我们可以使用nvcc -ptx和cuobjdump工具来查看生成的PTX和SASS汇编代码，确认确实生成了LDG.E.128这样的向量加载指令，而不是大量的LDG.E.32指令。

7.8 总结与进阶方向

本节课中，我们一起学习了如何逐步优化CUDA矩阵乘法内核：

朴素实现：基础功能，性能低下（~166 GFLOPS）。
全局内存合并：通过调整索引实现合并访问，性能提升约7倍（~1183 GFLOPS）。
共享内存分块：利用高速SRAM缓存数据，减少全局内存访问（~1600 GFLOPS）。
一维块分块：增加每个线程计算量，提高算术强度（~4800 GFLOPS）。
二维块分块：让每个线程计算一个小矩阵块，进一步优化计算与内存访问比（~9100 GFLOPS）。
向量化访问：使用float4进行向量加载/存储，最大化内存带宽利用率（~10800 GFLOPS，接近cuBLAS）。

通过这些步骤，我们见证了性能从百GFLOPS级别提升到万GFLOPS级别，深刻理解了内存层次结构、访问模式、算术强度以及指令优化对GPU内核性能的决定性影响。

进阶方向：

自动调优：像实际库一样，对块大小（BM, BN, BK, TM, TN等）进行自动搜索，找到特定硬件上的最优配置。
张量核心（Tensor Cores）：现代GPU（Volta架构及以后）配备了专门用于矩阵乘法的张量核心，可提供极高的吞吐量（用于FP16, BF16, INT8等精度）。CUDA提供了wmma（Warp Matrix Multiply Accumulate）命名空间来编程使用张量核心。这可以将性能再提升一个数量级。
异步拷贝与屏障：利用CUDA 11+的异步内存拷贝和屏障特性，进一步重叠计算与内存传输。
多级分块：结合寄存器、共享内存、L2缓存进行更复杂的分块策略。

希望本教程为你提供了优化CUDA内核的坚实基础和清晰路线图。鼓励你继续探索，尝试实现更高级的优化技术！

8：Triton入门 🚀

在本节课中，我们将学习Triton，这是一个基于Python的高级编程语言和编译器，旨在简化GPU编程。它将我们从繁琐的CUDA底层细节中解放出来，让我们能够用更简洁的语法实现高性能计算，例如之前学习的矩阵乘法分块优化。我们将通过向量加法和Softmax两个核心例子，来理解Triton的基本工作原理和语法。

Triton简介与设计哲学

上一节我们完成了矩阵乘法的核心优化。本节中，我们来看看Triton如何将这些优化抽象化，让我们用更简单的Python语法实现。

Triton与CUDA不同。安装PyTorch时，你会看到它包含了Triton。PyTorch在底层使用Triton来加速计算。Triton和CUDA一样快，但设计哲学有根本区别。

访问Triton官网或GitHub，可以看到完整的文档，其中最重要的是“Triton语言”部分，它列出了所有可用的操作。

Triton的设计灵感来源于一篇关于分块神经网络计算的论文。其核心思想可以概括为一个对比：

CUDA：标量程序，块级线程。你编写的内核（kernel）运行在单个线程（标量）级别，但必须显式地处理线程间的通信与协作（例如共享内存）。
Triton：块级程序，标量线程。你编写的程序运行在线程块（block）级别，而所有线程级别的操作和优化都由Triton编译器自动处理，你无需关心线程间如何通信。

这意味着Triton编译器为我们处理了大量底层细节。那么，我们能否跳过CUDA直接学习Triton？答案是否定的。Triton是构建在CUDA之上的抽象层，它利用了CUDA提供的底层优化。要确保正确应用优化或构建自己的高级抽象，理解CUDA的底层机制仍然是必要的。Triton帮助我们减少样板代码，而非取代对硬件执行模型的理解。

简而言之，Triton实现了与CUDA、CuBLAS相近的性能，但代码更简洁。接下来，我们将通过实例学习Triton的基础。

向量加法示例 ➕

现在，我们通过一个向量加法的例子，来具体看看如何在Triton中编写内核。与CUDA版本对比，你会发现它简洁许多。

首先，我们导入必要的库并设置环境。

import torch
import triton
import triton.language as tl
import random

# 设置随机种子以确保结果可复现
torch.manual_seed(123)
random.seed(123)

# 定义向量大小（约3300万个元素）
SIZE = 2 ** 25
n_elements = SIZE

我们创建两个随机初始化的向量 x 和 y，并准备一个输出张量 output。

# 在GPU上创建输入向量
x = torch.randn(n_elements, device='cuda')
y = torch.randn(n_elements, device='cuda')
# 创建输出张量
output = torch.empty_like(x)

以下是使用PyTorch进行向量加法的基准测试，作为性能对比的参考。

# PyTorch基准测试
def add_pytorch(x, y):
    return x + y

现在，我们来看核心的Triton内核函数。@triton.jit 装饰器告诉Triton编译这个函数。

@triton.jit
def add_kernel(
    x_ptr,  # 输入x的指针
    y_ptr,  # 输入y的指针
    output_ptr,  # 输出指针
    n_elements,  # 向量总元素数
    BLOCK_SIZE: tl.constexpr,  # 每个块处理的元素数，编译时常量
):
    # 获取当前程序（块）在网格中的索引
    pid = tl.program_id(axis=0)
    # 计算当前块负责的数据起始索引
    block_start = pid * BLOCK_SIZE
    # 生成当前块内所有线程的偏移量 (0, 1, ..., BLOCK_SIZE-1)
    offsets = block_start + tl.arange(0, BLOCK_SIZE)
    # 创建掩码，防止访问超出数组边界的数据
    mask = offsets < n_elements
    # 从全局内存加载数据到快速内存（由Triton优化）
    x = tl.load(x_ptr + offsets, mask=mask)
    y = tl.load(y_ptr + offsets, mask=mask)
    # 执行逐元素加法（由Triton优化）
    output = x + y
    # 将结果存回全局内存
    tl.store(output_ptr + offsets, output, mask=mask)

为了启动这个内核，我们需要配置执行网格。以下是一个辅助函数。

def add_triton(x, y):
    output = torch.empty_like(x)
    n_elements = output.numel()
    # 定义网格大小：需要多少个块来覆盖所有元素
    grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']),)
    # 启动内核
    add_kernel[grid](x, y, output, n_elements, BLOCK_SIZE=1024)
    return output

最后，我们验证Triton的结果与PyTorch是否一致。

# 验证结果
output_triton = add_triton(x, y)
output_pytorch = add_pytorch(x, y)
print(f"最大差异: {torch.max(torch.abs(output_triton - output_pytorch))}")
# 输出: 最大差异: 0.0

在这个例子中，我们看到了Triton内核的关键部分：tl.program_id 获取块ID，tl.arange 生成偏移，tl.load/tl.store 处理内存，以及使用掩码进行边界检查。所有线程级别的协同和内存优化都由Triton在幕后完成。

Softmax函数实现 🔥

理解了向量加法后，我们来看一个更复杂的例子：Softmax函数。这能展示Triton处理归约操作的能力。

首先，我们回顾Softmax的数学公式和数值稳定性的重要性。

Softmax函数将一个向量转换为概率分布，公式为：
softmax(x_i) = exp(x_i) / sum(exp(x_j))，对所有的 j 求和。

直接计算可能导致数值溢出（例如 exp(1000)）。因此，通常采用稳定版本：
softmax(x_i) = exp(x_i - max(x)) / sum(exp(x_j - max(x)))。

我们先在纯Python/C中理解其步骤：

找到向量中的最大值 max_val。
对每个元素计算 exp(x_i - max_val)。
计算所有 exp(x_i - max_val) 的总和 sum_exp。
每个元素的结果为 exp(x_i - max_val) / sum_exp。

在深度学习中，我们通常处理一个批次（batch）的向量，对每一行独立进行Softmax。接下来，我们看Triton如何实现。

我们设置输入，一个形状为 (batch_size, n) 的张量。

import torch.nn.functional as F

@triton.jit
def softmax_kernel(
    output_ptr, input_ptr, input_row_stride, output_row_stride, n_cols,
    BLOCK_SIZE: tl.constexpr
):
    # 行ID：每个块处理一行
    row_idx = tl.program_id(0)
    # 计算当前行在内存中的起始位置
    row_start_ptr = input_ptr + row_idx * input_row_stride
    # 计算输出行的起始位置
    output_row_start_ptr = output_ptr + row_idx * output_row_stride
    # 生成列偏移量
    col_offsets = tl.arange(0, BLOCK_SIZE)
    # 计算输入指针，准备加载数据
    input_ptrs = row_start_ptr + col_offsets
    # 掩码：确保不越界
    mask = col_offsets < n_cols
    # 将整行数据加载到快速内存
    row = tl.load(input_ptrs, mask=mask, other=-float('inf'))
    # 1. 求行最大值（归约操作）
    row_minus_max = row - tl.max(row, axis=0)
    # 2. 计算指数
    numerator = tl.exp(row_minus_max)
    # 3. 计算分母（指数和）
    denominator = tl.sum(numerator, axis=0)
    # 4. 计算Softmax输出
    softmax_output = numerator / denominator
    # 将结果存回输出内存
    output_ptrs = output_row_start_ptr + col_offsets
    tl.store(output_ptrs, softmax_output, mask=mask)

以下是启动内核和验证的函数。

def softmax(x):
    n_rows, n_cols = x.shape
    # 计算大于等于n_cols的最小的2的幂，作为块大小
    BLOCK_SIZE = triton.next_power_of_2(n_cols)
    # 限制最大块大小，并调整BLOCK_SIZE为2的幂
    num_warps = 4
    if BLOCK_SIZE >= 2048:
        num_warps = 8
    if BLOCK_SIZE >= 4096:
        num_warps = 16
    BLOCK_SIZE = min(BLOCK_SIZE, 1024)
    # 网格配置：每个行一个块
    grid = (n_rows,)
    # 分配输出内存
    y = torch.empty_like(x)
    # 启动内核
    softmax_kernel[grid](
        y, x,
        x.stride(0), y.stride(0),
        n_cols,
        num_warps=num_warps,
        BLOCK_SIZE=BLOCK_SIZE,
    )
    return y

# 测试与验证
torch.manual_seed(0)
x = torch.randn(1823, 781, device='cuda')
y_triton = softmax(x)
y_torch = torch.softmax(x, axis=1)
print(f"最大差异: {torch.max(torch.abs(y_triton - y_torch))}")
# 输出应为一个极小的数，例如 5.960464477539063e-08

在这个内核中，我们看到了Triton如何优雅地处理归约操作（tl.max, tl.sum）。我们以“行”为块单位，每行数据由一个线程块处理，Triton编译器自动处理了块内线程的协作以完成归约。

总结 📚

本节课中我们一起学习了Triton的基础知识。我们从Triton的设计哲学讲起，理解了它作为“块级程序，标量线程”的抽象，如何将开发者从繁琐的CUDA线程管理细节中解放。随后，我们通过向量加法和Softmax两个实例，逐步剖析了Triton内核的编写方法：

使用 @triton.jit 装饰器定义内核。
利用 tl.program_id、tl.arange 管理数据索引。
通过 tl.load 和 tl.store 高效搬运数据，并用掩码处理边界。
依赖Triton内置操作（如 tl.max、tl.sum、tl.exp）简洁地实现复杂计算。

Triton让我们能够用更接近Python思维的方式编写高性能GPU代码，同时保留了接近CUDA的性能。掌握它，你就能在保持开发效率的前提下，充分挖掘GPU的并行计算潜力。

10：PyTorch扩展

概述

在本节课中，我们将学习如何为PyTorch创建自定义的CUDA扩展。我们将通过一个简单的多项式激活函数示例，展示如何编写CUDA内核，将其编译为Python模块，并与PyTorch集成，最终对比其与原生PyTorch实现的性能差异。

扩展的必要性

上一章我们完成了Tri部分。现在我们将更深入地探讨Python和PyTorch方面。为PyTorch添加自定义CUDA扩展，可以针对特定用例提升运算速度。

项目文件结构

以下是本示例项目包含的文件及其简要说明。

README.md：包含关于不同类型、名称的描述以及我们将要查看的直观示例。
setup.py：用于编译独立PyTorch扩展的脚本。
polynomial.cu：包含我们将用于执行多项式运算的独立函数的CUDA内核代码。
polynomial.cpp：负责编译并将扩展绑定到PyTorch的脚本，以便在Python中使用。
benchmark.py：Python脚本本身，用于与原生PyTorch实现进行性能对比。

我们实现的操作非常简单：x² + x + 1。

解析CUDA内核 (`polynomial.cu`)

模板与限制符

首先查看CUDA内核文件。顶部有一个新概念：模板。在之前的Mamal部分曾简要提及。

模板本质上允许我们编写通用代码。在调用内核时，我们指定 scalar_t 类型。这意味着PyTorch将确保输入x和输出y属于此类型。PyTorch会自动识别并处理float、double或FP16等类型，并相应地进行编译。这是PyTorch内置的自定义类型，是此处最简便的默认选择。

接下来是 restrict 关键字。简而言之，它意味着内存访问不会重叠。我们在这里有输入x和输出y。我们只对x进行操作，然后将结果存储在y中。我们没有在相同内存位置进行混合操作或存储。因此，我们可以使用restrict，这将允许编译器对二进制代码进行积极的优化。

内核实现

在内核函数内部，我们简单地遍历x维度，使用典型的内核索引。然后执行平方、相加、再加一的操作。内核本身非常简单，主要是一些需要注意的新颜色和关键字。

template <typename scalar_t>
__global__ void polynomial_activation_kernel(
    const scalar_t* __restrict__ x,
    scalar_t* __restrict__ y,
    int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        scalar_t val = x[idx];
        y[idx] = val * val + val + 1;
    }
}

解析C++绑定代码 (`polynomial.cpp`)

自动类型推断与内核配置

向下滚动，我们看到一些C++语法。这里使用了 auto 关键字，它可以自动推断变量的类型。它会识别出这是一个PyTorch张量类型。auto让代码看起来更简洁。

然后是线程配置，我们使用每个块1024个线程，这是典型的最大值。接着计算块的数量：(numel + threads - 1) / threads。这些内容之前已经讲过，应该非常简单。

函数封装与Python绑定

接下来是关于如何调用内核的一些额外设置，即我们向它传递哪些参数。基于这些参数，我们返回一个输出。

最后是Python绑定部分。不必过于担心细节，相信这个过程即可。

PYBIND11_MODULE：这是一个宏，定义了Python模块的入口点。
TORCH_EXTENSION_NAME：这是PyTorch定义的宏，用于扩展模块的名称，通常由setup.py文件定义。
m.def()：此方法向模块添加一个新函数。第一个参数"polynomial_activation"是Python中函数的名称。第二个参数是指向要调用的C++函数的指针。最后一个参数是函数的文档字符串。

这是完整的CUDA脚本，本质上已经是最简单的形式，可以作为模板使用。

解析Python集成与基准测试 (`benchmark.py`)

自定义Autograd函数

在Python脚本中，我们导入编译好的polynomial_cuda函数。我们定义了一个类，使用torch.autograd.Function。默认情况下，当我们创建一个autograd函数时，必须包含forward和backward方法，两者都应为静态方法。我们添加装饰器@staticmethod。

在forward方法中，我们调用编译好的polynomial_activation函数。在backward方法中，我们暂时不支持，因此直接抛出NotImplementedError。

模块定义与实现选择

在主要的模块定义中，我们进行初始化。在forward方法中，我们根据implementation参数决定使用PyTorch原生实现还是我们的CUDA扩展。如果未指定实现，则提示错误。

基准测试逻辑

在主函数中，我们设置随机种子，在指定设备上创建正态分布的随机张量。我们分别使用PyTorch实现和CUDA扩展实现，并将它们移动到设备上。

基准测试的过程是：记录开始时间，运行指定次数的函数，使用cuda.synchronize()确保所有操作完成，记录结束时间，然后计算平均耗时（毫秒）。

编译与运行

编译扩展

要编译此扩展，只需运行python setup.py install。这将使用Ninja进行构建，并在当前目录下创建一个build文件夹。

执行与性能对比

编译完成后，我们可以运行Python脚本。它将打印出输入张量、CUDA扩展的输出结果，以及两种实现的平均运行时间。

例如，PyTorch内置函数可能平均耗时约10.47毫秒，而CUDA扩展可能仅需约0.243毫秒。通过计算 10.47 / 0.243 ≈ 43.1，我们得到了约43倍的加速比，这对于大型张量来说是非常显著的性能提升。

如果尝试调用.backward()方法，CUDA扩展版本会抛出错误，因为我们尚未实现反向传播，此时会回退到前向传播。

总结

本节课我们一起学习了如何为PyTorch创建自定义CUDA扩展。我们从编写一个简单的多项式激活函数CUDA内核开始，然后使用C++和pybind11将其绑定到Python，最后在PyTorch模块中集成并进行了性能基准测试。这个过程展示了如何通过自定义CUDA代码来显著提升特定运算的性能。你可以自由地以此模板为基础，添加自己的自定义研究代码，使其易于自己、他人或组织使用。这是我能找到的最简单易懂的编写和解释示例。

接下来，我们将进入一个最终项目，该项目将非常令人兴奋，帮助你从头理解神经网络，以及如何为性能优化它们，例如添加数据加载器等多种优化，以进行真实世界的训练运行。

11：第10章 (MNIST多层感知机) 🚀

概述

在本节课中，我们将学习如何从零开始训练一个MNIST多层感知机（MLP）。我们将从Python和PyTorch开始，逐步深入到NumPy，然后将其移植到C语言，最后利用CUDA进行加速。通过这个过程，我们将深入理解神经网络的前向传播、反向传播以及梯度下降等核心概念。

从PyTorch开始 🐍

上一节我们介绍了课程的整体结构，本节中我们来看看如何使用PyTorch快速搭建并训练一个MNIST MLP。

我们首先导入必要的库，并设置超参数。

import time
import numpy as np
import torch
import torch.nn as nn
import torch.optim as optim
from torch.utils.data import DataLoader
from torchvision import datasets, transforms

# 超参数
learning_rate = 1e-3
batch_size = 4
num_epochs = 3
train_size = 10000

接着，我们设置数据目录并初始化数据集和数据加载器。

# 设置数据目录
data_dir = './python/data'

# 启用TF32以使用张量核心加速
torch.backends.cuda.matmul.allow_tf32 = True

# 数据预处理
transform = transforms.Compose([
    transforms.ToTensor(),
    transforms.Normalize((0.1307,), (0.3081,))
])

# 初始化数据集
train_dataset = datasets.MNIST(root=data_dir, train=True, download=True, transform=transform)
test_dataset = datasets.MNIST(root=data_dir, train=False, download=True, transform=transform)

# 初始化数据加载器
train_loader = DataLoader(train_dataset, batch_size=batch_size, shuffle=True)
test_loader = DataLoader(test_dataset, batch_size=batch_size, shuffle=False)

然后，我们定义神经网络模型的结构。

class MLP(nn.Module):
    def __init__(self, input_size=784, hidden_size=256, output_size=10):
        super(MLP, self).__init__()
        self.fc1 = nn.Linear(input_size, hidden_size)
        self.relu = nn.ReLU()
        self.fc2 = nn.Linear(hidden_size, output_size)

    def forward(self, x):
        x = x.view(x.size(0), -1)  # 展平输入
        x = self.fc1(x)
        x = self.relu(x)
        x = self.fc2(x)
        return x

我们将模型转移到CUDA设备上，并定义损失函数和优化器。

# 初始化模型、损失函数和优化器
model = MLP().cuda()
criterion = nn.CrossEntropyLoss()
optimizer = optim.SGD(model.parameters(), lr=learning_rate)

最后，我们编写训练循环。

# 训练循环
for epoch in range(num_epochs):
    model.train()
    running_loss = 0.0
    correct = 0
    total = 0

    for batch_idx, (data, target) in enumerate(train_loader):
        data, target = data.cuda(), target.cuda()

        # 前向传播
        output = model(data)
        loss = criterion(output, target)

        # 反向传播和优化
        optimizer.zero_grad()
        loss.backward()
        optimizer.step()

        # 统计
        running_loss += loss.item()
        _, predicted = output.max(1)
        total += target.size(0)
        correct += predicted.eq(target).sum().item()

    # 打印每个epoch的统计信息
    print(f'Epoch [{epoch+1}/{num_epochs}], Loss: {running_loss/len(train_loader):.4f}, Accuracy: {100.*correct/total:.2f}%')

运行此脚本，我们可以在三个epoch后获得约90%的准确率。

深入NumPy实现 🔍

上一节我们使用PyTorch的高层API快速实现了训练，本节中我们来看看如何使用NumPy从零实现相同的网络，以理解其底层原理。

以下是核心步骤的概述：

初始化参数：随机初始化权重和偏置。
前向传播：计算线性变换和激活函数。
计算损失：使用交叉熵损失。
反向传播：计算梯度。
更新参数：使用梯度下降。

以下是关键函数的代码示例：

前向传播（线性层）

def linear_forward(x, w, b):
    return np.dot(x, w) + b

ReLU激活函数

def relu_forward(x):
    return np.maximum(0, x)

Softmax函数

def softmax(x):
    exp_x = np.exp(x - np.max(x, axis=1, keepdims=True))
    return exp_x / np.sum(exp_x, axis=1, keepdims=True)

交叉熵损失

def cross_entropy_loss(y_pred, y_true):
    m = y_true.shape[0]
    log_likelihood = -np.log(y_pred[range(m), y_true])
    loss = np.sum(log_likelihood) / m
    return loss

反向传播（线性层梯度）

def linear_backward(dout, x, w):
    dw = np.dot(x.T, dout)
    db = np.sum(dout, axis=0, keepdims=True)
    dx = np.dot(dout, w.T)
    return dx, dw, db

通过实现这些函数并组合成训练循环，我们可以用纯NumPy复现PyTorch的训练过程，准确率同样能达到约90%。

移植到C语言 ⚙️

上一节我们在NumPy中理解了所有操作，本节中我们将其移植到C语言，为后续的CUDA加速做准备。

C语言实现需要手动管理内存和循环。以下是核心数据结构和函数：

定义神经网络结构体

typedef struct {
    float* weights1;
    float* bias1;
    float* weights2;
    float* bias2;
    float* grad_weights1;
    float* grad_bias1;
    float* grad_weights2;
    float* grad_bias2;
    int input_size;
    int hidden_size;
    int output_size;
} NeuralNetwork;

矩阵乘法函数

void matmul(float* a, float* b, float* c, int m, int n, int k) {
    for (int i = 0; i < m; i++) {
        for (int j = 0; j < n; j++) {
            float sum = 0.0f;
            for (int p = 0; p < k; p++) {
                sum += a[i * k + p] * b[p * n + j];
            }
            c[i * n + j] = sum;
        }
    }
}

前向传播函数

void forward(NeuralNetwork* net, float* input, float* hidden, float* output, int batch_size) {
    // 第一层: input -> hidden
    matmul(input, net->weights1, hidden, batch_size, net->hidden_size, net->input_size);
    add_bias(hidden, net->bias1, batch_size, net->hidden_size);
    relu_forward(hidden, batch_size * net->hidden_size);
    
    // 第二层: hidden -> output
    matmul(hidden, net->weights2, output, batch_size, net->output_size, net->hidden_size);
    add_bias(output, net->bias2, batch_size, net->output_size);
    softmax(output, batch_size, net->output_size);
}

反向传播和参数更新
反向传播需要计算每一层的梯度，并按照链式法则传递。参数更新则遵循梯度下降公式：weight = weight - learning_rate * gradient。

通过完整的C语言实现，我们可以在CPU上运行MNIST训练，虽然速度较慢，但为理解CUDA并行化打下了坚实基础。

使用CUDA加速 ⚡

上一节我们完成了CPU上的C语言实现，本节中我们利用CUDA将其加速。

CUDA实现的核心是将计算密集的操作（如矩阵乘法）移植到GPU上并行执行。我们主要修改以下部分：

设备内存管理：使用cudaMalloc和cudaMemcpy在主机和设备间传输数据。
核函数编写：为矩阵乘法、ReLU、Softmax等操作编写并行核函数。
启动配置：合理设置线程块和网格大小以最大化GPU利用率。

一个简单的矩阵乘法核函数示例

__global__ void matmul_kernel(float* a, float* b, float* c, int m, int n, int k) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (row < m && col < n) {
        float sum = 0.0f;
        for (int i = 0; i < k; i++) {
            sum += a[row * k + i] * b[i * n + col];
        }
        c[row * n + col] = sum;
    }
}

主机端调用

// 设置启动配置
dim3 blockDim(16, 16);
dim3 gridDim((n + blockDim.x - 1) / blockDim.x, (m + blockDim.y - 1) / blockDim.y);

// 启动核函数
matmul_kernel<<<gridDim, blockDim>>>(d_a, d_b, d_c, m, n, k);

通过将前向传播、反向传播和参数更新中的所有关键操作都实现为CUDA核函数，我们可以显著提升训练速度。在相同的超参数下，CUDA版本可以在更短的时间内达到与CPU版本相近的准确率（约90%）。

总结 🎯

本节课中我们一起学习了如何从零开始实现并训练一个MNIST多层感知机。

我们首先使用PyTorch快速搭建了模型，理解了高级API的便捷性。然后，我们深入NumPy，从零实现前向传播、损失计算、反向传播和梯度下降，揭示了神经网络训练的核心原理。接着，我们将算法移植到C语言，处理了内存管理和循环优化，为性能提升做准备。最后，我们利用CUDA进行并行加速，将计算密集型任务分配给GPU，实现了显著的性能提升。

通过这个完整的项目，我们不仅掌握了MNIST分类任务，更深入理解了从高级框架到底层硬件加速的完整深度学习流水线。你可以在此基础上尝试优化CUDA核函数、调整网络结构或探索其他数据集。

12：下一步方向 🚀

在本节课中，我们将总结整个课程，并为希望继续深入学习的你指明方向。我们将介绍一些高级优化概念、推荐学习资源以及活跃的社区，帮助你进一步提升CUDA和GPU编程技能。

如果你已经坚持学到这里，请给自己一些鼓励。这基本上就是本课程的终点了。你做到了，做得很好。

我们接下来将快速浏览一些提示，为你指明继续深入的方向。你可能很难完全掌握所有内容，这完全可以理解。但如果你希望继续，这里有一些额外的资源可以支持你。

在课程的README文件中，有一个关于统一内存和内存架构的章节，我认为这很有用，你可能会感兴趣。但此刻我想主要介绍的是“深入探索”部分。如果你想更进一步，真正了解如何在CUDA和GPU编程（尤其是在深度学习领域）应用深度优化和先进技术，你可以这样做。

以下是几个关键方向：

量化技术

量化是指从高精度数据类型（如FP32）转换到低精度数据类型（如FP16或INT8）的过程。即使从FP32降到INT8，模型仍能保持相当好的性能和精度质量。实现这一点有一些特定的技巧。

很多技巧与数值范围有关。例如，如果你的数值范围被限制在-10到10之间，你就不必担心很多指数值的问题。如果权重初始化稳定且训练过程中数值不会超出这个范围，你可以直接将其作为精度上限，这将是它能达到的最大值。量化本质上就是这种将高精度数字转换为低精度数字的艺术，然后利用这些低精度数字进行非常快速的操作。

INT8运算比FP32快得多，不仅仅是四倍。我们在比较kblas与kblas_lt的32位与16位性能时已经看到了显著的提升。你可以想象INT8会更快，因为它只涉及整数运算，无需担心浮点数或小数点。

量化技术很酷，被广泛应用于当前模型中，例如GPT-4或Llama 4/5B。很多模型实际上都使用了量化，很可能是BF16、FP8甚至Float4。这很酷。

张量核心

张量核心非常棒。我虽然已经提到过，但不能不提。我之所以没有在本课程中详细覆盖，是因为这更像是一个入门课程。我试图在有限的课时内塞入尽可能多的内容，以便你能消化。如果你想继续深入，张量核心显然是下一步。

稀疏性

稀疏性是一个很酷的概念。你可以这样理解稀疏性：假设我有一个数组，它可能像这样：[0, 0, 0, 0, -7, 0, 0, 0, 0, 0, 0, 6, 0, 0, 0, 0]。这就是稀疏的含义：存在大量零值，偶尔出现一些代表重要信息的大数值。

这里的核心思想是，你可以用更少的内存来存储这些数据。这更多是关于内存和计算性能的优化，而不仅仅是质量提升。我们可以使用两个矩阵：一个存储非零值（如[-7, 6]），另一个存储这些值对应的坐标（如[4, 11]）。这样，我们只需要存储4个整数，而不是原来的16个整数，大大减少了存储需求。

想象一下，当你扩展到2D或3D结构时，你将节省数个数量级的内存，这可以非常高效。因此，在设计高性能神经网络时，需要考虑是否能利用稀疏性。如果外部编写神经网络架构的人（例如使用PyTorch时）鼓励稀疏性，并且模型在这方面运行良好，那么这对你来说就非常有利，会让你的工作更轻松。稀疏性是一种性能优化技巧，有机会就应该采用。

CUDA Mode 社区

最后，我决定压轴介绍的是cuda-mode。他们有自己的GitHub、YouTube频道和Discord服务器。这里包含的很多材料都超出了我的课程范围。我的课程更偏向视频辅助，但cuda-mode背后的社区非常棒。这里有真正优秀的工程师和研究人员，不断构建酷炫的东西，社区成员也非常活跃。这是一个绝佳的地方，我绝对推荐你去看看。他们有很多章节，例如Flash Attention、Cutlass、Triton、BEED Kernels、数据处理、张量核心等等。我推荐你加入他们的Discord服务器，你可以在这里找到。里面有很多很酷的小组，例如#beginners频道。超级活跃，比如今天的最后一条消息甚至是几小时前发布的，而这只是一个频道。往下翻，最后一条消息可能就在一小时前。

本节课中我们一起学习了CUDA编程课程的总结与进阶方向。我们回顾了量化、张量核心和稀疏性等高级优化概念，并介绍了《CUDA by Example》、分布式训练文章、MNIST CUDA、Micrograd CUDA、GPU Puzzles以及活跃的CUDA Mode社区等宝贵的学习资源和实践项目。希望这些内容能为你继续探索高性能GPU计算世界提供坚实的基础和明确的方向。

13：课程结束与后续资源 🎬

在本节课中，我们将回顾整个CUDA编程课程，并了解如何获取更多学习资源、加入社区以及联系课程讲师。

课程概述

在之前的课程中，我们系统性地学习了CUDA编程的核心概念，从GPU并行计算的基础到高性能计算的实践技巧。本节是课程的结尾部分，主要介绍课程结束后的学习路径和资源获取方式。

后续学习与社区

如果您喜欢本课程，可以在其他多个平台找到讲师和相关资源。

以下是主要的联系平台和社区：

YouTube：可以找到讲师的频道，观看更多相关视频内容。
X/Twitter：可以在X（原Twitter）上关注讲师，获取最新动态。
Discord：讲师拥有一个Discord服务器，里面聚集了许多志同道合的人。这个社区不仅讨论CUDA，也广泛涉及其他学习主题，鼓励协作与交流。
LinkedIn：可以在LinkedIn上联系讲师。

资源获取方式

上述所有平台的链接通常会在视频描述中提供。

如果描述中没有，也可以在下方描述中提到的GitHub代码仓库里找到相关链接。

课程总结

本节课中，我们一起学习了在CUDA编程课程结束后如何继续深入学习。我们了解了可以通过YouTube、X/Twitter、Discord和LinkedIn等多个平台联系讲师并加入学习社区，同时明确了相关资源的获取途径。希望本课程为您的高性能计算之旅奠定了坚实的基础。

posted @ 2026-03-26 08:20 布客飞龙II 阅读(0) 评论(0) 收藏举报

刷新页面返回顶部

龙哥盟

人最大的痛苦就是说一些自己都不相信的话。

Elliot-CUDA-编程笔记-全-

Elliot CUDA 编程笔记（全）

0：课程介绍 🚀

课程概述

什么是CUDA？

课程最终项目

创建课程的初衷

CUDA的应用场景

课程要求与前提

遇到问题怎么办？

代码与资源

课程范围说明

课程核心收获

额外资源

1：深度学习生态系统概述 🧠

研究与应用框架 🧪

生产与推理优化 🚀

底层硬件编程 ⚙️

边缘计算与嵌入式系统 📱

易用性与工具链 🛠️

编译器基础设施 🏗️

模型与数据集社区 🤗

2：CUDA环境搭建 🛠️

Windows系统配置

安装WSL与Ubuntu

配置Ubuntu环境

下载并安装CUDA工具包

配置环境变量

验证安装

编写并运行测试程序

Ubuntu原生系统安装（补充说明）

3：C/C++ 语言回顾

概述

学习资源

指针

示例 1：基础指针

示例 2：多级指针

示例 3：Void 指针

示例 4：NULL 指针

示例 5：指针与数组

示例 6：指针数组（模拟矩阵）

自定义类型

示例：自定义 size_t 和结构体

类型转换

宏与全局变量

编译器简介

Makefile

调试器 (GDB)

总结

4：GPU硬件入门 🚀

硬件类型概述

GPU发展简史

GPU为何适合深度学习？

CUDA编程核心概念

典型的CUDA程序流程

本章总结

5：编写你的第一个内核 🚀

概述：从硬件规格到软件抽象

CUDA基础：运行时与内存模型

CUDA层次结构：网格、块与线程 🧱

实战：向量加法内核 ✨

深入：矩阵乘法内核 ⚡

性能剖析：使用Nsight Compute 🔍

高级概念：原子操作与流 ⚛️

总结

6：CUDA API 🚀

概述 📋

访问官方文档 🔍

理解“不透明结构类型” ⚫

错误检查宏 🛡️

cuBLAS：CUDA基础线性代数子程序 🧮

cuBLAS的不同版本

代码示例：cuBLAS Sgemm 和 Hgemm

cuBLASLt 注意事项

性能对比实验

cuDNN：CUDA深度神经网络库 🧠

cuDNN的架构：图API与融合引擎

代码示例：cuDNN Tanh激活函数

示例：自定义 `size_t` 和结构体

解析CUDA内核 (`polynomial.cu`)

解析C++绑定代码 (`polynomial.cpp`)

解析Python集成与基准测试 (`benchmark.py`)