CUDA编程模型概述

CUDA是一种通用的并行计算平台和编程模型，是在C语言上扩展的的。借助于CUDA，可以像编写C语言程序一样实现并行算法。
CUDA编程模型提供了一个计算机架构抽象作为应用程序和其可用硬件之间的桥梁。下图说明了程序和编程模型实现之间的抽象结构。其中通信抽象是程序与编程模型实现之间的分界线，它通过专业的硬件原语和操作系统的编译器或库来实现。利用编程模型所编写的程序指定了程序的各组成部分是如何共享信息及相互协作的。编程模型从逻辑上提供了一个特定的计算机架构，通常它体现在编程语言或编程环境中。

除了与其他并行编程模型共有的抽象外，CUDA编程模型还利用GPU架构的计算能力提供了以下特有功能：
>一种通过层次结构在GPU中组织线程的方法；
>一种通过层次结构在GPU中访问内存的方法。

CUDA 编程结构
CUDA编程模型使用由C语言扩展生成的注释代码在异构计算系统中执行应用程序。在一个异构环境中包含多个CPU和GPU，每个GPU和CPU的内存都由一条PCI-Express总线分隔开，因此，需要注意区分主机和设备这两个内容：
>主机：CPU及其内存（主机内存）；
>设备：GPU及其内存（设备内存）。
为了清楚地指明不同地内存空间，主机内存中的变量名与设备内存中的变量名一般使用不同的前缀。
内核（kernel）是一个CUDA编程模型的一个重要组成部分，其代码在GPU上运行。多数情况下，主机可以独立地对设备进行操作。内核一旦被启动，管理权立刻返回给主机，释放CPU来执行由设备上运行地并行代码实现地额外的任务。CUDA编程模型主要是异步的，因此在GPU上进行的运算可以与主机-设备通信重叠。一个典型的CUDA程序包括并行代码互补的串行代码。如下图所示，串行代码在CPU上执行，而并行代码在GPU上执行。主机代码按照ANSI C标准进行编写，而设备代码使用CUDA C进行编写。NVIDIA的C编译器（nvcc）为主机和设备生成可执行代码。

一个典型的CUDA程序实现流程遵循以下模式：
>1.把数据从CPU内存拷贝到GPU内存；
>2.调用核函数对存储在GPU内存中的数据进行操作。
>3.将数据从GPU内存传送回到CPU内存。

内存管理
为了拥有充分的控制权并使系统达到最佳性能，CUDA运行时负责分配与释放设备内存，并且在主机内存和设备内存之间传输数据。C语言以及相应的针对内存操作的CUDA C函数如下表所示：

标准的C函数	CUDA C函数	标准的C函数	CUDA C函数
malloc	cudaMalloc	memset	cudaMemset
memcpy	cudaMemcpy	free	cudaFree

cudaMalloc与标准C语言的malloc函数几乎一样，只是此函数在GPU的内存里分配内存。
cudaMemcpy函数负责主机和设备之间的数据传输，其函数原型为：cudaError_t cudaMemcpy(void* dst,const void* src,size_t count, cudaMemcpyKind kind);此函数从src指向的源存储区复制一定数量的字节到dst指向的目标存储区。复制方向由kind指定，其中的kind有以下几种：
>cudaMemcpyHostToHost
>cudaMemcpyHostToDevice
>cudaMemcpyDeviceToHost
>cudaMemcpyDeviceToDevice
这个函数以同步方式执行，因为在cudaMemcpy函数返回以及传输操作完成之前主机应用程序是阻塞的。除了内核启动之外的CUDA调用都会返回一个错误的枚举类型cudaError_t。如果GPU内存分配成功，函数返回cudaSuccess，否则返回cudaErrorMemoryAllocation。
CUDA编程模型从GPU架构中抽象出一个内存层次结构。如下图所示，主要包含两部分：全局内存和共享内存。

线程管理
当核函数在主机端启动时，它的执行会移动到设备上，此时设备中会产生大量的线程，并且每个线程都执行由核函数指定的语句。一个两层的线程层次结构如下图所示，由线程块和线程块网格构成。

又一个内核所产生的所有线程统称为一个网格。同一网格中的所有线程共享相同的全局内存空间。一个网格由多个线程块构成，一个线程块包含一组线程，同一线程块内的线程协作可以通过同步、共享内存等方式来实现，而不同块内的线程不能协作。
线程依靠blockIdx(线程块在线程格内的索引)以及threadIdx(块内的线程索引)这两个坐标变量来区分彼此。这些变量是核函数中需要预初始化的内置变量。当执行一个核函数时，CUDA运行时为每个线程分配坐标变量blockIdx和threadIdx。基于这些坐标，我们可以将部分数据分配给不同的线程。该坐标变量是基于uint3定义的CUDA内置的向量类型，是一个包含3个无符号整数的结构，可以通过x、y、z三个字段来指定。
CUDA可以组织三维的网格和块。线程层次结构如上面的图所示，其结构是一个包含二维块的二维网格。网格和块的维度由blockDim(线程块的维度，用每个线程块中的线程数来表示)、gridDim(线程格的维度，用每个线程格中的线程数来表示)两个内置变量指定，它们是dim3类型的变量，是基于uint3定义的整数型向量，用来表示维度。当定义一个dim3类型的变量时，所有未指定的元素都被初始化为1.dim3类型变量中的每个组件可以通过它的x、y、z字段获得。
在CUDA程序中有两组不同的网格和块变量：手动定义的dim3数据类型和预定义的uint3数据类型。在主机端，作为内核调用的一部分，可以使用dim3数据类型定义一个网格和块的维度。当执行核函数时， CUDA运行时会生成相应的内置预初始化的网格、块和线程变量，它们在核函数内均可被访问到且为uint3类型。手动定义的dim3类型的网格和块变量仅在主机端可见，而uint3类型的内置预初始化的网格和块变量仅在设备端可见。
由于一个内核启动的网格和块的维数会影响性能，这一结构为程序员优化程序提供了一个额外的途径。对于一个给定的数据大小，确定网格和块尺寸的一般步骤为：确定块的个数、在已知数据大小和块大小的基础上计算网格维度。网格和块的维度存在几个限制因素：对于块大小的一个主要限制因素就是可利用的计算资源，如寄存器，共享内存等。某些限制可以通过查询GPU设备。

启动一个 CUDA 核函数
C语言函数调用语句：function_name(argument list);而CUDA内核调用是对C语言函数调用语句的延申，kernal_name <<<grid,block>>>(argument list);其中<<<>>>运算符内是核函数的执行配置。执行配置的第一个值是网络维度，也就是启动块的数目，第二个值是块维度，也就是每个块中线程的数目。通过指定网格和块的维度，可以进行内核中线程的数目以及内核中使用的线程布局的配置。
由于数据在全局内存中是线性存储的，因为可以用变量blockIdx.x和threadIdx.x来进行以下操作：在网格中标识一个唯一的线程、建立线程和数据元素之间的映射关系。不同于C语言的函数调用，所有的CUDA核函数的启动都是异步的。CUDA内核调用完成后，控制权立刻返回给CPU。可以通过调用以下函数来强制主机端程序等待所有的核函数执行结束：cudaError_t cudaMemcpy(void* dst,const void* src,size_t count,cudaMemcpykind kind);

编写核函数
核函数是在设备端执行的代码。在核函数中，需要为一个线程规定要进行的计算以及要进行的数据访问。当核函数被调用时，许多不同的CUDA线程并行执行同一个计算任务。核函数必须有一个void返回类型。
函数类型限定符指定了一个函数在主机上执行还是在设备上执行，以及可被主机调用还是被设备调用，下表总结了CUDA C程序中的函数类型限定符。

限定符	执行	调用	备注
global	在设备端执行	可从主机端调用，也可以从计算能力为3的设备中调用	必须有一个void返回类型
device	在设备端执行	仅能从设备端调用
host	在主机端执行	仅能从主机端调用	可以省略

__device__和__host__限定符可以一齐使用，这样函数可以同时在主机和设备端进行编译。
CUDA核函数的限制：只能访问设备内存、必须具有void返回类型、不支持可变数量的参数、不支持静态变量、显示异步行为。

验证核函数
除了许多可用的调试工具外，还有两个非常简单实用的方法可以验证核函数：首先，你可以在Fermi及更高版本的设备端的核函数中使用printf函数；其次，可以将执行参数设置为<<<1,1>>>,因此强制用一个块和一个线程执行核函数，这模拟了串行执行程序。

处理错误
由于许多CUDA调用是异步的，所有有时很难确定某个错误是由哪一步程序引起的。定义一个错误处理宏封装所有的CUDA API调用，这简化了错误检查过程：

#define CHECK(call)
{const cudaError_t error = call;if(error != cudaSuccess){printf("Error: %s:%d",__FILE__,__LINE__);printf("code:%d, reson:%s\n",error,cudaGetErrorString(error));exit(1);}
}

完成错误检查宏的定义之后，就可以在代码中使用宏。比如：CHECK(cudaMemcpy(d_C,gpuRef,nBytes,cudaMemcpyHostToDevice));如果内存拷贝或之前的异步操作产生了错误，这个宏会报告错误代码，并且输出一个可读信息，然后停止程序。再比如：CHECK(cudaDeviceSynchronize())会阻塞主机端线程的运行直到设备端所有的请求都结束，并确保最后的核函数启动部分不会出错。

给核函数计时

用 CPU 计时器计时
可以使用gettimeofday系统调用来创建一个CPU计时器，以获取系统的时钟时间，它需要包含sys/time.h头文件。

double cpuSecord(){struct timeval tp;gettimeofday(&tp,NULL);return ((double)tp.tv_sec + (double)tp.tv_usec*1.e-6);
}

这样可以用cpuSecond函数来测试核函数：

double iStart = cpuSecord();
kernel_name<<<grid,block>>>(argument list);
cudaDeviceSynchronize();
double iElaps = cpuSecond() - iStart;

用 nvprof 工具计时
自CUDA5.0以来，NVIDIA提供了一个名为nvprof的命令行分析工具，可以帮助从应用程序的CPU和GPU活动情况中获取时间线信息，其包括内核执行、内存传输以及CUDA API的调用。

组织并行线程

如果使用了合适的网络和块大小来正确地组织线程，那么可以对内核性能产生很大的影响。对于矩阵运算，传统的方法是在内核中使用一个包含二维网格与二维块的布局来组织线程，但是，这种传统的方法无法获得最佳性能。在矩形加法中使用由二维线程块构成的二维网络、由一维线程块构成的一维网格、由一维线程块构成的二维网格等布局可以取得不同性能的结果。
使用块和线程建立矩阵索引
在一个矩阵加法核函数中，一个线程通常被分配一个数据元素来处理。首先要完成的任务是使用块和线程索引从全局内存中访问指定的数据。通常情况下，对一个二维示例来说，需要管理线程和块索引、矩阵中给定点的坐标、全局线性内存中的偏移量三种索引。对于一个给定的线程，首先可以通过把线程和块索引映射到矩阵坐标上来获取线程块和县城索引的全局内存偏移量，然后将这些矩阵坐标映射到全局内存的存储单元中。
第一步，可以用以下公式把线程和块索引映射到矩阵坐标上：

    ix = threadIdx.x + blockIdx.x * blockDim.x;iy = threadIdx.y + blockIdx.y * blockDim.y;

第二步，可以用以下公式把矩阵坐标映射到全局内存中的索引/存储单元上：

    idx = iy * nx + ix;

其中，printThreadIndo函数被用于输出关于每个线程的以下信息：线程索引、块索引、矩阵坐标、线性全局内存偏移、相应元素的值。

设备管理

查询和管理GPU设备的两种方法：
>CUDA运行时API函数；
>NVIDIA系统管理界面（nvidia-smi）命令行实用程序。
使用运行时 API 查询 GPU 信息
在CUDA运行时API中由很多函数可以帮助管理这些设备。可以使用以下函数来查询关于GPU设备的所有信息：

cudaError_t cudaGetDeviceProperties(cudaDeviceProp* prop,int device);

cudaDeviceProp结构体返回GPU设备的属性。
确定最优 GPU
一些系统支持多GPU。在每个GPU都不同的情况下，选择性能最好的GPU运行核函数是非常重要的。通过比较GPU包含的多处理器的数量选出计算能力最佳的GPU。可以使用以下代码来选择计算能力最优的设备：

int numDevices = 0;
cudaGetDeviceCount(&numDevices);
if(numDevices > 1)
{int maxMultiprocessors = 0, maxDevice = 0;for(int device=0;device <numDevices; device++){cudaDeviceProp props;cudaGetDeviceProperties(&props,device);if(maxMultiprocessors <props.multiProcessorCount){maxMultiprocessors = props.multiProcessorCount;maxDevice = device;}}
}

使用 nvidia-smi 查询 GPU 信息
nvidia-smi是一个命令行工具，用于管理和监控GPU设备，并允许查询和修改设备状态。可以从命令行调用nvidia-smi。

CUDA C编程（二）CUDA编程模型相关推荐

Spring Webflux 响应式编程 (二) - WebFlux编程实战
第一章 Reactive Stream 第1节 jdk9的响应式流就是reactive stream,也就是flow.其实和jdk8的stream没有一点关系.说白了就一个发布-订阅模式,一共只有4 ...
php控制器面向对象编程,PHP 面向对象编程（2）
一些内建方法: class Person { public $isAlive = true; function __construct($name) { //这里我们创建了一个name的属性 $thi ...
2.CUDA 编程手册中文版---编程模型
2.编程模型更多精彩内容,请扫描下方二维码或者访问https://developer.nvidia.com/zh-cn/developer-program 来加入NVIDIA开发者计划本章通过概述 ...
一种基于CUDA标准的异构并行编程模型开发简介
一种基于CUDA标准的异构并行编程模型开发简介目录一.绪论 1.1研究背景及意义 1.2目标平台体系结构简介二.HPPA基本组成结构三.编译工具链开发 3.1 拆分工具HPCufe开发 3.2 ...
CUDA刷新器：CUDA编程模型
CUDA刷新器:CUDA编程模型 CUDA Refresher: The CUDA Programming Model CUDA,CUDA刷新器,并行编程这是CUDA更新系列的第四篇文章,它的目标是 ...
(CUDA 编程1).CUDA 线程执行模型分析（一）招兵 ------ GPU的革命
(CUDA 编程1).CUDA 线程执行模型分析(一)招兵 ------ GPU的革命作者:赵开勇来源:http://www.hpctech.com/2009/0818/198.html 序:或许 ...
linux运行并行计算cuda,并行化计算与CUDA编程
原标题:并行化计算与CUDA编程近年来,显卡的更新换代也比较快,NVIDIA今年的发布会也宣布了RTX3080TI即将到来.显卡的运算能力也越来越强. 很多人对显卡的了解可能源于游戏,但是你的显卡不 ...
【OpenCV CUDA】OpenCV和Cuda结合编程
一.利用OpenCV中提供的GPU模块目前,OpenCV中已提供了许多GPU函数,直接使用OpenCV提供的GPU模块,可以完成大部分图像处理的加速操作. 基本使用方法,请参考:http://www ...
推荐书籍：CUDA并行程序设计：GPU编程指南
过去的五年中,计算领域目睹了英伟达(NVIDIA)公司带来的变革.随后的几年,英伟达公司异军突起,逐渐成长为最知名的游戏硬件制造商之一.计算统一设备架构(Compute Unified Device ...
(CUDA 编程5).CUDA编程接口（一）------一十八般武器
(CUDA 编程5).CUDA编程接口(一)------一十八般武器作者:赵开勇来源:http://www.hpctech.com/2009/0818/203.html 子曰:工欲善其事,必先利其 ...

CUDA C编程（二）CUDA编程模型

CUDA编程模型概述

给核函数计时

组织并行线程

设备管理

CUDA C编程（二）CUDA编程模型相关推荐

最新文章

热门文章