01、cuda基本概念

前言

以下信息介绍来自樊哲勇的cuda书籍，仅自己学习记录

1、GPU硬件介绍

GPU（graphics processing unit），图形处理器，俗称显卡，但并不是显卡，显卡中包含GPU。

1）与CPU的区别

• CPU有更多晶体管，用于数据缓存和流程控制，只有少数几个逻辑计算单元，适合完成简单的逻辑计算

• GPU有数千个核心，适合大规模矩阵运算。
  

• 常见的GPU系列：
  

2）CPU+GPU 异构计算平台

• CPU：host ，主机

• GPU：device 设备

3）计算性能

• 浮点数运算峰值：

• 内存带宽（显存）

• 计算能力不等价计算性能，不是简单地正比关系
  

4）CUDA提供了两层API

• CUDA driver API(底层，灵活，不好用),

• cuda runtime API(高级，可读性强)，

两者性能无差别

5）nvidia-smi 检查设备信息

export CUDA_VISIBLE_DEVICES=1 设置使用的GPU编号

2、cuda 线程组织

1）nvcc 编译cuda文件 ，./a.out 执行

2）kernel 核函数

• 被限定词__global__来修饰
• 必须返回void

<<<grid, block>>>是标准写法，grid规定了block的数量，而block规定了线程的数量。

3）概念

• 一个线程束由32个连续的线程构成，选取id连续的thread，是CUDA执行的基本单位，不同的warp之间时独立执行的
• 一个线程块内的所有线程会在同一处理器内核上共享资源，所以块内的线程是有限制的
• 一个block中的线程可以协作，不同block的线程不可以
• 一个线程网格由若干个线程块组成
• 一个CUDA kernel对应一个grid

4）gridDim.x、blockDim.x && blockIdx.x、threadIdx.x

• gridDim.x/y/z：一个grid在指定方向上有多少数量的线程
• blockDim.x/y/z：一个block在指定方向上有多少数量的线程
• blockIdx.x/y/z：在一个grid中，当前线程在指定方向上的blockId
• threadIdx.x/y/z：在一个block中，当前线程在指定方向上的threadId
• block和thread都是最内层是x维，最外层是z维

3、函数执行空间标识符

函数前缀__device__、host、global

• host：在CPU中调用，由CPU执行，普通函数
• global：表示内核函数，由GPU执行
• device：表示一个由GPU中线程调用的函数

4、cuda 程序的错误检测

1）检查运行时api函数

• CHECK（）

2）检查核函数

• CHECK(cudaGetLastError());

• CHECK(cudaDeviceSynchronize());

3）检查内存错误

• cuda-memcheck [options] app_name [options]

5、GPU加速的关键

1）cuda时间计时

• timing 函数

• nvprof 性能分析工具

2）影响GPU加速的关键因素

• 数据传输比例

• 算术强度

• 并行规模

3）cuda 中数学函数库

经常使用才会记得

6、cuda 内存组织

1）内存组织简介

容量和延迟的关系

2）SM及其占有率

2.1 SM 组成

寄存器、共享内存、常量内存的缓存、纹理和表面内存的缓存、L1缓存、两个线程束调度器、执行核心

2.2 SM 的占有率

• 线程块大小
• 寄存器数量
• 共享内存数量

7、全局内存

1）合并和非合并访问

• 合并度：线程束请求的字节数和该请求导致的数据传输处理的字节数之比。

• 可以视为一种资源利用率的表征。

• 100%则为合并访问，否则为非合并访问

2）cudaMalloc()分配的内存首地址至少是256字节的整数倍。

3）几种常见的内存访问模式及其合并度

（1）顺序的合并访问

（2）乱序的合并访问

（3）不对齐的非合并访问(地址错位，必须要多一次地址访问)

（4）跨越式非合并访问（数据是横着存的，结果竖着取数据）

（5）广播式非合并访

问(例如，一次取了32字节的浮点数，但每次只用其中的一个浮点数，合并度为4/32=12.5%)

4）矩阵转置查看效果

8、共享内存

共享内存可以被程序员直接操控

1)作用

• 减少对全局内存的访问，实现高效的线程块内部通信；

• 也可以提高全局内存访问的合并度。

2）核函数对于共享内存的操作作用在所有线程上。

• __syncthreads() 是轻量级的，并且是以block 级别做同步。

3）静态共享内存和动态共享内存

• 静态共享内存
  shared 变量名可以以s_开头

• 动态共享内存
  <<<grid_size, block_size, sizeof(real) * block_size>>> 定义核函数的时候，第三个参数其实是共享内存的大小，被默认设为了0。

extern shared real s_[] 注意这里是数组，不能定义成指针

4）共享内存在物理上被分为32个同样宽度、能被同时访问的内存bank，每个bank的宽度为4B(只有kepler架构为8B)，地址模128的结果相同的属于同一层.

5）数组规约计算

9、原子函数

1）原子操作

不受其他线程影响，完成对某个（全局内存或者共享内存）数据的一套“读-写-改”操作。

2）规约的进一步计算

3）常见的原子函数

10、线程束基本函数和协作组

1）线程束

• 一个GPU由多个SM组成
• 一个SM上可以放多个线程块，不同线程块之间并行或者顺序执行
• 一个线程块分为多个线程束
• 一个线程束由32个连续的线程号组成
  从更细粒度来看，一个SM以一个线程束为单位产生、管理、调度、执行线程

2）伏特架构

• 之前

__synthreads()可以使线程块同步

每个warp只有一个程序计数器，需要注意分支发散问题。一些严重的分支发散会极大降低性能。

• 之后

_synwarp()线程束内同步

加入了独立线程调度机制后，可以允许线程的同步和通信，粒度更细，操作更灵活。会牺牲一点寄存器。

3）协作组

线程块级的协作组，

规约计算

11、cuda 流

cuda并行层次主要有两个：

• 核函数内部并行

• 核函数外部并行

1）cuda 流

概念：由主机发出的在一个设备中执行cuda操作（即和cuda有关的操作，如主机-设备数据传输和和函数执行）序列。

一个cuda流中各个操作的次序是由主机控制的，按照主机发布的次序执行。

默认流（空流）

流的定义、产生和销毁

2)默认流中重叠主机和设备计算

例子

3）非默认流cuda流重叠多个核函数的执行

4）非默认流cuda流重叠核函数的执行和数据传输

12、使用统一内存编程

1. 好处：

• 使编程更简单，不用手动copy内存。
• 可能提供更好的性能
• 允许超量分配

2）动态统一内存、静态统一内存

3. 实例分析

13、cuda标准库的使用