编程模型：

选择管理内存：大多数平台都需要一个程序来选择进行自动数据管理，方法是使用__managed__关键字注释__device__变量或使用新的cudaMallocManaged（）调用来分配数据。计算能力低于6.x的设备必须始终在堆上分配托管内存，无论是使用分配器还是声明全局存储。 将先前分配的内存与统一内存相关联或使统一内存系统管理CPU或GPU堆栈指针是不可能的。从CUDA 8.0开始，支持具有计算能力6.x的设备的系统，可以使用相同的指针从GPU代码和CPU代码访问分配有默认操作系统分配程序（例如malloc或new）的内存。 在这些系统上，统一内存是默认设置：不需要使用特殊的分配器或创建专门管理的内存池。使用cudaMallocManaged（）显式分配：统一内存通常是使用语义和语法上类似于标准CUDA分配器cudaMalloc（）的分配函数创建的。 功能描述如下：

cudaError_t cudaMallocManaged(void **devPtr,    size_t size,    unsigned int flags = 0);

cudaMallocManaged（）函数在GPU上分配托管内存的大小字节，并在devPtr中返回一个指针。 指针在所有GPU和系统中的CPU上都是有效的，尽管程序访问这个指针必须遵守统一内存编程模型的并发规则。 下面是一个简单的例子，展示了cudaMallocManaged（）的使用：

__global__ void printme(char *str) {    printf(str);}int main() {    // Allocate 100 bytes of memory, accessible to both Host and Device code    char *s;    cudaMallocManaged(&s, 100);    // Note direct Host-code use of "s"    strncpy(s, "Hello Unified Memory\n", 99);    // Here we pass "s" to a kernel without explicitly copying    printme << < 1, 1 >> >(s);    cudaDeviceSynchronize();    // Free as for normal CUDA allocations    cudaFree(s);    return 0;}

当cudaMalloc（）被cudaMallocManaged（）替换时，程序的行为在功能上是不变的; 但是，该程序应该继续消除显式内存副本并利用自动迁移。 此外，可以消除双指针（一个指向主机，一个指向设备内存）。

设备代码无法调用cudaMallocManaged（）。 所有托管内存必须从主机或全局范围分配（请参阅下一节）。 在内核中使用malloc（）分配设备堆将不会在托管内存空间中创建，因此CPU代码无法访问。全局范围管理变量使用__managed__：文件范围和全局范围的CUDA __device__变量也可以通过向声明添加新的__managed__注释来选择加入Unified Memory管理。 这些可以直接从主机或设备代码中引用，如下所示：

__device__ __managed__ int x[2];__device__ __managed__ int y;__global__ void kernel() {    x[1] = x[0] + y;}int main() {    x[0] = 3;    y = 5;    kernel << < 1, 1 >> >();    cudaDeviceSynchronize();    printf("result = %d\n", x[1]);    return 0;}

原始__device__内存空间的所有语义以及一些额外的统一内存特定约束都由托管变量继承。 有关详细信息，请参阅CUDA C编程指南中的Compilation with NVCC）。

请注意，标记为__constant__的变量也可能不会被标记为__managed__; 此注释仅保留给__device__变量。 常量内存必须在编译时静态设置，或像往常一样在CUDA中使用cudaMemcpyToSymbol（）