RayiYanuTTara-OpenCL-编程笔记-全-

RayiYanuTTara OpenCL 编程笔记（全）

001：OpenCL 1.2 高级概述 🚀

在本节课中，我们将学习OpenCL 1.2规范的高级概述。我们将了解OpenCL是什么、其核心模型（设备、内存、执行模型）以及主机API的基本概念。掌握这些基础知识对于后续深入学习OpenCL至关重要。

什么是OpenCL？🤔

OpenCL是一个用于异构计算的开放标准。在高层次上，它包含一个主机，负责向设备分派命令。这些设备是异构系统的一部分，例如GPU、CPU或Xeon Phi协处理器。主机的唯一目的是告诉设备如何操作，而设备则代表主机执行实际工作。

OpenCL标准简介 📜

OpenCL 1.0规范于2008年12月发布。在2013年7月，OpenCL 2.0临时规范发布。本课程主要基于OpenCL 1.2，其核心概念对于理解后续版本依然非常重要。

OpenCL规范包含几个部分：

• 核心规范：定义了所有符合OpenCL 1.2的实现都必须支持的功能。
• 嵌入式配置文件：为核心规范提供了一个功能子集，主要面向手持设备。
• 扩展：供应商可以添加的额外功能，未来可能被纳入核心规范。

OpenCL主要由两部分组成：

1. 主机API：这是应用程序员从主机端调用的接口，用于管理设备、内存和执行任务。
2. OpenCL C：一种基于C99的语言，用于编写在设备上运行的内核函数。

在主机API和设备编程语言之下，是理解OpenCL如何工作的核心模型。接下来，我们将逐一深入探讨这些模型。

设备模型 💻

设备模型描述了OpenCL如何看待一个计算设备。它不是一个具体的硬件，而是一个抽象的编程模型。

一个OpenCL设备内部可以看作由以下部分组成：

• 全局内存：所有处理单元都能访问的共享内存区域。主机也可以读写此内存。
• 常量内存：所有处理单元都能访问的只读内存区域，访问效率很高。
• 计算单元：设备的子部分，每个计算单元包含多个处理单元。
• 本地内存：每个计算单元内部私有的内存，由该计算单元内的所有处理单元共享。
• 处理单元：实际执行指令的最小单元。每个处理单元都附有私有内存，仅供自己访问。

重要提示：不要简单地将处理单元的数量等同于“处理器”的数量。正确的思考方式是：设备包含N个计算单元，每个计算单元包含M个处理单元。这种分组概念对算法设计至关重要。

内存模型 🧠

OpenCL设备包含多个具有不同特性和功能的内存区域，这与传统的C/C++内存视图有显著不同。

以下是各内存区域的关键特性：

• 全局内存：持久性存储，在多次内核执行之间保持数据。所有处理单元和主机都可访问。公式：所有PE可读/写。
• 常量内存：只读，所有处理单元可高效访问。非持久性。公式：所有PE只读。
• 本地内存：每个计算单元私有，由该单元内所有处理单元共享。用于高效协作。非持久性。公式：单个CU内的所有PE可读/写。
• 私有内存：每个处理单元私有，其他处理单元无法直接访问。非持久性。公式：单个PE私有。

理解这些内存区域的层次和访问特性是编写高性能OpenCL应用程序的基础。

执行模型 ⚙️

执行模型描述了工作如何在设备上实际运行，这是OpenCL中最核心也最具特色的部分。

内核与NDRange

OpenCL在设备上执行称为内核的特殊函数。内核调用有两个部分：

1. 普通的函数参数列表。
2. 控制并行度的执行参数。

主机并不直接执行内核，而是协调设备的执行。它告诉设备：“调用这个函数，这是参数，现在开始运行。” 最通用和有用的启动方式是NDRange。

NDRange的核心思想是：同一个内核函数会被调用很多次，每次调用的参数列表完全相同。调用次数由执行参数决定。

工作项、工作组与索引空间

为了更好地理解，可以将NDRange想象成一个N维的索引空间。主机创建这个空间，然后设备上的每次内核调用会从这个空间中取出一个唯一的索引（一个N元组）来执行。

与执行模型相关的关键概念定义如下：

• 工作项：针对特定索引的内核的一次调用。可以看作是循环中的一次迭代。
• 全局ID：工作项在全局索引空间中的唯一标识。
• 全局工作大小：工作项的总数。
• 工作组：将全局工作大小划分成更小的组。这是连接执行模型和设备模型的关键。
  • 工作组被调度到计算单元上执行。
  • 工作组内的工作项被映射到该计算单元内的处理单元上执行。
  • 工作组内的所有工作项共享该计算单元的本地内存，这使得它们能够高效协作。

工作组的大小是一个设备相关特性，可以通过查询API获得。虽然工作组可以是N维的（例如 2x2x2），但其总大小（各维度乘积）不能超过设备支持的最大工作组大小（一个标量值）。

执行流程图示

从几何角度理解：假设有一个 8x4 的2D全局工作空间，并指定 2x2 的工作组大小。OpenCL会自动将其划分为8个工作组。这些工作组被分配到可用的计算单元上执行。在每个计算单元内部，工作组中的工作项被进一步分配到各个处理单元上运行。

内核执行要点总结：

• 主机提供执行维度，创建索引空间。
• 内核参数可以是值，也可以是全局内存对象。
• 只有全局内存在内核调用之间是持久的。常量、本地和私有内存都是临时空间，每次内核启动都会重置。
• OpenCL实现在如何将工作项映射到处理单元以及如何调度工作方面有很大的灵活性。

主机API 🎛️

主机API是标准的另一部分，为主机提供了控制设备所需的所有接口。

平台、上下文与程序

• 平台：可以看作是一个OpenCL实现或“设备驱动”。它用于发现和管理可用的设备（例如AMD GPU平台、Intel CPU平台）。
• 上下文：为一个特定平台创建的容器，其中包含设备和内存对象。大多数OpenCL操作都与某个上下文相关。一个上下文不能跨多个平台。
• 程序：内核函数的集合。主机需要将OpenCL C源代码编译（或加载二进制文件）为程序，然后从中提取出具体的内核来执行。

异步命令队列

OpenCL采用异步执行模型以实现最佳性能。主机通过命令队列向设备发送命令。

基本流程如下：

1. 主机调用 clEnqueue* 函数（如 clEnqueueNDRangeKernel）将命令（如“运行内核”、“复制内存”）放入与特定设备关联的命令队列中。
2. 该调用会返回一个 CL事件对象，作为该命令的句柄，用于查询其完成状态。
3. 命令可以指定依赖关系，即必须等待之前的一个或多个命令（通过它们的事件对象标识）完成后才能开始执行。
4. 设备从队列中取出命令并执行。
5. 主机可以等待事件或查询事件状态来获知命令完成。

这种依赖关系模型使得主机能够构建复杂的工作流水线，并高效管理多个设备。

总结 📝

本节课我们一起学习了OpenCL 1.2的高级概述。我们首先了解了OpenCL作为异构计算标准的基本定位。然后，我们深入探讨了其三大核心模型：

1. 设备模型：将设备抽象为计算单元、处理单元及多级内存的层次结构。
2. 内存模型：区分了全局、常量、本地和私有内存，并理解了它们的访问特性和持久性。
3. 执行模型：掌握了NDRange、工作项、工作组和索引空间的概念，明白了内核如何被大规模并行执行。

最后，我们介绍了主机API的关键组件：平台、上下文、程序以及基于命令队列的异步执行机制。这些概念为你阅读OpenCL规范和实践编程奠定了坚实的基础。记住，OpenCL编程的核心思想是主机控制，设备执行。在接下来的课程中，我们将学习如何使用OpenCL C语言编写内核代码。

002：OpenCL C语言基础

在本节课中，我们将要学习OpenCL C语言的基础知识。OpenCL C是用于编写在OpenCL设备（如GPU）上运行的代码的语言。它基于C99标准，但为了适应并行计算设备，进行了一些修改和限制。我们将从数据类型、内存区域、向量操作等核心概念开始，逐步构建对OpenCL C的理解。

概述与讲师介绍

我是RayiYanuTara，一名并行编程专家，精通C++、OpenCL和Linux。我专注于将OpenCL应用于实际、有趣的软件项目，而不仅仅是研究代码。我开发了OpenCL中间件，旨在简化OpenCL的使用，同时不牺牲效率。我长期从事高性能计算工作，并可以提供算法设计或OpenCL适用性评估等方面的咨询。

OpenCL C并非传统的C99语言，它存在一些差异。例如，OpenCL C中没有函数指针，也不支持递归。函数调用可能会被内联。OpenCL C并非C语言的子集，它包含一些C语言没有的特性，但整体上非常接近。

OpenCL C是什么？

首先，OpenCL C是一种OpenCL设备编程语言。在之前的课程中，我们介绍了OpenCL的设备抽象模型。OpenCL C可以看作是C编程语言的修改版，专门用于在这些设备上编程。

它是OpenCL编程中的主要角色，你将花费大部分时间与OpenCL C代码打交道。它也是OpenCL规范的另一主要部分。规范的大部分内容都致力于定义OpenCL C。你不需要学习所有细节，我将展示它与你已知的C语言非常相似，你可以将规范用作参考。

本节内容概览

在本视频中，我们将讨论以下主题：

• 类型
• 内存区域
• 向量操作
• 结构体
• 函数
• 内核
  我们还将快速浏览一些值得单独一提的主题。我的目标是为你提供足够的背景知识，让你能够立即开始学习OpenCL规范。

数据类型

让我们从类型开始。这部分内容应该非常熟悉，我假设你在学习本课程前已经了解C语言。

以下是OpenCL C与C语言在类型上的主要区别：

1. 具体表示：OpenCL C提供了具体的数据表示。例如，有符号整数使用二进制补码表示。C语言规范并未规定这一点。
2. 固定大小：OpenCL C中的类型具有固定的大小。C语言则不然。
3. 向量类型和操作：OpenCL C提供了向量类型和操作，这是C语言所没有的，甚至在C++中也没有优雅的实现。
4. 图像类型：OpenCL C提供了图像类型，这是一种不透明类型。你无法直接访问其内存表示，需要使用其他函数从中提取信息。图像是OpenCL C中目前唯一的此类不透明类型。

OpenCL C的类型大部分是C语言中的类型。

以下是整型类型：

• char：8位整数
• short：16位整数
• int：32位整数
• long：64位整数
  每个类型都有对应的无符号版本（如 unsigned char）。有符号整数严格使用二进制补码表示。

一个有用的建议是，创建并使用以下类型定义，这可以确保主机和设备之间的兼容性，稍后你会明白这为什么重要。

typedef int int32_t;
typedef unsigned int uint32_t;
// ... 其他类型定义

OpenCL C中还有浮点类型：

• float：单精度（32位），遵循IEEE 754标准。
• double：双精度。
• half：半精度，可能在图形处理中有用，但不如其他类型精确。
  需要注意的是，double和half类型可能不被所有设备支持。因此，你需要检查你的设备是否具备相应的能力。

必须警惕主机和设备类型的差异。当谈论OpenCL C整数时，我指的是在设备上运行、有标准表示的东西。

假设在主机程序（C99）中声明了 int x，在设备（OpenCL C）中声明了 int y。我们能将x的值复制到y吗？答案绝对是不能，必须非常小心。在设备上，我们使用32位二进制补码表示整数。在主机上，我们甚至不知道整数的位数和表示方式。因此，在进行主机和设备数据交换时要格外小心。

这并非全是坏消息。在主机标准头文件中，你会看到类似 cl_float、cl_int 这样的类型定义。这就是中间件或库可以帮助你确保数据传输正确的地方。

还有一些类型仅限于设备使用。这意味着你无法将它们从主机复制到设备。例如 bool、size_t、ptrdiff_t 等类型。根据C语言规范，它们各有用途，但重要的是你不能在主机和设备之间传输它们。

内存区域

我假设你已经完全理解了我之前关于内存模型的讨论。处理单元可以访问称为私有、常量、本地和全局的内存空间。

在OpenCL C语言中，我们有关键字来指定这些内存区域：__global、__local、__private、__constant。这些概念现在应该完全熟悉了，你可能还不知道如何应用它们，但你已经知道设备模型包含这些内存区域。

让我们看一个声明示例：

__global int* ptr;

首先要认识到，这个关键字指定了内存区域。我们想指向哪里？这个指针应该被视为一个整体：__global int* 表示指向全局内存区域中的一个整数。

这不需要太多时间就能习惯，但你需要熟悉它，并开始思考：当我有一段内存时，它位于哪里？这对于有C/C++背景的人来说可能有点不自然，但在OpenCL C中你应该习惯它，因为你大部分时间都将花在思考内存的位置上。

这里有一个例子。我们有两个指针x和y，它们都指向全局内存空间中的某个整数。我们可以进行赋值 x = y; 吗？是的，我们可以让x指针指向y当前指向的位置。这几乎是标准的C语言，除了那些关键字。

现在考虑，如果我有一个指向全局内存的指针和一个指向私有内存的指针，我们不能将指针重新赋值指向另一个区域，因为这是两个不同的内存区域。所以它不仅仅是指向整数的指针，它是指向全局空间中的整数的指针。这应该很快会感觉自然，但我想指出，我们应该认识到我们正在处理许多不同的内存空间。这是OpenCL中一个非常重要的概念。

然而，我们仍然可以通过值来复制。这样做是可以的。这个操作在物理上对你的设备产生了影响：我们正在从私有内存中取出东西，并发出一个复制到全局内存空间的指令。这不再是一个简单的语句，它产生了影响，是有成本的，可能成本很高。对于这个特定的例子，我们稍后会详细讨论这些操作的实际相对成本，但我想让你清楚，我们实际上是在从多个内存空间的角度思考问题，你真的应该养成这个习惯。从私有内存复制到全局内存是完全可行的。

我们是在内存之间复制，而不是重新分配指针指向的位置。

关于类型就讲这么多，应该很自然了。接下来让我们谈谈向量。

向量类型与操作

当我们谈论处理器内的向量操作时，向量操作是最重要的事情之一。那么，我们如何真正获得良好的向量化代码呢？这是一个复杂的问题，有几种方法。

以下是几种向量化方法：

1. 使用库：问题在于我们把问题转移给了别人。
2. 让编译器完成：如果你有一个向量化/优化编译器，它可以查看代码的某些结构并提取出向量操作。
3. 使用C++类：你可以重载运算符，让代码看起来不错，但问题在于效率可能不会很高。
4. 使用内部函数：例如Intel提供的一些内部函数，但这会变得非常混乱。
5. 汇编语言：更糟糕，谁想写汇编呢？

我想指出，实际上获得向量化的方法很少，而OpenCL C的一个优点就是向量化是一件非常自然、非常容易做的事情。这是我告诉人们的一个原因：即使你不使用GPU，OpenCL C也有很多可提供的东西。如果你只是在CPU上编写代码，而现代CPU都有运行向量指令的能力，OpenCL在这方面很棒。与你使用任何其他方法手动操作相比，它在向量化你的代码方面做得非常好。

让我们看看这些向量类型。它们是“一等公民”。注意左侧的字母N，它代表一个数字，可以是2、3、4、8或16。

这意味着，例如，charN。我们可以去掉N。我们已经有了char类型，但我们还有char2（两个字符）、char3（三个字符）、char4（四个字符）、char8（八个字符）和char16（十六个字符）。这些类型我们可以直接访问。我们可以有有符号或无符号类型，没有unsigned char2，你需要使用类似uchar2这样的形式。

到目前为止看起来还不错。那么，当你进行向量操作时会发生什么？

如果我们有两个标量x和y，都是单精度32位浮点数，我们赋值给z。这和传统的C语言一样，你在这里没学到什么新东西。

但是当你对两个向量进行操作时会发生什么？我们有4宽度的向量x、y和z。当我们做加法时，我们将进行分量级的加法。

假设我们有一个通用的操作符，比如 x cross y。当执行这个操作时，结果z1将等于x1 cross y1，并且这个操作将按分量进行，直到zK等于xK cross yK。

这个“cross”操作符实际上可以是什么？它可以是加法、减法、乘法或任何关系运算符。需要注意的是，关系运算符可能看起来有点令人困惑，因为从数学上讲，你不能自然地将“小于”映射到向量上，你需要使用其他操作来确定向量的顺序。所以我想让你意识到，当你比较两个向量时，结果是一个按分量进行布尔操作的向量。注意这个陷阱。

如果我们混合操作会怎样？我们有一个标量和一个向量，这定义了吗？在这种情况下，x + y，其中x是标量，y是4宽度向量，结果会是什么？这定义了吗？

它是定义的，结果将是一个float4。当混合标量和向量时，标量将被填充扩展，这很自然。所以x将被填充为向量[X, X, X, X]，然后我们将其加到y上。注意这一点，因为这可能不是你想要的特性，编译器可能会愉快地接受这个可能是个错误。

总结一下向量操作：

• 标量与标量操作：我们知道。
• 向量与向量操作：按分量进行。
• 标量与向量操作：标量被填充扩展到相同大小，然后我们使用向量与向量操作。

现在谈谈寻址。如果我有一个向量，我想访问一个分量，我想解包那个向量。我会使用点号，然后调用这个分量。由于我们做编程，你肯定乐意使用十六进制，而且因为我们是从0到16，十六进制非常适合这个。所以 vec.s0 是第一个分量，vec.s8 是第九个分量，最后一个分量是 vec.sf。

这里需要注意，从向量解包分量的角度，就编译器能做的事情而言，你可能没有做到最好。但这将在另一节课中讨论。你可以解包这些向量。

哪些地址是有效的？如果你有一个2向量（例如float2或char2），你可以使用 x.s0 和 x.s1，其他的是错误的。对于3宽度向量，可以使用s0到s2。对于4宽度向量，可以使用s0到s3。对于8宽度向量，可以使用s0到s7。对于完整的16宽度向量，可以使用s0到sf。非常自然，只需记住如何使用这个寻址方案，进行向量操作就很简单了。

但更高层次的问题是，我们为什么想使用OpenCL向量类型？编译器真的会做得更好吗？它能做得多好？

从概念上讲，OpenCL C可以处理这两个语句：z = x + y。如果我们不使用向量，也许我们必须按分量进行，所以可能需要进行四次加法。假设我们的架构有向量指令。让我们把这个OpenCL C片段交给编译器，我们会得到什么？我们会得到设备ISA或汇编语言中的一个完全虚构的语句。假设这是一个虚构的向量指令，我们现在将取寄存器1和寄存器2，将它们相加，并存储到寄存器3。这个加法可以一步完成。假设其他加法需要四步，那么通过使用向量操作，我们确实获得了四倍的性能提升。

使用向量操作，你真的想使用向量操作，但传统上，用你想编写代码的语言实际获得向量操作一直非常困难。使用OpenCL向量类型有一些优势：

1. 清晰的沟通：你和编译器之间就向量操作达成了清晰的沟通，你们都同意这些是向量。
2. 简化代码：它开始看起来比使用内部函数好得多。
3. 获得优异性能：编译器在使用向量的上下文中可以非常好地进行向量化。

我们已经讲完了向量，关于它们需要知道的就是这些。如何使用它们完全是另一回事，但这就是你需要知道的全部。现在让我们谈谈结构体。

结构体与联合体

这是在C语言中聚合数据的典型方式。你可以在OpenCL C中定义结构体，也可以定义联合体，就像C语言一样。

然而，有很好的理由要避免使用OpenCL C结构体，原因与性能有关，另一节课将详细讨论这一点。但请相信我，你可能想避免使用它们。

此外，要非常小心数据交换。记住，主机和设备必须就表示方式达成一致。当我们复制数据时，从主机到设备或反之是一个二进制操作，结构体的二进制布局在主机和设备上必须相同。考虑到对齐，几乎不可能获得正确的布局。可以做到，但根据我的经验，这非常困难。例如，如果你尝试在OpenCL C和C中编译一个结构体，它会根据不同架构的对齐方式以不同的方式填充。你必须开始手动添加填充，这变得一团糟。正如我所说，由于性能原因，你可能也不想使用结构体。

函数

我们已经看到了基本类型，看到了称为向量的特殊类型，也看到了我们可以将东西聚合到结构体中。那么OpenCL C中的函数是怎么回事呢？它们就是普通的C函数，没什么特别的。但是，你不能递归，我在本课开始时提到过递归是被禁止的。函数可能会被编译器内联展开，这不一定会影响你，但你应该知道。

让我们看一个函数的例子。我们有一个float4 add(float4 x, float4 y)函数，这是完全有效的。我们可以把它改成添加一些指针。这也是可以的，但我要指出这里有一些隐含的东西，希望你还记得前几张幻灯片：私有内存空间是隐含的。所以我们实际上有一个add函数，它添加两个私有浮点数。如果我们想添加两个全局浮点数，我们必须添加另一个函数。

请记住，内存区域是函数签名的一部分。如果你正在编写需要以相同方式处理不同内存区域的代码，仅仅将指针作为参数，这可能会有点困难，或者至少非常烦人。请注意这一点。

内核

函数真的没什么可说的，它和你以前见过的一样。而内核，正如你从我对执行模型的讨论中回忆起来的那样，这真的是你调用在设备上执行工作的东西。

我希望我花时间向你解释执行模型能在这里得到回报，因为现在描述内核非常简单。

首先要理解的是，内核实际上是设备执行的入口点。我们熟悉C语言中的这个概念，事实上，我们有int main，它接受特定的参数列表，这是一种固定形式。但在这种情况下，内核只是一个像main一样的入口点，只不过我们可以随意命名这个主要部分，你可以认为在某种意义上有多个main函数。

内核由主机调用。主机将设置控制执行的调用参数，主机也将设置参数列表并提供它，然后它将发出一个主机API调用，说“好的，去执行这个，完成后告诉我”。这就是完整的图景，设备将运行那个函数。

内核参数是指向全局空间某物的指针，或者只是值。这里有几件事……这不完全准确，你会在幻灯片中看到，但你应该这样想。不准确之处其实有点傻……我们假装我没说过。所以内核参数是指向全局某物的指针，或者只是被复制的值。请记住，每个内核启动或函数的每次调用都将具有相同的指定值，它们有相同的参数列表，它们使用索引来弄清楚它们实际上在做什么。

让我们看一个例子，这是你看到的第一个OpenCL C程序，它非常令人失望，因为它所做的只是将两个数组相加。

我们将添加这两个数组。我们得到数组x和y作为输入，我们将结果相加，并按分量存储到z的相同位置。现在我们提供了数组的长度，原因是我们需要终止那些索引溢出、超出范围的工作项。我们通过这里的if语句来实现这一点。

所以首先要注意的是，这个if语句所做的就是说，如果我超出了范围，我就停止，什么都不做。否则，我就做加法。现在，这个索引部分正在调用get_global_id。还记得我说过我们有这些完全唯一的全局ID吗？这就是它的作用。这里的0只是ID维度的指示符。记住，我说过在调用中我们可以有多个维度。这是你完整的内核调用。

这是完全有效的，我要指出kernel void总是必需的。我们不能返回任何东西。我们只是调用它，设备没有地方返回，它必须为此使用全局变量。__kernel只是将其标记为一个特殊函数，告诉编译器“嘿，你可以调用我，我是设备的入口点”，然后它可以继续构建所有必需的代码。

让我们回顾一下我们的执行模型。记住执行有维度：全局工作大小、全局偏移、工作组大小。我花了这么多时间解释这些概念，因为当你实际编写内核时它们非常重要。有一些相关的函数可以访问这些信息。再次记住，参数列表是固定的，所以弄清楚我们在做什么的唯一方法是使用外部函数来询问运行这些函数的系统我们在哪里。

OpenCL C中的相关函数是get_global_id，其中n是维度。这将给我们第n维的工作项ID。我们还可以使用get_global_offset，我在之前的讨论中提到过偏移量，它是一个被加到所有索引上的数字。我们还有get_local_id，它只是告诉我们我在我的工作组中是哪个工作项。这些只是众多函数中的几个，我鼓励你去查看文档，看看其他的，但这些是你将要使用的几个。我想说，你最常使用的可能是get_global_id，你可能比使用其他函数更频繁地使用这个。

现在让我们看另一个例子。这实际上是同一个例子，但我想让你快速看一下我们再次使用这个函数的事实。我想强调，我们基于这个调用弄清楚了自己的位置。顺便说一下，我还要指出，我使用的是size_t，它是设备特定的，根据设备可能是32或64位，但这没关系，因为它不在我们被外部调用的列表中。我还要指出，我继续使用我建议你使用的类型定义，因为这会让你的生活轻松一点。

本地内存

现在让我们谈谈本地内存。这就是我们如何修改我们的执行概念，将本地内存引入到这里的图景中。本地内存由工作组内的工作项共享。本地内存可能是硬件中的物理存在，也可能不是。但因为它是一个模型，它并不总是以相同的权重映射到所有硬件。在GPU上，它将映射到物理的东西。所以当你可以的时候，你想使用本地内存。通常，本地内存是获得顶级性能的关键。

那么我们如何声明使用本地内存的东西呢？有几种方法可以做到这一点。

第一种方法是提供一个参数。你会看到这里我提供了一个指向本地内存的参数。记住本地内存……事实上，你可以在参数列表中这样做。记住我说过参数列表要么是指向全局某物的指针，要么是指向值的指针，要么是值本身。然后我在这里的讨论中有点磕巴，我不完全确定如何提及这个例外……当你调用函数来声明本地内存时……抱歉，当你在排队内核并要调用内核时，当你有一个本地值或变量时……问题是，我使用的术语存在于许多其他地方。当你有一个指向本地某物的指针参数时，主机一方的期望并不是你提供一个真正的指针，而是提供该内存应该有多大的大小。所以，你所能做的只是提供指向全局某物的指针或值，这仍然是正确的。

现在我使用了两个索引，我使用了本地信息。那么这是在做什么呢？让我们看一下。我获取一个全局索引，同样，如果我溢出了，我什么都不做。我在这里做的是将某些东西加载到暂存空间中。所以我取我的全局索引，将某些东西加载到暂存空间，然后暂存空间将从0到……哦，这不是我期望的，抱歉……它将从0到K，例如，其中K是工作组大小。所以这个索引是不同的。我将我的全局值加载到一个本地值中，然后我把它写回去。这不是最有趣的内核，我只是想向你展示如何使用本地内存。

还有另一种使用本地内存的方法，你可能更喜欢。

就是，它更有意义一些。当我可以的时候，我更喜欢这种方法。我只是直接说，嘿，我正在使用一个局部变量，这是暂存空间。但我要指出，即使它看起来是在单个工作项内部，实际上它是在工作组内所有工作项之间共享的。所以这在语法上表现为好像是一个工作项在执行这个，它没有执行任何东西，它只是一个分配。所以这是所有工作项都会看到的分配，在这个例子中它的大小固定为1024，一切都和之前的例子完全一样。所以这是做你刚才在另一个例子中看到的完全相同事情的另一种方法。

常量内存

常量内存是只读内存，由所有工作项共享。我们之前看到它读取速度非常快，但空间相对较小。在某些情况下它很有用，如果你想到查找表，你现在应该真正考虑常量内存，因为常量内存相对较小，如果你有一个巨大的查找表，它就没那么有用了。但我想让你记住，常量内存和本地内存类似。

你如何填充常量内存呢？你可以直接把它放在你的文件中，或者你可以把它作为一个参数传递，将一个指向全局内存区域的指针传递给它。

内核限制

内核有一些限制。首先要认识到的是，内核可能在设备上并发执行。当我说内核时，我指的是add内核和f内核，两个不同的内核（不是同一个内核中的工作项）。它们可能共享同一个设备，但它们不能合作，完全不能。任何让它们尝试合作的行为都是编程错误。你不应该这样做，这是未定义行为。

单个内核能做的事情确实有限。你能写一个像main那样做所有事情的函数吗？不，你不能。你可能需要启动几个内核来完成一个特定的工作。当你实际看到如何使用OpenCL编写算法时，你会看到这一点。现在，你只是看到了游戏规则，语法。

内核不能自己分配内存。所有东西都是固定就位的，然后它们只是运行，直到我之前描述的那个索引空间被清空。所以所有东西都设置好了，内核只是不断运行、运行、运行，直到一切完成，它们终止。它们不能分配内存，因为这会违反这种哲学。

内核属性

有一些属性，这些非常重要。属性是你在文档中可以看到的东西，你会在更高级的演示中看到。我所做的一切都是为了构建一个更高级的讨论，以便你能快速进入使用OpenCL编写非常快的代码。

你可以拥有的一个属性是向量化的提示，关于向量的宽度。这允许你向编译器传达一些它可能自己无法弄清楚的东西，关于你想象中向量化的方式。

另一件你可以做的事情（顺便说一下，还有其他一些属性，但最重要的两个）是reqd_work_group_size。这强制了工作组的大小，这对性能非常有用。你会看到在我编写的每一个试图获得最高速度的内核中，我都会使用这个。这将强制一个特定的工作组大小。发生的情况是，因为编译器知道工作组大小是64（64是AMD的一个特殊数字，不管它是什么），如果它看到工作组大小是64，它可以应用非常特殊和具体的优化，并在寄存器定位等有趣的事情上做得非常好，当我们讨论这些模型如何被转换（抱歉，是转换）成实际代码的实际方式时，我们会深入探讨。

快速主题

我们已经差不多讲完了OpenCL C。现在有一些快速主题，我不会详细讨论，但我想为你提供足够的背景知识，让你可以自己去学习。

OpenCL确实提供了对图像的支持。你可以加载图像，用它做点什么，你可以写图像。这些东西在OpenCL C中直接支持，你会看到规范中有相当一部分是关于图像的。我个人不喜欢OpenCL中的图像，并不是我不喜欢图形或者我对图像有问题，问题在于我认为它们代表了一个更深层次的缺陷。这是一个更高级的话题，如果你看到摘要笔记，你可以阅读我写的关于SPIR的文章中我对图像问题的看法。所以它们存在，但我不打算涵盖它们，我实际上从未使用过它们。它们不……图像有一些有趣的特性，但我做的是通用计算，我做的是非图像处理的事情。如果我在做图像处理，我会使用它们。现在，如果你想了解图像，可以在规范中阅读，关键是我在开始时提到的：它们是不透明类型，我们不知道它们的内存表示。这是出于性能原因，关于图像就讲这么多。

OpenCL C有内置函数。我们有许多内置函数可以直接调用。它不完全像一个标准库，但有点像标准库。我们有工作项函数，可以找出我们排队时的参数；有数学函数；有整数函数；有几何函数。你真的应该看看文档，我不会通过遍历每个函数来让你感到无聊。但那里有一些真正的好东西，比如叉积或点积，你有机会真正使用它们来获得良好的性能，因为这些实现很可能由供应商以该架构最快的方式实现，甚至可能调用特殊指令，我不知道，这取决于设备和实现。

同步。我的意思是，在OpenCL C中，你还没有看到，它真的是一种并行编程方法，同步非常重要。它也是一个非常复杂的话题，你必须等着我做另一个视频，我将专门讨论同步。OpenCL C提供了原子操作。OpenCL C中的并行性有点棘手，你真的会开始欣赏像我这样有免等待算法设计背景的人，因为这意味着我可以编写适合OpenCL C编程模型的算法。然而，如果你习惯使用锁或互斥锁，无论你怎么称呼它们，你在OpenCL中会非常困难。这并不是说你有困难是因为有什么你不知道的东西，而是因为OpenCL的设计是为了性能，为性能所做的一切几乎排除了锁的使用。所以你必须小心。

我把这些话题分成小块的原因是，在这一点上，你已经看到了这么多OpenCL C，现在我可以讨论执行模型了。我可以说，好吧，让我们谈谈工作组，谈谈工作项。事物如何映射？计算单元的接口是什么？计算单元如何看待事物？处理单元如何看待事物？一切如何关联？经过这两个小时的课程，你现在有足够的背景知识来欣赏所有额外的话题。我真的想做这些课程，它们对我来说是基础，它们为你提供了一些东西，现在我们可以进行一些真正有趣和严肃的讨论了。

我要在本次讨论中留给你的最后一件事是关于扩展。OpenCL C标准有扩展，可以向其中添加东西，这些是你可以通过使用#pragma指令启用的额外功能。和往常一样，对于#pragma或任何尚未完全进入规范的外部东西，要小心，运用良好的判断力，安全第一，注意你如何设计你的代码，因为你可能会发现你最喜欢的扩展消失了，或者使用扩展可能不是一个好主意，因为会导致供应商锁定。所以当你使用它时要意识到，你可以浏览提供的扩展，它们很有用，其中一些值得一看。

总结

首先，恭喜你，你现在理解了OpenCL及其相关模型，以及OpenCL C的基础知识。我鼓励你出去阅读规范，拿起来，随便翻翻，你现在拥有了一切真正理解正在发生的事情所需的东西。请记住，一切都需要一些时间才能消化，如果有什么不清楚的地方，或者我没有描述得尽可能好，请随时给我留言，我可能会制作另一个视频来修正你觉得我没有做好的地方，或者我可能会写一篇文章来帮助你更好地理解那个概念。所以我鼓励你订阅我的博客，你可以在这个视频的“关于”部分查看，我会在下一张幻灯片中发布链接。我也可以作为顾问提供帮助，包括帮助你实际获得良好的性能，如何构建你的程序以便重用，OpenCL的设计模式是什么。我正在谈论这些，向公众传达这些信息需要时间。但如果你雇我作为顾问，你今天和现在就能得到我的建议的好处。

我叫Aion，非常感谢你查看这些内容，请随时给我发邮件或查看我的网站。当有新视频发布到这个频道和那个网站时，我会发布通知。

非常感谢你的时间，晚安。

003：GPU架构详解 🚀

在本节课中，我们将深入探讨GPU的通用架构。虽然会以AMD 7970 GPU为例，但所涉及的核心概念同样适用于NVIDIA或其他厂商的GPU，甚至与CPU上的SIMD指令集原理相通。理解GPU的工作原理，将为我们后续讨论针对此独特架构优化的算法和数据结构奠定基础。

讲师简介

我是并行编程领域的专家，精通C++、OpenCL和Linux。我专注于后端服务器开发与性能压榨，致力于解决现实中的软件工程问题，并进行OpenCL中间件开发。我在高性能计算领域拥有多年经验。如果您需要相关咨询或帮助，可以联系我。

课程动机与背景

本次讲座的主题是GPU架构。我们首先思考一下，为何要将原本用于图形处理的GPU用于通用计算。

GPU通用计算（GPGPU）的发展颇具偶然性。硬件厂商和整个行业无意中催生了一个从未被设想过的GPU角色。GPU的架构最初由游戏渲染和像素处理驱动，但其构造的演变恰好允许我们进行非常有趣的计算。今天，你将看到的GPU架构，需要你意识到市场营销信息可能具有误导性。我们旨在拨开迷雾，理解GPU真正新颖和有趣之处，以便你能基于此做出明智的硬件选择。

OpenCL视角下的GPU架构

这里讨论的是OpenCL所呈现的GPU架构视图，我们只关注对OpenCL开发暴露的特性。

你应该还记得这个设备模型图。在GPU上，与之前模型不同的是，常量内存空间和全局内存空间是由硬件实际实现的。显卡上存在物理的全局内存和常量内存。

现在，让我们深入到计算单元内部。这与OpenCL的计算单元模型略有不同，我们稍作修订以展示GPU的实际工作方式。

在GPU的计算单元内部，有一组处理元件。本地内存和私有内存同样由硬件实现，这与CPU架构不同。私有内存虽然被所有处理元件共享（在硬件上作为一个寄存器文件池分配），但各个处理元件仍然无法互相访问彼此的私有内存。

此外，每个计算单元还有一个工作调度器。GPU的执行模型比较特殊：一个计算单元内所有处理元件的指令指针是锁定在一起的。

波前与锁步执行

让我们通过一个假想的低级汇编指令序列来追踪执行过程，以理解其含义。

假设有四个处理元件。它们都从各自的私有内存中读取寄存器R2和R1的值，然后执行相同的加法指令，并将结果写入各自的R3。你可以将此视为一个4宽度的SIMD指令。接下来，它们又一起执行乘法指令。即使在获取内存时，每个处理元件都是独立操作自己的数据，但所有处理元件必须执行相同的操作。

NVIDIA将这种锁步执行的集合称为“Warp”，AMD则称之为“Wavefront”（波前）。波前是来自同一工作组、指令指针锁定在一起执行的一组工作项。你可以将波前视为工作组中被切割出来、一起执行的最小单元。

分支与线程发散

既然指令指针被锁定，那么当遇到条件语句时会发生什么？

考虑一个条件赋值语句：if (a < b) f = x; else f = y;。每个工作项都有自己的a和b副本。

1. 理想情况：如果波前内所有工作项的条件判断结果一致（例如都为真），那么它们将一起执行f = x;，没有问题。
2. 线程发散：如果有一个工作项的条件判断为假，情况就复杂了。GPU的处理方式是：对于条件为真的工作项，执行赋值操作；对于条件为假的工作项，它也会“参与”执行这条指令，但会设置一个掩码，使其写操作无效（不产生副作用）。然后，当执行else分支的f = y;时，之前条件为假的工作项会执行赋值，而条件为真的工作项则掩码其写入。这被称为线程发散。

线程发散会导致性能下降，因为部分处理元件在“空转”。为了缓解这个问题，OpenCL提供了内置的select函数，它可能被编译成一条条件移动指令，效率更高。但请注意，如果条件分支涉及复杂的函数调用，使用select可能使代码难以阅读，此时编译器通常也能很好地处理if语句。

GPU内存层次与访问代价

现在我们来讨论GPU内存。让我们通过一个思想实验来理解内存访问的代价。假设每个处理元件每秒执行一条指令。

考虑一个简单的操作：z = x + y;，其中x和y从全局内存读取，结果z写回全局内存。

• 从全局内存读取两个4字节整数（共8字节）：耗时 57秒。
• 执行加法运算：耗时 1秒。
• 将4字节结果写回全局内存：耗时 28秒。
• 总耗时：86秒。其中只有1秒用于实际计算，效率仅为 1/86 ≈ 1.2%。

如果将数据换成8字节的长整型：

• 读取两个长整型（16字节）：耗时 114秒。
• 加法：1秒。
• 写回一个长整型（8字节）：耗时 57秒。
• 总耗时：172秒，计算效率降至 1/172 ≈ 0.6%。

以下是访问不同内存空间所需时间的对比表（以“1次操作=1秒”为基准）：

内存类型	访问32位值耗时	访问64位值耗时
全局内存	28.6 秒	57.2 秒
常量内存	1.0 秒	2.0 秒
私有内存	0.3 秒	0.6 秒
本地内存	0.5 秒	1.0 秒

注意：上表中64位值的访问时间只是简单地将32位值的时间翻倍，这在实际中是一个“谎言”。因为像AMD 7970这样的GPU，其内存合并访问优化是针对32字节步长设计的，64位访问的性能可能更差。

一个提升效率的技巧是：增加每次内存访问后执行的算术逻辑单元操作数量。如果内存IO开销相对固定，那么通过进行大量计算来“分摊”这个开销，就能提高ALU利用率。例如，在读取数据后执行一百万次操作，那么内存访问的开销就几乎可以忽略不计，ALU效率接近100%。这在整数分解、密码分析等计算密集型任务中很常见。

延迟隐藏与工作调度

内存访问延迟是我们的主要问题。为了实现100%的ALU利用率，我们需要让处理元件一直有工作可做，而不是空等数据。

关键思路是：让计算单元超载多个工作组。这是通过工作调度器实现的。

想象一个计算单元有一个工作池。最初只放入一个工作组（WG 0），那么当该工作组因内存请求而阻塞时，计算单元就闲置了。

如果我们向工作池中放入大量工作组（例如，容量为40个）。调度器会这样工作：

1. 从池中选取一个“就绪”的工作组（例如WG 68）开始执行。
2. 执行固定周期或直到遇到内存请求。
3. 当遇到内存请求时，将该工作组标记为“阻塞”并放回池中。
4. 立即从池中选取另一个“就绪”的工作组（例如WG 23）执行。
5. 如此反复。当某个工作组等待的内存数据到达时，它会被重新标记为“就绪”，并有机会被再次调度执行。

这样，通过在不同工作组之间快速切换，计算单元几乎总在执行有用的计算，而长内存延迟被其他工作组的工作所“隐藏”。这就是延迟隐藏的核心原理。

需要澄清的是，延迟隐藏的实际调度单位是波前，而非整个工作组。一个工作组可能包含多个波前。

占用率

每个计算单元的工作调度器有固定数量的波前槽位（例如AMD 7970是40个）。占用率是指实际可以同时驻留在计算单元上的波前数量与最大可能波前数量之比。

例如，如果最大波前槽位是40，而你的内核设计允许同时运行30个波前，那么占用率就是 30 / 40 = 75%。

占用率是一个重要的性能指标。高占用率通常意味着有更多的工作可以用于隐藏延迟。但请注意，如果内核本身的计算与内存访问比例极高（计算密集型），那么即使占用率低，也可能获得高性能，因为ALU本身已经很忙。

那么，是什么限制了占用率？主要是私有内存和本地内存的使用量。因为这些内存资源在所有处理元件/波前之间共享，且总量固定。

在AMD 7970上：

• 每个计算单元有 256 KB 私有内存（寄存器文件）。
• 每个计算单元有 64 KB 本地内存。

内核使用的这些资源越多，能同时驻留的波前就越少。

占用率计算示例

我们通常将工作组大小设置为波前大小（在7970上是64），以便于分析。

计算最大波前数量的公式如下：

最大波前数 = (可用内存资源总量) / (每个波前所需内存资源)

其中，每个波前所需内存资源 = 每个工作项所需内存 * 波前大小（工作项数）

示例：为了在AMD 7970上实现最大占用率（40个波前），每个工作项最多能使用多少私有内存？

每个工作项最大私有内存 = 总私有内存 / (最大波前数 * 波前大小)
                        = 256 KB / (40 * 64)
                        = (256 * 1024 字节) / 2560
                        ≈ 102.4 字节

这意味着，每个工作项大约只能使用 102字节 的私有内存（约25个32位变量），才能达到最大占用率。这包括了内核运行本身所需的基本寄存器。因此，为了高占用率，必须精打细算地使用私有变量。

本地内存的计算方式类似，只需将总容量替换为64 KB即可。

内存通道与合并访问

全局内存并非一个所有计算单元都能无冲突访问的统一池。实际上，它被划分为多个分区，每个分区连接一个内存通道。在AMD 7970上，有32个计算单元和12个内存通道。

每个全局内存地址都映射到特定的内存通道。当多个计算单元请求访问映射到同一通道的内存地址时，这些请求会被序列化，导致性能下降。

根据鸽巢原理，32个计算单元向12个通道发起请求，必然有通道收到多个请求。为了获得最佳性能，硬件厂商推荐使用合并访问模式。

合并访问是指：相邻的工作项访问相邻的全局内存地址。这种模式能让内存控制器最有效地工作，实现高带宽。

相反，随机或跨步很大的内存访问模式会降低性能，可能导致请求集中在少数通道上。

对于性能优化，我们的建议是：

1. 优先设计算法以实现合并访问。
2. 如果无法避免非合并访问，则尝试在内存操作之间插入足够多的ALU操作，利用延迟隐藏来减轻性能损失。
3. 最终，应依赖实际的基准测试来指导优化，因为理论性能可能与实际硬件行为有差异。

课程总结与展望

本节课中，我们一起深入学习了GPU的核心架构：

1. 锁步执行与波前：理解了GPU计算单元内处理元件如何以波前为单位，锁定指令指针一起执行。
2. 线程发散：认识了条件分支导致的性能问题及其原理。
3. 内存层次与代价：通过思想实验，直观感受了全局内存访问的巨大延迟，以及本地/私有内存的速度优势。
4. 延迟隐藏：掌握了通过超载工作组、利用工作调度器在多个波前间切换，以隐藏内存访问延迟的核心机制。
5. 占用率：学会了如何计算和分析占用率，并明白其受私有/本地内存资源限制。
6. 内存通道与合并访问：了解了全局内存通过多通道访问的物理现实，以及合并访问对性能的关键影响。

到目前为止，我们的学习路径是：OpenCL高级概述 -> OpenCL C语言基础 -> GPU架构。我们尚未深入讨论并行编程中工作项间的协作（如原子操作、同步）、如何具体测量和获取高性能、以及如何编写良好的OpenCL软件工程实践。这些主题将在未来的课程中陆续展开。

希望本课程为你理解GPU编程打下了坚实的基础。要真正掌握性能调优，还需要大量的实践和经验积累。

004：OpenCL简介 🚀

在本节课中，我们将要学习OpenCL的基础知识。我们将探讨OpenCL是什么，它如何工作，以及为什么它对现代高性能计算至关重要。通过本教程，你将了解OpenCL的核心概念、应用场景以及它与其他技术的关系。

什么是OpenCL？

OpenCL代表开放计算语言。它最初由苹果公司在2008年提出，并由包括NVIDIA、AMD、英特尔在内的多家大型公司共同开发其规范。OpenCL本质上是一个规范，而非特定的技术实现。该规范由Khronos集团维护，该集团也负责维护OpenGL等技术的规范。

由于OpenCL是一个规范，这意味着要实际使用它，必须有人根据规范实现相应的库、框架和资源。这与OpenGL的工作方式非常相似，它是一个开放标准，任何厂商都可以基于其硬件和软件编写自己的实现。实现者只需确保其实现符合规范的所有最低要求，即可拥有一个兼容的实现。

为什么需要OpenCL？

计算性能的重点已从时钟速度转向核心数量。过去，通过提高CPU的时钟速度就能轻松获得性能提升。而现在，人们开始寻求使用多核处理器，即在一个系统中集成多个CPU核心。

然而，编程范式并未跟上这一变化。我们一直专注于如何从单个CPU中榨取最大性能，编写的程序和算法也基于此，而没有充分考虑如何在多核环境下高效地分割算法和共享数据。OpenCL旨在解决这个问题，它像一种“粘合剂”，让你能够访问计算机中的所有硬件资源。

OpenCL是一个编程接口，其核心理念是：既然系统中有这么多闲置的计算资源，为什么不尝试利用它们呢？因此，OpenCL旨在支持通用目的的并行计算，而不仅仅是多媒体或图形应用，也可以用于科学计算等任务。

OpenCL的核心特点

上一节我们介绍了OpenCL的诞生背景和目标，本节中我们来看看它的几个核心特点。

• 设备无关性：OpenCL被设计为与设备无关。这意味着规范本身不规定它必须在何种设备上运行。只要硬件能够满足规范的要求，它就可以成为一个OpenCL设备。常见的设备包括CPU和GPU，但也可能是DSP芯片、FPGA或任何嵌入式处理器。
• 代码可移植性：作为一个开放标准，你的OpenCL代码应该能够在不同厂商的实现之间移植。只要目标平台的支持符合规范，你的代码就应该能正常工作，这与OpenGL的理念一致。
• 开放规范：规范由Khronos集团管理，没有单一公司控制它。这意味着你不必担心技术变得封闭或专有。虽然具体实现可能是专有的，但规范本身是开放的。

OpenCL的目标与能力

OpenCL旨在成为一个简洁、高效的API，用于访问系统中所有设备以进行通用、高性能的计算。它基于C99语言，增加了一些额外的数据类型、内置函数和限定符。你可以将其视为一个线程管理框架，它帮你处理创建、销毁线程以及锁等底层细节，让你无需操心。

它需要易于使用、轻量且高效，不应给系统带来显著负担。更重要的是，它需要提供一定的保证，例如在不同实现间能获得相同（或至少满足最低精度要求的）数值结果，并提供具有确定精度的数学函数。

那么，OpenCL可以用在哪些地方呢？它的应用非常广泛：

• 科学计算
• 图像和视频处理
• 医学成像
• 金融服务（如高速交易、金融模型分析与生成）

简而言之，OpenCL适用于任何数据并行且计算密集型的算法，即那些需要大量时间和计算资源才能完成的任务。

理解数据并行计算

当我们谈论数据并行计算时，指的是一个巨大的技术谱系。从粗粒度到细粒度，包括：

• 网格计算
• 在单个系统内使用MPI
• 使用Pthreads或OpenMP的标准线程模型
• 非常细粒度的SIMD（单指令多数据）并行，如SSE、AltiVec等向量引擎

在OpenCL的上下文中，我们主要关注后两种：类OpenMP/Pthread的线程模型以及SIMD。当然，这些模型可以混合使用，但OpenCL的核心在于高效处理这两种并行模式。

在OpenCL中，数据并行计算主要涉及两个方面：

1. 任务并行：可以看作是在单个系统内的粗粒度分发模型。例如，你的程序有多个任务（如处理多张图片），你可以使用OpenCL将这些任务分配给不同的CPU核心。
2. 数据并行：这是我们重点关注的。例如，我们有一个数字数组，希望对每个元素执行相同的操作（如取绝对值）。计算这个元素的绝对值不需要知道另一个元素的值，因此这是一个完美的数据并行任务。

让我们看一个更具体的例子：盒式滤波器（用于图像模糊）。算法是取一个像素周围一个“盒子”区域内的所有像素，计算它们的平均值，并将这个平均值赋给中心像素。然后在整个图像上滑动这个盒子，为每个像素生成新值。

关键在于，每个盒子区域的计算是独立的，因为它们从原始图像读取数据，并将结果写入一个新图像缓冲区。由于读写位置分离，我们无需担心同步问题。虽然这个算法本身可能不是最优的，但它清晰地展示了数据并行计算的特点：大量相同的、可独立执行的计算单元。

OpenCL与其他技术的关系

OpenCL被设计为与OpenGL协同工作。OpenGL是图形编程语言，而OpenCL则专注于数值计算。它们是“姐妹”技术，OpenCL可以轻松地与OpenGL共享数据缓冲区，因为两者使用相同的内存位表示。这意味着你可以用OpenCL进行数值计算，然后将结果无缝传递给OpenGL进行渲染，整个过程性能开销极低。

那么，OpenCL不能很好地处理哪些问题呢？

• 顺序问题：本身不具备并行性的顺序问题不适合用OpenCL，你无法从中获得想要的性能提升。
• 需要大量同步的计算：那些需要频繁来回通信、同步数据的计算可能也不适合OpenCL。虽然OpenCL提供了同步机制，但你更希望数据尽可能独立。
• 设备依赖限制：某些计算可能因特定设备的限制或需要特殊处理而无法使用OpenCL。

OpenCL并非“万能钥匙”，它旨在解决一个特定问题：如何尽可能简单、便携地利用计算机中的全部计算能力。

OpenCL与CUDA

OpenCL常被拿来与NVIDIA的CUDA进行比较。CUDA是一个强大、先进的GPGPU编程接口，极大地推动了GPU计算的主流化，让开发者能够用C/C++（加上一些特殊修饰符和语义）在GPU上运行代码。

两者主要区别在于：

• 专有性与开放性：CUDA不是设备无关的，它仅适用于NVIDIA的硬件，并由NVIDIA完全控制。而OpenCL是开放标准。
• 代码移植：使用CUDA意味着你将主要与NVIDIA硬件绑定。幸运的是，NVIDIA显卡销量很大。

但重要的是，你不必在CUDA和OpenCL之间二选一。实际上，从一种技术迁移到另一种并不困难，因为核心的计算部分（内核）非常相似，只需进行一些语言语义上的微小修改。你在一种技术上的投入，可以相对快速、容易地在另一种技术上获得回报。

如何获取与使用OpenCL？

OpenCL的第一个主要实现随Mac OS X 10.6 Snow Leopard（于2009年8月28日发布）推出，作为一个系统框架提供。这意味着对于开发者，它已经内置在系统中，无需用户额外安装。只要用户运行10.6或更高版本，你的应用就能使用它。

NVIDIA和AMD也在开发自己的实现。NVIDIA的测试版已经有一段时间了。这些实现也让你能在其他操作系统和平台上使用OpenCL。

什么样的计算机能运行OpenCL计算？

• 对于苹果电脑：任何运行Mac OS X 10.6或更高版本的Mac都具备OpenCL能力（因为它有CPU）。要使用GPU，则需要较新的机型，如Mac Pro、24英寸iMac、新款MacBook Pro等。
• 目前所有NVIDIA出货的显卡都支持OpenCL。

在Snow Leopard的架构中，OpenCL位于图形和媒体技术层，与QuickTime、Core Video、Core Image和OpenGL并列，因为它与OpenGL紧密结合良好。

为什么关注GPU？优势与挑战

人们如此关注OpenCL和GPU计算，是因为GPU是浮点计算的“怪兽”。它们专为数值计算设计，具有高度可扩展性。如果你的算法是数据并行且优化良好，它在GPU上能表现得非常出色，并随着更强大的GPU发布而持续扩展。

例如，比较常见的Core 2 Duo CPU（约45 GFLOPs）与高端的NVIDIA GTX 285显卡（近1 TFLOPs），后者浮点性能有数量级的提升，而价格并非遥不可及。

然而，在使用GPU时也需注意一些挑战：

• 数据传输瓶颈：数据需要通过PCIe总线从主机内存传输到显卡内存，这个速度（约3-4 GB/s）远慢于显卡内部的内存带宽。在传输少量数据（如16字节）的时间里，CPU可能已经完成了很多工作。因此，你需要确保计算任务足够繁重，以抵消数据传输的开销。
• 调试与错误处理：GPU上的调试和复杂错误处理不如CPU成熟，虽然正在改善，但目前仍有一定挑战。
• 数据组织：GPU对数据的组织方式有特定要求。为了获得最佳性能，你需要按照GPU喜欢的方式（而非人类思维中“整洁”的方式）来组织数据。遵守这些规则，你就能获得惊人的性能回报。

总结与演示

本节课中我们一起学习了OpenCL的基础知识。OpenCL是一个开放的、设备无关的规范，旨在让开发者能够轻松、便携地利用系统中所有计算设备（尤其是GPU）进行通用目的的并行计算。它简化了GPU编程，并与OpenGL无缝集成。

为了让你更直观地感受OpenCL的威力，请看一个实际演示：将一个用于计算生物分子静电特性的真实科学代码（APBS软件的一部分），从CPU移植到GPU。该计算在单个CPU核心上运行约60秒，在16个CPU线程上并行后缩短至约4.8秒，而在单个NVIDIA GTX 285 GPU上运行仅需约180毫秒，并且计算结果与CPU完全一致。这充分展示了OpenCL在加速数据并行计算方面的巨大潜力。

希望本教程能帮助你理解OpenCL的核心价值，并激发你探索如何在自己的项目中利用这项强大的技术。

005：OpenCL基础概念

在本节课中，我们将学习OpenCL的核心基础概念，包括其架构、关键对象模型、内存层次结构以及一个典型OpenCL程序的基本执行流程。我们将通过对比CPU与GPU的硬件差异，帮助你理解OpenCL的设计哲学。

概述

OpenCL是一个用于编写跨平台并行程序的框架。它允许程序在多种处理器（如CPU、GPU）上运行。本节将介绍OpenCL的基本对象、工作单元模型、内存空间以及编程步骤。

---

硬件对比：CPU与GPU

上一节我们介绍了OpenCL的跨平台特性。本节中，我们来看看支撑其性能的硬件基础，特别是CPU与GPU的核心差异。

从硬件层面看，CPU与GPU的设计目标不同。下图展示了两种处理单元的核心布局：

• CPU（如Core 2 Duo）：其芯片（Die）上有大量面积用于高速缓存（Cache）。这些缓存的主要目的是隐藏从主内存（RAM）读取数据的延迟。CPU将频繁使用的数据存储在缓存中，以加速访问。
• GPU：其芯片上用于缓存的空间非常少。绝大部分晶体管都用于数学计算和处理单元。例如，图中区域包括着色器处理单元、内存控制器、纹理处理单元以及硬件线程管理单元。GPU通过硬件管理大量线程，这是其能高效执行海量并行任务的关键。

这种设计差异意味着：CPU擅长处理复杂的、串行任务，并依赖缓存优化；而GPU则将资源倾注于并行计算能力，通过硬件级线程管理来隐藏内存访问延迟。

---

OpenCL核心对象

理解了硬件背景后，我们来看看OpenCL的软件抽象。OpenCL对象主要分为三大类。

计算设备

计算设备是指能够执行数据并行程序的处理器。例如：

• 一个四核CPU中，整个CPU是一个设备，每个核心可以视为一个计算单元。
• GPU当然也是计算设备。

多个计算设备可以组合成设备组。一个设备组可以是任何计算设备的混合（如CPU和GPU），也可以只包含同类设备（如多块GPU）。这是OpenCL与CUDA等框架的一个关键区别，它真正将CPU与GPU视为对等伙伴。

所有设备组都存在于一个主机中。主机通常是我们的桌面或笔记本电脑系统。一个主机可以包含多个设备组。

内存对象

内存对象是程序操作的数据容器，主要分为两类：

以下是内存对象的主要类型：

1. 数组：与C语言中的数组概念类似，通过指针访问。在CPU上，数组的读写通常有缓存支持；在GPU上，通常没有（但可通过特定方式利用纹理缓存进行优化）。
2. 图像：分为2D和3D图像。图像数据以高度优化的非线性格式存储，专为图形硬件设计。不能像普通数组那样用指针直接访问。读取图像数据会使用纹理缓存，这也可以被用来加速非图像数据的访问。

可执行对象

可执行对象是我们要运行的代码部分。

以下是可执行对象的组成部分：

1. 计算内核：这是一个数据并行函数，将在计算设备（CPU或GPU）上执行。它看起来像一个C函数，但带有一些特殊修饰符（如__global）。关键概念：一个内核函数一次只处理一个数据元素（一个工作项）。例如，一个对两个数组求和的核函数，每次调用只计算一个索引位置的和。

   __kernel void sum(__global const float* a,
                     __global const float* b,
                     __global float* answer) {
       int i = get_global_id(0); // 获取当前工作项的全局索引
       answer[i] = a[i] + b[i];  // 只处理一个元素
   }
   

2. 计算程序：这是一个计算内核和函数的集合。我们可以将多个相关的内核放在同一个程序源文件中。OpenCL使用即时编译（JIT），在运行时编译这些程序。在macOS中，这个编译器是LLVM，速度极快。程序也可以预编译。

---

工作项与工作组

现在我们已经了解了代码（内核）和数据（内存对象）。本节中，我们来看看OpenCL如何组织并行执行。

OpenCL将并行任务分解为工作项。每个工作项执行一次内核实例。在GPU上，一个工作项通常对应一个硬件线程。

工作项被分组为工作组。工作组的大小是我们可以控制的。这种划分对于优化至关重要，特别是在GPU上利用本地内存时。

整个问题空间的大小称为ND范围（N维范围）或全局大小。工作组的大小称为局部大小。在GPU上，局部大小必须能整除全局大小。在CPU上，局部大小通常为1。

工作维度可以是1维、2维或3维，这非常适合于科学计算中常见的多维数据问题。

每个工作项都知道自己在整个ND范围中的位置，可以通过内置函数查询：

• get_global_id(0)：获取当前工作项在第一维（X）的全局索引。
• get_local_id(0)：获取当前工作项在其所属工作组内的局部索引。

---

内核与内存地址空间

上一节我们介绍了工作组织方式。本节中，我们深入看看内核编程细节和至关重要的内存模型。

OpenCL内核基于C语言，并增加了以下特性：

• 2D/3D图像类型。
• 内置函数（如get_global_id）。
• 同步原语（如屏障）。
• 向量数据类型：这是非常有用的特性，允许将多个标量组合成一个单元进行处理（如float4）。有float2和c_float2等多种声明方式，后者可避免命名冲突。

GPU能够同时管理数千个线程，这得益于其硬件线程管理器和极快的上下文切换能力。启动大量线程的一个主要目的是隐藏内存访问延迟：当一个线程等待数据时，硬件可以立即切换到另一个就绪的线程执行。

OpenCL定义了四种内存地址空间，对于GPU编程尤为重要：

以下是四种内存地址空间：

1. 全局内存：GPU上最大的内存池（如512MB）。使用__global限定符。所有工作项都可读写，但速度相对较慢。
2. 常量内存：用于存储不会改变的数据（如常数、查询表）。访问速度比全局内存快，通常有缓存支持。
3. 本地内存：与一个计算单元（一个工作组在其上执行）关联的高速内存。由工作组内所有工作项共享，速度比全局内存快一个数量级以上（如可达每秒数百GB），但容量很小（通常为16KB或32KB）。用于工作项间通信和高速数据复用。
4. 私有内存：每个工作项私有的内存。用于函数内的局部变量。一个工作项无法访问另一个工作项的私有内存。

理解这些内存空间及其性能特征是优化OpenCL内核的关键。

---

OpenCL程序执行流程

最后，我们来看一个完整的OpenCL程序是如何从初始化到执行再到清理的。遵循以下五个主要步骤：

以下是执行OpenCL计算的五个主要步骤：

1. 初始化：选择计算设备（如GPU），创建与该设备关联的上下文，然后创建命令队列。命令队列用于向设备发送操作指令（如数据传输、内核执行）。
2. 资源分配：在设备上分配内存缓冲区（使用clCreateBuffer），并将主机上的数据复制到设备内存（使用clEnqueueWriteBuffer命令放入队列）。
3. 程序与内核创建：将内核源代码（作为字符串或从文件读取）创建为程序对象，然后构建（编译）该程序。接着，从已构建的程序中创建特定的内核对象。
4. 执行：为内核设置参数（指向之前分配的内存缓冲区），通过clEnqueueNDRangeKernel命令将内核执行任务放入命令队列。在此命令中，需要指定问题的全局大小（ND范围）和局部大小（工作组大小）。通常使用clFinish来等待所有命令执行完毕，然后使用clEnqueueReadBuffer将结果读回主机。
5. 清理：释放所有创建的对象（上下文、命令队列、程序、内核、内存缓冲区），做良好的内存管理公民。通常，步骤1-3（初始化和创建）只需执行一次，而步骤4（执行）可以重复多次。

关于工作组大小的提示：在GPU上，选择合适的工作组大小对性能至关重要。通常，2的幂次方（如64，128，256）是好的选择，并且需要匹配硬件的特性。

总结

本节课中我们一起学习了OpenCL的基础知识。我们从CPU与GPU的硬件差异出发，理解了OpenCL高性能的根源。然后，我们系统地学习了OpenCL的三大核心对象：计算设备、内存对象和可执行对象（内核与程序）。我们探讨了OpenCL如何通过工作项和工作组来组织并行任务，并深入了解了四种关键的内存地址空间及其对性能的影响。最后，我们梳理了一个典型OpenCL程序从初始化、资源分配、编译、执行到清理的完整流程。掌握这些概念，就为编写和优化OpenCL并行程序打下了坚实的基础。

006：构建OpenCL项目 🚀

概述

在本节课中，我们将学习如何构建一个OpenCL项目。我们将更详细地探讨一些关键函数和功能，了解如何设置计算，以及如何使用像Xcode这样的工具来构建项目、运行计算和创建在底层使用OpenCL的应用程序。我们首先会回答一些常见问题，然后通过一个简单的示例项目来演示整个过程。

---

常见问题解答

上一节我们介绍了OpenCL的基础概念，本节中我们来看看一些开发者常遇到的问题。

双精度运算

双精度运算对许多科学计算至关重要。在OpenCL 1.0规范中，双精度是可选扩展。这意味着硬件和实现都必须支持它。你可以通过查询设备信息来检查是否支持双精度。如果支持，你需要在执行任何双精度计算语句之前使用一个特定的#pragma指令。如果不这样做，规范定义其行为是未定义的，可能导致计算崩溃或编译错误。

核心概念：使用clGetDeviceInfo查询CL_DEVICE_EXTENSIONS来检查是否支持cl_khr_fp64扩展。

需要注意的是，即使支持，双精度运算也可能带来显著的性能损失。例如，在GTX 285显卡上，单精度浮点运算可达约1000 GFLOP/s，而双精度可能只有约90 GFLOP/s。如果你的硬件不支持双精度，可以考虑使用混合精度算法来模拟。

面向对象编程

OpenCL基于C语言，本身不支持面向对象的概念。这意味着你不能将复杂的对象（如C++或Objective-C对象）直接传递到OpenCL内核中。

核心概念：内核参数必须是C语言的内置类型或OpenCL支持的扩展类型（如标量、向量、结构体）。结构体可以包含这些类型。

你可以从面向对象的程序中调用OpenCL例程，但需要先将对象数据转换为OpenCL能理解的原始类型数组（如float*），传递给内核处理，然后再将结果转换回对象。对于计算密集型任务，这种数据格式转换的开销通常可以忽略不计。

工作组大小与维度

工作组大小是性能调优的关键。全局工作组大小是你的问题规模（如一个包含16个元素的数组）。本地工作组大小必须是全局大小的整数因子，并且在CPU上必须为1（因为CPU上线程间同步开销极大）。

在GPU上，确定最佳本地工作组大小通常需要实验。它不应小于硬件的基本处理单元大小（NVIDIA的Warp是32个线程，AMD的Wavefront大小类似）。通常，2的幂次方或其组合是较好的选择。有时，为了对齐到2的幂次方，可能需要在数据末尾填充一些无操作的元素。

OpenCL支持将问题划分为一维、二维或三维（NDRange）。这主要是为了方便思考和索引，例如将图像处理视为二维问题，或将网格计算视为三维问题。目前没有证据表明不同维度划分会带来性能差异，选择哪种取决于哪种方式更符合你对问题的建模。

核心概念：

• size_t global_work_size = 16; // 总工作项数
• size_t local_work_size = 2; // 每个工作组的工作项数，必须能整除global_work_size
• 在CPU上：local_work_size = 1;
• 在GPU上：local_work_size 通常为 32, 64, 128, 256 等值进行试验。

适用的问题类型

许多典型的科学计算问题都适合用OpenCL/GPU处理，并且可以高效实现，例如：

• 快速傅里叶变换
• 基础线性代数子程序
• LAPACK
• 蒙特卡洛模拟
• 偏微分方程

然而，并非所有算法都能在GPU上达到最优。关键通常在于数据布局。GPU对数据的存取模式有特定偏好，如果能将数据组织成合适的格式，性能会非常出色。有时，这可能意味着需要重构算法或创建临时数据结构。

另一个重要概念是：复杂的计算不需要在单个内核调用中完成。你可以将其分解为多个内核或多个队列调用。例如，共轭梯度算法包含多个步骤（如SAXPY操作、矩阵向量乘法），每个步骤可以是一个独立的内核，按顺序在队列中执行。

需要注意的是，一旦任务进入命令队列，就无法中途终止。因此，对于有提前退出条件（如收敛检查）的迭代算法，一种策略是：在CPU端设置固定迭代次数（如200次），执行完成后，将单个标量结果（如残差）读回CPU检查，再决定是否继续迭代。

---

关键OpenCL函数详解

在进入实际项目之前，让我们更深入地了解一些构建OpenCL程序时会用到的核心函数。

设备发现与信息查询

要使用OpenCL，首先需要发现可用的计算设备。

• clGetDeviceIDs: 此函数用于获取设备列表。其第二个参数device_type至关重要，你可以指定寻找CL_DEVICE_TYPE_CPU、CL_DEVICE_TYPE_GPU、CL_DEVICE_TYPE_ACCELERATOR或CL_DEVICE_TYPE_ALL等。
• clGetDeviceInfo: 获取设备的具体信息，例如供应商名称(CL_DEVICE_VENDOR)、全局内存大小(CL_DEVICE_GLOBAL_MEM_SIZE)、最大工作组大小(CL_DEVICE_MAX_WORK_GROUP_SIZE)以及支持的扩展列表(CL_DEVICE_EXTENSIONS)。这是你检查双精度支持等功能的途径。

程序构建与编译

OpenCL内核在运行时编译（即时编译）。

• clBuildProgram: 编译链接着色器源代码创建程序对象。如果编译失败，你需要检查构建日志。
• clGetProgramBuildInfo: 在clBuildProgram之后调用，特别是当构建失败时，用于获取构建日志(CL_PROGRAM_BUILD_LOG)，其中包含了编译器错误和警告信息，对于调试内核代码语法错误非常有用。

内存管理

数据需要在主机（CPU）内存和设备（如GPU）内存之间移动。

• clCreateBuffer: 在设备上分配内存缓冲区。创建时可以指定内存标志，例如：
  • CL_MEM_READ_ONLY: 内核只能读取此缓冲区。
  • CL_MEM_WRITE_ONLY: 内核只能写入此缓冲区。
  • CL_MEM_READ_WRITE: 内核可读写。
  • 使用CL_MEM_USE_HOST_PTR或CL_MEM_ALLOC_HOST_PTR可以在特定情况下优化主机与设备间的数据传输，但在GPU上频繁通过PCIe总线访问主机内存会非常慢，应避免。
• clEnqueueWriteBuffer: 将数据从主机内存写入设备缓冲区。关键参数blocking_write：
  • 设为CL_TRUE（阻塞）：函数会等待数据复制完成才返回。确保数据在计算开始前已就位。
  • 设为CL_FALSE（非阻塞）：函数立即返回，复制操作在后台进行。如果内核紧接着启动，可能会读取到不完整或旧数据。
• clEnqueueReadBuffer: 将计算结果从设备缓冲区读回主机内存。同样有blocking_read参数，通常使用阻塞读取以确保在后续处理前已获得完整结果。

内核执行

设置并启动内核进行计算。

• clSetKernelArg: 为内核函数设置参数。你需要传递缓冲区对象（cl_mem）或标量值。
• clEnqueueNDRangeKernel: 将内核执行命令放入命令队列。你需要指定全局工作大小（global_work_size）和可选的本地工作大小（local_work_size）。如果本地工作大小为NULL，OpenCL实现会尝试选择一个值。
• clFinish: 阻塞主机程序，直到命令队列中的所有命令都执行完毕。在读取结果之前调用clFinish可以确保所有计算已经完成。

---

实战：一个简单的OpenCL Xcode项目

现在，让我们通过一个实际的Xcode项目来整合上述概念。这是一个极简的示例，目的是展示OpenCL项目的结构和基本流程。

项目概述

本项目实现了一个简单的向量加法（A + B = C）。它首先尝试在GPU上运行，如果找不到GPU，则回退到CPU。代码有详细注释，结构清晰。

以下是项目的主要步骤：

1. 定义问题与主机数据：在main函数中，定义问题大小（例如32），并在主机上分配并初始化输入数组A和B。
2. 调用OpenCL运行函数：将数据和控制权传递给run_cl函数。
3. 设备发现与选择：
   • 尝试获取GPU设备。
   • 如果失败，则回退到CPU设备。
   • 使用clGetDeviceInfo打印设备信息。
4. 创建上下文和命令队列：为选定的设备创建上下文和命令队列。
5. 创建与构建程序：
   • 从磁盘文件（.cl后缀）读取内核源代码。
   • 调用clBuildProgram编译程序。
   • （可选）此处可添加错误检查，使用clGetProgramBuildInfo获取编译日志。
6. 创建内核对象：从已构建的程序中，通过内核函数名（"add"）创建内核对象。
7. 分配设备内存：
   • 为输入数组A和B创建CL_MEM_READ_ONLY缓冲区。
   • 为输出数组C创建CL_MEM_WRITE_ONLY缓冲区。
   • 使用clEnqueueWriteBuffer（阻塞方式）将A和B的数据传输到设备。
   • 调用clFinish确保数据传输完成。
8. 设置内核参数：使用clSetKernelArg将三个设备缓冲区设置为内核add的参数。
9. 执行内核：
   • 设置全局工作大小为问题大小（32）。
   • 本地工作大小设为NULL，让OpenCL自行决定。
   • 调用clEnqueueNDRangeKernel将内核放入队列。
   • 调用clFinish等待内核执行完毕。
10. 读取结果：使用clEnqueueReadBuffer（阻塞方式）将结果从设备缓冲区C读回主机内存。
11. 清理资源：释放OpenCL对象（缓冲区、内核、程序、队列、上下文）。
12. 验证结果：控制权返回main函数，打印结果数组C。

内核代码 (add.cl)

内核文件非常简单，定义了加法操作。

__kernel void add(__global const float* a,
                  __global const float* b,
                  __global float* c)
{
    // 获取当前工作项的全局ID
    int gid = get_global_id(0);
    // 执行加法
    c[gid] = a[gid] + b[gid];
}

关于内核代码管理的讨论

在项目中，内核代码通常以两种方式管理：

1. 外部文件（如本例）：优点是与主机代码分离，更清晰，易于编辑和调试。缺点是如果希望保护内核代码知识产权，需要额外处理（如预编译为二进制）。
2. 内嵌为C字符串：将内核源代码作为字符串常量写在主机代码中。这样做可以一定程度上混淆代码，但会使主机代码变得冗长且难以维护内核逻辑。

关于编译时机：

• 即时编译（JIT）：本例采用的方式。优点是可以针对运行时发现的特定硬件进行优化，可能获得最佳性能。编译开销通常在程序初始化阶段，对于长时间运行或计算密集型的应用来说可以接受。
• 预编译：OpenCL支持将内核预编译为二进制格式并直接加载。这可以保护知识产权并减少运行时编译开销。但缺点是失去了针对最终用户硬件进行特定优化的机会。

运行项目

在Xcode中构建并运行该项目，你将在控制台看到类似以下输出：

Device: NVIDIA GeForce GTX 285 by NVIDIA Corporation
Results: 2.0, 4.0, 6.0, ... , 64.0

这表示程序成功在GPU上执行了向量加法。

---

总结与资源

本节课我们一起学习了如何构建一个完整的OpenCL项目。我们从解答常见问题开始，涵盖了双精度支持、面向对象编程的局限性、工作组大小调优以及OpenCL适用的科学计算问题类型。接着，我们深入探讨了设备查询、程序构建、内存管理和内核执行等关键函数。最后，我们通过一个简单的Xcode示例项目，一步步演示了从主机代码编写、内核开发到项目运行的全过程。

延伸阅读：

1. 稀疏矩阵向量乘法：NVIDIA的Nathan Bell等人发表了一篇关于在GPU上实现稀疏矩阵向量乘法的优秀论文，详细讨论了存储格式和优化技巧，非常值得一读。
2. 混合精度算法：有一份在线演示文稿详细介绍了如何使用混合精度算法在支持有限精度的设备上模拟更高精度的计算，同时避免巨大的性能损失。

希望本教程能帮助你入门OpenCL项目开发。如果你有任何问题或评论，欢迎通过Macresearch.org网站或电子邮件提出。下次课程中，我们计划探讨数据布局、内存访问模式以及Warp/Wavefront等更深入的主题。

007：内存布局与访问 🧠

在本节课中，我们将学习OpenCL编程中至关重要的一个主题：内存布局与访问。理解如何高效地组织数据并访问内存，是充分发挥GPU性能的关键。我们将从GPU架构的视角出发，探讨数据对齐、合并访问、共享内存以及如何避免性能瓶颈。

---

概述

本节课我们将深入探讨GPU的内存架构，特别是NVIDIA硬件上的内存组织方式。我们将学习线程、线程块、warp等核心概念，并理解数据在全局内存和共享内存中的布局。通过一个矩阵转置的实例，我们将看到如何应用这些知识来优化内存访问，从而显著提升程序性能。

---

GPU架构简介

上一节我们介绍了OpenCL的基本执行模型。本节中，我们来看看GPU的物理架构，特别是NVIDIA GPU的组成，这有助于我们理解内存访问模式背后的原因。

GPU的计算单元是分层组织的。在NVIDIA的术语中（对应OpenCL的工作项和工作组）：

• 线程 对应一个 工作项，执行一个内核实例。
• 线程块 对应一个 工作组，是线程的集合。

以GTX 285显卡为例，其核心架构如下：

• 线程处理集群：显卡上最粗粒度的计算单元，共有10个。
• 流多处理器：每个TPC包含多个SM。GTX 285有30个SM。
• 流处理器：每个SM包含8个流处理器（也称为“核心”）。整个显卡共有240个这样的核心。
• 特殊功能单元：每个SM有2个SFU，用于处理超越函数（如sin, cos, sqrt）。
• 双精度单元：每个SM有1个，用于双精度浮点计算。
• 共享内存：每个SM有16KB的超高速内存，供该SM内的所有流处理器共享，用于线程块内线程间的数据交换。

---

线程执行与Warp

理解了硬件组成后，我们来看看线程是如何在SM上调度和执行的。这直接关系到我们如何编写高效的内核。

每个SM可以并发执行多个线程块。线程块在硬件上被进一步细分为更小的调度单元，称为 Warp。在NVIDIA硬件上：

• 一个 Warp 包含 32 个线程。
• 一个Warp还可以分为两个 Half-Warp，每个包含16个线程。16这个数字在内存访问中尤为重要。

Warp中的32个线程以 锁步 方式执行相同的指令。这意味着如果代码中存在条件分支（如if-else），导致Warp内部分线程执行不同路径，那么所有路径将被串行化执行，从而造成性能损失。这种现象称为 分支发散，应尽量避免。

---

内存层次与访问模式

上一节我们了解了线程的执行方式。本节中我们来看看数据在GPU内存层次结构中的流动，特别是如何高效地从全局内存加载数据。

GPU拥有不同层次的内存，其速度和容量各不相同：

1. 全局内存：容量大，但延迟高（慢）。
2. 共享内存：位于SM上，容量小（如16KB），但速度极快，堪比寄存器。
3. 寄存器：速度最快，但数量有限。

为了隐藏全局内存的高延迟，GPU采用 大规模多线程 策略。当一个Warp因等待数据而暂停时，硬件会迅速切换到另一个就绪的Warp执行。因此，拥有足够多“在途”的线程有助于保持计算单元忙碌。

从全局内存加载数据时，最关键的概念是 合并访问。

---

合并访问与对齐

以下是访问全局内存时需要注意的几种情况：

• 未对齐访问：线程0的起始地址不是硬件加载大小（如64字节）的整数倍。这会导致低效的多次内存事务。
• 交叉访问：线程访问的内存地址不连续。硬件无法将其识别为一次大块加载，会导致多次串行加载。
• 部分线程访问：并非所有线程都参与加载。同样无法形成合并访问。
• 理想情况：合并且对齐的访问：一个Warp（或Half-Warp）的所有线程访问一段连续的、对齐的内存地址。硬件可以将其合并为一次大的内存事务（如一次加载64字节），这是最高效的方式。

公式：对于Half-Warp（16个线程）和float类型（4字节），一次理想的合并加载大小是：
16 threads * 4 bytes/thread = 64 bytes

---

共享内存与存储体冲突

我们已经知道如何高效地将数据读入芯片。接下来，我们看看如何利用共享内存来优化数据重用和访问模式。

共享内存被组织成多个 存储体。在当前的硬件上，通常有16个存储体。多个线程可以同时访问不同的存储体，从而实现极高的内存带宽。

然而，如果多个线程试图访问同一个存储体中的不同数据地址，就会发生 存储体冲突，导致对这些地址的访问被串行化，降低性能。

以下是共享内存访问模式的例子：

• 无冲突访问：每个线程访问不同存储体中的连续地址（步长为1）。这是最佳情况。
• 广播：所有线程读取同一个地址的数据。这是一种特殊情况，硬件会进行广播，不会导致冲突。
• 存储体冲突：多个线程访问同一个存储体中的不同地址（例如，步长为2的访问可能导致2路冲突）。这会导致串行化，应尽量避免。

核心思想：将数据从全局内存合并加载到共享内存后，线程可以在共享内存中自由地、快速地以任意模式访问所需数据，而无需担心全局内存访问的延迟和合并问题。

---

实战案例：矩阵转置

现在，让我们将前面学到的所有概念应用到一个实际例子中：矩阵转置。这个例子完美展示了共享内存和合并访问的价值。

矩阵转置操作（B[y][x] = A[x][y]）在直接使用全局内存时面临一个问题：要么读取是合并的但写入是交叉的，要么反之。这会导致低效的内存访问。

优化策略如下：

1. 合并读取：让一个线程块的所有线程从全局内存中合并读取一个数据块（例如，一块16x16的矩阵数据）到共享内存。每个线程负责读取一个元素。
2. 共享内存中转：数据现在位于高速的共享内存中。线程在共享内存中交换数据（进行转置操作）。由于共享内存访问速度快，且我们通过填充等手段可以避免存储体冲突，这一步开销很小。
3. 合并写入：线程块的所有线程将转置后的数据从共享内存合并写入到全局内存的输出矩阵中。

通过这个“全局内存 -> 共享内存 -> 全局内存”的流程，我们确保了在全局内存层面的读取和写入都是高效的合并访问，从而大幅提升性能。

代码概念：

// 伪代码示意内核函数中的关键步骤
__kernel void transpose(__global float* input, __global float* output, __local float* block) {
    int local_x = get_local_id(0);
    int local_y = get_local_id(1);
    int global_in_x = get_group_id(0) * BLOCK_SIZE + local_x;
    int global_in_y = get_group_id(1) * BLOCK_SIZE + local_y;

    // 1. 合并读取：从输入矩阵读取到共享内存块
    block[local_y * (BLOCK_SIZE+1) + local_x] = input[global_in_y * WIDTH + global_in_x];
    barrier(CLK_LOCAL_MEM_FENCE); // 确保块内所有线程完成读取

    // 2. 在共享内存中交换索引（转置）
    int swapped_local_x = local_y;
    int swapped_local_y = local_x;

    // 3. 计算输出全局坐标
    int global_out_x = get_group_id(1) * BLOCK_SIZE + swapped_local_x;
    int global_out_y = get_group_id(0) * BLOCK_SIZE + swapped_local_y;

    // 4. 合并写入：从共享内存写入到输出矩阵
    output[global_out_y * WIDTH + global_out_x] = block[swapped_local_y * (BLOCK_SIZE+1) + swapped_local_x];
}

注意：共享内存数组block的宽度被填充了1（BLOCK_SIZE+1），这是为了避免在从共享内存读取转置数据时发生存储体冲突。

---

总结

本节课我们一起深入学习了OpenCL中内存布局与访问的核心知识。我们首先了解了NVIDIA GPU的基本架构，包括SM、Warp等概念。然后，我们重点探讨了如何实现高效的合并内存访问以利用全局内存带宽，以及如何利用共享内存来优化数据重用和访问模式。最后，通过矩阵转置的案例，我们看到了如何将这些理论应用于实践，通过“全局->共享->全局”的流程，将低效的交叉访问转化为高效的合并访问。

记住这些关键点：对齐、合并、利用共享内存、避免分支发散和存储体冲突。掌握它们是编写高性能OpenCL内核的基础。在接下来的课程中，我们将通过更多实际的内核优化示例来巩固这些概念。

008：问题与解答 🎥

在本节课中，我们将专注于解答观众提出的两个核心问题：GPU硬件架构的术语解释，以及共享内存中“存储体冲突”的详细原理。我们将通过清晰的图示和简单的例子，帮助你理解这些关键概念。

---

GPU硬件架构与术语解析 🧩

上一节我们介绍了GPU编程的基本概念，本节中我们来看看GPU硬件的具体组织方式，并澄清一些容易混淆的术语。

下图展示了一个典型的GPU（以NVIDIA 10系列架构为例）内部结构：

以下是其层级结构的分解：

• GPU：整个黑色区域代表一个GPU芯片。
• 线程处理集群：每个深灰色方块代表一个线程处理集群。在10系列架构中，一个GPU包含10个TPC。
• 流多处理器：每个浅灰色方块代表一个流多处理器。每个TPC包含3个SM，因此一个GPU总共有 30个SM。
• 核心/流处理器：每个橙色小方块代表一个核心（也称为流处理器或标量处理器）。每个SM包含8个这样的核心，因此一个GPU总共有 240个核心。

重要提示：这里的“核心”概念与CPU的“核心”不同。GPU的“核心”主要指执行单精度浮点运算和整数运算的算术逻辑单元。而一个CPU核心则包含ALU、FPU、缓存、内存控制器等多种功能单元。这种术语差异有时是出于市场宣传，但理解其本质区别对编程至关重要。

此外，每个SM还包含其他专用硬件单元：

• 双精度单元：每个SM有1个，整个GPU共30个。
• 特殊功能单元：每个SM有2个，用于快速计算超越函数（如sin, cos）和倒数平方根等。整个GPU共60个。
• 本地/共享内存：用于SM内线程（或工作组）之间的数据共享。

---

深入理解存储体冲突 🔄

在上一节关于共享内存的讨论中，我们提到了“存储体冲突”会严重影响性能，并以矩阵转置为例说明了可能发生冲突的情况。本节我们将通过一个更详细的图示来彻底解释其原因。

首先，快速回顾共享内存的关键特性：

• 容量通常为16 KB。
• 分为 16个存储体，每个存储体1 KB。
• 每个存储体能提供32位（4字节）宽的数据访问。
• 关键规则：连续的32位字被分配到连续的存储体中。例如，地址0的字在存储体0，地址1的字在存储体1，……，地址16的字又回到存储体0，依此类推。
• 存储体冲突：当同一个半线程束（16个线程）中的两个或更多线程同时访问同一个存储体中的不同数据时，这些访问会被序列化（即一个一个执行），从而导致性能下降。唯一的例外是当所有线程访问同一个存储体中的完全相同地址的数据时，这属于“广播”操作，不会冲突。

现在，我们来看矩阵转置中导致冲突的具体场景。假设我们有一个包含32个元素（0-31）的数组，每个元素4字节。一个完整的线程束（32线程）处理所有数据，但硬件以半线程束（16线程）为单位调度。

第一步：将数据从全局内存写入共享内存
我们进行合并访问，每个线程写入共享内存中连续且对齐的位置。此时没有冲突。

• 元素 0 -> 存储体 0
• 元素 1 -> 存储体 1
• ...
• 元素 15 -> 存储体 15
• 元素 16 -> 存储体 0 （因为只有16个存储体，所以取模16）
• 元素 17 -> 存储体 1
• ...

第二步：从共享内存读取以进行转置
这才是问题所在。为了转置，线程的读取模式发生了变化：

• 线程0 读取 元素0 （存储体0）
• 线程1 读取 元素16 （存储体0）
• 线程2 读取 元素32 （存储体0）
• ...

你会发现，同一个半线程束内的多个线程（线程0、1、2...）试图同时访问不同但都位于存储体0的数据。根据上述规则，这构成了存储体冲突，硬件必须将这些访问序列化，严重拖慢速度。

解决方案：填充共享内存数组
解决方法是给共享内存数组增加填充。通常只需填充一个元素。

修改后的写入步骤：

• 我们声明共享内存数组的大小为 17个元素（而非16个），多出的一个位置作为填充，不使用。
• 写入数据时，我们仍然从全局内存合并读取，但写入共享内存时故意错位：
  • 全局元素0 -> 共享位置0 (存储体0)
  • 全局元素1 -> 共享位置1 (存储体1)
  • ...
  • 全局元素15 -> 共享位置15 (存储体15)
  • 共享位置16 留空（填充）
  • 全局元素16 -> 共享位置17 (存储体1，因为17 mod 16 = 1)
  • 全局元素17 -> 共享位置18 (存储体2)
  • ...

修改后的读取步骤：

• 线程0 读取共享位置0 （存储体0）
• 线程1 读取共享位置17 （存储体1）
• 线程2 读取共享位置34 （存储体2）
• ...

通过填充，原本会冲突的访问被分散到了不同的存储体，从而避免了存储体冲突。这种技术在处理需要非连续或跨步访问共享内存的算法（如矩阵转置、卷积等）时非常常见。

---

总结与资源 📚

本节课中我们一起学习了两个重点：

1. GPU架构术语：理解了GPU由TPC、SM和核心构成的多层级结构，并明确了GPU“核心”与CPU“核心”的功能差异。
2. 存储体冲突：通过矩阵转置的详细例子，剖析了共享内存中存储体冲突产生的根本原因（多个线程同时访问同一存储体的不同数据），并掌握了通过数组填充这一有效方法来避免冲突。

希望这些解释能让之前感到困惑的概念变得清晰。要深入掌握OpenCL，持续的实践和查阅资料至关重要。

延伸学习资源：

• 本系列所有视频和资料均可在 Macresearch.org/opencl 找到。
• 网站还提供其他优秀教程，如Jerry McCormack的“Cocoa for Scientists”系列。
• NVIDIA官方定期举办在线研讨会（Webinars），讲解CUDA、OpenCL等GPU编程技术，会后会提供视频，是很好的学习资源。

下一节课，我们将兑现承诺，通过一个实际的应用示例，把到目前为止学到的所有知识——内核编写、优化、利用本地/共享内存及填充技巧——融合贯通。敬请期待！

---

注：部分提及的推广内容（如主机服务商、亚马逊商店）已按教程要求省略，仅保留核心学习资源链接。

009：共享内存内核优化

在本节课中，我们将学习如何利用OpenCL的共享内存来优化内核性能。我们将通过一个源自真实科学计算程序（APBS）的代码示例，深入探讨共享内存的概念、如何通过协作加载数据来减少全局内存访问延迟，以及如何使用同步屏障来协调工作组内的工作项。通过对比优化前后的性能，你将直观地看到共享内存带来的巨大性能提升。

---

概述与背景

在之前的课程中，我们介绍了共享内存的基本概念和用途。本节课程将把这些概念付诸实践，通过一个具体的计算示例来展示如何利用共享内存进行内核优化。

这个示例计算生物分子中每个网格点的静电势能，其核心是对所有原子进行求和计算。在串行计算中，我们可以采用原子中心或网格中心的方法。但在并行环境中，特别是GPU上，网格中心的方法是更优的选择，因为它避免了数据竞争，也无需使用锁或归约操作。

核心计算与初始实现

计算的核心是为每个网格点累加所有原子的贡献值。公式上，这类似于库仑定律的扩展：

网格点值 += 函数(原子坐标, 原子电荷, 原子半径, 网格点坐标)

在CPU上，一个简单的串行实现是双层循环：

for (每个网格点 i) {
    for (每个原子 j) {
        // 计算原子j对网格点i的贡献并累加
    }
}

当我们将此计算移植到GPU时，最直接的思路是将外层循环（网格点迭代）映射为全局工作项（NDRange）。每个工作项（对应一个网格点）独立地循环遍历所有原子。初始的、未优化的内核代码如下所示：

__kernel void unoptimized_kernel(__global float* gridData,
                                 __global float* atomData,
                                 int numAtoms) {
    int gid = get_global_id(0); // 当前网格点ID
    float gridValue = 0.0f;
    // 假设atomData中按顺序存储了所有原子的x, y, z, charge, radius
    for (int atom = 0; atom < numAtoms; ++atom) {
        float dx = atomData[atom*5 + 0] - gridX[gid];
        float dy = atomData[atom*5 + 1] - gridY[gid];
        float dz = atomData[atom*5 + 2] - gridZ[gid];
        float charge = atomData[atom*5 + 3];
        float radius = atomData[atom*5 + 4];
        // ... 进行计算并累加到 gridValue ...
    }
    gridData[gid] = gridValue;
}

这种实现存在明显的性能问题：每个工作项都需要反复从全局内存中读取原子数据（坐标、电荷等）。尽管硬件可能会检测到相邻工作项在读取相同地址并进行合并访问，但在原子循环中，这种加载操作仍然是序列化的，会成为性能瓶颈。

共享内存优化策略

为了解决上述性能问题，我们引入共享内存。基本思路是：让一个工作组（Work-Group）内的所有工作项协作，将一大块原子数据从全局内存一次性加载到快速的共享内存中，然后所有工作项再从共享内存中读取数据进行计算。

我们首先在kernel中声明共享内存：

__local float sharedAtomData[5 * LOCAL_SIZE];

这里，LOCAL_SIZE是工作组的大小（例如64）。我们分配了5倍于工作组大小的浮点数空间，用于连续存储原子的X、Y、Z坐标、电荷和半径数据。

上一节我们介绍了优化策略，本节中我们来看看具体的实现步骤。以下是优化后内核中内层循环的主体结构：

1. 协作加载数据：工作组以LOCAL_SIZE为步长，分批处理原子。在每一批中，每个工作项负责将特定原子的数据从全局内存拷贝到共享内存的指定位置。
2. 同步屏障：确保工作组内所有工作项都完成数据拷贝后，才能进行下一步计算。
3. 共享内存计算：所有工作项从共享内存中读取当前批次的原子数据，并行完成各自网格点的部分累加计算。
4. 再次同步：确保所有工作项都使用完当前批次的共享内存数据后，才能加载下一批数据，防止数据被覆盖。

对应的内核代码片段如下：

for (int atomBase = 0; atomBase < numAtoms; atomBase += LOCAL_SIZE) {
    int loadLimit = min(LOCAL_SIZE, numAtoms - atomBase);
    int localIdx = get_local_id(0);

    // 步骤1: 协作加载数据到共享内存
    if (localIdx < loadLimit) {
        int globalAtomIdx = atomBase + localIdx;
        sharedAtomData[localIdx + 0*LOCAL_SIZE] = atomData[globalAtomIdx*5 + 0]; // X
        sharedAtomData[localIdx + 1*LOCAL_SIZE] = atomData[globalAtomIdx*5 + 1]; // Y
        sharedAtomData[localIdx + 2*LOCAL_SIZE] = atomData[globalAtomIdx*5 + 2]; // Z
        sharedAtomData[localIdx + 3*LOCAL_SIZE] = atomData[globalAtomIdx*5 + 3]; // Charge
        sharedAtomData[localIdx + 4*LOCAL_SIZE] = atomData[globalAtomIdx*5 + 4]; // Radius
    }

    // 步骤2: 等待所有工作项完成加载
    barrier(CLK_LOCAL_MEM_FENCE);

    // 步骤3: 从共享内存进行计算
    for (int i = 0; i < loadLimit; ++i) {
        float dx = sharedAtomData[i + 0*LOCAL_SIZE] - gridX[gid];
        float dy = sharedAtomData[i + 1*LOCAL_SIZE] - gridY[gid];
        float dz = sharedAtomData[i + 2*LOCAL_SIZE] - gridZ[gid];
        float charge = sharedAtomData[i + 3*LOCAL_SIZE];
        float radius = sharedAtomData[i + 4*LOCAL_SIZE];
        // ... 进行计算并累加到 gridValue ...
    }

    // 步骤4: 等待所有工作项完成计算，再开始下一批次
    barrier(CLK_LOCAL_MEM_FENCE);
}

同步屏障的重要性

同步屏障（barrier）在此优化中至关重要。第一个屏障确保了在计算开始前，所有需要的原子数据都已安全地驻留在共享内存中。如果没有这个屏障，执行得快的工作项可能会读到尚未被其他工作项加载的无效数据。

第二个屏障确保了在覆盖共享内存以加载下一批原子数据之前，所有工作项都已经完成了对当前批次数据的计算。省略这个屏障可能导致数据竞争和计算结果错误。

性能对比与总结

通过实际运行示例代码，我们可以得到清晰的性能对比：

• CPU单线程：约 25-32 秒
• CPU多线程（16核，OpenMP）：约 2.5 秒 （比单线程快约10倍）
• GPU未优化内核：约 1.2 秒 （已比单线程CPU快约20倍）
• GPU共享内存优化内核：约 0.125 秒 （比未优化GPU内核快约10倍，比单线程CPU快约200倍）

本节课中我们一起学习了如何利用OpenCL共享内存进行内核优化。关键点包括：

1. 识别出内核中频繁访问的只读数据（本例中的原子数据）。
2. 使用共享内存作为高速缓存，由工作组协作加载数据块。
3. 正确使用同步屏障来协调工作组内工作项的执行顺序，保证数据一致性。
4. 通过将全局内存访问转换为共享内存访问，显著降低了内存延迟，从而极大提升了内核性能。

这个示例充分表明，理解硬件特性（如内存层次结构）并据此设计算法，是释放GPU强大并行计算能力的关键。你可以下载并修改附带的代码，尝试调整工作组大小或注释掉屏障语句，以更深入地观察其影响。

---

注意：如果您的系统只有一块显卡，运行长时间计算的GPU内核可能导致显示界面暂时无响应，因为图形命令队列可能被计算任务阻塞。建议在测试时调小问题规模，或使用专用于计算的第二块GPU。

010：并行与异构计算入门 🚀

概述

在本节课中，我们将要学习并行与异构计算的基本概念。我们将探讨为什么传统的单核处理器性能提升遇到了瓶颈，以及如何通过并行计算和异构系统（特别是结合CPU和GPU）来应对这一挑战。我们还将介绍相关的术语，并展望未来的编程模型发展趋势。

---

为什么需要并行计算？ ⚡

在2005年，Herb Sutter提出了著名的“免费午餐已经结束”的观点。这意味着硬件厂商不能再仅仅依赖提升时钟频率、优化指令级并行性或增加缓存大小来获得显著的性能提升。物理限制和功耗问题成为了主要瓶颈。

因此，性能提升的关键转向了并行性。通过让多个计算核心同时执行独立的任务，我们可以在不显著增加功耗的前提下，大幅提升计算能力。

核心概念：并行性是指计算中各个部分相互独立，因此可以同时执行，从而缩短整体计算时间。

---

关键术语定义 📖

在深入讨论之前，让我们明确几个核心术语的定义，以确保我们在同一基础上进行交流。

并行性 vs. 并发性

• 并行性：这是一种计算属性，指的是计算的各个部分相互独立，因此可以同时执行。其核心目标是提升性能。
  • 示例：计算 A = B + C 和 D = E * F 这两个赋值语句是独立的，可以并行执行。
• 并发性：这是一种逻辑上的编程抽象，允许构建可以相互通信的多个任务。它不强制要求任务同时执行。
  • 示例：在单核处理器上运行两个线程。处理器通过时间片轮转交替执行它们，从逻辑上看它们是并发运行的，但并非物理上的并行。

简单来说：并行关乎物理执行（同时发生），而并发关乎逻辑结构（可能同时发生）。

异构计算

异构计算是指一个系统由两个或更多在结构上存在显著差异的计算引擎组成。

• 典型例子：传统的CPU和GPU组合。
  • CPU：专为低延迟、复杂控制流和顺序任务优化，拥有大容量缓存。
  • GPU：专为高吞吐量、数据并行任务设计，采用大规模多线程架构。

融合架构

这是AMD提出的愿景，旨在将CPU和GPU（可能还有其他计算单元）集成到同一块硅芯片上。这种设计旨在实现高性能、低功耗，并简化编程模型。

核心优势：CPU和GPU可以共享内存，极大减少了数据拷贝和通信开销。

---

异构世界的硬件展望 🖥️

上一节我们介绍了基本概念，本节中我们来看看硬件是如何演进来支持异构计算的。

计算平台的演进

计算平台的发展大致经历了三个阶段：

1. 单核时代：依赖提升时钟频率和指令级并行性。最终受限于功耗和物理极限。
2. 多核时代：通过增加处理器核心数量来提升性能。但面临缓存一致性带来的可扩展性挑战和软件并行化的难题。
3. 异构时代（当前）：我们正处在这个时代的开端。通过结合不同架构的计算单元（如CPU和GPU），利用各自的优势。GPU因其在数据并行任务上的极高能效而成为关键。

为什么GPU如此重要？

以下是GPU的一些关键优势：

• 极高的计算吞吐量：现代GPU能提供每秒数万亿次浮点运算（TeraFLOPS）的性能。
• 高内存带宽：专用显存提供远超系统内存的带宽。
• 优异的能效比：在单位功耗和单位面积上，GPU能提供比CPU高得多的计算性能。

公式示例：性能提升 ≈ (GPU计算单元数量) × (并行任务数量)

为什么仍然需要CPU？

尽管GPU强大，但CPU不可或缺，原因如下：

• 处理串行和任务并行工作负载：擅长处理低延迟、分支密集型的标量代码。
• 庞大的软件生态：需要支持现有的操作系统（如Windows、Linux）和成千上万的应用程序。
• 控制与协调：负责系统管理、I/O操作以及为GPU准备和调度任务。

结论：未来的方向不是用GPU取代CPU，或将CPU变成巨型GPU，而是让两者紧密协同，各自处理最擅长的任务。

AMD融合架构示例

传统的PC系统中，CPU和GPU是独立的设备，通过PCIe总线连接，并拥有各自独立的内存。数据交换需要显式的拷贝操作。

在AMD的融合架构中，CPU和GPU被集成到同一块芯片（APU）上：

• 共享系统内存：CPU和GPU可以直接访问同一块内存，消除了大量数据拷贝。
• 更低延迟通信：芯片内部的通信延迟远低于通过PCIe总线。
• 更简化的编程模型：为未来实现更统一的内存访问模型（如GPU直接引用CPU指针）奠定了基础。

注意：初代融合APU的GPU性能可能不及高端独立显卡，但其在简化编程和降低特定场景延迟方面优势明显。未来会有面向不同市场（从移动设备到数据中心）的、性能更强的融合产品。

---

异构世界的软件挑战 💻

上一节我们探讨了硬件趋势，本节中我们来看看随之而来的软件编程挑战和可能的解决方案。

未来的性能提升完全取决于软件，即我们如何有效地为这些异构系统编程。

并行化方法

主要有两种分解问题以实现并行的方法：

1. 任务并行
任务并行侧重于将问题分解为多个不同的、可以独立或按依赖关系执行的任务。

以下是任务并行的关键特点：

• 任务间可能存在依赖。
• 需要任务间通信。
• 核心是负载均衡：动态地将任务分配给空闲的计算核心，以最大化资源利用率。

相关技术示例：Intel TBB, Apple GCD, OpenMP tasks, Microsoft TPL。

2. 数据并行
数据并行侧重于对大量独立的数据元素应用相同的操作。

以下是数据并行的关键特点：

• 核心是同时处理大量数据元素。
• 元素间通常独立，但也可支持局部通信（如粒子系统模拟中邻近粒子的交互）。
• 非常适合GPU加速。

相关技术示例：OpenCL, CUDA, OpenMP (SIMD), Microsoft Accelerator。

编织并行：未来的方向

大多数现实世界的应用程序（如现代游戏）同时包含任务并行和数据并行的元素。这种混合模式被称为“编织并行”。

示例：在一个游戏场景中：

• AI逻辑、用户输入处理：属于任务并行。
• 粒子系统、物理模拟：属于数据并行。

未来的编程模型需要自然地支持这种任务并行与数据并行的结合，并能智能地将不同部分调度到最适合的计算单元（CPU或GPU）上执行。

OpenCL的角色

OpenCL是一个低级别的编程模型，支持在CPU、GPU和其他加速器上执行数据并行任务。它是实现异构计算的重要工具之一。

OpenCL与CUDA：两者都是GPGPU编程语言，设计用于在GPU上高效运行数据并行任务，共享许多特性。OpenCL的优势在于其跨平台和跨厂商的开放性。

当前局限与未来：目前，OpenCL提供了源码可移植性，但难以实现性能可移植性（即同一段代码在CPU和GPU上都能高效运行）。未来的发展（如更高级的容器类型、编译器优化）将致力于改善这一点，并更好地与任务并行框架集成，以支持“编织并行”。

---

总结 🎯

本节课中我们一起学习了并行与异构计算的基础知识。

我们首先了解了为什么需要从串行计算转向并行计算。然后，我们明确了并行性、并发性和异构计算等关键术语。接着，我们探讨了硬件向异构时代演进的趋势，特别是AMD融合架构如何将CPU和GPU的优势结合。最后，我们审视了随之而来的软件挑战，介绍了任务并行和数据并行两种方法，并指出结合两者的“编织并行”模型与OpenCL等工具将是应对异构系统编程的关键。

异构计算的时代刚刚开始，如何为其编写高效、简洁的程序是留给整个行业的核心问题。

011：OpenCL简介 🚀

在本节课中，我们将要学习OpenCL的基础概念、核心架构以及一个简单的编程流程。OpenCL是一个用于编写在异构平台（如CPU、GPU）上运行程序的框架，它允许开发者利用多种设备的并行计算能力来加速应用程序。

---

什么是OpenCL？🤔

OpenCL是一个基于平台的编程模型。在这个模型中，有一个主机（通常是传统的CPU），它通过某种总线（如PCI Express或HyperTransport）连接到一个或多个计算设备。这些设备可以是另一个CPU、GPU或其他类型的加速器。

从OpenCL的视角看，一个设备被视为一组计算单元的集合。每个计算单元又进一步划分为多个处理元素，这些元素以单指令多数据的方式执行指令。

OpenCL的执行模型基于内核的概念。内核是可执行代码的基本单位，类似于C语言中的函数。它定义了一个可以从主机程序调用的入口点，用于在设备上执行计算任务。

---

OpenCL执行模型 ⚙️

OpenCL主要支持两种执行模型：数据并行和任务并行。

数据并行模型

数据并行模型是当前最高效的模型，尤其适合GPU类设备。在这个模型中，计算被定义在一个N维的计算域中。

• 全局工作项：计算域中的每个独立元素称为一个工作项。所有工作项可以并行执行。
• 工作组：工作项可以被分组为工作组。工作组在一个SIMD核心上执行。
• 关键特性：工作组内的工作项可以共享本地内存，并且可以进行同步。这使得工作组内的通信变得非常高效。然而，不同工作组之间的工作项不能直接进行快速通信。

数据并行示例：
假设我们要处理一个1024x1024的图像，这就是我们的全局问题维度。我们将为图像中的每个像素（即每个工作项）启动一个内核实例。

传统CPU上的标量乘法循环代码：

for (int i = 0; i < N; i++) {
    result[i] = inputA[i] * inputB[i];
}

在OpenCL内核中，循环被隐式化，通过get_global_id函数获取每个工作项的索引：

kernel void vec_mul(global float* A, global float* B, global float* result) {
    int id = get_global_id(0);
    result[id] = A[id] * B[id];
}

任务并行模型

在任务并行模型中，内核仅使用单个工作项执行。这对于在CPU设备上运行本机编译的代码、集成现有C/C++库或利用OpenCL的队列模型进行任务调度非常有用。

---

OpenCL内存模型 🧠

OpenCL明确地暴露了设备（尤其是GPU）的内存层次结构，这与传统的C语言内存模型不同。

以下是OpenCL的主要内存空间：

• 主机内存：由主机（CPU）管理的内存。
• 全局内存：所有工作项和工作组都可以访问的内存区域。访问延迟较高。
• 常量内存：全局可见的只读内存，通常放置在高速缓存中。
• 本地内存：在工作组内共享的内存。访问速度非常快，但需要程序员显式管理。
• 私有内存：每个工作项私有的内存。

关键点：在OpenCL中，所有内存管理都必须是显式的。这意味着数据在主机内存、全局内存和本地内存之间的移动需要通过明确的API调用或内核中的加载/存储操作来完成。

---

OpenCL框架与对象 🏗️

使用OpenCL编程主要涉及创建和管理一系列对象。

核心对象

以下是使用OpenCL框架时需要了解的核心对象：

• 平台：代表一个特定的OpenCL实现（如AMD或NVIDIA的实现）。
• 设备：系统中的计算单元（如CPU或GPU）。可以通过API查询设备的能力。
• 上下文：将一组设备和内存对象关联在一起的对象。它定义了这些内存对象之间的一致性模型。
• 命令队列：与特定设备和上下文关联，用于向设备提交命令（如内核执行、内存传输）。
• 内存对象：包括缓冲区（一维内存块）和图像（用于优化图像访问的特殊不透明类型）。
• 程序对象：封装了OpenCL C源代码或二进制代码，包含一个内核列表。
• 内核对象：代表程序中的一个内核函数，可以为其设置参数并排队执行。
• 事件对象：用于处理命令之间的依赖关系和同步，因为大多数clEnqueue*命令都是异步执行的。

基本工作流程

一个典型的OpenCL应用程序遵循以下步骤：

1. 查询并选择平台和设备。
2. 创建上下文和命令队列。
3. 创建内存对象（缓冲区/图像）。
4. 创建程序对象，编译OpenCL C源代码，并从中提取内核对象。
5. 为内核对象设置参数。
6. 将内核执行命令放入命令队列。
7. 排队进行内存读写操作，以传输数据。
8. 使用事件来同步命令的执行。

---

OpenCL C语言简介 📝

OpenCL C是基于C99的语言，用于编写在内核中执行的代码。它包含一些限制和许多针对并行计算的扩展。

主要特性

• 基于C99：但不支持递归等特性。
• 并行工作项函数：例如get_global_id、get_local_id、get_group_id，用于获取工作项在计算域中的位置信息。
• 地址空间限定符：如__global、__local、__constant、__private，用于指定变量的存储位置。
• 向量类型：支持长度为2、4、8、16的向量（如float4、int8），并提供了丰富的向量操作函数。
• 同步原语：如barrier，用于工作组内的同步。
• 大量内置函数：包括数学函数、几何函数、图像读写函数等。

内核示例

以下是一个简单的向量加法内核：

kernel void vec_add(__global const float* a,
                    __global const float* b,
                    __global float* result) {
    int gid = get_global_id(0);
    result[gid] = a[gid] + b[gid];
}

---

一个完整的简单示例 🔧

让我们将以上概念整合到一个完整的、简化的主机端代码流程中，实现向量加法。

步骤概述：

1. 获取平台和设备。
2. 创建上下文和命令队列。
3. 创建输入和输出的缓冲区。
4. 创建程序，编译内核源代码。
5. 创建内核对象并设置其参数。
6. 执行内核。
7. 将结果读回主机。

简化代码流程：

// 1. 获取平台和设备ID
cl_platform_id platform;
cl_device_id device;
clGetPlatformIDs(1, &platform, NULL);
clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/rayiyanutara-ocl-prog/img/6ae83480be5553072b13ec1e6a5e6ada_12.png)

// 2. 创建上下文和命令队列
cl_context context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL);
cl_command_queue queue = clCreateCommandQueue(context, device, 0, NULL);

// 3. 创建缓冲区
cl_mem bufA = clCreateBuffer(context, CL_MEM_READ_ONLY, size, NULL, NULL);
cl_mem bufB = clCreateBuffer(context, CL_MEM_READ_ONLY, size, NULL, NULL);
cl_mem bufResult = clCreateBuffer(context, CL_MEM_WRITE_ONLY, size, NULL, NULL);

// 4. 创建并构建程序
const char* source = "kernel void vec_add(...) {...}";
cl_program program = clCreateProgramWithSource(context, 1, &source, NULL, NULL);
clBuildProgram(program, 1, &device, NULL, NULL, NULL);

// 5. 创建内核并设置参数
cl_kernel kernel = clCreateKernel(program, "vec_add", NULL);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufA);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &bufB);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &bufResult);

// 6. 执行内核
size_t global_work_size = N;
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_work_size, NULL, 0, NULL, NULL);

// 7. 读取结果（阻塞方式，隐含了等待内核完成）
clEnqueueReadBuffer(queue, bufResult, CL_TRUE, 0, size, host_result_ptr, 0, NULL, NULL);

---

总结 📚

本节课中我们一起学习了OpenCL的基础知识。我们了解到OpenCL是一个用于异构计算的开放标准框架，它通过数据并行和任务并行模型，允许程序利用CPU、GPU等多种设备的计算能力。

我们探讨了OpenCL的平台模型、执行模型（特别是基于工作项和工作组的数据并行）以及显式的内存层次结构。我们还介绍了OpenCL框架中的核心对象（如上下文、队列、缓冲区、内核）和基本工作流程。最后，我们简要了解了用于编写内核的OpenCL C语言，并通过一个向量加法的例子串联了整个编程过程。

OpenCL提供了对硬件底层的控制能力，虽然编程模型稍显复杂，但对于具有大量并行性且对性能有要求的应用程序，它能带来显著的加速效果。在接下来的课程中，我们将深入探讨GPU架构、OpenCL C语言细节以及性能优化技巧。

012：从图形处理器到通用计算

在本节课中，我们将要学习GPU架构的基础知识，特别是从图形处理到通用计算的演变过程。我们将探讨GPU设计的核心理念，以及它与CPU的根本区别。

什么是GPU？

从计算程序员的视角来看，GPU是一个为处理像素而优化的多核处理器。它拥有宽SIMD核心，以实现高度的数据并行性，其核心设计侧重于吞吐量而非延迟。

图形处理基础

上一节我们介绍了GPU的基本定义，本节中我们来看看它的起源——图形处理。

假设我们正在处理单个像素的数据。我们可能有一个定义在缓冲区中的光源、像素的颜色和法线等属性。我们可以从输入寄存器中读取这些值，计算光源如何影响像素，并生成输出颜色。这个过程很简单。

当然，我们可以同时处理多个像素。在图形处理中，四边形是大多数GPU处理的基本单位。四边形是渲染API绘制到屏幕上的多边形的一部分。光栅化硬件会根据DirectX或OpenGL传递的多边形输入流来生成四边形以供硬件处理。

当我们想要生成输出结果时，会执行一个片段程序。这个程序使用GLSL或HLSL等语言编写。重要的是，程序被写成独立的形式，它只与单个像素相关，完全不考虑像素如何组合在一起。程序的输出仅取决于与该特定像素相关的输入，并且没有像素间的通信。这使得硬件能够进行非常高效的并行处理，而无需进行依赖性分析。

早期的通用GPU计算就是以这种方式进行的，使用屏幕多边形来生成大型矩形工作负载。

此外，对于每个多边形，我们可能运行一个略有不同的片段程序。因此，除了需要执行顶点数据处理、光栅化等操作外，GPU在线程级并行方面也有很高的要求，这也影响了其设计。

SIMD执行与分支

上一节我们了解了GPU如何处理图形像素，本节中我们来看看这种处理方式如何塑造了其核心执行模型——SIMD。

像素数据是图形工作负载的主要组成部分，因此影响了硬件的设计方式。我们仍然以四边形为例。编写代码时，为每个像素单独考虑程序是合理的，因为这是我们生成数据的层面。我们通常不知道像素在屏幕上的具体位置或其与其他像素的关系。

这意味着硬件层面的数据并行执行是从一个无依赖性的程序中推断出来的。我们知道每个像素可以独立并行执行，但同时它们执行的是相同的程序（或至少是同一小组像素），因为它们运行的是相同的着色器代码。GPU的前端光栅化和工作调度硬件会将来自同一多边形的四个相邻像素打包成一个四边形，并努力将来自相似多边形或单个较大多边形的四边形打包在一起，以便执行相同的着色器代码。

由此我们首先可以看到，我们可以将这些像素映射到一个SIMD引擎上。这意味着单指令，多数据。所有四个像素将同时执行相同的指令和程序。这类似于x86处理器的SSE扩展。在AMD的GPU硬件上，这被称为一个波前。

这带来了很高的ALU密度。但当代码中出现分支时，问题就来了。现在，多个独立的程序在单个SIMD单元上一起执行，这意味着它们都执行那条单一指令。如果分支走向一致，则没有问题。如果不一致，我们可能会看到以下情况：

首先，从内存中进行的第一次收集操作将是单个指令中每个像素的内存地址集合。然后，我们遇到一个条件判断。这将生成一个掩码。如果掩码指示所有分支走向相同，硬件可能会执行一个“全部投票”指令，将其转换为分支。但如果像本页所示，并非所有像素都满足条件，我们不希望它们全部执行。然而，由于这是SIMD指令，无法避免执行，指令仍会为所有四个通道发出。因此，硬件会使用计算出的掩码来屏蔽掉不满足条件的像素，只输出满足条件的像素的结果。然后，当执行到else部分时，我们将反转掩码，剩余的像素将由后续的条件赋值指令处理。最后，返回语句将再次为所有四个通道执行。

因此，原本看似四个操作（读取、测试、赋值、返回）的着色器，变成了六个操作：条件赋值、反转第一个掩码、条件赋值和返回。这对于四边形中的所有四个像素，乃至整个64宽的波前中分组在一起的所有像素都是如此。特别需要注意的是执行轨迹中的间隙，这些是ALU被浪费的地方，SIMD通道在执行指令但没有写入有用结果。显然，if-else块越大、条件嵌套越深、波前越宽、不同通道间的计算分歧越大，SIMD引擎的整体利用率就越低，我们获得的数据并行执行效率也就越低。因此，尽管硬件由于这些宽SIMD引擎和后续讨论的效率而具有非常高的峰值吞吐量，但实际的平均吞吐量会因这种分歧而低得多。

SIMD指令的必要性

上一节我们看到了分支对SIMD执行效率的影响，本节中我们来探讨一个相关问题：SIMD执行是否一定需要SIMD指令？

像素着色器的处理方式（这也是我将其作为从像素逐步引入的原因）是，我们在不使用向量指令的情况下编写了SIMD程序。向量指令可以由硬件实时生成（自动掩码）或由编译器生成。在x86世界中，人们习惯于使用SSE intrinsic等指令。思考这些指令的难度会使开发变得繁琐，并且需要相当高的技巧。我们必须考虑如何手动打包不同的操作、收集/分散寄存器，然后对它们发出收集或分散指令。这意味着我们还需要手动处理分支掩码。我们必须仔细地手工编写所有这些代码，虽然可能，但对于非常宽的真正向量编程来说，这有点像一门艺术。

另一方面，显式向量指令也有优势。以这种通道方式编写SIMD代码，存在让开发者误以为每个通道是独立分支的风险，误以为着色器或OpenCL内核的每个实例是独立分支的。虽然在硬件上可能看起来如此，但实际上并非如此。如果你想获得良好的性能，这一点必须考虑进去。

因此，对于宽SIMD硬件，编程方式各有利弊。以这种通道方式编程似乎更容易让程序员理解。对于当前的AMD GPU，掩码由硬件控制，并隐含在中间语言中。

这对计算有何意义？

这为何对计算很重要？这正是本次讲座的目的。传统上，图形代码的着色器相对较短，且处理的是较大的三角形，因此随着时间的推移，分支分歧的影响会减弱。总体而言，分支分歧的水平通常不会很高，而且在这些情况下也难以控制，因为你无法清楚地知道哪些像素会映射到同一个SIMD引擎。即使你知道有一个大三角形，你也不知道摄像机会离它多近。因此，你只能有限地规避分支分歧。

而在OpenCL代码中，你可以精确定义执行空间，选择哪些工作项执行哪些工作，以及这些工作项在网格中相对于其他工作项的位置。因此，你可以选择如何构建算法来避免这种分歧。这意味着，如果你的OpenCL代码出现分支分歧，那是你的责任，是你决定让它这样的，希望这是因为该算法需要如此。本系列后续的讲座更适合解释这在优化方面的实际意义。

吞吐量计算

上一节我们讨论了SIMD和分支，本节中我们来看看GPU架构的另一个主要可见方面——为吞吐量计算而设计。

这意味着什么？假设我们有一个四宽的SIMD向量正在执行指令，并且它发生了停顿。比如，它正在等待浮点数加法完成所需的周期数。一个具有复杂控制逻辑的快速时钟、乱序执行的CPU会尝试用同一指令流中的其他不依赖于停顿指令结果的指令来覆盖这个停顿周期（因此称为乱序执行），或者通过流水线中的数据前传来让下一条指令比从寄存器读取更早地执行。所有这些都需要复杂的硬件，这些硬件会占用本可用于ALU的空间和功耗。

因此，GPU采取了一种略有不同的方法。在SIMD架构上，我们通常在多个周期上运行一个更宽的向量。这意味着我们可以使用相同的指令来覆盖停顿时间。图中向量第一个四分之一的停顿时间被第二、第三和第四个四分之一执行相同指令所覆盖。这也有减少指令解码带宽的好处。因此，一旦你看到这些周期中ALU在做什么，停顿时间就会短得多。

在5870 GPU上，我们在16宽的硬件SIMD单元上运行一个64宽的向量，以实现这种四周期设计。在这个非常简单的例子中，我们完全覆盖了停顿，ALU保持忙碌，达到了100%的利用率。

当然，那是一个非常简单的案例。如果停顿时间更长呢？比如，我们正在等待内存请求的返回数据，而编译器无法在访问和使用之间插入其他指令来覆盖。即使使用我们的宽向量，也会出现间隙。我们可以执行更宽的向量，将向量宽度加倍，这可能就足够了。但向量可以越来越宽，然而由于我们之前看到的SIMD分支分歧问题，这只会增加低效率，降低ALU的利用率。因此，这可能不是正确的解决方案。

另一种方案是，我们可以从另一个正在运行的线程中“插入”另一条指令。在这里，我将一个线程定义为一个最终在同一个SIMD单元上的波前。我们可以看到，线程B的指令0填补了线程A指令0和1之间的间隙。在每个SIMD单元上运行的波前越多，我们就能以这种方式覆盖更多的延迟。如果我们有足够多的波前和ALU指令，就几乎不需要让SIMD单元等待纹理单元返回内存数据。显然，有些情况下这是不可行的，但在大多数情况下，你几乎可以做到这一点。

这意味着单个波前执行需要更长的时间。波前第一个四分区的第一条指令和第二条指令之间现在不是间隔四个周期，而是八个周期。你插入的波前越多，这个间隙就变得越长。但关键在于，我们不是试图最小化单个线程在硬件上的执行延迟，而是试图最大化整个线程集的吞吐量。这意味着我们努力提高利用率，而不关心第一个线程完成需要多长时间。

这样做的效果是，我们原本只有一个SIMD引擎，但由于有许多线程来覆盖延迟，我们可以在下一个周期送入下一个波前或波前的另一部分。反过来，一旦我们获得了第一组像素正在进行的任何计算结果，我们就可以将其反馈回顶部。我们可以用尽可能多的像素重复这个过程，以保持流水线尽可能接近满负荷运行。

但这里需要注意的是，我们现在有多个像素。虽然它们实际上并不都在ALU流水线中，但它们正在等待进入，你当然不希望它们在准备进入之前出现延迟。这意味着我们需要为每个等待使用ALU的波前存储状态。每个在运行中的线程都有自己的一组寄存器状态，覆盖其使用的所有寄存器。而且，这组寄存器状态必须覆盖向量的整个宽度。因此，这种寄存器状态在设备上占用空间，并且会随着状态集的数量和向量的宽度而扩展。

GPU通过维护一个大的寄存器池来工作，这些寄存器只是根据每个波前的资源需求来通用地分配给状态集。大多数多线程CPU都有多个相当精确的架构状态集副本，因此往往有一个固定的最大数量。但关键点是，你想要的波前越多，并行度越高，你需要的状态集就越多。

内存层次结构

上一节我们讨论了如何通过多线程隐藏延迟，本节中我们来看看这对内存访问意味着什么。

我们不太关心延迟，因为我们的目标是吞吐量。这是GPU与CPU的另一个不同之处。CPU希望尽可能减少内存读取的平均访问时间，其缓存层次结构就是为此设计的。你希望尽可能从最接近的缓存层级读取，因为那具有最少的周期延迟。这是因为CPU内部状态较少，且试图以低延迟实现高吞吐。

但由于GPU的设计已经考虑了覆盖延迟，我们可以从一个略有不同的角度来看待这个问题。我们现在运行着大量线程。正因为有大量线程，它们会访问大量内存地址。因此，我们现在必须最大化可用内存带宽，以满足所有这些并行性，而不是试图最小化延迟。

一种方法是GPU使用高带宽、高延迟的内存接口，如GDDR。但即便如此，减少数据流量也是值得的。在大量线程的情况下，所需的数据总量可能非常庞大。这就是纹理缓存和本地内存的用武之地。这两者都是为了支持工作项之间的数据重用而设计的。

在我们的架构中，本地内存区域被称为本地数据共享。这意味着它们支持空间局部性，而非时间局部性。这些缓存和程序员控制的共享内存区域通过允许数据仅通过主内存接口复制一次，然后由不同工作项以不同路径重用来实现数据传输的减少。

纹理缓存特别为此设计，它可以自动对数据强制执行二维结构，以高效捕获四边形和同一多边形上像素可能从输入缓冲区访问的紧密相关数据。

对于本地数据区域，我们有更多的控制权。与SIMD分支结构类似，高效使用此内存、共享数据以减少全局内存需求是开发者的责任。

本地内存相对于缓存的另一个好处是，它可以更高效地构建在芯片上。因为它不需要自动缓存所需的标签查找和管理结构。因此，对于相同数量的晶体管，可以实现更高的容量和带宽。

硬件设计权衡

上一节我们介绍了内存层次结构，本节中我们来看看这些设计选择如何影响硬件的整体布局。

你应该记得我之前说过，每个SIMD单元共享一个指令解码块。通过在多个ALU之间以SIMD方式共享此硬件，我们最小化了控制逻辑消耗的硅片面积。随着我们增加状态集的数量，这些指令解码步骤的尺寸保持相对固定，但我们增加了执行块的尺寸，将更多的硅片面积专用于有用的计算。在这种情况下，执行块要大得多。我们之所以能做到这一点，不仅是因为我们可以扩展SIMD引擎，还可以增加状态集的数量。因此，这个图的垂直维度实际上是系统中的状态量，水平维度是SIMD向量的宽度。我们获得了更多的指令解码硬件复用。

这就是为什么当前的GPU拥有非常宽的SIMD引擎和大量的状态集。随着状态集数量的增加，不仅寄存器状态量会增加，本地存储和缓存的需求也会增长。因此，仍然会存在一些权衡，基于芯片上非ALU硬件的数量与可实现的吞吐量之间。但总会有一个点，纯粹的数据并行性不再足够。

正如我们之前看到的，并非所有多边形都需要运行相同的片段程序，特别是在进行计算代码时。如果我们精心设计代码，了解它将如何启动，比如让每个工作组以不同的方式分支，那么就没有分支分歧的问题，因为这些代码集是完全独立的。这意味着我们可以达到一个点，在硬件上拥有多个核心。根据工作负载，设计更多但更窄的SIMD核心可能更高效。具有高度依赖性和控制流的程序，可能在不依赖宽SIMD的架构上运行得更好。而最大的ALU密度将通过一个执行单一指令流的巨大SIMD单元实现，但那将没有延迟隐藏能力。

因此，你添加的每个状态集或核心都会降低ALU密度，但好处是提高了ALU的利用率。这种权衡的落点很大程度上取决于工作负载。GPU的权衡之所以如此，是因为图形工作负载具有非常高的并行性。

本质上，从通用计算程序员的角度看，GPU就是一组宽SIMD核心的集合，每个核心承载多个程序状态，并交织许多线程以覆盖流水线延迟。这意味着我们可以将GPU置于现有CPU核心设计的设计空间中来思考。GPU设计并非在抽象意义上完全不同，它只是这个设计空间中的一个点，是权衡以特定方式落定的结果。

GPU在处理器设计空间中的位置

上一节我们讨论了硬件设计的权衡，本节中我们通过几个简单的例子，将GPU置于更广阔的处理器设计空间中来理解其定位。

这些例子略有简化，但应该能传达基本思想。例如，AMD的Phenom II X4核心。众所周知，每个核心有一个状态集，每个核心运行一个硬件线程，并使用四宽SIMD（实际上有多个SIMD流水线）。

Intel的i7采取了略有不同的方法，Pentium 4也做过类似的事情。我们现在有一个稍大的核心，但它包含两个状态集。因此，它们可以在同一核心上同时交错执行在不同ALU上的指令。这有助于增加可用的指令级并行性，而不是仅仅尝试从长指令序列中提取。但这确实是以更复杂的控制硬件为代价的。

Sun的UltraSPARC T2（Niagara 2）以及最初的Niagara采用了一种有趣的方法。它有八个核心，每个核心有八个状态集。这是为标量网络处理设计的，因此它完全放弃了数据并行性，转而追求高度的线程级并行性。

当然，AMD的GPU设计可以说是这个频谱的极端。考虑到GPU工作负载的高数据并行性，这正是我们所期望的。市场上的竞争性GPU选择了相当相似的设计点。例如，NVIDIA的GTX 480有15个核心和16宽SIMD，同样有大量的状态集，在如何提取指令级并行方面实现方式略有不同，但在设计空间中的总体位置是相似的。

因此，如果我们考虑GPU和CPU设计是否会融合的问题（忽略将它们放在同一芯片上的融合，而是真正的架构融合），我们真正要问的是工作负载是什么，以及它们如何被最佳执行。当前的CPU工作负载在GPU上执行得并不好，反之亦然。在它们融合之前，设计将始终处于设计空间的不同点上。

AMD Radeon HD 5870架构细节

上一节我们将GPU置于更广阔的设计空间中，本节中我们来具体看一个GPU实现的技术细节——AMD Radeon HD 5870。

从计算程序员的角度来看，大多数相同的细节也适用于该系列中更低端的设备以及6000系列。总体而言，它是一个2.72 TFLOP（单精度）和略低于550 GFLOP（双精度）的架构，拥有21.5亿个晶体管。特别要注意的是，本地数据共享的带宽为每时钟周期2560字节，约合2.1 TB/s。这比L1缓存的带宽高一倍，原因如前所述，与缓存相比，设计可编程寻址的共享内存效率更高。它比外部内存带宽高出一个数量级。

设备上并发波前的数量为496，但这是基于调度硬件可以处理的数量。实际上，并发波前的数量将取决于每个波前使用多少寄存器（总共5.24 MB通用寄存器）以及本地数据共享（640 KB），这些资源根据每个波前和工作组的需求进行分配。

当内核分派命令处理器解码内核并生成一定数量的波前到设备时（或者如果是处理光栅化像素数据，则生成四边形或打包在一起），如果SIMD单元有空闲资源，并且整个工作组恰好有足够的资源适合该空间，它将被分配给该硬件。整个芯片有20个SIMD引擎（可以理解为20个核心），八个GDDR5内存库、全交叉开关和L2缓存。此外，还有用于全局内存原子操作的写组合缓存和读-修改-写缓存。

一个重要的设计特点是，它遵循宽松的全局内存一致性模型，需要栅栏指令来确保写入的可见性。这与x86设计（主要是强一致性）不同。这确保了SIMD引擎和纹理单元能够保持高度的数据并行性，而无需复杂的硬件支持。

一个不寻常的特点是，硬件有一个基于子句的执行模型。在流内核分析器中查看反汇编时，你会看到这一点。这里的绿色代码是控制流代码，由顶部的两个标量单元执行。每个单元交错执行来自所有线程的控制流程序。因此，它们之间可以交错执行多达496个波前的控制流。

示例的第一行是一个“ALU子句开始”指令，这意味着定序器将通过将该子句的内容分派给当前波前状态所在的任何SIMD单元来执行该指令。然后，该标量单元将执行该子句，而定序器则去执行其他波前的控制流子句代码。其他SIMD单元也将执行它们自己的子句。在子句执行结束时，会生成一个谓词值用于继续控制流。这里它执行一个跳转指令，这可能意味着它必须生成一个分支掩码，或者根据谓词的投票结果实际进行分支。掩码的生成将由硬件自动完成。

当遇到纹理子句时，它将被发送到内存控制器，纹理单元与SIMD引擎相连。当数据获取正在执行时，其他控制流程序的其他子句可以在SIMD引擎上执行。这增加了设备上可能的并行执行量，并有助于覆盖延迟。如果有大量子句可供执行，这种开销是最小的。但如果你的程序没有太多的线程级并行性，没有发出足够的波前来占用设备，那么开销可能会变得显著。因此，在为峰值性能调整代码时，查看这些指令块以获得代码可能如何执行的概览是很重要的。

SIMD引擎与处理单元

上一节我们概述了5870的顶层架构，本节中我们深入看看其核心——SIMD引擎。

在之前的图表中，我们看到了20个SIMD引擎。这是一个简化图。一个SIMD引擎包含16个处理单元，构成物理上的SIMD硬件。这些处理单元将在波前上跨ALU子句执行单条指令，持续8个周期（每周期执行一部分）。因此，应用程序中的每个波前都位于给定的SIMD引擎上，其所有寄存器和LDS数据都维护在那里，移动这些数据将是低效的。每个引擎可以同时执行来自多个内核的波前，在那些指令上交错来自不同子句的波前。当然，SIMD引擎中的通道将根据我们讨论的掩码机制被启用或禁用。

在SIMD引擎块内，有两个对计算程序有用的主要组件：第一个是本地数据共享，它允许工作组内的工作项共享数据；另一个是处理单元（有时容易混淆地称为流处理器），它执行内核中的ALU子句指令。如前所述，有16个处理单元。

本地数据共享包含32个存储体。SIMD引擎中的16个处理单元每个周期可以从LDS的任意地址请求读取或写入两个32位字（至少读取是任意地址，写入由于指令中的槽位数量限制，可以是基地址和步长）。在LDS存储体发生冲突时，单元将检测到冲突并在后续周期重新发出读/写请求，直到所有请求完成。原子操作使用图中底部附近与每个存储体关联的一排小型、仅支持整数操作的ALU来执行。这样，非返回型的LDS原子操作完全在LDS内部处理，独立于SIMD引擎的其余部分。这使得具有大量原子操作的代码能够实现非常高的指令吞吐率。由于直接在LDS中完成锁定等操作非常高效，但这并不意味着浮点原子操作会更复杂。

LDS的带宽远高于外部内存，并且比缓存更高效。因此，每个SIMD引擎每时钟周期可以实现1 KB的带宽，单指令延迟为8个周期（注意，单指令延迟并不意味着一个周期，实际上是八个）。所有这些都可以通过LDS进行流水线处理。这是吞吐量执行的另一个例子。我们可以通过下一条指令可能到达的时间来轻松覆盖这个峰值周期延迟。

SIMD引擎的另一个特性是处理单元。5870架构的处理单元是一个由五个ALU组成的集群，它们操作一个超长指令字（VLIW）包。这意味着，不是执行单条指令，而是执行一个由编译器打包的、具有一组已知依赖关系的五操作指令包。其中四个ALU是相同的，基本上可以执行任何32位浮点操作，或者可以组合起来执行更复杂的操作。第五个ALU也能执行大多数基本操作，此外还包含一个可以生成超越函数（如sin/cos）的特殊功能单元。

VLIW设计意味着编译器必须生成指令包才能充分利用设备。但通过将这项工作转移到编译器，硬件可以非常密集地封装ALU。我们再次进行了权衡，通过减少控制硬件来追求高峰值吞吐量。这就是为什么该设计具有非常高的峰值性能。当然，你永远无法完全达到峰值，因为你永远无法获得100%的封装率。同样，这可以在ISA代码中看到。分析工具会为此提供线索。

一个特性是，它可以在单个包中共同发出某些依赖操作集，这有助于增加编译器可以实现的封装率，不一定需要五个完全独立的操作。例如，可以发出一个点积4操作，它需要在通道之间进行加法操作，这可以作为一个包实现。每个通道还可以执行24位整数操作，要执行32位整数乘法，则必须组合通道。有不同方式可以组合通道，编译器知道这些方式。

最后一个特性是全局数据共享。我最初提到过它，在大的架构图中可以看到。它的大小为64 KB，是LDS的两倍，基本功能相同。重要的是，它连接到设备上的整个SIMD引擎集合。目前，在OpenCL或DirectCompute中，它并未完全作为内存暴露。在DirectX中，它通过特定方式暴露。你可以想象在OpenCL中暴露它。它支持更快的原子计数器。DirectX中可见的追加缓冲区功能（允许你从每个工作项创建带有可选输出的紧凑输出数据集）就是使用这个来加速的。这些特性允许加速特定结构，尽管它们可能不作为通用内存可用。它提供25个时钟周期的延迟，远低于全局内存，并且每个时钟周期可以处理八个工作项请求。

总结与问答

本节课中我们一起学习了GPU架构的基本原理，结合处理器设计空间进行了探讨，并概述了AMD Radeon HD 5870架构中影响计算应用的基本特性。后续讲座将深入探讨GPU优化以及这些架构特性的实际使用方式。

以下是问答环节的部分内容整理：

• 问：波前中的工作项可以访问纹理缓存中的随机像素吗？这会导致性能问题吗？
  答： 可以访问随机像素。如果访问是聚集的，会导致性能问题。如果访问能落在缓存行内，会比全局内存访问快。具体性能取决于缓存结构和存储体冲突情况，与LDS类似但更复杂。

• 问：第五个ALU中的特殊功能单元（超越函数）是原生版本还是完全精度的OpenCL函数？
  答： 是原生版本。如果OpenCL中的原生函数生成了它们，那就是这些。由该通道生成的超越函数不是完全精度的，它们是快速近似，对大多数图形代码足够好，对计算代码是否足够取决于你的需求。完全精度需要执行更长的指令序列。

• 问：获取设备完全占用的最佳方式是什么？工作组大小应该是16或64的倍数吗？
  答： 这是一个复杂的问题。一个基本点是，工作组大小应该是64的倍数。如果工作组只有32个项，你只能获得50%的占用率。实际的完全占用取决于许多因素：每个波前使用的寄存器数量、每个工作组所需的本地内存量、工作组大小、屏障数量、分支分歧程度等。这是一个需要综合权衡的复杂问题。

• 问：OpenCL中是否有方法查询设备的波前大小？
  答： 没有直接的方法。在OpenCL 1.1中，有一个查询可以获取“首选工作组大小倍数”，我们的运行时通常会返回64。但这不能保证，你可以为你可能运行的设备创建一个查找表。

• 问：能否解释一下“完全合并访问”？
  答： 这指的是全局内存访问。最简单的情况是，你的波前将在一个操作中发出内存请求，每个通道访问连续的、对齐的地址上的浮点数或4维向量。这是内存访问的绝对峰值。如果访问未对齐或是随机收集，效率就会下降，性能取决于是否是缓存访问。

• 问：波前和工作组在OpenCL术语中是什么关系？
  答： 工作组是一组工作项。如果高效编程，工作组应该是一个或多个波前（即一组或多组64个工作项）。你希望工作组大小是64的倍数。一个工作组可能只有一个波前大小，也可能包含多个。如果包含多个，那么当你发出屏障指令时，同步将发生在这多个波前之间。如果工作组只有一个波前，那么屏障指令没有意义，会被优化掉。

• 问：如果内核需要很多寄存器，不同的工作组还会切换以隐藏延迟吗？寄存器上下文会被换出到RAM吗？
  答： 在当前设计中，寄存器上下文永远不会存储到RAM中。硬件不会自动将寄存器换出到RAM。一个工作组会在一个核心上运行直至完成。这就是为什么你不能保证调度中工作组之间的全局同步，也不推荐全局同步，因为你无法控制任何给定时间设备上实际有哪些工作组。然而，如果你的内核使用了大量寄存器，编译器可以生成将寄存器数据移动到位于VRAM中的“私有内存”再移回的代码，这称为“溢出”。这样，编译器可以减少内核所需的寄存器数量，但硬件不会自动溢出。

013：N体算法实现教程 🚀

概述

在本教程中，我们将深入学习OpenCL编程，重点讨论N体算法的实现。我们将从OpenCL基础概念回顾开始，接着介绍简化主机端编程的StandardCL库，最后详细剖析N体算法的内核代码和主机端代码实现。

---

OpenCL基础概念回顾

OpenCL为混合CPU/GPU架构提供了一套平台和运行时层，用于管理跨多个设备的并发操作执行。它包含一个C语言扩展，用于实际编程设备（如GPU）。最重要的是，OpenCL API是一个具有广泛行业支持的平台和设备无关的API。

OpenCL API的基本结构包括：

• 语言规范：用于编程设备（如GPU）的C语言扩展。
• 平台API：提供查询系统可用资源并为其设置计算层的例程。
• 运行时API：在主机端用于管理内核对象、内存对象以及在OpenCL设备上执行内核。

执行模型

执行模型分为两部分：

• 内核：代表将在OpenCL设备上运行的可执行代码，支持数据并行和任务并行编程模型。
• 主机程序：在主机端执行，负责内存管理以及通过命令队列在一个或多个设备上管理这些内核的执行。

内核编程的C语言扩展

基于ISO C99，并有一些限制和扩展以支持并行性：

• 向量数据类型：映射到许多设备（如支持向量操作的GPU）的类型。
• 工作项和工作组：可以将其视为线程，线程被分组到工作组中，并提供了内置函数来帮助内核管理执行。
• 同步内置函数：由于编程模型需要跨线程的细粒度并行化，因此需要同步能力。
• 地址空间限定符：反映了混合系统中分布式内存元素的复杂内存模型。
• 大量内置函数：例如许多数学内置函数。

数据并行性表达

理解内核编程首先需要理解如何用这些内核表达数据并行性。本质上，你需要定义一个N维计算域（N可以是1、2或3维）。计算域中的每个元素称为一个工作项，可以将其视为一个线程。

在这个N维域中，有一个全局维度，定义了将并行执行的工作项或线程的总数。核心思想是每个工作项或线程执行相同的内核。

工作项被分组为工作组，在同步和共享内存访问方面具有特殊属性。有一个局部维度定义工作组的大小。通常，工作组内的线程在同一计算单元上一起执行，可以访问相同的共享本地内存，并且能够同步。相反，在不同工作组中执行的线程不会被同步。

主机端执行模型

主机程序中的所有内容都集中在一个上下文中。上下文内包括：

• 要执行内核的设备集合。
• 程序对象（OpenCL支持即时编译模型）。
• 内核本身（已编译和链接的可执行文件）。
• 内存对象（支持跨设备管理内存的缓冲区）。
• 命令队列（每个设备一个），用于排队内核、内存/数据传输操作以及同步操作。

内存模型

OpenCL支持不同层次的内存类型，以反映混合系统中分布式内存的性质：

• 全局内存
• 常量内存
• 本地内存
• 私有内存

OpenCL具有完全宽松的内存一致性模型，这意味着程序员必须显式管理内存内的数据传输。

主机端同步

同步变得非常重要，尤其是在使用多个设备时。程序员有责任同步事件（包括内核执行和数据传输），可以通过命令队列或显式阻塞主机端的某些事件来实现。

---

使用StandardCL简化主机端编程

上一节我们回顾了OpenCL的核心概念，本节中我们来看看如何简化主机端编程。OpenCL提供了对内核执行和数据移动的非常明确且平台/设备独立的控制。在实践中，OpenCL有很多步骤（尤其是设置上下文）本质上是样板代码，每次都会重复执行。

StandardCL的理念是基于典型用例提供简化接口，并以类似UNIX的风格构建。我们将使用StandardCL来简化后续代码剖析中的主机端代码。其优点是允许我们专注于概念，而不会迷失在大量底层语法中。

需要指出的是，此API在访问OpenCL提供的全部功能方面没有限制，API调用本身非常接近原始的OpenCL调用。

以下是我们将使用的一些简化API调用：

获取计算层上下文

使用OpenCL，你必须查询可用平台，选择平台，获取该平台的所有可用设备，为每个设备创建上下文，然后为每个设备创建命令队列。而StandardCL提供了默认的、开箱即用的上下文。只需包含 standardcl.h 并链接库，你就拥有了一个包含所有CPU和GPU设备的就绪上下文。

管理内核

OpenCL支持即时编译模型，因此你必须管理程序文本、原始字符串、创建程序、构建程序并创建内核。StandardCL提供了简化的方法，本质上可以通过两次调用获取内核：使用 cl_open 打开包含内核代码的文件获取句柄，然后使用该句柄按名称查询你想要构建和链接的内核。结果是你获得了一个准备就绪的OpenCL内核。

内存管理

OpenCL要求使用不透明的内存缓冲区，并且必须在命令队列中排队读写缓冲区命令来传输数据，这与大多数C程序员习惯的内存管理方式不同。StandardCL提供了 cl_malloc，它以一种与 malloc 非常相似的语义分配内存，只是分配的内存实际上可以在OpenCL设备之间共享。cl_malloc 实际上是在为你创建和管理这些内存缓冲区。

管理事件执行

使用OpenCL，你需要在命令队列中排队操作并管理结果事件。使用StandardCL，你可以使用 cl_msync 和 cl_fork。cl_msync 用于将数据同步到设备或从设备同步回主机。cl_fork 用于在选定的设备上执行内核。cl_wait 是一个同步调用，用于等待已排队的事件完成。cl_msync 和 cl_fork 都有阻塞和非阻塞的变体，这在开始使用多个设备以实现并发操作时非常重要。

---

N体算法代码剖析

现在，我们进入本教程的核心部分：N体算法实现的代码剖析。

算法概述

N体算法是一个相对简单的算法，通常用于展示许多加速协处理器的性能。它模拟了N个粒子在某种粒子-粒子相互作用下的运动（例如，万有引力）。该算法的计算复杂度为 O(n²)，这意味着如果将模拟中的粒子数量加倍，计算负载将增加四倍。这对于研究带宽与计算比率问题非常有用。

算法有两个主要步骤：

1. 计算每个粒子上的力：通过对系统中所有其他粒子的相互作用贡献求和来确定。这是算法的 O(n²) 部分。
2. 更新粒子位置和速度：使用基本的牛顿动力学在某个小时间步长内更新。

整个未优化的算法可以用几十行C代码编写。

OpenCL程序结构

OpenCL实现由两部分组成：

• 内核代码：编译后在GPU上运行，执行实际计算。
• 主机代码：不进行有意义的计算，但处理初始化和簿记任务，并协调OpenCL设备上的操作（主要是内存管理和内核执行本身）。

内核代码实现

内核代码的目标是提供一个可理解的、合理标准的实现，并尝试使用OpenCL的良好实践。需要记住OpenCL内核代码的上下文：内核将在索引空间内的每个工作项上执行。在这个应用中，我们有一个简单的一维索引空间，工作项的数量等于系统中的粒子数。

内核代码将为系统中的每个N粒子调用一次，其任务是使用牛顿力学更新一个粒子的位置和速度。

以下是内核代码的关键部分解析：

内核原型

__kernel void nbody(
    float dt,
    __global float4* oldPos,
    __global float4* newPos,
    __global float4* vel,
    __local float4* pblock)

• __kernel 限定符表示这是一个内核函数。
• __global 限定符表示这些指针指向全局内存。
• __local float4* pblock 将用作每个工作组的本地缓存。

大小和索引确定

float4 dt4 = (float4)(dt, dt, dt, 0.0f*dt);
int gti = get_global_id(0);
int ti = get_local_id(0);
int n = get_global_size(0);
int nt = get_local_size(0);
int nb = n / nt;

• get_global_id(0) 获取全局ID。
• get_local_id(0) 获取局部ID（相对于本地工作组）。
• get_global_size(0) 和 get_local_size(0) 获取全局索引空间和本地工作组的大小。
• nb 计算我们分块方案中的块数。

实际计算
内核首先读取粒子的位置和速度，并将加速度清零。然后开始循环遍历块。

协作读取填充缓存

pblock[ti] = oldPos[j*nt + ti];
barrier(CLK_LOCAL_MEM_FENCE);

内核缓存一个粒子位置，但依赖于工作组中的其他工作项或线程执行相同的操作来加载缓存。然后使用 barrier 等待工作组内的其他线程赶上。

循环遍历缓存的粒子位置并计算力

for (int j = 0; j < nt; j++) {
    float4 p2 = pblock[j];
    float4 d = p2 - p;
    float invr = rsqrt(d.x*d.x + d.y*d.y + d.z*d.z + eps);
    float f = p2.w * invr * invr * invr;
    a += f * d;
}

这里实现了与C代码相同的力计算，但使用了OpenCL的向量数据类型和内置函数（如 rsqrt）。

更新粒子位置和速度

vel[gti] += dt4 * a;
newPos[gti] = oldPos[gti] + dt4 * vel[gti];

使用 float4 数据类型，可以同时更新X、Y和Z分量。

主机代码实现

主机代码负责初始化、内存分配、内核设置、数据传输和协调内核执行。

初始化和参数设置

#include "standardcl.h"
int nParticle = 8192;
int nStep = 100;
int nBu = 20;
int nThread = 64;

分配共享内存

float4* pos1 = (float4*)cl_malloc(nParticle * sizeof(float4));
float4* pos2 = (float4*)cl_malloc(nParticle * sizeof(float4));
float4* vel = (float4*)cl_malloc(nParticle * sizeof(float4));

使用 cl_malloc 分配主机和GPU之间可共享的内存。

构建和创建OpenCL内核

void* clp = cl_open("nbody.cl");
cl_kernel krn = cl_kernel(clp, "nbody");

cl_open 返回已编译的OpenCL内核代码的句柄，cl_kernel 按名称提取内核。

设置计算域和内核参数

cl_ndrange ndr = cl_ndrange_1d(0, nParticle, nThread);
cl_arg(krn, 0, dt);
cl_arg_global(krn, 4, vel);
cl_arg_local(krn, 5, nThread * sizeof(float4));

设置一维计算域，并设置内核参数。注意，参数2和3（位置数组）稍后动态设置，以实现双缓冲方案。

数据传输到设备

cl_msync(0, pos1, nParticle * sizeof(float4), CL_MSYNC_TO_DEVICE);
cl_msync(0, vel, nParticle * sizeof(float4), CL_MSYNC_TO_DEVICE);

使用 cl_msync 将数据同步到GPU。

执行循环

for (int s = 0; s < nStep; s += nBu) {
    for (int b = 0; b < nBu; b++) {
        cl_arg_global(krn, 2, pos1);
        cl_arg_global(krn, 3, pos2);
        cl_fork(0, krn, &ndr, CL_EVENT_NOWAIT);
        // 交换pos1和pos2以进行双缓冲
        swap(pos1, pos2);
    }
    cl_wait(0);
    cl_msync(0, pos1, nParticle * sizeof(float4), CL_MSYNC_TO_HOST);
    // 输出结果
    nbody_output(...);
}

在循环中，动态设置位置数组参数，使用 cl_fork 非阻塞地排队内核执行，然后使用 cl_wait 等待所有排队的内核完成，最后使用 cl_msync 将数据同步回主机。

清理资源

cl_close(clp);
free(pos1); free(pos2); free(vel);

编译代码

编译主机程序时，需要提供OpenCL头文件路径和StandardCL头文件路径，并链接相应的库。内核代码文件（.cl）需要与可执行文件一起携带，因为它是在运行时使用即时编译方法编译的。

---

扩展到多GPU支持

上一节我们完成了单GPU的N体算法实现，本节中我们来看看如何修改代码以支持多个设备（以两个GPU为例）。

修改内核代码

内核代码的修改相对较小，主要是添加一个额外的参数来指向“远程”粒子位置（即另一个GPU负责更新的粒子）。

内核原型修改

__kernel void nbody(
    float dt,
    __global float4* oldPos,
    __global float4* newPos,
    __global float4* vel,
    __local float4* pblock,
    __global float4* oldPosRemote) // 新增参数

计算力的循环分为两部分
首先，循环遍历本地粒子位置（oldPos）计算力。然后，再循环遍历远程粒子位置（oldPosRemote）计算力，并将两部分力累加。

修改主机代码

主机代码的修改更为显著，涉及同步和内存管理。

分配半尺寸数组

float4* pos1A = (float4*)cl_malloc(nParticle/2 * sizeof(float4));
float4* pos1B = (float4*)cl_malloc(nParticle/2 * sizeof(float4));
// ... 为pos2和vel分配类似的半尺寸数组

因为要将粒子分配到两个GPU上，所以需要分配半尺寸的数组。

分割数据
将完整的粒子位置和速度数组分割到这些半尺寸数组中。

设置计算域
计算域的全局大小现在是原来的一半。

cl_ndrange ndr = cl_ndrange_1d(0, nParticle/2, nThread);

数据传输到两个GPU

cl_msync(0, pos1A, (nParticle/2)*sizeof(float4), CL_MSYNC_TO_DEVICE);
cl_msync(1, pos1B, (nParticle/2)*sizeof(float4), CL_MSYNC_TO_DEVICE);
// ... 同步速度数组

执行内核并设置参数
需要为两个GPU分别排队内核执行，并正确设置参数（包括指向本地和远程位置数组的指针）。

GPU间数据交换
这是多设备编程中新引入的关键步骤。在一个GPU更新了其负责的粒子位置后，另一个GPU需要这些更新后的位置来计算力。

// 将更新后的位置从两个GPU同步回主机
cl_msync(0, pos2A, ..., CL_MSYNC_TO_HOST | CL_EVENT_NOWAIT);
cl_msync(1, pos2B, ..., CL_MSYNC_TO_HOST | CL_EVENT_NOWAIT);
// 等待数据传输完成
cl_wait(0); cl_wait(1);
// 将交换后的位置数据同步回GPU
cl_msync(0, pos2B, ..., CL_MSYNC_TO_DEVICE | CL_EVENT_NOWAIT); // GPU 0 获取 GPU 1 的数据
cl_msync(1, pos2A, ..., CL_MSYNC_TO_DEVICE | CL_EVENT_NOWAIT); // GPU 1 获取 GPU 0 的数据

合并数据回完整数组（可选）
为了便于使用现有的辅助函数，可以将半尺寸数组合并回完整的数组。

---

总结

在本教程中，我们一起深入学习了OpenCL编程，并完成了N体算法的完整实现。我们从OpenCL的核心概念和编程模型回顾开始，了解了其执行模型、内存模型和同步机制。接着，我们引入了StandardCL库来简化主机端繁琐的样板代码，使我们可以更专注于算法逻辑。

随后，我们详细剖析了N体算法的内核代码和主机端代码，理解了如何将串行算法转化为并行内核，以及主机端如何管理内存、设置参数和协调执行。最后，我们探讨了如何将实现扩展到多GPU环境，这涉及到内核参数的调整、数据的分割以及关键的设备间数据交换步骤。

通过本教程，你应该对使用OpenCL进行异构计算编程有了更扎实的理解，并掌握了实现一个典型计算密集型算法（N体模拟）的基本方法。

014：真实世界OpenCL应用

在本节课中，我们将学习如何构建一个多路视频流处理应用。我们将探讨从视频解码、数据上传、GPU处理到最终显示的完整流程，并重点关注如何利用OpenCL运行时API（特别是C++ API）来优化性能，实现实时处理。

---

应用概述

我们计划构建一个能够同时打开、解码、处理并最终合并显示多个视频流的应用。核心目标是在保证实时性的前提下，高效利用GPU资源。整个处理流程将形成一个循环：解码 -> 上传至GPU -> 处理 -> 显示。

解码视频

上一节我们介绍了应用的整体目标，本节中我们来看看第一步：视频解码。解码的任务是将压缩的视频格式转换为可供处理的单帧二维图像。

以下是解码环节的关键点：

• 异步执行：解码操作应独立于主处理循环运行，以避免阻塞主线程。主循环只需请求获取最新解码完成的帧。
• 帧缓冲：建议解码多帧并缓冲，以应对可能出现的I/O延迟（如从硬盘读取数据），确保后续帧处理的流畅性。

上传数据与DMA

解码完成后，我们需要将帧数据从CPU主内存传输到GPU的显存中。OpenCL提供了简单的API来完成此操作。

在OpenCL C++ API中，写入图像对象的调用示例如下：

cl::CommandQueue queue;
cl::Image2D image;
std::vector<cl::Event> events;
// ... 初始化 queue, image 等
queue.enqueueWriteImage(image, CL_TRUE, origin, region, row_pitch, slice_pitch, host_ptr, &events, nullptr);

其中，CL_TRUE参数表示这是一个阻塞写入。调用将等待数据传输完成才返回。

为了提升性能，我们可以利用GPU的直接内存访问引擎进行异步传输。DMA允许数据在系统内存和GPU显存之间直接传输，无需消耗CPU或GPU的计算周期。要启用异步DMA传输，需要在创建缓冲区或图像时使用特殊标志，并合理安排程序逻辑，避免在数据传输后立即依赖该数据。

利用DMA优化：单缓冲与双缓冲

简单的单缓冲流程是：解码 -> 上传 -> 处理 -> 显示。这会导致处理步骤必须等待上传完成，影响性能。

我们可以重组循环顺序：处理并显示第N帧 -> 解码第N+1帧 -> 上传第N+1帧。这样，上传操作可以利用“处理与显示”的时间窗口并行执行，缓解了实时性压力。

更优的方案是采用双缓冲。我们同时处理两个帧（A和B）。流程如下：

1. 处理并显示已上传的帧A。
2. 解码下一帧A，并开始异步上传。
3. 与此同时，处理并显示已上传的帧B。
4. 解码下一帧B，并开始异步上传。

这样，上传操作（如下一帧A）与处理操作（如帧B）完全并行。整体性能不再取决于“处理+解码+上传”的总时间，而是取决于“处理与解码”和“上传”两者中耗时更长的部分，从而显著提升了效率。

视频处理与事件同步

现在数据已在GPU上，我们可以进行视频处理。OpenCL中，处理通常通过编写在内核中、作用于每个像素的算法来完成。我们创建一个与帧尺寸（像素数）相同的执行域来启动内核。

设置内核参数时需注意：许多视频滤镜的参数（如色彩校正）一旦设定便很少更改。在OpenCL中，内核参数是其状态的一部分，因此可以只设置一次，后续执行时直接调用，无需重复设置。

执行内核的API调用如下：

cl::Kernel kernel;
cl::CommandQueue queue;
std::vector<cl::Event> wait_events;
cl::Event out_event;
// ... 初始化 kernel, 设置参数等
queue.enqueueNDRangeKernel(kernel, cl::NullRange, cl::NDRange(width, height), cl::NullRange, &wait_events, &out_event);

最后两个参数至关重要：

• wait_events：一个事件向量，列出了当前操作所依赖的所有前置操作事件。运行时将等待这些事件完成后，才执行当前操作。
• out_event：一个输出事件，用于标识当前操作的完成状态。此事件可作为后续操作的依赖项。

通过这种方式，我们可以构建一个操作依赖图，让OpenCL运行时自动处理同步，而无需在主线程中显式等待每个步骤完成（例如使用queue.finish()或event.wait()），这对于性能至关重要。

性能剖析：事件分析

在异步执行模式下进行性能剖析需要特殊支持。OpenCL提供了事件剖析功能。需要在创建命令队列时启用剖析，然后可以从事件对象中查询四个关键时间戳：

1. CL_PROFILING_COMMAND_QUEUED：命令入队时间。
2. CL_PROFILING_COMMAND_SUBMIT：命令提交给设备的时间。
3. CL_PROFILING_COMMAND_START：命令在设备上开始执行的时间。
4. CL_PROFILING_COMMAND_END：命令在设备上执行结束的时间。

通过分析START和END的时间差，我们可以精确测量内核执行或数据上传的耗时，从而定位性能瓶颈。

显示与OpenCL-OpenGL互操作

所有处理步骤完成后，我们得到了最终的输出帧，它位于GPU显存中。为了将其显示到屏幕，我们需要将此帧数据传递给显示API。

为了避免将数据从GPU显存读回CPU主内存再传递给显示API造成的性能损耗，我们可以使用OpenCL-OpenGL互操作。这允许OpenCL和OpenGL共享数据（如纹理），但请注意，这是OpenCL规范中的可选扩展，需要查询设备是否支持。

启用互操作的关键步骤：

1. 查询设备扩展，确认支持cl_khr_gl_sharing。
2. 在创建OpenCL上下文时，设置相应的属性以开启与OpenGL的互操作。
3. 使用特殊的API，从已有的OpenGL纹理创建OpenCL图像对象，使两者共享同一块内存。

   clCreateFromGLTexture2D(context, CL_MEM_READ_WRITE, GL_TEXTURE_2D, 0, textureId, &err);
   

4. 在OpenCL操作共享纹理/图像前后，需要进行“获取”和“释放”操作：
   • enqueueAcquireGLObjects(): OpenCL获取对共享对象的使用权。
   • 执行OpenCL处理内核。
   • enqueueReleaseGLObjects(): OpenCL释放使用权，交还给OpenGL进行渲染显示。

在本应用中，只有最终合成的输出流需要创建为这种共享图像，并在处理前后进行获取和释放操作。

总结与问答回顾

本节课我们一起学习了构建实时多路视频处理应用的完整流程。我们深入探讨了利用OpenCL进行高效编程的几个核心方面：通过异步解码和双缓冲DMA上传来优化数据准备；利用事件依赖图实现处理步骤的异步执行与同步；使用事件剖析进行性能分析；以及通过OpenCL-OpenGL互操作实现零拷贝显示，最大化性能。

在问答环节，我们额外讨论了：

• 场景过渡效果：可通过随时间更新内核参数（如划像位置）来实现。
• 性能剖析：事件时间戳基于高分辨率计时器，反映操作在设备上的实际执行时间。
• 顺序执行队列：多数设备默认按入队顺序执行，但显式设置事件依赖是良好实践，尤其对于多设备或乱序队列。
• OpenCL vs OpenGL着色器：若已有高性能GLSL着色器，可结合互操作使用。但OpenCL在语言表达力和对硬件特性的控制上通常更具优势，适合新开发。
• 示例代码：AMD OpenCL SDK中包含查询扩展、互操作等大量示例，可供参考。

---

本节课中，我们从理论到实践，详细剖析了一个真实世界OpenCL应用的关键技术与优化策略。希望这些内容能帮助你更好地理解和运用OpenCL进行高性能计算开发。

015：OpenCL设备分割扩展

在本节课中，我们将要学习OpenCL的一个扩展功能——设备分割。这个功能允许我们将一个物理计算设备（如多核CPU）在逻辑上分割成多个独立的OpenCL设备，从而实现对硬件资源的更精细控制。我们将了解扩展的类型，并通过一个具体的并行向量填充案例来演示设备分割的实际应用。

OpenCL扩展概述

在深入设备分割之前，我们首先需要了解OpenCL扩展的通用知识。OpenCL标准定义了三种不同类型的扩展，它们为开发者提供了超越核心规范的功能。

以下是OpenCL的三种扩展类型：

1. KHR扩展：这是由OpenCL工作组正式批准的标准扩展。与OpenCL核心规范一样，实现KHR扩展需要通过Khronos组织的一系列一致性测试，以确保其跨平台兼容性。
2. EXT扩展：这类扩展由至少两个或更多的工作组成员共同开发，但尚未获得整个工作组的正式批准。因此，它没有强制的一致性测试要求，不同厂商的实现可能存在细微差异。
3. 厂商扩展：这类扩展由单个硬件或软件供应商开发，通常也只在该供应商的平台或设备上得到支持。使用这类扩展会牺牲代码的可移植性，因此应谨慎使用，通常用于调试等非生产环境。

所有扩展的官方文档都可以在Khronos的OpenCL注册网站上找到，确保了规范的公开透明。

设备分割扩展简介

上一节我们介绍了OpenCL扩展的通用分类，本节中我们来看看一个具体的KHR/EXT扩展——设备分割。

设备分割扩展的核心思想是，允许开发者将一个包含多个计算单元（例如，CPU的多个核心）的物理设备，在逻辑上分割成多个独立的OpenCL设备。例如，一个四核CPU在默认情况下会被OpenCL视为一个单一设备。通过设备分割，我们可以将其划分为两个、四个甚至更多个逻辑设备。

这种能力带来了显著的优势：

• 资源预留与控制：你可以为高优先级任务预留特定的核心。
• 负载与缓存优化：在NUMA（非统一内存访问）架构系统中，你可以根据内存亲和性来分割设备，确保相关任务组在共享缓存的核心上运行，从而提高缓存复用率，减少内存访问延迟。
• 保证执行进度：在某些并行算法（如我们后面将看到的管道模式）中，将工作组分派到独立的设备上可以避免因OpenCL运行时工作组调度顺序不确定而导致的死锁问题。

目前，设备分割扩展主要支持在多核CPU和IBM Cell Broadband Engine处理器上使用。AMD在其CPU实现中支持此功能，并且同时兼容AMD和Intel的处理器。未来，这项技术也有可能扩展到GPU等其他设备上。

应用案例：并行流填充算法

为了具体说明设备分割的用途，我们来看一个实际的编程案例：并行流填充算法。这个案例改编自一个实际的图像处理问题（Impac2流处理）。

问题描述：我们有一个输入字节流。每当在流中遇到一个值为 0xFF 的字节时，就需要在该字节之后插入一个值为 0x00 的填充字节。

顺序实现：用C语言实现这个算法非常简单。我们遍历输入数组，将每个字节写入输出数组。如果当前字节是 0xFF，则在写入该字节后，再向输出数组写入一个 0x00。这里的关键是，我们需要维护两个索引：一个用于输入数组（每次递增1），另一个用于输出数组（根据是否遇到 0xFF 递增1或2），以确保输出顺序与输入顺序严格一致。

直接并行化的挑战：如果我们尝试用OpenCL内核直接并行化这个循环（例如，每个工作项处理一个输入字节），就会遇到问题。虽然每个工作项都能正确判断是否需要填充并计算自己在输出数组中的位置，但OpenCL不保证工作项的执行顺序。因此，位置靠后的工作项可能先于位置靠前的工作项执行，导致最终输出数组中的元素顺序混乱，尽管填充本身是正确的。

解决方案：基于管道模式的并行算法

为了解决执行顺序问题，我们需要重新设计算法。上一节我们看到了直接并行的弊端，本节中我们引入管道模式作为解决方案。

管道模式是一种经典的并行计算模式。在这个模式中，我们将计算任务分解为多个阶段（或“管道工”），每个阶段处理数据的一个块，并将处理结果（或元数据）传递给下一个阶段。

我们将输入数据分成连续的块。每个计算单元（例如，一个CPU核心）负责处理一个数据块。处理过程包括：

1. 计算本数据块内需要插入的 0x00 的数量。
2. 计算本数据块在最终输出数组中的起始写入偏移量。

关键点：第2步的偏移量取决于前面所有数据块的写入长度。因此，每个计算单元必须等待前一个单元计算出其累积偏移量后，才能开始自己的写入操作。这就形成了一个管道：单元N等待单元N-1的信号，处理完后通知单元N+1。

使用独立设备保证进度：如果我们使用OpenCL的默认设备，并将所有工作组提交到同一个命令队列，OpenCL运行时无法保证工作组的启动和执行顺序。这可能导致死锁（例如，负责处理后面数据块的工作组先开始执行并空等，而负责处理前面数据块的工作组迟迟得不到调度）。通过设备分割，我们将每个核心创建为一个独立的OpenCL设备和命令队列。这样，我们可以显式地将每个数据块的处理任务提交到对应的核心设备上。操作系统调度器会保证每个核心上的任务都能获得执行时间，从而确保整个管道能够向前推进，避免死锁。

技术实现：结合设备分割与原生内核

为了实现上述管道方案，我们将结合使用设备分割和OpenCL的另一个功能——原生内核。

原生内核允许我们在OpenCL设备（特别是CPU设备）上直接运行普通的C/C++函数，而不仅仅是OpenCL C内核。这对于集成现有库、进行复杂调试或实现需要操作系统调用的机制（如更高级的线程同步）非常有用。

以下是实现步骤的概要：

1. 初始化与设备分割：

   • 查询平台和设备。
   • 选择支持设备分割的CPU设备。
   • 使用 clCreateSubDevices 函数和 CL_DEVICE_PARTITION_EQUALLY 属性将设备分割为多个子设备（例如，每个核心一个子设备）。
   • 为每个子设备创建独立的命令队列和上下文。

2. 准备数据与通信：

   • 创建输入/输出缓冲区。
   • 创建“邮箱”缓冲区数组，用于在管道阶段间传递累积偏移量。每个子设备对应一个邮箱。

3. 提交原生内核任务：

   • 我们将处理每个数据块的C++函数作为原生内核。
   • 遍历所有子设备（计算单元），为每个单元准备参数：
     • 输入/输出缓冲区指针（通过OpenCL内存对象传递）。
     • 邮箱缓冲区指针。
     • 数据块大小、块ID等。
   • 使用 clEnqueueNativeKernel 将每个任务异步提交到对应的子设备命令队列中。每个任务会收到一个事件对象。

4. 内核函数逻辑：

   • 函数内部，首先根据块ID读取对应的输入数据块。
   • 如果块ID不是0，则忙等待或通过更高级的机制轮询对应的邮箱，直到前一个阶段写入其累积偏移量。
   • 计算本块的输出偏移量（前一块的偏移 + 本块基础大小 + 本块中填充的 0x00 数量）。
   • 立即将本块计算出的新累积偏移量写入下一个阶段的邮箱，以通知其可以开始工作。
   • 将本块处理后的数据（包括可能的填充）写入输出缓冲区的正确位置。

5. 同步与完成：

   • 主机程序等待所有子设备命令队列返回的事件完成。
   • 所有事件完成后，输出缓冲区中即为顺序正确的填充后数据。

通过这种设计，我们既利用了多核并行处理数据块内部的计算，又通过邮箱通信保证了数据块间写入全局内存的顺序。设备分割确保了每个管道阶段都有独立的硬件资源，从而保证了执行进度。

性能与总结

本节课中我们一起学习了OpenCL设备分割扩展及其在一个并行流填充算法中的应用。

• 核心概念：设备分割允许将物理设备划分为多个逻辑设备，实现资源隔离和进度保证。
• 关键模式：管道模式结合邮箱通信，解决了并行计算中的顺序依赖问题。
• 实现工具：我们使用了原生内核来在CPU上执行复杂的C++逻辑，并利用设备分割创建了独立的执行上下文。

即使在这样一个计算密度不高（主要是条件判断和内存写入）的案例中，在一台四核机器上我们也观察到了 2.3 到 2.4 倍的性能加速。这充分展示了通过精细的硬件控制和并行算法设计，即使在CPU上也能利用OpenCL获得显著的性能提升。如果算法中包含更多的向量化计算，性能收益将更加可观。

设备分割的用途不仅限于此，它还在NUMA系统优化、实时任务预留等场景中发挥着重要作用。随着OpenCL生态的发展，我们期待看到更多设备支持这一强大的扩展功能。

016：平滑粒子流体动力学 (SPH) 🚀

在本节课中，我们将学习如何使用OpenCL实现平滑粒子流体动力学（SPH）方法，这是一种用于模拟流体（如水）行为的计算技术。我们将从流体力学的基本概念开始，逐步深入到SPH的数学原理，最后解析一个完整的OpenCL模拟程序。

---

概述：什么是流体？ 💧

流体是我们日常生活中常见的物质状态，主要包括液体（如水、油）和气体（如空气）。从物理学的角度看，流体遵循纳维-斯托克斯（Navier-Stokes）方程所描述的规律。本节课我们将专注于不可压缩的纳维-斯托克斯方程，它适用于描述在常规速度和温度下，像水这样的流体行为。

流体运动主要受三种力支配：重力、压力和粘性力。压力差驱动流体从高压区流向低压区，而粘性力则描述了流体的“粘稠度”，它影响流体各部分运动的协调性。

---

纳维-斯托克斯方程 📐

纳维-斯托克斯方程是流体动力学的核心偏微分方程组。对于不可压缩流体，我们主要关注两个方程：

1. 运动方程：描述了速度场随时间的变化。
   [
   \rho \left( \frac{\partial \mathbf{v}}{\partial t} + \mathbf{v} \cdot \nabla \mathbf{v} \right) = -\nabla p + \mu \nabla^2 \mathbf{v} + \rho \mathbf{g}
   ]

   • ρ 是密度。
   • v 是速度向量。
   • p 是压力。
   • μ 是动力粘度系数。
   • g 是重力加速度向量。
   • ∇ 是梯度算子，∇² 是拉普拉斯算子。

   方程左侧是惯性项（包含对流加速度 v·∇v），右侧依次是压力梯度、粘性力和重力。

2. 质量连续性方程：对于不可压缩流体，密度恒定，该方程简化为速度场的散度为零。
   [
   \nabla \cdot \mathbf{v} = 0
   ]
   这表示流体在运动过程中体积保持不变。

为了得到适用于粒子模拟的方程，我们沿粒子路径取物质导数，从而得到单个粒子 i 的运动方程：
[
\frac{d \mathbf{v}_i}{dt} = \mathbf{g} - \frac{1}{\rho} \nabla p + \frac{\mu}{\rho} \nabla^2 \mathbf{v}
]
我们的目标就是求解这个方程。

---

平滑粒子流体动力学 (SPH) 方法 🔬

上一节我们介绍了描述流体运动的方程，本节中我们来看看如何使用SPH方法数值求解这些方程。SPH是一种无网格的粒子方法，它将流体离散为一系列相互作用的粒子。

SPH的核心思想是使用平滑核函数 W 来估算空间任意点的场量（如密度、压力）。核函数定义了粒子间相互作用的范围和权重：距离越近，影响越大；超过相互作用半径 h，影响为零。

以下是SPH方法中对关键物理量的近似公式：

密度近似
粒子的密度通过对邻近粒子的质量进行加权求和来近似：
[
\rho_i \approx \sum_j m_j W(|\mathbf{r}_i - \mathbf{r}_j|, h)
]

压力梯度近似
压力梯度项近似为：
[
-\frac{1}{\rho_i} \nabla p \approx -\sum_j m_j \left( \frac{p_i}{\rho_i^2} + \frac{p_j}{\rho_j^2} \right) \nabla W(|\mathbf{r}_i - \mathbf{r}_j|, h)
]
其中压力 p 由状态方程给出：p = k(ρ - ρ₀)，ρ₀ 是静止密度，k 是常数。

粘性力近似
粘性力项近似为：
[
\frac{\mu}{\rho_i} \nabla^2 \mathbf{v} \approx \frac{\mu}{\rho_i} \sum_j m_j \frac{\mathbf{v}_j - \mathbf{v}_i}{\rho_j} \nabla^2 W(|\mathbf{r}_i - \mathbf{r}_j|, h)
]
这个项会使邻近粒子的速度趋于一致。

常用的平滑核函数
以下是程序中使用的具体核函数形式（r = |r_i - r_j|）：

• 标量核函数 W:
  [
  W(r, h) = \frac{315}{64\pi h^9} (h^2 - r²⁾3 \quad (0 \le r \le h)
  ]
• 梯度 ∇W (一个向量):
  [
  \nabla W(r, h) = \frac{45}{\pi h^6} (h - r)^2 \frac{\mathbf{r}_j - \mathbf{r}_i}{r}
  ]
• 拉普拉斯算子 ∇²W (一个标量):
  [
  \nabla^2 W(r, h) = \frac{45}{\pi h^6} (h - r)
  ]

---

SPH模拟算法步骤 ⚙️

综合以上公式，SPH模拟每一时间步的计算流程如下：

1. 计算密度和压力：对每个粒子 i，使用密度近似公式计算 ρ_i，进而计算 p_i。
2. 计算压力加速度：对每个粒子 i，使用压力梯度近似公式计算由压力产生的加速度。
3. 计算粘性加速度：对每个粒子 i，使用粘性力近似公式计算由粘性产生的加速度。
4. 计算总加速度：将重力、压力加速度和粘性加速度相加，得到粒子的总加速度 a_i。
5. 时间积分：使用显式欧拉法更新粒子的速度和位置。
   [
   \mathbf{v}_i^{new} = \mathbf{v}_i + \Delta t \cdot \mathbf{a}_i
   ]
   [
   \mathbf{r}_i^{new} = \mathbf{r}_i + \Delta t \cdot \mathbf{v}_i^{new}
   ]
6. 处理边界条件：确保粒子不会超出模拟区域（如盒子）。

一个朴素的实现需要计算所有粒子对之间的相互作用，复杂度为 O(N²)，效率低下。

---

高效邻居搜索与OpenCL实现 🖥️

为了提升性能，我们需要优化邻居搜索过程。基本思路是将空间划分为边长为 2h 的体素（voxel）。这样，任意粒子可能相互作用的邻居粒子，只可能位于其所在的体素及相邻的26个体素之内。

以下是优化后的算法步骤及其对应的OpenCL内核：

1. 体素化与排序

   • 内核: hash_particles：计算每个粒子所在的体素索引（哈希值），并将体素ID和粒子ID作为键值对存储。
   • 内核: sort 与 sort_postpass：对键值对（体素ID， 粒子ID）进行排序。排序后，属于同一体素的粒子在内存中连续排列。然后重排粒子的位置和速度数据以匹配新的顺序。

2. 构建空间索引

   • 内核: index 与 index_postpass：构建 grid_cell_index，这是一个数组，其下标是体素ID，内容是该体素内第一个粒子在排序后数组中的索引。这允许我们快速定位任意体素中的所有粒子。

3. 构建邻居图

   • 内核: find_neighbors：对于每个粒子，在其所在的 2x2x2 体素区域内搜索潜在邻居。为了在只采样部分邻居（如32个）时避免偏差，搜索起点是随机的，并且搜索方向交替进行。最终将选中的邻居粒子ID存入 neighbor_map。

4. 执行物理计算

   • 内核: compute_density_pressure：根据 neighbor_map，每个粒子仅与其32个邻居相互作用，计算密度和压力。
   • 内核: compute_acceleration：计算压力加速度和粘性加速度，并合成总加速度。
   • 内核: integrate：积分更新速度和位置，并实施边界碰撞。

程序中的缓冲区
主要的OpenCL缓冲区包括：

• pos / vel: 粒子的位置和速度 (float4类型，便于硬件访问)。
• sorted_pos / sorted_vel: 排序后的位置和速度。
• particle_index: 存储（体素ID， 粒子ID）键值对，用于排序。
• grid_cell_index: 体素到粒子索引的映射。
• neighbor_map: 每个粒子的邻居列表。

关于互操作性 (Interop)
理想情况下，应使用OpenCL-DirectX互操作，使得GPU上的计算缓冲区可以直接被图形API用于渲染，避免在主机和设备间复制数据的开销。虽然示例代码中包含了相关设置，但当前版本并未启用此优化。

---

总结与问答要点 📝

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

1. 流体基础：了解了流体的定义和支配其运动的力（压力、粘性、重力）。
2. 纳维-斯托克斯方程：学习了描述不可压缩流体运动的偏微分方程组及其物质导数形式。
3. SPH方法：掌握了使用平滑核函数将连续方程离散化为粒子相互作用形式的原理和关键近似公式。
4. OpenCL实现：剖析了一个高效的SPH模拟程序结构，包括空间划分、邻居搜索、多内核协作的完整流程。

常见问题解答：

• 密度是否恒定？ 在不可压缩流体中，整体密度恒定，但局部密度可以有微小变化，正是这种变化产生了压力波。
• 使用单精度还是双精度？ 本程序使用单精度(float)以追求性能。双精度(double)精度更高但速度更慢。
• 能模拟湍流吗？ SPH方法更适合模拟层流。模拟湍流通常需要更专门的技术和极高的分辨率。
• 粒子数量上限？ 受限于GPU显存。在AMD HD 5870级别的GPU上，约32K粒子可保证实时模拟。
• 为什么用缓冲区而非图像内存？ 开发时测试发现性能差异不大。未来OpenCL的缓存缓冲区特性可能是更好的选择。

鼓励大家下载并研究示例代码，进一步探索和修改。希望本教程对你有所帮助！

017：图像卷积优化

在本节课中，我们将以图像卷积算法为例，学习在GPU和CPU上使用OpenCL进行性能优化的多种技术。我们将从理解AMD GPU硬件架构开始，逐步探讨如何通过数据重用、缓存利用、循环展开等策略，将算法性能提升数十倍。

AMD GPU架构概述

上一节我们介绍了课程目标，本节中我们来看看优化工作的基础——理解底层硬件架构。这对于后续的优化至关重要。

AMD GPU（以Evergreen系列，如Radeon HD 5870为例）的核心是流处理器阵列。HD 5870拥有20个SIMD核心。每个SIMD核心是一个单指令多数据单元，意味着它在每个时钟周期执行一条指令，但该指令会同时作用于多个数据元素。

该GPU的峰值单精度浮点性能为2.72 TFLOPs，双精度性能为其五分之一（544 GFLOPs），全局内存带宽为153.6 GB/s。优化目标通常是达到这些峰值性能的70%-80%。

每个SIMD核心包含两种缓存：

• 本地数据共享：对应OpenCL中的 __local 内存（或CUDA中的共享内存），大小为32 KB。
• 纹理缓存：对应OpenCL中的图像对象，大小为8 KB。

OpenCL到硬件的映射

理解了硬件基础后，我们来看看OpenCL的抽象概念是如何映射到这些硬件单元上的。

以下是OpenCL内存层次结构与AMD GPU硬件的对应关系：

• 私有内存：映射到寄存器。
• 工作项：映射到流处理器中的单个ALU线程。
• 计算单元：映射到一个SIMD核心。
• 本地内存：映射到SIMD核心的32 KB本地数据共享。
• 全局内存：映射到板载的显存（如1GB或2GB）。
• 常量内存：映射到常量缓存（主要来自全局内存）。
• 图像：数据通过L2缓存，最终缓存在8 KB的纹理缓存中。

一个关键概念是波前。在AMD GPU上，一个波前包含64个工作项。硬件会以16个ALU为一组，分4个时钟周期执行同一条指令，从而完成这64个工作项的计算。因此，为了获得最佳性能，工作组大小应是64的倍数（如64、128、256）。如果一个工作组内的线程执行路径出现分歧（例如if-else分支），GPU将串行执行所有路径，可能导致性能下降。

卷积算法简介

在深入GPU优化之前，我们先简要回顾一下将要优化的算法——图像卷积。

卷积算法用于处理图像等信号。对于2D图像，算法使用一个滤波器（或掩膜）在输入图像上滑动。在每个输出像素位置，将滤波器权重与对应的输入图像像素值相乘并求和，结果即为该输出像素的值。

对于一个 N x N 的滤波器，计算每个输出像素需要读取 N x N 个输入像素和 N x N 个滤波器权重。算法核心是一个双重嵌套循环：

for (int i = -filterRadius; i <= filterRadius; i++) {
    for (int j = -filterRadius; j <= filterRadius; j++) {
        sum += inputImage[y+i][x+j] * filter[i+filterRadius][j+filterRadius];
    }
}
outputImage[y][x] = sum;

该算法的特点是数据重用性高：当滤波器移动一个像素时，大部分输入数据与上一个位置重叠。优化将重点利用这一特性。

CPU优化技术

尽管本课重点在GPU，但相同的OpenCL内核也可在CPU上运行并获得良好性能。以下是几种简单的CPU优化方法：

以下是三种关键的CPU优化策略：

1. 循环展开：手动或通过编译器指令减少循环控制开销。
2. 编译时常量：利用OpenCL内核在运行时编译的特性，通过 -D 选项传递滤波器大小等参数作为宏定义，使编译器能进行更积极的优化（如自动循环展开）。
3. 使用向量类型：使用 float4、int4 等类型。这能生成SSE/AVX指令，并结合OpenCL的自动多线程，性能可能超越传统的OpenMP实现。

实验表明，通过向量化等技术，OpenCL内核在CPU上的性能可以达到传统多线程OpenMP实现的两倍。

GPU优化实战

现在，我们进入核心部分，探讨如何在GPU上优化卷积算法。我们将以一个4K x 4K输出图像、16 x 16滤波器的案例进行测试。初始的朴素实现耗时约1511毫秒。

优化1：使用本地内存实现数据重用

朴素实现直接从全局内存读取数据，而全局内存延迟高且未缓存。第一个优化是利用本地内存手动缓存数据。

我们创建一个工作组（例如8x8，共64个线程）来处理一块输出图像。这块输出图像所需的输入图像区域比输出区域大（因为有滤波器重叠）。工作组的所有线程协作将这块所需的输入数据从全局内存加载到快速的本地内存中。然后，每个线程在计算自己的输出像素时，都从本地内存中读取数据。

效果：每个输出像素需要从全局内存读取的数据量从256次大幅降至约8.3次。执行时间从 1511毫秒 降至 359毫秒。

优化2：增大工作组尺寸

既然使用本地内存有效，那么增大工作组可以重用更多数据。

我们将工作组大小从8x8增加到16x16。这样，每个输出像素需要从全局内存读取的数据量进一步降至约3.7次。执行时间从 359毫秒 降至 182毫秒。

优化3：使用图像对象自动缓存

手动管理本地内存需要额外代码。OpenCL的图像对象提供了自动缓存机制。

将输入数据声明为 image2d_t 类型并使用采样器访问。硬件会自动通过纹理缓存缓存数据。对于8x8和16x16工作组，执行时间分别约为 346毫秒 和 207毫秒，与手动管理本地内存效果相近。

优化4：将滤波器放入常量内存

我们回过头优化滤波器数据的读取。滤波器数据较小且被所有线程重复读取。

将滤波器指针声明为 __constant。当多个线程访问常量内存中的同一地址时，访问会被优化。在朴素实现上应用此优化，时间从1511毫秒降至 1375毫秒。

优化5：循环展开与向量化加载

现在关注计算效率。原始的双重循环有控制流开销，且每次读取一个float。

我们将内层循环展开4次，并一次性读取4个输入值和4个滤波器值（使用float4）。这减少了循环开销，并且128位宽的内存访问通常比32位更高效。

效果：仅应用此优化，时间降至 401毫秒。如果滤波器也使用float4访问，时间进一步降至 389毫秒。

优化6：组合优化与陷阱

组合优化4和5（常量内存+向量化），时间反而增加至 680毫秒。原因是：将滤波器声明为 __constant float* 并访问float4时，编译器需要生成额外的ALU指令来计算所需的float在float4中的位置，增加了计算开销。

解决方案：直接将滤波器声明为 __constant float4* 数组。这样编译器可以直接进行向量化访存和计算，消除了额外的索引计算。应用此优化后，时间显著降至 346毫秒。

优化7：利用编译时常量进行终极优化

OpenCL允许在运行时编译内核时传递宏定义。我们可以将滤波器宽度等参数作为编译时常量传入。

这使得编译器能够进行彻底的静态优化，例如完全展开循环、预计算偏移量等。结合所有最佳实践（使用本地内存、大工作组、float4向量化、__constant float4*滤波器），最终执行时间从最初的1511毫秒降至惊人的 25毫秒（使用LDS）和 63毫秒（使用图像），实现了近60倍的性能提升。

总结与性能回顾

本节课中我们一起学习了针对图像卷积算法的多层次OpenCL优化技术。

以下是优化路径与性能提升的总结：

• 起点：朴素实现，1511毫秒。
• 数据重用：使用本地内存或图像缓存，降至~180-200毫秒。
• 内存访问优化：使用float4向量化读取和__constant float4*滤波器，降至~70-90毫秒。
• 编译器优化：利用编译时常量进行激进优化，最终达到25-63毫秒。

优化过程揭示了性能瓶颈的转移：最初是内存带宽受限，通过缓存优化解决后，可能变为ALU计算受限或指令调度受限，需要采取不同的优化策略。理解硬件架构（如波前、缓存层次）是进行有效优化的基础。