GPGPU学习笔记（一）

GPGPU概述

主要内容：相关术语与主流GPU产品介绍，并通过GEMM与分块GEMM，理解GPU代码

GPGPU(general purpose graphics unit，通用图形处理器)

并行体系结构——弗林分类法
指令流：单个程序计数器产生的指令序列
数据流：指令所需的数据及其访问地址的序列（输入、中间与输出数据）

1. 单指令流单数据流（SISD）
2. 单指令流多数据流（SIMD）：采用一条指令对多个数据进行操作，如向量处理器
3. 多指令流单数据流（MISD）：多条指令（处理单元）处理单条数据流，如脉动阵列 脉动阵列说明
4. 多指令流多数据流（MIMD）

NVIDIA GPU产品：GeForce（家庭娱乐）、Quadro（专业用途显卡）、Tesla（专业的GPGPU，专注数据计算而不是图形显示，无显示输出接口）

NVIDIA GPGPU产品（Tesla）的架构进化：Tesla、Fermi、Kepler、Maxwell、Pascal、Volta、Turing、Ampere
Fermi架构：第一个为HPC应用提供所需功能的架构
Kepler架构：为高性能科学计算设计的，双精度浮点运算能力高且强调功耗比
Maxwell架构：支持统一的虚拟内存技术，允许CPU直接访问显存和GPU访问主存
Pascal架构：增强神经网络方面的适用性
Volta架构：对流多处理器(streaming multiprocessor，SM)进行了重新设计；为深度学习设计了张量核(tensor core)；每个线程都有独立的程序计数器(PC)和堆栈，实现了独立的线程调度（在此之前每个线程束（warp）内的32threads共用一个PC）
Turing架构：引入用于加速光线追踪的RT(ray-tracing)核
Ampere架构：相较于Volta架构，将张量核性能提升2.5倍

AMD GPGPU产品：GCN架构与RDNA架构 介绍

GCN（graphics core next）架构采用GCN单元，也称为CU（compute unit），每个CU内部有4个SIMD阵列，每个SIMD阵列有16个ALU（对应于streaming processor，SP）
第四代GCN架构——Polaris架构
第五代GCN架构——Vega架构：高带宽缓存控制器、下一代计算单元（Next-generation compute unit，NCU）、高级像素引擎、新一代几何渲染引擎

单指令多线程（single instruction multiple threads，SIMT）：每个线程输入数据不同，但需要执行的指令相同
CUDA与OpenCL编程模型术语对SIMT计算模型进行了封装
CUDA：线程网格（thread grid）、线程块（thread block）、线程（thread）
OpenCL：N维网格（NDRange）、工作组（work-group）、工作项（work-item）
网格 <—> 执行内核函数启用的所有线程
线程块 <—> 分配到SM上执行的单位
线程 <—> 每个SP执行的单位

CUDA编程——GEMM

主机端代码与设备端代码
主机端：数据复制、GPGPU启动、数据写回

设备端：内核函数（kernel）


主机端启动kernel时利用<<< >>>向GPGPU传了两个参数，分别为gridDim、blockDim，均为三维参数（x,y,z），上图中gridDim=（3，2，1）、blockDim=（4，3，1）（注意：采用先排列再排行的行优先编号）
gridDim指明了主机端调用kernel时唤醒的所有线程（网格规模）
blockDim指明了分配到SM的线程单位（线程块规模）
索引线程：blockIdx、threadIdx，均为三维参数（x,y,z）
blockIdx指明了当前线程在网格中的哪个线程块
threadIdx指明了当前线程在线程块中的具体位置
假设结果矩阵C(M*N)与图2-3中线程规模一致
那么结果矩阵d_C中的一个结果索引（row,col）的计算为
row = threadIdx.x + blockIdx.x * blockDim.x
col = threadIdx.y + blockIdx.y * blockDim.y
其中，blockDim.x = 4，blockDim.y = 3，blockIdx按照图2-3中上半部分的（3,2,1）规模给出块在网格中的定位
因此，结果矩阵规模为（12,6,1）
注意，结果矩阵的存储也是按照行优先策略

分块GEMM：利用数据局部性原理减少访存次数


分块前，计算每个结果需要读取一行一列，计算结果矩阵中的同行、同列元素时，会重复读取同一行（同一列）以列数（行数）次
改为分块乘后，以线程块大小为单位搬运数据，每次搬运两块，计算得出一块部分和，循环N/BLOCK_SIZE次并累加部分和（代码2-7中9-20行），这样做实现了一次读取数据被BLOCK_SIZE次复用
分块前，需要从global memory读取2 * M * N * K次（读取A需要M * N * K，读取B需要M * N * K）
分块后，读取次数减少到2 * M * N * K/BLOCK_SIZE
两次线程同步__synthreads()，第一次是为了等待线程块内线程完成读取并缓存在shared memory再进行计算（防止读后写，而本应该读取新值），第二次是为了等待所有线程计算完部分和之后再对shared memory做下一轮读取（防止写后读，而本应该读取旧值）
问题：第二次同步应该放在内循环的外部吧，由17行换到18行后，作为外循环的结尾