基本的GPU内存操作
GPU内存
虽然GPU具有强大的算力，但GPU不能单独工作，需要与CPU一起并作为CPU的协处理器才能工作。CPU与GPU分别具有独立的内存系统，见下图。CPU端也称为Host端，CPU内存称为Host（主机）内存;GPU端也成为Device（设备）端，其内存称为Device内存。一般情况下，如果我们要在GPU端进行计算，就需要把待处理的数据拷贝到到Device内存中，待数据处理完成之后，还需要把计算结果拷贝到Host端做进一步的处理，比如存储到硬盘中或者打印到显示器上。这一小节主要介绍如何在GPU端分配与释放内存以及如何在CPU与GPU之间进行数据的拷贝。
把前面小节中向量加的main函数代码拷贝到下面：
int main(void) {
size_t size = N * sizeof(int);
int *h_a, *h_b;
int *d_a, *d_b, *d_c;
h_a = (int *)malloc(size);
h_b = (int *)malloc(size);
...
cudaMalloc((void **)&d_a, size);
cudaMalloc((void **)&d_b, size);
cudaMalloc((void **)&d_c, size);
cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
vectoradd<<<Grid, block>>>(d_a, d_b, d_c, N);
cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
free(h_a); free(h_b);
return 0;
}
从这段代码中我们可以看到，总共声明了两类不同的指针，int *h_a, *h_b;以及int *d_a, *d_b, *d_c;。前者是Host指针，指向host端内存；后者是Device指针，指向Device端内存。Device端的指针可以在主机端调用，但是不能在主机端解引用。主机端指针同样不能在设备端解引用。这里的设备端指针并不是位于设备端，而是指向设备端的内存。指针仍然是在主机端上，所以主机端可以使用这个指针，但是不能够解引用指向设备端内存的指针。
内存分配与释放
我们知道在CPU端代码中，内存的分配、初始化以及释放可以调用下面的函数实现：
malloc(size_t)
memset(void *, int, size_t)
free(void*)
相应的CUDA也提供了丰富的API进行内存管理与操作，其内存的分配、初始化以及释放的API如下：
cudaMalloc(void**, size_t)
cudaMemset(void*, int, size_t)
cudaFree(void*)
其使用方法与CPU中的相应函数类似，更加具体的参数以及使用参加官方的文档：CUDA Documentation。cudaMalloc()分配的是线性内存，对应的释放内存的API是cudaFree()。线性内存也可以采用cudaMallocPitch()以及CUDAMalloc3D()来分配。这两个函数更加推荐用于2D以及3D数组的分配，这样可以保证内存的对齐要求。设备端内存在对齐访问的时候有更高的效率，这点会在后面详细GPU内存管理中进行介绍。
内存拷贝
不同方式的内存分配对应不同方式的内存拷贝的API。比如采用cudaMalloc分配的内存可以采用下面的CUDA API来在CPU与GPU之间传输数据：
cudaMemcpy ( void* dst, const void* src, size_t count, cudaMemcpyKind kind )
其中dst代表目的内存地址，src代表源内存地址，count代表需要拷贝的内存大小（bytes），kind代表数据拷贝的方向，必须是cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDevice以及cudaMemcpyDefault之一。可以看出该函数既可以实现在CPU与GPU内存之间的拷贝，也可以实现GPU内部之间的数据拷贝。需要注意的是cudaMemcpyDefault只能在支持统一虚拟地址（UVA）的系统上实现。在调用cudaMemcpy的时候，如果dst以及src的指针与拷贝的方向不一致时将会导致错误。相应的，采用 cudaMallocPitch()以及cudaMalloc3D()分配的内存，可以采用 cudaMemcpy2D()以及cudaMemcpy3D()来传输数据。
cudaMalloc是在设备端动态的分配内存，类似于CPU代码，我们也可以直接在设备端声明一个数据，这里需要用到__device__以及__constant__等标识符。下面的一段代码展示了不同访问device内存的方式：
__constant__ float constData[256];
float data[256];
cudaMemcpyToSymbol(constData, data, sizeof(data));
cudaMemcpyFromSymbol(data, constData, sizeof(data));
__device__ float devData;
float value = 3.14f;
cudaMemcpyToSymbol(devData, &value, sizeof(float));
__device__ float* devPointer;
float* ptr;
cudaMalloc(&ptr, 256 * sizeof(float));
cudaMemcpyToSymbol(devPointer, &ptr, sizeof(ptr));
需要指出的是cudaMemcpy是阻塞式的API，也就是CPU端代码在调用该API时，只有当该API完成拷贝之后，CPU才能继续处理后面的任务。这有一个好处就是保证了计算结果已经完全从GPU端拷贝到了CPU。同时CUDA也提供了非阻塞拷贝的API：cudaMemcpyAsync()， 非阻塞拷贝也称为异步拷贝，指的是该API在拷贝完成之前就返回，使得CPU可以继续处理后续的代码。异步拷贝API使得CPU与GPU之间的数据拷贝与CPU计算的并发称为可能。如果该API与CUDA中流（Stream）相结合使用，也可以实现数据的拷贝与GPU计算进行并发执行，这一点会在流与并发这一部分进行介绍。
