CUDA基础知识总结

CUDA学习

CUDA 入门基础知识

CPU

CPU(Central Processing Unit)是一块超大规模的集成电路,是一台计算机的运算核心(Core)和控制核心( Control Unit)。它的功能主要是解释计算机指令以及处理计算机软件中的数据。

CPU与内部存储器(Memory)和输入/输出(I/O)设备合称为电子计算机三大核心部件。CPU主要包括运算器(算术逻辑运算单元,ALU,Arithmetic Logic Unit)、控制单元(CU, Control Unit)、寄存器(Register)、和高速缓冲存储器(Cache)及实现它们之间联系的数据(Data)、控制及状态的总线(Bus)。简单来说就是:计算单元、控制单元和存储单元。CPU遵循的是冯诺依曼架构,其核心就是:存储程序,顺序执行。

因为CPU的架构中需要大量的空间去放置存储单元和控制单元,相比之下计算单元只占据了很小的一部分,所以它在大规模并行计算能力上极受限制,而更擅长于逻辑控制。

GPU

显卡(Video card,Graphics card)全称显示接口卡,又称显示适配器,是计算机最基本配置、最重要的配件之一。显卡是电脑进行数模信号转换的设备,承担输出显示图形的任务。具体来说,显卡接在电脑主板上,它将电脑的数字信号转换成模拟信号让显示器显示出来,同时显卡还是有图像处理能力,可协助CPU工作,提高整体的运行速度。在科学计算中,显卡被称为显示加速卡。

原始的显卡一般都是集成在主板上,只完成最基本的信号输出工作,并不用来处理数据。显卡也分为独立显卡和集成显卡。一般而言,同期推出的独立显卡的性能和速度要比集成显卡好、快。
在这里插入图片描述

CPU和GPU

在没有GPU之前,基本上所有的任务都是交给CPU来做的。有GPU之后,二者就进行了分工,CPU负责逻辑性强的事物处理和串行计算,GPU则专注于执行高度线程化的并行处理任务(大规模计算任务)。GPU并不是一个独立运行的计算平台,而需要与CPU协同工作,可以看成是CPU的协处理器,因此当我们在说GPU并行计算时,其实是指的基于CPU+GPU的异构计算架构。在异构计算架构中,GPU与CPU通过PCIe总线连接在一起来协同工作,CPU所在位置称为为主机端(host),而GPU所在位置称为设备端(device)。

在这里插入图片描述

GPU包括更多的运算核心,其特别适合数据并行的计算密集型任务,如大型矩阵运算,而CPU的运算核心较少,但是其可以实现复杂的逻辑运算,因此其适合控制密集型任务。另外,CPU上的线程是重量级的,上下文切换开销大,但是GPU由于存在很多核心,其线程是轻量级的。因此,基于CPU+GPU的异构计算平台可以优势互补,CPU负责处理逻辑复杂的串行程序,而GPU重点处理数据密集型的并行计算程序,从而发挥最大功效。GPU无论发展得多快,都只能是替CPU分担工作,而不是取代CPU。

GPU中有很多的运算器ALU和很少的缓存cache,缓存的目的不是保存后面需要访问的数据的,这点和CPU不同,而是为线程thread提高服务的。如果有很多线程需要访问同一个相同的数据,缓存会合并这些访问,然后再去访问DRAM。

CUDA介绍

2006年,NVIDIA公司发布了CUDA(Compute Unified Device Architecture),是一种新的操作GPU计算的硬件和软件架构,是建立在NVIDIA的GPUs上的一个通用并行计算平台和编程模型,它提供了GPU编程的简易接口,基于CUDA编程可以构建基于GPU计算的应用程序,利用GPUs的并行计算引擎来更加高效地解决比较复杂的计算难题。它将GPU视作一个数据并行计算设备,而且无需把这些计算映射到图形API。操作系统的多任务机制可以同时管理CUDA访问GPU和图形程序的运行库,其计算特性支持利用CUDA直观地编写GPU核心程序。

CUDA提供了对其它编程语言的支持,如C/C++,Python,Fortran等语言。只有安装CUDA才能够进行复杂的并行计算。主流的深度学习框架也都是基于CUDA进行GPU并行加速的,几乎无一例外。还有一个叫做cudnn,是针对深度卷积神经网络的加速库。

CUDA在软件方面组成有:一个CUDA库、一个应用程序编程接口(API)及其运行库(Runtime)、两个较高级别的通用数学库,即CUFFT和CUBLAS。CUDA改进了DRAM的读写灵活性,使得GPU与CPU的机制相吻合。另一方面,CUDA提供了片上(on-chip)共享内存,使得线程之间可以共享数据。应用程序可以利用共享内存来减少DRAM的数据传送,更少的依赖DRAM的内存带宽。

CUDA架构

CUDA的架构中引入了主机端(host)和设备(device)的概念。CUDA程序中既包含host程序,又包含device程序。同时,host与device之间可以进行通信,这样它们之间可以进行数据拷贝。

主机(Host):将CPU及系统的内存(内存条)称为主机。

设备(Device):将GPU及GPU本身的显示内存称为设备。

动态随机存取存储器(DRAM):Dynamic Random Access Memory,最为常见的系统内存。DRAM只能将数据保持很短的时间。为了保持数据,DRAM使用电容存储,所以必须隔一段时间刷新(refresh)一次,如果存储单元没有被刷新,存储的信息就会丢失。(关机就会丢失数据)

典型的CUDA程序的执行流程如下:

  1. 分配host内存,并进行数据初始化;
  2. 分配device内存,并从host将数据拷贝到device上;
  3. 调用CUDA的核函数在device上完成指定的运算;
  4. 将device上的运算结果拷贝到host上;
  5. 释放device和host上分配的内存。

线程层次结构

CUDA执行流程中最重要的一个过程是调用CUDA的核函数来执行并行计算,kernel是C

在CUDA程序构架中,主机端代码部分在CPU上执行,是普通的C代码;当遇到数据并行处理的部分,CUDA 就会将程序编译成GPU能执行的程序,并传送到GPU,这个程序在CUDA里称做核(kernel)。

设备端代码部分在GPU上执行,此代码部分在kernel上编写(.cu文件)。kernel用__global__符号声明,在调用时需要用<<<grid, block>>>来指定kernel要执行及结构。

CUDA是通过函数类型限定词区别在host和device上的函数,主要的三个函数类型限定词如下:

global:在device上执行,从host中调用(一些特定的GPU也可以从device上调用),返回类型必须是void,不支持可变参数参数,不能成为类成员函数。注意用__global__定义的kernel是异步的,这意味着host不会等待kernel执行完就执行下一步。
device:在device上执行,单仅可以从device中调用,不可以和__global__同时用。
host:在host上执行,仅可以从host上调用,一般省略不写,不可以和__global__同时用,但可和__device__同时使用,此时函数会在device和host都编译。

网格 grid
kernel在device上执行时,实际上是启动很多线程,一个kernel所启动的所有线程称为一个网格(grid),同一个网格上的线程共享相同的全局内存空间。grid是线程结构的第一层次。

线程块 block
网格又可以分为很多线程块(block),一个block里面包含很多线程。各block是并行执行的,block间无法通信,也没有执行顺序。block的数量限制为不超过65535(硬件限制)。第二层次。

grid和block都是定义为dim3类型的变量,dim3可以看成是包含三个无符号整数(x,y,z)成员的结构体变量,在定义时,缺省值初始化为1。grid和block可以灵活地定义为1-dim,2-dim以及3-dim结构。

CUDA中,每一个线程都要执行核函数,每一个线程需要kernel的两个内置坐标变量(blockIdx,threadIdx)来唯一标识,其中blockIdx指明线程所在grid中的位置,threaIdx指明线程所在block中的位置。它们都是dim3类型变量。

一个线程在block中的全局ID,必须还要知道block的组织结构,这是通过线程的内置变量blockDim来获得。它获取block各个维度的大小。对于一个2-dim的block(D_x,D_y),线程 (x,y) 的ID值为(x+y∗D_x),如果是3-dim的block (D_x,D_y,D_z),线程(x,y,z)的ID值为(x+y∗D_x+z∗D_x∗D_y) 。另外线程还有内置变量gridDim,用于获得grid各个维度的大小。

每个block有包含共享内存(Shared Memory),可以被线程块中所有线程共享,其生命周期与线程块一致。
每个thread有自己的私有本地内存(Local Memory)。此外,所有的线程都可以访问全局内存(Global Memory),还可以访问一些只读内存块:常量内存(Constant Memory)和纹理内存(Texture Memory)。

线程 thread
一个CUDA的并行程序会被以许多个threads来执行。数个threads会被群组成一个block,同一个block中的threads可以同步,也可以通过shared memory通信。

线程束 warp
GPU执行程序时的调度单位,SM的基本执行单元。目前在CUDA架构中,warp是一个包含32个线程的集合,这个线程集合被“编织在一起”并且“步调一致”的形式执行。同一个warp中的每个线程都将以不同数据资源执行相同的指令,这就是所谓 SIMT架构(Single-Instruction, Multiple-Thread,单指令多线程)。
在这里插入图片描述

CUDA内存模型

SP:最基本的处理单元,streaming processor,也称为CUDA core。最后具体的指令和任务都是在SP上处理的。GPU进行并行计算,也就是很多个SP同时做处理。

SM:GPU硬件的一个核心组件是流式多处理器(Streaming Multiprocessor)。SM的核心组件包括CUDA核心、共享内存、寄存器等。SM可以并发地执行数百个线程。一个block上的线程是放在同一个流式多处理器(SM)上的,因而,一个SM的有限存储器资源制约了每个block的线程数量。在早期的NVIDIA 架构中,一个block最多可以包含 512个线程,而在后期出现的一些设备中则最多可支持1024个线程。一个kernel可由多个大小相同的block同时执行,因而线程总数应等于每个块的线程数乘以块的数量。

综上,SM要为每个block分配shared memory,而也要为每个warp中的线程分配独立的寄存器。所以SM的配置会影响其所支持的线程块和线程束并发数量。所以kernel的grid和block的配置不同,性能会出现差异。还有,由于SM的基本执行单元是包含32个线程的warp,所以block大小一般要设置为32的倍数。

CUDA编程模型

CUDA编程模型将CPU作为主机(Host),GPU作为协处理器或者设备(Device)。在一个系统中可以存在一个主机和若干设备。CPU负责进行逻辑性强的事务处理和串行计算,GPU则专注于执行高度线程化的并行处理任务。CPU和GPU拥有各自独立的存储地址空间:主机端的内存和设备端的显存。

CUDA对内存的操作同一般的C程序相同,而对显存则需要调用CUDA API中的存储管理函数。

一旦确定了程序中的并行部分,就可以考虑把这部分计算工作交给GPU。运行在GPU上的CUDA并行计算函数成为kernel(内核函数)。

一个完整的CUDA程序是由一系列的设备端kernel函数并行步骤和主机端的串行处理步骤共同组成。这些处理步骤会按照程序中相应语句的顺序依次执行,满足顺序一致性。

内核函数

运行在GPU上的内核函数必须通过__global__函数类型限定,并且只能在主机端代码调用。在调用时,必须声明内核函数的执行参数。

__global__ void VecAdd(float* A,float* B,float* C)
{
 
}
 
int main()
{//
   VecAdd<<<1,N>>>(A,B,C);//1代表block数,N代表每个block中的thread数
}

在设备端运行的kernel函数时并行执行的,每个线程均执行kernel函数的指令,每一个线程有自己的block ID和thread ID,block ID和thred ID只能在kernel中通过内建变量访问。

需要注意的时内建变量不需要自己定义,由设备中的专用寄存器提供。因此,内建变量时只读的,并且只能在GPU端的kernel函数中使用。

函数在最前有一个修饰符“global”,这个修饰符告诉编译器,被修饰的函数应该编译为在GPU而不是在CPU上运行,所以这个函数将被交给编译设备代码的编译器——NVCC编译器来处理,其他普通的函数或语句将交给主机编译器处理。

注:关于threadIdx.x的说明。

CUDA中的线程(thread)是设备中并行运算结构中的最小单位,类似于主机中的线程的概念,thread可以以一维、二维、三维的形式组织在一起,threadIdx.x表示的是thread在x方向的索引号,还可能存在thread在y和z方向的索引号threadIdx.y和threadIdx.z。
一维、二维或三维的thread组成一个线程块(Block),一维、二维或三维的线程块(Block)组合成一个线程块网格(Grid),线程块网格(Grid)可以是一维或二维的。通过网格块(Grid)->线程块(Block)->线程(thread)的 顺序可以定位到每一个并且唯一的线程。

函数会被GPU上的多个线程同时执行一次,线程间彼此没有通信,相互独立。到底会有多少个线程来分别执行核函数,是在“<<< >>>”符号里定义的。“<<< >>>”表示运行时配置符号,在本程序中的定义是<<<1,N>>>,表示分配了一个线程块(Block),每个线程块有分配了N个线程,“<<<>>>”中的 参数并不是传递给设备代码的参数,而是定义主机代码运行时如何启动设备代码。以上定义的这些线程都是一个维度上的,可以通过thredaIdx.x来获取执行当前计算任务的线程的ID号。

线程结构

kernel以线程网格(Grid)的形式组织,每个线程网格由若干个线程块(block)组成,而每个线程块又由若干个线程(thread)组成,但是block之间不能相互通信,也没有执行顺序,但是同一个block里面的线程可以相互通信,在一个block里面,有shared memory,并通过栅栏同步保证线程间能够正确地共享数据。具体来说,可以在kernel函数中通过设置_syncthreads()函数实现。

内建变量使用了dim3类型,所谓dim3是基于uint3定义的矢量类型

对于一维的block,线程的threadID就是threadIdx.x

对于二维的block(Dx,Dy),线程的threadID就是threadIdx.x+threadIdx.y*Dx

对于三维的block(Dx,Dy,Dz)的三维block,线程的threadID是threadIdx.x+threadIdx.yDx+threadIdx.zDx*Dy

硬件映射

前面讲的线程结构,编程模型,都只是逻辑模型,这种逻辑模型还需要映射到硬件上。

kernel实际上是在block上运行的,映射到硬件上就是在SM(流多处理器)中执行,而block中的每一个thread则放到SP上执行,但是一个SM上面可以同时有多个block。

在实际运行中,block会被分割为更小的线程束(warp)。线程束的大小由硬件的计算能力决定,在telsa架构的CPU中,一个线程束是由连续的32个线程组成。例如每个block中,ID为0-31的线程为一束,ID为32-63的为第二束。

为什么是32,这是因为每发射一个warp指令,SM中的8个SP会将这个指令执行4遍。值得注意的是,warp是一个由硬件决定的概念,在抽象的CUDA模型中见不到,但是仍有相当的影响

执行模型

CUDA采用了SIMT(Single Instruction,Multiple Thread)执行模型,即单指令多线程模型。需要注意的是,如果在SIMT中要控制单个线程的某个行为,即需要用到分支,这会降低效率,因为一个warp指令,在个线程中是相同的,执行时间就是单个线程的执行时间,但是如果线程执行有差异,进入SP需要决定某些指令需不需要执行,执行时间就变成了各个执行时间之和,开发CUDA程序时,应尽量避免分支。

  • 7
    点赞
  • 20
    收藏
    觉得还不错? 一键收藏
  • 0
    评论

“相关推荐”对你有帮助么?

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值