Reference

CUDA 存储单元架构

CUDA 各存储单元架构如下

CUDA 各存储单元的对比如下

CUDA 函数关系与 device 端变量访问权限如下

主存（Host Memory）是 host 端的内存，其可分为可分页内存（Pageable Memory）和锁定内存（Page-Locked Memory / Pinned Memory） 两种

可分页内存由 malloc()、new() 等操作系统 API 在 host 端分配与释放的，该内存是可以换页的，即内存页可以被置换到磁盘中，普通的 C/C++ 程序使用的内存就是该内存

锁定内存是由 CUDA 函数 cudaMallocHost()、cudaFree() 分配与释放的，其一大特点是操作系统不会对这块内存进行分页与交换操作，能够确保该内存始终驻留在物理内存中，不会被分配到低速的虚拟内存

同时，由于 GPU 知道锁定内存的物理地址，因此可以通过 DMA（Direct Memory Acess）技术直接在 host 端和 device 端进行通信，速率更快

cudaMallocHost() 和 cudaFree() 的函数原型如下：

__host__ cudaError_t cudaMallocHost(void **ptr, size_t size)
__host__ cudaError_t cudaFreeHost(void *ptr)

关键在于 cudaMallocHost() 函数第一个参数的二维指针，如下例

double *host_data = NULL;
size_t size = sizeof(double) * 1024;
cudaMallocHost((void**) &host_data, size);

host_data 是存储在 host 端上的指针变量，其要存储的值是内存地址，现在想要利用 cudaMallocHost() 在 host 端申请了一个大小为 1024 的 double 型一维数组

由于该函数的无法返回在 host 申请的首地址，那么就需要利用参数来传递这个地址，也就是存储在 host 端中 host_data 这个指针变量的地址 &host_data

当 cudaMallocHost() 执行完后，会向 host_data 这个指针变量中写入一个地址值，这个地址值就是在 host 端所申请的数组首地址

寄存器与本地内存

寄存器（Registers）是速度最快的存储单元，位于 GPU 的计算单元（Streaming Multiprocessor，SM）上，当核函数启动后，这些位于计算单元上的寄存器会被分配给指定的线程使用

在核函数中，没有特殊声明的自动变量与数组都是存放在寄存器里，这些变量是每个线程私有的，一旦线程执行结束，寄存器变量就会失效

本地内存（Local Memory）在硬件中没有特定的存储单元，其是从全局内存上虚拟出来的地址空间，因此针对它的访问速度与全局内存是相近的

本地内存是为寄存器无法满足存储需求的情况而设计的，其与寄存器相似，是线程私有的，当寄存器不够用时，就会使用本地内存来代替这部分存储空间

此外，当出现以下几种情况时，编译器会将变量放到内存空间

• 编译期间无法确定值的本地数组
• 消耗太多寄存器的较大的结构体或数组
• 超过寄存器限制的变量

全局内存（Global Memory）是 GPU 中空间最大、最基础的内存，任意 SM 都可以在整个程序的生命周期中获取其状态

某种意义上，常说的 GPU 显存就是指全局内存，其是核函数输入数据和写入结果的唯一来源

在 CUDA 中，使用 cudaMalloc() 和 cudaFree() 函数可以申请和释放 GPU 显存，函数原型如下：

__host__ __device__ cudaError_t cudaMalloc(void **devPtr, size_t size)
__host__ __device__ cudaError_t cudaFree(void *devPtr)

其使用方法与 cudaMallocHost() 和 cudaFreeHost() 类似，例如：

double *device_data = NULL;
size_t size = sizeof(double) * 1024;
cudaMalloc((void**) &device_data, size);

使用 cudaMemset() 可以对 device 端申请的显存进行初始化，其类似于 memset() 函数，函数原型如下：

__host__ cudaError_t cudaMemset (void* devPtr, int value, size_t count)

需要注意的是，与 memset() 类似，其是以字节为单位来进行赋值的，需要使用十六进制来进行赋值，因此一般使用该函数将申请的空间置为全 0 或全 -1

对于 CUDA 程序，除了 host 端和 device 端的内存申请、释放，以及核函数在 device 端的执行外，另一关键的步骤就是数据在 host 端和 device 端上的传输

CUDA 使用 cudaMemcpy() 函数将资源内存复制到目标内存中，类似于 C 语言中的 memcpy()，函数原型如下：

__host__ cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, cudaMemcpyKind kind)

其输入参数有四个：

• *dst：指向用于存储复制内容的目标数组
• *src：指向要复制内容的数据源
• conut：要复制的数据大小，以 Byte 为单位
• kind：复制的方向，从 host 端复制到 device 端为 cudaMemcpyHostToDevice，从 device 端复制到 host 端为 cudaMemcpyDeviceToHost

下述代码给出了一个使用 cudaMemcpy() 传输资源的实例

// 申请大小
const int size = sizeof(double) * 100;
// 申请host内存
double *h_x = (double*) malloc(size);
// 申请device显存
double *d_x;
cudaMalloc((void **)&d_x, size);
// host到device
cudaMemcpy(d_x, h_x, size, cudaMemcpyHostToDevice);
// device到host
cudaMemcpy(h_x, d_x, size, cudaMemcpyDeviceToHost);

需要注意的是，该函数是一个同步函数，在未完成数据转移操作前会锁死并一直占有 CPU 控制权，因此无需再添加 cudaDeviceSynchronize() 同步函数

使用 __device__ 修饰符可以定义 device 端的全局变量，其与 C/C++ 的全局变量声明位置相同，只能在类和函数外声明

在 host 端无法直接访问 __device__ 修饰的变量，只能通过 cudaMemcpyToSymbol() 和 cudaMemcpyFromSymbol() 函数来传递或获取变量值，函数原型如下：

__host__ cudaError_t cudaMemcpyToSymbol(const void* symbol, const void* src, size_t count, size_t offset = 0, cudaMemcpyKind kind = cudaMemcpyHostToDevice)
__host__ cudaError_t cudaMemcpyFromSymbol(void* dst, const void* symbol, size_t count, size_t offset = 0, cudaMemcpyKind kind = cudaMemcpyDeviceToHost)

下述代码给出了一个使用全局变量的实例

#include <stdio.h>
#define N 5

// 声明device端全局变量
__device__ int constant_a[N];

// 输出全局变量核函数
__global__ void print_constant() {
    int index = blockDim.x * blockIdx.x + threadIdx.x;
    if (index >= N)
        return;
    printf("%d ",constant_a[index]);
}

// 修改全局变量核函数
__global__ void update_constant() {
    int index = blockDim.x * blockIdx.x + threadIdx.x;
    if (index >= N)
        return;
    constant_a[index] = 1024;
}

int main() {
    int h_a[N] = {0,1,2,3,4};

// 修改device端全局变量
    cudaMemcpyToSymbol(constant_a, h_a, sizeof(h_a));

printf("修改前的全局变量：");
    print_constant<<<1, 16>>>();

cudaDeviceSynchronize();

// 获取全局变量
    update_constant<<<1, 16>>>();
    cudaMemcpyFromSymbol(h_a, constant_a, sizeof(h_a));
    printf("\n修改后的全局变量：");
    for (int i = 0; i < N; i++) {
        printf("%d ", h_a[i]);
    }
    printf("\n");

cudaDeviceSynchronize();

return 0;
}

固定内存（Constant Memory）类似于本地内存，没有特定的存储单元的，只是全局内存的虚拟地址

与本地内存不同的是，其范围是全局的，所有核函数均可见，同时其是只读的，因此一定程度上简化了缓存管理，硬件无需管理复杂的回写策略

固定内存中的变量使用 __constant__ 来修饰，即 device 端的常量

由于其是只读的，因此必须在 host 端使用 cudaMemcpyToSymbol() 函数来进行初始化赋初值，且一经赋值后就无法再更改

当想要在 host 端读取常量值时，需要使用 cudaMemcpyFromSymbol() 来获取其值

下述代码给出了一个使用全局常量的实例

#include <stdio.h>
#define N 5

// 声明device端常量
__device__ int constant_a[N];

// 输出常量核函数
__global__ void print_constant() {
    int index = blockDim.x * blockIdx.x + threadIdx.x;
    if (index >= N)
        return;
    printf("%d ",constant_a[index]);
}

int main() {
    int h_a[N] = {0,1,2,3,4};

// 初始化device端常量
    cudaMemcpyToSymbol(constant_a, h_a, sizeof(h_a));

printf("通过核函数读取常量：");
    print_constant<<<16, 1>>>();

cudaDeviceSynchronize();

// 获取全局变量
    printf("\n通过host端读取常量：");
    cudaMemcpyFromSymbol(h_a, constant_a, sizeof(h_a));
    for (int i = 0; i < N; i++) {
        printf("%d ", h_a[i]);
    }
    printf("\n");

cudaDeviceSynchronize();

return 0;
}

共享内存（Shared Memory）的访问延迟仅次于寄存器，其可以被一个 block 中的所有线程访问，从而实现 block 内线程间的低开销通信

若线程频繁对某个数据进行读写操作，可以设置将该数据常驻共享内存，从而提高代码运行效率

块内共享变量

共享内存中的变量使用 __shared__ 来修饰，其能被一个 block 中的所有线程访问，因为被称为块内共享变量

共享变量只能在 __device__ 函数或者 __global__ 函数内被声明，不能跨过一个 block，因此，某个 block 中的共享变量是无法被其他 block 所访问到的

此外，由于共享内存中的数据可以被一个 block 中的所有线程访问，那么当多个线程对同一个共享变量进行操作时，需要对线程进行同步操作，从而避免竞争的发生，常使用 __syncthreads() 来控制线程同步

共享变量的一个典型应用是用来优化矩阵乘法，详见：NVIDIA CUDA2023春训营（五）CUDA 向量加法与矩阵乘法

纹理内存（Texture Memory）实际上也是全局内存的一部分，从读取性能的角度来说，其与固定内存类似

但与固定内存不同的是，它有自己专属的只读 Cache，这个 Cache 在进行浮点运算时十分有用

纹理内存实质上是针对 2D 空间局部性的优化策略，要获取 2D 数据时，就可以使用纹理内存来获取

统一内存（Unifled Memory）是 CUDA 6.0 引入的，其避免了编写程序时在 host 和 device 上进行内存分配与数据传输的麻烦

统一内存定义了一个托管内存（Managed Memory）来共同管理 host 和 device 中的内存，使得 host 端和 device 端都可以看到共同的地址空间，无需再使用 cudaMemcpy() 函数进行资源传递

在使用统一内存时，分配空间是在 host 端和 device 端的全局内存上各自申请了一块空间，只是可以使用一个变量来共同维护

统一内存资源分配

统一内存中的变量使用 __managed__ 来修饰，需要在类和函数外声明

此外，CUDA 中还可使用 cudaMallocManaged() 函数在 host 端上分配统一内存，但使用完毕后需要使用 cudaFree() 进行资源释放，函数原型如下：

__host__ cudaError_t cudaMallocManaged(void **devPtr, size_t size, unsigned int flags = cudaMemAttachGlobal)

其第三个参数是一个标志，默认值为 cudaMemAttachGlobal，代表内存可由任何设备上的任何流访问

两种统一内存分配方式的分配行为相同，不同的是由于 cudaMallocManaged() 函数第三个参数的限制，使用 cudaMallocManaged() 函数分配的统一内存可能会受到 cudaStreamAttachMemAsync() 的限制

如下给出了一个使用统一内存的实例

#include <stdio.h>
#define N 64

// 第一种分配方式
__managed__ int arr_managed1[N];

__global__ void add1(int a, int b) {
    arr_managed1[threadIdx.x] = a + b;
}

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

int main() {

printf("第一种分配方式运行核函数：\n");

add1<<< 1, N>>>(100, 100);
    cudaDeviceSynchronize();

for(int i = 0; i < N; i++)
        printf("%d ", arr_managed1[i]);
    printf("\n");


    // 第二种分配方式
    int *arr_managed2;
    cudaMallocManaged(&arr_managed2, sizeof(int) * N);

printf("第二种分配方式运行核函数：\n");
    add2<<< 1, N>>>(arr_managed2, 200, 100);
    cudaDeviceSynchronize();

for(int i = 0; i < N; i++)
        printf("%d ", arr_managed2[i]);
    printf("\n");

// 释放第二种分配方式申请的统一内存资源
    cudaFree(arr_managed2); 

return 0;
}