CUDA GPU编程指南：内存的申请，释放与拷贝

前言

​ 在CPU上运行的程序，会将数据存放在内存中，在C++中直接使用new delete 或者malloc free就可以管理内存的申请和释放。

​ 在GPU上运行的程序，会将数据放入显存中，需要CUDA提供的API来管理内存，并且在显存中存放的数据，是不能被CPU直接访问的。

主机内存（host）和设备内存（device）

​ 主机内存指在CPU上运行的代码所使用的内存，正常C/C++代码中的使用的变量，无论是在堆还是栈上，使用的都是主机内存。

​ 使用设备内存指在GPU上运行的代码可以使用到的内存。

​ 需要注意的是，核函数和device函数中是不能访问和读取主机内存的。这意味着，在调用核函数时，将数据（主机内存）指针当作参数传入核函数，在核函数中是无法读取到数据的，因为参数的指针指向的是主机内存的地址，而在核函数中读写会被当做设备内存来解析。

​ 因此需要掌握在CPU代码中如何管理设备内存的方法。

在CPU代码中管理设备内存

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

​ CUDA提供cudaMalloc 和 cudaFree 来管理设备内存的申请和释放，用法与C语言中的malloc 和 free相似，只是传递数据的方式不同。

​ 使用时需要注意cudaMalloc 的第一个参数是一个二级指针，所以传参的时候一定要注意。

    int *arry = 0;
    int arrySize = 100;
    //申请一个长度为100的int数组
    cudaMalloc((void**)&arry, size * sizeof(int));
    //释放内存
    cudaFree(arry);

​ 如上述代码，完成了一个int数组的内存的申请与释放。

数据拷贝

​ 上文提到，在核函数中无法读写主机内存中的数据，那么同理，在CPU上运行的代码，也无法读写在设备内存中的数据。

​ 但是在实际应用场景中，一定会需要到两边数据的交互，那么怎么实现？

​ CUDA提供了数据拷贝函数，以供CPU操作设备内存的数据。

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

​ dst: 要拷贝到的目标地址

​ src: 要拷贝的源地址

​ count：要拷贝的长度，单位是字节。

​ kind：拷贝数据的方式。

​ cudaMencpykind的定义是：

enum __device_builtin__ cudaMemcpyKind
{
    cudaMemcpyHostToHost          =   0,      /**< Host   -> Host */
    cudaMemcpyHostToDevice        =   1,      /**< Host   -> Device */
    cudaMemcpyDeviceToHost        =   2,      /**< Device -> Host */
    cudaMemcpyDeviceToDevice      =   3,      /**< Device -> Device */
    cudaMemcpyDefault             =   4       /**< Direction of the transfer is inferred from the pointer values. Requires unified virtual addressing */
};

​ cudaMemcpyHostToHost ：主机 拷贝到 主机
​ cudaMemcpyHostToDevice ：主机拷贝到设备
​ cudaMemcpyDeviceToHost ：设备拷贝到主机
​ cudaMemcpyDeviceToDevice ：设备拷贝到设备
​ cudaMemcpyDefault ：拷贝方式通过指针的地址来推断。

​ 综上所述，cudaMemcpy提供了多种数据拷贝的方式。

​ 使用时需要注意的一点是，在调用数据拷贝函数的时候，CUDA会自动等待同步，即等到GPU中的任务执行完成之后才会真正的开始拷贝数据。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

__global__ void printArray(int* array, const int* size)
{
    printf("device memory array : ");
    for (int i = 0; i < *size; i++)
    {
        printf(" %d", array[i]);
    }
    printf("\n");

    for (int i = 0; i < *size; i++)
    {
        array[i] = 5 - i;
    }

}

int main()
{
    const int array_size = 5;

    int array_a[array_size] = { 1, 2, 3, 4, 5 };
    int *array_size_device;
    int* array_a_device;
    //申请设备内存
    cudaMalloc(&array_size_device,sizeof(int));
    cudaMalloc(&array_a_device, sizeof(int) * array_size);
    //将主机内存上的数据拷贝到设备内存
    cudaMemcpy(array_size_device, &array_size, sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(array_a_device, array_a, sizeof(int) * array_size,cudaMemcpyHostToDevice);
    //调用核函数，打印数组的值，并改变数组的值
    printArray <<<1, 1 >>> (array_a_device, array_size_device);
    int e = cudaGetLastError();
    //将设备内存上的数据拷贝回主机内存
    cudaMemcpy(array_a, array_a_device,  sizeof(int) * array_size, cudaMemcpyDeviceToHost);

    printf("host memory arry: ");
    //打印主机内存上的数组
    for (int i = 0; i < array_size; i++)
    {
        printf(" %d", array_a[i]);
    }
    printf("\n");

    cudaFree(array_a_device);
    return 0;
}

​ 运行结果：