陕西网站建设方案,软文推广网,国内个人网站,wordpress主页打不开本文主要内容为介绍CUDA编程前的一些基础知识
参考资料#xff1a;
高升博客 《CUDA C编程权威指南》 以及 CUDA官方文档
文章、讲解视频同步更新公众《AI知识物语》#xff0c;B站#xff1a;出门吃三碗饭
1#xff1a;并行计算
并行程序可以分为 指令并行#xff1…本文主要内容为介绍CUDA编程前的一些基础知识
参考资料
高升博客 《CUDA C编程权威指南》 以及 CUDA官方文档
文章、讲解视频同步更新公众《AI知识物语》B站出门吃三碗饭
1并行计算
并行程序可以分为 指令并行一般应用在管理系统比如淘宝交易系统每时每刻都有很多人在同时使用后台需要能够并行处理这些请求。 数据并行一般应用大规模数据计算大量数据使用相同的计算程序计算
CUDA非常适合数据并行计算。
数学并行第一步 数据按线程划分 1 块划分把一整块数据切成小块每个小块随机的划分给一个线程每个块的执行顺序随机。
2 周期划分线程按照顺序处理相邻的数据块每个线程处理多个数据块比如我们有五个线程线程1执行块1线程2执行块2……线程5执行块5线程1执行块6 2计算机架构
佛林分类法Flynn’s Taxonomy根据指令和数据进入CPU的方式分类分为以下四类 1分别以数据和指令进行分析
单指令单数据SISD传统串行计算机386 单指令多数据SIMD并行架构比如向量机所有核心指令唯一但是数据不同现在CPU基本都有这类的向量指令 多指令单数据MISD少见多个指令围殴一个数据 多指令多数据MIMD并行架构多核心多指令异步处理多个数据流从而实现空间上的并行MIMD多数情况下包含SIMD就是MIMD有很多计算核计算核支持SIMD
2为了提高并行的计算能力我们要从架构上实现下面这些性能提升
降低延迟 提高带宽 提高吞吐量
延迟是指操作从开始到结束所需要的时间一般用微秒计算延迟越低越好。 带宽是单位时间内处理的数据量一般用MB/s或者GB/s表示。 吞吐量是单位时间内成功处理的运算数量一般用gflops来表示十亿次浮点计算吞吐量和延迟有一定关系都是反应计算速度的一个是时间除以运算次数得到的是单位次数用的时间–延迟一个是运算次数除以时间得到的是单位时间执行次数–吞吐量
3异构架构
CPU我们可以把它看做一个指挥者主机端host而完成大量计算的GPU是我们的计算设备device
1上面这张图能大致反应CPU和GPU的架构不同。
左图CPU4个ALU主要负责逻辑计算1个控制单元Control1个DRAM内存1个Cache缓存
右图GPU绿色小方块是ALU我们注意红色框内的部分SM这一组ALU公用一个Control单元和Cache这个部分相当于一个完整的多核CPU但是不同的是ALU多了control部分变小可见计算能力提升了控制能力减弱了所以对于控制逻辑复杂的程序一个GPU的SM是没办法和CPU比较的但是对了逻辑简单数据量大的任务GPU更搞笑并且注意一个GPU有好多个SM而且越来越多
2主机代码在主机端运行**被编译成主机架构的机器码设备端的在设备上执行被编译成设备架构的机器码所以主机端的机器码和设备端的机器码是隔离的自己执行自己的没办法交换执行。
主机端代码主要是控制设备完成数据传输等控制类工作设备端主要的任务就是计算。 因为当没有GPU的时候CPU也能完成这些计算只是速度会慢很多所以可以把GPU看成CPU的一个加速设备。
3CPU和GPU线程的区别
CPU线程是重量级实体操作系统交替执行线程线程上下文切换花销很大 GPU线程是轻量级的GPU应用一般包含成千上万的线程多数在排队状态线程之间切换基本没有开销。 CPU的核被设计用来尽可能减少一个或两个线程运行时间的延迟而GPU核则是大量线程最大幅度提高吞吐量
4CUDA编程结构
一个完整的CUDA应用可能的执行顺序如下图 从host的串行到调用核函数核函数被调用后控制马上归还主机线程也就是在第一个并行代码执行时很有可能第二段host代码已经开始同步执行了。
5内存管理 1HostCPU 通过 cudaMalloc cudaMemcpycudaMemsetcudaFree等方式与DeviceGPU进行内存管理 2Device空间里面有Grid其由许多Block以及GlobalMemory组成Grid的大小就是其包含的Block数量 3Block里面包含许多Thread和Shared Memory共享空间Block大小等于其包含的线程数量
6线程管理
1当内核函数开始执行如何组织GPU的线程就变成了最主要的问题了我们必须明确一个核函数只能有一个grid一个grid可以有很多个块每个块可以有很多的线程这种分层的组织结构使得我们的并行过程更加自如灵活 一个线程块block中的线程可以完成下述协作
1同步 2共享内存 不同块内线程不能相互影响他们是物理隔离的
2每个线程都执行同样的一段串行代码为了让这段相同的代码对应不同的数据首先第一步就是让这些线程彼此区分开才能对应到相应从线程使得这些线程也能区分自己的数据。如果线程本身没有任何标记那么没办法确认其行为。 依靠下面两个内置结构体确定线程标号 blockIdx线程块在线程网格内的位置索引 threadIdx线程在线程块内的位置索引
上面这两个是坐标当然我们要有同样对应的两个结构体来保存其范围也就是blockIdx中三个字段的范围threadIdx中三个字段的范围
blockDim gridDim
7核函数
核函数就是在CUDA模型上诸多线程中运行的那段串行代码这段代码在设备上运行用NVCC编译产生的机器码是GPU的机器码所以我们写CUDA程序就是写核函数第一步我们要确保核函数能正确的运行产生正切的结果第二优化CUDA程序的部分无论是优化算法还是调整内存结构线程结构都是要调整核函数内的代码来完成这些优化的。
我们一直把我们的CPU当做一个控制者运行核函数要从CPU发起那么我们开始学习如何启动一个核函数
1核函数调用
kernel_name4,8(argument list);这个三个尖括号’grid,block’内是对设备代码执行的线程结构的配置或者简称为对内核进行配置 上面这条指令的线程布局是4个块每个块内分配调用8个线程 我们的核函数是同时复制到多个线程执行的上文我们说过一个对应问题多个计算执行在一个数据肯定是浪费时间所以为了让多线程按照我们的意愿对应到不同的数据就要给线程一个唯一的标识由于设备内存是线性的基本市面上的内存硬件都是线性形式存储数据的我们观察上图可以用threadIdx.x 和blockIdx.x 来组合获得对应的线程的唯一标识后面我们会看到threadIdx和blockIdx能组合出很多不一样的效果
kernel_name1,32(argument list); // 调用1个块kernel_name32,1(argument list);// 调用32个块
