Skip to main content

CUDA指令和架构


A100显卡上的tensorcore有自己的私有寄存器吗，微架构探索 https://zhuanlan.zhihu.com/p/620257581

乘影GPGPU架构文档手册v2.01.pdf

https://zhuanlan.zhihu.com/p/166180054

https://www.tinyedi.com/cuda_learning/#pipeline

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

https://zhuanlan.zhihu.com/p/486224812

image.pngimage.png

image.pngimage.png

虽然异步屏障仍然是 Hopper 编程模型的一部分，但 Hopper 添加了一种新形式的屏障，称为Asynchronous Transaction Barrier。Asynchronous Transaction Barrier与异步屏障非常相似。它也是一个拆分屏障，但它不仅计算线程到达，还计算事务。 Hopper 包含一个用于写入共享内存的新命令，该命令传递要写入的数据和事务计数。事务计数本质上是一个字节计数。异步事务屏障将阻塞等待命令的线程，直到所有生产者线程都执行了到达，并且所有事务计数的总和达到预期值。异步事务屏障是用于异步内存副本或数据交换的强大新原语。

Predicate  有时候代指指令前加的形如 @P0, @!P3 等的predication，PTX文档里也把这个叫guard
predicate。有时候也会特指那个predicate register，如 P0, P6, PT ，反正都差不多这个意思，
不混淆就行。在之前的文章也提到过，在现有的所有架构中，每个指令都有4bit的编码来指定每
个predicate，3bit用来指定索引（所以每线程有2^3=8个predicate register P0~P7 ，其中 P7=PT
为恒 True ），1bit表示是否取反。如果是 @PT ，那就会在反汇编中省去不显示。那如果是 @!
PT 呢？嗯，大家自己想……
Predicate是控制某个线程是否执行某个指令的两种方式之一，另一种就是conditional branch。两
者的区别在于用predicate时，可以让warp内的所有线程名义上走同一路径而省去跳转的开销，
从而也避免了可能的divergence。因为branch的latency比较长，还涉及到instruction cache的问
题，一般很短的分支是不太愿意跳转的。当然branch在出现divergence的时候，内部也有一个
mask，表明当前这个thread是否active，但是用户不能直接修改这个mask。PTX中可以通过warp
vote或是load特殊寄存器 %lanemask_* 之类的方法获得当前warp内的mask情况。

Cache control指令：这个主要意义是我知道我将要访问的某个元素位于某个cache line，但
是又不确定具体要哪一个值，所以没法先load。所以可以先把整个cacheline放到cache里，
等到要用的时候从cache里取。这样load的latency可以更好的被隐藏。PTX有相关的控制指
令，但是我在Turing上测试过几次，只测到过TLB的latency被消除，实际访问latency好像
没变，存疑。

Control code主要有reuse、read barrier、write barrier、wait barrier、yield hint、stall count等几个
域。reuse是唯一能在反汇编文本里看到的。它可以有限的解决一些GPR的bank conflict问题，也
许同时还能减少GPR的读写，节省一点功耗。关于barrier的几个主要控制thread内的依赖性问
题，这会影响指令发射的仲裁过程。yield和stall count主要影响的是warp调度的逻辑。

memory类的指令是性能优化的重点中的重点。绝大部分未经优化的程序都会是memory bound，
多数实际应用优化完了还是memory bound。所以，通过合理的选择相应指令，搭配合适的内存
排布，从而更好的隐藏内存访问的latency，或者是减少相应访问的开销，是性能优化中的主要
课题之一。

SASS里并没有直接的条件跳转指令，所有的条件跳转是用predicate实现的。

每个指令能否发射还要满足相应的依赖关系和资源需求。比如指令LDG.E R6, [R2] ;首先需要等待之前写入R[2:3]的指令完成，其次需要当前memory IO的queue还有空位，否则指令也无法下发。还有一些指令可能有conflict的情况，比如shared memory的bank conflict，register的bank conflict，atomic的地址conflict，constant memory在同一warp内地址不统一的conflict等等，这些都有可能导致指令re-issue（甚至cache miss也可能导致指令replay）。这些情况会不会重复占用dispatch port发射带宽

当前CUDA的所有架构都没有乱序执行（Out of order）

Register Reuse Cache : 编译器hint，明确一个寄存器的值会被下面的指令使用到，可以放到cache里面去，避免再次read

Wait Dependency Barrier : 有6bit，每个bit表示是否需要等待对应的dependency barrier。每个线程有6个dependency barrier，每个barrier都可以被后面的Read或Write操作设置上。设置wait dependency barrier是等待依赖的其中一种方式。SASS里面还有一个对应的指令，如DEPBAR.LE SB0, 0x0, {2,1} ;

Read dependency barrier : 有3bit，表示需要设置的6个barrier中对应的索引（0~5，对应barrier 1-6，如果不需要设置barrier，就设置为0b111）。Read dependency barrier主要是一些指令不会在一开始就把所有操作数读进去，所以需要hold住GPR的值，防止后面的指令在它读取其内容之前把GPR改掉。使用Read dependency barrier的主要是memory类的指令，但是一些转换指令如F2I/I2F之类好像偶尔也能见到。

Write dependency barrier : 与read dependency很类似，也是3bit，后面跟barrier索引。注意Read和Write两者用的dependency barrier资源是一样的，也都是上面wait的那6个。Write dependency barrier比较好理解，就是某个指令要把操作结果保存到某个GPR或是predicate中，使用barrier进行保序可以防止出现data race。不过这主要针对的是不定长latency指令。如果一个指令的latency是确定的（或者有不太长的上限），那用后面提到的stall cycle停足够长时间就可以保证没有race。

Yield hint flag : 是1bit。如果Yield，就表示下一个cycle会优先发射其他warp的指令。也许yield这个bit就是stall count的高位。只是假如这个bit不为0，那stall的cycle会>16，相当于warp被切换的概率也会大大增加。

Stall count : 有4bit，表示当前指令后需要stall指令发射的cycle数，然后再决定是不是要继续发射。

image.pngimage.png

CUDA性能优化原则

  • 寻找并行化顺序代码的方法。

  • 最大限度地减少主机和设备之间的数据传输。

  • 调整内核启动配置以最大限度地提高设备利用率。Adjust kernel launch configuration to maximize device utilization.

  • 确保全局内存访问被合并。

  • 尽可能减少对全局内存的冗余访问。 Minimize redundant accesses to global memory whenever possible.

  • 避免同一扭曲内的线程进行长序列的发散执行。Avoid long sequences of diverged execution by threads within the same warp.

TMA

TMA 允许应用程序在全局内存和共享内存之间以及同一集群中不同 SM 的共享内存区域之间双向传输 1D 和最多 5D 张量（请参阅线程块集群）。此外，对于从共享内存到全局内存的写入，它允许指定按元素归约操作，例如 add/min/max 以及按位和/或最常见的数据类型。

memcpy_async的API有三个变形, 分别是group,barrier,pipeline.

  • group: 搭配cooperative_groups::wait(group)使用,注意不是 group.sync(), 从逻辑上说，相当于动态为group插入了一个异步线程.cooperative_groups::wait只是让当前线程等待所有异步线程完成,并不会有group.sync()的效果.
  • barrier: 搭配barrier.arrive_and_wait()使用,从实现角度讲,相当于动态增大了barrier的count,这个count只有在async_memcpy完成时才会下降
  • pipeline: 约等于线程安全的queue
Back to top