本 节 书 摘 来 自 华 章 出 版 社 《CUDA高性能并行计算》 一 书 中 的 第3章，第3.4节， 作 者 CUDA for Engineers: An Introduction to High-Performance Parallel Computing[美] 杜安·斯托尔蒂（Duane Storti）梅特·尤尔托卢（Mete Yurtoglu） 著，苏统华 项文成 李松泽 姚宇鹏 孙博文 译 ， 更 多 章 节 内 容 可 以 访 问 云 栖 社 区 “华 章 计 算 机” 公 众 号 查 看。
3.4 简化操作流程
上面所述的标准操作流是主流的工作方式，然而其中的部分过于死板和烦琐，因此一些NVIDIA专家一起努力提供了一个可替代的流式方案，叫作统一内存（unified memory）。这个方法打破了主机内存和设备内存的围墙，因此你可以只用一个可以从主机端和设备端共同访问的数组（至少看起来是这样的）。
3.4.1 统一内存和托管数组
有一个好消息是统一内存把你从创建一个数组的不同拷贝（在主机和设备中）和显式的CPU和GPU间传输数据中解放出来。你可以建立一个设备和主机都能访问的托管数组来代替以前的工作。实际上，数组中的数据仍然需要在主机和设备间传输，但是CUDA系统会调度并执行这些操作，而不需要你去管理。注意托管内存数组会带来一个开发效率和执行效率的折衷。让系统去管理数据传输可能会简化并加速你的应用开发进程，但是自动管理的数据传输不能保证会像你手工管理得一样好。在开发过程中，你可能会发现编码的瓶颈在数据传输部分，这时决定进行显式的分配内存和传输数据是明智的。关于这个话题的更多内容我们会在后面详细讨论性能时提及，但现在我们只关心最基本的统一内存知识并实现一个托管内存版本的距离计算应用。
关于统一内存的一些重要事项：
统一内存对系统有特殊的要求。包括计算能力3.0以上的GPU和64位的Linux或者Windows系统[1]（目前OS X系统不支持统一内存）。
另一个可能的巨大收益是对于那些喜欢面向对象C++编程的读者，统一内存更适合于处理传输结构数据并有助于避免深拷贝（deep copy）问题[2]。
目前统一内存访问还是从软件层面进行模拟，当GPU转成硬件层面实现时，托管数组的性能损耗可能会降低。
3.4.2 使用cudaMallocManaged()实现的距离应用
当我们并行化dist_v2来创建dist_v2_cuda时，我们遵循了主流的范式：创建一些额外的数组并显式地使用cudaMemcpy()函数在主机和设备间通过PCIe总线进行数据的输入和输出。
这里我们来看一下流式实现的细节。它可以使得一个数组在设备和主机上都能访问统一内存。CUDA的C语言拓展中最关键的部分是函数cudaMallocManaged()，我们将使用这个函数来为托管内存分配空间（在定义了指针以后，与cudaMalloc()函数使用方法类似）。
我们再次简化讨论的内容，直接浏览一个使用统一内存优化的距离应用的代码，代码参见代码清单3.7，其中有一些修改之处值得注意：
distanceArray()函数的代码量被大大减少了，因为我们不需要再分配内存并显式地进行数据传输，所有必须做的事情就是启动一个计算用的核函数（并且调用cudaDeviceSynchronize()保证在返回之前计算已经完成）。
只在一个地方创建了所有相关的数组。在main()函数中，定义in和out指针。
使用cudaMallocManaged()函数为数组分配统一内存，然后在计算完成后使用cudaFree()函数来释放内存。
为了简化介绍，我们选择将所有的代码放在一个文件中。这一次我们创建一个.cu文件，而不是像第1章中那样使用不包含CUDA语言拓展的只包含纯C代码的.cpp文件。
注意每个函数的修饰符：scale()在CPU执行的for循环内调用，来生成输入数据。合适的修饰符是__host__但是我们可以不添加它，因为它是默认的标识符。
distance()函数将从核函数中调用并在GPU上执行，因此需要添加__device__标识。
distanceKernel()如所有的核函数一样，是从主机上调用而在设备上执行，因此需要使用修饰符__global__，并且设置返回类型是void。
又到了需要编译并运行应用的时候，我们可以检查运算结果是否正确。（Linux下使用的Makefile文件参见代码清单3.2，也可以用来构建这个项目。）我们在第22行中包含一个printf()语句来在控制台上打印结果，你也可以通过debug工具来观察结果。
在Linux下，也可以使用cuda-gbd命令来观察托管数组中的值，就像观察任何一个变量的值一样（使用print out [0]@64）。
使用Windows下的Visual Studio，你可以在Locals窗口中看到，但是在那里只能查看一个值。你可以使用如下的步骤查看其他的值：
1.在第47行设置一个断点，正好处于释放内存之前。
2.开始调试（或按F5功能键）。
3.选择DEBUGQuickWatch（或按下Shift+F9组合键）。
4.当QuickWatch窗口打开时，输入out, 64。
5.点击名称旁边的三角形按钮来打开out数组中的成员列表。
关于CUDA专用的调试工具，请参见附录D。
