异步¶
由于CPU与GPU的异步执行性,可以实现异步函数。GPU执行计算密集型任务时,CPU将返回,例如去处理网络IO等。
warp¶
warp是SM级别的,单个线程是SP级别的,SP每次运行一个线程。 每个SM的core数量决定其一个始终周期内可以运行的warp指令。 单个warp指令可以在任何给定的时钟周期内处理,但需要32个core才能完成(并且可能需要多个时钟周期才能完成,具体取决于指令)。 SM将block“解包”为warp,并在SM内部资源(例如core和SFU)可用时调度warp指令。
时间延迟(Latency)¶
指令延迟¶
数据类型 | 字长(bits) | 操作 | 延迟(cycles) | 备注 |
---|---|---|---|---|
整数 | + | 4 | ||
整数 | 32 | * | 16 | |
整数 | /,mod | 特别昂贵 | 尽可能避免,或使用位操作替代 | |
位操作 | 4 | |||
比较 | 4 | |||
max,min | 4 | |||
类型转换 | 4 | |||
浮点 | +,*+ | 4 | ||
浮点 | / | 36 | ||
倒数 | 16 | |||
平方根倒数 | 16 | |||
浮点 | 平方根 | 32 | 实现:平方根倒数+倒数,对0和无穷大能正确得到结果;因此它处理一个warp花费32个时钟周期。 | |
__log(x) | 16 | |||
__sin(x) | 32 | |||
__cos(x) | 32 | |||
__exp(x) | 32 | |||
__fdividef(x,y) | 20 | 除法 | ||
- | - | 其他函数 | 更多时钟周期 | 因为实现为多个指令的组合 |
- | - | 内存指令 | 4 | 包括从共享或全局内存中读写的任何指令,SM使用4个时钟周期来发射warp的一个内存指令 |
- | - | __syncthreads() | 4 | 同步 |
有时候,编译器不得不额外插入转换指令,从而引入额外的执行周期:
对操作数为char或short的函数操作,其操作数通常转换为int;
在单精度浮点数计算中,将双精度浮点数常量(不使用任何类型后缀定义)作为输入值;
在数学函数的双精度版本,将单精度浮点数变量作为输入参数。
后两种情况可以通过以下方式避免:
对单精度浮点数常量,使用f后缀定义,比如3.141592653f
对数学函数的单精度版本,也使用f后缀定义,比如sinf().
访存延迟¶
内存类型 | 延迟 |
---|---|
L1/Shared Memory | 10-20 |
Global Memory | 400-600 |
架构¶
每个SM都有L1 cache,所有的SM共享L2 cache; L1 cache:用于缓存local memory,包括临时寄存器溢出。 L2 cache:用于缓存global memory。 缓存行为(例如,读操作同时使用L1和L2 cache,还是仅仅使用L2 cache)可以使用load或者store指令修饰符来部分配置,3.5以上计算能力设备允许通过编译选项配置。
对于计算能力3.x的设备
on-chip memory(unified data cache)可以在shared memory和L1 cache之间分配(使用cudaFuncSetCacheConfig()/cuFuncCacheConfig()):
(1)48 KB作为shared memory,16 KB作为L1 cache; (2)32 KB作为shared memory,32 KB作为L1 cache。 对于计算能力5.x的设备 on-chip memory(unified data cache)可以在shared memory和L1 cache之间分配(使用cudaFuncSetCacheConfig()/cuFuncCacheConfig()): (1)48 KB作为shared memory,16 KB作为L1 cache; (2)32 KB作为shared memory,32 KB作为L1 cache。
设备内存访问¶
内存分布对指令吞吐量的影响:根据一个warp内的线程和内存地址的分布,可能会出现内存访问指令被发起(re-issue)多次的情况。例如,全局内存地址越分散,指令吞吐量就越小。 指令吞吐量:执行一条指令所需的时钟数。
全局内存(global memory / device memory)¶
内存事务(memory transactions):内存事务指一次内存访问。可以以32字节、64字节、128字节访问。 合并访问(coalesces the memory accesses):一个warp内的线程的内存访问会被合并。通常,内存事物越多,线程束的指令利用率越低,指令吞吐量越小。例如,32字节的内存,如果每个线程分配成4字节访问,那么指令吞吐量会降低为1/8倍。
1个warp包含32线程,全局内存的缓存行(cache line)是128字节。 如果内存事务为128字节,那么L1和L2缓存都会被用到; 如果内存事务为32字节,那么只有L2缓存会被用到; L2缓存只能降低访存次数(over-fetch),例如,对地址分散的内存的访问。 如果每个线程对全局内存的访问超过4个字节,那么1个warp的内存访问请求会被分割成128字节的内存访问: 32 * 8 / 128 = 2次内存请求,每次请求为1/2个warp,如果每个线程请求8字节; 32 * 16 / 128 = 4次内存请求,每次请求为1/4个warp,如果每个线程请求16字节。
对于整个kernel生命周期内的只读数据,可以使用__ldg()函数(可能是load global的缩写)缓存在read-only data cache中。编译器也会使用__ldg()进行这方面的尝试,当使用const和__restrict__限定符标记这部分数据指针时,会增加编译器探测read-only条件(进而进行上述缓存优化)的概率。
内存对齐会增加cache概率,减少访存(memory transaction)次数
共享内存(shared memory)¶
bank conflict:存储体/存储池冲突
不同计算能力(Compute Capabilities)比较¶
3.x | 5.x | 6.x | 7.x | |
---|---|---|---|---|
L1 cache | L1/shared memory | L1/texture | ||
Shared Memory | (Max) 48 KB | 64 KB | ||
Texture Memory | - | |||
L2 cache | ||||
Constant cache | ||||
CUDA cores / MP | 192 | 128 | 64 (6.0), 128 (6.1, 6.2) | |
Special Function Units | 32 | 32 | ||
warp Schedulers | 4 | 4 |
D. 执行环境与内存模型¶
D.2. 执行环境¶
D.2.1. 父子线程格(Parent and child grids)¶
Parent grid中的线程可以创建grid,称为线程格嵌套(grids nesting) 嵌套grids会隐式同步。