NVIDIA CUDA并行编程语言及其矢量相加实例——一文带你快速入门

张小殊.

已于 2023-10-25 13:42:59 修改

阅读量4.7k

点赞数 28

分类专栏： # 并行编程模型-入门文章标签：人工智能 c++ linux bash 并行编程 CUDA编程

于 2023-10-10 19:36:56 首次发布

本文链接：https://blog.csdn.net/qq_46009046/article/details/133753993

版权

并行编程模型-入门专栏收录该内容

9 篇文章 7 订阅

订阅专栏

　　✍️CUDA 是 2007 年由NVIDIA 公司推出的只能运行在本公司各种型号 GPU 上的并行编程语言，使用扩展的 C 语言来进行 GPU 编程。自 2007 年 CUDA 1.0 版本诞生后，由于大大降低了 GPU 通用编程的难度，因此大量的研究者尝试利用 GPU 加速各个领域的算法。此后 CUDA 版本快速迭代，通用计算能力越来越强，今天 CUDA 已经被广泛应用于人工智能计算系统，这篇文章可以带大家快速入门。

　　🥳先沾上目录：

概述&背景

CUDA Runtime API与CUDA Driver API

　　CUDA并行编程的官方文档可以参考以下链接：

https://docs.nvidia.com/cuda/index.htmlhttps://docs.nvidia.com/cuda/index.html

概述&背景

　　CUDA的软件堆栈由以下三层构成：CUDA Library、CUDA Runtime API、CUDA Driver API，如图1所示，CUDA的核心是CUDA C语言，它包含对C语言的最小扩展集和一个运行时库，使用这些扩展和运行时库的源文件必须通过nvcc编译器进行编译。

图1 CUDA软件栈示意图

　　CUDA C语言编译得到的只是GPU端代码，而要管理分配GPU资源，在GPU上分配显存并启动内核函数，就必须借助CUDA运行时的API（Runtime API)或者CUDA驱动API（Driver API）来实现。在一个程序中只能使用CUDA运行时API与CUDA驱动API中的一种，不能混和使用。

CUDA Runtime API与CUDA Driver API

　　CUDA运行时API（Runtime API）和CUDA驱动API（Driver API）提供了实现设备管理（Device management)，上下文管理（Context management），存储器管理费用（Memory Control），代码块管理（Code Module management），执行控制（Excution Control），纹理索引管理（Texture Reference management）与OpenGL和Direct3D的互操作性（Interoperity with OpenGL and Direct3D）的应用程序接口。

　　 CUDARuntime API在CUDA Driver API 的基础上进行了封装，隐藏了一些实现细节，编程更加方便，代码更加简洁。CUDARuntime API被打包放在CUDAArt包里，其中的函数都有CUDA 前缀。CUDA运行时没有专门的初始化函数，它将在第一次调用函数时自动完成初始化。对使用运行时函数的CUDA程序测试时要避免将这段初始化的时间计入。CUDARuntime API的编程较为简洁，通常都会用这种API进行开发。

　　CUDA Driver API是一种基于句柄的底层接口（式多对象通过句柄被引用），可以加载二进制或汇编形式的内核函数模块，指定参数，并启动计算。CUDA Driver API的编程复杂，但有时能通过直接操作硬件的执行实行一些更加复杂的功能键，或者获得更高的性能。由于它使用的设备端代码是二进制或者汇编代码，因此可以在各种语言中调用。CUDA Driver API被放在nvCUDA包里，所有函数前缀为cu。

CUDA Library

　　CUDA Library（CUDA函数库），目前CUDA中有cuBLAS、cuSPARSE、cuFFT、cuDNN、视频编解码与图像处理库和CUDPP等函数库。CUBLAS库是一个基本的矩阵与向量运算库，提供了与BLAS相似的接口，可以用于简单的矩阵计算，也可以作为基础构建更加复杂的函数包，如LAPACK等，cuBLAS操作的数据也存储在显存中，同样需要封装后才能替代BLAS中的函数。cuSPARSE是一个线性代数库，内含很多通用地稀疏线性代数函数，这些函数支持一系列稠密和稀疏的数据格式。CUFFT是利用GPU进行傅立叶变换的函数库，提供了与广泛使用的FFTW库相似的接口。不同的是FFTW操作的数据存储在内在中，而CUFFT操作的数据存储在显存，不能直接相互取代，必须加入显存与内存之间的数据交换，进行封装后才能替代FFTW库。NVIDIA CUDA Deep Neural Network （cuDNN）库是一个用于深度神经网络的 GPU 加速原语库。cuDNN 为标准例程（如前向和后向卷积、池化、规范化和激活层）提供了高度调优的实现。CUDPP为提供了很多基本的常州用的并行操作灵敏，如排序、搜索等，可以作为基本组件快速地搭建出并行计算程序。调用上述函数库使得程序员无须按照硬件特性设计复杂的算法就能获得很高的性能，大大缩短开发时间。

CUDA编程模型

　　CUDA是一种基于NVIDIA CUDA生态的并行计算平台和编程语言。你可以像编写C或C++语言程序一样实现算法和程序的并行性并在在NVIDIA GPU上运行，本文我们将了解如何编写一个CDUA并行程序。

CUDA编程结构

　　通常我们在编写的C/C++程序时，都是按照既定的问题解决方式设计相应的算法然后用计算机语言去实现这种逻辑功能，然而，对于一些问题可能涉及大量的重复计算，例如，假如我们要计算两个长度为N的向量序列V1、V2的和，使用C语言编写程序时需要定义循环体实现向量序列V1、V2中对应元素的和，循环体中的算法是相同的，每一次的求和过程并不依赖于前一次的结果，因此，若能将串行实现的循环体改变成并行执行的程序将会大大提高算法的运行速度，即，使用CUDA编程语言实现相同逻辑算法的并行执行。

　　相较于串行代码，在CUDA编程模型中引入了主机端和设备端的概念，其中CPU作为主机端（Host），GPU作为设备端（Device）。CPU负责任务的调度，数据的传输、逻辑处理以及运算量少的计算， GPU硬件主要通过“CUDA核”进行并行计算。将在 CPU 上执行的代码称为主机代码，在 GPU 上运行的代码称为设备代码。设备端代码又称为核函数（kernel）。CUDA写的设备端源文件一般以“.cu”为后缀，Host端的串行程序和Device端的并行程序可以各自独立运行，如图2所示。GPU程序可以异步执行，当CUDA程序中在GPU开始执行后，程序的流程控制权立刻交还给Host端串行程序，即CPU可以在GPU进行大规模并行运算时进行串行运算，提高异构设备的运行效率。

图2 CUDA程序执行流程

　　在实际的并行代码开发中，在主机端，开发者通常基于应用程序数据以及GPU的性能设计如何让实现算法的功能，而在设备端，开发者将大量需要并行的任务通过编写内核函数实现多线程并行计算。CUDA程序实现流程基本可以总结为以下三点：

把数据从CPU内存拷贝到GPU内存；
调用和核函数对储存在GPU内存中的数据进行操作；
将数据从GPU内存传送回CPU内存；

　　对CDUA并行编程程序的执行流程有一个清晰的认识后，我们便可以了解如何编写一个CUDA编写程序，从上述执行流程中可以看到，代码的设计主要涉及主机端和设备端的数据传输、并行内核程序的编写和调度等，下面将分别进行讲述。

内存和线程管理

　　前文提到，在CDUA并行程序的执行流程中很关键的一步就是数据在主机端和设备端之间的传输，在CUDA中提供了相应的API去实现这种功能，但是，数据在不同的设备上进行传输之前，还需要进行一些“准备工作”，这便是开辟内存，

　　在编写C/C++语言的过程中，往往会涉及到需要开辟动态内存来进行一些数据的存储，进行CUDA并行编程时，同样需要开辟动态内存来进行数据存储，相较于串行程序不同的是，并行程序涉及了主机端和设备端，而它们的动态内存开辟方式也是不同的，如下表1所示，列出了标准的C函数以及相对应的CUDA并行编程模型设备端内存操作的API函数。

表1 内存申请和释放函数

标准C函数（Host端）	CUDA（device端）
Malloc()	cudaMalloc()
free()	cudaFree()

　　从上述表格中可以看到。CUDA并行编程语言的内存分配和管理与标准C函数基本相同，只是前面需要加上cdua前缀。下面我们具体的来看一下上述表格中内存分配和释放API的一个实例，假设要为一个有N个浮点类型元素的数组分配内存，主机的代码变量名命名以h_为前缀。而设备端的代码变量名命名以d_为前缀，主机端和设备端的分配方式如下代码所示：

size_t size = N * sizeof(float);
float *h_A= NULL, *d_A = NULL;
float *h_A = (float *)malloc(size);
cudaMalloc((void **)&d_A, size);

　　为主机端和设备端的数据分配内存后。我们便可以借用CUDA提供的API实现主机端的数据传输到设备端预先分配好的内存中，从而在设备端对数据进行计算。这一操作通过cudaMemcpy进行实现，其定义如下所示：

cudaError_t cduaMemcpy(void *dist, const void* src,size_t count, cudaMemcpyKind kind)

　　其中此函数从src指向的源储存区复制一定数量的字节到dist指定的目标储存区，复制的方向由kind指定，其中的kind共有四种，表2给出了数据传输和通信对应kind类型。

表2 数据传输和通信对应kind类型

数据通信方向	kind
主机端到设备端	cudaMemcpyHostToDevice
设备端到主机端	cudaMemcpyDeviceToHost
主机端到主机端	cudaMemcpyHostToHost
设备端到设备端	cudaMemcpyDeviceToDevice

　　如上表2所示。将数据从主机端传到设备端和从设备端传到主机端是不同的kind。因此在具体使用的过程中，注意不能混淆顺序。数据从主机端传到设备端和从设备端传到主机端的具体代码示例如下所示：

cudaMemcpy(d_A,h_A,size,hipMemcpyHostToDevice);
cudaMemcpy(h_A,d_A,size,hipMemcpyDeviceToHost);

　　上述代码的第一句将h_A的数据从CPU端传输到设备端的d_A上，而第二句将将d_A的数据从GPU端传输到设备端的h_A上，通过以上的数据传输API，我们便可以很容易的控制数据在主机端与设备端的通信和传输。另外，在一个并行CUDA程序中，有关内存的操作除了内存分配以及数据传输，程序的最后一定不能忘记释放在程序中申请的内存空间，其中，主机端和设备端的内存释放如下代码所示：

cudaFree(d_A);
free(h_A);

　　对内存分配和传输有一个清晰的认识后，我们接着来看在设备端的计算任务如何实现和硬件资源的映射以及内核程序的编写和设计。

　　CUDA编程模型使用GPU的众核实现并行运算。在CUDA编程模型中，通过众多并行执行的细粒度线程来执行计算任务，CUDA的线程组织分三层结构：最外层是线程网格（Grid）、中间层为线程块（block）、最内层为线程（Thread），如图3所示。一个Grid包含多个Block，这些Blocks的组织方式可以是一维、二维或者三维。任何一个Block包含有多个线程，这些线程的组织方式也可以是一维、二维或者三维。因此定义ThreadIdx、blockIdx、blockDim和gridDim为uint3类型, uint3是一个包含了三个整数的整型向量类型。ThreadIdx表示一个线程的索引（一个线程的ID），blockIdx是一个线程块的索引ID，blockDim表示线程块的大小。gridDim表示网格的大小，即一个网格中有多少个线程块。CUDA 核函数能够识别如下两种索引的特殊变量：执行核函数的线程索引和线程所在的线程块索引。当线程组织为一维结构时这两个变量分别为 threadIdx.x 和 blockIdx.x。

图3 CUDA线程的组织结构

　　CUDA中每一个线程都有一个唯一标识ThreadIdx，从 0 开始，这个ID随着线程组织结构形式的变化而变化，每个线程块也会被分配一个索引，也从 0 开始。例如Block是一维的，Thread也是一维时ThreadIdx的计算公式为：

$eq?ThreadIdx.x%3D%20blockIdx.x%20*blockDim.x%20+%20threadIdx.x$

　　在CUDA编程模型中，采用了两级并行机制，分别是block层和thread层，对应的block层映射到SM，thread层映射到SP或者CUDA Core，block内的thread可以通过共享存储器和同步机制进行交互。一般情况下，kernel中设置的block数量和thread数量分别大于GPU硬件中SM和CUDA Core数量。Block和硬件的映射关系如图3-3所示。从图中可以看出，一个GPU拥有的SM越多，执行时就有更多的block处于并行计算，那么计算速度就更快，因此这种映射具有良好的并行可扩展性。

　　总之，GPU工作时的最小单位是 thread，多个 thread 可以组成一个 block，但每一个 block 所能包含的 thread 数目是有限的。执行相同程序的多个 block，可以组成grid。

核函数的启动和编写

　　在CUDA并行编程中，一个kernel函数对应一个Grid，每个Grid根据求解问题规模配置不同的Block数量和Thread数量，下面列出了一个kernel函数的示例，这个例子中Grid由128个Block组成，每个Block有N个线程。__global__ 关键字表明函数将在 GPU 上运行, 使用其定义的函数需要返回 void 类型。Input表示需要从内存传输到GPU显存上的参数,output表示需要从GPU显存传输到内存的参数。

__global__void KernelFunction( float *input, float *output）{
每个线程执行的并行计算程序；
}
int main（）
{
KernelFunction<<<128,N>>>（input,output）；7.	}

　　在上述代码块中，kernel 是函数名，“<<< >>>”中参数的告诉系统使用什么样的网格启动核函数(设置线程)。 Param是内存与GPU显存交互的数据。在主机端调用核函数采用如下的形式为：

kernel<<<Dgrid,Dblock >>>(param list);

Dgrid： int型或者dim3类型(x,y,z)。用于定义一个grid中的block是如何组织的。其中int型表示为1维结构。

Dblock： int型或者dim3类型(x,y,z)。用于定义一个block中的thread是如何组织的。其中int型表示为1维结构。

　　在核函数实际运行过程中，Block会被分为更小的warp（线程束）。一般情况下warp的值为32。在硬件实际运行程序时，线程的数量以warp为单位开启，在设计一个算法的block线程数量时必须要考虑其影响。表3说明了CUDA编程概念和GPU硬件之间对应关系。

表3 CUDA编程模型概念和GPU硬件之间对应关系

CUDA概念	GPU硬件
网格Grid	GPU
线程块Block	流多处理器SM
线程Thread	流处理器SP（CUDA core）
Warp	一组同时执行的32个线程

CUDA并行矢量相加实例　　

　　下面我们来考虑一个简单的例子。假如我们要实现两个大小为numElements的向量相加。如果使用CPU进行串行程序设计。其代码如下所示：

void vectorAdd(float *h_A,float  *h_B,float  *h_C,int numElements)
{　
　　for(int i = 0; i < numElements; i++)
　　h_C[i] = h_A[i] + h_B[i];	 
}

　　如果使用GPU进行并行程序设计。核函数如下所示：

__global__ void vectorAdd(float *d_A,float  *d_B,float  *d_C,int numElements)
{
    int i = cudaBlockIdx_x * cudaBlockDim_x + cudaThreadIdx_x;
    if(i<numElements)
    {
    d_C[i] = d_A[i] + d_B[i];
　　}
}

　　观察上述两个函数的代码。我们可以发现在GPU上设计的核函数没有了循环体，向量的索引通过不同线程的全局索引来进行确定，因此。当我们所开辟出来的线程数大于所要计算的数组元素时，我们需要用if语句来进行判断。否则会导致数组越界的错误。

　　根据前几节中的内容，我们为此核函数编写CPU端程序，使之成为一个完整的CUDA并行代码，代码如下所示：

#include <stdio.h>

// For the CUDA runtime routines (prefixed with "cuda_")
#include <cuda_runtime.h>

/**
 * CUDA Kernel Device code
 *
 * Computes the vector addition of A and B into C. The 3 vectors have the same
 * number of elements numElements.
 */
__global__ void
vectorAdd(const float *A, const float *B, float *C, int numElements)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;

    if (i < numElements)
    {
        C[i] = A[i] + B[i];
    }
}

/**
 * Host main routine
 */
int
main(void)
{
    // Error code to check return values for CUDA calls
    cudaError_t err = cudaSuccess;

    // Print the vector length to be used, and compute its size
    int numElements = 50000;
    size_t size = numElements * sizeof(float);
    printf("[Vector addition of %d elements]\n", numElements);

    // Allocate the host input vector A
    float *h_A = (float *)malloc(size);

    // Allocate the host input vector B
    float *h_B = (float *)malloc(size);

    // Allocate the host output vector C
    float *h_C = (float *)malloc(size);

    // Verify that allocations succeeded
    if (h_A == NULL || h_B == NULL || h_C == NULL)
    {
        fprintf(stderr, "Failed to allocate host vectors!\n");
        exit(EXIT_FAILURE);
    }

    // Initialize the host input vectors
    for (int i = 0; i < numElements; ++i)
    {
        h_A[i] = rand()/(float)RAND_MAX;
        h_B[i] = rand()/(float)RAND_MAX;
    }

    // Allocate the device input vector A
    float *d_A = NULL;
    err = cudaMalloc((void **)&d_A, size);

    if (err != cudaSuccess)
    {
        fprintf(stderr, "Failed to allocate device vector A (error code %s)!\n", cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    // Allocate the device input vector B
    float *d_B = NULL;
    err = cudaMalloc((void **)&d_B, size);

    if (err != cudaSuccess)
    {
        fprintf(stderr, "Failed to allocate device vector B (error code %s)!\n", cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    // Allocate the device output vector C
    float *d_C = NULL;
    err = cudaMalloc((void **)&d_C, size);

    if (err != cudaSuccess)
    {
        fprintf(stderr, "Failed to allocate device vector C (error code %s)!\n", cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    // Copy the host input vectors A and B in host memory to the device input vectors in
    // device memory
    printf("Copy input data from the host memory to the CUDA device\n");
    err = cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);

    if (err != cudaSuccess)
    {
        fprintf(stderr, "Failed to copy vector A from host to device (error code %s)!\n", cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    err = cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    if (err != cudaSuccess)
    {
        fprintf(stderr, "Failed to copy vector B from host to device (error code %s)!\n", cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    // Launch the Vector Add CUDA Kernel
    int threadsPerBlock = 256;
    int blocksPerGrid =(numElements + threadsPerBlock - 1) / threadsPerBlock;
    printf("CUDA kernel launch with %d blocks of %d threads\n", blocksPerGrid, threadsPerBlock);
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);
    err = cudaGetLastError();

    if (err != cudaSuccess)
    {
        fprintf(stderr, "Failed to launch vectorAdd kernel (error code %s)!\n", cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    // Copy the device result vector in device memory to the host result vector
    // in host memory.
    printf("Copy output data from the CUDA device to the host memory\n");
    err = cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    if (err != cudaSuccess)
    {
        fprintf(stderr, "Failed to copy vector C from device to host (error code %s)!\n", cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    // Verify that the result vector is correct
    for (int i = 0; i < numElements; ++i)
    {
        if (fabs(h_A[i] + h_B[i] - h_C[i]) > 1e-5)
        {
            fprintf(stderr, "Result verification failed at element %d!\n", i);
            exit(EXIT_FAILURE);
        }
    }

    printf("Test PASSED\n");

    // Free device global memory
    err = cudaFree(d_A);

    if (err != cudaSuccess)
    {
        fprintf(stderr, "Failed to free device vector A (error code %s)!\n", cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    err = cudaFree(d_B);

    if (err != cudaSuccess)
    {
        fprintf(stderr, "Failed to free device vector B (error code %s)!\n", cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    err = cudaFree(d_C);

    if (err != cudaSuccess)
    {
        fprintf(stderr, "Failed to free device vector C (error code %s)!\n", cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    // Free host memory
    free(h_A);
    free(h_B);
    free(h_C);

    // Reset the device and exit
    // cudaDeviceReset causes the driver to clean up all state. While
    // not mandatory in normal operation, it is good practice.  It is also
    // needed to ensure correct operation when the application is being
    // profiled. Calling cudaDeviceReset causes all profile data to be
    // flushed before the application exits
    err = cudaDeviceReset();

    if (err != cudaSuccess)
    {
        fprintf(stderr, "Failed to deinitialize the device! error=%s\n", cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    printf("Done\n");
    return 0;
}

　　上述代码便是实现两个大小为numElements的向量相加的完整并行代码。编写完程序后。我们便可以对这个代码进行编译和执行，具体的编译和运行指令为nvcc cuda_vectorAdd.cu -o cuda_vectorAdd和./ cuda_vectorAdd。

张小殊.

关注

28
点赞
踩
32

收藏

觉得还不错? 一键收藏
打赏
1
评论
NVIDIA CUDA并行编程语言及其矢量相加实例——一文带你快速入门

CUDA的软件堆栈由以下三层构成：如图1所示，CUDA的核心是CUDA C语言，它包含对C语言的最小扩展集和一个运行时库，使用这些扩展和运行时库的源文件必须通过nvcc编译器进行编译。图1 CUDA软件栈示意图CUDA C语言编译得到的只是GPU端代码，而要管理分配GPU资源，在GPU上分配显存并启动内核函数，就必须借助CUDA运行时的API或者CUDA驱动APIDriver API）来实现。在一个程序中只能使用CUDA运行时API与。
复制链接

扫一扫