Cuda实战-02 向量加法与内存分配

概述

        本节通过一个向量加法的实例程序探讨如何进行有效的内存分配，案例中定义了三个向量VectorA，VectorB，VectorC。VectorA和VectorB为待计算的两个向量，VectorC为结果向量。VectorA初始值为(1, 2, 3, 4, 5)，VectorB初始值为(1, 1, 1, 1, 1)，VectorC初始值为(0, 0, 0, 0, 0)。经过正确运算后，结果向量VectorC的值应该为(2, 3, 4, 5, 6)。
        整个过程中，涉及了许多的内存分配和释放操作，得重点关注。

程序

//C++标准库
#include <iostream>
//Cuda运行库
#include <cuda_runtime.h>

using namespace std;

//核函数，GPU执行向量加法
__global__ void add(int* VectorA, int* VectorB, int* VectorC){
    int index = threadIdx.x;
    VectorC[index] = VectorA[index] + VectorB[index];
}

//主函数，CPU执行数据初始化
int main(){
    //CPU中定义三个向量
    int Vector_A[5] = { 1, 2, 3, 4, 5 };
    int Vector_B[5] = { 1, 1, 1, 1, 1 };
    int Vector_C[5] = { 0, 0, 0, 0, 0 };
    //定义GPU中的三个向量，D代表Device
    int* D_Vector_A, * D_Vector_B, * D_Vector_C;
    
    //计算占用空间
    int size = 5 * sizeof(int);
    //GPU分配内存
    cudaMalloc((void**)&D_Vector_A, size);
    cudaMalloc((void**)&D_Vector_B, size);
    cudaMalloc((void**)&D_Vector_C, size);
    //将数据从CPU复制到GPU
    cudaMemcpy(D_Vector_A, Vector_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(D_Vector_B, Vector_B, size, cudaMemcpyHostToDevice);
    
    //启用GPU执行任务
    add << <1, 5 >> > (D_Vector_A, D_Vector_B, D_Vector_C);
    //将GPU中运算结果复制到CPU
    cudaMemcpy(Vector_C, D_Vector_C, size, cudaMemcpyDeviceToHost);
    //将运算结果打印出来
    cout << "运算结果为：" << endl << "(";
    for (int i = 0; i < 5; i++) cout << Vector_C[i] << "    ";
    cout << ")" << endl;
    //检查一下是否有错误提示
    cout << cudaGetErrorString(cudaGetLastError()) << endl;
    
    //释放GPU内存
    cudaFree(D_Vector_A);
    cudaFree(D_Vector_B);
    cudaFree(D_Vector_C);

    return 0;
}

输出

讲解

        CUDA编程的内存管理与C语言的类似，需要程序员显式管理主机和设备之间的数据移动。随着CUDA版本的升级，NVIDIA正系统地实现主机和设备内存空间的统一，但对于大多数应用程序来说，仍需要手动移动数据。对于CUDA内存管理来说，工作重点在于如何使用CUDA函数来显式地管理内存和数据移动，主要是两个方面：分配和释放设备内存；在主机和设备之间传输数据。为了达到最优性能，CUDA提供了在主机端准备设备内存的函数，并且显式地向设备传输数据和从设备中获取数据。
        CUDA编程模型假设了一个包含一个主机和一个设备的异构系统，每一个异构系统都有自己独立的内存空间。核函数在设备内存空间中运行，CUDA运行时提供函数以分配和释放设备内存。在实际操作中，可以使用cudaError_t cudaMalloc(void **devPtr, size_t count)函数分配全局内存，这个函数在设备上分配了count字节的全局内存，并用devptr指针返回该内存的地址。所分配的内存支持任何变量类型，包含整型、浮点类型变量、布尔类型等。如果cudaMalloc函数执行失败则返回cudaErrorMemoryAllocation。在已分配的全局内存中的值不会被清除。
        你需要用从主机上传输的准备特定数据来填充所分配的全局内存，可用下列函数将其初始化：cudaError_t cudaMemcpy(void *devPtr, void *value, size_t count),这个函数用存储在CPU内存中value的值来填充设备内存地址devPtr处开始的count个字节。一旦一个应用程序不再使用已分配的全局内存，那么可以使用以下代码释放该内存空间：cudaError_t cudaFree(void *devPtr);这个函数释放了devPtr指向的全局内存，该内存必须在此前使用了一个设备分配函数（cudaMalloc）来进行分配。否则，它将返回一个错误cudaErrorInvalidDevicePointer。如果地址空间已经被释放，那么cudaFree也返回一个错误。设备内存的分配和释放操作成本较高，所以应用程序应重利用设备内存，以减少对整体性能的影响。
        一旦分配好了全局内存，就可以使用cudaError_t cudaMemcpy(void *dst,const void *src,size_t count,enum cudaMemcpyKind kind)函数从主机向设备传输数据，这个函数从内存weizhisrc复制了count字节到内存位置dst。变量kind指定了复制的方向，可以有下列取值；cudaMemcpyHostToHost、cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost、cudaMemcpyDeviceToDevice。
        CUDA运行时允许使用如下指令直接分配固定主机内存：cudaError_t cudaMallocHost(void **devPtr,size_t count);这个函数分配了count字节的主机内存，这些内存是页面锁定的并且对设备来说是可访问的。由于固定内存能被设备直接访问，所以它能用比分页内存高得多的带宽进行读写。然而，分配过多的固定内存可能会降低主机系统的性能，因为它减少了用于存储虚拟内存数据的可分页内存的数量，其中分页内存对主机系统是可用的。固定主机内存通过下列指令来释放：cudaError_t cudaFreeHost(void *ptr);