CUDA C编程入门-硬件实现

  NVIDIA GPU架构是建立在一个可扩展的多线程流处理器(SMs)阵列之上的。当一个在主机CPU的CUDA程序调用一个核函数时,网格的blocks分配给运行容量空闲的多核处理器计算。线程块的线程同时在多核处理器上执行,多个线程块能同时在一个多核处理器上执行。当线程块结束时,新的线程块会加载到空出的多核处理器上。

  一个多核处理器被设计成能同时执行数以百计的线程。为了管理这样大容量的线程,GPU使用一个在SIMT架构描述的叫SIMT(单个指令,多个线程)的独特架构。在单个线程中,指令是流水的去利用指令级别的并行,和通过同时发生的硬件多线程(细节在硬件多线程)的线程级别并行一样。不像CPU核心,指令被按序发射,没有分支预测和推测执行。

  SIMT架构和硬件多线程描述流处理器是对所有设备的共同的架构特性。计算能力为1.x、2.x、3.x和5.0章节分别地提供计算能力为1.x、2.x、3.x和5.0的细节。

  NVIDIA GPU架构使用小端表示。

4.1 SIMT架构

  多核处理器以叫做warps的32个并行线程为一组来创建、管理、安排和执行线程。组成warps的单个线程在程序的同一的程序地址开始,但是,都拥有自己的指令地址计数器和寄存器状态,因此,可以自由地分支和独立地执行。术语warp源于weaving,第一个并行的技术。一个half-warp不是warp的前半个就是后半个。一个四分之一的warp不是第一、第二、第三和第四的四分之一warp。

  当给定一个多核处理器一个或者多个线程块去执行,会把线程块划分成warps,每个warp通过一个warp调度程序安排执行。一个线程块划分成warp的方式总是相同的,每个warp包含连续的线程,递增的线程ID,第一个warp含有ID为0的线程。线程层次这章节描述在线程块的线程ID怎样与线程索引相关。

  一个warp在一个时刻执行一个指令,当一个warp的32个线程执行的路径一致时,此时实现全部的效率。如果warp的线程分出依赖数据的条件分支,warp执行这个分支路径,线程不会在其它分支执行,当所有的分支完成后,线程集中回到相同的执行路径。分支只会在一个warp内出现。不同的warp独立地执行,不管它们是执行在相同的或者不同的代码路径上。

  SIMT架构类似于SIMD向量组织,单个指令控制多个处理中的元素。SIMD向量组织的一个关键的区别是会提供SIMD宽度给软件。而SIMT指令指定单个线程的执行和分支行为。与SIMD相比,SIMT是程序员能写出独立的具体数量的线程的线程级别的并行代码和线程相应的数据并行代码。为了正确的目的,程序员能基本上忽视SIMT的行为,然而,注意在warp的线程很少有分支的代码会提升实际执行的性能。实际上,这相当于传统代码中的cache lines的角色:为了正确性,cache line大小能安全的忽视,而为了最佳性能,必须在代码结构中考虑到。向量组织,另一方面,要求软件合并加载进向量,并手工地管理分支。

  注意:

  warp的线程,在warp当前执行路径的线程叫活跃线程,而不在当前执行路径的叫不活跃线程。线程能不活跃是因为warp中的其他的线程早就退出,或者因为它们在与warp当前执行的路径的不同的执行路径,或者是因为它们是线程数不是warp大小的倍数的线程块的最后的线程。

  如果一个被warp中多个线程执行写数据到全局或者共享内存的同一位置单元的非原子指令,系列化的写的引起该单元变化的数目依赖于设备的计算能力,且那个线程执行最后的写是未定义的。

  如果一个被warp中多个线程执行读、修改和写数据到全局内存的同一单元的原子指令,每个读/修改、写那个单元,它们所有都被系列化,但它们的执行顺序是未定义的。

4.2 硬件多线程

  在warp整个生命周期里,每个被多核处理器执行的warp的上下文(程序计数器、寄存器等)依靠在芯片上。因此,从一个执行的上下文转换到另一个没有代价,在整个指令发射时间,warp调度程序选择一个有线程准备执行下一个指令的warp,发射这些线程的指令。

  尤其地,每个处理器有一系列被warp划分的32位寄存器,一个被线程块划分的并行的数据缓存或者共享内存。

  对一个给定的核函数,可以在处理器上驻留和处理的线程块和warp的数量取决于在处理器可用的寄存器的数量和被核函数使用的共享内存。这里也有在一个处理器上驻留的线程块的最大数量和驻留的warps最大数量。这些同样限制于寄存器的数量和处理器上可用的共享内存,是附录F给定的设备计算能力的函数。如果在每个处理器上没有足够的可用的寄存器或者共享内存去处理至少一个线程块,核函数加载将失败。

  一个线程块的warp的数量如下:

  • ceil(T/Wsize,1)
  • T是每个线程块线程的数量。
  • Wsize是warp大小,等于32。
  • ceil(x,y)等于x去最靠近y倍数的整数。

  寄存器的数量和一个线程块可分配的共享内存的总数是在CUDA软件开发包的CUDA占用计算器记录。

CUDA C编程入门-硬件实现

时间: 2024-10-12 08:10:22

CUDA C编程入门-硬件实现的相关文章

CUDA C编程入门-介绍

CUDA C编程入门-介绍 1.1.从图形处理到通用并行计算 在实时.高清3D图形的巨大市场需求的驱动下,可编程的图形处理单元或者GPU发展成拥有巨大计算能力的和非常高的内存带宽的高度并行的.多线程的.多核处理器.如图1和图2所示. 图 1 CPU和GPU每秒的浮点计算次数 图 2 CPU和GPU的内存带宽 在CPU和GPU之间在浮点计算能力上的差异的原因是GPU专做密集型计算和高度并行计算-恰好是图形渲染做的-因此设计成这样,更多的晶体管用于数据处理而不是数据缓存和流控制,如图3所示. 图 3

CUDA C编程入门

CUDA C编程入门系列文章是CUDA C Programming guide的翻译,同时也会加入一些个人的理解.由于刚刚接触CUDA编程,对此领域不是很熟悉,翻译的质量和正确与否不能保证,不过如果读者发现哪里有误,欢迎指正. CUDA C编程入门文档结构: 第一章:介绍-关于CUDA的总体介绍. 第二章:编程模型-概括CUDA编程模型. 第三章:编程接口-描述CUDA编程的C语言的接口. 第四章:硬件实现-描述GPU的硬件实现. 第五章:编程指引-给出一些指导,怎样才能发挥CPU最大的性能.

CUDA C编程入门-编程接口(3.2)CUDA C运行时

在cudart库里实现了CUDA C运行时,应用可以链接静态库cudart.lib或者libcudart.a,动态库cudart.dll或者libcudart.so.动态链接cudart.dll或者libcudart.so的应用需要把CUDA的动态链接库(cudart.dll或者libcudart.so)包含到应用的安装包里. CUDA所有的运行时函数都以cuda为前缀. 在异构编程这一章节提到,CUDA编程模型假设系统是由一个自带各自的内存的主机和设备组成.设备内存这一小节概述用于管理设备内存

CUDA C编程入门-编程模型

这章节介绍CUDA编程模型的主要的概念. 2.1.kernels(核函数) CUDA C扩展了C语言,允许程序员定义C函数,称为kernels(核函数).并行地在N个CUDA线程中执行N次. 使用__global__说明符声明一个核函数,调用使用<<<...>>>,并且指定执行的CUDA线程数目.执行的每个线程都有一个独一的ID,在核函数中可以通过变量threadIdx获取. 例子,两个向量的加,A加B,并把结果存入C,A.B和C的长度为N. __global__ vo

CUDA C编程入门-编程接口

CUDA C给熟悉C编程语言的人提供一个简单的途径去编写在设备(GPU)上执行的代码. 由一个最小的C语言的扩展集和运行时库组成. 核心的语言扩展在编程模型这一章节已经介绍过了.允许程序员定义核函数并且使用一些新的语法指定核函数每次运行时的grid和block的维数.可以在C语言扩展这个章节里找到扩展的完整描述.所有的含有这些扩展的源代码都需要使用nvcc编译,nvcc的概述可以查看使用nvcc编译这一小节. 在CUDA C运行这一小节介绍运行时.运行时提供在主机执行的用于分配和回收设备内存.设

CUDA C编程入门-编程接口(3.5)模式转换

GPUs有一个显示输出,输出到一个叫主表面的DRAM内存中,这个表面被用于刷新输出给用户看的显示设备.当用户通过改变显示器的分辨率或者位深(使用NVIDIA控制面版或者Windows显示器控制面版),开始一个显示器模式选择时,主表面需要的内存数量会改变.例如,如果以后改变显示器分辨率从1280x1024x32位到1600x1200x32位,系统必须增加7.68MB而不是5.24MB的内存给主表面(带反走样的全屏的图形应用可能要求更多的显示内存给主表面).在Windows上,其它事件,包括加载一个

CUDA C编程入门-编程接口(3.4)计算模型

在Windows Server 2008和后来的版本或者Linux上运行的Tesla解决方案,能通过NVIDIA系统管理界面(nvidia-smi,作为发行的驱动的一个工具)设置系统中的任何设备运行于下面三种计算模式中的一种 : 默认计算模式:多主机线程能同时使用设备(使用运行时API,在这个设备上调用cudaSetDevice(),或者使用驱动API,使当前的上下文与设备关联). 处理器独占的计算模式:仅创建一个跨系统所有处理器CUDA上下文,同时对于一个处理器的尽可能多的线程. 处理器和线程

CUDA C编程入门-编程接口(3.3)版本和兼容性

有两个版本号是开发者开发CUDA应用时需要关心的:计算能力-描述产品规格和计算设备的特性和CUDA驱动API的版本-描述驱动API和运行时支持的特性. 在驱动头文件的宏CUDA_VERSION可以获取驱动API的版本.可以允许开发者检查他们的应用是否要求一个比现有版本更新的版本.这很重要,因为驱动API是向后兼容的,意味着针对特别版本编译的应用.插件和库(包括C运行时)能继续在后来发行版本的设备驱动上运行,如图11所示.驱动不是向前兼容的,意味着针对特别版本编译的应用.插件和库(包括C运行时)不

OpenGL ES编程入门资源集合

OpenGL ES 2.0中文手册:http://www.dreamingwish.com/articlelist/category/opengl-es-api 里边讲解了部分API的功能,作为基本的参考. OpenGL ES2.0 渲染管线:http://codingnow.cn/opengles/1504.html OpenGL ES2.0 绘制三角形:http://codingnow.cn/opengles/1514.html 一个OpenGL ES2.0的例子:http://blog.c