CUDA学习笔记

1、cuda存储器模型 各类存储器的的特点

CUDA存储器模型:
在这里插入图片描述
在这里插入图片描述

GPU片内:register,shared memory;
板载显存:local memory,constant memory, texture memory ,global memory;
host 内存: host memory, pinned memory.

1)register: 访问延迟极低;
每个线程占有的register有限,编程时不要为其分配过多私有变量;
对每个线程来说,寄存器都是线程私有的–这与CPU中一样。如果寄存器被消耗完,数据将被存储在本地存储器(localmemory)。
Localmemory对每个线程也是私有的,但是localmemory中的数据是被保存在显存中,而不是片内的寄存器或者缓存中,因此速度很慢。
线程的输入和中间输出变量将被保存在寄存器或者本地存储器中。
2)local memory:寄存器被使用完毕,数据将被存储在局部存储器中;
3)shared memory:访问速度与寄存器相似;
实现线程间通信的延迟最小;
保存公用的计数器或者block的公用结果;
用于线程间通信的共享存储器。共享存储器是可以被同一block中的所有thread访问的可读写存储器。
访问共享存储器几乎和访问寄存器一样快速,是实现线程间通信的延迟最小的方法。
共享存储器可以实现许多不同的功能,如用于保存共用的计数器(例如计算循环次数)或者block的公用结果。
4)global memory:存在于显存中,也称为线性内存(显存可以被定义为线性存储器或者CUDA数组);
使用的是普通的显存,无缓存,可读写,速度慢
整个网格中的任意线程都能读写全局存储器的任意位置,并且既可以从CPU访问,也可以从GPU访问。
cudaMalloc()函数分配,cudaFree()函数释放,cudaMemcpy()进行主机端与设备端的数据传输;

2、cuda执行模型 线程组织结构、与任务/数据的映射关系

Cuda执行模型
SIMT和SIMD架构:
相似性:两者都将相同的指令广播给多个执行单元来实现并行。
区别:SIMD要求同一向量的所有元素在一个同步组中一起执行,而SIMT则允许同一线程束的多个线程独立执行。
尽管一个线程束中所有线程在相同的程序地址同时开始执行,但是单独的线程仍有可能有不同的行为。SIMT确保可以编写独立的线程级并行代码、标量线程、以及用于协调线程的数据并行代码。
SIMT模型包含3个SIMD所不具备的关键特征:
每个线程都有自己的指令地址计数器;
每个线程都有自己的寄存器状态;
每个线程都可以有一个独立的执行路径
Fermi架构,Kepler架构
GPU架构:GPU架构是围绕一个流式多处理器(SM)的扩展阵列搭建的

线程组织结构
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述

网格 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,单指令多线程)。一个warp中的线程必然在同一个block中,如果block所含线程数目不是warp大小的整数倍,那么多出的那些thread所在的warp中,会剩余一些inactive(待用)的thread
Warp divergence:因为所有同一个warp中的thread必须执行相同的指令,那么如果这些线程在遇到控制流语句时,如果进入不同的分支,那么同一时刻除了正在执行的分支之外,其余分支都被阻塞了,十分影响性能。这类问题就是warp divergence。
在这里插入图片描述
在这里插入图片描述

SM(Streaming Multiprocessor): 流式多处理器,每个SM最多占用1536个thread,8个block

Grid:由一个kernel启动所产生的所有线程统称为一个线程网格(Grid)。
同一线程网格中的所有线程共享同全局内存空间。一个网格有多个线程块(Block)构成,一个线程块包含一组线程,同一线程块内的线程协同可以通过“同步”和“共享内存”的方式来实现。不同线程块内的线程不能协作。
在一个网格中,我们通过以下两个坐标变量来定位一个线程,
(1)blockIdx:线程块在线程网格中ID号
(2)threadIdx:线程在线程块内的ID号
这些坐标变量是kernel函数中需要预初始化的内置变量。

当执行有一个核函数时,CUDA Runtime 为每个线程分配坐标变量blockIdx和threadIdx。基于这些坐标,我们将数据分配到不同的GPU线程上,然后并行处理所有的数据。

坐标变量blocIdx和threadIdx都是基于unit3定义的CUDA内置的向量类型,分别包含3个无符号的整数结构,可以通过x,y,z三个元素来进行索引。
在这里插入图片描述

任务/数据的映射关系

3、根据具体任务会线程配置

这个我没写

4、warp 数量的计算、会计算发生control divergence 的warp数量

在2里面

5、DRAM burst 的访问、是否合并内存访问等

现代DRAM系统设计为始终可访问在burst模式下。突发字节被传输到处理器,但当访问不到连续位置时,将被丢弃。
当warp的所有线程执行load指令时,如果所有访问的位置都属于同一个burst段,则只会发出一个DRAM请求,并且访问完全合并。
当访问的位置分布在突发段时边界:凝聚失败,发出多个DRAM请求,通道未完全合并。
在这里插入图片描述
在这里插入图片描述

6、原子操作、吞吐量的计算

原子操作是指不会被线程调度机制打断的操作;这种操作一旦开始,就一直运行到结束,
原子操作的关键概念
-由单个硬件指令对存储器位置地址执行的读-修改-写操作
-读取旧值,计算新值,并将新值写入该位置
-硬件确保在当前原子操作完成之前,没有其他线程可以在同一位置执行另一个读-修改-写操作
-试图在同一位置执行原子操作的任何其他线程通常会被保留在队列中
-所有线程在同一位置连续执行原子操作
并发度=吞吐量∗延迟
直方图
吞吐量=1/延迟
在这里插入图片描述

7、占用率的概念与计算

占用率是指每个多处理器(Streaming Multiprocessor,SM)的活动线程束(warps)数量与实际的活动warps数量的比率。
占用率=每个SM中活跃线程束的数量/每个SM中最大的线程束的数量

8、CPU与GPU区别

CPU:强大的算术逻辑单元,大型的缓存,复杂的控制逻辑,低延时,适合处理复杂的逻辑分支判断运算
GPU:简单大量的计算单元,小型的缓存,简单的控制逻辑,吞吐量大,适合处理大规模的数据并发运算

9、共享存储的特点及其bank 冲突

在这里插入图片描述

10、矩阵乘法、卷积、直方图、直方图均衡化、并行规约、扫描算法

向量加法host code
在这里插入图片描述
在这里插入图片描述
矩阵乘法
在这里插入图片描述
在这里插入图片描述
直方图
在这里插入图片描述

在这里插入图片描述

在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
卷积

在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
并行归约
在这里插入图片描述
扫描算法
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述

  • 2
    点赞
  • 15
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
在现有省、市港口信息化系统进行有效整合基础上,借鉴新 一代的感知-传输-应用技术体系,实现对码头、船舶、货物、重 大危险源、危险货物装卸过程、航管航运等管理要素的全面感知、 有效传输和按需定制服务,为行政管理人员和相关单位及人员提 供高效的管理辅助,并为公众提供便捷、实时的水运信息服务。 建立信息整合、交换和共享机制,建立健全信息化管理支撑 体系,以及相关标准规范和安全保障体系;按照“绿色循环低碳” 交通的要求,搭建高效、弹性、高可扩展性的基于虚拟技术的信 息基础设施,支撑信息平台低成本运行,实现电子政务建设和服务模式的转变。 实现以感知港口、感知船舶、感知货物为手段,以港航智能 分析、科学决策、高效服务为目的和核心理念,构建“智慧港口”的发展体系。 结合“智慧港口”相关业务工作特点及信息化现状的实际情况,本项目具体建设目标为: 一张图(即GIS 地理信息服务平台) 在建设岸线、港口、港区、码头、泊位等港口主要基础资源图层上,建设GIS 地理信息服务平台,在此基础上依次接入和叠加规划建设、经营、安全、航管等相关业务应用专题数据,并叠 加动态数据,如 AIS/GPS/移动平台数据,逐步建成航运管理处 "一张图"。系统支持扩展框架,方便未来更多应用资源的逐步整合。 现场执法监管系统 基于港口(航管)执法基地建设规划,依托统一的执法区域 管理和数字化监控平台,通过加强对辖区内的监控,结合移动平 台,形成完整的多维路径和信息追踪,真正做到问题能发现、事态能控制、突发问题能解决。 运行监测和辅助决策系统 对区域港口与航运业务日常所需填报及监测的数据经过科 学归纳及分析,采用统一平台,消除重复的填报数据,进行企业 输入和自动录入,并进行系统智能判断,避免填入错误的数据, 输入的数据经过智能组合,自动生成各业务部门所需的数据报 表,包括字段、格式,都可以根据需要进行定制,同时满足扩展 性需要,当有新的业务监测数据表需要产生时,系统将分析新的 需求,将所需字段融合进入日常监测和决策辅助平台的统一平台中,并生成新的所需业务数据监测及决策表。 综合指挥调度系统 建设以港航应急指挥中心为枢纽,以各级管理部门和经营港 口企业为节点,快速调度、信息共享的通信网络,满足应急处置中所需要的信息采集、指挥调度和过程监控等通信保障任务。 设计思路 根据项目的建设目标和“智慧港口”信息化平台的总体框架、 设计思路、建设内容及保障措施,围绕业务协同、信息共享,充 分考虑各航运(港政)管理处内部管理的需求,平台采用“全面 整合、重点补充、突出共享、逐步完善”策略,加强重点区域或 运输通道交通基础设施、运载装备、运行环境的监测监控,完善 运行协调、应急处置通信手段,促进跨区域、跨部门信息共享和业务协同。 以“统筹协调、综合监管”为目标,以提供综合、动态、实 时、准确、实用的安全畅通和应急数据共享为核心,围绕“保畅通、抓安全、促应急"等实际需求来建设智慧港口信息化平台。 系统充分整合和利用航运管理处现有相关信息资源,以地理 信息技术、网络视频技术、互联网技术、移动通信技术、云计算 技术为支撑,结合航运管理处专网与行业数据交换平台,构建航 运管理处与各部门之间智慧、畅通、安全、高效、绿色低碳的智 慧港口信息化平台。 系统充分考虑航运管理处安全法规及安全职责今后的变化 与发展趋势,应用目前主流的、成熟的应用技术,内联外引,优势互补,使系统建设具备良好的开放性、扩展性、可维护性。

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

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

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

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

抵扣说明:

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

余额充值