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

2023年7月26日 2742点热度 0人点赞 2条评论

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;
}

​ 运行结果:

image-20230725213456780

大脸猫

这个人虽然很勤快,但什么也没有留下!

文章评论

  • RsYes

    好文,就是有多处 typo:现存 -> 显存,和函数 -> 核函数。

    2023年7月27日