CUDA 6.0 统一内存寻址

统一寻址（Unified Memory）：

可直接访问CPU内存、GPU显存，不需要手动拷贝数据。

CUDA 6在现有的内存池结构上增加了一个统一内存系统，程序员可以直接访问任何内存/显存资源，或者在合法的内存空间内寻址，而不用管涉及到的到底是内存还是显存。

CUDA 6的数据拷贝由程序员的手动转移，变成自动执行，因此，它仍然受制于PCI-E的带宽和延迟。

NVIDIA的统一内存寻址

上面这段代码显示了“统一寻址”的一个例子。两段代码都从磁盘中载入文件，并对文件中数据进行排序操作。它们仅有的不同在于，GPU的版本启动的kernel函数（并在启动后进行了同步的操作），在统一寻址中使用新的API cudaMallocManaged（）对载入文件分配内存。

我们仅仅只分配一次内存，这个数据指针对于host端和device端都是可用的。我们可以直接从文件中读取数据到这片内存，并将指针直接传入CUDA核函数。

kernel函数执行完后，我们可以直接从CPU端获取数据。CUDA的运行环境隐藏了所有的复杂性，自动将数据转移到需要的地方。

CUDA 4.0开始就支持统一虚拟寻址Unified Virtual Addressing(UVA)了，Unified Memory依赖于UVA，但他们不同。UVA为系统中所有内存提供虚拟的单一的虚拟内存地址，不论是设备内存，主机内存或是片上共享内存。它允许cudaMemcpy的使用，不管输入和输出参数在哪。UVA能够使用“Zero-Copy”  memory， 一种pinned host memory，设备端能够通过PCI-Express直接获取，不需要memcpy。Zero-Copy提供了一些统一内存的便利性，但性能并不好，因为它总是和PCI-Express的低带宽和高延迟相关的。

例：减少深度复制

统一寻址的关键在于通过减少GPU kernel函数获取数据使深度复制的需求，简化异构计算内存模型。从CPU到GPU传递数据结构和指针需要深度复制。如下图：

以下面这个结构体dataElem为例

<pre name="code" class="plain">struct dataElem {
    int prop1;
    int prop2;
    char *name;
}

为了再设备上使用这个结构体，我们需要复制结构体本身和它的成员，然后复制结构体指针指向的所有数据，接着跟新这个结构体副本中的所有指针。这导致了以下复杂代码，这仅仅只是为kernel函数传递一个数据元素。

<pre name="code" class="plain">void launch(dataElem *elem) {
    dataElem *d_elem;
    char *d_name;

    int namelen = strlen(elem->name)+1;
   
    //Allocate storage for struct and name
    cudaMalloc(&d_elem, sizeof(dataElem));
    cudaMalloc(&d_name, namelen);

    //Copy up each piece separately, including new "name" pointer value
    cudaMemcpy(d_name, elem, sizeof(dataElem), cudaMemcpyHostToDevice);
    cudaMemcpy(d_name, elem-<name, namelen, cudaMemcpyHostToDevice);
    cudaMemcpy(&(d_elem->name), &d_name, sizeof(char*), cudaMemcpyHostToDevice);

    //Finally we can launch our kernel, but CPU & GPU use differernt copies of "elem"
    Kernel<<<...>>>(d_elem);
}

主机端需要大量额外的代码来处理CPU和GPU端的数据。使用统一寻址可大大简化这一过程：

<pre name="code" class="plain">void launch(dataElem *elem) {
    Kernel<<<...>>>(elem);
}

这仅仅是在降低代码复杂性上的进步，看另一个例子。

例：CPU/GPU共享链表

链表是非常常用的数据结构，但是由于它本质上是有指针组成的嵌套的数据结构，使得内存空间的传递非常复杂。没有统一寻址，CPU和GPU间的链表的共享是不易管理的。唯一的选择是在Zero-copy memory（pinned host memory）中分配内存，这意味着GPU端数据获取受制于PCI-express的性能。通过在Unified memory中分配链表数据，设备端可以在GPU上可以以设备端最好的性能正常的跟踪指针。这个程序可以保证一个单链表，链表中的元素可以从设备端或主机端增加或删除。

把代码中复杂的数据结构移植到GPU原本是件令人生畏的事，但是统一寻址使它变得简单。

C++统一寻址

统一寻址在C++数据结构中显得尤其闪耀。C++通过使用拷贝构造函数的类简化深度复制问题。拷贝构造函数创建一个类的对象，给它的成员分配空间，并把他们的值复制给另一个对象。C++也允许new和delete内存管理运算符的重载。这意味着我们可以创建一个基类，我们称之为Managed，在里面使用cudaMallocManaged（）重载新运算符，下面给出代码

<pre name="code" class="plain">class Managed {
public:
  void *operator new(size_t len) {
    void *ptr;
    cudaMallocManaged(&ptr, len);
    return ptr;
  }

  void operator delete(void *ptr) {
    cudaFree(ptr);
  }
};

接着，我们可以让String类继承Managed类，实现复制构造函数为我们的string副本分配unified memory。

<pre name="code" class="plain">//Deriving from "Managed" allows pass-by-reference"
class String : public Managed {
    int length;
    char *data;

public:
    //Unified mamory copy constructor allows pass-by-value
    String(const String &s) {
        length = s.length();
        cudaMallocManged(&data, length);
        memcpy(data, s.data, length);
    }

    //...
};

类似的，使dataElem类继承Managed

// Note “managed” on this class, too.
// C++ now handles our deep copies
class dataElem : public Managed {
public:
  int prop1;
  int prop2;
  String name;
};

</pre>通过这些盖面，C++类在unified memory中分配他们的内存，并自动处理深度复制。我们可以在unified memory中分配一个dataElem就像C++中的对象<p></p><p></p><pre name="code" class="plain">dataElem *data = new dataElem;

注意，你需要确保每一个类在其树中都继承于Managed，否则，你的内存映射就会有漏洞。如果你倾向于简化所有unified memory，你可以全局性的重载new和delete，但这仅在没有只存在CPU数据时有用，因为数据会有不必要的迁移。

现在我们想kernel函数传递对象时就有了选择；想普通的C++，我们可以传递值或传递引用。

// Pass-by-reference version
__global__ void kernel_by_ref(dataElem &data) { ... }

// Pass-by-value version
__global__ void kernel_by_val(dataElem data) { ... }

int main(void) {
  dataElem *data = new dataElem;
  ...
  // pass data to kernel by reference
  kernel_by_ref<<>>(*data);

  // pass data to kernel by value -- this will create a copy
  kernel_by_val<<>>(*data);
}