#pragma unroll，减少循环控制的开销，如条件判断、分支跳转、计数器更新；提高指令级并行．
__builtin_assume(cond) 向编译器断言表达式 cond 在此时为 true，编译器可能启用特定优化
__restrict__ 应用于指针声明，向编译器承诺 __restrict__ 修饰的指针是数据唯一的访问途径。免除 aliasing 导致保守优化．
__const__ 将数据放在常量内存，支持广播机制
__ldg() 通过纹理内存或者只读内存缓存
编译器选项
-prec-div=false 禁用高精度除法对除法进行加速，但是牺牲精度
-ftz=true
-use_fast_math 近似计算数学函数，精度换速度
-maxrregcount=N 限制线程寄存器使用数量
-Xptxas -v 输出编译信息
-Xptxas -dlcm=<> 控制 L1 缓存的行为，影响内存访问性能
- ca (Cache All)
- cg (Cache Global)
- cs (Cache Streaming)