内置类型和函数
线程同步
线程调度
存储模型
重访
原子函数
Global和device函数
1、尽量少用递归
2、不要用静态变量
3、少用malloc
4、小心通过指针实现的函数调用
向量数据类型
同时适用于host和device代码
通过函数make_<type_name>构造
这些向量类型都是从基本的整型和浮点型derived而来。这些类型均为数据结构，并且可以通过x,y,z,w分别访问他们的1st,2nd,3rd和4th component. 它们均可通过make_<type_name>格式的构造函数构造。比如：
int2 make_int2(int x, int y);
创建了一个int2类型的向量（x，y）;
结构体的对齐要求如下：
数学函数
sqrt rsqrt
exp log
sin cos tan sincos
asin acos atan2
trunc ceil floor
Intrinsic function 内函函数
仅面向Device设备端
更快 单精度降低
以__为前缀，如：
__exp, __log, __sin, __pow，…
线程同步
块内线程可以同步
通过___syncthreads创建一个barrier栅栏（线程块内的同步）
每个线程在调用点等待块内所有线程执行到这个地方，然后所有线程继续执行后续指令
1、要求线程的执行时间尽量接近，为什么这点很重要？
类似木桶只有一个短板，多数人等一个人浪费时间。
2、为什么只在一个块内部进行同步？
全局同步开销比较大。
3、__syncthreads()会导致线程暂停吗？
会，而且有可能导致死锁。
线程调度
G80架构
16SMs 每个SM有8个SPs 共计128个SPs
SP：Streaming processing
每个SM驻扎多达768个线程。 允许共计768个线程在这里存储，但不一定在执行。（通过context存起来 交替执行 隐藏延迟）
总过同时执行12,288个线程。
GT200架构
Warp 线程束 块内的一组线程
（一个Warp内的线程是天然同步的，是在硬件层面进行的保证）
运行于同一个SM
线程调度的基本单位
threadIdx值连续0-31 32-63 …
一个实现细节–理论上
warpSize
同一个SM上调度3块block时的warp配置（即，SM调度block时会按照warpSize的大小将block切分成对应的warp进行调度）
！！在任何一个时刻，一个SM上只有一个warp正在执行。。？？
每个warp含32个线程，但每个SM只有8个SPs。如何分配？
32个线程，每次分配8个线程给SM，然后依次执行完32个线程的第一个指令。接着按照同样的步骤做。
问题：一个kernnel包含1次对global memory的读操作（200 cycles）和4次独立的multiples/adds操作独立于访存，需要多少个warp才可以隐藏内存延迟？（设multiples/adds一次需要4个指令周期）
解答： m/a操作 16个周期，需要覆盖200个周期。200/16=12.5 ceil(12.5)=13 需要13个warps
寄存器Registers
每个线程专用
快速 片上 可读写
增加kernel的register用量会导致什么结果？
性能可能变高 可能变低，例如：
寄存器 G80
每个SM
有768 threads
共8K个寄存器
每个线程可以分到多少个寄存器？
8K / 768
每个线程最多分10个寄存器
如果每个线程must用11个寄存器，每个block包含256个线程。
每个block要用256*11个寄存器。每个block需要2816个寄存器。也就是一个SM拥有8K个寄存器 8K/2816 = 2 最多装两个block。 也就是最多有256 * 2 / 32 = 16个warp，warp个数变少，意味着低效。且多余的寄存器浪费了。
局部存储器Local Memory
存储于global memory(这玩意儿竟然是在全局内存吗？wtf)
作用域是每个thread
用于存储自动变量数组
通过常量索引访问
共享存储器 Shared Memory
每个块
快速 片上 读写
全速随机访问（用户可编程cache）
G80 共享存储器
每个SM包含：
至多8个blocks
16KB共享存储器
每个block分配多少KB？
如果每个block用5KB，一个SM可以驻扎多少block？
16KB / 5KB = 3
全局存储器 Global Memory
长延时（100个周期）
片外 可读写
随机访问影响性能
Host主机端可读写
GT200 带宽150 GB/s 容量4 GB
常量存储器 Constant Memory
短延时 高带宽 当所有线程访问同一位置时只读
存储于Global memory 但是有缓存
Host主机端可读写
容量 64KB
变量的生命周期
Global and constant 变量
Host可以通过以下函数访问：
cudaGetSymbolAddress()
cudaGetSymbolSize()
cudaMemcpyToSymbol()
cudaMemcpyFromSymbol()
Constants变量必须声明在函数外声明
