4. 统一寻址内存模型

在过去，主机（host）和设备（device）内存地址是相互独立的，无法直接访问对方的数据。这种特性给CPU程序的移植带来了困难，开发者需要时刻关注数据的可访问性和正确性。幸运的是，MXMACA引入了统一寻址内存（Unified Addressing Memory，UA）技术，它使得主机和设备之间的内存访问变得更加灵活。统一内存允许主机和设备共享同一内存地址空间，从而在保证程序运行效率和正确性的同时，大大降低了程序移植的复杂性。关于统一寻址内存API的信息和限制，参见《沐曦通用GPU MXMACA运行时API参考》。

../_images/image3.png — 图 4.1 统一寻址内存示意图

4.1. 使用方法

统一寻址内存是一种可被系统中任何处理器访问的单一内存地址空间。该技术使得应用程序能够分配一段可由CPU和GPU同时访问的内存空间。使用统一寻址内存，我们只需使用 mcMallocManaged 函数来替代传统的 mcMalloc 、 malloc 或 new 等内存分配函数。通过该函数分配的内存将返回一个可由任何处理器访问的指针，从而实现了内存的统一访问。

以下代码片段为统一寻址内存地址的使用示例：

void foo(int size)
{
    void *data;
    mcMallocManaged(&data, size);

    cpu_func1(data, size);

    gpu_func2<<<...>>>(data, size);
    mcDeviceSynchronize();

    cpu_func3(data, size);

    mcFree(data);
}

首先，我们使用 mcMallocManaged 分配内存。这将分配一个统一寻址内存空间，返回的指针data可以被CPU和GPU访问。然后，我们可以通过 host 函数 cpu_func1、cpu_func3 和 device 函数 gpu_func2 对申请的内存空间进行访问或操作。

这一技术能大大简化复杂结构体在host和device之间的数据搬运过程。

如以下代码片段所示， myStruct 结构体通过统一寻址内存的方式申请后，可以轻松的在GPU和CPU间来回使用。软件开发工程师无需再考虑如何在host和device间拷贝复杂的数据结构。

struct MyStruct
{
    char *myText;
    int myValue;
};

__global__ void Kernel(MyStruct *elem)
{
    printf("On device: name=%s, value=%d\n", elem->myText, elem->myValue);

    elem->myText[0] = 'd';
    elem->myValue++;
}

int main(void)
{
    MyStruct *e;
    mcMallocManaged((void **)&e, sizeof(MyStruct));

    e->myValue = 10;
    mcMallocManaged((void **)&(e->myText), sizeof(char) * (strlen("hello") + 1));
    strcpy(e->myText, "hello");

    Kernel<<<1, 1>>>(e);
    mcDeviceSynchronize();

    printf("On host: myText=%s, myValue=%d\n", e->myText, e->myValue);

    mcFree(e->myText);
    mcFree(e);

    mcDeviceReset();
}

4.2. 技术实现

使用 mcMallocManaged 申请统一寻址内存时，驱动会将申请的内存空间放置于GPU内存中，并将物理地址映射为虚拟地址返回给应用程序。当host需访问或修改这一内存地址时，CPU通过host端页表解析出物理地址，并通过PCIe直接获取或修改这一数据。

备注

host和device不能同时修改同一统一寻址内存地址，否则无法保证数据一致性。
host访问这一地址空间后，数据仍将驻留于GPU内存中。在GPU应用中，CPU对数据的访问并不频繁。这一特性将避免数据在host和device间频繁转移，能提高程序运行效率。

4.3. 使用示例（C++）

4.3.1. 示例1

在C++中构建和销毁类的实例时，常使用 new() 和 delete() 函数。可以通过全局重载这两个函数的方式，实现将程序中所有的实例均放置于统一内存空间中，例如：

void* operator new(size_t size)
{
    void *ptr;
    mcMallocManaged(&ptr, size);
    return ptr;
}
void operator delete(void* ptr)
{
    mcDeviceSynchronize();
    mcFree(ptr);
}

4.3.2. 示例2

也可以使用C++类的继承特性，将统一内存技术用于某些特定的类。以下示例中，MyClass类将继承Managed类，并使用其中的 new() 和 delete() 函数。因此MyClass类中的 data_ 将位于统一内存空间中。

class Managed
{
public:
    void *operator new(size_t len)
    {
        void *ptr;
        mcMallocManaged(&ptr, len);
        return ptr;
    }
    void operator delete(void *ptr)
    {
        mcDeviceSynchronize();
        mcFree(ptr);
    }
};
template <class T>
class MyClass : public Managed
{
    int size_;
    T *__restrict__ data_;

    ...

    //- Construct with given size
    explicit List(const int len)
    {
        this->size_ = len;
        this->data_ = nullptr;
        if (len > 0)
            this->data_ = new T[this->size_];
    }

    ...
}