前言 本章别离从硬件层面和软件层面对 CUDA 编程模型进行描绘。主要讨论 GPU 的并行核算是如安在硬件上完成的，CUDA 中的模块了解以及 CPU和 GPU 之间的交互，指令的同步。这部分内容比较抽象和单调，期望咱们耐性看完。

欢迎重视大众号CV技能攻略，专心于核算机视觉的技能总结、最新技能盯梢、经典论文解读、CV招聘信息。

核算机视觉入门1v3辅导班

本教程制止转载。一起，本教程来自常识星球【CV技能攻略】更多技能教程，可参加星球学习。

一 GPU 架构与异构并行核算

什么是异构并行核算

最初的核算机只包括中央处理器，为了处理越来越杂乱的图形核算，GPU 营运而生，因其数据众多的轻量级线程，十分合适处理大规模异构并行核算。

下图所示是一个典型的异构并行架构，包括一个 CPU及其内存 和一个 GPU及其内存，GPU 设备端经过 PCIe 总线与根据 CPU 主机端进行交互。一个异构并行运用包括主机代码和设备代码，别离运行在主机端和设备端。运用由 CPU 初始化，在设备端进行数据运算前，CPU 负责管理设备端的环境，代码和数据。咱们称 host 为 CPU 及其内存，device 为 GPU 及其内存。

CUDA 教程（二）CUDA 模型概述

CPU 核算合适处理操控密集型使命，GPU 核算合适处理包括数据并行的核算密集型使命。在 CPU 上履行串行部分或使命并行部分，在 GPU 上履行数据密集型并行部分，这种异构并行架构使得核算能力能够充分被运用。

NVIDIA GPU 显卡架构开展历程

  • Tesla（特斯拉）2008年，运用于早期的 CUDA 系列显卡芯片中，并不是真实意义上的 GPU 芯片。
  • Fermi（费米）2010年，是第一个完整的 GPU 核算架构。首款可支撑与同享存储结合纯 cache 层次的 GPU 架构，支撑 ECC(Error Correcting Code) 的 GPU 架构。
  • Kepler（开普勒）2012年，Fermi 的优化版。
  • Maxwell（麦克斯韦）2014年，初次支撑实时的动态大局光照作用，
  • Pascal（帕斯卡）2016年，GPU 将处理器和数据集成在同一个程序包内，以完成更高的核算功率。
  • Volta（伏打）2017年，初次将一个 CUDA 内核拆分为FP32 和 INT32 两部分，初次支撑混合精度运算，提高了5倍于 Pascal 核算速度，还增加了专用于深度学习的 Tensor Core 张量单元。
  • Turing（图灵）2018年，增加了 RT Core 专用光线追寻处理器，将实时光线追寻运算加快至上一代架构的 25 倍，并能以高出 CPU 30 多倍的速度进行电影作用的最终帧烘托。去掉了对 FP64 核算的支撑。
  • Ampere（安培）2020年，重新支撑 FP64，新增异步拷贝指令能够从 global memory 中将数据直接加载到 SM shared memory，降低中心寄存器堆（RF）的需求。新增 BF16 数据类型，专为深度学习优化。

二 CUDA 编程模型

CUDA 是一个通用并行核算渠道和编程模型，如下图所示，CUDA 渠道能够经过 CUDA 加快库、编译器指令、运用程序编程接口或编程言语接口来运用。后边的章节咱们会要点讲解 CUDA C 以及 PyCUDA 的编程。

CUDA 教程（二）CUDA 模型概述

CUDA 软件体系

CUDA 提供了两层 API 来调用底层 GPU 硬件

  • CUDA 驱动 API (CUDA Driver API)
    是一种根据句柄的底层接口,大多数对象经过句柄被引用,其函数前缀均为cu，在调用 Driver API 前必须进行初始化，再创建 CUDA 上下文，该上下文相关到特定设备并成为主机线程的当时上下文，经过加载 PTX 汇编形式 或 二进制对象形式 的内核，然后发动内核核算。Driver API 能够经过直接操作硬件履行一些杂乱的功能，但其编程较为杂乱，难度较大。
  • CUDA 运行时 API (CUDA Runtime API)
    Runtime API 对 Driver API 进行了必定的封装，隐藏了部分完成细节，因而运用起来更为便利，因而咱们更多运用的是 Runtime API。Runtime API 没有专门的初始化函数，它将在第一次调用运行时函数时主动完成初始化。运用时，一般需要包括头文件 cuda_runtime.h，其函数前缀均为cuda。

如下图所示

CUDA 教程（二）CUDA 模型概述

Runtime API 和 Driver API 之间没有显着的性能差距，这两种 API 不能混合运用，只用独自运用其一。

CUDA 函数库 (CUDA Libraries)

CUDA 提供了几个较为成熟的高效函数库，能够直接调用这些库函数进行核算，常见的包括

  • CUFFT：运用 CUDA 进行傅立叶变换的函数库
  • CUBLAS：运用 CUDA 进行加快的完整标准矩阵与向量的运算库
  • CUDPP：并行操作函数库
  • CUDNN：运用CUDA进行深度卷积神经网络

CUDA 运用程序 (CUDA Application)

CUDA 程序包括在 host 上运行的主机代码和在 device 上运行的设备代码，设备代码会在编译时经过 CUDA nvcc 编译器从主机代码中分离，再转换成 PTX(ParallelThread Execution) 汇编言语，由 GPU 并行线程履行，主机代码由 CPU 履行。如下图所示

CUDA 教程（二）CUDA 模型概述

履行流程如下

  • 分配 host 内存，并进行数据初始化（CPU初始化）
  • 分配 device 内存，并从 host 将数据拷贝到 device 上（GPU初始化）
  • 调用 CUDA 的核函数在 device 上完成指定的运算（GPU并行运算）
  • 将 device上的运算结果拷贝到 host 上（将GPU结果传回CPU）
  • 开释 device 和 host 上分配的内存（初始化清空）

CUDA 硬件结构

  • SP（Streaming Processor）也称为 CUDA core，是最基本的处理单元，最终详细的指令和使命都是在 SP 上处理的。GPU 进行并行核算，也便是很多个 SP 一起做处理。
  • SM（Streaming Multiprocessor）多个 SP 加上其他资源组成一个 SM，也叫 GPU 大核，其他资源如包括warp scheduler，register，shared memory 等。SM能够看做GPU的心脏（类似 CPU 核心）。每个 SM 都具有 register 和 shared memory，CUDA 将这些资源分配给一切驻留在 SM 中的线程，但资源十分有限，SM 结构如下图所示。

CUDA 教程（二）CUDA 模型概述

每个 SM 包括的 SP 数量根据 GPU 架构而不同，如 Fermi 架构 GF100 是 32 个，GF10X 是 48 个，Kepler 架构都是 192 个，Maxwell 都是128 个。

CUDA 教程（二）CUDA 模型概述

在软件逻辑上是一切 SP 是并行核算的，但是物理上并不是，比如只有 8 个 SM 却有 1024 个线程块需要调度处理，由于有些会处于挂起，就绪等其他状况，这有关 GPU 的线程调度，后续章节会展开讨论。

三 了解 kernel, thread, block , grid 与 warp

CUDA 线程模型

线程是程序履行的最基本单元，CUDA 的并行核算经过不计其数个线程的并行履行来完成。下图为 GPU 的线程结构

CUDA 教程（二）CUDA 模型概述

CUDA的线程模型从小往大依次是

  • Thread，线程，并行的基本单位

  • Block，线程块，互相合作的线程组，线程块有如下几个特点：

  • 以1维、2维或3维安排

  • 允许彼此同步

  • 能够经过同享内存快速沟通数据

  • Grid，网格，由一组 Block 组成

  • 同享大局内存

  • 以1维、2维安排

kernel

kernel 是在 device 上线程中并行履行的函数，是软件概念，核函数用__global__符号声明，并用 <<<grid, block>>> 履行配置语法指定内核调用的 CUDA 线程数，每个 kernel 的 thread 都有一个仅有的线程 ID，能够经过内置变量在内核中访问。block 一旦被分配好 SM，该 block 就会一直驻留在该 SM 中，直到履行结束。一个 SM 能够一起具有多个 blocks。

warp

warp 是 SM 的基本履行单元，也称线程束，一个 warp 有 32 个并行的 thread， SM 旨在一起履行数百个 thread，为了管理如此很多的线程，采用了 SIMT （Single-Instruction, Multiple-Thread：单指令，多线程）的架构，也便是一个 warp 中的一切 thread 一次履行一条公共指令，并且每个thread会运用各自的data履行该指令。

一个块中的 warp 总数核算如下

CUDA 教程（二）CUDA 模型概述

对应下图

CUDA 教程（二）CUDA 模型概述

从硬件视点来看，一切的 thread 以一维形式安排，每个 thread 都有个仅有的 ID，所以作为补全整数倍的 thread 在地点的 warp 中为 inactive 状况，会额外消耗 SM 资源，所以要设定 block 中的 thread 一般为32的倍数。

下面从硬件视点和软件视点解说 CUDA 的线程模型

软件

硬件

描绘

Thread

SP

每个线程由每个线程处理器（SP）履行

Block

SM

线程块由多核处理器（SM）履行

Grid

Device

一个 kernel 由一个 grid 来履行，一次只能在一个 GPU 上履行

线程索引

确定线程的仅有索引，以 2D grid 和 2D block 的状况为例。

咱们要核算的数值矩阵在内存中是 row-major（行主序） 线性存储的，如下图

CUDA 教程（二）CUDA 模型概述

将 thread 和 block 索引映射到矩阵坐标

ix = threadIdx.x + blockIdx.x * blockDim.x

iy = threadIdx.y + blockIdx.y * blockDim.y

idx = iy * nx + ix

下图为 block 和 thread 索引，矩阵坐标以及线性地址之间的联系

CUDA 教程（二）CUDA 模型概述

在实践运用中，常常会多一维 grid， 那便是三维状况的索引，如下图所示，设 (gridDim.x,gridDim.y) = (2,3), (blockDim.x,blockDim.y) = (4,2)，咱们以 thread_id(3,1) block_id(0,1) 为例

CUDA 教程（二）CUDA 模型概述

能够得到

ix = threadIdx.x + blockIdx.x * blockDim.x = 3 + 0 * 4 = 3

iy = threadIdx.y + blockIdx.y * blockDim.y = 1 + 1 * 2 = 3

coordinate(3,3)

global index: idx = iy * blockDim.x * gridDim.x + ix = 3 * 4 * 2 + 3 = 27

本教程制止转载。一起，本教程来自常识星球【CV技能攻略】更多技能教程，可参加星球学习。

欢迎重视大众号CV技能攻略，专心于核算机视觉的技能总结、最新技能盯梢、经典论文解读、CV招聘信息。

核算机视觉入门1v3辅导班

【技能文档】《从零建立pytorch模型教程》122页PDF下载

QQ沟通群：444129970。群内有大佬负责解答咱们的日常学习、科研、代码问题。

其它文章

本来Transformer便是一种图神经网络，这个概念你清楚吗？

3D方针检测中点云的稀疏性问题及解决方案

一文深度分析扩散模型终究学到了什么？

OpenMMLab教程【零】OpenMMLab介绍与装置

代码实战：YOLOv5完成钢材外表缺点检测

TensorRT教程（六）运用Python和C++部署YOLOv5的TensorRT模型

常识蒸馏的搬迁学习运用

超全汇总 | 核算机视觉/主动驾驶/深度学习材料合集！

高精度语义地图构建的一点考虑

一文看尽深度学习中的各种注意力机制

点云切割练习哪家强？监督，弱监督，无监督还是半监督？

姿态估量端到端新方案 | DirectMHP：用于全规模视点2D多人头部姿势估量

深度了解变分自编码器(VAE) | 从入门到通晓

核算机视觉入门1v3辅导班

核算机视觉沟通群

用于超大图画的练习战略：Patch Gradient Descent

CV小常识讨论与分析（5）究竟什么是Latent Space?

【免费送书活动】关于语义切割的亿点考虑

新方案：从过错中学习，点云切割中的自我规范化层次语义表示

经典文章：Transformer是如何进军点云学习领域的？

CVPR 2023 Workshop | 首个大规模视频全景切割竞赛

如何更好地应对下流小样本图画数据？不平衡数据集的建模的技巧和策

Transformer沟通群

U-Net在2022年相关研讨的论文推荐

用少于256KB内存完成边缘练习，开销不到PyTorch千分之一

PyTorch 2.0 重磅发布：一行代码提速 30%

Hinton 最新研讨：神经网络的未来是前向-前向算法

聊聊核算机视觉入门