Alex_McAvoy

想要成为渔夫的猎手

NVIDIA CUDA2023春训营(四)CUDA 存储单元

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() 的函数原型如下:

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

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

1
2
3
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 显存,函数原型如下:

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

其使用方法与 cudaMallocHost()cudaFreeHost() 类似,例如:

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

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

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

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

资源传递

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

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

1
__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() 传输资源的实例

1
2
3
4
5
6
7
8
9
10
11
// 申请大小
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() 函数来传递或获取变量值,函数原型如下:

1
2
__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)

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

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
#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() 来获取其值

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

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
#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() 进行资源释放,函数原型如下:

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

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

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

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

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
#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;
}
感谢您对我的支持,让我继续努力分享有用的技术与知识点!