企业官网建设流程全解析

统一内存和系统内存

统一虚拟地址空间——CPU内存和每个GPU的内存都在单一虚拟地址空间内拥有独立的地址范围

统一内存——一种CUDA功能，可实现托管内存，并能自动在CPU和GPU之间迁移

受限统一内存——一种具有部分限制的统一内存范式

完全统一内存——全面支持统一内存功能

硬件一致性下的完全统一内存——利用硬件能力全面支持统一内存

统一内存提示——用于指导特定分配中统一内存行为的API

页锁定主机内存——非页可访问的系统内存，某些CUDA操作需要使用

映射内存——一种不同于统一内存的机制，允许内核直接访问主机内存

异构管理内存（HMM）——Linux内核的一项特性，可实现对完整统一内存的软件一致性

地址转换服务（ATS）——当GPU通过NVLink芯片间互连（C2C）连接到CPU时可用的硬件特性，为完整统一内存提供硬件一致性

统一虚拟地址空间(UVA)

在单个操作系统进程内，所有主机内存和系统中所有GPU的全局内存均使用同一个虚拟地址空间，无论是通过CUDA API（如cudaMalloc、cudaMallocHost）还是系统分配API（如new、malloc、mmap）进行内存分配，所有内存分配都位于该虚拟地址空间内。

简单来说统一虚拟地址空间是CUDA提供的一种内存寻址机制，其作用便是让CPU和GPU使用同一虚拟地址空间来管理内存地址。

作用

其主要有以下三点作用

自动识别指针指向哪里的内存--cudaPointerGetAttributes()

cudaError_t cudaPointerGetAttributes( cudaPointerAttributes* attributes, const void* ptr );

attributes:输出参数，用于保存查询结果

ptr:输入参数，需要查询的内存指针

cudaPointerAttributes结构体

struct cudaPointerAttributes { cudaMemoryType type; int device; void* devicePointer; void* hostPointer; };

type:内存类型

枚举值	含义	常见分配方式
cudaMemoryTypeUnregistered	未注册的普通主机内存	malloc()
cudaMemoryTypeHost	CUDA 注册或分配的主机内存	cudaMallocHost()
cudaMemoryTypeDevice	GPU 设备内存	cudaMalloc()
cudaMemoryTypeManaged	CUDA 托管内存	cudaMallocManaged()

device:关联的GPU设备编号

devicePointer:当前GPU可以用来访问该内存的指针，无对用地址时为NULL

hostPointer:CPU可以用来访问该内存的指针，无对用地址时为NULL

#include <iostream> #include <cuda_runtime.h> void checkPointer(void* ptr) { cudaPointerAttributes attr{}; cudaError_t err = cudaPointerGetAttributes(&attr, ptr); if (err != cudaSuccess) { std::cout << "查询失败:" << cudaGetErrorString(err) << std::endl; return; } std::cout << "Memory type:" << attr.type << std::endl; std::cout << "Device:" << attr.device << std::endl; std::cout << "Device pointer:" << attr.devicePointer << std::endl; std::cout << "Host pointer:" << attr.hostPointer << std::endl; std::cout << std::endl; } int main() { int* h_ptr; int* d_ptr; int* m_ptr; // 普通 CPU 内存 h_ptr = (int*)malloc(1024); // GPU 显存 cudaMalloc(&d_ptr, 1024); // 托管内存 cudaMallocManaged(&m_ptr, 1024); checkPointer(h_ptr); checkPointer(d_ptr); checkPointer(m_ptr); free(h_ptr); cudaFree(d_ptr); cudaFree(m_ptr); return 0; }

自动判断拷贝方向--cudaMemcpyDefault

在没有UVA时，使用cudaMemcpy通常需要明确指定数据传输方向

cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); cudaMemcpy(h_A, d_A, size, cudaMemcpyDeviceToHost);

而有UVA后，通常可以写成

cudaMemcpy(d_A, h_A, size, cudaMemcpyDefault); cudaMemcpy(h_A, d_A, size, cudaMemcpyDefault);

CUDA根据指针信息判断数据传输方向

简化部分CPU/GPU指针的使用

在支持UVA的情况下，CPU和GPU在某些内存分配方式下，可以使用相同的指针值访问同一块内存，不需要分别维护两个指针。

#include <cuda_runtime.h> #include <iostream> __global__ void kernel(int* ptr) { ptr[0] = 100; } int main() { int* ptr; cudaHostAlloc(&ptr, sizeof(int), cudaHostAllocMapped); // CPU写入 ptr[0] = 10; // GPU使用同一个指针 kernel<<<1, 1>>>(ptr); cudaDeviceSynchronize(); // CPU读取GPU修改后的值 std::cout << ptr[0] << std::endl; cudaFreeHost(ptr); }

内存分配方式	CPU/GPU 能否直接使用同一指针
malloc()	一般不能
cudaMalloc()	CPU 一般不能直接解引用
cudaHostAlloc()	在 UVA 支持下通常可以
cudaMallocManaged()	可以

统一内存

统一内存是CUDA的一项内存特性，它允许在CPU或GPU通过同一个指针访问由CUDA管理的内存，从而减少开发者手动管理数据传输的工作。

因此我们可以简单的认为，统一内存就是管理内存

在某些系统上，统一内存必须显式分配。可以通过以下几种方式显式分配管理内存：

cudaMallocManaged

cudaMallocFromPoolAsync，并使用“allocType”设置为cudaMemAllocationTypeManaged创建的内存池
带有“__managed__”指示符的全局变量

统一内存范式

统一内存由于各个操作系统、硬件等有所不同，因此可以通过cudaDeviceGetAttribute查询一下属性来确定可用的统一形式（具体用法可查看CUDA C++入门和重要函数讲解--GPU属性）

查询相关属性，可以使用cudaDeciveProp的成员函数--cudaDevAttrConcurrentManagedAccess:1表示完全支持统一内存，0表示仅有限支持

cudaDevAttrPageableMemoryAccess:1表示所有系统内存均支持统一内存，0表示仅显式分配为管理内存的部分才支持统一内存

cudaDevAttrPageableMemoryAccessUsesHostPageTables:指示CPU/GPU一致性机制：1 表示硬件实现，0 表示软件实现

#include <cuda_runtime.h> #include <iostream> int main() { int DevAttrConcurrentManagedAccess; cudaDeviceGetAttribute( &DevAttrConcurrentManagedAccess, cudaDevAttrConcurrentManagedAccess, 0 ); int DevAttrPageableMemoryAccess; cudaDeviceGetAttribute( &DevAttrPageableMemoryAccess, cudaDevAttrPageableMemoryAccess, 0 ); int DevAttrPageableMemoryAccessUsesHostPageTables; cudaDeviceGetAttribute( &DevAttrPageableMemoryAccessUsesHostPageTables, cudaDevAttrPageableMemoryAccessUsesHostPageTables, 0 ); std::cout << "是否完全支持统一内存:" << DevAttrConcurrentManagedAccess << std::endl; std::cout << "是否所有系统内存均支持统一内存:" << DevAttrPageableMemoryAccess << std::endl; std::cout << "CPU/GPU一致性机制实现方式:" << DevAttrPageableMemoryAccessUsesHostPageTables << std::endl; }

统一内存操作共有四种模式：
- 完全支持显式管理内存分配
- 完全支持所有分配并采用软件一致性机制
- 完全支持所有分配并采用硬件一致性机制
- 有限支持统一内存

当完全支持可用时，可以要求显式分配内存，或者所有系统内存都可能隐式地统一为共享内存。当所有内存都隐式统一时，一致性机制可以是软件或硬件实现。Windows 和部分 Tegra 设备对统一内存的支持有限。

下面将cudaDevAttrConcurrentManagedAccess、cudaDevAttrPageableMemoryAccess和cudaDevAttrPageableMemoryAccessUsesHostPageTables简化为[1]、[2]、[3]

模式	[1]	[2]	[3]	判定结果
1	0	不关心	不关心	Limited unified memory support（受限统一内存支持）
2	1	0	不关心	Full support for explicit managed memory allocations（显式托管内存完全支持）
3	1	1	0	Full support for all allocations with software coherence（所有分配都支持，软件一致性）
4	1	1	1	Full support for all allocations with hardware coherence（所有分配都支持，硬件一致性）

模式	能否使用cudaMallocManaged()	普通malloc/new内存能否被GPU直接纳入统一内存支持	一致性方式	使用特点	典型理解
1	可以	不可以	无完整一致性支持	只能用托管内存，且 CPU/GPU 同时访问受限，通常要严格同步	能用统一内存，但限制较多
2	可以	不可以	对托管内存提供完整支持	cudaMallocManaged()分配的内存支持更完整的迁移/访问；普通主机内存仍不行	托管内存很好用，但仅限托管内存
3	可以	可以	软件一致性	不仅托管内存，普通系统内存也能纳入统一内存支持；一致性主要靠软件机制维护	范围更大，但一致性更多靠软件
4	可以	可以	硬件一致性	托管内存和普通系统内存都支持，且由硬件提供更强的一致性支持	最完整、最理想的统一内存支持

全面支持统一内存功能

通常来说，工业上CUDA使用的环境都是linux，而大多数linux都全面支持统一内存，即模式2、3和4

对于全面支持统一内存功能的系统通常有如下特点：

- 统一内存通常在处理器首次触达其所在内存空间时进行分配
- 当统一内存被当前处理器以外的其他处理器使用时，通常会被迁移
- 统一内存的迁移或访问以内存页（软件一致性）或缓存行（硬件一致性）为粒度进行
- 允许超分配：应用程序可分配超过GPU物理可用内存的统一内存

- 分配和迁移行为可能与上述描述有所不同

有限支持统一内存功能

在 Windows（包括Windows子系统--WSL）上，仅提供统一内存功能的有限子集。在这些系统中，可使用受管理的内存，但 CPU 与 GPU 之间的迁移行为有所不同。

对于有限支持统一内存功能的系统通常有如下特点：

- 统一内存首先分配在 CPU 的物理内存中
- 统一内存以比虚拟内存页更大的粒度进行迁移
- 当GPU开始执行时，统一内存才会被迁移至GPU
- 在GPU活动期间，CPU不能访问统一内存
- 当GPU完成同步后，统一内存将被迁移回CPU
- 不允许对GPU内存进行超分配
- 仅CUDA显式分配为统一内存的内存才被统一处理

Page-Locked Host Memory(页锁定主机内存)

页锁定主机内存，也叫Pinned Memory（固定内存、锁页内存）。它是一种特殊的CPU内存，主要用于提高CPU和GPU之间的数据传输效率，并支持真正的异步数据传输。

其常使用的相关API如下（部分API可参考CUDA C++入门和重要函数讲解--内存管理）:

API	作用
cudaMallocHost()	分配页锁定主机内存
cudaHostAlloc()	分配页锁定主机内存，并可指定额外属性
cudaFreeHost()	释放 CUDA 分配的页锁定内存
cudaHostRegister()	将已经分配的普通内存注册为页锁定内存
cudaHostUnregister()	解除已有内存的页锁定注册
cudaHostGetDevicePointer()	获取映射主机内存对应的 GPU 可访问指针

需要注意的是，页锁定主机内存并不是越多越好，虽然该内存可以快速的传输数据，但是分配过多会减少操作系统可以灵活调度的内存数量，从而导致整理性能的降低。

其实这个页锁定主机内存在实际中使用相当多，比如我们之前的项目，需要通过摄像头采集图像，采集好的图像，通常储存在CPU中，因此就可以通过页锁定主机内存来传输预处理后的图像给GPU，然后进行推理。

nvcc:CUDA编译器

CUDA编译器nvcc是NVIDIA提供的一套工具链，用于编译CUDA C/C++代码，这里主要讲解在构建应用程序时常用的nvcc使用方式和相关细节。

更加详细的文档请看：

1. Introduction — NVIDIA CUDA Compiler Driver 13.4 documentation

CUDA源文件和头文件

使用 nvcc 编译的源文件通常是主机代码（在 CPU 上执行）和设备代码（在 GPU 上执行）的组合。nvcc 接受常见的 C/C++ 源文件扩展名：.c、.cpp、.cc、.cxx 用于仅包含主机代码的文件，以及 .cu 用于包含设备代码或混合主机与设备代码的文件。包含设备代码的头文件通常采用 .cuh 扩展名，以区别于仅包含主机代码的头文件（如 .h、.hpp、.hh、.hxx 等）。

简单来说，.cu和.cuh分别是内核执行代码的源文件和头文

nvcc中代码文件的后缀名和作用

源代码文件

GPU编译过程中的文件

NVCC编译工作流

NVCC基础用法

NVCC的通用结构如下：nvcc [编译选项] 源文件... [-o 输出文件]

nvcc:调用CUDA编译器驱动

编译选项:指定编译标准、优化、架构等

-Ipath--指定包含目录(大写的i)、-Lpath--指定包含库路径、-lpath--链接其他库

源文件:要编译的一个或多个文件

-o:指定输出文件名称和路径

观察编译过程

指定C++标准(-std)和GPU架构目标(-arch)

通过使用-std=C++17来指定使用C++17标准

通过使用-arch=sm_120来指定使用CC为12的机器代码架构目标。

优化级别(-o)

调试与错误检查相关选项

性能优化

-maxrregcount=64:限制寄存器数量为64

-Xptxas=-v:查看GPU资源使用情况

-res-usage：编译后打印资源使用报告，包括每个内核函数分配的寄存器、共享内存、常量内存和局部内存的数量

-opt-info=inline：打印内联函数的相关信息

浮点运算精度相关参数

Xptxas

Xptxas会讲参数传递给PTX汇编共ptaxs，其之所以重要，是因为其主要用法可以查看GPU资源使用、检查寄存器溢出、控制设备代码优化，以及调整缓存策略。

其常用格式为:-Xptxas=参数，部分用法在上面已经提及

从上表可以看出，Xptxas支持多个ptxas参数同时使用，不同参数之间使用都好隔开即可

默认缓存策略

正如前面文章所说的，L1缓存更加靠近SM，因此其访问延迟通常更低，因此更加适合在同一个线程块内的线程反复读取相同或附近的数据。而L2缓存容量更大，同时由多个SM共享，因此更加适合多个线程块或多个SM访问同一批数据。通过默认缓存策略便可以控制在加载或者储存是对于L1和L2的缓存策略。

控制加载时的缓存策略(-dlcm)

-Xptxas=-dlcm=ca -Xptxas=-dlcm=cg

参数值	含义	用法
ca	Cache at All Levels	加载时允许使用 L1、L2 等缓存层级
cg	Cache at Global Level	加载时绕过 L1，使用 L2 及以下缓存层级

控制储存时的缓存策略(-dscm)

-Xptxas=-dscm=wb -Xptxas=-dscm=cg

参数值	含义	直观理解
wb	Write-back	使用常规写回缓存策略
cg	Cache at Global Level	存储时绕过 L1，缓存于 L2 及以下层级