CUDA 5 ---- GPU架构(Fermi、Kepler)

GPU架构

SM(Streaming Multiprocessors)是GPU架构中非常重要的部分,GPU硬件的并行性就是由SM决定的。

以Fermi架构为例,其包含以下主要组成部分:

  • CUDA cores
  • Shared Memory/L1Cache
  • Register File
  • Load/Store Units
  • Special Function Units
  • Warp Scheduler

GPU中每个SM都设计成支持数以百计的线程并行执行,并且每个GPU都包含了很多的SM,所以GPU支持成百上千的线程并行执行,当一个kernel启动后,thread会被分配到这些SM中执行。大量的thread可能会被分配到不同的SM,但是同一个block中的thread必然在同一个SM中并行执行。

CUDA采用Single Instruction Multiple Thread(SIMT)的架构来管理和执行thread,这些thread以32个为单位组成一个单元,称作warps。warp中所有线程并行的执行相同的指令。每个thread拥有它自己的instruction address counter和状态寄存器,并且用该线程自己的数据执行指令。

SIMT和SIMD(Single Instruction, Multiple Data)类似。二者都通过将同样的指令广播给多个执行官单元来实现并行。一个主要的不同就是,SIMD要求所有的vector element在一个统一的同步组里同步的执行,而SIMT允许线程们在一个warp中独立的执行。SIMT有三个SIMD没有的主要特征:

  • 每个thread拥有自己的instruction address counter
  • 每个thread拥有自己的状态寄存器
  • 每个thread可以有自己独立的执行路径

一个block只会由一个SM调度,block一旦被分配好SM,该block就会一直驻留在该SM中,直到执行结束。一个SM可以同时拥有多个block。下图显示了软件硬件方面的术语:

需要注意的是,大部分thread只是逻辑上并行,并不是所有的thread可以在物理上同时执行。这就导致,同一个block中的线程可能会有不同步调。

并行thread之间的共享数据回导致竞态:多个线程请求同一个数据会导致未定义行为。CUDA提供了API来同步同一个block的thread以保证在进行下一步处理之前,所有thread都到达某个时间点。尽管如此,我们是没有什么原子操作来保证block内部的同步的。

同一个warp中的thread可以以任意顺序执行,active warps被SM资源限制。当一个warp空闲时,SM就可以调度驻留在该SM中另一个可用warp。在并发的warp之间切换是没什么消耗的,因为硬件资源早就被分配到所有thread和block,所以该新调度的warp的状态已经存储在SM中了。

SM可以看做GPU的心脏,寄存器和共享内存是SM的稀缺资源。CUDA将这些资源分配给所有驻留在SM中的thread。因此,这些有限的资源就使每个SM中active warps有非常严格的限制,也就限制了并行能力。所以,掌握部分硬件知识,有助于CUDA性能提升。

Fermi架构

Fermi是第一个完整的GPU计算架构。

  • 512个accelerator cores即所谓CUDA cores(包含ALU和FPU)
  • 16个SM,每个SM包含32个CUDA  core
  • 六个384位 GDDR5 DRAM,支持6GB global on-board memory
  • GigaThread engine(图左侧)将thread blocks分配给SM调度
  • 768KB L2 cache
  • 每个SM有16个load/store单元,允许每个clock cycle为16个thread(即所谓half-warp,不过现在不提这个东西了)计算源地址和目的地址
  • Special function units(SFU)用来执行sin cosine 等
  • 每个SM两个warp scheduler两个instruction dispatch unit,当一个block被分配到一个SM中后,所有该block中的thread会被分到不同的warp中。
  • Fermi(compute capability 2.x)每个SM同时可处理48个warp共计1536个thread。

每个SM由一下几部分组成:

  • 执行单元(CUDA cores)
  • 调度分配warp的单元
  • shared memory,register file,L1 cache

Kepler 架构

Kepler相较于Fermi更快,效率更高,性能更好。

  • 15个SM
  • 6个64位memory controller
  • 192个单精度CUDA cores,64个双精度单元,32个SFU,32个load/store单元(LD/ST)
  • 增加register file到64K
  • 每个Kepler的SM包含四个warp scheduler、八个instruction dispatchers,使得每个SM可以同时issue和执行四个warp。
  • Kepler K20X(compute capability 3.5)每个SM可以同时调度64个warp共计2048个thread。

Dynamic Parallelism

Dynamic Parallelism是Kepler的新特性,允许GPU动态的启动新的Grid。有了这个特性,任何kernel内都可以启动其它的kernel了。这样直接实现了kernel的递归以及解决了kernel之间数据的依赖问题。也许D3D中光的散射可以用这个实现。

Hyper-Q

Hyper-Q是Kepler的另一个新特性,增加了CPU和GPU之间硬件上的联系,使CPU可以在GPU上同时运行更多的任务。这样就可以增加GPU的利用率减少CPU的闲置时间。Fermi依赖一个单独的硬件上的工作队列来从CPU传递任务给GPU,这样在某个任务阻塞时,会导致之后的任务无法得到处理,Hyper-Q解决了这个问题。相应的,Kepler为GPU和CPU提供了32个工作队列。

不同arch的主要参数对比

时间: 2024-12-28 12:06:02

CUDA 5 ---- GPU架构(Fermi、Kepler)的相关文章

Pascal GPU 架构详解

写在前面:本文假定读者有一定 CUDA 基础.如果你对 GPU, sm_60/sm_61,CUDA 这些名词感到陌生,可以看我之前写的博客<CUDA 从入门到精通>.  1. 前言 Nvidia 在今年的 GTC( GPU Technology Conference ) 上高调宣布了 Pascal 架构--专门针对每瓦性能优化的新架构,采用 16nm 工艺.接着发布了该系列的扛鼎之作 P100 及其装进机箱后的产品 DGX-1: Tesla P100 采用顶级大核心 GP100 GP100 参

CUDA范例精解通用GPU架构-(2)其实写个矩阵相乘并不是那么难

程序代码及图解析: #include <iostream> #include "book.h" __global__ void add( int a, int b, int *c ) { *c = a + b; } int main( void ) { int c; int *dev_c; HANDLE_ERROR( cudaMalloc( (void**)&dev_c, sizeof(int) ) ); add<<<1,1>>>

ARM GPU 架构简介

1. 架构 2. 开发流程 3. Mali GPU Linux 内核设备驱动程序 Mali GPU DDK 的 Linux 版本包含在内核中运行的以下三个组件: 1)设备驱动程序: 它是最重要的组件,提供对 Mali-200 或 Mali-400 GPU 的低级访问.其主要功能如下: ?对 Mali GPU 硬件的访问 ?中断处理 ?低级内存管理 2) 统一内存提供程序 (UMP): 它是重要的辅助组件为,可以各种不同方式使用以实现驱动程序堆栈中的零拷贝操作. ?通过安全 ID 访问分配的 UM

vs2015+caffe+python3.5编译环境的搭建

修改build_win.cmd如下: @echo off @setlocal EnableDelayedExpansion :: Default values if DEFINED APPVEYOR ( echo Setting Appveyor defaults if NOT DEFINED MSVC_VERSION set MSVC_VERSION=14 if NOT DEFINED WITH_NINJA set WITH_NINJA=1 if NOT DEFINED CPU_ONLY se

[论文笔记] CUDA Cuts: Fast Graph Cuts on the GPU

Paper:V. Vineet, P. J. Narayanan. CUDA cuts: Fast graph cuts on the GPU. In Proc. CVPR Workshop, 2008. 原文出处: http://lincccc.blogspot.tw/2011/03/cuda-cuts-fast-graph-cuts-on-gpu_03.html 问题概述:Graph cut是一种十分有用和流行的能量优化算法,在计算机视觉领域普遍应用于前背景分割(Image segmenta

四 GPU 并行编程的存储系统架构

前言 在用 CUDA 对 GPU 进行并行编程的过程中,除了需要对线程架构要有深刻的认识外,也需要对存储系统架构有深入的了解. 这两个部分是 GPU 编程中最为基础,也是最为重要的部分,需要花时间去理解吸收,加深内功. 了解 GPU 存储系统架构的意义 CUDA 编程架构的设计思路本身也就是让程序员去使用缓存,而不是让缓存像 CPU 编程结构那样对程序员透明. 通过对所使用存储结构的优化,能够让程序的并行后的效果得到很大提高. 因此,这个问题是需要我们在开发全程中考虑的. 第一层:寄存器 每个流

Ubuntu12.04 之 CUDA 编程 (二) ~~~ GPU 程序加速

关于 Ubuntu12.04 下 CUDA5.5 的安装请参看如下链接 Ubuntu-12.04 安装 CUDA-5.5 关于 Ubuntu12.04 下 CUDA5.5 程序的运行请参看如下链接 Ubuntu12.04 之 CUDA 编程 (一) --- GPU 运行程序 1.程序的并行化 前一篇文章讲到了如何利用 CUDA5.5 在 GPU 中运行一个程序.通过程序的运行,我们看到了 GPU 确实可以作为一个运算器,但是,我们在前面的例子中并没有正真的发挥 GPU 并行处理程序的能力,也就是

win10+anaconda+cuda配置dlib,使用GPU对dlib的深度学习算法进行加速(以人脸检测为例)

在计算机视觉和机器学习方向有一个特别好用但是比较低调的库,也就是dlib,与opencv相比其包含了很多最新的算法,尤其是深度学习方面的,因此很有必要学习一下.恰好最近换了一台笔记本,内含一块GTX1060的显卡,可以用来更快地跑深度学习算法.以前用公司HP的工作站配置过dlib,GPU是Quadro K420,用dlib自带的人脸识别算法(ResNet)测试过,相比较1060的速度确实要快上很多.dlib.cuda和cudnn的版本经常会更新,每次重新配置环境会遇到一些问题,在这里记下来吧.

(转)移动GPU三种主流架构优缺点浅析

导读: GPU是Graphic Processor Unit的简称,顾名思义就是图形处理器. GPU的概念最早是从图形工作站发展而来,从90年代的个人电脑普及开始,GPU迎来了其大发展的时代. 在90年代中期,桌面GPU经历了2D到3D的跨越,从此3D图形渲染取代2D成为PC游戏的主流 1. 移动GPU与桌面GPU移动GPU相对桌面GPU只能算是小弟弟. 移动GPU的劣势主要表现在理论性能和带宽. 移动GPU受限于芯片的面积,能耗以及成本所以必须牺牲部分性能和带宽来求得性价比和电池续航力的平衡.