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;
       运行结果：
                                    AI绘图：使用Stable Diffusion Prompt(提示词)语法
                                    AI绘图：使用Stable Diffusion WebUI绘制第一幅AI图
                                    华硕路由器RT-AC88U增加温度控制风扇
                                    CUDA GPU编程指南：运行时的错误处理
                                    CUDA GPU编程指南：统一内存
                华硕路由器RT-AC88U增加温度控制风扇AI绘图：使用Stable Diffusion WebUI绘制第一幅AI图AI绘图：使用Stable Diffusion Prompt(提示词)语法            
                                    使用QT 实现等值线图、Contour图的图表着色功能
                                    QT多线程之QtConcurrent::run()
                                    FreeCAD源码分析之文件打开
                                    CUDA GPU编程指南：运行时的错误处理
                                    VTK的可视化流水线（渲染管线）
                    lct 发布于 1 周前（10月30日）
                    鼠标划过便利贴窗口事是否弹出菜单  ---这个配置项对于 “列表”类型便笺不生效，对普通类型便笺则生...