Synopsis
CUDA Runtime API：函数以cuda开头。
CUDA Driver API：函数以cu开头。
延迟（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 |
性能指南¶
《CUDA中文手册》第5章