4

文章目录

4.4 图形处理器

这儿没写

4.4.1 GPU编程

  • CPU程序员的挑战不只是在GPU上获得出色的性能,还要协调系统处理器与GPU上的计算调度、系统存储器与GPU存储器之间的数据传输。
  • 此外,在本节后面将看到,
    • GPU几乎拥有所有可以由编程环境捕获的并行类型:
      • 多线程、MIMD、SIMD,甚至还有指令级并行。

NVIDIA开发与C类似的语言和编程环境,通过克服异质计算及多种并行带来的双重挑战来提高GPU程序员的效率

  • 系统名称CUDA,“计算统一设备体系结构”( Compute Unified Device Architecture)
  • CUDA为系统处理器(主机)生成C/C++,为GPU(设备,也就是CUDA中的D)生成C和C++方言
  • 一种类似的编程语言是OPENGL,几家公司共同开发,为多种平台提供一种与供应商无关的语言

NVIDIA认为,所有这些并行形式的统一主题就是CUDA线程

  • 以这种最低级别的并行作为编程原型,编译器和硬件可以将数千CUDA线程聚合在一起,利用CPU中的各种并行类型:
    • 多线程、MIMD、SIMD和指令级并行。
  • 因此, NVIDIA将CUDA编程模型定义为“单指令多线程”(SIMT)。
  • 这些线程进行了分块,执行时32个线程为一组,称线程块,马上会明白其原因
  • 将执行整个线程块的硬件称为多线程SIMD处理器

几个细节

  • 区分GPU(设备)的功能与系统处理器(主机)的功能
    • CUDA用__device__或__global__,用__host__后者
  • 被声明为__device__或 __global__functions的CUDA变量被分配给GPU存储
    • 供所有多线程SIMD处理器访问
  • 对于在GPU上运行的函数name进行扩展函数调用的语法为:
    • name<<<dimGrid, dimblock>>>(…parameter 1ist …)
    • dimGrid和dimBlock规定代码的大小(块)和块的大小(线程)
  • 块识别符(blockIdx)和每个块的线程识别符( threadIdx)外
    • CUDA还为每个块的线程数提供了一个关键字(blockDim),它来自上一个细节中dimBlock

DAXPY的C代码

4

CUDA版本

  • 在一个多线程SIMD处理器中启动n个线程,每个向量元素一个线程,
    • 每个线程块256个CUDA线程
  • GPU功能首先根据块ID、每个块的线程数以及线程ID来计算相应的元素索引 i。
  • 只要索引没超出数组范围
    • 就执行乘和加
      4

对比C码和CUDA,可看出一种用于实现数据并行CUDA代码并行化的共同模式

  • C中有一循环的所有迭代都
    • 与其他迭代相独立,
    • 可将此循环转为并行代码,
    • 其中每个循环迭代都是个独立线程
  • 曾提到,向量化编译器也要求循环的迭代间无相关性,这种相关被称为循环间相关,4.5节介绍
  • 程序员通过明确指定
    • 网格大小及每个SIMD处理器中的线程数
    • dimGrid和线程块吗？
    • 明确指出CUDA中的并行
  • 由于为每个元素都分配了一个线程
    • so在向存储器中写结果时不要在线程间实行同步

行执行和线程管理由GPU硬件负责,而不是由应用程序或OS完成

  • 为简化硬件处理的排程,CUDA要求线程块能够按任意顺序独立执行
  • 尽管不同的线程块可用
    • 全局存储器中的原子存储器操作进行协调
    • 但它们之间不能直接通信

许多GPU硬件概念在CUDA中不是非常明显。

  • 从程序员生产效率角度来看,是好事,但大多数程序员使用GPU而不是CPU来提高性能。
  • 重视性能的程序员在用CUDA编写程序时须时刻惦记着GPU硬件。
  • 他们知道需要将控制流中的32个线程分为一组,
    • 以从多线程SIMD处理器中获得最佳性能,
    • 并在每个多线程SIMD处理器中另外创建许多线程,
    • 以隐藏访问DRAM的延迟,稍后将解释其原因。
  • 它们还需要将数据地址保持在一个或一些存储器块的局部范围内,
    • 以获得所期望的存储器性能。

和许多并行系统一样,CUDA在生产效率和性能间进行折中

  • 提供一些本身固有的功能,让程序员能显式控制硬件
  • 一方面是生产效率,另一方面是使程序员能够表达硬件所能完成的所有操作,在并行计算中,这两方面经常会发生竞争
  • 了解编程语言在这著名的生产效率与性能大战中如何发展,
    • 了解CUDA是否能够在其他GPU或者其他类型的体系结构中变得普及
    • 都将是非常有意义的

4.4.2 NVIDIA GPU计算结构

  • 上文提到的罕见传统可帮助解释为什么GPU有自己的体系结构类型,为什么有与CPU独立的专门术语
  • 理解GPU的一个障碍就是术语,有些词汇的名称甚至可能导致误解
  • 克服这一障碍的难度很大,这章经过多次重写就是例证
  • 为让读者既能理解GPU的体系结构,又能学习许多采用非传统定义的GPU术语,最终的解决方案是用CUDA术语来描述软件,
    • 而在开始时使用更具描述性的术语来介绍硬件,有时还会借用 OpenCL用的术语
  • 在用我们的术语解释GPU体系结构之后,再将它们对应到 NVIDIA GPU的官方术语

表4-5左至右列出

  • 本节使用的一些更具描述性的术语、主流计算中的最接近术语、官方 NVIDIA GPU术语,以及这些术语的简短描述。
  • 本节的后续部分将使用该表左侧描述性术语来解释GPU的徽体系结构特征。

4

4

以 NVIDIA系统为例,它们是GPU体系结构的代表。

  • 具体,将用上面CUDA并行编程语言的术语,以Fermi体系结构为例(见4.7节)

和向量体系结构一样,GPU只能很好地解决数据级并行

  • 这两种类型都拥有集中-分散数据传送和遮罩寄存器,GPU处理器的寄存器比向量处理器更多
  • 由于它们没有一种接近的标量处理器,所以GPU有时会在运行时以硬件实现一些功能
    • 而向量计算机通常是在编译时用软件来实现这些功能的
  • 与大多数向量体系结构不同,GPU还依靠单个多线程SIMD处理器中的多线程来隐藏存储器延迟(见2和3章)。
  • 想为向量体系结构和GPU编写出高效代码,还需要考虑SIMD操作分组。

网格是在GPU上运行、由一组线程块构成的代码。

  • 表4-5：网格与向量化循环、线程块与循环体(已经进行了条带挖掘,所以它是完整的计算循环)之间的相似
  • 希望把两个向量乘在一起,每个向量8192
  • 图4-8给出了这个示例与前两个GPU术语之间的关系
  • 执行所有8192个元素乘法的GPU代码被称为网格(或向量化循环)
  • 为了将它分解为更便于管理的大小,网格由线程块(或向量化循环体)组成,每个线程块最多512个元素
  • 注意,一条SIMD指令次执行32个元素
  • 由于向量中有8192个元素,所以这个示例中有16个线程块(16=8192÷512)。
  • 网络和线程块是在GPU硬件中实现的编程抽象,帮助程序员组织自己的CUDA代码。
    • 线程块类似于一个向量长度为32的条带挖掘向量循环。

线程块调度程序将线程块指定给执行该代码的处理器,这种处理器为多线程SIMD处理器

  • 线程块调度程序与向量体系结构中的控制处理器有相似
  • 它决定了该循环所需要的线程块数,在完成循环之前,一直将它们分配给不同的多线程SIMD处理器
  • 示例将16个线程块发送给多线程SDMD处理器,计算这个循环的所有8192个元素。

图4-8

4

  • 网格(可向量化循环)、线程块(SIMD基本块)和SIMD指令线程
    • 与
    • 向量——向量乘法的对应
    • 向量长度为8192
  • 每个SIMD指令线程的每条指令计算32个元素
    • 示例中
    • 每个线程块包含16个SIMD指令线程
    • 网格包含16个线程块
  • 硬件线程块调度程序将线程块指定给多线程SIMD处理器
    • 硬件线程调度程序选择某个SIMD指令线程来运行一个SIMD处理器中的每个时钟周期
  • 只有同一线程块中的SIMD线程可以通过本地存储器通信
  • (Tesla代的GPU,每个线程块可同时执行的最大SIMD线程数为16,后来 Fermi一代的GPU为32)

图4-9显示多线程SIMD处理器简化框图

  • 与向量处理器类似,但它有许多并行功能单元都是深度流水化的,而不像向量处理器一样只有小部分如此。
  • 图4-8中的示例中,向每个多线程SIMD处理器分配这些向量的512个元素进行处理
  • SIMD处理器都有独立PC的完整处理器,用线程进行编程(见3章)

4

GPU硬件包含一组用来执行线程块网络(向量化循环体)的多线程SIMD处理器

  • GPU是一个由多线程SIMD处理器组成的多处理器

Fermi体系结构的前四种实现拥有7、11、14或15个多线程SMD处理器;

  • 未来的版本可能仅有2个或4个。
  • 为了在拥有不同个多线程SIMD处理器的GPU型号之间实现透明的可伸缩功能,
  • 线程块调度程序将线程块(向量化循环主体)指定给多线程SIMD处理器
  • 图4-10给出了Fermi体系结构的GTX480实现的平面图

4

硬件创建、管理、调度和执行的机器对象是SIMD指令线程

  • 它是个包含专用SIMD指令的传统线程
  • 这些SIMD指令线程有其自己的PC,运行在多线程SIMD处理器上
  • SIMD线程调度程序包括一个记分板,哪些SIMD指令线程已经做好运行准备,然后将它们发送给分发单元,以在多线程SIMD处理器上运行。
  • 与传统多线程处理器中的硬件线程调度程序相同(见3章),就是对SIMD指令线程进行调度
  • GPU硬件有两硬件调度程序
    • 线程块调度程序,将线程块(向量化循环体)分配给多线程SDMD处理器,确保线程块被分配给其局部存储器拥有相应数据的处理器,
    • SIMD处理器内部的SIMD线程调度程序,由它来调度应当何时运行SIMD指令线程。

这些线程的SIMD指令的宽度为32,所以这个示例中每个SIMD指令线程将执行32个元素运算。

  • 本例中,线程块将包含512÷32=16SIMD线程(图4-8)。

由于线程由SIMD指令组成,所以SIMD处理器必须拥有并行功能单元来执行运算。

  • 我们称之为SIMD车道,它们与4.2节的向量车道很类似

每个SIMD处理器中车道数在各代GPU中不同

  • 对Fermi,宽度为32的SIMD指令线程被映射到16个物理SIMD车道
    • SIMD指令线程中的每条SIMD指令要两时钟周期完成
  • 每个SIMD指令线程在锁定步骤执行,仅在开始时进行调度。
  • 将SIMD处理器类比为向量处理器,可说它有16个车道,向量长度为32,钟鸣为2个时钟周期。
  • 用术语“SIMD处理器”,而不是“向量处理器”,就是因为这种既宽且浅的本质,前者的描述性更强一些

根据定义,SIMD指令的线程是独立的,

  • SIMD线程调度程序可选择任何已准备就绪的SIMD指令线程,
    • 不需町着线程序列中的下一条SIMD指令。
  • SIMD线程调度程序包括一个记分板(见3章),
    • 用于跟踪多达48个SIMD线程,
    • 以了解哪个SIMD指令已做好运行准备。
  • 之所以需要这个记分板,是因为存储器访问指令占用的时钟周期数可能无法预测,
    • 比如存储器组的冲突就可能导致这一现象。
  • 图4-11给出的SIMD线程调度程序在不同时间以不同顺序选取SIMD指令线程。
  • 假定GPU应用程序拥有很多的SIMD指令线程,
    • 因此,实施多线程既可隐藏到DRAM的延迟,
    • 又可提高多线程SIMD处理器的使用率。
  • 但为防止损失,最近的 NVIDIA Femi GPU包含了一个L2缓存(见4.7节)。

4

继续探讨向量乘法示例,

  • 每个多线程SIMD处理器须将两个向量的32个元素从存储器载入寄存器中,通过读、写寄存器来执行乘法,
    • 然后将乘积从寄存器存回存储器中。
  • 为保存这些存储器元素,SIMD处理器拥有32768个32位寄存器,给人以深刻印象。
  • 就像向量处理器样,从逻辑上在向量车道之间划分这些寄存器,这里自然是在SIMD车道之同划分。
  • 每个SIMD线程被限制为不超过64个寄存器,所以我们可以认为一个SIMD线程最多拥有64个向量寄存器,
    • 每个向量寄存器有32个元素,每个元素的宽度为32位。
    • (双精度浮点用两个相邻的32位寄存器,
    • 所以另一种意见是每个SIMD线程拥有32个各包括32个元素的向量寄存器,每个宽度为64位。)

由于Femi拥有16个物理SIMD车道,各包含2048个寄存器。

  • (GPU没有尝试根据位来设计硬件寄存器,使其拥有许多读取端口和写入端口,而是像向量处理器一样,使用较简单的存储器结构,但将它们划分为组,以获得足够的带宽。)
  • 每个CUDA线程获取每个向量寄存器中的一个元素。
  • 为了用16个SIMD车道处理每个SMD指令线程的32个元素,线程块的CUDA线程可以共同使用2048个寄存器的一半。

为了能够执行许多个SIMD指令线程,

  • 需要在创建SIMD指令线程时在每个SIMD处理器上动态分配一组物理寄存器,并在退出SIMD线程时加以释放。

注意,CUDA线程就是SIMD指令线程的垂直抽取,

  • 与ー个SMD车道上执行的元素相对应。
  • 要当心,CUDA线程与POSX线程完全不同;
    • 不能从CUDA线程进行任意系统调用。

现在可看看GPU指令是什么样。

参考链接

相关文章：

猜你喜欢
相关资源
相似解决方案
热门标签
Java Python linux javascript Mysql C# Docker 算法 前端 SpringBoot Redis Vue spring 设计模式 .net core .net kubernetes c++ 数据库 数据结构 大数据 js 机器学习 微服务 Android Go 程序员 面试 JVM ASP.net core 云原生 人工智能 后端 PHP git CSS golang k8s Nginx Django mybatis 深度学习 多线程 React 架构 devops 爬虫 云计算 Spring Boot LeetCode