我之前的介绍性文章“CUDA C++更简单的介绍”介绍了CUDA编程的基础，展示了如何编写一个简单的程序，将两个数组的数字分配给 GPU 可访问的内存，然后在 GPU 上把它们相加起来。为此，我向您介绍了统一内存，它使分配和访问数据变得非常容易，这些数据可以由系统中任何处理器(CPU或GPU)上运行的代码使用。
我用一些简单的“练习”完成了这篇文章，其中一个练习鼓励您在最近基于 Pascal 的GPU上运行，看看会发生什么。(我希望读者能尝试一下，并对结果发表评论，你们中的一些人就这么做了!)我提出这个建议有两个原因。首先，NVIDIA Titan X 和 NVIDIA Tesla P100 等 Pascal GPU 是第一个包含页面迁移引擎的 GPU ，该引擎支持统一内存页面故障和迁移。第二个原因是，它提供了一个很好的机会来学习更多关于统一内存的知识。
快速GPU，快速内存，对吧?
没错! 但让我们看看。首先，我将重新打印运行在两个 NVIDIA Kepler GPU上的结果(一个在我的笔记本电脑上，一个在服务器上)。
现在让我们试着在一个非常快的 Tesla P100 加速器上运行，其基于 Pascal GP100 GPU。
> 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: 比在我笔记本电脑的基于 Kepler 的 GeForce GPU 上运行要慢。不过，不要气馁;我们可以解决这个问题。为了理解这一点，我将告诉你更多关于统一内存的知识。
下面是 add_grid.cu 的完整代码
#include <iostream>
#include <math.h>
// 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<<<numBlocks, blockSize>>>(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;
}
分配和初始化内存的代码在第 19-27 行。
什么是统一内存？
统一内存是系统中任何处理器都可以访问的一个内存地址空间(参见图1)。这种硬件/软件技术允许应用程序分配可以从运行在 CPU 或 GPU 上的代码读取或写入的数据。分配统一内存就像调用 malloc() 或 new 一样简单，只需要调用 cudaMallocManaged() 即可，这个分配函数返回一个可以从任何处理器访问的指针(如下的 ptr )。
cudaError_t cudaMallocManaged(void** ptr, size_t size);当运行在 CPU 或 GPU 上的代码以这种方式访问分配的数据(通常称为 CUDA 托管数据)时，CUDA 系统软件和/或硬件负责将内存页迁移到正进行访问的处理器的内存中。这里的重点是 Pascal GPU 架构是第一个通过其页面迁移引擎支持虚拟内存页面故障和迁移的硬件架构。基于 Kepler 和 Maxwell 体系结构的老式 GPU 则支持更受限的统一内存形式。
当 Kepler GPU 调用了 cudaMallocManaged() 时发生了什么？
在使用 Pascal 之前的 GPU 的系统上，比如 Tesla K80，调用 cudaMallocManaged() 在 GPU 设备上分配 size 字节的托管内存，当调用发出时，该设备将处于活动状态。在内部，驱动程序还为分配所覆盖的所有页面设置页表条目，以便系统知道这些页面驻留在该 GPU 上。
因此，在我们的例子中，运行在 Tesla K80 GPU( Kepler 架构)上，x和y最初都完全驻留在 GPU 内存中。然后在从第 6 行开始的循环中，CPU 遍历两个数组，分别将它们的元素初始化为 1.0f 和 2.0f 。由于页面最初驻留在设备内存中，所以对于它写入的每个数组页面，CPU上都会发生页面错误，接着 GPU 驱动程序将页面从设备内存迁移到 CPU 内存。循环之后，两个数组的所有页面都驻留在 CPU 内存中。
在初始化 CPU 上的数据之后，程序启动 add() 内核，将 x 的元素添加到 y 的元素中。
add<<<1, 256>>>(N, x, y);在 Pascal 之前的 GPU 上，在启动内核时，CUDA 运行时必须将之前迁移到主机内存或另一个 GPU 的所有页面迁移回运行内核的设备的设备内存。这是由于这些旧的 GPU 不能分页错误，所以所有数据必须驻留在 GPU 上，以防内核访问它(即使它不会访问)。这意味着在每次内核启动时都有潜在的迁移开销。
这就是我在 K80 或 Macbook Pro 上运行程序时发生的情况。但是，请注意，分析器将内核运行时与迁移时间分开显示，因为迁移发生在内核运行之前。
==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
当 Pascal GPU 调用了 cudaMallocManaged() 时发生了什么？
在 Pascal 和后来的 GPU 上，当 cudaMallocManaged() 返回时，托管内存可能不会被物理分配;它可能只在访问(或预取)时填充。换句话说，在 GPU 或 CPU 访问页和页表条目之前，可能不会创建它们。页面可以在任何时候迁移到任何处理器的内存中，驱动程序使用启发式来维护数据局部性并防止过多的页面错误。(注意:应用程序可以使用 cudaMemAdvise() 指导驱动程序，并使用 cudaMemPrefetchAsync() 显式地迁移内存，如本文所述)。
与 Pascal 之前的 GPU 不同，Tesla P100 支持硬件页面故障和迁移。所以在这种情况下，运行时不会在运行内核之前自动将所有页面复制回 GPU。内核启动时没有任何迁移开销，当它访问任何缺少的页面时，GPU 会停止访问线程的执行，然后页面迁移引擎会在恢复线程之前将页面迁移到该设备上。
这意味着当我在 Tesla P100 (2.1192 ms)上运行程序时，迁移的成本包含在内核运行时中。在这个内核中，数组中的每个页面都是由 CPU 所写入的，然后由 GPU 上的 CUDA 内核访问，导致内核等待大量的页面迁移。这就是为什么在像 Tesla P100 这样的 Pascal GPU 上，分析器测量的内核时间更长。让我们看看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 不干预的情况下(可能多次)执行更多的数据计算。这个简单代码中的迁移开销是由 CPU 初始化数据而 GPU 只使用一次数据造成的。有几种不同的方法可以消除或更改迁移开销，从而更准确地度量向量 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 P100 GPU 的分析文件中看到了两个内核:
==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: 12add 内核现在运行得快得多:25.8us，约相当于 500 GB/s。下面是计算带宽的方法。
Bandwidth = Bytes / Seconds = (3 * 4,194,304 bytes * 1e-9 bytes/GB) / 25.8e-6s = 488 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 的统一内存分析输出，它显示了从主机到设备总共 8 MB 的页面错误，对应于第一次运行 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<<<numBlocks, blockSize>>>(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 页面错误，并且由于预取，设备传输的主机仅显示为 4 个 2MB 的传输。
现在我们已经在 P100 上运行得很快了，让我们将它添加到上次的结果表中。
并发性说明
请记住，您的系统有多个处理器同时运行 CUDA 应用程序的各个部分:一个或多个 CPU 和一个或多个 GPU 。即使在我们的简单示例中，也有一个 CPU 线程和一个 GPU 执行上下文。因此，在访问任何处理器上的托管分配时，我们都必须小心，以确保没有竞争条件。
不可能同时从 CPU 和计算能力低于6.0的 GPU 访问托管内存。这是因为 Pascal 之前的 GPU 缺少硬件页面故障，因此不能保证一致性。在这些 GPU 上，当内核运行时从 CPU 的访问将导致段错误。
在 Pascal 和后来的 GPU 上，CPU 和 GPU 可以同时访问托管内存，因为它们都可以处理页面错误;但是，由应用程序开发人员来确保不存在由同时访问引起的竞争条件。
在我们的简单示例中，我们在内核启动后调用 cudaDeviceSynchronize()。这确保在 CPU 试图从托管内存指针读取结果之前内核运行已经完成。否则，CPU 可能会读取无效数据(在 Pascal 或更高版本上)，或者获得段错误(在 Pascal 之前的 GPU 上)。
在 Pascal 和后来的 GPU 上使用统一内存的好处
从 Pascal GPU 架构开始，通过 49 位虚拟寻址和随需应变的页面迁移，统一内存功能得到了显著改进。49 位虚拟地址足以使 GPU 访问整个系统内存以及系统中所有 GPU 的内存。页面迁移引擎允许 GPU 线程在非驻留内存访问上出错，因此系统可以根据需要将页面从系统中的任何位置迁移到 GPU 的内存中，以实现高效的处理。
换句话说，统一内存透明地启用了对 GPU 内存的过度订阅，允许对任何使用统一内存进行分配的代码(例如cudaMallocManaged())进行核外计算。它“只是工作”，没有任何修改的应用程序，无论是运行在一个 GPU 或多个 GPU 上。
另外，Pascal 和 Volta GPU 支持全系统的原子内存操作。这意味着您可以从多个 GPU 对系统中的任何地方的值进行原子操作。这对于编写高效的多 GPU 协作算法非常有用。
对于使用稀疏模式访问数据的应用程序，请求分页尤其有用。在某些应用程序中，不知道特定处理器将访问哪些特定内存地址。没有硬件页面故障的话，应用程序只能预加载整个数组，或者承受高延迟的设备外访问的代价(也称为“零拷贝”)。但是页面错误意味着只需要迁移内那些核访问的页面。
何去何从？
我希望这篇文章能够帮助您继续学习 CUDA 编程，并希望您有兴趣学习更多的知识，并在自己的计算中应用CUDA C++。如果您有问题或评论，请使用下面的评论部分联系我们。
有关统一内存预取和使用提示(cudaMemAdvise())的更多信息，请参见本文在 Pascal 上使用统一内存超越GPU内存限制。如果您想了解 CUDA中 使用 cudaMalloc 和 cudaMemcpy 的显式内存管理，请参阅旧文章CUDA C/ C++的简单介绍。
我们计划用更多的CUDA编程材料来继续这篇文章，但是为了让您现在保持忙碌，您可以继续阅读一系列更老的介绍性文章。
- 如何在CUDA C++中实现性能指标
- 如何在CUDA C++中查询设备属性和处理错误
- 如何在CUDA C++中优化数据传输
- 如何在CUDA C++中重叠数据传输
- 如何在CUDA C++中有效地访问全局内存
- 在CUDA C++中使用共享内存
- 一个有效的矩阵转置在CUDA C++
- CUDA C++中的有限差分方法，第1部分
- CUDA C++中的有限差分方法，第2部分
注解：
- 从技术上讲，这是一种简化。在具有 Pascal 之前的 GPU 的多 GPU 系统上，如果某些 GPU 禁用了对等(Peer 2 Peer访问，那么内存将被分配，因此它最初驻留在 CPU 上。
- 严格地说，您可以使用 cudaStreamAttachMemAsync() 来限制分配到特定 CUDA 流的可见性。这允许驱动程序只迁移附加到内核启动时所在流的页面。默认情况下，托管分配附加到所有流，因此任何内核启动都会触发迁移。阅读更多CUDA编程指南。
- 设备属性 concurrentManagedAccess 告诉 GPU 是否支持硬件页面迁移及其支持的并发访问功能。值 1 表示支持。目前，它只支持 Pascal 和运行在 64 位 Linux 上的更新的 GPU。
统一内存和 NVLink 是 CUDA® 程序员的强大组合。统一内存为您提供一个指向数据的指针，并在 CPU 和 GPU 之间自动迁移数据。在拥有 NVLink 连接的 CPU 和 GPU 的机器上，使用 80GB/s 或更高的带宽，这意味着 GPU 内核将能够以与CPU 相同的带宽访问主机系统内存中的数据——比 PCIe 快得多。应用程序的主机和设备部分将能够更高效地共享数据，并在共享数据结构上进行协作操作，支持更大的问题规模将比以往任何时候都更容易。
668
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
