4.1. 统一内存

4.1.1.1.1. 以文件为后备的统一内存

由于完全支持 CUDA 统一内存的系统允许设备访问主机进程所拥有的任何内存，因此设备可以直接访问以文件为后备的内存。

这里展示上一节中最初示例的一个修改版本，使用以文件为后备的内存，从 GPU 打印直接从输入文件读取的字符串。在以下示例中，内存以物理文件为后备，但该示例也适用于以内存为后备的文件。

__global__ void kernel(const char* type, const char* data) {
  static const int n_char = 8;
  printf("%s - first %d characters: '", type, n_char);
  for (int i = 0; i < n_char; ++i) printf("%c", data[i]);
  printf("'\n");
}

void test_file_backed() {
  int fd = open(INPUT_FILE_NAME, O_RDONLY);
  ASSERT(fd >= 0, "Invalid file handle");
  struct stat file_stat;
  int status = fstat(fd, &file_stat);
  ASSERT(status >= 0, "Invalid file stats");
  char* mapped = (char*)mmap(0, file_stat.st_size, PROT_READ, MAP_PRIVATE, fd, 0);
  ASSERT(mapped != MAP_FAILED, "Cannot map file into memory");
  kernel<<<1, 1>>>("file-backed", mapped);
  ASSERT(cudaDeviceSynchronize() == cudaSuccess,
    "CUDA failed with '%s'", cudaGetErrorString(cudaGetLastError()));
  ASSERT(munmap(mapped, file_stat.st_size) == 0, "Cannot unmap file");
  ASSERT(close(fd) == 0, "Cannot close file");
}

请注意，不具备hostNativeAtomicSupported属性的系统（请参见主机原生原子操作），包括启用了 Linux HMM 的系统，不支持对以文件为后备的内存进行原子访问。

4.1.1.1.2. 使用统一内存进行进程间通信（IPC）

许多应用程序倾向于每个进程管理一个 GPU，但仍然需要使用统一内存，例如用于超额订用，并从多个 GPU 访问它。

CUDA IPC（请参见进程间通信）不支持托管内存：不能通过本节讨论的任何机制共享此类内存的句柄。在完全支持 CUDA 统一内存的系统上，系统分配的内存能够用于 IPC。一旦与其他进程共享了对系统分配的内存的访问权限，同样的编程模型便适用，类似于以文件为后备的统一内存。

有关在 Linux 下创建可用于 IPC 的系统分配内存的各种方式的更多信息，请参见以下参考资料：

• 使用 MAP_SHARED 的 mmap

• POSIX IPC API

• Linux memfd_create。

请注意，无法使用这种技术在不同主机及其设备之间共享内存。

4.1.1.2.1. 内存分页与页大小

要更好地理解统一内存对性能的影响，了解虚拟寻址、内存页和页大小很重要。本小节尝试定义所有必要的术语，并解释为什么分页对性能很重要。

目前支持统一内存的所有系统都使用虚拟地址空间：这意味着应用程序使用的内存地址表示一个虚拟位置，该位置可能会被映射到内存实际驻留的物理位置。

目前支持的所有处理器，包括 CPU 和 GPU，还使用内存分页。由于所有系统都使用虚拟地址空间，因此存在两种类型的内存页：

• 虚拟页：这是操作系统按进程跟踪的一块固定大小的连续虚拟内存，可以将其映射到物理内存中。请注意，虚拟页与映射相关联：例如，可以使用不同的页大小将同一个虚拟地址映射到物理内存中。

• 物理页：这是处理器的主内存管理单元（MMU）支持的一块固定大小的连续内存，虚拟页可以映射到其中。

目前，所有 x86_64 CPU 使用的默认物理页大小为 4KiB。Arm CPU 支持多种物理页大小，包括 4KiB、16KiB、32KiB 和 64KiB，具体取决于 CPU 型号。最后，NVIDIA GPU 支持多种物理页大小，但更倾向于使用 2MiB 或更大的物理页。请注意，这些大小在未来的硬件中可能会改变。

虚拟页的默认页大小通常与物理页大小相对应，但只要操作系统和硬件支持，应用程序就可以使用不同的页大小。通常，受支持的虚拟页大小必须是 2 的幂，并且是物理页大小的倍数。

跟踪虚拟页到物理页映射的逻辑实体称为页表，而将具有给定虚拟大小的某个虚拟页映射到物理页的每个映射称为页表项（PTE）。所有受支持的处理器都为页表提供专用缓存，以加速虚拟地址到物理地址的转换。这些缓存称为转换后备缓冲区（TLB）。

应用程序性能调优有两个重要方面：

• 虚拟页大小的选择，

• 系统是提供一张由 CPU 和 GPU 共用的合并页表，还是为每个 CPU 和 GPU 分别提供独立页表。

4.1.1.2.2. 从主机直接访问统一内存

某些设备在硬件上支持主机对驻留于 GPU 的统一内存进行相干的读取、存储和原子访问。这些设备的属性cudaDevAttrDirectManagedMemAccessFromHost设置为 1。请注意，所有硬件相干系统都会为通过 NVLink 连接的设备设置此属性。在这些系统上，主机可以直接访问驻留于 GPU 的内存，而不会发生缺页异常和数据迁移。请注意，对于 CUDA 托管内存，必须使用位置类型为cudaMemLocationTypeHost的cudaMemAdviseSetAccessedBy提示，才能启用这种不发生缺页异常的直接访问，请参见下面的示例。

__global__ void write(int *ret, int a, int b) {
  ret[threadIdx.x] = a + b + threadIdx.x;
}

__global__ void append(int *ret, int a, int b) {
  ret[threadIdx.x] += a + b + threadIdx.x;
}

void test_malloc() {
  int *ret = (int*)malloc(1000 * sizeof(int));
  // for shared page table systems, the following hint is not necesary
  cudaMemLocation location = {.type = cudaMemLocationTypeHost};
  cudaMemAdvise(ret, 1000 * sizeof(int), cudaMemAdviseSetAccessedBy, location);

  write<<< 1, 1000 >>>(ret, 10, 100);            // pages populated in GPU memory
  cudaDeviceSynchronize();
  for(int i = 0; i < 1000; i++)
      printf("%d: A+B = %d\n", i, ret[i]);        // directManagedMemAccessFromHost=1: CPU accesses GPU memory directly without migrations
                                                  // directManagedMemAccessFromHost=0: CPU faults and triggers device-to-host migrations
  append<<< 1, 1000 >>>(ret, 10, 100);            // directManagedMemAccessFromHost=1: GPU accesses GPU memory without migrations
  cudaDeviceSynchronize();                        // directManagedMemAccessFromHost=0: GPU faults and triggers host-to-device migrations
  free(ret);
}

__global__ void write(int *ret, int a, int b) {
  ret[threadIdx.x] = a + b + threadIdx.x;
}

__global__ void append(int *ret, int a, int b) {
  ret[threadIdx.x] += a + b + threadIdx.x;
}

void test_managed() {
  int *ret;
  cudaMallocManaged(&ret, 1000 * sizeof(int));
  cudaMemLocation location = {.type = cudaMemLocationTypeHost};
  cudaMemAdvise(ret, 1000 * sizeof(int), cudaMemAdviseSetAccessedBy, location);  // set direct access hint

  write<<< 1, 1000 >>>(ret, 10, 100);            // pages populated in GPU memory
  cudaDeviceSynchronize();
  for(int i = 0; i < 1000; i++)
      printf("%d: A+B = %d\n", i, ret[i]);        // directManagedMemAccessFromHost=1: CPU accesses GPU memory directly without migrations
                                                  // directManagedMemAccessFromHost=0: CPU faults and triggers device-to-host migrations
  append<<< 1, 1000 >>>(ret, 10, 100);            // directManagedMemAccessFromHost=1: GPU accesses GPU memory without migrations
  cudaDeviceSynchronize();                        // directManagedMemAccessFromHost=0: GPU faults and triggers host-to-device migrations
  cudaFree(ret); 

write核函数完成后，ret将在 GPU 内存中创建并初始化。接下来，CPU 将访问ret，随后append核函数将再次使用同一块ret内存。根据系统架构和对硬件相干性的支持情况，此代码将表现出不同的行为：

• 在具有directManagedMemAccessFromHost=1的设备上：CPU 对托管缓冲区的访问不会触发任何迁移；数据将保持驻留于 GPU 内存中，后续的任何 GPU 核函数都可以继续直接访问它，而不会引发缺页异常或迁移

• 在具有directManagedMemAccessFromHost=0的设备上：CPU 对托管缓冲区的访问将发生缺页异常并启动数据迁移；任何尝试首次访问同一数据的 GPU 核函数都会发生缺页异常，并将页迁移回 GPU 内存。

4.1.1.2.3. 主机原生原子操作

某些设备（包括硬件相干系统中通过 NVLink 连接的设备）支持对驻留于 CPU 的内存进行硬件加速的原子访问。这意味着，对主机内存的原子访问不必通过缺页异常来模拟。对于这些设备，属性cudaDevAttrHostNativeAtomicSupported设置为 1。

4.1.1.2.4. 原子访问和同步原语

CUDA 统一内存支持主机线程和设备线程可用的所有原子操作，使所有线程都能通过并发访问同一共享内存位置来协作。libcu++库提供了许多针对主机线程与设备线程之间的并发使用而调优的异构同步原语，其中包括cuda::atomic、cuda::atomic_ref、cuda::barrier、cuda::semaphore等许多原语。

在软件相干系统上，不支持从设备对文件后备的主机内存进行原子访问。以下示例代码在硬件相干系统上有效，但在其他系统上会出现未定义行为：

#include <cuda/atomic>

#include <cstdio>
#include <fcntl.h>
#include <sys/mman.h>

#define ERR(msg, ...) { fprintf(stderr, msg, ##__VA_ARGS__); return EXIT_FAILURE; }

__global__ void kernel(int* ptr) {
  cuda::atomic_ref{*ptr}.store(2);
}

int main() {
  // this will be closed/deleted by default on exit
  FILE* tmp_file = tmpfile64();
  // need to allocate space in the file, we do this with posix_fallocate here
  int status = posix_fallocate(fileno(tmp_file), 0, 4096);
  if (status != 0) ERR("Failed to allocate space in temp file\n");
  int* ptr = (int*)mmap(NULL, 4096, PROT_READ | PROT_WRITE, MAP_PRIVATE, fileno(tmp_file), 0);
  if (ptr == MAP_FAILED) ERR("Failed to map temp file\n");

  // initialize the value in our file-backed memory
  *ptr = 1;
  printf("Atom value: %d\n", *ptr);

  // device and host thread access ptr concurrently, using cuda::atomic_ref
  kernel<<<1, 1>>>(ptr);
  while (cuda::atomic_ref{*ptr}.load() != 2);
  // this will always be 2
  printf("Atom value: %d\n", *ptr);

  return EXIT_SUCCESS;
}

在软件相干系统上，对统一内存的原子访问可能引发缺页异常，从而导致显著的延迟。请注意，在这些系统上，并非所有对 CPU 内存的 GPU 原子操作都会出现这种情况：nvidia-smi -q | grep "Atomic Caps Outbound"列出的操作可能避免缺页异常。

在硬件相干系统上，主机与设备之间的原子操作不需要缺页异常，但仍可能由于会导致任何内存访问发生缺页异常的其他原因而发生缺页异常。

4.1.1.2.5. 使用统一内存时 Memcpy()/Memset() 的行为

cudaMemcpy*()和cudaMemset*()接受任何统一内存指针作为参数。

对于cudaMemcpy*()，以cudaMemcpyKind指定的方向是一种性能提示；如果任一参数是统一内存指针，该提示可能对性能产生更大的影响。

因此，建议遵循以下性能建议：

• 当统一内存的物理位置已知时，使用准确的cudaMemcpyKind提示。

• 优先使用cudaMemcpyDefault，而不是不准确的cudaMemcpyKind提示。

• 始终使用已填充（已初始化）的缓冲区：避免使用这些 API 来初始化内存。

• 如果两个指针都指向系统分配的内存，应避免使用cudaMemcpy*()：改为启动核函数或使用std::memcpy等 CPU 内存复制算法。

4.1.1.2.6. 统一内存的内存分配器概述

对于完全支持 CUDA 统一内存的系统，可以使用各种不同的分配器来分配统一内存。下表概述了部分分配器及其各自的特性。请注意，本节中的所有信息都可能在未来的 CUDA 版本中发生变化。

表不同分配器的统一内存支持概述展示了几种分配器在语义上的差异，可以考虑使用这些分配器来分配可由多个处理器（包括主机和设备）同时访问的数据。有关cudaMemPoolCreate的更多详细信息，请参见内存池一节；有关cuMemCreate的更多详细信息，请参见虚拟内存管理一节。

在设备内存作为 NUMA 域向系统公开的硬件相干系统上，可以使用numa_alloc_on_node等特殊分配器将内存固定到给定的 NUMA 节点，该节点可以是主机或设备。主机和设备都可以访问此内存，并且此内存不会迁移。同样，mbind可用于将内存固定到给定的一个或多个 NUMA 节点，并且可以使文件后备内存在首次被访问之前就放置到给定的一个或多个 NUMA 节点上。

以下内容适用于所分配内存用于共享的分配器：

• mmap等系统分配器允许使用MAP_SHARED标志在进程之间共享内存。CUDA 支持此功能，可用于在连接到同一主机的不同设备之间共享内存。但是，目前不支持用它在多个主机以及多个设备之间共享内存。有关详细信息，请参见使用统一内存进行进程间通信（IPC）。

• 若要在多个主机上通过网络访问统一内存或其他 CUDA 内存，请查阅所用通信库的文档，例如NCCL、NVSHMEM、OpenMPI、UCX等。

4.1.1.2.7. 局部性域和统一内存

不支持将托管内存分配局部化。局部性域不是统一内存的迁移目的地；cudaMemAllocationTypeManaged不能与局部化分配结合使用。需要在局部性域中放置内存的应用程序应改用局部化的页锁定内存池或局部化的 VMM 分配。

有关整体编程模型，请参见局部性域。

4.1.1.2.8. 访问计数器迁移

在硬件相干系统上，访问计数器功能会跟踪 GPU 对位于其他处理器上的内存的访问频率。这是确保将内存页移动到最频繁访问这些页的处理器的物理内存中所必需的。它可以指导 CPU 与 GPU 之间以及对等 GPU 之间的迁移，这一过程称为访问计数器迁移。

从 CUDA 12.4 开始，系统分配的内存支持访问计数器。请注意，文件后备内存不会根据访问情况迁移。对于系统分配的内存，可以通过使用cudaMemAdviseSetAccessedBy提示，以相应的设备 id 指定设备，来启用访问计数器迁移。如果访问计数器已启用，可以使用设置为主机的cudaMemAdviseSetPreferredLocation来防止迁移。默认情况下，cudaMallocManaged基于缺页异常后迁移的机制进行迁移。[7]

驱动程序也可能使用访问计数器，以便在抖动缓解或内存超额订用场景中提高效率。

4.1.1.2.9. 避免从 CPU 频繁写入驻留于 GPU 的内存

如果主机访问统一内存，缓存未命中可能会在主机与设备之间产生比预期更多的流量。许多 CPU 架构要求所有内存操作（包括写入）都经过缓存层次结构。如果系统内存驻留于 GPU，这意味着 CPU 对该内存的频繁写入可能导致缓存未命中，从而在将实际值写入所请求的内存范围之前，先将数据从 GPU 传输到 CPU。在软件相干系统上，这可能引入额外的缺页异常，而在硬件相干系统上，这可能导致 CPU 操作之间的延迟更高。因此，为了与设备共享主机生成的数据，可以考虑写入驻留于 CPU 的内存，并由设备直接读取这些值。下面的代码展示了如何使用统一内存实现这一点。

  size_t data_size = sizeof(int);
  int* data = (int*)malloc(data_size);
  // ensure that data stays local to the host and avoid faults
  cudaMemLocation location = {.type = cudaMemLocationTypeHost};
  cudaMemAdvise(data, data_size, cudaMemAdviseSetPreferredLocation, location);
  cudaMemAdvise(data, data_size, cudaMemAdviseSetAccessedBy, location);

  // frequent exchanges of small data: if the CPU writes to CPU-resident memory,
  // and GPU directly accesses that data, we can avoid the CPU caches re-loading
  // data if it was evicted in between writes
  for (int i = 0; i < 10; ++i) {
    *data = 42 + i;
    kernel<<<1, 1>>>(data);
    cudaDeviceSynchronize();
    // CPU cache potentially evicted data here
  }
  free(data);

  int* data;
  size_t data_size = sizeof(int);
  cudaMallocManaged(&data, data_size);
  // ensure that data stays local to the host and avoid faults
  cudaMemLocation location = {.type = cudaMemLocationTypeHost};
  cudaMemAdvise(data, data_size, cudaMemAdviseSetPreferredLocation, location);
  cudaMemAdvise(data, data_size, cudaMemAdviseSetAccessedBy, location);

  // frequent exchanges of small data: if the CPU writes to CPU-resident memory,
  // and GPU directly accesses that data, we can avoid the CPU caches re-loading
  // data if it was evicted in between writes
  for (int i = 0; i < 10; ++i) {
    *data = 42 + i;
    kernel<<<1, 1>>>(data);
    cudaDeviceSynchronize();
    // CPU cache potentially evicted data here
  }
  cudaFree(data);

4.1.1.2.10. 利用对系统内存的异步访问

如果应用程序需要与主机共享设备上工作的结果，有几种可行的选项：

1. 设备将其结果写入驻留于 GPU 的内存，使用cudaMemcpy*传输该结果，然后主机读取传输过来的数据。

2. 设备将其结果直接写入驻留于 CPU 的内存，然后主机读取该数据。

3. 设备写入驻留于 GPU 的内存，然后主机直接访问该数据。

如果在向主机传输结果或主机访问结果的同时，可以在设备上调度独立的工作，则优先选择方案 1 或 3。如果在主机访问结果之前，设备一直没有工作可做，则方案 2 可能更合适。这是因为，除非使用许多主机线程来读取数据，否则设备写入的带宽通常高于主机读取的带宽。

void exchange_explicit_copy(cudaStream_t stream) {
  int* data, *host_data;
  size_t n_bytes = sizeof(int) * 16;
  // allocate receiving buffer
  host_data = (int*)malloc(n_bytes);
  // allocate, since we touch on the device first, will be GPU-resident
  cudaMallocManaged(&data, n_bytes);
  kernel<<<1, 16, 0, stream>>>(data);
  // launch independent work on the device
  // other_kernel<<<1024, 256, 0, stream>>>(other_data, ...);
  // transfer to host
  cudaMemcpyAsync(host_data, data, n_bytes, cudaMemcpyDeviceToHost, stream);
  // sync stream to ensure data has been transferred
  cudaStreamSynchronize(stream);
  // read transferred data
  printf("Got values %d - %d from GPU\n", host_data[0], host_data[15]);
  cudaFree(data);
  free(host_data);
}

void exchange_device_direct_write(cudaStream_t stream) {
  int* data;
  size_t n_bytes = sizeof(int) * 16;
  // allocate receiving buffer
  cudaMallocManaged(&data, n_bytes);
  // ensure that data is mapped and resident on the host
  cudaMemLocation location = {.type = cudaMemLocationTypeHost};
  cudaMemAdvise(data, n_bytes, cudaMemAdviseSetPreferredLocation, location);
  cudaMemAdvise(data, n_bytes, cudaMemAdviseSetAccessedBy, location);
  kernel<<<1, 16, 0, stream>>>(data);
  // sync stream to ensure data has been transferred
  cudaStreamSynchronize(stream);
  // read transferred data
  printf("Got values %d - %d from GPU\n", data[0], data[15]);
  cudaFree(data);
}

void exchange_host_direct_read(cudaStream_t stream) {
  int* data;
  size_t n_bytes = sizeof(int) * 16;
  // allocate receiving buffer
  cudaMallocManaged(&data, n_bytes);
  // ensure that data is mapped and resident on the device
  cudaMemLocation device_loc = {};
  cudaGetDevice(&device_loc.id);
  device_loc.type = cudaMemLocationTypeDevice;
  cudaMemAdvise(data, n_bytes, cudaMemAdviseSetPreferredLocation, device_loc);
  cudaMemAdvise(data, n_bytes, cudaMemAdviseSetAccessedBy, device_loc);
  kernel<<<1, 16, 0, stream>>>(data);
  // launch independent work on the GPU
  // other_kernel<<<1024, 256, 0, stream>>>(other_data, ...);
  // sync stream to ensure data may be accessed (has been written by device)
  cudaStreamSynchronize(stream);
  // read data directly from host
  printf("Got values %d - %d from GPU\n", data[0], data[15]);
  cudaFree(data);

最后，在上面的显式复制示例中，可以使用主机或设备核函数来显式执行此传输，而不是使用cudaMemcpy*来传输数据。对于连续数据，优先使用 CUDA 复制引擎，因为复制引擎执行的操作可以与主机和设备上的工作重叠。cudaMemcpy*和cudaMemPrefetchAsyncAPI 可能会使用复制引擎，但不能保证cudaMemcpy*API 调用会使用复制引擎。出于同样的原因，对于足够大的数据，显式复制优于主机直接读取：如果主机和设备所执行的工作都没有使其各自的内存系统达到饱和，复制引擎就可以在主机和设备执行工作的同时并发执行传输。

复制引擎通常既用于主机与设备之间的传输，也用于通过 NVLink 连接的系统中对等设备之间的传输。由于复制引擎的总数有限，某些系统中cudaMemcpy*的带宽可能低于使用设备显式执行传输时的带宽。在这种情况下，如果传输位于应用程序的关键路径上，则可能更适合采用基于设备的显式传输。

4.1.3.3.1. 流回调

CPU 可以在流回调内部合法地访问托管数据，前提是 GPU 上没有其他可能访问托管数据的流处于活动状态。此外，后面没有任何设备工作的回调可以用于同步：例如，在回调内部向条件变量发出信号；否则，CPU 访问仅在回调执行期间有效。有几个重要的注意事项：

1. 当 GPU 处于活动状态时，始终允许 CPU 访问非托管映射内存中的数据。

2. 当 GPU 正在运行任何核函数时，即使该核函数不使用托管数据，也会将 GPU 视为处于活动状态。如果核函数可能使用数据，则禁止访问

3. 对于 GPU 之间对托管内存的并发访问，除了适用于多 GPU 访问非托管内存的约束之外，没有其他约束。

4. 对于并发执行的 GPU 核函数访问托管数据，没有任何约束。

请注意，最后一点允许 GPU 核函数之间出现竞争，这与目前非托管 GPU 内存的情况相同。从 GPU 的角度看，托管内存的工作方式与非托管内存完全相同。以下代码示例说明了这些要点：

int main() {
    cudaStream_t stream1, stream2;
    cudaStreamCreate(&stream1);
    cudaStreamCreate(&stream2);
    int *non_managed, *managed, *also_managed;
    cudaMallocHost(&non_managed, 4);    // Non-managed, CPU-accessible memory
    cudaMallocManaged(&managed, 4);
    cudaMallocManaged(&also_managed, 4);
    // Point 1: CPU can access non-managed data.
    kernel<<< 1, 1, 0, stream1 >>>(managed);
    *non_managed = 1;
    // Point 2: CPU cannot access any managed data while GPU is busy,
    //          unless concurrentManagedAccess = 1
    // Note we have not yet synchronized, so "kernel" is still active.
    *also_managed = 2;      // Will issue segmentation fault
    // Point 3: Concurrent GPU kernels can access the same data.
    kernel<<< 1, 1, 0, stream2 >>>(managed);
    // Point 4: Multi-GPU concurrent access is also permitted.
    cudaSetDevice(1);
    kernel<<< 1, 1 >>>(managed);
    return  0;
}

4.1.3.3.2. 与流关联的托管内存允许更细粒度的控制

统一内存建立在流独立性模型之上，允许 CUDA 程序显式地将托管分配与 CUDA 流关联。这样，程序员便可根据核函数是否启动到指定流中来指示核函数对数据的使用情况。这使得可以根据程序特定的数据访问模式来实现并发。用于控制此行为的函数是：

cudaError_t cudaStreamAttachMemAsync(cudaStream_t stream,
                                     void *ptr,
                                     size_t length=0,
                                     unsigned int flags=0);

cudaStreamAttachMemAsync()函数将从 ptr 开始的 length 字节内存与指定流关联。这样，只要 stream 中的所有操作都已完成，就允许 CPU 访问该内存区域，而不论其他流是否处于活动状态。实际上，这将活动 GPU 对托管内存区域的独占所有权约束在每个流的活动范围内，而不是整个 GPU 的活动范围内。最重要的是，如果某个分配未与特定流关联，那么它对所有正在运行的核函数都可见，而不论这些核函数位于哪个流中。这是cudaMallocManaged()分配或__managed__变量的默认可见性；因此，在简单情况下，规则是只要有任何核函数正在运行，CPU 就不能访问这些数据。

以下示例展示了如何显式地将y与主机可访问性关联，从而使 CPU 在任何时候都能够访问。（请注意，核函数调用之后没有cudaDeviceSynchronize()。）现在，在 GPU 上运行的核函数对y的访问将产生未定义的结果。

__device__ __managed__ int x, y=2;
__global__  void  kernel() {
    x = 10;
}
int main() {
    cudaStream_t stream1;
    cudaStreamCreate(&stream1);
    cudaStreamAttachMemAsync(stream1, &y, 0, cudaMemAttachHost);
    cudaDeviceSynchronize();          // Wait for Host attachment to occur.
    kernel<<< 1, 1, 0, stream1 >>>(); // Note: Launches into stream1.
    y = 20;                           // Success – a kernel is running but “y”
                                      // has been associated with no stream.
    return  0;
}

4.1.3.3.3. 关于多线程主机程序的一个更详细的示例

cudaStreamAttachMemAsync()的主要用途是利用 CPU 线程实现独立任务的并行执行。在这样的程序中，CPU 线程通常会为其生成的所有工作创建自己的流，因为使用 CUDA 的 NULL 流会导致线程之间产生依赖关系。托管数据默认对任何 GPU 流都具有全局可见性，这可能使得多线程程序难以避免 CPU 线程之间的相互影响。因此，使用函数cudaStreamAttachMemAsync()将线程的托管分配与该线程自己的流关联，而且这种关联通常在线程的整个生命周期内都不会改变。这样的程序只需增加一次对cudaStreamAttachMemAsync()的调用，就可以使用统一内存进行数据访问：

// This function performs some task, in its own , in its own private stream and can be run in parallel
void run_task(int *in, int *out, int length) {
    // Create a stream for us to use.
    cudaStream_t stream;
    cudaStreamCreate(&stream);
    // Allocate some managed data and associate with our stream.
    // Note the use of the host-attach flag to cudaMallocManaged();
    // we then associate the allocation with our stream so that
    // our GPU kernel launches can access it.
    int *data;
    cudaMallocManaged((void **)&data, length, cudaMemAttachHost);
    cudaStreamAttachMemAsync(stream, data);
    cudaStreamSynchronize(stream);
    // Iterate on the data in some way, using both Host & Device.
    for(int i=0; i<N; i++) {
        transform<<< 100, 256, 0, stream >>>(in, data, length);
        cudaStreamSynchronize(stream);
        host_process(data, length);    // CPU uses managed data.
        convert<<< 100, 256, 0, stream >>>(out, data, length);
    }
    cudaStreamSynchronize(stream);
    cudaStreamDestroy(stream);
    cudaFree(data);
}

在此示例中，分配与流之间的关联只建立一次，然后主机和设备都会反复使用这些数据。与在主机和设备之间显式复制数据相比，得到的代码要简单得多，尽管结果相同。

函数cudaMallocManaged()指定了 cudaMemAttachHost 标志，这会创建一个最初对设备端执行不可见的分配。（默认分配对所有流上的所有 GPU 核函数都可见。）这可确保在分配数据之后、为特定流获取数据之前的这段时间内，不会与另一个线程的执行发生意外的相互影响。

如果没有此标志，而恰好有另一个线程启动的核函数正在运行，那么新分配将被视为正在由 GPU 使用。这可能影响该线程在能够将新分配的数据显式附加到私有流之前，从 CPU 访问这些数据的能力。因此，为了让线程之间能够安全地保持独立，应在分配时指定此标志。

另一种方法是在将分配附加到流之后，在所有线程之间设置一个进程范围的屏障。这将确保所有线程在启动任何核函数之前都完成其数据与流的关联，从而避免这一风险。在销毁流之前还需要第二个屏障，因为销毁流会使分配恢复其默认可见性。提供cudaMemAttachHost标志既是为了简化这一过程，也是因为并非总能在需要的位置插入全局屏障。

4.1.3.3.4. 与流关联的统一内存的数据移动

在未设置concurrentManagedAccess的设备上，Memcpy()/Memset() 对与流关联的统一内存的行为有所不同，适用以下规则：

如果指定了cudaMemcpyHostTo*且源数据是统一内存，那么，如果主机在复制流中能够以相干方式访问该数据，就会从主机访问该数据(1)；否则，将从设备访问该数据。当指定了cudaMemcpy*ToHost且目标是统一内存时，类似的规则适用于目标。

如果指定了cudaMemcpyDeviceTo*且源数据是统一内存，那么将从设备访问该数据。设备必须能够在复制流中以相干方式访问源数据(2)；否则，将返回错误。当指定了cudaMemcpy*ToDevice且目标是统一内存时，类似的规则适用于目标。

如果指定了cudaMemcpyDefault，那么，当无法在复制流中从设备以相干方式访问统一内存时(2)，或者当数据的首选位置为cudaCpuDeviceId且能够在复制流中从主机以相干方式访问它时(1)，将从主机访问统一内存；否则，将从设备访问它。

对统一内存使用cudaMemset*()时，必须能够在用于cudaMemset()操作的流中从设备以相干方式访问该数据(2)；否则，将返回错误。

通过cudaMemcpy*或cudaMemset*从设备访问数据时，该操作所在的流被视为在 GPU 上处于活动状态。在此期间，如果 GPU 的设备属性concurrentManagedAccess的值为零，则 CPU 对与该流关联的数据或具有全局可见性的数据的任何访问都会导致段错误。程序必须进行适当的同步，以确保在从 CPU 访问任何关联数据之前，该操作已完成。

1. 能够在给定流中从主机以相干方式访问，意味着该内存既不具有全局可见性，也不与给定流关联。

2. 能够在给定流中从设备以相干方式访问，意味着该内存具有全局可见性或与给定流关联。