【CUDA编程】系列博客参考NVIDIA官方文档“CUDA C++ Programming Guide（v12.6）”。
本文为原创文章，未经本人允许，禁止转载。转载请注明出处。
1.Device Memory
CUDA编程模型假定系统由host和device组成，每个都有自己独立的内存。kernel在device内存中运行，因此运行时提供了分配、释放和复制device内存的功能，以及在host内存和device内存之间传输数据。
device内存可以通过线性内存（linear memory）或CUDA数组（CUDA arrays）来分配。
CUDA数组为不透明的内存布局，专门为texture获取进行了优化，主要用于图像处理等需要高效内存访问的应用。
线性内存在单一且统一的地址空间中分配，允许通过指针进行随机访问。地址空间的大小取决于CPU和所用GPU的计算能力。
“x86_64 (AMD64)”、“POWER (ppc64le)”、“ARM64”是三种不同的CPU架构。表中40bit表示地址的位数为40个比特位，即该地址空间共有$2^{40}$个地址，一个地址通常对应一个字节的内存空间（即1B），那么这个地址空间对应的内存空间就是$2^{40} \times 1\text{B} = 1 \text{TB}$。47bit对应128TB，48bit对应256TB，49bit对应512TB。
注意：
在计算能力为5.3（Maxwell）及更早的device上，CUDA驱动程序会创建一个未提交的40位虚拟地址保留，以确保内存分配的指针落入支持的范围内。此保留显示为保留的虚拟内存，但在程序实际分配内存之前不占用任何物理内存。
线性内存通常使用cudaMalloc()来分配，使用cudaFree()来释放。host内存和device内存之间的数据传输通常使用cudaMemcpy()。下面是一个一维向量相加的例子。
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
47
48
49
50
51
52
// Device code
__global__ void VecAdd(float* A, float* B, float* C, int N)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < N)
C[i] = A[i] + B[i];
}
// Host code
int main()
{
int N = ...;
size_t size = N * sizeof(float);
// Allocate input vectors h_A and h_B in host memory
float* h_A = (float*)malloc(size);
float* h_B = (float*)malloc(size);
float* h_C = (float*)malloc(size);
// Initialize input vectors
...
// Allocate vectors in device memory
float* d_A;
cudaMalloc(&d_A, size);
float* d_B;
cudaMalloc(&d_B, size);
float* d_C;
cudaMalloc(&d_C, size);
// Copy vectors from host memory to device memory
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
// Invoke kernel
int threadsPerBlock = 256;
int blocksPerGrid =
(N + threadsPerBlock - 1) / threadsPerBlock;
VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
// Copy result from device memory to host memory
// h_C contains the result in host memory
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
// Free device memory
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
// Free host memory
...
}
二维数组和三维数组的线性内存分配可以使用cudaMallocPitch()和cudaMalloc3D()，因为它们可以确保分配的内存经过适当的填充，以满足device内存访问时的对齐要求，从而在内存访问和数据传输时获得最佳性能。二维数据和三维数据的数据传输分别使用cudaMemcpy2D()和cudaMemcpy3D()。
下面是一个遍历二维数组（数组元素为float类型）的一个例子：
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
// Host code
int width = 64, height = 64;
float* devPtr;
size_t pitch;
cudaMallocPitch(&devPtr, &pitch,
width * sizeof(float), height);
MyKernel<<<100, 512>>>(devPtr, pitch, width, height);
// Device code
__global__ void MyKernel(float* devPtr,
size_t pitch, int width, int height)
{
for (int r = 0; r < height; ++r) {
float* row = (float*)((char*)devPtr + r * pitch);
for (int c = 0; c < width; ++c) {
float element = row[c];
}
}
}
要解释上述代码，我们先来简单的了解下内存对齐。内存对齐指的是数据在内存中的起始地址应该是某个特定字节数的倍数，这个倍数通常由处理器的架构决定。例如，32字节对齐意味着数据的起始地址应该是32的倍数（比如十六进制的地址0x1000，对应十进制为4096，便是32的整倍数，而下一个地址0x1001就不是32的整倍数）。内存对齐有助于以最高效率读取和处理数据。如果二维数组的每一行不能满足这种对齐要求，CUDA可能会在每行末尾添加一些额外的字节（即padding），使下一行数据的起始地址符合对齐要求。而填充后每一行在内存中的实际宽度（单位是字节）就是pitch。如下图所示：
说回上述代码，cudaMallocPitch()函数：
1
2
3
4
5
6
__host__cudaError_t cudaMallocPitch (
void **devPtr, //返回一个指向二维数组首元素（即(0,0)）的指针
size_t *pitch, //返回pitch值
size_t width, //指定二维数组的width（以字节为单位）
size_t height //指定二维数组的height
)
上述代码的MyKernel函数并没有使用多线程，而是只用了一个线程。二维数组中获取元素地址的方法：
1
T* pElement = (T*)((char*)BaseAddress + Row * pitch) + Column;
下面是一个遍历三维数组（数组元素为float类型）的一个例子：
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
// Host code
int width = 64, height = 64, depth = 64;
cudaExtent extent = make_cudaExtent(width * sizeof(float),
height, depth); //定义三维数组的尺寸
cudaPitchedPtr devPitchedPtr;
cudaMalloc3D(&devPitchedPtr, extent);
MyKernel<<<100, 512>>>(devPitchedPtr, width, height, depth);
// Device code
__global__ void MyKernel(cudaPitchedPtr devPitchedPtr,
int width, int height, int depth)
{
//MyKernel只用了1个线程
char* devPtr = devPitchedPtr.ptr; //三维数组起始地址的指针
size_t pitch = devPitchedPtr.pitch; //每一行的pitch（单位为字节）
size_t slicePitch = pitch * height;
for (int z = 0; z < depth; ++z) {
char* slice = devPtr + z * slicePitch;
for (int y = 0; y < height; ++y) {
float* row = (float*)(slice + y * pitch);
for (int x = 0; x < width; ++x) {
float element = row[x];
}
}
}
}
为避免分配过多内存从而影响系统性能，我们应该根据问题的规模和实际需求来设置内存分配参数。如果内存分配失败（比如内存不足），建议使用其他类型的内存（比如
cudaMallocHost()、cudaHostRegister()等）或者返回一个错误信息。如果我们的应用程序因为某些原因无法请求内存分配，建议使用cudaMallocManaged()。
通过运行时API访问全局变量的各种方法：
1
2
3
4
5
6
7
8
9
10
11
12
13
__constant__ float constData[256]; //创建位于constant内存的全局变量
float data[256];
cudaMemcpyToSymbol(constData, data, sizeof(data)); //将数据从host内存复制到constant内存
cudaMemcpyFromSymbol(data, constData, sizeof(data)); //将数据从constant内存复制到host内存
__device__ float devData; //创建位于global内存的全局变量
float value = 3.14f;
cudaMemcpyToSymbol(devData, &value, sizeof(float));
__device__ float* devPointer;
float* ptr;
cudaMalloc(&ptr, 256 * sizeof(float));
cudaMemcpyToSymbol(devPointer, &ptr, sizeof(ptr)); //从global内存到global内存
global内存、constant内存参见：Memory Hierarchy。
对于通过__device__或__constant__声明的全局变量，cudaGetSymbolAddress()可以获取它们的地址，cudaGetSymbolSize()用于获取分配给全局变量的内存大小。
