1. Gpu 适合大规模并行计算场景 cpu更适合串行任务优化&线程数量少
  2. gpu结构：global memory (显存) + 流处理器（里面有很多 cuda core），通过PCIE与cpu连接
    1. sm(流处理器)
    2. 存储 globalmem->L2->L1/SMEM/C -> 寄存器
    3. gpu之间可以通过nv link(40GB/s)，避免pcie慢速
  3. cuda编程：主函数: host(cpu)->kernel(gpu)->host(单线程)
    1. Nvcc编译，__global____标识gpu程序入口
    2. 显存分配cudaMalloc cudaMemset cudaFree
    3. 数据从cpu端->显存，__host__ cudaMemcpy() 异步拷贝：cudaMemoryAsync
    4. 现有技术： 从ssd直接读入显存，从另一个服务器的gpu通过网卡直接读入本机gpu，不需要经过cpu处理
    5. 运行模型：1.Grid，包含很多block，2.block(单个流处理器中)包含很多thread
    6. <<<Grid, Block>>> 在main函数中指定grid，block个数(并发使用，如果sm间没有数据交互，优先使用SM(也就是Grid尽量大)) thread_idx = blockidx.x * blockdim.x + threadidx.x
    7. 推荐尽量多的使用线程，而不是根据sm数量分配线程

cuda优化

cudnn：深度学习相关库

cublas：矩阵相乘

cuSPARSE: 稀疏矩阵计算

NCCL: 自动化处理多GPU数据传输

TensorRT: 重要

TensorRT

  1. Fuse network layers(FP16 与 INT8) 1.5-4倍加速比
  2. Onnx/tensor模型文件->build(自动调优)->deploy

优化手段：

  1. vec相乘->矩阵相乘

  2. 使用评估工具，内存带宽瓶颈 or 计算瓶颈

内存瓶颈
  1. warp是内存访问单位，32B为一个块，抓进来的warp有一部分数据是程序不需要的，存在一部分的访问浪费
  2. 主要避免：非连续的内存分配，这样读入数据会浪费大量的内存带宽
  3. Texture 纹理缓存，只读的
  4. shared memory 用户可控，可以用于多个线程间的数据同步
延迟受限
  1. 需要足够多的warp隐藏延迟（重点）
  2. 使用低延迟指令代替高延迟指令
  3. Occupancy calculator 延迟计算工具，occpancy不是越大越好
指令优化
  1. 指令计算误用，int32->double性能会急剧下降
  2. 分支语句优化，把所有的if放到一个warp，所有else放到一个warp，避免一个warp既走if又走else
  3. 需要注意 使用float而不是double，如1.0f
  4. 特殊函数，cos，平方等，精度要求不高的操作 __func()计算速度要高于普通func
数据传输优化

  1. 尽量避免cpu->gpu数据搬移

  2. 小量数据传输进行打包，批量传输

  3. async方式的数据传输，cpu需要使用pinned(不会分页) memory，即不适用malloc, 数据传输效率提升70%左右