GPU异构计算基础知识

article/2025/9/1 20:19:11

CUDA Toolkit Documentation (nvidia.com)

  • host CPU和内存 (host memory)
  • Device GPU和显存 (device memory)

SIMT模型与SIMD模型的区别

SIMD(Single Instruction Multi Data，单指令多数据)模型要求同一个向量中的所有元素要在统一的同步组中一起执行，而SIMT则允许属于同一个线程束的多个线程独立执行，这几个线程可以有不同的行为。因此SIMT(Single Instruction Multi Thread)允许线程级并发，也就是在统一线程束下的线程可以同时做不同的事情。

SIMT有SIMD所不具备的三个特征：

  1. 每个线程都有自己的指令地址计数器
  2. 每个线程都有自己的寄存器状态
  3. 每个线程可以有一个独立的执行路径

硬件基础

系统每一bit的输入与输出，包括磁盘、网络控制器、键盘、鼠标、USB设备以及GPU的输入输出，都要通过芯片组。直到最近，芯片组被一分为二：一个是连接大多数外围设备和系统的“南桥”，另一个是包含图形总线(加速图形端口，后被PCIe接口取代)和内存控制器(通过前端总线与内存相连)的“北桥”。

外部数据总线是中央处理器CPU(Central Processing Unit)的一部分，是CPU与外部数据传输的通道。外部数据总线一次可传输二进制数据的位数越大，CPU与外部交换数据的能力越强。

GPU内存控制器总是和GPU集成，它是在一个与CPU内存控制器完全不同的约束集下设计。

在这里插入图片描述
在这里插入图片描述
在这里插入图片描述

CPU的缺陷

CPU 与 GPU 架构的一个主要区别就是 CPU 与 GPU 映射寄存器的方式。 CPU 通过使用寄存器重命名和栈来执行多线程。为了运行个新任务, CPU 需要进行上下文切换,将当前所有寄存器的状态保存到栈(系统内存)上,然后从栈中恢复当前需要执行的新线程上次的执行状态。这些操作通常需要花费上百个 CPU 时钟周期。如果在 CPU 上开启过多的线程,时间儿乎都将花费在上下文切换过程中寄存器内容的换进/换出操作上。因此,如果在 CPU 开启过多的线程,有效工作的吞吐量将会快速降低。
然而, GPU 却恰恰相反。 GPU 利用多线程隐藏了内存获取与指令执行带来的延迟。因此,在 GPU 上开启过少的线程反而会因为等待内存事务使 GPU 处于闲置状态。此外, GPU 也不使用寄存器重命名的机制,而是致力于为每一个线程都分配真实的寄存器。因此,当需要上下文切换时,所需要的操作就是将指向当前寄存器组的选择器(或指针)更新,以指向下一个执行的线程束的寄存器组,因此儿乎是零开销。
一个线程束即同时调度的一组线程。在当前的硬件中,一个线程束包含32个线程。因此,在一个 SM 中,每次换进/换出、调度都是32个线程同时执行。
每个 SM 能调度若干个线程块。在 SM 层,线程块即若干个独立线程束的逻辑组。编译时会计算出每个内核线程需要的寄存器数目。所有的线程块都具有相同的大小,并拥有已知数目的线程,每个线程块需要的寄存器数目也就是已知和固定的。因此, GPU 就能为在硬件上调度的线程块分配固定数目的寄存器组。

线程层次

在这里插入图片描述

每32线程分为一个warp

GigaThread/Grid：GPU级别的调度单位，共享L2级别的Cache

SM(Streaming Multiprocessor)/block：SM调度block(线程块)，共享L1级别的Cache

Warp/Threads：一个Warp包含32个hread，Warp是CUDA core级别(一个thread对应一个CUDA core，但一个CUDA core可以通过分时复用的方式实现多个thread)的调度单位，允许同一个warp中的thread读取其他thread的值，共享L0级别的Cache

每一个warp共享一个contex执行上下文 取指译码器

  • Thread : sequential execution unit

    • 所有线程执行相同的核函数
    • 并行执行

  • Block : a group of threads

    • 执行在同一个SM(Streaming Multiprocessor)
    • 同一个Block中的线程共享一块Shared Memory（L1级别的Cache）
    • CUDA不支持超过1024个线程的线程块

  • Grid : a collection of thread blocks

    • 一个Grid当中的Block可以在多个SM中执行，一个SM可以有很多Block

每个SM所能容纳的block数量受到以下两方面限制：

  • 每个SM可提供的寄存器空间固定。
    如果每个线程所需的寄存器过多，则每个SM能够调度的线程数量就受到限制。
  • 每个SM能够调度的线程束的数量有限制。

CUDA的执行流程

  1. 加载核函数
  2. 将Grid分配到一个Device(GPU)
  3. 根据<<<…>>>内的第一个参数，Giga thread eigine将block分配到SM中。一个Block内的线程一定会在同一个SM内，一个SM可以有很多Block。
  4. 根据<<<…>>>内的第二个参数，Warp调度器将调用线程
  5. Warp调度器为了提高运行效率，会将每32个线程分为一组，称作一个warp。block(x,y,z)按先X方向，再Y方向，再Z方向优先组合成warp
  6. 每个warp会被分配到32个Core上运行。

CPU与GPU交互、多GPU间交互原理

Page-locked or pinned memory transfers attain the highest bandwidth between the host and the device.

Pinned memory is allocated using the cudaHostAlloc() functions in the Runtime API.

cudaHostRegister() 可以将host memory变为锁页内存.

计算和内存拷贝异步执行

cudaMemcpy()是主机内存与设备内存间阻塞式数据传输的API，只有当数据传输完成后才会将控制权返回主机线程。

cudaMemcpyAsync()是主机锁页内存与设备内存间非阻塞式数据传输的API，会立即将控制器返回主机线程。比cudaMemcpy()多一个参数Stream ID.

  • 锁页主机内存：GPU可以直接访问的CPU内存；
  • 命令缓冲区：由CUDA驱动程序写入命令，GPU从此缓冲区读取命令并控制其执行；
  • CPU/GPU同步：指的是CPU如何跟踪GPU的进度。

CPU/GPU并发

CUDA内核的所有启动都是异步的：CPU通过将命令写入命令缓冲区来请求启动内核，然后直接返回，而不检查GPU进度。内存复制也可以采用异步方式，这使CPU/GPU并发以及可能使内存复制与内核处理并发执行。

命令缓冲区和同步位置都位于锁页主机内存上，被CUDA驱动程序和GPU使用，以跟踪相应进程。主要的GPU操作后都跟着一个将新进度值写入共享同步位置的命令。如下图所示，进度值将一直为3,直到GPU完成当前命令的执行并将4写人同步位置中为止。

在这里插入图片描述

CUDA既隐式又显式地暴露了这些硬件功能。上下文范围的同步通过简单调用cuCtxSynchronize()、cudaThreadSynchronize()等函数来检查GPU请求的最近同步值，并且一直等待，直到同步位置获得该值。例如，在下图中，如果由CPU写人的命令8之后紧接着cuCtxSynchronize()或cudaThreadSynchronize()函数，则驱动程序将一直等待，直到共享同步值大于或等于8为止。

在这里插入图片描述

CUDA 事件则更明确地暴露了这些硬件能力。cuEventRecord()函数的作用是将一个命令加入队列使得一个新的同步值写人共享同步位置中,cuEventQuery()和 cuEventSynchronize()则分别用于检查和等待这个事件的同步值。早期版本的 CUDA 只是简单地轮询共享的同步位置,反复地读内存,直到等待准则满足为止。但是这种方法代价很大,且只有当应用程序不必等待太久时才有用(即同步位置不一定要被读取很多遍,就可以因等待标准已经得到满足而退出)。对大多数应用程序来说,基于中断的方案( CUDA 公开称为“阻塞同步”)更好,因为它们使 CPU 等待线程挂起,直到 GPU 发出中断信号为止。驱动程序将 GPU 中断映射到一个特定的线程同步原语平台,如Win32事件或 Linux 的信号。
通过指定 CU_CTX _BLOCKING_SYNC 到 cuCtxCreate()或指定 cudaDeviceBlockingSync 到 cudaSetDeviceFlags()，应用程序可以强制使用上下文范围的同步进人阻塞状态。然而,使用阻塞的 CUDA 事件(指定 CU_EVENT_BLOCKING_SYNC 到 cuEventCreate()或指定 cudaEventBlockingSync 到 cudaEventCreate() )更可取，因为它们粒度更细且可以与任何类型的 CUDA 上下文进行无缝互操作。
敏锐的读者可能会关注 CPU 和 GPU 在不使用原子操作或其他同步原语的情况下读取和写人这种共享的内存位置。但由于 CPU 只读取共享位置,竞争条件并不是关注点。这样,最坏的情况是 CPU 读取了一个“过时的”值,从而导致其等待时间要比实际的长。

CPU接口与内部GPU同步

GPU 可能包含多个引擎,以使内核执行和内存复制并发进行。在这种情况下, GPU 的驱动程序将写人一些命令,这些命令被分发到同时运行的不同引擎中。每个引擎都有自己的命令缓冲区和共享同步值。图2-28显示的是两个复制引擎和一个计算引擎并行工作时的情形。其中,主机接口负责读取命令并将其调度到相应引擎。在图2-28中,一个主机到设备的内存复制操作和两个相关操作(一个内核启动和一个设备到主机的内存复制)已被提交给硬件。依据 CUDA 编程抽象模型,这些操作都是在同一个流上进行的。这个流就像是一个 CPU 线程,在内存复制之后接受内核启动的提交。因此, CUDA 驱动程序为了实现内部 GPU 同步必须将命令插入主机接口的命令流中。

在这里插入图片描述

如上图所示，主机接口在协调流同步需求上起着核心作用。例如,在完成所需的内存复制之前,内核是启动不了的。此时,DMA单元可以停止给既定引擎传递命令,直到同步位置获得一个特定值为止。此操作与 CPU / GPU 同步类似,但 GPU 是对它内部进行同步的。

此硬件机制之上的软件抽象模型是一个 CUDA 流。它在该操作中就像 CPU 线程一样都是串行排队,为了并发执行需要多个流。由于引擎间共享命令缓冲区,应用程同的流里以软件的方式流水线化它们的请求。

GPU间同步

由于图2-26-图2-28中的同步位置都是在主机内存上,所以它们可以被系统中的任何一个 GPU 访问。其结果是,在CUDA4.0中,英伟达能够在cudaStreamWaitEvent()和 cuStreamWaiEvent() 函数的形式中添加 GPU 之间的同步。这些 API调用导致驱动程序为主机接口将等待命令插人当前 GPu 的命令缓冲区中,使得 GPU 一直等待,直到事件的给定同步值被写入为止。从CUDA4.0开始,事件不一定会被等待中的同一个 GPU用信号唤醒。流原先只能在单个 GPU 的硬件单元之间同步执行,现在已提升到可以在 GPU 之间同步执行了。


http://chatgpt.dhexx.cn/article/Vyoc7pa1.shtml

相关文章

GPU硬件架构以及运行机制笔记

GPU硬件架构以及运行机制笔记

本文是对向往大神的文章的一个笔记。 想阅读原文章移步博客园搜索向往 原文章比较长&#xff0c;这是一个精简和自己的一些理解 这篇文章要带着下面的问题去看 1、GPU是如何与CPU协调工作的&#xff1f; 2、GPU也有缓存机制吗&#xff1f;有几层&#xff1f;它们的速度差异多…
阅读更多...
4. CUDA编程手册中文版---硬件实现

4. CUDA编程手册中文版---硬件实现

第四章 硬件实现 更多精彩内容&#xff0c;请扫描下方二维码或者访问https://developer.nvidia.com/zh-cn/developer-program 来加入NVIDIA开发者计划 NVIDIA GPU 架构围绕可扩展的多线程流式多处理器 (SM: Streaming Multiprocessors) 阵列构建。当主机 CPU 上的 CUDA 程序调…
阅读更多...
SMIT介绍

SMIT介绍

System Management Interface Tool(系统管理界面工具) 软件安装与维护&#xff08;Sofeware Installation and Maintenance&#xff09; 软件许可管理&#xff08;Sofeware License Management&#xff09; 版本管理&#xff08;Manage Editions&#xff09; 设备管理&#…
阅读更多...
GPU硬件结构和编程模型(源于nvidia的CUDA文档)

GPU硬件结构和编程模型(源于nvidia的CUDA文档)

GPU的硬件结构 GPU通过一个可扩展的多线程流式多处理器(SMs)构建。一个multiprocessor可以在同一时间处理上百个线程。为了管理这些线程&#xff0c;使用一个特殊的结构SIMT。利用单线程中指令级的并行&#xff0c;以及同步硬件多线程实现的广泛线程级并行性。 SIMT Architec…
阅读更多...
实例分割文献阅读笔记（一）SimT

实例分割文献阅读笔记（一）SimT

阅读 SimT: Handling Open-set Noise for Domain Adaptive Semantic Segmentation 原作者知乎文章链接&#xff1a;知乎文章 GitHub链接&#xff1a;开源数据 SimT (CVPR22)&#xff1a;为了解决域自适应&#xff08;包含UDA和SFDA&#xff09;任务中目标域数据伪标签中存在…
阅读更多...
第三章 SIMT 内核：指令和寄存器数据流

第三章 SIMT 内核：指令和寄存器数据流

在本章和下一章中&#xff0c;我们将研究现代 GPU 的架构和微架构。 我们将对 GPU 架构的讨论分为两部分&#xff1a;(1) 在本章中研究实现计算部分的 SIMT 内核&#xff0c;然后 (2) 在下一章中研究内存系统。 在其传统的图形渲染角色中&#xff0c;GPU 访问数据集&#xff0…
阅读更多...
从GPU编程到SIMT核心

从GPU编程到SIMT核心

本文转自&#xff1a;从GPU编程到SIMT核心 - 知乎 (zhihu.com) 1、前言&本文重点 在 GPGPU 显得愈发重要的今天&#xff0c;仅凭 nVidia, AMD 提供的编程接口来了解 GPU 未免显得太单薄了些。时至今日&#xff0c; GPU 内部如何执行一条指令的对程序员来说依然是透明的、…
阅读更多...
并行计算范式-SIMD vs SIMT vs SMT: What’s the Difference Between Parallel Processing Models?

并行计算范式-SIMD vs SIMT vs SMT: What’s the Difference Between Parallel Processing Models?

Modern processor architectures utilize various execution models. Out of these, two are most popular: SIMD (Single Instruction Multiple Data) and SIMT (Single Instruction Multiple Threads). There’s also SMT (Simultaneous Multithreading), but that’s someth…
阅读更多...
SIMD＜SIMT＜SMT: NVIDIA GPU的并行机制

SIMD＜SIMT＜SMT: NVIDIA GPU的并行机制

原文出处&#xff1a; SIMD < SIMT < SMT: parallelism in NVIDIA GPUs 目录 1、概述 1.1、SIMD 2、SIMD vs SIMT 2.1 单指令、多套寄存器组 2.2 单指令、多个数据访问单元 2.3 单指令、多种运算逻辑路径 3、SIMD vs SIMT 3.1 GPU通过多thread来实现高thro…
阅读更多...
关于GPU一些笔记（SIMT方面)

关于GPU一些笔记（SIMT方面)

GPU组成 《计算机组成原理 — GPU 图形处理器》已经大概说明出GPU一般都是由比CPU多的core组成&#xff0c;而每个core 相当于一个单独线程进行计算&#xff0c;并且可以同时触发执行相同的单一指令但是每个计算单元数据不同(称之为SIMD)的指令执行。在英伟达GPU中 core一般称…
阅读更多...
如何理解GPU中的SIMT（单指令流多线程模型）

如何理解GPU中的SIMT（单指令流多线程模型）

随着设备尺寸逐渐变小&#xff0c;使得时钟频率很难有大的提升&#xff0c;人们开始寻找更有效的架构。为了提高能源效率&#xff0c;需要引入支持向量运算的硬件和减少数据的移动。 当下的架构通常是CPUGPU的&#xff0c;CPU在未来一段时间不会完全被GPU所取代&#xff0c;因…
阅读更多...
mysql怎么设置主键唯一性约束_MySQL主键约束和唯一性约束

mysql怎么设置主键唯一性约束_MySQL主键约束和唯一性约束

MySQL主键约束和唯一性约束都是索引&#xff0c;它们的区别是&#xff1a; 主键字段可以确保唯一性&#xff0c;但主键字段不能为NULL. 唯一性约束可以确保唯一性&#xff0c;但唯一性约束的字段可以为NULL 唯一性约束对含有NULL的记录不起作用&#xff0c;即可以重复加入含有N…
阅读更多...
mysql唯一性约束的作用_sql唯一约束有什么用

mysql唯一性约束的作用_sql唯一约束有什么用

SQL中唯一约束的作用是保证每个记录中都有一个唯一的标识&#xff0c;使得该列上没有相同的两个记录值&#xff1b;其中表的主键就是一个唯一性约束&#xff0c;不过主键只能有一个&#xff0c;所以如果其他列的数据不允许重复的话&#xff0c;就可以建立唯一性约束。 SQL中唯一…
阅读更多...
mysql多字段唯一约束_mysql多字段唯一约束

mysql多字段唯一约束_mysql多字段唯一约束

MySQL唯一约束(Unique Key)要求该列唯一，允许为空，但只能出现一个空值。唯一约束可以确保一列或者几列不出现重复值。 在创建表时设置唯一约束 在定义完列之后直接使用 UNIQUE 关键字指定唯一约束，语法规则如下： UNIQUE 创建数据表 tb_dept2，指定部门的名称唯一，输入的 S…
阅读更多...
mysql 修改唯一约束_mysql如何修改唯一性约束跟主键约束

mysql 修改唯一约束_mysql如何修改唯一性约束跟主键约束

一、如何添加主键约束和唯一性约束 1、添加主键约束 执行语法： alter table tableName add primarykey(column_name);#千万别忘了加()，同时要保证表结构中没有其他的主键，因为一个表中只能有一个主键。 2、添加唯一性约束 执行语法： alter table tableName addunique(colum…
阅读更多...
在mysql中怎么样添加唯一约束_mysql怎么添加唯一约束？

在mysql中怎么样添加唯一约束_mysql怎么添加唯一约束？

方法：1、创建表时，使用“CREATE TABLE 表名(字段名 数据类型 UNIQUE);”语句来添加；2、修改表时，使用“ALTER TABLE 表名 ADD CONSTRAINT 唯一约束名 UNIQUE(列名);”语句来添加。 (推荐教程：mysql视频教程) MySQL 唯一约束(Unique Key)是指所有记录中字段的值不能重复出现…
阅读更多...
添加唯一约束

添加唯一约束

– 1.添加唯一约束 – 方式1&#xff1a;创建表时指定 use mydb1; create table user1( id int, phone_number varchar(20)unique – 指定唯一约束 ); insert into user1 values(1001,‘123’); – insert into user1 values(1001,‘123’); --不是唯一会报错 – 在MySQL中&…
阅读更多...
SQL Server 2012 唯一约束（定义唯一约束、删除唯一约束）

SQL Server 2012 唯一约束（定义唯一约束、删除唯一约束）

文章目录 准备知识定义唯一约束使用SSMS工具定义唯一约束使用SQL方式定义唯一约束方式一&#xff1a;在创建数据表的时候定义唯一约束方式二&#xff1a;修改数据表定义唯一约束 删除唯一约束使用SSMS工具删除唯一约束方式一&#xff1a;在对象资源管理器中删除唯一约束方式二&…
阅读更多...
数据库----------唯一约束、默认约束、零填充约束

数据库----------唯一约束、默认约束、零填充约束

目录 1.唯一约束&#xff08;Unique&#xff09; 1.概念 2.语法 3.添加唯一约束 4.删除唯一约束 2.默认约束(default) 1.概念 2.语法 3.添加默认约束 4.删除默认约束 3.零填充约束&#xff08;zerofill&#xff09;了解即可 1.概念 2.操作 3.删除 1.唯一约束&…
阅读更多...
window连接远程桌面快捷键

window连接远程桌面快捷键

1、使用window R 打开运行界面&#xff0c;输入mstsc按回车 2、或者按window键打开开始界面&#xff0c;在搜索框输入mstsc按回车 3、效果
阅读更多...
推荐文章

Copyright @ 2022~2023