CUDA从入门到放弃（十）：统一寻址编程 Unified Memory Programming

CUDA从入门到放弃（十）：统一寻址编程 Unified Memory Programming

1 统一寻址 General description of unified memory

统一寻址为所有处理器（CPU和GPU等）提供了一个统一的内存池，使得它们能够使用各自的内存操作访问这块内存。这一特性有助于简化GPU编程，提升编程效率和程序性能。统一寻址减少了数据复制，优化了数据访问速度，并允许GPU处理超出其物理内存容量的数据。

获取CUDA统一寻址主要有两种方法：

• 系统分配的内存：使用系统API在主机上分配的内存，包括栈变量、全局/文件作用域变量、malloc()/mmap()、线程局部变量等。
• 通过CUDA API显式分配统一寻址：例如，使用cudaMallocManaged()等函数分配的内存，这种方式适用于更多的系统，并且其性能可能优于系统分配的内存。

在编程过程中，程序员可以关注应用的并行化，而数据的移动和迁移则可以作为性能优化的一部分来考虑。统一寻址的使用使得数据的物理位置对程序透明，从而简化了内存管理的复杂性。

1-1 系统要求 System Requirements

在不支持统一寻址的系统上，应用程序的行为是未定义的。CUDA应用程序通过以下属性检查系统对统一寻址的支持程度，确保跨不同支持级别的系统可移植性：

• pageableMemoryAccess：在支持CUDA统一寻址且所有线程可访问系统分配内存和CUDA托管内存的系统上，此属性设为1。这些系统包括NVIDIA Grace Hopper、IBM Power9 + Volta，以及启用HMM的现代Linux系统等。
• concurrentManagedAccess：在具有完整CUDA托管内存支持的系统上，此属性设为1。若设为0，则CUDA托管内存对统一寻址的支持有限。

1-2 编程模型 Programming Model

使用CUDA统一寻址，无需在主机和设备间进行单独分配和显式传输。程序可以通过以下方式来分配统一寻址：

• 系统分配API：在具有完整CUDA统一寻址支持的系统上，可以通过主机进程的任何系统分配（如C的malloc()函数、C++的new运算符、POSIX的mmap等）来进行。
• CUDA托管内存分配API：通过cudaMallocManaged() API进行分配，其语法与cudaMalloc()相似。
• 声明__managed__变量：使用__managed__关键字声明的变量，其语义与__device__变量相似。

No unified memory

__global__ void write_value(int* ptr, int v) {
  *ptr = v;
}

int main() {
  int* d_ptr = nullptr;
  // Does not require any unified memory support
  cudaMalloc(&d_ptr, sizeof(int));
  write_value<<<1, 1>>>(d_ptr, 1);
  int host;
  // Copy memory back to the host and synchronize
  cudaMemcpy(&host, d_ptr, sizeof(int),
             cudaMemcpyDefault);
  printf("value = %d\n", host); 
  cudaFree(d_ptr); 
  return 0;
}

System-Allocated Memory（malloc）

__global__ void write_value(int* ptr, int v) {
  *ptr = v;
}

int main() {
  // Requires System-Allocated Memory support
  int* ptr = (int*)malloc(sizeof(int));
  write_value<<<1, 1>>>(ptr, 1);
  // Synchronize required
  // (before, cudaMemcpy was synchronizing)
  cudaDeviceSynchronize();
  printf("value = %d\n", *ptr); 
  free(ptr); 
  return 0;
}

System-Allocated Memory（Stack）

__global__ void write_value(int* ptr, int v) {
  *ptr = v;
}

int main() {
  // Requires System-Allocated Memory support
  int value;
  write_value<<<1, 1>>>(&value, 1);
  // Synchronize required
  // (before, cudaMemcpy was synchronizing)
  cudaDeviceSynchronize();
  printf("value = %d\n", value);
  return 0;
}

Managed Memory（cudaMallocManaged）

__global__ void write_value(int* ptr, int v) {
  *ptr = v;
}

int main() {
  int* ptr = nullptr;
  // Requires CUDA Managed Memory support
  cudaMallocManaged(&ptr, sizeof(int));
  write_value<<<1, 1>>>(ptr, 1);
  // Synchronize required
  // (before, cudaMemcpy was synchronizing)
  cudaDeviceSynchronize();
  printf("value = %d\n", *ptr); 
  cudaFree(ptr); 
  return 0;
}

Managed Memory（__managed__ ）

__global__ void write_value(int* ptr, int v) {
  *ptr = v;
}

// Requires CUDA Managed Memory support
__managed__ int value;

int main() {
  write_value<<<1, 1>>>(&value, 1);
  // Synchronize required
  // (before, cudaMemcpy was synchronizing)
  cudaDeviceSynchronize();
  printf("value = %d\n", value);
  return 0;
}

这些示例展示了在GPU上合并两个数字并返回结果数组的过程。不使用统一寻址时，需分别为主机和设备分配内存，并进行显式复制。而使用统一寻址后，GPU可直接访问主机内存，无需额外分配和复制，从而简化程序并减小其大小。此外，当使用系统分配或托管内存时，相应分配方式也需作调整。

1-2-1 系统分配内存的分配API

在支持CUDA统一寻址的系统上，所有内存都是统一的，包括通过系统分配API（如malloc()、mmap()、C++的new()运算符）分配的内存以及CPU上的自动变量、线程局部变量、全局变量等。这些内存可能在首次访问时才会分配物理内存。物理内存的选择取决于首次访问的线程类型（CPU或GPU）。

CUDA提供了统一寻址提示和预取API，如cudaMemAdvise和cudaMemPrefetchAsync，可用于系统分配的内存。

__global__ void printme(char *str) {  
  printf(str);  
}  
  
int main() {  
  char *s = (char*)malloc(100);  
  strncpy(s, "Hello Unified Memory\n", 99);  
  printme<<<1, 1>>>(s);  
  cudaDeviceSynchronize();  
  cudaFree(s);  
  return 0;  
}

1-2-2 CUDA托管内存的分配API：cudaMallocManaged()

在支持CUDA托管内存的系统上，可以使用

host cudaError_t cudaMallocManaged(void **devPtr, size_t size);

分配统一寻址，该内存可以被CPU和GPU并发访问。与cudaMalloc()相似，但分配的内存可以在主机和设备间自动迁移。使用完毕后，通过cudaFree()释放内存。

示例，使用cudaMallocManaged()分配内存，并在GPU内核中直接访问它：

__global__ void printme(char *str) {  
  printf(str);  
}  
  
int main() {  
  char *s;  
  cudaMallocManaged(&s, 100);  
  strncpy(s, "Hello Unified Memory\n", 99);  
  printme<<<1, 1>>>(s);  
  cudaDeviceSynchronize();  
  cudaFree(s);  
  return 0;  
}

1-2-3 使用__managed__的全局作用域托管变量

__managed__变量无需手动分配或复制，它简化了程序，并确保了CPU和GPU都可以直接访问这些变量。同时，为了__managed__变量的正确操作，需要一个有效的CUDA上下文。如果当前设备没有上下文，访问__managed__变量会触发上下文创建。

在完全支持CUDA统一寻址的系统上，文件作用域或全局作用域的变量不能被设备代码直接访问。但是，指向这些变量的指针可以作为参数传递给内核。

System Allocator示例：

__global__ void write_value(int* ptr, int v) {
  *ptr = v;
}

int main() {
  // Requires System-Allocated Memory support
  int value;
  write_value<<<1, 1>>>(&value, 1);
  // Synchronize required
  // (before, cudaMemcpy was synchronizing)
  cudaDeviceSynchronize();
  printf("value = %d\n", value);
  return 0;
}

Managed示例：

__global__ void write_value(int* ptr, int v) {
  *ptr = v;
}

// Requires CUDA Managed Memory support
__managed__ int value;

int main() {
  write_value<<<1, 1>>>(&value, 1);
  // Synchronize required
  // (before, cudaMemcpy was synchronizing)
  cudaDeviceSynchronize();
  printf("value = %d\n", value);
  return 0;
}

CUDA的__managed__变量隐含了__device__，并等价于__managed__ device，这也是允许的。标记为__constant__的变量不能标记为__managed__。

1-2-4 统一寻址与映射内存的区别

统一寻址确保所有类型的内存访问在所有系统上均受支持，而CUDA映射内存则不保证这一点。尽管如此，CUDA映射内存支持的可移植内存操作集比统一寻址更广泛，可应用于更多系统。

1-2-5 指针属性 Pointer Attributes

CUDA程序可以通过调用cudaPointerGetAttributes()并检查指针属性值是否为cudaMemoryTypeManaged来判断一个指针是否指向CUDA托管内存分配。该API返回cudaMemoryTypeHost表示通过cudaHostRegister()注册的系统分配内存，返回cudaMemoryTypeUnregistered表示CUDA未知的系统分配内存。
指针属性并不表示内存驻留在何处，而是表示内存是如何分配或注册的。

示例：定义了一个kind函数，用于根据指针属性判断内存类型，并在main函数中创建不同类型的指针，然后调用check_pointer函数来检查每个指针所指向的内存类型。

char const* kind(cudaPointerAttributes a, bool pma, bool cma) {
    switch(a.type) {
    case cudaMemoryTypeHost: return pma?
      "Unified: CUDA Host or Registered Memory" :
      "Not Unified: CUDA Host or Registered Memory";
    case cudaMemoryTypeDevice: return "Not Unified: CUDA Device Memory";
    case cudaMemoryTypeManaged: return cma?
      "Unified: CUDA Managed Memory" : "Not Unified: CUDA Managed Memory";
    case cudaMemoryTypeUnregistered: return pma?
      "Unified: System-Allocated Memory" :
      "Not Unified: System-Allocated Memory";
    default: return "unknown";
    }
}

void check_pointer(int i, void* ptr) {
  cudaPointerAttributes attr;
  cudaPointerGetAttributes(&attr, ptr);
  int pma = 0, cma = 0, device = 0;
  cudaGetDevice(&device);
  cudaDeviceGetAttribute(&pma, cudaDevAttrPageableMemoryAccess, device);
  cudaDeviceGetAttribute(&cma, cudaDevAttrConcurrentManagedAccess, device);
  printf("Pointer %d: memory is %s\n", i, kind(attr, pma, cma));
}

__managed__ int managed_var = 5;

int main() {
  int* ptr[5];
  ptr[0] = (int*)malloc(sizeof(int));
  cudaMallocManaged(&ptr[1], sizeof(int));
  cudaMallocHost(&ptr[2], sizeof(int));
  cudaMalloc(&ptr[3], sizeof(int));
  ptr[4] = &managed_var;

  for (int i = 0; i < 5; ++i) check_pointer(i, ptr[i]);
  
  cudaFree(ptr[3]);
  cudaFreeHost(ptr[2]);
  cudaFree(ptr[1]);
  free(ptr[0]);
  return 0;
}

1-2-6 运行时检测统一寻址支持级别 Runtime detection of Unified Memory Support Level

以下示例展示了如何在运行时检测统一寻址的支持级别：

int main() {
  int d;
  cudaGetDevice(&d);

  int pma = 0;
  cudaDeviceGetAttribute(&pma, cudaDevAttrPageableMemoryAccess, d);
  printf("Full Unified Memory Support: %s\n", pma == 1? "YES" : "NO");
  
  int cma = 0;
  cudaDeviceGetAttribute(&cma, cudaDevAttrConcurrentManagedAccess, d);
  printf("CUDA Managed Memory with full support: %s\n", cma == 1? "YES" : "NO");

  return 0;
}

1-2-7 GPU内存超额预定 GPU Memory Oversubscription

统一寻址允许应用分配超过单一处理器内存大小的数组，实现跨多个处理器的数据处理。

1-2-8 性能提示 Performance Hints

统一寻址性能提示用于向CUDA提供额外信息，以优化内存访问性能。这些提示不影响应用逻辑，仅影响性能，可按需使用。

1-2-8-1 数据预取 Data Prefetching

cudaMemPrefetchAsync API 是一个异步的、按流排序的 API，它可以将数据迁移到更接近指定处理器的位置。在数据预取的过程中，仍然可以访问这些数据。预取操作不会开始，直到流中先前的所有操作都已完成，并且会在流中的任何后续操作之前完成。

cudaError_t cudaMemPrefetchAsync(const void *devPtr,  
                                 size_t count,  
                                 int dstDevice,  
                                 cudaStream_t stream);

System Allocator 示例:

void test_prefetch_sam(cudaStream_t s) {
  char *data = (char*)malloc(N);
  init_data(data, N);                                     // execute on CPU
  cudaMemPrefetchAsync(data, N, myGpuId, s);              // prefetch to GPU
  mykernel<<<(N + TPB - 1) / TPB, TPB, 0, s>>>(data, N);  // execute on GPU
  cudaMemPrefetchAsync(data, N, cudaCpuDeviceId, s);      // prefetch to CPU
  cudaStreamSynchronize(s);
  use_data(data, N);
  free(data);
}

Managed 示例:

void test_prefetch_managed(cudaStream_t s) {
  char *data;
  cudaMallocManaged(&data, N);
  init_data(data, N);                                     // execute on CPU
  cudaMemPrefetchAsync(data, N, myGpuId, s);              // prefetch to GPU
  mykernel<<<(N + TPB - 1) / TPB, TPB, 0, s>>>(data, N);  // execute on GPU
  cudaMemPrefetchAsync(data, N, cudaCpuDeviceId, s);      // prefetch to CPU
  cudaStreamSynchronize(s);
  use_data(data, N);
  cudaFree(data);
}

1-2-8-2 数据使用提示 Data Usage Hints

当多个处理器同时访问相同数据时，可以使用 cudaMemAdvise 来提示系统如何访问 [devPtr, devPtr + count) 范围内的数据：

cudaError_t cudaMemAdvise(const void *devPtr,
                          size_t count,
                          enum cudaMemoryAdvise advice,
                          int device);

这些策略包括：

• cudaMemAdviseSetReadMostly：表示数据主要被读取，较少写入，有助于优化读写带宽。
• cudaMemAdviseSetPreferredLocation：设置数据的首选存储位置，通常为特定设备的物理内存，以减少不必要的数据迁移。
• cudaMemAdviseSetAccessedBy：表明数据将由特定设备频繁访问，以优化内存映射。

这些策略可以通过相应的 Unset 值来取消。这些提示有助于 CUDA 优化内存访问和迁移，从而提高性能。

cudaMemAdviseSetReadMostly 示例：

void test_advise_managed(cudaStream_t stream) {
  char *dataPtr;
  size_t dataSize = 64 * TPB;  // 16 KiB
  // Allocate memory using cudaMallocManaged
  // (malloc may be used on systems with full CUDA Unified memory support)
  cudaMallocManaged(&dataPtr, dataSize);
  // Set the advice on the memory region
  cudaMemAdvise(dataPtr, dataSize, cudaMemAdviseSetReadMostly, myGpuId);
  int outerLoopIter = 0;
  while (outerLoopIter < maxOuterLoopIter) {
    // The data is written to in the outer loop on the CPU
    init_data(dataPtr, dataSize);
    // The data is made available to all GPUs by prefetching.
    // Prefetching here causes read duplication of data instead
    // of data migration
    for (int device = 0; device < maxDevices; device++) {
      cudaMemPrefetchAsync(dataPtr, dataSize, device, stream);
    }
    // The kernel only reads this data in the inner loop
    int innerLoopIter = 0;
    while (innerLoopIter < maxInnerLoopIter) {
      mykernel<<<32, TPB, 0, stream>>>((const char *)dataPtr, dataSize);
      innerLoopIter++;
    }
    outerLoopIter++;
  }
  cudaFree(dataPtr);
}

1-2-8-3 查询托管内存上的数据使用属性

程序可以使用 API 查询 CUDA 托管内存上通过 cudaMemAdvise 或 cudaMemPrefetchAsync 设置的内存范围属性。这些属性包括：

• cudaMemRangeAttributeReadMostly：表示内存范围是否主要被读取。
• cudaMemRangeAttributePreferredLocation：表示内存的首选位置是 GPU 还是 CPU。
• cudaMemRangeAttributeAccessedBy：列出访问该内存范围的设备列表。
• cudaMemRangeAttributeLastPrefetchLocation：显示内存范围最后一次被显式预取的位置。

这些属性可以通过 cudaMemRangeGetAttribute 函数查询，也可使用 cudaMemRangeGetAttributes 函数同时查询多个属性。

cudaMemRangeGetAttribute(void *data,
                         size_t dataSize,
                         enum cudaMemRangeAttribute attribute,
                         const void *devPtr,
                         size_t count);

2 具有完整CUDA统一寻址支持的设备上的统一寻址 Unified Memory on devices with full CUDA Unified Memory support

2-1 系统分配内存示例

具有完整CUDA统一寻址支持的系统允许设备访问与设备交互的主机进程所拥有的任何内存。

__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");
}

Malloc 调用示例

void test_malloc() {
  const char test_string[] = "Hello World";
  char* heap_data = (char*)malloc(sizeof(test_string));
  strncpy(heap_data, test_string, sizeof(test_string));
  kernel<<<1, 1>>>("malloc", heap_data);
  ASSERT(cudaDeviceSynchronize() == cudaSuccess,
    "CUDA failed with '%s'", cudaGetErrorString(cudaGetLastError()));
  free(heap_data);
}

Managed 调用示例

void test_managed() {
  const char test_string[] = "Hello World";
  char* data;
  cudaMallocManaged(&data, sizeof(test_string));
  strncpy(data, test_string, sizeof(test_string));
  kernel<<<1, 1>>>("managed", data);
  ASSERT(cudaDeviceSynchronize() == cudaSuccess,
    "CUDA failed with '%s'", cudaGetErrorString(cudaGetLastError()));
  cudaFree(data);
}

Stack variable 调用示例

void test_stack() {
  const char test_string[] = "Hello World";
  kernel<<<1, 1>>>("stack", test_string);
  ASSERT(cudaDeviceSynchronize() == cudaSuccess,
    "CUDA failed with '%s'", cudaGetErrorString(cudaGetLastError()));
}

File-scope static variable 调用示例

void test_static() {
  static const char test_string[] = "Hello World";
  kernel<<<1, 1>>>("static", test_string);
  ASSERT(cudaDeviceSynchronize() == cudaSuccess,
    "CUDA failed with '%s'", cudaGetErrorString(cudaGetLastError()));
}

Global-scope variable 调用示例

const char global_string[] = "Hello World";

void test_global() {
  kernel<<<1, 1>>>("global", global_string);
  ASSERT(cudaDeviceSynchronize() == cudaSuccess,
    "CUDA failed with '%s'", cudaGetErrorString(cudaGetLastError()));
}

Global-scope extern variable 调用示例

// declared in separate file, see below
extern char* ext_data;

void test_extern() {
  kernel<<<1, 1>>>("extern", ext_data);
  ASSERT(cudaDeviceSynchronize() == cudaSuccess,
    "CUDA failed with '%s'", cudaGetErrorString(cudaGetLastError()));
}

2-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");
}

2-1-2 使用统一寻址进行进程间通信（IPC）

目前，使用IPC与统一寻址可能对性能有显著影响。CUDA IPC不支持管理内存共享。在支持完整CUDA统一寻址的系统上，系统分配的内存支持IPC。一旦系统分配的内存被共享给其他进程，就适用与文件支持的统一寻址类似的统一寻址编程模型。

2-2 性能调优

为获得统一寻址的良好性能，需了解系统分页机制和如何减少页面错误，理解保持数据靠近处理器的机制，并考虑根据系统内存传输粒度调整应用。

2-2-1 内存分页和页面大小

统一寻址系统使用虚拟地址空间和内存分页。了解虚拟页面大小的选择和系统是否提供CPU和GPU共用的页表对性能调优至关重要。

选择合适的页面大小，小页面大小减少内存碎片，但可能增加TLB未命中，大页面大小可能导致更多内存碎片，但TLB未命中减少。GPU上的TLB未命中通常比CPU更昂贵。

硬件一致性系统提供了CPU和GPU共用的页表，而软件一致性系统为CPU和GPU分别维护页表。硬件一致性系统在CPU和GPU频繁访问同一内存页时提供更好的性能。

某些设备支持从主机直接访问GPU内存，无需页面错误和数据迁移。这要求使用特定的内存使用提示来启用。

部分设备支持对主机内存的原子操作，无需页面错误。这些设备上的属性表明支持主机本地原子操作。

3 不具有完整CUDA统一寻址支持的设备上的统一寻址 Unified Memory on devices without full CUDA Unified Memory support

计算能力低于6.0的设备或Windows平台支持CUDA管理内存v1.0，但对数据迁移和一致性以及内存超额预定的支持有限。下面描述了如何在这些平台上使用和管理内存。

3-1 数据迁移和一致性

低于6.0计算能力的GPU不支持按需将托管数据移动到GPU。引入了新的GPU页面错误机制，通过系统范围的虚拟地址空间，页面错误提供了无需在每次内核启动前同步所有托管内存的好处。

3-2 GPU内存超额预分配

低于6.0计算能力的设备无法分配超过GPU物理内存大小的托管内存。

3-3 多GPU

在这些系统上，托管分配对所有GPU自动可见，通过GPU的点对点功能。托管内存的行为类似于未托管内存，当前活动的设备是物理分配的主场，但其他GPU可以通过PCIe总线以降低的带宽访问内存。

3-4 一致性和并发

在计算能力低于6.0的设备上，CPU和GPU不能同时访问托管内存，因为无法保证一致性。