一、nvprof分析线程束和内存读写
（1）线程束占用率分析
线程束占用率：nvprof --metrics achieved_occupancy
（2）内存读写分析
内核数据读取效率：nvprof --metrics gld_throughput
程序对设备内存带宽利用率：nvprof --metrics gld_efficiency
二、GPU动态并行
1. 动态并行概念
（1）在GPU上动态创建内核并同步内核执行
（2）通过动态并行技术可以在内核运行时设置线程模型配置
（3）动态并行技术有效减少GPU、CPU间执行权限的切换和数据传输
2. 动态并行特点
（1）线程网格、线程块、线程模型配置以及内核函数启动方式在动态并行中仍然适用
（2）动态并行中内核执行分为两类：父内核执行（父线程、父线程块、父线程网格）、子内核执行（子线程、子线程块、子线程网格）
（3）父内核执行由主机程序发起、子内核执行由父内核执行发起
（4）父内核与子内核共享全局（global）和常量（constant）内存
（5）父内核与子内核具有各自独立的本地（local）和共享（shared）内存
三、GPU内存结构
CUDA内存模型将独立的主机内存和GPU设备内存作为整体形成完整的内存层次结构
（1）非可编程内存
（2）可编程内存
1. 寄存器
（1）速度最快
（2）内核函数中无修饰符的自动变量
（3）数组索引为常量或在编译时有确定数值时，数组可以保存在寄存器中
（4）寄存器中数据为每个线程独有，并具有与内核函数相同的生命周期
（5）查看寄存器使用：nvcc --resource-usage 源程序
四、寄存器溢出
内核函数使用的寄存器数量超过硬件限制时，数据会被保存到线程的本地内存（local memory）中
1. 使用控制
（1）内核函数 __launch_bounds__
（2）nvcc编译参数-maxrregcount，（设置__launch_bounds__时该参数被忽略）
五、本地内存和共享内存
本地内存
（1）本地内存由每个线程独有，延迟比寄存器大
（2）寄存器溢出时会被保存到本地内存
（3）内核编译时无法确定索引的数组保存在本地内存中
（4）结构体和大数组保存在本地内存中
（5）无法保存在寄存器中的其他数据
共享内存
（1）__shared__修饰，低延迟、高带宽
（2）可被线程块中所有线程访问
（3）具有与其线程块相同的生命周期
（4）是一种线程间通信机制
（5）对共享内存的访问必须要做同步处理，__syncthreads()
（6）流处理器中的L1缓存和共享内存 共享 片上的64K存储区域
（7）动态配置共享内存：cudaFuncSetCacheConfig
六、常量内存
（1）常量内存（constant memeory）是GPU设备上的内存区域，每个流处理器有独立的常量内存
（2）常量变量必须由__constant__关键字修饰
（3）常量变量必须在全局域中声明
（4）常量内存对程序中所有的内核可见
（5）内核函数只能读取常量内存中的数据
（6）常量内存大小为64K
初始化
（1）常量内存在主机程序中初始化：cudaMemoryToSymbol
七、全局内存
概念
（1）全局内存是GPU上容量最大、延迟最大、使用最多的内存空间
（2）全局内存具有与程序相同的生命周期
（3）全局内存可被所有流处理器访问
（4）全局内存首字节地址必须是32字节、64字节、128字节的整数倍
初始化
（1）使用__device__关键字静态声明全局内存
（2）主机代码中使用cudaMalloc动态声明全局内存、cudaFree释放全局内存
