CUDA编程之统一内存的介绍

我之前的介绍文章，“ 更容易介绍 CUDA C ++ ”介绍了 CUDA 编程的基本知识，它演示了如何编写一个简单的程序，在内存中分配两个可供 GPU 访问的数字数组，然后将它们加在 GPU 上。为此，我向您介绍了统一内存，这使得分配和访问系统中任何处理器上运行的代码都可以使用的数据变得非常容易， CPU 或 GPU 。

图 1 。统一内存是可从系统中的任何处理器访问的单个内存地址空间。

我以几个简单的“练习”结束了这篇文章，其中一个练习鼓励您运行最近基于 Pascal 的 GPU ，看看会发生什么。（我希望读者能尝试一下并对结果发表评论，你们中的一些人也这样做了！）。我建议这样做有两个原因。首先，因为 PascalMIG 如 NVIDIA Titan X 和 NVIDIA Tesla P100 是第一个包含页 GPUs 定额引擎的 GPUs ，它是统一内存页错误处理和 MIG 比率的硬件支持。第二个原因是它提供了一个很好的机会来学习更多的统一内存。

快 GPU ，快内存…对吗？

正确的！但让我们看看。首先，我将重新打印在两个 NVIDIA 开普勒 GPUs 上运行的结果（一个在我的笔记本电脑上，一个在服务器上）。

现在让我们尝试在一个非常快的 Tesla P100 加速器上运行，它基于 pascalgp100GPU 。

> nvprof ./add_grid ... Time(%) Time Calls Avg Min Max Name 100.00% 2.1192ms 1 2.1192ms 2.1192ms 2.1192ms add(int, float*, float*)

嗯，这低于 6gb / s ：比在我的笔记本电脑基于开普勒的 GeForceGPU 上运行慢。不过，别灰心，我们可以解决这个问题的。为了理解这一点，我将告诉你更多关于统一内存的信息。

下面是要添加的完整代码，以供参考_网格. cu 从上次开始。

#include  #include  // CUDA kernel to add elements of two arrays __global__ void add(int n, float *x, float *y) { int index = blockIdx.x * blockDim.x + threadIdx.x; int stride = blockDim.x * gridDim.x; for (int i = index; i < n; i += stride) y[i] = x[i] + y[i]; } int main(void) { int N = 1<<20; float *x, *y; // Allocate Unified Memory -- accessible from CPU or GPU cudaMallocManaged(&x, N*sizeof(float)); cudaMallocManaged(&y, N*sizeof(float)); // initialize x and y arrays on the host for (int i = 0; i < N; i++) { x[i] = 1.0f; y[i] = 2.0f; } // Launch kernel on 1M elements on the GPU int blockSize = 256; int numBlocks = (N + blockSize - 1) / blockSize; add<<>>(N, x, y); // Wait for GPU to finish before accessing on host cudaDeviceSynchronize(); // Check for errors (all values should be 3.0f) float maxError = 0.0f; for (int i = 0; i < N; i++) maxError = fmax(maxError, fabs(y[i]-3.0f)); std::cout << "Max error: " << maxError << std::endl; // Free memory cudaFree(x); cudaFree(y); return 0; }

对 27-19 行的内存进行初始化。

什么是统一内存？

统一内存是可从系统中的任何处理器访问的单个内存地址空间（请参见图 1 ）。这种硬件/软件技术允许应用程序分配可以从 CPU s 或 GPUs 上运行的代码读取或写入的数据。分配统一内存非常简单，只需将对malloc()或new的调用替换为对cudaMallocManaged()的调用，这是一个分配函数，返回可从任何处理器访问的指针（以下为ptr）。

cudaError_t cudaMallocManaged(void** ptr, size_t size);

当在 CPU 或 GPU 上运行的代码访问以这种方式分配的数据（通常称为 CUDA 管理 数据）， CUDA 系统软件和/或硬件负责将 MIG 额定内存页分配给访问处理器的内存。这里重要的一点是， PascalGPU 体系结构是第一个通过页面 MIG 比率引擎对虚拟内存页错误处理和 MIG 比率提供硬件支持的架构。基于更旧的 kezbr 架构和更为统一的 kezbr 形式的支持。

当我打电话给cudaMallocManaged()时，开普勒会发生什么？

在具有 pre-PascalGPUs 的系统上，如 Tesla K80 ，调用 cudaMallocManaged（） 会分配 size 字节的托管内存 在 GPU 设备上 ，该内存在调用 1 时处于活动状态。在内部，驱动程序还为分配覆盖的所有页面设置页表条目，以便系统知道这些页驻留在 GPU 上。

所以，在我们的例子中，在 Tesla K80GPU （开普勒架构）上运行， x 和 y 最初都完全驻留在 GPU 内存中。然后在第 6 行开始的循环中， CPU 逐步遍历两个数组，分别将它们的元素初始化为 1.0f 和 2.0f 。由于这些页最初驻留在设备存储器中，所以它写入的每个数组页的 CPU 上都会发生一个页错误， GPU 驱动程序 MIG 会将设备内存中的页面分配给 CPU 内存。循环之后，两个数组的所有页都驻留在 CPU 内存中。

在初始化 CPU 上的数据之后，程序启动 add（） 内核，将 x 的元素添加到 y 的元素中。

add<<<1, 256>>>(N, x, y);

在 pre-PascalGPUs 上，启动一个内核后， CUDA 运行时必须 MIG 将以前 MIG 额定为主机内存或另一个 GPU 的所有页面重新评级到运行内核 2 的设备内存。由于这些旧的 GPUs 不能出现分页错误，所有数据都必须驻留在 GPU 以防万一 上，内核访问它（即使它不会访问）。这意味着每次启动内核时都可能存在 MIG 定额开销。

当我在 K80 或 macbookpro 上运行程序时，就会发生这种情况。但是请注意，探查器显示的内核运行时间与 MIG 定额时间是分开的，因为 MIG 定额发生在内核运行之前。

==15638== Profiling application: ./add_grid ==15638== Profiling result: Time(%) Time Calls Avg Min Max Name 100.00% 93.471us 1 93.471us 93.471us 93.471us add(int, float*, float*) ==15638== Unified Memory profiling result: Device "Tesla K80 (0)" Count Avg Size Min Size Max Size Total Size Total Time Name 6 1.3333MB 896.00KB 2.0000MB 8.000000MB 1.154720ms Host To Device 102 120.47KB 4.0000KB 0.9961MB 12.00000MB 1.895040ms Device To Host Total CPU Page faults: 51

当我调用cudaMallocManaged()时， Pascal 上会发生什么？

在 Pascal 和更高版本的 GPUs 上， cudaMallocManaged（） 返回时可能不会物理分配托管内存；它只能在访问（或预取）时填充。换言之，在 GPU 或 CPU 访问页和页表项之前，可能无法创建它们。页面可以在任何时候对任何处理器的内存进行 cudaMemPrefetchAsync（） 速率，驱动程序使用启发式来维护数据的局部性并防止过多的页面错误 3 。（注意：应用程序可以使用 cudaMemAdvise（） 指导驱动程序，并使用 MIG 显式地 MIG 对内存进行速率调整，如 这篇博文描述了 ）。

与 pre-PascalGPUs 不同， Tesla P100 支持硬件页错误和 MIG 比率。所以在这种情况下，运行库在运行内核之前不会自动将 全部的 页面复制回 GPU 。内核在没有任何 MIG 定额开销的情况下启动，当它访问任何缺失的页时， GPU 会暂停访问线程的执行，页面 MIG 定额引擎 MIG 会在恢复线程之前对设备的页面进行评级。

这意味着当我在 Tesla P100 （ 2 。 1192ms ）上运行程序时， MIG 定额的成本包含在内核运行时中。在这个内核中，数组中的每一页都由 CPU 写入，然后由 GPU 上的 CUDA 内核访问，导致内核等待大量的页 MIG 配额。这就是为什么分析器在像 Tesla P100 这样的 PascalGPU 上测量的内核时间更长。让我们看看 P100 上程序的完整 nvprof 输出。

==19278== Profiling application: ./add_grid ==19278== Profiling result: Time(%) Time Calls Avg Min Max Name 100.00% 2.1192ms 1 2.1192ms 2.1192ms 2.1192ms add(int, float*, float*) ==19278== Unified Memory profiling result: Device "Tesla P100-PCIE-16GB (0)" Count Avg Size Min Size Max Size Total Size Total Time Name 146 56.109KB 4.0000KB 988.00KB 8.000000MB 860.5760us Host To Device 24 170.67KB 4.0000KB 0.9961MB 4.000000MB 339.5520us Device To Host 12 - - - - 1.067526ms GPU Page fault groups Total CPU Page faults: 36

如您所见，存在许多主机到设备页面错误，降低了 CUDA 内核的吞吐量。

我该怎么办？

在实际应用中， GPU 可能会在数据上执行更多的计算（可能多次），而不需要 CPU 来接触它。这个简单代码中的 MIG 定额开销是由于 CPU 初始化数据， GPU 只使用一次。有几种不同的方法可以消除或更改 MIG 比率开销，从而更准确地测量 vector add 内核的性能。

将数据初始化移动到另一个 CUDA 内核中的 GPU 。

多次运行内核，查看平均和最小运行时间。

在运行内核之前，将数据预取到 GPU 内存。

我们来看看这三种方法。

初始化内核中的数据

如果我们将初始化从 CPU 移到 GPU ，则add内核不会出现页面错误。这里有一个简单的 CUDA C ++内核来初始化数据。我们可以用启动这个内核来替换初始化x和y的主机代码。

__global__ void init(int n, float *x, float *y) { int index = threadIdx.x + blockIdx.x * blockDim.x; int stride = blockDim.x * gridDim.x; for (int i = index; i < n; i += stride) { x[i] = 1.0f; y[i] = 2.0f; } }

当我这样做时，我在 Tesla P100GPU 的配置文件中看到两个内核：

==44292== Profiling application: ./add_grid_init ==44292== Profiling result: Time(%) Time Calls Avg Min Max Name 98.06% 1.3018ms 1 1.3018ms 1.3018ms 1.3018ms init(int, float*, float*) 1.94% 25.792us 1 25.792us 25.792us 25.792us add(int, float*, float*) ==44292== Unified Memory profiling result: Device "Tesla P100-PCIE-16GB (0)" Count Avg Size Min Size Max Size Total Size Total Time Name 24 170.67KB 4.0000KB 0.9961MB 4.000000MB 344.2880us Device To Host 16 - - - - 551.9940us GPU Page fault groups Total CPU Page faults: 12

add内核现在运行得更快： 25 . 8us ，相当于接近 500gb / s 。

带宽=字节/秒=（ 3 * 4194304 字节* 1e-9 字节/ GB ）/ 25 . 8e-6s = 488 [UNK] GB / s

（要了解如何计算理论带宽和实现的带宽，请参阅这个帖子。）仍然存在设备到主机页错误，但这是由于在程序末尾检查 CPU 结果的循环造成的。

运行多次

另一种方法是只运行内核多次，并查看探查器中的平均时间。为此，我需要修改错误检查代码，以便正确报告结果。以下是在 Tesla P100 上 100 次运行内核的结果：

==48760== Profiling application: ./add_grid_many ==48760== Profiling result: Time(%) Time Calls Avg Min Max Name 100.00% 4.5526ms 100 45.526us 24.479us 2.0616ms add(int, float*, float*) ==48760== Unified Memory profiling result: Device "Tesla P100-PCIE-16GB (0)" Count Avg Size Min Size Max Size Total Size Total Time Name 174 47.080KB 4.0000KB 0.9844MB 8.000000MB 829.2480us Host To Device 24 170.67KB 4.0000KB 0.9961MB 4.000000MB 339.7760us Device To Host 14 - - - - 1.008684ms GPU Page fault groups Total CPU Page faults: 36

最短的内核运行时间只有 24 . 5 微秒，这意味着它可以获得超过 500GB / s 的内存带宽。我还包括了来自nvprof的统一内存分析输出，它显示了从主机到设备总共 8MB 的页面错误，对应于第一次运行add时通过页面错误复制到设备上的两个 4MB 数组（x和y）。

预取

第三种方法是在初始化后使用统一内存预取将数据移动到 GPU 。 CUDA 为此提供了cudaMemPrefetchAsync()。我可以在内核启动之前添加以下代码。

 // Prefetch the data to the GPU int device = -1; cudaGetDevice(&device); cudaMemPrefetchAsync(x, N*sizeof(float), device, NULL); cudaMemPrefetchAsync(y, N*sizeof(float), device, NULL); // Run kernel on 1M elements on the GPU int blockSize = 256; int numBlocks = (N + blockSize - 1) / blockSize; saxpy<<>>(N, 1.0f, x, y);

现在当我在 Tesla P100 上评测时，我得到以下输出。

==50360== Profiling application: ./add_grid_prefetch ==50360== Profiling result: Time(%) Time Calls Avg Min Max Name 100.00% 26.112us 1 26.112us 26.112us 26.112us add(int, float*, float*) ==50360== Unified Memory profiling result: Device "Tesla P100-PCIE-16GB (0)" Count Avg Size Min Size Max Size Total Size Total Time Name 4 2.0000MB 2.0000MB 2.0000MB 8.000000MB 689.0560us Host To Device 24 170.67KB 4.0000KB 0.9961MB 4.000000MB 346.5600us Device To Host Total CPU Page faults: 36

在这里，您可以看到内核只运行了一次，运行时间为 26 。 1us ，与前面显示的 100 次运行中最快的一次相似。您还可以看到，不再报告任何 GPU 页错误，主机到设备的传输显示为四个 2MB 的传输，这要归功于预取。

现在我们已经让它在 P100 上运行得很快，让我们将它添加到上次的结果表中。

关于并发性的注记

请记住，您的系统有多个处理器同时运行 CUDA 应用程序的部分：一个或多个 CPU 和一个或多个 GPUs 。即使在我们这个简单的例子中，也有一个 CPU 线程和一个 GPU 执行上下文，因此在访问任何一个处理器上的托管分配时都要小心，以确保没有竞争条件。

从计算能力低于 6 。 0 的 CPU 和 GPUs 同时访问托管内存是不可能的。这是因为 pre-Pascal GPUs 缺少硬件页面错误，所以不能保证一致性。在这些 GPUs 上，内核运行时从 CPU 访问将导致分段错误。

在 Pascal 和更高版本的 GPUs 上， CPU 和 GPU 可以同时访问托管内存，因为它们都可以处理页错误；但是，由应用程序开发人员来确保不存在由同时访问引起的争用条件。

在我们的简单示例中，我们在内核启动后调用了 cudaDeviceSynchronize（） 。这可以确保内核在 CPU 尝试从托管内存指针读取结果之前运行到完成。否则， CPU 可能会读取无效数据（在 Pascal 和更高版本上），或获得分段错误（在 pre-Pascal GPUs ）。

Pascal 及更高版本上统一内存的好处 GPUs

从 PascalGPU 体系结构开始，通过 49 位虚拟寻址和按需分页 GPU 比率，统一内存功能得到了显著改善。 49 位虚拟地址足以使 GPUs 访问整个系统内存加上系统中所有 GPUs 的内存。页面 MIG 比率引擎允许 GPU 线程在非驻留内存访问时出现故障，因此系统可以根据需要从系统中的任何位置对 MIG 的内存中的页面进行 MIG 分级，以实现高效处理。

允许使用统一内存 cudaMallocManaged（） 对统一内存进行分配。无论是在一个 GPU 上运行还是在多个 GPU 上运行，它都不会对应用程序进行任何修改。

另外， Pascal 和 VoltaGPUs 支持系统范围的原子内存操作。这意味着您可以对系统中任何地方的多个 GPUs 值进行原子操作。这对于编写高效的 multi-GPU 协作算法非常有用。

请求分页对于以稀疏模式访问数据的应用程序尤其有利。在某些应用程序中，不知道特定处理器将访问哪些特定内存地址。如果没有硬件页面错误，应用程序只能预加载整个阵列，或者承受设备外访问的高延迟成本（也称为“零拷贝”）。但是页面错误意味着只有内核访问的页面需要被 MIG 评级。

关于作者

Mark Harris 是 NVIDIA 杰出的工程师，致力于 RAPIDS 。 Mark 拥有超过 20 年的 GPUs 软件开发经验，从图形和游戏到基于物理的模拟，到并行算法和高性能计算。当他还是北卡罗来纳大学的博士生时，他意识到了一种新生的趋势，并为此创造了一个名字： GPGPU （图形处理单元上的通用计算）。

审核编辑：郭婷