CUDA编程基础【内存模型和线程束】

最新推荐文章于 2024-02-23 22:49:01 发布

Enjoyment0

最新推荐文章于 2024-02-23 22:49:01 发布

阅读量209

点赞数

文章标签： c语言

本文链接：https://blog.csdn.net/Enjoyment0/article/details/132766536

版权

该博客内容从B站https://www.bilibili.com/video/BV1jX4y1w7Um/?spm_id_from=333.999.0.0&vd_source=4a380242eeb6f602ee5572294c438a28记录得来。
代码上传至github供以后学习使用。

内存模型（Memory Model）

内存类型和速度

Global Memory，全局内存
速度：普通，读写。大小：显存大小（11GB etc.）
Constant Memory，常量内存
速度：很快，只读。大小：一般64KB，16bits寻址
Shared Memory，共享内存
速度：快，读写。大小：2080Ti有48KB
Local Memory，本地内存（其实是全局内存）
速度：普通，读写。大小：可用内存/SM数量/SM最大常驻线程数
Texture Memory，纹理内存
速度：快，只读。
Register，寄存器
速度：最快，读写

CUDA内存模型
在这里插入图片描述

Global Memory 2 ----- Memory Transfer
在这里插入图片描述
CPU到GPU之间的传输经过PCIe，速度很慢8GB/s
GPU到GPU Memory之间的传输，使用GDDR5，速度很快144GB/s
所以应该尽量避免GPU与CPU之间的数据传输

Global Memory 3 —Pinned Memory(页锁定内存)
在这里插入图片描述
Pinned Memory依旧是Host Memory。它的使用场景是代替malloc、new，然后再进行cudaMemcpy
虽然Pinned Memory可以允许GPU直接访问，但是这种操作效率是低效的
正确的使用场景是适合大数据传输的中间储存，不适合频繁的读写操作

Global Memory 4 —Zero Copy Memory(零拷贝内存)
Zero Copy Memory：即内存的复制过程，没有CPU的参与。直接由GPU和内存条操作也即DMA。
Zero Copy Memory实质上就是Pinned Memory映射到device的空间地质，本质上是等家的。
由cudaMallocHost、cudaHostAlloc实现分配，cudaFreeHost释放。
cudaHostGetDevicePointer可以获取分配的Pinned Memory映射到device的地址。
在频繁读写上，Zero Copy Memory的速度比device memory速度差很多。
在这里插入图片描述 Global Memory 5 —Unified Memory(统一内存)
由统一内存系统管理的内存。Unified Virtual Addressing(统一虚拟内存)将CPU和GPU的内存看作为一个整体进行管理和使用。分配的内存可以CPU/GPU直接访问。
区别：在Linux下内存的位置将会选择在device memory上。因此设备访问速度会快，而host访问会慢。在Windows下，将会以Zero Copy Memory的形式存在host memory上。
在这里插入图片描述由cudaMallocManaged分配，cudaFree释放

Constant Memory
使用__constant__修饰符修饰一下即可。cudaDeviceSynchronize()

线程束（warp）

SIMT(Single Instruction Multiple Thread).CUDA执行模型。
SIMD(Single Instruction Multiple Data).多在CPU的流指令集使用。

一个Stream Multiprocessors有两个Block(warp)
在这里插入图片描述线程束，是指thread warp，也就是新这里32个core组成的SM block
注意：thread warp是逻辑上的线程束，硬件执行时的抽象，并不能直接指硬件。

Warps和thread blocks
任意时刻，任务的执行都是以warp为单位的，warp_size的大小通常都是32。也就是启动n个线程，需要ceil(n/warp_size)个线程束。

线程束中的线程有两种状态：
active：活跃，线程启动后的状态
inactive：不活跃，线程未启动，或者等待时的状态（分支问题）

warp类型根据调度器来决定其状态有：
Selected warp：选中
Eligible warp：准备好待执行
Stalled warp：没准备好

一个block需要的warp数量 = ceil (T / warp_size) T= block中的线程数
对于一个warp中启动的线程束，如果小于32，则生于部分线程为inactive状态。
inactive的线程依旧消耗了资源。

对于gradDim = m，blockDim = n的人物，总共需要的warp数为: m * ceil (n / warp_size)。而实际上物理warp数有限(64SM，64 Core per SM， total warp = 68 * 2 = 136)。则选中的条件为：32个cpre空闲，并且指令参数准备就绪。
block内的warps公用shared memory，通常block内的所有warps需要一定程序上的同步和并行
grid内的blocks，通常是独立、并行的。每个block之间使用喜爱唔呼独立的shared memory
因此设计gridDim、blockDim时，考虑block内通常可以互通数据、并行、同步数据等。而block间相互独立
block的大小也尽量是warp_size的倍数，尽量避免warp_size无法填满造成的inactive线程消耗。