cuda_cuda tensorflow - CSDN
cuda 订阅
CUDA(Compute Unified Device Architecture),是显卡厂商NVIDIA推出的运算平台。 CUDA™是一种由NVIDIA推出的通用并行计算架构,该架构使GPU能够解决复杂的计算问题。 它包含了CUDA指令集架构(ISA)以及GPU内部的并行计算引擎。 开发人员可以使用C语言来为CUDA™架构编写程序,C语言是应用最广泛的一种高级编程语言。所编写出的程序可以在支持CUDA™的处理器上以超高性能运行。CUDA3.0已经开始支持C++和FORTRAN。 展开全文
CUDA(Compute Unified Device Architecture),是显卡厂商NVIDIA推出的运算平台。 CUDA™是一种由NVIDIA推出的通用并行计算架构,该架构使GPU能够解决复杂的计算问题。 它包含了CUDA指令集架构(ISA)以及GPU内部的并行计算引擎。 开发人员可以使用C语言来为CUDA™架构编写程序,C语言是应用最广泛的一种高级编程语言。所编写出的程序可以在支持CUDA™的处理器上以超高性能运行。CUDA3.0已经开始支持C++和FORTRAN。
信息
外文名
CUDA
类    别
运算平台
适用领域
计算机
中文名
统一计算设备架构
组    成
ISA、GPU
推出者
NVIDIA
CUDA应用范围
计算行业正在从只使用CPU的“中央处理”向CPU与GPU并用的“协同处理”发展。为打造这一全新的计算典范,NVIDIA™(英伟达™)发明了CUDA(Compute Unified Device Architecture,统一计算设备架构)这一编程模型,是想在应用程序中充分利用CPU和GPU各自的优点。该架构已应用于GeForce™(精视™)、ION™(翼扬™)、Quadro以及Tesla GPU(图形处理器)上,对应用程序开发人员来说,这是一个巨大的市场。在消费级市场上,几乎每一款重要的消费级视频应用程序都已经使用CUDA加速或很快将会利用CUDA来加速,其中不乏Elemental Technologies公司、MotionDSP公司以及LoiLo公司的产品。 在科研界,CUDA一直受到热捧。例如,CUDA现已能够对AMBER进行加速。AMBER是一款分子动力学模拟程序,全世界在学术界与制药企业中有超过60,000名研究人员使用该程序来加速新药的探索工作。在金融市场,Numerix以及CompatibL针对一款全新的对手风险应用程序发布了CUDA支持并取得了18倍速度提升。Numerix为近400家金融机构所广泛使用。CUDA的广泛应用造就了GPU计算专用Tesla GPU的崛起。全球财富五百强企业已经安装了700多个GPU集群,这些企业涉及各个领域,例如能源领域的斯伦贝谢与雪佛龙以及银行业的法国巴黎银行。随着微软Windows 7与苹果Snow Leopard操作系统的问世,GPU计算必将成为主流。在这些全新的操作系统中,GPU将不仅仅是图形处理器,它还将成为所有应用程序均可使用的通用并行处理器。
收起全文
  • CUDA编程方法

    2020-07-03 11:28:53
    学懂YOLOv3目标检测原理 读懂C语言实现的Darknet源码
  • CUDA——"从入门到放弃

    千次阅读 多人点赞 2020-04-18 11:11:35
    1. 知识准备 1.1 中央处理器(CPU) 中央处理器(CPU,Central Processing Unit)是一块超大规模的集成电路,是一台计算机的运算核心(Core)和控制核心( Control Unit)。它的功能主要是解释计算机指令以及...

     转载:https://www.jianshu.com/p/34a504af8d51

    1. 知识准备

    1.1 中央处理器(CPU)

    中央处理器(CPU,Central Processing Unit)是一块超大规模的集成电路,是一台计算机的运算核心(Core)和控制核心( Control Unit)。它的功能主要是解释计算机指令以及处理计算机软件中的数据。
    中央处理器主要包括运算器(算术逻辑运算单元,ALU,Arithmetic Logic Unit)和高速缓冲存储器(Cache)及实现它们之间联系的数据(Data)、控制及状态的总线(Bus)。它与内部存储器(Memory)和输入/输出(I/O)设备合称为电子计算机三大核心部件。

    CPU的结构主要包括运算器(ALU, Arithmetic and Logic Unit)、控制单元(CU, Control Unit)、寄存器(Register)、高速缓存器(Cache)和它们之间通讯的数据、控制及状态的总线

    简单来说就是:计算单元、控制单元和存储单元,架构如下图所示:

    CPU微架构示意图

    什么?架构记不住?来,我们换种表示方法:

    CPU微架构示意图(改)

    嗯,大概就是这个意思。

    从字面上我们也很好理解,计算单元主要执行算术运算、移位等操作以及地址运算和转换存储单元主要用于保存运算中产生的数据以及指令等控制单元则对指令译码,并且发出为完成每条指令所要执行的各个操作的控制信号

    所以一条指令在CPU中执行的过程是这样的:读取到指令后,通过指令总线送到控制器(黄色区域)中进行译码,并发出相应的操作控制信号;然后运算器(绿色区域)按照操作指令对数据进行计算,并通过数据总线将得到的数据存入数据缓存器(大块橙色区域)。过程如下图所示:

    CPU执行指令图

    是不是有点儿复杂?没关系,这张图完全不用记住,我们只需要知道,CPU遵循的是冯诺依曼架构,其核心就是:存储程序,顺序执行

    讲到这里,有没有看出问题,没错——在这个结构图中,负责计算的绿色区域占的面积似乎太小了,而橙色区域的缓存Cache和黄色区域的控制单元占据了大量空间。

    高中化学有句老生常谈的话叫:结构决定性质,放在这里也非常适用。

    因为CPU的架构中需要大量的空间去放置存储单元(橙色部分)和控制单元(黄色部分),相比之下计算单元(绿色部分)只占据了很小的一部分,所以它在大规模并行计算能力上极受限制,而更擅长于逻辑控制。

    另外,因为遵循冯诺依曼架构(存储程序,顺序执行),CPU就像是个一板一眼的管家,人们吩咐的事情它总是一步一步来做。但是随着人们对更大规模与更快处理速度的需求的增加,这位管家渐渐变得有些力不从心。

    于是,大家就想,能不能把多个处理器放在同一块芯片上,让它们一起来做事,这样效率不就提高了吗?

    没错,GPU便由此诞生了。


    1.2 显卡

    显卡(Video card,Graphics card)全称显示接口卡,又称显示适配器,是计算机最基本配置、最重要的配件之一。显卡作为电脑主机里的一个重要组成部分,是电脑进行数模信号转换的设备,承担输出显示图形的任务。显卡接在电脑主板上,它将电脑的数字信号转换成模拟信号让显示器显示出来,同时显卡还是有图像处理能力,可协助CPU工作,提高整体的运行速度。对于从事专业图形设计的人来说显卡非常重要。 民用和军用显卡图形芯片供应商主要包括AMD(超微半导体)Nvidia(英伟达)2家。现在的top500计算机,都包含显卡计算核心。在科学计算中,显卡被称为显示加速卡

    为什么GPU特别擅长处理图像数据呢?这是因为图像上的每一个像素点都有被处理的需要,而且每个像素点处理的过程和方式都十分相似,也就成了GPU的天然温床。
     

    GPU微架构示意图

    从架构图我们就能很明显的看出,GPU的构成相对简单,有数量众多的计算单元和超长的流水线,特别适合处理大量的类型统一的数据。

    再把CPU和GPU两者放在一张图上看下对比,就非常一目了然了。

    GPU的工作大部分都计算量大,但没什么技术含量,而且要重复很多很多次。

    但GPU无法单独工作,必须由CPU进行控制调用才能工作。CPU可单独作用,处理复杂的逻辑运算和不同的数据类型,但当需要大量的处理类型统一的数据时,则可调用GPU进行并行计算。

    借用知乎上某大佬的说法,就像你有个工作需要计算几亿次一百以内加减乘除一样,最好的办法就是雇上几十个小学生一起算,一人算一部分,反正这些计算也没什么技术含量,纯粹体力活而已;而CPU就像老教授,积分微分都会算,就是工资高,一个老教授资顶二十个小学生,你要是富士康你雇哪个?

    注:GPU中有很多的运算器ALU和很少的缓存cache,缓存的目的不是保存后面需要访问的数据的,这点和CPU不同,而是为线程thread提高服务的。如果有很多线程需要访问同一个相同的数据,缓存会合并这些访问,然后再去访问dram。

    可爱的你如果对CUDA硬件有更多的兴趣,可移步NVIDIA中文官网进一步学习。
     

    1.3 内存

    内存是计算机中重要的部件之一,它是与CPU进行沟通的桥梁。计算机中所有程序的运行都是在内存中进行的,因此内存的性能对计算机的影响非常大。内存(Memory)也被称为内存储器,其作用是用于暂时存放CPU中的运算数据,以及与硬盘外部存储器交换的数据。只要计算机在运行中,CPU就会把需要运算的数据调到内存中进行运算,当运算完成后CPU再将结果传送出来,内存的运行也决定了计算机的稳定运行。 内存是由内存芯片、电路板、金手指等部分组成的。

    1.4 显存

    显存,也被叫做帧缓存,它的作用是用来存储显卡芯片处理过或者即将提取的渲染数据。如同计算机的内存一样,显存是用来存储要处理的图形信息的部件。

    1.5 显卡、显卡驱动、CUDA之间的关系

    显卡:(GPU)主流是NVIDIA的GPU,深度学习本身需要大量计算。GPU的并行计算能力,在过去几年里恰当地满足了深度学习的需求。AMD的GPU基本没有什么支持,可以不用考虑。

    驱动:没有显卡驱动,就不能识别GPU硬件,不能调用其计算资源。但是呢,NVIDIA在Linux上的驱动安装特别麻烦,尤其对于新手简直就是噩梦。得屏蔽第三方显卡驱动。下面会给出教程。

    CUDA:是NVIDIA推出的只能用于自家GPU的并行计算框架。只有安装这个框架才能够进行复杂的并行计算。主流的深度学习框架也都是基于CUDA进行GPU并行加速的,几乎无一例外。还有一个叫做cudnn,是针对深度卷积神经网络的加速库。

    查看显卡驱动信息(以实验室服务器为例)

    ssh ubuntu@192.168.1.158
    

    输入服务器密码登陆
    然后,进入cuda

    cd /usr/local/cuda-8.0/samples/1_Utilities/deviceQuery
    

    运行其中的可执行文件

    ./deviceQuery
    

    得到如下信息

    ./deviceQuery Starting...
    
     CUDA Device Query (Runtime API) version (CUDART static linking)
    
    Detected 4 CUDA Capable device(s)
    
    Device 0: "GeForce GTX 1080 Ti"
      CUDA Driver Version / Runtime Version          9.0 / 8.0
      CUDA Capability Major/Minor version number:    6.1
      Total amount of global memory:                 11171 MBytes (11713708032 bytes)
      (28) Multiprocessors, (128) CUDA Cores/MP:     3584 CUDA Cores
      GPU Max Clock rate:                            1620 MHz (1.62 GHz)
      Memory Clock rate:                             5505 Mhz
      Memory Bus Width:                              352-bit
      L2 Cache Size:                                 2883584 bytes
      Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
      Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
      Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 2048 layers
      Total amount of constant memory:               65536 bytes
      Total amount of shared memory per block:       49152 bytes
      Total number of registers available per block: 65536
      Warp size:                                     32
      Maximum number of threads per multiprocessor:  2048
      Maximum number of threads per block:           1024
      Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
      Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
      Maximum memory pitch:                          2147483647 bytes
      Texture alignment:                             512 bytes
      Concurrent copy and kernel execution:          Yes with 2 copy engine(s)
      Run time limit on kernels:                     No
      Integrated GPU sharing Host Memory:            No
      Support host page-locked memory mapping:       Yes
      Alignment requirement for Surfaces:            Yes
      Device has ECC support:                        Disabled
      Device supports Unified Addressing (UVA):      Yes
      Device PCI Domain ID / Bus ID / location ID:   0 / 2 / 0
      Compute Mode:
         < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >
    
    Device 1: "GeForce GTX 1080 Ti"
      CUDA Driver Version / Runtime Version          9.0 / 8.0
      CUDA Capability Major/Minor version number:    6.1
      Total amount of global memory:                 11172 MBytes (11715084288 bytes)
      (28) Multiprocessors, (128) CUDA Cores/MP:     3584 CUDA Cores
      GPU Max Clock rate:                            1620 MHz (1.62 GHz)
      Memory Clock rate:                             5505 Mhz
      Memory Bus Width:                              352-bit
      L2 Cache Size:                                 2883584 bytes
      Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
      Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
      Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 2048 layers
      Total amount of constant memory:               65536 bytes
      Total amount of shared memory per block:       49152 bytes
      Total number of registers available per block: 65536
      Warp size:                                     32
      Maximum number of threads per multiprocessor:  2048
      Maximum number of threads per block:           1024
      Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
      Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
      Maximum memory pitch:                          2147483647 bytes
      Texture alignment:                             512 bytes
      Concurrent copy and kernel execution:          Yes with 2 copy engine(s)
      Run time limit on kernels:                     No
      Integrated GPU sharing Host Memory:            No
      Support host page-locked memory mapping:       Yes
      Alignment requirement for Surfaces:            Yes
      Device has ECC support:                        Disabled
      Device supports Unified Addressing (UVA):      Yes
      Device PCI Domain ID / Bus ID / location ID:   0 / 3 / 0
      Compute Mode:
         < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >
    
    Device 2: "GeForce GTX 1080 Ti"
      CUDA Driver Version / Runtime Version          9.0 / 8.0
      CUDA Capability Major/Minor version number:    6.1
      Total amount of global memory:                 11172 MBytes (11715084288 bytes)
      (28) Multiprocessors, (128) CUDA Cores/MP:     3584 CUDA Cores
      GPU Max Clock rate:                            1620 MHz (1.62 GHz)
      Memory Clock rate:                             5505 Mhz
      Memory Bus Width:                              352-bit
      L2 Cache Size:                                 2883584 bytes
      Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
      Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
      Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 2048 layers
      Total amount of constant memory:               65536 bytes
      Total amount of shared memory per block:       49152 bytes
      Total number of registers available per block: 65536
      Warp size:                                     32
      Maximum number of threads per multiprocessor:  2048
      Maximum number of threads per block:           1024
      Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
      Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
      Maximum memory pitch:                          2147483647 bytes
      Texture alignment:                             512 bytes
      Concurrent copy and kernel execution:          Yes with 2 copy engine(s)
      Run time limit on kernels:                     No
      Integrated GPU sharing Host Memory:            No
      Support host page-locked memory mapping:       Yes
      Alignment requirement for Surfaces:            Yes
      Device has ECC support:                        Disabled
      Device supports Unified Addressing (UVA):      Yes
      Device PCI Domain ID / Bus ID / location ID:   0 / 130 / 0
      Compute Mode:
         < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >
    
    Device 3: "GeForce GTX 1080 Ti"
      CUDA Driver Version / Runtime Version          9.0 / 8.0
      CUDA Capability Major/Minor version number:    6.1
      Total amount of global memory:                 11172 MBytes (11715084288 bytes)
      (28) Multiprocessors, (128) CUDA Cores/MP:     3584 CUDA Cores
      GPU Max Clock rate:                            1620 MHz (1.62 GHz)
      Memory Clock rate:                             5505 Mhz
      Memory Bus Width:                              352-bit
      L2 Cache Size:                                 2883584 bytes
      Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
      Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
      Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 2048 layers
      Total amount of constant memory:               65536 bytes
      Total amount of shared memory per block:       49152 bytes
      Total number of registers available per block: 65536
      Warp size:                                     32
      Maximum number of threads per multiprocessor:  2048
      Maximum number of threads per block:           1024
      Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
      Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
      Maximum memory pitch:                          2147483647 bytes
      Texture alignment:                             512 bytes
      Concurrent copy and kernel execution:          Yes with 2 copy engine(s)
      Run time limit on kernels:                     No
      Integrated GPU sharing Host Memory:            No
      Support host page-locked memory mapping:       Yes
      Alignment requirement for Surfaces:            Yes
      Device has ECC support:                        Disabled
      Device supports Unified Addressing (UVA):      Yes
      Device PCI Domain ID / Bus ID / location ID:   0 / 131 / 0
      Compute Mode:
         < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >
    > Peer access from GeForce GTX 1080 Ti (GPU0) -> GeForce GTX 1080 Ti (GPU1) : Yes
    > Peer access from GeForce GTX 1080 Ti (GPU0) -> GeForce GTX 1080 Ti (GPU2) : No
    > Peer access from GeForce GTX 1080 Ti (GPU0) -> GeForce GTX 1080 Ti (GPU3) : No
    > Peer access from GeForce GTX 1080 Ti (GPU1) -> GeForce GTX 1080 Ti (GPU0) : Yes
    > Peer access from GeForce GTX 1080 Ti (GPU1) -> GeForce GTX 1080 Ti (GPU2) : No
    > Peer access from GeForce GTX 1080 Ti (GPU1) -> GeForce GTX 1080 Ti (GPU3) : No
    > Peer access from GeForce GTX 1080 Ti (GPU2) -> GeForce GTX 1080 Ti (GPU0) : No
    > Peer access from GeForce GTX 1080 Ti (GPU2) -> GeForce GTX 1080 Ti (GPU1) : No
    > Peer access from GeForce GTX 1080 Ti (GPU2) -> GeForce GTX 1080 Ti (GPU3) : Yes
    > Peer access from GeForce GTX 1080 Ti (GPU3) -> GeForce GTX 1080 Ti (GPU0) : No
    > Peer access from GeForce GTX 1080 Ti (GPU3) -> GeForce GTX 1080 Ti (GPU1) : No
    > Peer access from GeForce GTX 1080 Ti (GPU3) -> GeForce GTX 1080 Ti (GPU2) : Yes
    
    deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 9.0, CUDA Runtime Version = 8.0, NumDevs = 4, Device0 = GeForce GTX 1080 Ti, Device1 = GeForce GTX 1080 Ti, Device2 = GeForce GTX 1080 Ti, Device3 = GeForce GTX 1080 Ti
    Result = PASS

    大家可以在自己PC或者工作机上尝试一下。

    再啰嗦两句

    GPU就是用很多简单的计算单元去完成大量的计算任务,纯粹的人海战术。这种策略基于一个前提,就是小学生A和小学生B的工作没有什么依赖性,是互相独立的。

    但有一点需要强调,虽然GPU是为了图像处理而生的,但是我们通过前面的介绍可以发现,它在结构上并没有专门为图像服务的部件,只是对CPU的结构进行了优化与调整,所以现在GPU不仅可以在图像处理领域大显身手,它还被用来科学计算、密码破解、数值分析,海量数据处理(排序,Map-Reduce等),金融分析等需要大规模并行计算的领域。

    所以GPU也可以认为是一种较通用的芯片。

     

    2. CUDA软件构架

    CUDA是一种新的操作GPU计算的硬件和软件架构,它将GPU视作一个数据并行计算设备,而且无需把这些计算映射到图形API。操作系统的多任务机制可以同时管理CUDA访问GPU和图形程序的运行库,其计算特性支持利用CUDA直观地编写GPU核心程序。目前Tesla架构具有在笔记本电脑、台式机、工作站和服务器上的广泛可用性,配以C/C++语言的编程环境和CUDA软件,使这种架构得以成为最优秀的超级计算平台。

    CUDA软件层次结构


    CUDA在软件方面组成有:一个CUDA库、一个应用程序编程接口(API)及其运行库(Runtime)、两个较高级别的通用数学库,即CUFFT和CUBLAS。CUDA改进了DRAM的读写灵活性,使得GPU与CPU的机制相吻合。另一方面,CUDA提供了片上(on-chip)共享内存,使得线程之间可以共享数据。应用程序可以利用共享内存来减少DRAM的数据传送,更少的依赖DRAM的内存带宽。

     

    3. 编程模型

    CUDA程序构架分为两部分:HostDevice。一般而言,Host指的是CPUDevice指的是GPU。在CUDA程序构架中,主程序还是由CPU来执行,而当遇到数据并行处理的部分,CUDA 就会将程序编译成GPU能执行的程序,并传送到GPU。而这个程序在CUDA里称做(kernel)。CUDA允许程序员定义称为核的C语言函数,从而扩展了C语言,在调用此类函数时,它将由N个不同的CUDA线程并行执行N次,这与普通的C语言函数只执行一次的方式不同。执行核的每个线程都会被分配一个独特的线程ID,可通过内置的threadIdx变量在内核中访问此ID。在 CUDA 程序中,主程序在调用任何GPU内核之前,必须对核进行执行配置,即确定线程块数和每个线程块中的线程数以及共享内存大小。

    3.1 线程层次结构

    在GPU中要执行的线程,根据最有效的数据共享来创建块(Block),其类型有一维、二维或三维。在同一个块内的线程可彼此协作,通过一些共享存储器来共享数据,并同步其执行来协调存储器访问。一个块中的所有线程都必须位于同一个处理器核心中。因而,一个处理器核心的有限存储器资源制约了每个块的线程数量。在早期的NVIDIA 架构中,一个线程块最多可以包含 512个线程,而在后期出现的一些设备中则最多可支持1024个线程。一般GPU程序线程数目是很多的,所以不能把所有的线程都塞到同一个块里。但一个内核可由多个大小相同的线程块同时执行,因而线程总数应等于每个块的线程数乘以块的数量。这些同样维度和大小的块将组织为一个一维或二维线程块网格(Grid)。具体框架如下图所示。

    线程块网格

    NOTICE:

    线程(Thread)
    一般通过GPU的一个核进行处理。(可以表示成一维,二维,三维,具体下面再细说)。
    线程块(Block)

    1. 由多个线程组成(可以表示成一维,二维,三维,具体下面再细说)。
    2. 各block是并行执行的,block间无法通信,也没有执行顺序。
    3. 注意线程块的数量限制为不超过65535(硬件限制)。

    线程格(Grid)
    由多个线程块组成(可以表示成一维,二维,三维,具体下面再细说)。
    线程束
    在CUDA架构中,线程束是指一个包含32个线程的集合,这个线程集合被“编织在一起”并且“步调一致”的形式执行。在程序中的每一行,线程束中的每个线程都将在不同数据上执行相同的命令。

    从硬件上看

    SP:最基本的处理单元,streaming processor,也称为CUDA core。最后具体的指令和任务都是在SP上处理的。GPU进行并行计算,也就是很多个SP同时做处理。
    SM:多个SP加上其他的一些资源组成一个streaming multiprocessor。也叫GPU大核,其他资源如:warp scheduler,register,shared memory等。SM可以看做GPU的心脏(对比CPU核心),register和shared memory是SM的稀缺资源。CUDA将这些资源分配给所有驻留在SM中的threads。因此,这些有限的资源就使每个SM中active warps有非常严格的限制,也就限制了并行能力。

    从软件上看

    thread:一个CUDA的并行程序会被以许多个threads来执行。
    block:数个threads会被群组成一个block,同一个block中的threads可以同步,也可以通过shared memory通信。
    grid:多个blocks则会再构成grid。
    warp:GPU执行程序时的调度单位,目前cuda的warp的大小为32,同在一个warp的线程,以不同数据资源执行相同的指令,这就是所谓 SIMT。

    3.2 存储器层次结构

    CUDA设备拥有多个独立的存储空间,其中包括:全局存储器、本地存储器、共享存储器、常量存储器、纹理存储器和寄存器,如图

    CUDA设备上的存储器

    NOTICE:

    主机(Host)
    将CPU及系统的内存(内存条)称为主机。
    设备(Device)
    将GPU及GPU本身的显示内存称为设备。
    动态随机存取存储器(DRAM)
    DRAM(Dynamic Random Access Memory),即动态随机存取存储器,最为常见的系统内存DRAM只能将数据保持很短的时间。为了保持数据,DRAM使用电容存储,所以必须隔一段时间刷新(refresh)一次,如果存储单元没有被刷新,存储的信息就会丢失。 (关机就会丢失数据)

    CUDA线程可在执行过程中访问多个存储器空间的数据,如下图所示其中:

    • 每个线程都有一个私有的本地存储器。
    • 每个线程块都有一个共享存储器,该存储器对于块内的所有线程都是可见的,并且与块具有相同的生命周期。
    • 所有线程都可访问相同的全局存储器。
    • 此外还有两个只读的存储器空间,可由所有线程访问,这两个空间是常量存储器空间和纹理存储器空间。全局、固定和纹理存储器空间经过优化,适于不同的存储器用途。纹理存储器也为某些特殊的数据格式提供了不同的寻址模式以及数据过滤,方便Host对流数据的快速存取。
    存储器的应用层次

    3.3 主机(Host)和设备(Device)

    如下图所示,CUDA假设线程可在物理上独立的设备上执行,此类设备作为运行C语言程序的主机的协处理器操作。内核在GPU上执行,而C语言程序的其他部分在CPU上执行(即串行代码在主机上执行,而并行代码在设备上执行)。此外,CUDA还假设主机和设备均维护自己的DRAM,分别称为主机存储器和设备存储器。因而,一个程序通过调用CUDA运行库来管理对内核可见的全局、固定和纹理存储器空间。这种管理包括设备存储器的分配和取消分配,还包括主机和设备存储器之间的数据传输。

     

    4. CUDA软硬件

    4.1 CUDA术语

    由于CUDA中存在许多概念和术语,诸如SM、block、SP等多个概念不容易理解,将其与CPU的一些概念进行比较,如下表所示。

    CPU GPU 层次
    算术逻辑和控制单元 流处理器(SM) 硬件
    算术单元 批量处理器(SP) 硬件
    进程 Block 软件
    线程 thread 软件
    调度单位 Warp 软件

    4.2 硬件利用率

    当为一个GPU分配一个内核函数,我们关心的是如何才能充分利用GPU的计算能力,但由于不同的硬件有不同的计算能力,SM一次最多能容纳的线程数也不尽相同,SM一次最多能容纳的线程数量主要与底层硬件的计算能力有关,如下表显示了在不同的计算能力的设备上,每个线程块上开启不同数量的线程时设备的利用率。

    计算能力 每个线程块的线程数 1.0 1.1 1.2 1.3 2.0 2.1 3.0
    64 67 50 50 50 33 33 50
    96 100 100 75 75 50 50 75
    128 100 100 100 100 67 67 100
    192 100 100 94 94 100 100 94
    96 100 100 100 100 100 100 10
    ... ... ... ...            

    查看显卡利用率 (以实验室服务器为例)
    输入以下命令

    nvidia-smi
    Thu Aug 23 21:06:36 2018       
    +-----------------------------------------------------------------------------+
    | NVIDIA-SMI 384.130                Driver Version: 384.130                   |
    |-------------------------------+----------------------+----------------------+
    | GPU  Name        Persistence-M| Bus-Id        Disp.A | Volatile Uncorr. ECC |
    | Fan  Temp  Perf  Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |
    |===============================+======================+======================|
    |   0  GeForce GTX 108...  Off  | 00000000:02:00.0 Off |                  N/A |
    | 29%   41C    P0    58W / 250W |      0MiB / 11171MiB |      0%      Default |
    +-------------------------------+----------------------+----------------------+
    |   1  GeForce GTX 108...  Off  | 00000000:03:00.0 Off |                  N/A |
    | 33%   47C    P0    57W / 250W |      0MiB / 11172MiB |      0%      Default |
    +-------------------------------+----------------------+----------------------+
    |   2  GeForce GTX 108...  Off  | 00000000:82:00.0 Off |                  N/A |
    | 36%   49C    P0    59W / 250W |      0MiB / 11172MiB |      0%      Default |
    +-------------------------------+----------------------+----------------------+
    |   3  GeForce GTX 108...  Off  | 00000000:83:00.0 Off |                  N/A |
    | 33%   46C    P0    51W / 250W |      0MiB / 11172MiB |      1%      Default |
    +-------------------------------+----------------------+----------------------+
                                                                                   
    +-----------------------------------------------------------------------------+
    | Processes:                                                       GPU Memory |
    |  GPU       PID   Type   Process name                             Usage      |
    |=============================================================================|
    |  No running processes found                                                 |
    +-----------------------------------------------------------------------------+

     

    5. 并行计算

    5.1 并发性

    CUDA将问题分解成线程块的网格,每块包含多个线程。快可以按任意顺序执行。不过在某个时间点上,只有一部分块处于执行中。一旦被调用到GUP包含的N个“流处理器簇(SM)”中的一个上执行,一个块必须从开始到结束。网格中的块可以被分配到任意一个有空闲槽的SM上。起初,可以采用“轮询调度”策略,以确保分配到每一个SM上的块数基本相同。对绝大多数内核程序而言,分块的数量应该是GPU中物理SM数量的八倍或更多倍。

    以一个军队比喻,假设有一支由士兵(线程)组成的部队(网格)。部队被分成若干个连(块),每个连队由一位连长来指挥。按照32名士兵一个班(一个线程束),连队又进一步分成若干个班,每个班由一个班长来指挥。

    基于GPU的线程视图

    要执行某个操作,总司令(内核程序/ 主机程序)必须提供操作名称及相应的数据。每个士兵(线程)只处理分配给他的问题中的一小块。在连长(负责一个块)或班长(负责一个束)的控制下,束与束之间的线程或者一个束内部的线程之间,要经常地交换数据。但是,连队(块)之间的协同就得由总司令(内核函数/ 主机程序)来控制。

    5.2 局部性

    对于GPU程序设计,程序员必须处理局部性。对于一个给定的工作,他需要事先思考需要哪些工具或零件(即存储地址或数据结构),然后一次性地把他们从硬件仓库(全局内存)可能把与这些数据相关的不同工作都执行了,避免发生“取来--存回--为了下一个工作再取”。

    5.3 缓存一致性

    GPU与CPU在缓存上的一个重要差别就是“缓存一致性”问题。对于“缓存一致”的系统,一个内存的写操作需要通知所有核的各个级别的缓存。因此,无论何时,所有的处理器核看到的内存视图是完全一样的。随着处理器中核数量的增多,这个“通知”的开销迅速增大,使得“缓存一致性”成为限制一个处理器中核数量不能太多的一重要因素。“缓存一致”系统中最坏的情况是,一个内存操作会强迫每个核的缓存都进行更新,进而每个核都要对相邻的内存单元写操作。

    相比之下,非“缓存一致”系统不会自动地更新其他核的缓存。它需要由程序员写清楚每个处理器核输出的各自不同的目标区域。从程序的视角看,这支持一个核仅负责一个输出或者一个小的输出集。通常,CPU遵循“缓存一致性”原则,而GPU则不是。故GPU能够扩展到一个芯片内具有大数量的核心(流处理器簇)。

    5.4 弗林分类法

    根据弗林分类法,计算机的结构类型有:

    SIMD--单指令,多数据
    MIMD--多指令,多数据
    SISD--单指令,单数据
    MISD--多指令,单数据

    5.5 分条 / 分块

    CUDA提供的简单二维网格模型。对于很多问题,这样的模型就足够了。如果在一个块内,你的工作是线性分布的,那么你可以很好地将其他分解成CUDA块。由于在一个SM内,最多可以分配16个块,而在一个GPU内有16个(有些是32个)SM,所以问题分成256个甚至更多的块都可以。实际上,我们更倾向于把一个块内的元素总数限制为128、256、或者512,这样有助于在一个典型的数据集内划分出更多数量的块。

    5.6 快速傅氏变换(FFT)

    FFT: FFT(Fast Fourier Transformation)是离散傅氏变换(DFT)的快速算法。即为快速傅氏变换。它是根据离散傅氏变换的奇、偶、虚、实等特性,对离散傅立叶变换的算法进行改进获得的。

    由于不是刚需,这里不展开讲。好奇的你可以点击楼下时光机,通过下面的教程进行学习。
    FFT(最详细最通俗的入门手册)

    5.7 CUDA计算能力的含义

    体现GPU计算能力的两个重要特征:
    1)CUDA核的个数;
    2)存储器大小。
    描述GPU性能的两个重要指标: :
    1)计算性能峰值;
    2)存储器带宽。

    参考
    1.CUDA计算能力的含义
    2.CUDA GPUs

     

    6. 实践

    6.1 Ubuntu 系统下环境搭建

    6.1.1 系统要求

    要搭建 CUDA 环境,我们需要自己的计算机满足以下这三个条件:
    1. 有至少一颗支持 CUDA 的 GPU(我的是GeForece GT 650M)
    2. 有满足版本要求的 gcc 编译器和链接工具
    3. 有 NVIDIA 提供的 CUDA 工具包(点击神奇的小链接下载)

    6.1.2 准备工作

    下面,我们一步一步来验证自己的系统是否满足安装要求。
    Step 1: 验证计算机是否拥有至少一颗支持 CUDA 的 GPU
    打开终端(Ctrl + Alt + T),键入以下命令:

    lspci | grep -i nvidia
    

    可以看到以下内容(结果因人而异,与具体的GPU有关)


     

    看到这个就说明至少有一颗支持 CUDA 的 GPU,可以进入下一步了。

    Step 2: 验证一下自己操作系统的版本
    键入命令:

    lsb_release -a
    
    No LSB modules are available.
    Distributor ID: Ubuntu
    Description:    Ubuntu 16.04.4 LTS
    Release:    16.04
    Codename:   xenial
    

    更多信息请移步Ubuntu查看版本信息

    Step 3: 验证 gcc 编译器的版本
    键入命令:

    gcc --version
    

    或者

    gcc -v
    

    得到如下信息

    gcc (Ubuntu 5.4.0-6ubuntu1~16.04.10) 5.4.0 20160609
    Copyright (C) 2015 Free Software Foundation, Inc.
    This is free software; see the source for copying conditions.  There is NO
    warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
    

    Step 4: 验证系统内核版本
    键入命令:

    uname -r
    

    得到如下信息

    对照官方提供的对各种 Linux 发行版的安装要求进行安装

    6.1.3 搭建 CUDA 环境

    Step 1: 安装 CUDA 工具包
    在前面几项验证都顺利通过以后就来到最关键的一步。首先下载对应自己系统版本的 CUDA 工具包(以CUDA Toolkit 9.2 为例),然后进入到安装包所在目录:

    sudo dpkg -i cuda-repo-ubuntu1604-9-2-local_9.2.148-1_amd64.deb
    
    sudo apt-key add /var/cuda-repo-<version>/7fa2af80.pub
    
    sudo apt-get update
    
    sudo apt-get install cuda
    

    NOTICE:

    Other installation options are available in the form of meta-packages. For example, to install all the library packages, replace "cuda" with the "cuda-libraries-9-2" meta package. For more information on all the available meta packages click here.

    此时静静地等待安装完成。不出意外,一段时间后安装完成了。
    Step 2: 设置环境变量
    首先在 PATH 变量中加入 /usr/local/cuda-9.2/bin,在Terminal中执行:

    export PATH=/usr/local/cuda-9.2/bin:$PATH
    

    然后在 LD_LIBRARY_PATH 变量中添加 /usr/local/cuda-9.2/lib64,执行:

    export  LD_LIBRARY_PATH=/usr/local/cuda-9.2/lib64:$LD_LIBRARY_PATH
    

    Step 3: 验证环境搭建是否成功
    首先执行命令:

    nvcc -V
    

    关于测试...聪明的你一定想起来了,我们前面是讲过怎么做的。
    对,没错,就在1.5小节,话不多说,自行上翻吧。

    看到通过测试,到这里,64位 Ubuntu 16.04 系统下 CUDA 环境搭建就完成了。

    6.2 CUDA编程

    6.2.1 核函数

    1. 在GPU上执行的函数通常称为核函数。
    2. 一般通过标识符__global__修饰,调用通过<<<参数1,参数2>>>,用于说明内核函数中的线程数量,以及线程是如何组织的。
    3. 以线程格(Grid)的形式组织,每个线程格由若干个线程块(block)组成,而每个线程块又由若干个线程(thread)组成。
    4.是以block为单位执行的。
    5. 叧能在主机端代码中调用。
    6. 调用时必须声明内核函数的执行参数。
    7. 在编程时,必须先为kernel函数中用到的数组或变量分配好足够的空间,再调用kernel函数,否则在GPU计算时会发生错误,例如越界或报错,甚至导致蓝屏和死机。

    看完基本知识,装好CUDA以后,就可以开始写第一个CUDA程序了:

    #include <cuda_runtime.h>
     
    int main(){
    printf("Hello world!\n");
    }
    

    慢着,这个程序和C有什么区别?用到显卡了吗?

    答:没有区别,没用显卡。如果你非要用显卡干点什么事情的话,可以改成这个样子:

    /*
     * @file_name HelloWorld.cu  后缀名称.cu
     */
    
    #include <stdio.h>
    #include <cuda_runtime.h>  //头文件
    
    //核函数声明,前面的关键字__global__
    __global__ void kernel( void ) {
    }
    
    int main( void ) {
        //核函数的调用,注意<<<1,1>>>,第一个1,代表线程格里只有一个线程块;第二个1,代表一个线程块里只有一个线程。
        kernel<<<1,1>>>();
        printf( "Hello, World!\n" );
        return 0;
    }
    

    6.2.2 dim3结构类型

    1. dim3是基于uint3定义的矢量类型,相当亍由3个unsigned int型组成的结构体。uint3类型有三个数据成员unsigned int x; unsigned int y; unsigned int z;
    2. 可使用于一维、二维或三维的索引来标识线程,构成一维、二维或三维线程块。
    3. dim3结构类型变量用在核函数调用的<<<,>>>中。
    4. 相关的几个内置变量
      4.1. threadIdx,顾名思义获取线程thread的ID索引;如果线程是一维的那么就取threadIdx.x,二维的还可以多取到一个值threadIdx.y,以此类推到三维threadIdx.z
      4.2. blockIdx,线程块的ID索引;同样有blockIdx.xblockIdx.yblockIdx.z
      4.3. blockDim,线程块的维度,同样有blockDim.xblockDim.yblockDim.z
      4.4. gridDim,线程格的维度,同样有gridDim.xgridDim.ygridDim.z
    5. 对于一维的block,线程的threadID=threadIdx.x
    6. 对于大小为(blockDim.x, blockDim.y)的 二维block,线程的threadID=threadIdx.x+threadIdx.y*blockDim.x
    7. 对于大小为(blockDim.x, blockDim.y, blockDim.z)的 三维 block,线程的threadID=threadIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y
    8. 对于计算线程索引偏移增量为已启动线程的总数。如stride = blockDim.x * gridDim.x; threadId += stride

    6.2.3 函数修饰符

    1.__global__,表明被修饰的函数在设备上执行,但在主机上调用。

    2.__device__,表明被修饰的函数在设备上执行,但只能在其他__device__函数或者__global__函数中调用。

    6.2.4 常用的GPU内存函数

    cudaMalloc()
    1. 函数原型: cudaError_t cudaMalloc (void **devPtr, size_t size)
    2. 函数用处:与C语言中的malloc函数一样,只是此函数在GPU的内存你分配内存。
    3. 注意事项:
    3.1. 可以将cudaMalloc()分配的指针传递给在设备上执行的函数;
    3.2. 可以在设备代码中使用cudaMalloc()分配的指针进行设备内存读写操作;
    3.3. 可以将cudaMalloc()分配的指针传递给在主机上执行的函数;
    3.4. 不可以在主机代码中使用cudaMalloc()分配的指针进行主机内存读写操作(即不能进行解引用)。

    cudaMemcpy()
    1. 函数原型:cudaError_t cudaMemcpy (void *dst, const void *src, size_t count, cudaMemcpyKind kind)
    2. 函数作用:与c语言中的memcpy函数一样,只是此函数可以在主机内存和GPU内存之间互相拷贝数据。
    3. 函数参数:cudaMemcpyKind kind表示数据拷贝方向,如果kind赋值为cudaMemcpyDeviceToHost表示数据从设备内存拷贝到主机内存。
    4. 与C中的memcpy()一样,以同步方式执行,即当函数返回时,复制操作就已经完成了,并且在输出缓冲区中包含了复制进去的内容。
    5. 相应的有个异步方式执行的函数cudaMemcpyAsync(),这个函数详解请看下面的流一节有关内容。

    cudaFree()
    1. 函数原型:cudaError_t cudaFree ( void* devPtr )
    2. 函数作用:与c语言中的free()函数一样,只是此函数释放的是cudaMalloc()分配的内存。
    下面实例用于解释上面三个函数

    #include <stdio.h>
    #include <cuda_runtime.h>
    __global__ void add( int a, int b, int *c ) {
        *c = a + b;
    }
    int main( void ) {
        int c;
        int *dev_c;
        //cudaMalloc()
        cudaMalloc( (void**)&dev_c, sizeof(int) );
        //核函数执行
        add<<<1,1>>>( 2, 7, dev_c );   
        //cudaMemcpy()
        cudaMemcpy( &c, dev_c, sizeof(int),cudaMemcpyDeviceToHost ) ;
        printf( "2 + 7 = %d\n", c );
        //cudaFree()
        cudaFree( dev_c );
     
        return 0;
    }
    

    6.2.5 GPU内存分类

    全局内存
    通俗意义上的设备内存。

    共享内存
    1. 位置:设备内存。
    2. 形式:关键字__shared__添加到变量声明中。如__shared__ float cache[10]
    3. 目的:对于GPU上启动的每个线程块,CUDA C编译器都将创建该共享变量的一个副本。线程块中的每个线程都共享这块内存,但线程却无法看到也不能修改其他线程块的变量副本。这样使得一个线程块中的多个线程能够在计算上通信和协作。

    常量内存
    1. 位置:设备内存
    2. 形式:关键字__constant__添加到变量声明中。如__constant__ float s[10];。
    3. 目的:为了提升性能。常量内存采取了不同于标准全局内存的处理方式。在某些情况下,用常量内存替换全局内存能有效地减少内存带宽。
    4. 特点:常量内存用于保存在核函数执行期间不会发生变化的数据。变量的访问限制为只读。NVIDIA硬件提供了64KB的常量内存。不再需要cudaMalloc()或者cudaFree(),而是在编译时,静态地分配空间。
    5. 要求:当我们需要拷贝数据到常量内存中应该使用cudaMemcpyToSymbol(),而cudaMemcpy()会复制到全局内存。
    6. 性能提升的原因:
    6.1. 对常量内存的单次读操作可以广播到其他的“邻近”线程。这将节约15次读取操作。(为什么是15,因为“邻近”指半个线程束,一个线程束包含32个线程的集合。)
    6.2. 常量内存的数据将缓存起来,因此对相同地址的连续读操作将不会产生额外的内存通信量。

    纹理内存
    1. 位置:设备内存
    2. 目的:能够减少对内存的请求并提供高效的内存带宽。是专门为那些在内存访问模式中存在大量空间局部性的图形应用程序设计,意味着一个线程读取的位置可能与邻近线程读取的位置“非常接近”。如下图:

    3. 纹理变量(引用)必须声明为文件作用域内的全局变量。
    4. 形式:分为一维纹理内存 和 二维纹理内存。
    4.1. 一维纹理内存
    4.1.1. 用texture<类型>类型声明,如texture<float> texIn
    4.1.2. 通过cudaBindTexture()绑定到纹理内存中。
    4.1.3. 通过tex1Dfetch()来读取纹理内存中的数据。
    4.1.4. 通过cudaUnbindTexture()取消绑定纹理内存。
    4.2. 二维纹理内存
    4.2.1. 用texture<类型,数字>类型声明,如texture<float,2> texIn
    4.2.2. 通过cudaBindTexture2D()绑定到纹理内存中。
    4.2.3. 通过tex2D()来读取纹理内存中的数据。
    4.2.4. 通过cudaUnbindTexture()取消绑定纹理内存。

     

    固定内存
    1. 位置:主机内存。
    2. 概念:也称为页锁定内存或者不可分页内存,操作系统将不会对这块内存分页并交换到磁盘上,从而确保了该内存始终驻留在物理内存中。因此操作系统能够安全地使某个应用程序访问该内存的物理地址,因为这块内存将不会破坏或者重新定位。
    3. 目的:提高访问速度。由于GPU知道主机内存的物理地址,因此可以通过“直接内存访问DMA(Direct Memory Access)技术来在GPU和主机之间复制数据。由于DMA在执行复制时无需CPU介入。因此DMA复制过程中使用固定内存是非常重要的。
    4. 缺点:使用固定内存,将失去虚拟内存的所有功能;系统将更快的耗尽内存。
    5. 建议:对cudaMemcpy()函数调用中的源内存或者目标内存,才使用固定内存,并且在不再需要使用它们时立即释放。
    6. 形式:通过cudaHostAlloc()函数来分配;通过cudaFreeHost()释放。
    7. 只能以异步方式对固定内存进行复制操作。

    原子性
    1. 概念:如果操作的执行过程不能分解为更小的部分,我们将满足这种条件限制的操作称为原子操作。
    2. 形式:函数调用,如atomicAdd(addr,y)将生成一个原子的操作序列,这个操作序列包括读取地址addr处的值,将y增加到这个值,以及将结果保存回地址addr

    6.2.6 常用线程操作函数

    同步方法__syncthreads(),这个函数的调用,将确保线程块中的每个线程都执行完__syscthreads()前面的语句后,才会执行下一条语句。

    使用事件来测量性能
    1. 用途:为了测量GPU在某个任务上花费的时间。CUDA中的事件本质上是一个GPU时间戳。由于事件是直接在GPU上实现的。因此不适用于对同时包含设备代码和主机代码的混合代码设计。
    2. 形式:首先创建一个事件,然后记录事件,再计算两个事件之差,最后销毁事件。如:

    cudaEvent_t start, stop;
    cudaEventCreate( &start );
    cudaEventCreate( &stop );
    cudaEventRecord( start, 0 );
    //do something
    cudaEventRecord( stop, 0 );
    float   elapsedTime;
    cudaEventElapsedTime( &elapsedTime,start, stop );
    cudaEventDestroy( start );
    cudaEventDestroy( stop );

    6.2.7 流

    1. 扯一扯:并发重点在于一个极短时间段内运行多个不同的任务;并行重点在于同时运行一个任务。
    2. 任务并行性:是指并行执行两个或多个不同的任务,而不是在大量数据上执行同一个任务。
    3. 概念:CUDA流表示一个GPU操作队列,并且该队列中的操作将以指定的顺序执行。我们可以在流中添加一些操作,如核函数启动,内存复制以及事件的启动和结束等。这些操作的添加到流的顺序也是它们的执行顺序。可以将每个流视为GPU上的一个任务,并且这些任务可以并行执行。
    4. 硬件前提:必须是支持设备重叠功能的GPU。支持设备重叠功能,即在执行一个核函数的同时,还能在设备与主机之间执行复制操作。
    5. 声明与创建:声明cudaStream_t stream;,创建cudaSteamCreate(&stream);。
    6. cudaMemcpyAsync():前面在cudaMemcpy()中提到过,这是一个以异步方式执行的函数。在调用cudaMemcpyAsync()时,只是放置一个请求,表示在流中执行一次内存复制操作,这个流是通过参数stream来指定的。当函数返回时,我们无法确保复制操作是否已经启动,更无法保证它是否已经结束。我们能够得到的保证是,复制操作肯定会当下一个被放入流中的操作之前执行。传递给此函数的主机内存指针必须是通过cudaHostAlloc()分配好的内存。(流中要求固定内存)
    7. 流同步:通过cudaStreamSynchronize()来协调。
    8. 流销毁:在退出应用程序之前,需要销毁对GPU操作进行排队的流,调用cudaStreamDestroy()
    9. 针对多个流:
      9.1. 记得对流进行同步操作。
      9.2. 将操作放入流的队列时,应采用宽度优先方式,而非深度优先的方式,换句话说,不是首先添加第0个流的所有操作,再依次添加后面的第1,2,…个流。而是交替进行添加,比如将a的复制操作添加到第0个流中,接着把a的复制操作添加到第1个流中,再继续其他的类似交替添加的行为。
      9.3. 要牢牢记住操作放入流中的队列中的顺序影响到CUDA驱动程序调度这些操作和流以及执行的方式。

    TIPS:

    1. 当线程块的数量为GPU中处理数量的2倍时,将达到最优性能。
    2. 核函数执行的第一个计算就是计算输入数据的偏移。每个线程的起始偏移都是0到线程数量减1之间的某个值。然后,对偏移的增量为已启动线程的总数。

    6.2.8 这是一个栗子

    我们尝试用一个程序来比较cuda/c在GPU/CPU的运行效率,来不及了,快上车。
    这是一个CUDA程序,请保存文件名为“文件名.cu”,在你的PC或者服务器上运行。

    #include "cuda_runtime.h"
    #include "device_launch_parameters.h"
     
    #include <stdio.h>
    #include <time.h>
     
    #define N (1024*1024)
    #define M (10000)
    #define THREADS_PER_BLOCK 1024
     
    void serial_add(double *a, double *b, double *c, int n, int m)
    {
        for(int index=0;index<n;index++)
        {
            for(int j=0;j<m;j++)
            {
                c[index] = a[index]*a[index] + b[index]*b[index];
            }
        }
    }
     
    __global__ void vector_add(double *a, double *b, double *c)
    {
        int index = blockIdx.x * blockDim.x + threadIdx.x;
            for(int j=0;j<M;j++)
            {
                c[index] = a[index]*a[index] + b[index]*b[index];
            }
    }
     
    int main()
    {
        clock_t start,end;
     
        double *a, *b, *c;
        int size = N * sizeof( double );
     
        a = (double *)malloc( size );
        b = (double *)malloc( size );
        c = (double *)malloc( size );
     
        for( int i = 0; i < N; i++ )
        {
            a[i] = b[i] = i;
            c[i] = 0;
        }
     
        start = clock();
        serial_add(a, b, c, N, M);
     
        printf( "c[%d] = %f\n",0,c[0] );
        printf( "c[%d] = %f\n",N-1, c[N-1] );
     
        end = clock();
     
        float time1 = ((float)(end-start))/CLOCKS_PER_SEC;
        printf("CPU: %f seconds\n",time1);
     
        start = clock();
        double *d_a, *d_b, *d_c;
     
     
        cudaMalloc( (void **) &d_a, size );
        cudaMalloc( (void **) &d_b, size );
        cudaMalloc( (void **) &d_c, size );
     
     
        cudaMemcpy( d_a, a, size, cudaMemcpyHostToDevice );
        cudaMemcpy( d_b, b, size, cudaMemcpyHostToDevice );
     
        vector_add<<< (N + (THREADS_PER_BLOCK-1)) / THREADS_PER_BLOCK, THREADS_PER_BLOCK >>>( d_a, d_b, d_c );
     
        cudaMemcpy( c, d_c, size, cudaMemcpyDeviceToHost );
     
     
        printf( "c[%d] = %f\n",0,c[0] );
        printf( "c[%d] = %f\n",N-1, c[N-1] );
     
     
        free(a);
        free(b);
        free(c);
        cudaFree( d_a );
        cudaFree( d_b );
        cudaFree( d_c );
     
        end = clock();
        float time2 = ((float)(end-start))/CLOCKS_PER_SEC;
        printf("CUDA: %f seconds, Speedup: %f\n",time2, time1/time2);
     
        return 0;
    }
    

    效率对比
    我们通过修改count的值并且加大循环次数来观察变量的效率的差别。

    运行结果:


    可见在数据量大的情况下效率还是相当不错的。

     

    7. GPU or FPGA

    GPU优势
    1.从峰值性能来说,GPU(10Tflops)远远高于FPGA(<1TFlops);

    2.GPU相对于FPGA还有一个优势就是内存接口, GPU的内存接口(传统的GDDR5,最近更是用上了HBM和HBM2)的带宽远好于FPGA的传统DDR接口(大约带宽高4-5倍);

    3.功耗方面,虽然GPU的功耗远大于FPGA的功耗,但是如果要比较功耗应该比较在执行效率相同时需要的功耗。如果FPGA的架构优化能做到很好以致于一块FPGA的平均性能能够接近一块GPU,那么FPGA方案的总功耗远小于GPU,散热问题可以大大减轻。反之,如果需要二十块FPGA才能实现一块GPU的平均性能,那么FPGA在功耗方面并没有优势。

    4.FPGA缺点有三点:
    第一,基本单元的计算能力有限。为了实现可重构特性,FPGA 内部有大量极细粒度的基本单元,但是每个单元的计算能力(主要依靠LUT 查找表)都远远低于CPU 和GPU 中的ALU模块。
    第二,速度和功耗相对专用定制芯片(ASIC)仍然存在不小差距。
    第三,FPGA 价格较为昂贵,在规模放量的情况下单块FPGA 的成本要远高于专用定制芯片。最后谁能胜出, 完全取决于FPGA架构优化能否弥补峰值性能的劣势。

    5.个人更推荐: CPU+FPGA的组合模式; 其中FPGA用于整形计算,cpu进行浮点计算和调度,此组合的拥有更高的单位功耗性能和更低的时延。最后更想GPU稳定开放,发挥其长处, 达到真正的物美价廉!

    FPGA优势
    人工智能目前仍处于早期阶段,未来人工智能的主战场是在推理环节,远没有爆发。未来胜负尚未可知,各家技术路线都有机会胜出。目前英伟达的GPU在训练场景中占据着绝对领导地位,但是在未来,专注于推理环节的FPGA必将会发挥巨大的价值。

    FPGA和GPU内都有大量的计算单元,因此它们的计算能力都很强。在进行神经网络运算的时候,两者的速度会比CPU快很多。但是GPU由于架构固定,硬件原生支持的指令也就固定了,而FPGA则是可编程的。其可编程性是关键,因为它让软件与终端应用公司能够提供与其竞争对手不同的解决方案,并且能够灵活地针对自己所用的算法修改电路。

    在平均性能方面,GPU逊于FPGA,FPGA可以根据特定的应用去编程硬件,例如如果应用里面的加法运算非常多就可以把大量的逻辑资源去实现加法器,而GPU一旦设计完就不能改动了,所以不能根据应用去调整硬件资源。
    目前机器学习大多使用SIMD架构,即只需一条指令可以平行处理大量数据,因此用GPU很适合。但是有些应用是MISD,即单一数据需要用许多条指令平行处理,这种情况下用FPGA做一个MISD的架构就会比GPU有优势。 所以,对于平均性能,看的就是FPGA加速器架构上的优势是否能弥补运行速度上的劣势。如果FPGA上的架构优化可以带来相比GPU架构两到三个数量级的优势,那么FPGA在平均性能上会好于GPU。

    在功耗能效比方面,同样由于FPGA的灵活性,在架构优化到很好时,一块FPGA的平均性能能够接近一块GPU,那么FPGA方案的总功耗远小于GPU,散热问题可以大大减轻。 能效比的比较也是类似,能效指的是完成程序执行消耗的能量,而能量消耗等于功耗乘以程序的执行时间。虽然GPU的功耗远大于FPGA的功耗,但是如果FPGA执行相同程序需要的时间比GPU长几十倍,那FPGA在能效比上就没有优势了;反之如果FPGA上实现的硬件架构优化得很适合特定的机器学习应用,执行算法所需的时间仅仅是GPU的几倍或甚至于接近GPU,那么FPGA的能效比就会比GPU强。

    在峰值性能比方面,虽然GPU的峰值性能(10Tflops)远大于FPGA的峰值性能(<1Tflops),但针对特定的场景来讲吞吐量并不比GPU差。

     

    8. 深度学习的三种硬件方案:ASIC,FPGA,GPU

    8.1 对深度学习硬件平台的要求

    要想明白“深度学习”需要怎样的硬件,必须了解深度学习的工作原理。首先在表层上,我们有一个巨大的数据集,并选定了一种深度学习模型。每个模型都有一些内部参数需要调整,以便学习数据。而这种参数调整实际上可以归结为优化问题,在调整这些参数时,就相当于在优化特定的约束条件

    • 矩阵相乘(Matrix Multiplication)——几乎所有的深度学习模型都包含这一运算,它的计算十分密集。

    • 卷积(Convolution)——这是另一个常用的运算,占用了模型中大部分的每秒浮点运算(浮点/秒)。

    • 循环层(Recurrent Layers )——模型中的反馈层,并且基本上是前两个运算的组合。

    • All Reduce——这是一个在优化前对学习到的参数进行传递或解析的运算序列。在跨硬件分布的深度学习网络上执行同步优化时(如AlphaGo的例子),这一操作尤其有效。

    除此之外,深度学习的硬件加速器需要具备数据级别和流程化的并行性、多线程和高内存带宽等特性。 另外,由于数据的训练时间很长,所以硬件架构必须低功耗。 因此,效能功耗比(Performance per Watt)是硬件架构的评估标准之一。

    CNN在应用中,一般采用GPU加速,请解释为什么GPU可以有加速效果,主要加速算法的哪一个部分?

    这里默认gpu加速是指NVIDIA的CUDA加速。CPU是中央处理单元,gpu是图形处理单元,gpu由上千个流处理器(core)作为运算器。执行采用单指令多线程(SIMT)模式。相比于单核CPU(向量机)流水线式的串行操作,虽然gpu单个core计算能力很弱,但是通过大量线程进行同时计算,在数据量很大是会活动较为可观的加速效果。

    具体到cnn,利用gpu加速主要是在conv(卷积)过程上。conv过程同理可以像以上的向量加法一样通过cuda实现并行化。具体的方法很多,不过最好的还是利用fft(快速傅里叶变换)进行快速卷积。NVIDIA提供了cufft库实现fft,复数乘法则可以使用cublas库里的对应的level3的cublasCgemm函数。

    GPU加速的基本准则就是“人多力量大”。CNN说到底主要问题就是计算量大,但是却可以比较有效的拆分成并行问题。随便拿一个层的filter来举例子,假设某一层有n个filter,每一个需要对上一层输入过来的map进行卷积操作。那么,这个卷积操作并不需要按照线性的流程去做,每个滤波器互相之间并不影响,可以大家同时做,然后大家生成了n张新的谱之后再继续接下来的操作。既然可以并行,那么同一时间处理单元越多,理论上速度优势就会越大。所以,处理问题就变得很简单粗暴,就像NV那样,暴力增加显卡单元数(当然,显卡的架构、内部数据的传输速率、算法的优化等等也都很重要)。

    GPU主要是针对图形显示及渲染等技术的出众,而其中的根本是因为处理矩阵算法能力的强大,刚好CNN中涉及大量的卷积,也就是矩阵乘法等,所以在这方面具有优势。

    机器学习的算法一定得经过gpu加速吗?

    不一定。只有需要大量浮点数计算,例如矩阵乘法,才需要GPU加速。 用CNN对图像进行分类就是一个需要大量浮点数计算的典型案例,通常需要GPU加速

    对于ASICFPGA分布式计算,这里不再展开讲,有兴趣的小伙伴可以,自行学习。不过....说不定某天博主心情好,就会梳理一下这几种硬件方案在端到端上应用的区别了。

    菜鸟入门教程就到这里了,聪明的你一定不满足这个入门教程,如有兴趣进一步学习CUDA编程,可移步NVIDIA官方的课程平台CUDA ZONE(PS:中文网站,英文课程)


     

     

    展开全文
  • 什么是CUDA

    千次阅读 2018-10-26 18:09:38
    CUDA(百度百科版) CUDA(Compute Unified Device Architecture),是显卡厂商NVIDIA推出的运算平台。 CUDA™是一种由NVIDIA推出的通用并行计算架构,该架构使GPU能够解决复杂的计算问题。 它包含了CUDA指令集架构...

    CUDA(百度百科版)

    CUDA(Compute Unified Device Architecture),是显卡厂商NVIDIA推出的运算平台。 CUDA™是一种由NVIDIA推出的通用并行计算架构,该架构使GPU能够解决复杂的计算问题。 它包含了CUDA指令集架构(ISA)以及GPU内部的并行计算引擎。

    CUDA(网民有理版)

    CUDA呢简单来说就是GPU通用运算的一种编程框架。我这么说一定是不好理解了。
       GPU通用运算的含义简单地说就是让显卡的GPU去帮助cpu去干事儿。通用,是相对于专用而言,GPU用于游戏加速,那就是专用处理器,在普通软件的运用中,GPU的计算资源一般是闲置的,而承担运算工作的是CPU,CPU就是典型的通用计算处理器。所以在游戏的时候cpu也要工作。这就是为什么一般游戏都有个cpu的要求。用gpu通用运算是有好处的。一般来说衡量运算能力使用浮点运算速度。(浮点运算简单地说就是 加减乘除)cpu要远远低于gpu 就拿I7 来说 i7 950的浮点运算能力基本上才相当于GTX580gpu的四十分之一。这么强的能力不用 那就浪费了。所以大家就开使动脑子打gpu的主意。 最开始ATI提出了stream框架 来调用GPU协助cpu工作 但是没有之后Nvida、推出的CUDA有名。当然除了这俩还有其它的框架。 stream和 CUDA只适用于 各自的品牌 还不能完全称得上通用。
      **OpenCL 和 DirectCompute 就是两个大家全能用的通用框架。**前者时apple所倡导的 后者是微软所开发的。这就是为什么 iso6系统的safari要比之前的都要快很多 这就是为什么ie10要比ie9快很多的原因 因为它们都可以调用了gpu来帮助运算。加快速度。前者被苹果低调的使用 后者被微软写入到DirectX11中 成为了标准。
      再说回CUDA,CUDA是Compute Unified Device Architecture的缩写。我也不大会翻译 就叫cuda好了 这东西 用他的好处就是能调用GPU的平行运算能力 和强大的浮点运算能力。在视频解码。或者处理游戏中物理特效的时候有非常好的效果(如果用cpu 去算现在游戏所运用的物理特效 那cpu就会累死)其实日常用处很少。也就游戏和视频编码用得上。
      cuda的用处还有很多 其实用在消费级的地方cuda很少大部分用在了特殊领域了。 很多超级计算中心 比如计算天气预告需要进行海量的数据计算 就应用得到gpu的运算能力了。

    CUDA(自我总结版)

    CUDA就是一个软件编程的框架,将很多底层通用的代码集成了,有了这个框架我们就可以减少编程的工作量,非常方便我们调用GPU,所以CUDA是一个方便我们编写代码调用GPU的框架。

    展开全文
  • CUDA安装和测试

    万次阅读 2018-03-04 15:45:11
    转载:http://blog.csdn.net/u012235003/article/details/54575758一、前提ubuntu系统:14.04 不要安装任何系统补丁和进行系统升级 ubantu14.04(这是比较稳定的版本,推荐使用) cuda7.5 GPU:GeForce 820M...

    转载:http://blog.csdn.net/u012235003/article/details/54575758

    一、前提

    ubuntu系统:14.04 
    不要安装任何系统补丁和进行系统升级 
    ubantu14.04(这是比较稳定的版本,推荐使用) 
    cuda7.5 
    GPU:GeForce 820M(只要是支持cuda的就好) 
    查看GPU是否支持cuda https://developer.nvidia.com/cuda-gpus

    二、推荐安装前最好先看一遍官方手册

    三、安装cuda前的准备工作

    1.检查自己的GPU是否是CUDA-capable 
    在终端中输入:

    lspci | grep -i nvidia
    • 1

    ,会显示自己的NVIDIA GPU版本信息 
    去CUDA的官网查看自己的GPU版本是否在CUDA的支持列表中 
    https://developer.nvidia.com/cuda-gpus 
    2.检查自己的Linux版本是否支持 CUDA 
    Ubuntu 14.04是没问题的 
    3.检查自己的系统中是否装了gcc 
    在终端中输入:

    gcc  --version
    • 1

    可以查看自己的gcc版本信息 
    如果你装的是cuda8.0,而且你的ubuntu是16.04,那么你的GCC版本就会是5.0以上, 
    cuda8.0是不支持GCC5.0以上的,所以需要降级。而且,ubuntu16.04只能安装cuda8.0

    4.检查是否安装了kernel header和 package development 
    在终端中输入:

    uname -r       
    • 1

    可以查看自己的kernel版本信息 
    在终端中输入:

    sudo apt-get install linux-headers-$(uname -r)
    • 1

    可以安装对应kernel版本的kernel header和package development 
    5.禁用 nouveau 
    终端中运行:

     lsmod | grep nouveau
    • 1

    ,如果有输出则代表nouveau正在加载。 
    Ubuntu的nouveau禁用方法: 
    在 /etc/modprobe.d 中创建文件 blacklist-nouveau.conf , 
    创建文件方法:进入到/etc/modprobe.d目录下 
    终端命令:

    sudo  touch  blacklist-nouveau.conf
    sudo  chmod a+w+r blacklist-nouveau.conf  (给文件可读可写的权限)
    vim blacklist-nouveau.conf
    • 1
    • 2
    • 3

    在文件中输入一下内容

    blacklist nouveau
    options nouveau modeset=0
    • 1
    • 2

    按esc 在按输入 :wq (注意前面的:也要输入)保存退出 
    打开终端,运行命令:

    sudo update-initramfs -u
    • 1

    重启ubuntu 
    设置完毕可以再次运行

     lsmod | grep nouveau
    • 1

    检查是否禁用成功,如果运行后没有任何输出,则代表禁用成功。

    6.下载cuda 
    https://developer.nvidia.com/cuda-downloads 
    按自己电脑要求选择对应版本。 
    这个cuda有两个安装方式:一个是runfile,另一个是deb。 
    这里我们选择第一个runfile下载。 
    下载后放在一个目录下,我这里为 /home/tony/cuda。

    7.安装cuda的一些依赖库 
    因为这个依赖的原因,导致我重装了好几次系统,缺少这些依赖会无法安装成功。

    sudo apt-get install freeglut3-dev build-essential libx11-dev libxmu-dev libxi-dev libgl1-mesa-glx libglu1-mesa libglu1-mesa-dev
    • 1

    如无法安装,请看cuda安装错误解决篇

    四、安装cuda

    1.重启电脑 
    在进入到登录界面时候,按住Ctrl+Alt+F1,进入到text mode,登录账号 
    2.关闭图形界面 
    终端命令:

    sudo service lightdm stop
    • 1

    3.切换到cuda文件目录 
    cd到下载好的cuda目录,例如我的cuda包在/home/tony/cuda中 
    再ls查看cuda名字 
    4.给cuda可执行的权限

     sudo chmod a+x cuda_7.5.18_linux.run
    • 1

    5.安装步骤

    sudo sh cuda_7.5.18_linux.run
    • 1

    (a)会先有个阅读声明,一直按D即可,然后accept 
    (b)第一个选项install nvidia accelerated Graphics Driver (y) 
    (c)第二个选项install the OpenGL libraries 
    双显卡(指的是集成显卡+独立显卡)选n,单卡(如果没有集成显卡,只有一个或多个支持GPU的显卡)选y。双显卡选y的话,会出现黑屏或登录界面无限循环的问题。 
    (d)后面的选项都是yes,或者按默认路径即直接按回车即可。 
    6.若出现下列显示 
    上面无警告和报错

    Driver :Installed
    Toolkit :Installed in /usr/local/cuda-7.5
    Samples :Installed in /home/tony
    • 1
    • 2
    • 3

    即安装暂时成功。

    7.输入

    sudo service lightdm start 
    • 1

    重新启动图形化界面 
    Alt + ctrl +F7,返回到图形化登录界面,输入密码登录。 
    如果能够成功登录,则表示不会遇到循环登录的问题,基本说明cuda的安装成功了。 
    8.搭配环境 
    终端输入

    echo 'export PATH=/usr/local/cuda-7.5/bin:$PATH' >> ~/.bashrc
    echo 'export LD_LIBRARY_PATH=/usr/local/cuda-7.5/lib64:$LD_LIBRARY_PATH' >> ~/.bashrc
    source ~/.bashrc
    • 1
    • 2
    • 3

    重启ubuntu

    五、检查cuda是否安装成功

    1.检查路径 ~/dev 下 有无存在名为 nvidia* (以nvidia开头)的多个文件(device files) 
    若无,安装错误,见解决篇。 
    2.检查 CUDA Toolkit是否安装成功 
    终端输入 :

    nvcc -V
    • 1

    会输出CUDA的版本信息(V要大写) 
    3.编译samples例子 
    进入到Samples安装目录,然后在该目录下终端输入make,等待十来分钟。 
    4.编译完成后测试 
    可以在Samples里面找到bin/x86_64/linux/release/目录,并切换到该目录 
    运行deviceQuery程序,sudo ./deviceQuery 
    查看输出结果,重点关注最后一行,Pass表示通过测试 
    这里写图片描述 
    运行bandwidthTest程序,sudo ./bandwidthTest 
    查看输出结果,显示结果为PASS表示通过测试 
    这里写图片描述 
    5.若完成上述步骤,cuda完成安装!


    展开全文
  • CUDA 入门教程

    万次阅读 多人点赞 2018-01-09 17:56:07
    CUDA从入门到精通(零):写在前面 在老板的要求下,本博主从2012年上高性能计算课程开始接触CUDA编程,随后将该技术应用到了实际项目中,使处理程序加速超过1K,可见基于图形显示器的并行计算对于追求速度的...

    CUDA从入门到精通(零):写在前面

    在老板的要求下,本博主从2012年上高性能计算课程开始接触CUDA编程,随后将该技术应用到了实际项目中,使处理程序加速超过1K,可见基于图形显示器的并行计算对于追求速度的应用来说无疑是一个理想的选择。还有不到一年毕业,怕是毕业后这些技术也就随毕业而去,准备这个暑假开辟一个CUDA专栏,从入门到精通,步步为营,顺便分享设计的一些经验教训,希望能给学习CUDA的童鞋提供一定指导。个人能力所及,错误难免,欢迎讨论。

     

    PS:申请专栏好像需要先发原创帖超过15篇。。。算了,先写够再申请吧,到时候一并转过去。


    CUDA从入门到精通(一):环境搭建

    NVIDIA于2006年推出CUDA(Compute Unified Devices Architecture),可以利用其推出的GPU进行通用计算,将并行计算从大型集群扩展到了普通显卡,使得用户只需要一台带有Geforce显卡的笔记本就能跑较大规模的并行处理程序。

     

    使用显卡的好处是,和大型集群相比功耗非常低,成本也不高,但性能很突出。以我的笔记本为例,Geforce 610M,用DeviceQuery程序测试,可得到如下硬件参数:

    计算能力达48X0.95 = 45.6 GFLOPS。而笔记本的CPU参数如下:

    CPU计算能力为(4核):2.5G*4 = 10GFLOPS,可见,显卡计算性能是4核i5 CPU的4~5倍,因此我们可以充分利用这一资源来对一些耗时的应用进行加速。

     

    好了,工欲善其事必先利其器,为了使用CUDA对GPU进行编程,我们需要准备以下必备工具:

    1. 硬件平台,就是显卡,如果你用的不是NVIDIA的显卡,那么只能说抱歉,其他都不支持CUDA。

    2. 操作系统,我用过windows XP,Windows 7都没问题,本博客用Windows7。

    3. C编译器,建议VS2008,和本博客一致。

    4. CUDA编译器NVCC,可以免费免注册免license从官网下载CUDA ToolkitCUDA下载,最新版本为5.0,本博客用的就是该版本。

    5. 其他工具(如Visual Assist,辅助代码高亮)

     

    准备完毕,开始安装软件。VS2008安装比较费时间,建议安装完整版(NVIDIA官网说Express版也可以),过程不必详述。CUDA Toolkit 5.0里面包含了NVCC编译器、设计文档、设计例程、CUDA运行时库、CUDA头文件等必备的原材料。

    安装完毕,我们在桌面上发现这个图标:

    不错,就是它,双击运行,可以看到一大堆例程。我们找到Simple OpenGL这个运行看看效果:

      点右边黄线标记处的Run即可看到美妙的三维正弦曲面,鼠标左键拖动可以转换角度,右键拖动可以缩放。如果这个运行成功,说明你的环境基本搭建成功。

    出现问题的可能:

    1. 你使用远程桌面连接登录到另一台服务器,该服务器上有显卡支持CUDA,但你远程终端不能运行CUDA程序。这是因为远程登录使用的是你本地显卡资源,在远程登录时看不到服务器端的显卡,所以会报错:没有支持CUDA的显卡!解决方法:1. 远程服务器装两块显卡,一块只用于显示,另一块用于计算;2.不要用图形界面登录,而是用命令行界面如telnet登录。

    2.有两个以上显卡都支持CUDA的情况,如何区分是在哪个显卡上运行?这个需要你在程序里控制,选择符合一定条件的显卡,如较高的时钟频率、较大的显存、较高的计算版本等。详细操作见后面的博客。

    好了,先说这么多,下一节我们介绍如何在VS2008中给GPU编程。

    CUDA从入门到精通(二):第一个CUDA程序

    书接上回,我们既然直接运行例程成功了,接下来就是了解如何实现例程中的每个环节。当然,我们先从简单的做起,一般编程语言都会找个helloworld例子,而我们的显卡是不会说话的,只能做一些简单的加减乘除运算。所以,CUDA程序的helloworld,我想应该最合适不过的就是向量加了。

    打开VS2008,选择File->New->Project,弹出下面对话框,设置如下:

    之后点OK,直接进入工程界面。

    工程中,我们看到只有一个.cu文件,内容如下:

    1. #include "cuda_runtime.h"  
    2. #include "device_launch_parameters.h"  
    3.   
    4. #include <stdio.h>  
    5.   
    6. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);  
    7.   
    8. __global__ void addKernel(int *c, const int *a, const int *b)  
    9. {  
    10.     int i = threadIdx.x;  
    11.     c[i] = a[i] + b[i];  
    12. }  
    13.   
    14. int main()  
    15. {  
    16.     const int arraySize = 5;  
    17.     const int a[arraySize] = { 1, 2, 3, 4, 5 };  
    18.     const int b[arraySize] = { 10, 20, 30, 40, 50 };  
    19.     int c[arraySize] = { 0 };  
    20.   
    21.     // Add vectors in parallel.  
    22.     cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);  
    23.     if (cudaStatus != cudaSuccess) {  
    24.         fprintf(stderr, "addWithCuda failed!");  
    25.         return 1;  
    26.     }  
    27.   
    28.     printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",  
    29.         c[0], c[1], c[2], c[3], c[4]);  
    30.   
    31.     // cudaThreadExit must be called before exiting in order for profiling and  
    32.     // tracing tools such as Nsight and Visual Profiler to show complete traces.  
    33.     cudaStatus = cudaThreadExit();  
    34.     if (cudaStatus != cudaSuccess) {  
    35.         fprintf(stderr, "cudaThreadExit failed!");  
    36.         return 1;  
    37.     }  
    38.   
    39.     return 0;  
    40. }  
    41.   
    42. // Helper function for using CUDA to add vectors in parallel.  
    43. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)  
    44. {  
    45.     int *dev_a = 0;  
    46.     int *dev_b = 0;  
    47.     int *dev_c = 0;  
    48.     cudaError_t cudaStatus;  
    49.   
    50.     // Choose which GPU to run on, change this on a multi-GPU system.  
    51.     cudaStatus = cudaSetDevice(0);  
    52.     if (cudaStatus != cudaSuccess) {  
    53.         fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
    54.         goto Error;  
    55.     }  
    56.   
    57.     // Allocate GPU buffers for three vectors (two input, one output)    .  
    58.     cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));  
    59.     if (cudaStatus != cudaSuccess) {  
    60.         fprintf(stderr, "cudaMalloc failed!");  
    61.         goto Error;  
    62.     }  
    63.   
    64.     cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));  
    65.     if (cudaStatus != cudaSuccess) {  
    66.         fprintf(stderr, "cudaMalloc failed!");  
    67.         goto Error;  
    68.     }  
    69.   
    70.     cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));  
    71.     if (cudaStatus != cudaSuccess) {  
    72.         fprintf(stderr, "cudaMalloc failed!");  
    73.         goto Error;  
    74.     }  
    75.   
    76.     // Copy input vectors from host memory to GPU buffers.  
    77.     cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);  
    78.     if (cudaStatus != cudaSuccess) {  
    79.         fprintf(stderr, "cudaMemcpy failed!");  
    80.         goto Error;  
    81.     }  
    82.   
    83.     cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);  
    84.     if (cudaStatus != cudaSuccess) {  
    85.         fprintf(stderr, "cudaMemcpy failed!");  
    86.         goto Error;  
    87.     }  
    88.   
    89.     // Launch a kernel on the GPU with one thread for each element.  
    90.     addKernel<<<1, size>>>(dev_c, dev_a, dev_b);  
    91.   
    92.     // cudaThreadSynchronize waits for the kernel to finish, and returns  
    93.     // any errors encountered during the launch.  
    94.     cudaStatus = cudaThreadSynchronize();  
    95.     if (cudaStatus != cudaSuccess) {  
    96.         fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);  
    97.         goto Error;  
    98.     }  
    99.   
    100.     // Copy output vector from GPU buffer to host memory.  
    101.     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);  
    102.     if (cudaStatus != cudaSuccess) {  
    103.         fprintf(stderr, "cudaMemcpy failed!");  
    104.         goto Error;  
    105.     }  
    106.   
    107. Error:  
    108.     cudaFree(dev_c);  
    109.     cudaFree(dev_a);  
    110.     cudaFree(dev_b);  
    111.       
    112.     return cudaStatus;  
    113. }  
    #include "cuda_runtime.h"
    #include "device_launch_parameters.h"
    
    #include <stdio.h>
    
    cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);
    
    __global__ void addKernel(int *c, const int *a, const int *b)
    {
        int i = threadIdx.x;
        c[i] = a[i] + b[i];
    }
    
    int main()
    {
        const int arraySize = 5;
        const int a[arraySize] = { 1, 2, 3, 4, 5 };
        const int b[arraySize] = { 10, 20, 30, 40, 50 };
        int c[arraySize] = { 0 };
    
        // Add vectors in parallel.
        cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);
        if (cudaStatus != cudaSuccess) {
            fprintf(stderr, "addWithCuda failed!");
            return 1;
        }
    
        printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",
            c[0], c[1], c[2], c[3], c[4]);
    
        // cudaThreadExit must be called before exiting in order for profiling and
        // tracing tools such as Nsight and Visual Profiler to show complete traces.
        cudaStatus = cudaThreadExit();
        if (cudaStatus != cudaSuccess) {
            fprintf(stderr, "cudaThreadExit failed!");
            return 1;
        }
    
        return 0;
    }
    
    // Helper function for using CUDA to add vectors in parallel.
    cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)
    {
        int *dev_a = 0;
        int *dev_b = 0;
        int *dev_c = 0;
        cudaError_t cudaStatus;
    
        // Choose which GPU to run on, change this on a multi-GPU system.
        cudaStatus = cudaSetDevice(0);
        if (cudaStatus != cudaSuccess) {
            fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
            goto Error;
        }
    
        // Allocate GPU buffers for three vectors (two input, one output)    .
        cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
        if (cudaStatus != cudaSuccess) {
            fprintf(stderr, "cudaMalloc failed!");
            goto Error;
        }
    
        cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
        if (cudaStatus != cudaSuccess) {
            fprintf(stderr, "cudaMalloc failed!");
            goto Error;
        }
    
        cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
        if (cudaStatus != cudaSuccess) {
            fprintf(stderr, "cudaMalloc failed!");
            goto Error;
        }
    
        // Copy input vectors from host memory to GPU buffers.
        cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
        if (cudaStatus != cudaSuccess) {
            fprintf(stderr, "cudaMemcpy failed!");
            goto Error;
        }
    
        cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
        if (cudaStatus != cudaSuccess) {
            fprintf(stderr, "cudaMemcpy failed!");
            goto Error;
        }
    
        // Launch a kernel on the GPU with one thread for each element.
        addKernel<<<1, size>>>(dev_c, dev_a, dev_b);
    
        // cudaThreadSynchronize waits for the kernel to finish, and returns
        // any errors encountered during the launch.
        cudaStatus = cudaThreadSynchronize();
        if (cudaStatus != cudaSuccess) {
            fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
            goto Error;
        }
    
        // Copy output vector from GPU buffer to host memory.
        cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
        if (cudaStatus != cudaSuccess) {
            fprintf(stderr, "cudaMemcpy failed!");
            goto Error;
        }
    
    Error:
        cudaFree(dev_c);
        cudaFree(dev_a);
        cudaFree(dev_b);
        
        return cudaStatus;
    }
     可以看出,CUDA程序和C程序并无区别,只是多了一些以"cuda"开头的一些库函数和一个特殊声明的函数:
    1. __global__ void addKernel(int *c, const int *a, const int *b)  
    2. {  
    3.     int i = threadIdx.x;  
    4.     c[i] = a[i] + b[i];  
    5. }  
    __global__ void addKernel(int *c, const int *a, const int *b)
    {
        int i = threadIdx.x;
        c[i] = a[i] + b[i];
    }

    这个函数就是在GPU上运行的函数,称之为核函数,英文名Kernel Function,注意要和操作系统内核函数区分开来。

    我们直接按F7编译,可以得到如下输出:

    1. 1>------ Build started: Project: cuda_helloworld, Configuration: Debug Win32 ------    
    2. 1>Compiling with CUDA Build Rule...    
    3. 1>"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\\bin\nvcc.exe"  -G   -gencode=arch=compute_10,code=\"sm_10,compute_10\" -gencode=arch=compute_20,code=\"sm_20,compute_20\"  --machine 32 -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\bin"    -Xcompiler "/EHsc /W3 /nologo /O2 /Zi   /MT  "  -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\\include" -maxrregcount=0   --compile -o "Debug/kernel.cu.obj" kernel.cu      
    4. 1>tmpxft_000000ec_00000000-8_kernel.compute_10.cudafe1.gpu    
    5. 1>tmpxft_000000ec_00000000-14_kernel.compute_10.cudafe2.gpu    
    6. 1>tmpxft_000000ec_00000000-5_kernel.compute_20.cudafe1.gpu    
    7. 1>tmpxft_000000ec_00000000-17_kernel.compute_20.cudafe2.gpu    
    8. 1>kernel.cu    
    9. 1>kernel.cu    
    10. 1>tmpxft_000000ec_00000000-8_kernel.compute_10.cudafe1.cpp    
    11. 1>tmpxft_000000ec_00000000-24_kernel.compute_10.ii    
    12. 1>Linking...    
    13. 1>Embedding manifest...    
    14. 1>Performing Post-Build Event...    
    15. 1>copy "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\\bin\cudart*.dll" "C:\Users\DongXiaoman\Documents\Visual Studio 2008\Projects\cuda_helloworld\Debug"    
    16. 1>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\\bin\cudart32_50_35.dll    
    17. 1>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\\bin\cudart64_50_35.dll    
    18. 1>已复制         2 个文件。    
    19. 1>Build log was saved at "file://c:\Users\DongXiaoman\Documents\Visual Studio 2008\Projects\cuda_helloworld\cuda_helloworld\Debug\BuildLog.htm"    
    20. 1>cuda_helloworld - 0 error(s), 105 warning(s)    
    21. ========== Build: 1 succeeded, 0 failed, 0 up-to-date, 0 skipped ==========    
    1>------ Build started: Project: cuda_helloworld, Configuration: Debug Win32 ------  
    1>Compiling with CUDA Build Rule...  
    1>"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\\bin\nvcc.exe"  -G   -gencode=arch=compute_10,code=\"sm_10,compute_10\" -gencode=arch=compute_20,code=\"sm_20,compute_20\"  --machine 32 -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\bin"    -Xcompiler "/EHsc /W3 /nologo /O2 /Zi   /MT  "  -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\\include" -maxrregcount=0   --compile -o "Debug/kernel.cu.obj" kernel.cu    
    1>tmpxft_000000ec_00000000-8_kernel.compute_10.cudafe1.gpu  
    1>tmpxft_000000ec_00000000-14_kernel.compute_10.cudafe2.gpu  
    1>tmpxft_000000ec_00000000-5_kernel.compute_20.cudafe1.gpu  
    1>tmpxft_000000ec_00000000-17_kernel.compute_20.cudafe2.gpu  
    1>kernel.cu  
    1>kernel.cu  
    1>tmpxft_000000ec_00000000-8_kernel.compute_10.cudafe1.cpp  
    1>tmpxft_000000ec_00000000-24_kernel.compute_10.ii  
    1>Linking...  
    1>Embedding manifest...  
    1>Performing Post-Build Event...  
    1>copy "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\\bin\cudart*.dll" "C:\Users\DongXiaoman\Documents\Visual Studio 2008\Projects\cuda_helloworld\Debug"  
    1>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\\bin\cudart32_50_35.dll  
    1>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\\bin\cudart64_50_35.dll  
    1>已复制         2 个文件。  
    1>Build log was saved at "file://c:\Users\DongXiaoman\Documents\Visual Studio 2008\Projects\cuda_helloworld\cuda_helloworld\Debug\BuildLog.htm"  
    1>cuda_helloworld - 0 error(s), 105 warning(s)  
    ========== Build: 1 succeeded, 0 failed, 0 up-to-date, 0 skipped ==========  

    可见,编译.cu文件需要利用nvcc工具。该工具的详细使用见后面博客。

    直接运行,可以得到结果图如下:

    如果显示正确,那么我们的第一个程序宣告成功!

    刚入门CUDA,跑过几个官方提供的例程,看了看人家的代码,觉得并不难,但自己动手写代码时,总是不知道要先干什么,后干什么,也不知道从哪个知识点学起。这时就需要有一本能提供指导的书籍或者教程,一步步跟着做下去,直到真正掌握。

    一般讲述CUDA的书,我认为不错的有下面这几本:

    初学者可以先看美国人写的这本《GPU高性能编程CUDA实战》,可操作性很强,但不要期望能全看懂(Ps:里面有些概念其实我现在还是不怎么懂),但不影响你进一步学习。如果想更全面地学习CUDA,《GPGPU编程技术》比较客观详细地介绍了通用GPU编程的策略,看过这本书,可以对显卡有更深入的了解,揭开GPU的神秘面纱。后面《OpenGL编程指南》完全是为了体验图形交互带来的乐趣,可以有选择地看;《GPU高性能运算之CUDA》这本是师兄给的,适合快速查询(感觉是将官方编程手册翻译了一遍)一些关键技术和概念。

    有了这些指导材料还不够,我们在做项目的时候,遇到的问题在这些书上肯定找不到,所以还需要有下面这些利器:

    这里面有很多工具的使用手册,如CUDA_GDB,Nsight,CUDA_Profiler等,方便调试程序;还有一些有用的库,如CUFFT是专门用来做快速傅里叶变换的,CUBLAS是专用于线性代数(矩阵、向量计算)的,CUSPASE是专用于稀疏矩阵表示和计算的库。这些库的使用可以降低我们设计算法的难度,提高开发效率。另外还有些入门教程也是值得一读的,你会对NVCC编译器有更近距离的接触。

    好了,前言就这么多,本博主计划按如下顺序来讲述CUDA:

    1.了解设备

    2.线程并行

    3.块并行

    4.流并行

    5.线程通信

    6.线程通信实例:规约

    7.存储模型

    8.常数内存

    9.纹理内存

    10.主机页锁定内存

    11.图形互操作

    12.优化准则

    13.CUDA与MATLAB接口

    14.CUDA与MFC接口

    前面三节已经对CUDA做了一个简单的介绍,这一节开始真正进入编程环节。

    首先,初学者应该对自己使用的设备有较为扎实的理解和掌握,这样对后面学习并行程序优化很有帮助,了解硬件详细参数可以通过上节介绍的几本书和官方资料获得,但如果仍然觉得不够直观,那么我们可以自己动手获得这些内容。

    以第二节例程为模板,我们稍加改动的部分代码如下:

    1. // Add vectors in parallel.  
    2. cudaError_t cudaStatus;  
    3. int num = 0;  
    4. cudaDeviceProp prop;  
    5. cudaStatus = cudaGetDeviceCount(&num);  
    6. for(int i = 0;i<num;i++)  
    7. {  
    8.     cudaGetDeviceProperties(&prop,i);  
    9. }  
    10. cudaStatus = addWithCuda(c, a, b, arraySize);  
    // Add vectors in parallel.
    cudaError_t cudaStatus;
    int num = 0;
    cudaDeviceProp prop;
    cudaStatus = cudaGetDeviceCount(&num);
    for(int i = 0;i<num;i++)
    {
    	cudaGetDeviceProperties(&prop,i);
    }
    cudaStatus = addWithCuda(c, a, b, arraySize);

    这个改动的目的是让我们的程序自动通过调用cuda API函数获得设备数目和属性,所谓“知己知彼,百战不殆”。

    cudaError_t 是cuda错误类型,取值为整数。

    cudaDeviceProp为设备属性结构体,其定义可以从cuda Toolkit安装目录中找到,我的路径为:C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\include\driver_types.h,找到定义为:

    1. /** 
    2.  * CUDA device properties 
    3.  */  
    4. struct __device_builtin__ cudaDeviceProp  
    5. {  
    6.     char   name[256];                  /**< ASCII string identifying device */  
    7.     size_t totalGlobalMem;             /**< Global memory available on device in bytes */  
    8.     size_t sharedMemPerBlock;          /**< Shared memory available per block in bytes */  
    9.     int    regsPerBlock;               /**< 32-bit registers available per block */  
    10.     int    warpSize;                   /**< Warp size in threads */  
    11.     size_t memPitch;                   /**< Maximum pitch in bytes allowed by memory copies */  
    12.     int    maxThreadsPerBlock;         /**< Maximum number of threads per block */  
    13.     int    maxThreadsDim[3];           /**< Maximum size of each dimension of a block */  
    14.     int    maxGridSize[3];             /**< Maximum size of each dimension of a grid */  
    15.     int    clockRate;                  /**< Clock frequency in kilohertz */  
    16.     size_t totalConstMem;              /**< Constant memory available on device in bytes */  
    17.     int    major;                      /**< Major compute capability */  
    18.     int    minor;                      /**< Minor compute capability */  
    19.     size_t textureAlignment;           /**< Alignment requirement for textures */  
    20.     size_t texturePitchAlignment;      /**< Pitch alignment requirement for texture references bound to pitched memory */  
    21.     int    deviceOverlap;              /**< Device can concurrently copy memory and execute a kernel. Deprecated. Use instead asyncEngineCount. */  
    22.     int    multiProcessorCount;        /**< Number of multiprocessors on device */  
    23.     int    kernelExecTimeoutEnabled;   /**< Specified whether there is a run time limit on kernels */  
    24.     int    integrated;                 /**< Device is integrated as opposed to discrete */  
    25.     int    canMapHostMemory;           /**< Device can map host memory with cudaHostAlloc/cudaHostGetDevicePointer */  
    26.     int    computeMode;                /**< Compute mode (See ::cudaComputeMode) */  
    27.     int    maxTexture1D;               /**< Maximum 1D texture size */  
    28.     int    maxTexture1DMipmap;         /**< Maximum 1D mipmapped texture size */  
    29.     int    maxTexture1DLinear;         /**< Maximum size for 1D textures bound to linear memory */  
    30.     int    maxTexture2D[2];            /**< Maximum 2D texture dimensions */  
    31.     int    maxTexture2DMipmap[2];      /**< Maximum 2D mipmapped texture dimensions */  
    32.     int    maxTexture2DLinear[3];      /**< Maximum dimensions (width, height, pitch) for 2D textures bound to pitched memory */  
    33.     int    maxTexture2DGather[2];      /**< Maximum 2D texture dimensions if texture gather operations have to be performed */  
    34.     int    maxTexture3D[3];            /**< Maximum 3D texture dimensions */  
    35.     int    maxTextureCubemap;          /**< Maximum Cubemap texture dimensions */  
    36.     int    maxTexture1DLayered[2];     /**< Maximum 1D layered texture dimensions */  
    37.     int    maxTexture2DLayered[3];     /**< Maximum 2D layered texture dimensions */  
    38.     int    maxTextureCubemapLayered[2];/**< Maximum Cubemap layered texture dimensions */  
    39.     int    maxSurface1D;               /**< Maximum 1D surface size */  
    40.     int    maxSurface2D[2];            /**< Maximum 2D surface dimensions */  
    41.     int    maxSurface3D[3];            /**< Maximum 3D surface dimensions */  
    42.     int    maxSurface1DLayered[2];     /**< Maximum 1D layered surface dimensions */  
    43.     int    maxSurface2DLayered[3];     /**< Maximum 2D layered surface dimensions */  
    44.     int    maxSurfaceCubemap;          /**< Maximum Cubemap surface dimensions */  
    45.     int    maxSurfaceCubemapLayered[2];/**< Maximum Cubemap layered surface dimensions */  
    46.     size_t surfaceAlignment;           /**< Alignment requirements for surfaces */  
    47.     int    concurrentKernels;          /**< Device can possibly execute multiple kernels concurrently */  
    48.     int    ECCEnabled;                 /**< Device has ECC support enabled */  
    49.     int    pciBusID;                   /**< PCI bus ID of the device */  
    50.     int    pciDeviceID;                /**< PCI device ID of the device */  
    51.     int    pciDomainID;                /**< PCI domain ID of the device */  
    52.     int    tccDriver;                  /**< 1 if device is a Tesla device using TCC driver, 0 otherwise */  
    53.     int    asyncEngineCount;           /**< Number of asynchronous engines */  
    54.     int    unifiedAddressing;          /**< Device shares a unified address space with the host */  
    55.     int    memoryClockRate;            /**< Peak memory clock frequency in kilohertz */  
    56.     int    memoryBusWidth;             /**< Global memory bus width in bits */  
    57.     int    l2CacheSize;                /**< Size of L2 cache in bytes */  
    58.     int    maxThreadsPerMultiProcessor;/**< Maximum resident threads per multiprocessor */  
    59. };  
    /**
     * CUDA device properties
     */
    struct __device_builtin__ cudaDeviceProp
    {
        char   name[256];                  /**< ASCII string identifying device */
        size_t totalGlobalMem;             /**< Global memory available on device in bytes */
        size_t sharedMemPerBlock;          /**< Shared memory available per block in bytes */
        int    regsPerBlock;               /**< 32-bit registers available per block */
        int    warpSize;                   /**< Warp size in threads */
        size_t memPitch;                   /**< Maximum pitch in bytes allowed by memory copies */
        int    maxThreadsPerBlock;         /**< Maximum number of threads per block */
        int    maxThreadsDim[3];           /**< Maximum size of each dimension of a block */
        int    maxGridSize[3];             /**< Maximum size of each dimension of a grid */
        int    clockRate;                  /**< Clock frequency in kilohertz */
        size_t totalConstMem;              /**< Constant memory available on device in bytes */
        int    major;                      /**< Major compute capability */
        int    minor;                      /**< Minor compute capability */
        size_t textureAlignment;           /**< Alignment requirement for textures */
        size_t texturePitchAlignment;      /**< Pitch alignment requirement for texture references bound to pitched memory */
        int    deviceOverlap;              /**< Device can concurrently copy memory and execute a kernel. Deprecated. Use instead asyncEngineCount. */
        int    multiProcessorCount;        /**< Number of multiprocessors on device */
        int    kernelExecTimeoutEnabled;   /**< Specified whether there is a run time limit on kernels */
        int    integrated;                 /**< Device is integrated as opposed to discrete */
        int    canMapHostMemory;           /**< Device can map host memory with cudaHostAlloc/cudaHostGetDevicePointer */
        int    computeMode;                /**< Compute mode (See ::cudaComputeMode) */
        int    maxTexture1D;               /**< Maximum 1D texture size */
        int    maxTexture1DMipmap;         /**< Maximum 1D mipmapped texture size */
        int    maxTexture1DLinear;         /**< Maximum size for 1D textures bound to linear memory */
        int    maxTexture2D[2];            /**< Maximum 2D texture dimensions */
        int    maxTexture2DMipmap[2];      /**< Maximum 2D mipmapped texture dimensions */
        int    maxTexture2DLinear[3];      /**< Maximum dimensions (width, height, pitch) for 2D textures bound to pitched memory */
        int    maxTexture2DGather[2];      /**< Maximum 2D texture dimensions if texture gather operations have to be performed */
        int    maxTexture3D[3];            /**< Maximum 3D texture dimensions */
        int    maxTextureCubemap;          /**< Maximum Cubemap texture dimensions */
        int    maxTexture1DLayered[2];     /**< Maximum 1D layered texture dimensions */
        int    maxTexture2DLayered[3];     /**< Maximum 2D layered texture dimensions */
        int    maxTextureCubemapLayered[2];/**< Maximum Cubemap layered texture dimensions */
        int    maxSurface1D;               /**< Maximum 1D surface size */
        int    maxSurface2D[2];            /**< Maximum 2D surface dimensions */
        int    maxSurface3D[3];            /**< Maximum 3D surface dimensions */
        int    maxSurface1DLayered[2];     /**< Maximum 1D layered surface dimensions */
        int    maxSurface2DLayered[3];     /**< Maximum 2D layered surface dimensions */
        int    maxSurfaceCubemap;          /**< Maximum Cubemap surface dimensions */
        int    maxSurfaceCubemapLayered[2];/**< Maximum Cubemap layered surface dimensions */
        size_t surfaceAlignment;           /**< Alignment requirements for surfaces */
        int    concurrentKernels;          /**< Device can possibly execute multiple kernels concurrently */
        int    ECCEnabled;                 /**< Device has ECC support enabled */
        int    pciBusID;                   /**< PCI bus ID of the device */
        int    pciDeviceID;                /**< PCI device ID of the device */
        int    pciDomainID;                /**< PCI domain ID of the device */
        int    tccDriver;                  /**< 1 if device is a Tesla device using TCC driver, 0 otherwise */
        int    asyncEngineCount;           /**< Number of asynchronous engines */
        int    unifiedAddressing;          /**< Device shares a unified address space with the host */
        int    memoryClockRate;            /**< Peak memory clock frequency in kilohertz */
        int    memoryBusWidth;             /**< Global memory bus width in bits */
        int    l2CacheSize;                /**< Size of L2 cache in bytes */
        int    maxThreadsPerMultiProcessor;/**< Maximum resident threads per multiprocessor */
    };

    后面的注释已经说明了其字段代表意义,可能有些术语对于初学者理解起来还是有一定困难,没关系,我们现在只需要关注以下几个指标:

    name:就是设备名称;

    totalGlobalMem:就是显存大小;

    major,minor:CUDA设备版本号,有1.1, 1.2, 1.3, 2.0, 2.1等多个版本;

    clockRate:GPU时钟频率;

    multiProcessorCount:GPU大核数,一个大核(专业点称为流多处理器,SM,Stream-Multiprocessor)包含多个小核(流处理器,SP,Stream-Processor)

    编译,运行,我们在VS2008工程的cudaGetDeviceProperties()函数处放一个断点,单步执行这一函数,然后用Watch窗口,切换到Auto页,展开+,在我的笔记本上得到如下结果:

    可以看到,设备名为GeForce 610M,显存1GB,设备版本2.1(比较高端了,哈哈),时钟频率为950MHz(注意950000单位为kHz),大核数为1。在一些高性能GPU上(如Tesla,Kepler系列),大核数可能达到几十甚至上百,可以做更大规模的并行处理。

    PS:今天看SDK代码时发现在help_cuda.h中有个函数实现从CUDA设备版本查询相应大核中小核的数目,觉得很有用,以后编程序可以借鉴,摘抄如下:

    1. // Beginning of GPU Architecture definitions  
    2. inline int _ConvertSMVer2Cores(int major, int minor)  
    3. {  
    4.     // Defines for GPU Architecture types (using the SM version to determine the # of cores per SM  
    5.     typedef struct  
    6.     {  
    7.         int SM; // 0xMm (hexidecimal notation), M = SM Major version, and m = SM minor version  
    8.         int Cores;  
    9.     } sSMtoCores;  
    10.   
    11.     sSMtoCores nGpuArchCoresPerSM[] =  
    12.     {  
    13.         { 0x10,  8 }, // Tesla Generation (SM 1.0) G80 class  
    14.         { 0x11,  8 }, // Tesla Generation (SM 1.1) G8x class  
    15.         { 0x12,  8 }, // Tesla Generation (SM 1.2) G9x class  
    16.         { 0x13,  8 }, // Tesla Generation (SM 1.3) GT200 class  
    17.         { 0x20, 32 }, // Fermi Generation (SM 2.0) GF100 class  
    18.         { 0x21, 48 }, // Fermi Generation (SM 2.1) GF10x class  
    19.         { 0x30, 192}, // Kepler Generation (SM 3.0) GK10x class  
    20.         { 0x35, 192}, // Kepler Generation (SM 3.5) GK11x class  
    21.         {   -1, -1 }  
    22.     };  
    23.   
    24.     int index = 0;  
    25.   
    26.     while (nGpuArchCoresPerSM[index].SM != -1)  
    27.     {  
    28.         if (nGpuArchCoresPerSM[index].SM == ((major << 4) + minor))  
    29.         {  
    30.             return nGpuArchCoresPerSM[index].Cores;  
    31.         }  
    32.   
    33.         index++;  
    34.     }  
    35.   
    36.     // If we don't find the values, we default use the previous one to run properly  
    37.     printf("MapSMtoCores for SM %d.%d is undefined.  Default to use %d Cores/SM\n", major, minor, nGpuArchCoresPerSM[7].Cores);  
    38.     return nGpuArchCoresPerSM[7].Cores;  
    39. }  
    40. // end of GPU Architecture definitions  
    // Beginning of GPU Architecture definitions
    inline int _ConvertSMVer2Cores(int major, int minor)
    {
        // Defines for GPU Architecture types (using the SM version to determine the # of cores per SM
        typedef struct
        {
            int SM; // 0xMm (hexidecimal notation), M = SM Major version, and m = SM minor version
            int Cores;
        } sSMtoCores;
    
        sSMtoCores nGpuArchCoresPerSM[] =
        {
            { 0x10,  8 }, // Tesla Generation (SM 1.0) G80 class
            { 0x11,  8 }, // Tesla Generation (SM 1.1) G8x class
            { 0x12,  8 }, // Tesla Generation (SM 1.2) G9x class
            { 0x13,  8 }, // Tesla Generation (SM 1.3) GT200 class
            { 0x20, 32 }, // Fermi Generation (SM 2.0) GF100 class
            { 0x21, 48 }, // Fermi Generation (SM 2.1) GF10x class
            { 0x30, 192}, // Kepler Generation (SM 3.0) GK10x class
            { 0x35, 192}, // Kepler Generation (SM 3.5) GK11x class
            {   -1, -1 }
        };
    
        int index = 0;
    
        while (nGpuArchCoresPerSM[index].SM != -1)
        {
            if (nGpuArchCoresPerSM[index].SM == ((major << 4) + minor))
            {
                return nGpuArchCoresPerSM[index].Cores;
            }
    
            index++;
        }
    
        // If we don't find the values, we default use the previous one to run properly
        printf("MapSMtoCores for SM %d.%d is undefined.  Default to use %d Cores/SM\n", major, minor, nGpuArchCoresPerSM[7].Cores);
        return nGpuArchCoresPerSM[7].Cores;
    }
    // end of GPU Architecture definitions

    可见,设备版本2.1的一个大核有48个小核,而版本3.0以上的一个大核有192个小核!

    前文说到过,当我们用的电脑上有多个显卡支持CUDA时,怎么来区分在哪个上运行呢?这里我们看一下addWithCuda这个函数是怎么做的。

    1. cudaError_t cudaStatus;  
    2.   
    3. // Choose which GPU to run on, change this on a multi-GPU system.  
    4. cudaStatus = cudaSetDevice(0);  
    5. if (cudaStatus != cudaSuccess) {  
    6.     fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
    7.     goto Error;  
    8. }  
    cudaError_t cudaStatus;
    
    // Choose which GPU to run on, change this on a multi-GPU system.
    cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        goto Error;
    }

    使用了cudaSetDevice(0)这个操作,0表示能搜索到的第一个设备号,如果有多个设备,则编号为0,1,2...。

    再看我们本节添加的代码,有个函数cudaGetDeviceCount(&num),这个函数用来获取设备总数,这样我们选择运行CUDA程序的设备号取值就是0,1,...num-1,于是可以一个个枚举设备,利用cudaGetDeviceProperties(&prop)获得其属性,然后利用一定排序、筛选算法,找到最符合我们应用的那个设备号opt,然后调用cudaSetDevice(opt)即可选择该设备。选择标准可以从处理能力、版本控制、名称等各个角度出发。后面讲述流并发过程时,还要用到这些API。

    如果希望了解更多硬件内容可以结合http://www.geforce.cn/hardware获取。

    多线程我们应该都不陌生,在操作系统中,进程是资源分配的基本单元,而线程是CPU时间调度的基本单元(这里假设只有1个CPU)。

    将线程的概念引申到CUDA程序设计中,我们可以认为线程就是执行CUDA程序的最小单元,前面我们建立的工程代码中,有个核函数概念不知各位童鞋还记得没有,在GPU上每个线程都会运行一次该核函数。

    但GPU上的线程调度方式与CPU有很大不同。CPU上会有优先级分配,从高到低,同样优先级的可以采用时间片轮转法实现线程调度。GPU上线程没有优先级概念,所有线程机会均等,线程状态只有等待资源和执行两种状态,如果资源未就绪,那么就等待;一旦就绪,立即执行。当GPU资源很充裕时,所有线程都是并发执行的,这样加速效果很接近理论加速比;而GPU资源少于总线程个数时,有一部分线程就会等待前面执行的线程释放资源,从而变为串行化执行。

    代码还是用上一节的吧,改动很少,再贴一遍:

    1. #include "cuda_runtime.h"           //CUDA运行时API  
    2. #include "device_launch_parameters.h"     
    3. #include <stdio.h>  
    4. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);  
    5. __global__ void addKernel(int *c, const int *a, const int *b)  
    6. {  
    7.     int i = threadIdx.x;  
    8.     c[i] = a[i] + b[i];  
    9. }  
    10. int main()  
    11. {  
    12.     const int arraySize = 5;  
    13.     const int a[arraySize] = { 1, 2, 3, 4, 5 };  
    14.     const int b[arraySize] = { 10, 20, 30, 40, 50 };  
    15.     int c[arraySize] = { 0 };  
    16.     // Add vectors in parallel.  
    17.     cudaError_t cudaStatus;  
    18.     int num = 0;  
    19.     cudaDeviceProp prop;  
    20.     cudaStatus = cudaGetDeviceCount(&num);  
    21.     for(int i = 0;i<num;i++)  
    22.     {  
    23.         cudaGetDeviceProperties(&prop,i);  
    24.     }  
    25.     cudaStatus = addWithCuda(c, a, b, arraySize);  
    26.     if (cudaStatus != cudaSuccess)   
    27.     {  
    28.         fprintf(stderr, "addWithCuda failed!");  
    29.         return 1;  
    30.     }  
    31.     printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",c[0],c[1],c[2],c[3],c[4]);  
    32.     // cudaThreadExit must be called before exiting in order for profiling and  
    33.     // tracing tools such as Nsight and Visual Profiler to show complete traces.  
    34.     cudaStatus = cudaThreadExit();  
    35.     if (cudaStatus != cudaSuccess)   
    36.     {  
    37.         fprintf(stderr, "cudaThreadExit failed!");  
    38.         return 1;  
    39.     }  
    40.     return 0;  
    41. }  
    42. // 重点理解这个函数  
    43. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)  
    44. {  
    45.     int *dev_a = 0; //GPU设备端数据指针  
    46.     int *dev_b = 0;  
    47.     int *dev_c = 0;  
    48.     cudaError_t cudaStatus;     //状态指示  
    49.   
    50.     // Choose which GPU to run on, change this on a multi-GPU system.  
    51.     cudaStatus = cudaSetDevice(0);  //选择运行平台  
    52.     if (cudaStatus != cudaSuccess)   
    53.     {  
    54.         fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
    55.         goto Error;  
    56.     }  
    57.     // 分配GPU设备端内存  
    58.     cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));  
    59.     if (cudaStatus != cudaSuccess)   
    60.     {  
    61.         fprintf(stderr, "cudaMalloc failed!");  
    62.         goto Error;  
    63.     }  
    64.     cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));  
    65.     if (cudaStatus != cudaSuccess)   
    66.     {  
    67.         fprintf(stderr, "cudaMalloc failed!");  
    68.         goto Error;  
    69.     }  
    70.     cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));  
    71.     if (cudaStatus != cudaSuccess)   
    72.     {  
    73.         fprintf(stderr, "cudaMalloc failed!");  
    74.         goto Error;  
    75.     }  
    76.     // 拷贝数据到GPU  
    77.     cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);  
    78.     if (cudaStatus != cudaSuccess)   
    79.     {  
    80.         fprintf(stderr, "cudaMemcpy failed!");  
    81.         goto Error;  
    82.     }  
    83.     cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);  
    84.     if (cudaStatus != cudaSuccess)   
    85.     {  
    86.         fprintf(stderr, "cudaMemcpy failed!");  
    87.         goto Error;  
    88.     }  
    89.     // 运行核函数  
    90. <span style="BACKGROUND-COLOR: #ff6666"><strong>    addKernel<<<1, size>>>(dev_c, dev_a, dev_b);</strong>  
    91. </span>    // cudaThreadSynchronize waits for the kernel to finish, and returns  
    92.     // any errors encountered during the launch.  
    93.     cudaStatus = cudaThreadSynchronize();   //同步线程  
    94.     if (cudaStatus != cudaSuccess)   
    95.     {  
    96.         fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);  
    97.         goto Error;  
    98.     }  
    99.     // Copy output vector from GPU buffer to host memory.  
    100.     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);      //拷贝结果回主机  
    101.     if (cudaStatus != cudaSuccess)   
    102.     {  
    103.         fprintf(stderr, "cudaMemcpy failed!");  
    104.         goto Error;  
    105.     }  
    106. Error:  
    107.     cudaFree(dev_c);    //释放GPU设备端内存  
    108.     cudaFree(dev_a);  
    109.     cudaFree(dev_b);      
    110.     return cudaStatus;  
    111. }  
    #include "cuda_runtime.h"			//CUDA运行时API
    #include "device_launch_parameters.h"	
    #include <stdio.h>
    cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);
    __global__ void addKernel(int *c, const int *a, const int *b)
    {
        int i = threadIdx.x;
        c[i] = a[i] + b[i];
    }
    int main()
    {
        const int arraySize = 5;
        const int a[arraySize] = { 1, 2, 3, 4, 5 };
        const int b[arraySize] = { 10, 20, 30, 40, 50 };
        int c[arraySize] = { 0 };
        // Add vectors in parallel.
        cudaError_t cudaStatus;
    	int num = 0;
    	cudaDeviceProp prop;
    	cudaStatus = cudaGetDeviceCount(&num);
    	for(int i = 0;i<num;i++)
    	{
    		cudaGetDeviceProperties(&prop,i);
    	}
    	cudaStatus = addWithCuda(c, a, b, arraySize);
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "addWithCuda failed!");
            return 1;
        }
        printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",c[0],c[1],c[2],c[3],c[4]);
        // cudaThreadExit must be called before exiting in order for profiling and
        // tracing tools such as Nsight and Visual Profiler to show complete traces.
        cudaStatus = cudaThreadExit();
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaThreadExit failed!");
            return 1;
        }
        return 0;
    }
    // 重点理解这个函数
    cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)
    {
        int *dev_a = 0;	//GPU设备端数据指针
        int *dev_b = 0;
        int *dev_c = 0;
        cudaError_t cudaStatus;		//状态指示
    
        // Choose which GPU to run on, change this on a multi-GPU system.
        cudaStatus = cudaSetDevice(0);	//选择运行平台
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
            goto Error;
        }
        // 分配GPU设备端内存
        cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMalloc failed!");
            goto Error;
        }
        cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMalloc failed!");
            goto Error;
        }
        cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMalloc failed!");
            goto Error;
        }
        // 拷贝数据到GPU
        cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMemcpy failed!");
            goto Error;
        }
        cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMemcpy failed!");
            goto Error;
        }
        // 运行核函数
    <span style="BACKGROUND-COLOR: #ff6666"><strong>    addKernel<<<1, size>>>(dev_c, dev_a, dev_b);</strong>
    </span>    // cudaThreadSynchronize waits for the kernel to finish, and returns
        // any errors encountered during the launch.
        cudaStatus = cudaThreadSynchronize();	//同步线程
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
            goto Error;
        }
        // Copy output vector from GPU buffer to host memory.
        cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);		//拷贝结果回主机
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMemcpy failed!");
            goto Error;
        }
    Error:
        cudaFree(dev_c);	//释放GPU设备端内存
        cudaFree(dev_a);
        cudaFree(dev_b);    
        return cudaStatus;
    }

    红色部分即启动核函数的调用过程,这里看到调用方式和C不太一样。<<<>>>表示运行时配置符号,里面1表示只分配一个线程组(又称线程块、Block),size表示每个线程组有size个线程(Thread)。本程序中size根据前面传递参数个数应该为5,所以运行的时候,核函数在5个GPU线程单元上分别运行了一次,总共运行了5次。这5个线程是如何知道自己“身份”的?是靠threadIdx这个内置变量,它是个dim3类型变量,接受<<<>>>中第二个参数,它包含x,y,z 3维坐标,而我们传入的参数只有一维,所以只有x值是有效的。通过核函数中int i = threadIdx.x;这一句,每个线程可以获得自身的id号,从而找到自己的任务去执行。

    CUDA从入门到精通(六):块并行

     

    同一版本的代码用了这么多次,有点过意不去,于是这次我要做较大的改动大笑,大家要擦亮眼睛,拭目以待。

    块并行相当于操作系统中多进程的情况,上节说到,CUDA有线程组(线程块)的概念,将一组线程组织到一起,共同分配一部分资源,然后内部调度执行。线程块与线程块之间,毫无瓜葛。这有利于做更粗粒度的并行。我们将上一节的代码改为块并行版本如下:

    下节我们介绍块并行。

    1. #include "cuda_runtime.h"  
    2. #include "device_launch_parameters.h"  
    3. #include <stdio.h>  
    4. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);  
    5. __global__ void addKernel(int *c, const int *a, const int *b)  
    6. {  
    7. <span style="BACKGROUND-COLOR: #ff0000">    int i = blockIdx.x;  
    8. </span>    c[i] = a[i] + b[i];  
    9. }  
    10. int main()  
    11. {  
    12.     const int arraySize = 5;  
    13.     const int a[arraySize] = { 1, 2, 3, 4, 5 };  
    14.     const int b[arraySize] = { 10, 20, 30, 40, 50 };  
    15.     int c[arraySize] = { 0 };  
    16.     // Add vectors in parallel.  
    17.     cudaError_t cudaStatus;  
    18.     int num = 0;  
    19.     cudaDeviceProp prop;  
    20.     cudaStatus = cudaGetDeviceCount(&num);  
    21.     for(int i = 0;i<num;i++)  
    22.     {  
    23.         cudaGetDeviceProperties(&prop,i);  
    24.     }  
    25.     cudaStatus = addWithCuda(c, a, b, arraySize);  
    26.     if (cudaStatus != cudaSuccess)   
    27.     {  
    28.         fprintf(stderr, "addWithCuda failed!");  
    29.         return 1;  
    30.     }  
    31.     printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",c[0],c[1],c[2],c[3],c[4]);  
    32.     // cudaThreadExit must be called before exiting in order for profiling and  
    33.     // tracing tools such as Nsight and Visual Profiler to show complete traces.  
    34.     cudaStatus = cudaThreadExit();  
    35.     if (cudaStatus != cudaSuccess)   
    36.     {  
    37.         fprintf(stderr, "cudaThreadExit failed!");  
    38.         return 1;  
    39.     }  
    40.     return 0;  
    41. }  
    42. // Helper function for using CUDA to add vectors in parallel.  
    43. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)  
    44. {  
    45.     int *dev_a = 0;  
    46.     int *dev_b = 0;  
    47.     int *dev_c = 0;  
    48.     cudaError_t cudaStatus;  
    49.   
    50.     // Choose which GPU to run on, change this on a multi-GPU system.  
    51.     cudaStatus = cudaSetDevice(0);  
    52.     if (cudaStatus != cudaSuccess)   
    53.     {  
    54.         fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
    55.         goto Error;  
    56.     }  
    57.     // Allocate GPU buffers for three vectors (two input, one output)    .  
    58.     cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));  
    59.     if (cudaStatus != cudaSuccess)   
    60.     {  
    61.         fprintf(stderr, "cudaMalloc failed!");  
    62.         goto Error;  
    63.     }  
    64.     cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));  
    65.     if (cudaStatus != cudaSuccess)   
    66.     {  
    67.         fprintf(stderr, "cudaMalloc failed!");  
    68.         goto Error;  
    69.     }  
    70.     cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));  
    71.     if (cudaStatus != cudaSuccess)   
    72.     {  
    73.         fprintf(stderr, "cudaMalloc failed!");  
    74.         goto Error;  
    75.     }  
    76.     // Copy input vectors from host memory to GPU buffers.  
    77.     cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);  
    78.     if (cudaStatus != cudaSuccess)   
    79.     {  
    80.         fprintf(stderr, "cudaMemcpy failed!");  
    81.         goto Error;  
    82.     }  
    83.     cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);  
    84.     if (cudaStatus != cudaSuccess)   
    85.     {  
    86.         fprintf(stderr, "cudaMemcpy failed!");  
    87.         goto Error;  
    88.     }  
    89.     // Launch a kernel on the GPU with one thread for each element.  
    90.  <span style="BACKGROUND-COLOR: #ff0000">   addKernel<<<size,1 >>>(dev_c, dev_a, dev_b);  
    91. </span>    // cudaThreadSynchronize waits for the kernel to finish, and returns  
    92.     // any errors encountered during the launch.  
    93.     cudaStatus = cudaThreadSynchronize();  
    94.     if (cudaStatus != cudaSuccess)   
    95.     {  
    96.         fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);  
    97.         goto Error;  
    98.     }  
    99.     // Copy output vector from GPU buffer to host memory.  
    100.     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);  
    101.     if (cudaStatus != cudaSuccess)   
    102.     {  
    103.         fprintf(stderr, "cudaMemcpy failed!");  
    104.         goto Error;  
    105.     }  
    106. Error:  
    107.     cudaFree(dev_c);  
    108.     cudaFree(dev_a);  
    109.     cudaFree(dev_b);      
    110.     return cudaStatus;  
    111. }  
    #include "cuda_runtime.h"
    #include "device_launch_parameters.h"
    #include <stdio.h>
    cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);
    __global__ void addKernel(int *c, const int *a, const int *b)
    {
    <span style="BACKGROUND-COLOR: #ff0000">    int i = blockIdx.x;
    </span>    c[i] = a[i] + b[i];
    }
    int main()
    {
        const int arraySize = 5;
        const int a[arraySize] = { 1, 2, 3, 4, 5 };
        const int b[arraySize] = { 10, 20, 30, 40, 50 };
        int c[arraySize] = { 0 };
        // Add vectors in parallel.
        cudaError_t cudaStatus;
    	int num = 0;
    	cudaDeviceProp prop;
    	cudaStatus = cudaGetDeviceCount(&num);
    	for(int i = 0;i<num;i++)
    	{
    		cudaGetDeviceProperties(&prop,i);
    	}
    	cudaStatus = addWithCuda(c, a, b, arraySize);
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "addWithCuda failed!");
            return 1;
        }
        printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",c[0],c[1],c[2],c[3],c[4]);
        // cudaThreadExit must be called before exiting in order for profiling and
        // tracing tools such as Nsight and Visual Profiler to show complete traces.
        cudaStatus = cudaThreadExit();
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaThreadExit failed!");
            return 1;
        }
        return 0;
    }
    // Helper function for using CUDA to add vectors in parallel.
    cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)
    {
        int *dev_a = 0;
        int *dev_b = 0;
        int *dev_c = 0;
        cudaError_t cudaStatus;
    
        // Choose which GPU to run on, change this on a multi-GPU system.
        cudaStatus = cudaSetDevice(0);
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
            goto Error;
        }
        // Allocate GPU buffers for three vectors (two input, one output)    .
        cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMalloc failed!");
            goto Error;
        }
        cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMalloc failed!");
            goto Error;
        }
        cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMalloc failed!");
            goto Error;
        }
        // Copy input vectors from host memory to GPU buffers.
        cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMemcpy failed!");
            goto Error;
        }
        cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMemcpy failed!");
            goto Error;
        }
        // Launch a kernel on the GPU with one thread for each element.
     <span style="BACKGROUND-COLOR: #ff0000">   addKernel<<<size,1 >>>(dev_c, dev_a, dev_b);
    </span>    // cudaThreadSynchronize waits for the kernel to finish, and returns
        // any errors encountered during the launch.
        cudaStatus = cudaThreadSynchronize();
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
            goto Error;
        }
        // Copy output vector from GPU buffer to host memory.
        cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMemcpy failed!");
            goto Error;
        }
    Error:
        cudaFree(dev_c);
        cudaFree(dev_a);
        cudaFree(dev_b);    
        return cudaStatus;
    }

    和上一节相比,只有这两行有改变,<<<>>>里第一个参数改成了size,第二个改成了1,表示我们分配size个线程块,每个线程块仅包含1个线程,总共还是有5个线程。这5个线程相互独立,执行核函数得到相应的结果,与上一节不同的是,每个线程获取id的方式变为int i = blockIdx.x;这是线程块ID。

    于是有童鞋提问了,线程并行和块并行的区别在哪里?

    线程并行是细粒度并行,调度效率高;块并行是粗粒度并行,每次调度都要重新分配资源,有时资源只有一份,那么所有线程块都只能排成一队,串行执行。

    那是不是我们所有时候都应该用线程并行,尽可能不用块并行?

    当然不是,我们的任务有时可以采用分治法,将一个大问题分解为几个小规模问题,将这些小规模问题分别用一个线程块实现,线程块内可以采用细粒度的线程并行,而块之间为粗粒度并行,这样可以充分利用硬件资源,降低线程并行的计算复杂度。适当分解,降低规模,在一些矩阵乘法、向量内积计算应用中可以得到充分的展示。

    实际应用中,常常是二者的结合。线程块、线程组织图如下所示。

    多个线程块组织成了一个Grid,称为线程格(经历了从一位线程,二维线程块到三维线程格的过程,立体感很强啊)。

    好了,下一节我们介绍流并行,是更高层次的并行。

    前面我们没有讲程序的结构,我想有些童鞋可能迫不及待想知道CUDA程序到底是怎么一个执行过程。好的,这一节在介绍流之前,先把CUDA程序结构简要说一下。

    CUDA程序文件后缀为.cu,有些编译器可能不认识这个后缀的文件,我们可以在VS2008的Tools->Options->Text Editor->File Extension里添加cu后缀到VC++中,如下图:

    一个.cu文件内既包含CPU程序(称为主机程序),也包含GPU程序(称为设备程序)。如何区分主机程序和设备程序?根据声明,凡是挂有“__global__”或者“__device__”前缀的函数,都是在GPU上运行的设备程序,不同的是__global__设备程序可被主机程序调用,而__device__设备程序则只能被设备程序调用。

    没有挂任何前缀的函数,都是主机程序。主机程序显示声明可以用__host__前缀。设备程序需要由NVCC进行编译,而主机程序只需要由主机编译器(如VS2008中的cl.exe,Linux上的GCC)。主机程序主要完成设备环境初始化,数据传输等必备过程,设备程序只负责计算。

    主机程序中,有一些“cuda”打头的函数,这些都是CUDA Runtime API,即运行时函数,主要负责完成设备的初始化、内存分配、内存拷贝等任务。我们前面第三节用到的函数cudaGetDeviceCount(),cudaGetDeviceProperties(),cudaSetDevice()都是运行时API。这些函数的具体参数声明我们不必一一记下来,拿出第三节的官方利器就可以轻松查询,让我们打开这个文件:

    打开后,在pdf搜索栏中输入一个运行时函数,例如cudaMemcpy,查到的结果如下:

    可以看到,该API函数的参数形式为,第一个表示目的地,第二个表示来源地,第三个参数表示字节数,第四个表示类型。如果对类型不了解,直接点击超链接,得到详细解释如下:

    可见,该API可以实现从主机到主机、主机到设备、设备到主机、设备到设备的内存拷贝过程。同时可以发现,利用该API手册可以很方便地查询我们需要用的这些API函数,所以以后编CUDA程序一定要把它打开,随时准备查询,这样可以大大提高编程效率。

    好了,进入今天的主题:流并行。

    前面已经介绍了线程并行和块并行,知道了线程并行为细粒度的并行,而块并行为粗粒度的并行,同时也知道了CUDA的线程组织情况,即Grid-Block-Thread结构。一组线程并行处理可以组织为一个block,而一组block并行处理可以组织为一个Grid,很自然地想到,Grid只是一个网格,我们是否可以利用多个网格来完成并行处理呢?答案就是利用流。

    流可以实现在一个设备上运行多个核函数。前面的块并行也好,线程并行也好,运行的核函数都是相同的(代码一样,传递参数也一样)。而流并行,可以执行不同的核函数,也可以实现对同一个核函数传递不同的参数,实现任务级别的并行。

    CUDA中的流用cudaStream_t类型实现,用到的API有以下几个:cudaStreamCreate(cudaStream_t * s)用于创建流,cudaStreamDestroy(cudaStream_t s)用于销毁流,cudaStreamSynchronize()用于单个流同步,cudaDeviceSynchronize()用于整个设备上的所有流同步,cudaStreamQuery()用于查询一个流的任务是否已经完成。具体的含义可以查询API手册。

    下面我们将前面的两个例子中的任务改用流实现,仍然是{1,2,3,4,5}+{10,20,30,40,50} = {11,22,33,44,55}这个例子。代码如下:

    1. #include "cuda_runtime.h"  
    2. #include "device_launch_parameters.h"  
    3. #include <stdio.h>  
    4. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);  
    5. __global__ void addKernel(int *c, const int *a, const int *b)  
    6. {  
    7.     int i = blockIdx.x;  
    8.     c[i] = a[i] + b[i];  
    9. }  
    10. int main()  
    11. {  
    12.     const int arraySize = 5;  
    13.     const int a[arraySize] = { 1, 2, 3, 4, 5 };  
    14.     const int b[arraySize] = { 10, 20, 30, 40, 50 };  
    15.     int c[arraySize] = { 0 };  
    16.     // Add vectors in parallel.  
    17.     cudaError_t cudaStatus;  
    18.     int num = 0;  
    19.     cudaDeviceProp prop;  
    20.     cudaStatus = cudaGetDeviceCount(&num);  
    21.     for(int i = 0;i<num;i++)  
    22.     {  
    23.         cudaGetDeviceProperties(&prop,i);  
    24.     }  
    25.     cudaStatus = addWithCuda(c, a, b, arraySize);  
    26.     if (cudaStatus != cudaSuccess)   
    27.     {  
    28.         fprintf(stderr, "addWithCuda failed!");  
    29.         return 1;  
    30.     }  
    31.     printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",c[0],c[1],c[2],c[3],c[4]);  
    32.     // cudaThreadExit must be called before exiting in order for profiling and  
    33.     // tracing tools such as Nsight and Visual Profiler to show complete traces.  
    34.     cudaStatus = cudaThreadExit();  
    35.     if (cudaStatus != cudaSuccess)   
    36.     {  
    37.         fprintf(stderr, "cudaThreadExit failed!");  
    38.         return 1;  
    39.     }  
    40.     return 0;  
    41. }  
    42. // Helper function for using CUDA to add vectors in parallel.  
    43. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)  
    44. {  
    45.     int *dev_a = 0;  
    46.     int *dev_b = 0;  
    47.     int *dev_c = 0;  
    48.     cudaError_t cudaStatus;  
    49.   
    50.     // Choose which GPU to run on, change this on a multi-GPU system.  
    51.     cudaStatus = cudaSetDevice(0);  
    52.     if (cudaStatus != cudaSuccess)   
    53.     {  
    54.         fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
    55.         goto Error;  
    56.     }  
    57.     // Allocate GPU buffers for three vectors (two input, one output)    .  
    58.     cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));  
    59.     if (cudaStatus != cudaSuccess)   
    60.     {  
    61.         fprintf(stderr, "cudaMalloc failed!");  
    62.         goto Error;  
    63.     }  
    64.     cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));  
    65.     if (cudaStatus != cudaSuccess)   
    66.     {  
    67.         fprintf(stderr, "cudaMalloc failed!");  
    68.         goto Error;  
    69.     }  
    70.     cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));  
    71.     if (cudaStatus != cudaSuccess)   
    72.     {  
    73.         fprintf(stderr, "cudaMalloc failed!");  
    74.         goto Error;  
    75.     }  
    76.     // Copy input vectors from host memory to GPU buffers.  
    77.     cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);  
    78.     if (cudaStatus != cudaSuccess)   
    79.     {  
    80.         fprintf(stderr, "cudaMemcpy failed!");  
    81.         goto Error;  
    82.     }  
    83.     cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);  
    84.     if (cudaStatus != cudaSuccess)   
    85.     {  
    86.         fprintf(stderr, "cudaMemcpy failed!");  
    87.         goto Error;  
    88.     }  
    89. <span style="BACKGROUND-COLOR: #ff6666">  cudaStream_t stream[5];  
    90.     for(int i = 0;i<5;i++)  
    91.     {  
    92.         cudaStreamCreate(&stream[i]);   //创建流  
    93.     }  
    94. </span>    // Launch a kernel on the GPU with one thread for each element.  
    95. <span style="BACKGROUND-COLOR: #ff6666">  for(int i = 0;i<5;i++)  
    96.     {  
    97.         addKernel<<<1,1,0,stream[i]>>>(dev_c+i, dev_a+i, dev_b+i);    //执行流  
    98.     }  
    99.     cudaDeviceSynchronize();  
    100. </span>    // cudaThreadSynchronize waits for the kernel to finish, and returns  
    101.     // any errors encountered during the launch.  
    102.     cudaStatus = cudaThreadSynchronize();  
    103.     if (cudaStatus != cudaSuccess)   
    104.     {  
    105.         fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);  
    106.         goto Error;  
    107.     }  
    108.     // Copy output vector from GPU buffer to host memory.  
    109.     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);  
    110.     if (cudaStatus != cudaSuccess)   
    111.     {  
    112.         fprintf(stderr, "cudaMemcpy failed!");  
    113.         goto Error;  
    114.     }  
    115. Error:  
    116. <span style="BACKGROUND-COLOR: #ff6666">  for(int i = 0;i<5;i++)  
    117.     {  
    118.         cudaStreamDestroy(stream[i]);   //销毁流  
    119.     }  
    120. </span>    cudaFree(dev_c);  
    121.     cudaFree(dev_a);  
    122.     cudaFree(dev_b);      
    123.     return cudaStatus;  
    124. }  
    #include "cuda_runtime.h"
    #include "device_launch_parameters.h"
    #include <stdio.h>
    cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);
    __global__ void addKernel(int *c, const int *a, const int *b)
    {
        int i = blockIdx.x;
        c[i] = a[i] + b[i];
    }
    int main()
    {
        const int arraySize = 5;
        const int a[arraySize] = { 1, 2, 3, 4, 5 };
        const int b[arraySize] = { 10, 20, 30, 40, 50 };
        int c[arraySize] = { 0 };
        // Add vectors in parallel.
        cudaError_t cudaStatus;
    	int num = 0;
    	cudaDeviceProp prop;
    	cudaStatus = cudaGetDeviceCount(&num);
    	for(int i = 0;i<num;i++)
    	{
    		cudaGetDeviceProperties(&prop,i);
    	}
    	cudaStatus = addWithCuda(c, a, b, arraySize);
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "addWithCuda failed!");
            return 1;
        }
        printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",c[0],c[1],c[2],c[3],c[4]);
        // cudaThreadExit must be called before exiting in order for profiling and
        // tracing tools such as Nsight and Visual Profiler to show complete traces.
        cudaStatus = cudaThreadExit();
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaThreadExit failed!");
            return 1;
        }
        return 0;
    }
    // Helper function for using CUDA to add vectors in parallel.
    cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)
    {
        int *dev_a = 0;
        int *dev_b = 0;
        int *dev_c = 0;
        cudaError_t cudaStatus;
    
        // Choose which GPU to run on, change this on a multi-GPU system.
        cudaStatus = cudaSetDevice(0);
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
            goto Error;
        }
        // Allocate GPU buffers for three vectors (two input, one output)    .
        cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMalloc failed!");
            goto Error;
        }
        cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMalloc failed!");
            goto Error;
        }
        cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMalloc failed!");
            goto Error;
        }
        // Copy input vectors from host memory to GPU buffers.
        cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMemcpy failed!");
            goto Error;
        }
        cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMemcpy failed!");
            goto Error;
        }
    <span style="BACKGROUND-COLOR: #ff6666">	cudaStream_t stream[5];
    	for(int i = 0;i<5;i++)
    	{
    		cudaStreamCreate(&stream[i]);	//创建流
    	}
    </span>    // Launch a kernel on the GPU with one thread for each element.
    <span style="BACKGROUND-COLOR: #ff6666">	for(int i = 0;i<5;i++)
    	{
    		addKernel<<<1,1,0,stream[i]>>>(dev_c+i, dev_a+i, dev_b+i);	//执行流
    	}
    	cudaDeviceSynchronize();
    </span>    // cudaThreadSynchronize waits for the kernel to finish, and returns
        // any errors encountered during the launch.
        cudaStatus = cudaThreadSynchronize();
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
            goto Error;
        }
        // Copy output vector from GPU buffer to host memory.
        cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMemcpy failed!");
            goto Error;
        }
    Error:
    <span style="BACKGROUND-COLOR: #ff6666">	for(int i = 0;i<5;i++)
    	{
    		cudaStreamDestroy(stream[i]);	//销毁流
    	}
    </span>    cudaFree(dev_c);
        cudaFree(dev_a);
        cudaFree(dev_b);    
        return cudaStatus;
    }

    注意到,我们的核函数代码仍然和块并行的版本一样,只是在调用时做了改变,<<<>>>中的参数多了两个,其中前两个和块并行、线程并行中的意义相同,仍然是线程块数(这里为1)、每个线程块中线程数(这里也是1)。第三个为0表示每个block用到的共享内存大小,这个我们后面再讲;第四个为流对象,表示当前核函数在哪个流上运行。我们创建了5个流,每个流上都装载了一个核函数,同时传递参数有些不同,也就是每个核函数作用的对象也不同。这样就实现了任务级别的并行,当我们有几个互不相关的任务时,可以写多个核函数,资源允许的情况下,我们将这些核函数装载到不同流上,然后执行,这样可以实现更粗粒度的并行。

    好了,流并行就这么简单,我们处理任务时,可以根据需要,选择最适合的并行方式。

    我们前面几节主要介绍了三种利用GPU实现并行处理的方式:线程并行,块并行和流并行。在这些方法中,我们一再强调,各个线程所进行的处理是互不相关的,即两个线程不回产生交集,每个线程都只关注自己的一亩三分地,对其他线程毫无兴趣,就当不存在。。。。

    当然,实际应用中,这样的例子太少了,也就是遇到向量相加、向量对应点乘这类才会有如此高的并行度,而其他一些应用,如一组数求和,求最大(小)值,各个线程不再是相互独立的,而是产生一定关联,线程2可能会用到线程1的结果,这时就需要利用本节的线程通信技术了。

    线程通信在CUDA中有三种实现方式:

    1. 共享存储器;

    2. 线程 同步;

    3. 原子操作;

    最常用的是前两种方式,共享存储器,术语Shared Memory,是位于SM中的特殊存储器。还记得SM吗,就是流多处理器,大核是也。一个SM中不仅包含若干个SP(流处理器,小核),还包括一部分高速Cache,寄存器组,共享内存等,结构如图所示:

    从图中可看出,一个SM内有M个SP,Shared Memory由这M个SP共同占有。另外指令单元也被这M个SP共享,即SIMT架构(单指令多线程架构),一个SM中所有SP在同一时间执行同一代码。

    为了实现线程通信,仅仅靠共享内存还不够,需要有同步机制才能使线程之间实现有序处理。通常情况是这样:当线程A需要线程B计算的结果作为输入时,需要确保线程B已经将结果写入共享内存中,然后线程A再从共享内存中读出。同步必不可少,否则,线程A可能读到的是无效的结果,造成计算错误。同步机制可以用CUDA内置函数:__syncthreads();当某个线程执行到该函数时,进入等待状态,直到同一线程块(Block)中所有线程都执行到这个函数为止,即一个__syncthreads()相当于一个线程同步点,确保一个Block中所有线程都达到同步,然后线程进入运行状态。

    综上两点,我们可以写一段线程通信的伪代码如下:

    1. //Begin  
    2. if this is thread B  
    3.      write something to Shared Memory;  
    4. end if  
    5. __syncthreads();  
    6. if this is thread A  
    7.     read something from Shared Memory;  
    8. end if  
    9. //End  
    //Begin
    if this is thread B
         write something to Shared Memory;
    end if
    __syncthreads();
    if this is thread A
        read something from Shared Memory;
    end if
    //End

    上面代码在CUDA中实现时,由于SIMT特性,所有线程都执行同样的代码,所以在线程中需要判断自己的身份,以免误操作。

    注意的是,位于同一个Block中的线程才能实现通信,不同Block中的线程不能通过共享内存、同步进行通信,而应采用原子操作或主机介入。

    对于原子操作,如果感兴趣可以翻阅《GPU高性能编程CUDA实战》第九章“原子性”。

    本节完。下节我们给出一个实例来看线程通信的代码怎么设计。

    接着上一节,我们利用刚学到的共享内存和线程同步技术,来做一个简单的例子。先看下效果吧:

    很简单,就是分别求出1~5这5个数字的和,平方和,连乘积。相信学过C语言的童鞋都能用for循环做出同上面一样的效果,但为了学习CUDA共享内存和同步技术,我们还是要把简单的东西复杂化(^_^)。

    简要分析一下,上面例子的输入都是一样的,1,2,3,4,5这5个数,但计算过程有些变化,而且每个输出和所有输入都相关,不是前几节例子中那样,一个输出只和一个输入有关。所以我们在利用CUDA编程时,需要针对特殊问题做些让步,把一些步骤串行化实现。

    输入数据原本位于主机内存,通过cudaMemcpy API已经拷贝到GPU显存(术语为全局存储器,Global Memory),每个线程运行时需要从Global Memory读取输入数据,然后完成计算,最后将结果写回Global Memory。当我们计算需要多次相同输入数据时,大家可能想到,每次都分别去Global Memory读数据好像有点浪费,如果数据很大,那么反复多次读数据会相当耗时间。索性我们把它从Global Memory一次性读到SM内部,然后在内部进行处理,这样可以节省反复读取的时间。

    有了这个思路,结合上节看到的SM结构图,看到有一片存储器叫做Shared Memory,它位于SM内部,处理时访问速度相当快(差不多每个时钟周期读一次),而全局存储器读一次需要耗费几十甚至上百个时钟周期。于是,我们就制定A计划如下:

    线程块数:1,块号为0;(只有一个线程块内的线程才能进行通信,所以我们只分配一个线程块,具体工作交给每个线程完成)

    线程数:5,线程号分别为0~4;(线程并行,前面讲过)

    共享存储器大小:5个int型变量大小(5 * sizeof(int))。

    步骤一:读取输入数据。将Global Memory中的5个整数读入共享存储器,位置一一对应,和线程号也一一对应,所以可以同时完成。

    步骤二:线程同步,确保所有线程都完成了工作。

    步骤三:指定线程,对共享存储器中的输入数据完成相应处理。

    代码如下:

    1. #include "cuda_runtime.h"  
    2. #include "device_launch_parameters.h"  
    3.   
    4. #include <stdio.h>  
    5.   
    6. cudaError_t addWithCuda(int *c, const int *a, size_t size);  
    7.   
    8. __global__ void addKernel(int *c, const int *a)  
    9. {  
    10.     int i = threadIdx.x;  
    11. <span style="font-size:24px;"><strong>  extern __shared__ int smem[];</strong>  
    12. </span>   smem[i] = a[i];  
    13.     __syncthreads();  
    14.     if(i == 0)  // 0号线程做平方和  
    15.     {  
    16.         c[0] = 0;  
    17.         for(int d = 0; d < 5; d++)  
    18.         {  
    19.             c[0] += smem[d] * smem[d];  
    20.         }  
    21.     }  
    22.     if(i == 1)//1号线程做累加  
    23.     {  
    24.         c[1] = 0;  
    25.         for(int d = 0; d < 5; d++)  
    26.         {  
    27.             c[1] += smem[d];  
    28.         }  
    29.     }  
    30.     if(i == 2)  //2号线程做累乘  
    31.     {  
    32.         c[2] = 1;  
    33.         for(int d = 0; d < 5; d++)  
    34.         {  
    35.             c[2] *= smem[d];  
    36.         }  
    37.     }  
    38. }  
    39.   
    40. int main()  
    41. {  
    42.     const int arraySize = 5;  
    43.     const int a[arraySize] = { 1, 2, 3, 4, 5 };  
    44.     int c[arraySize] = { 0 };  
    45.     // Add vectors in parallel.  
    46.     cudaError_t cudaStatus = addWithCuda(c, a, arraySize);  
    47.     if (cudaStatus != cudaSuccess)   
    48.     {  
    49.         fprintf(stderr, "addWithCuda failed!");  
    50.         return 1;  
    51.     }  
    52.     printf("\t1+2+3+4+5 = %d\n\t1^2+2^2+3^2+4^2+5^2 = %d\n\t1*2*3*4*5 = %d\n\n\n\n\n\n", c[1], c[0], c[2]);  
    53.     // cudaThreadExit must be called before exiting in order for profiling and  
    54.     // tracing tools such as Nsight and Visual Profiler to show complete traces.  
    55.     cudaStatus = cudaThreadExit();  
    56.     if (cudaStatus != cudaSuccess)   
    57.     {  
    58.         fprintf(stderr, "cudaThreadExit failed!");  
    59.         return 1;  
    60.     }  
    61.     return 0;  
    62. }  
    63.   
    64. // Helper function for using CUDA to add vectors in parallel.  
    65. cudaError_t addWithCuda(int *c, const int *a,  size_t size)  
    66. {  
    67.     int *dev_a = 0;  
    68.     int *dev_c = 0;  
    69.     cudaError_t cudaStatus;  
    70.   
    71.     // Choose which GPU to run on, change this on a multi-GPU system.  
    72.     cudaStatus = cudaSetDevice(0);  
    73.     if (cudaStatus != cudaSuccess)   
    74.     {  
    75.         fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
    76.         goto Error;  
    77.     }  
    78.   
    79.     // Allocate GPU buffers for three vectors (two input, one output)    .  
    80.     cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));  
    81.     if (cudaStatus != cudaSuccess)   
    82.     {  
    83.         fprintf(stderr, "cudaMalloc failed!");  
    84.         goto Error;  
    85.     }  
    86.   
    87.     cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));  
    88.     if (cudaStatus != cudaSuccess)   
    89.     {  
    90.         fprintf(stderr, "cudaMalloc failed!");  
    91.         goto Error;  
    92.     }  
    93.     // Copy input vectors from host memory to GPU buffers.  
    94.     cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);  
    95.     if (cudaStatus != cudaSuccess)   
    96.     {  
    97.         fprintf(stderr, "cudaMemcpy failed!");  
    98.         goto Error;  
    99.     }  
    100.     // Launch a kernel on the GPU with one thread for each element.  
    101. <span style="font-size:24px;"><strong>    addKernel<<<1, size, size * sizeof(int), 0>>>(dev_c, dev_a);</strong>  
    102. </span>  
    103.     // cudaThreadSynchronize waits for the kernel to finish, and returns  
    104.     // any errors encountered during the launch.  
    105.     cudaStatus = cudaThreadSynchronize();  
    106.     if (cudaStatus != cudaSuccess)   
    107.     {  
    108.         fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);  
    109.         goto Error;  
    110.     }  
    111.   
    112.     // Copy output vector from GPU buffer to host memory.  
    113.     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);  
    114.     if (cudaStatus != cudaSuccess)   
    115.     {  
    116.         fprintf(stderr, "cudaMemcpy failed!");  
    117.         goto Error;  
    118.     }  
    119.   
    120. Error:  
    121.     cudaFree(dev_c);  
    122.     cudaFree(dev_a);      
    123.     return cudaStatus;  
    124. }  
    #include "cuda_runtime.h"
    #include "device_launch_parameters.h"
    
    #include <stdio.h>
    
    cudaError_t addWithCuda(int *c, const int *a, size_t size);
    
    __global__ void addKernel(int *c, const int *a)
    {
        int i = threadIdx.x;
    <span style="font-size:24px;"><strong>	extern __shared__ int smem[];</strong>
    </span>	smem[i] = a[i];
    	__syncthreads();
    	if(i == 0)	// 0号线程做平方和
    	{
    		c[0] = 0;
    		for(int d = 0; d < 5; d++)
    		{
    			c[0] += smem[d] * smem[d];
    		}
    	}
    	if(i == 1)//1号线程做累加
    	{
    		c[1] = 0;
    		for(int d = 0; d < 5; d++)
    		{
    			c[1] += smem[d];
    		}
    	}
    	if(i == 2)	//2号线程做累乘
    	{
    		c[2] = 1;
    		for(int d = 0; d < 5; d++)
    		{
    			c[2] *= smem[d];
    		}
    	}
    }
    
    int main()
    {
        const int arraySize = 5;
        const int a[arraySize] = { 1, 2, 3, 4, 5 };
        int c[arraySize] = { 0 };
        // Add vectors in parallel.
        cudaError_t cudaStatus = addWithCuda(c, a, arraySize);
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "addWithCuda failed!");
            return 1;
        }
    	printf("\t1+2+3+4+5 = %d\n\t1^2+2^2+3^2+4^2+5^2 = %d\n\t1*2*3*4*5 = %d\n\n\n\n\n\n", c[1], c[0], c[2]);
        // cudaThreadExit must be called before exiting in order for profiling and
        // tracing tools such as Nsight and Visual Profiler to show complete traces.
        cudaStatus = cudaThreadExit();
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaThreadExit failed!");
            return 1;
        }
        return 0;
    }
    
    // Helper function for using CUDA to add vectors in parallel.
    cudaError_t addWithCuda(int *c, const int *a,  size_t size)
    {
        int *dev_a = 0;
        int *dev_c = 0;
        cudaError_t cudaStatus;
    
        // Choose which GPU to run on, change this on a multi-GPU system.
        cudaStatus = cudaSetDevice(0);
        if (cudaStatus != cudaSuccess) 
        {
            fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
            goto Error;
        }
    
        // Allocate GPU buffers for three vectors (two input, one output)    .
        cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
        if (cudaStatus != cudaSuccess) 
        {
            fprintf(stderr, "cudaMalloc failed!");
            goto Error;
        }
    
        cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
        if (cudaStatus != cudaSuccess) 
        {
            fprintf(stderr, "cudaMalloc failed!");
            goto Error;
        }
        // Copy input vectors from host memory to GPU buffers.
        cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
        if (cudaStatus != cudaSuccess) 
        {
            fprintf(stderr, "cudaMemcpy failed!");
            goto Error;
        }
        // Launch a kernel on the GPU with one thread for each element.
    <span style="font-size:24px;"><strong>    addKernel<<<1, size, size * sizeof(int), 0>>>(dev_c, dev_a);</strong>
    </span>
        // cudaThreadSynchronize waits for the kernel to finish, and returns
        // any errors encountered during the launch.
        cudaStatus = cudaThreadSynchronize();
        if (cudaStatus != cudaSuccess) 
        {
            fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
            goto Error;
        }
    
        // Copy output vector from GPU buffer to host memory.
        cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
        if (cudaStatus != cudaSuccess) 
        {
            fprintf(stderr, "cudaMemcpy failed!");
            goto Error;
        }
    
    Error:
        cudaFree(dev_c);
        cudaFree(dev_a);    
        return cudaStatus;
    }
    从代码中看到执行配置<<<>>>中第三个参数为共享内存大小(字节数),这样我们就知道了全部4个执行配置参数的意义。恭喜,你的CUDA终于入门了!

    入门后的进一步学习的内容,就是如何优化自己的代码。我们前面的例子没有考虑任何性能方面优化,是为了更好地学习基本知识点,而不是其他细节问题。从本节开始,我们要从性能出发考虑问题,不断优化代码,使执行速度提高是并行处理的唯一目的。

    测试代码运行速度有很多方法,C语言里提供了类似于SystemTime()这样的API获得系统时间,然后计算两个事件之间的时长从而完成计时功能。在CUDA中,我们有专门测量设备运行时间的API,下面一一介绍。

    翻开编程手册《CUDA_Toolkit_Reference_Manual》,随时准备查询不懂得API。我们在运行核函数前后,做如下操作:

    1. cudaEvent_t start, stop;<span style="white-space: pre;">    </span>//事件对象  
    2. cudaEventCreate(&start);<span style="white-space: pre;">    </span>//创建事件  
    3. cudaEventCreate(&stop);<span style="white-space: pre;">     </span>//创建事件  
    4. cudaEventRecord(start, stream);<span style="white-space: pre;"> </span>//记录开始  
    5. myKernel<<<dimg,dimb,size_smem,stream>>>(parameter list);//执行核函数  
    6.   
    7. cudaEventRecord(stop,stream);<span style="white-space: pre;">   </span>//记录结束事件  
    8. cudaEventSynchronize(stop);<span style="white-space: pre;"> </span>//事件同步,等待结束事件之前的设备操作均已完成  
    9. float elapsedTime;  
    10. cudaEventElapsedTime(&elapsedTime,start,stop);//计算两个事件之间时长(单位为ms)  
    cudaEvent_t start, stop;	//事件对象
    cudaEventCreate(&start);	//创建事件
    cudaEventCreate(&stop);		//创建事件
    cudaEventRecord(start, stream);	//记录开始
    myKernel<<<dimg,dimb,size_smem,stream>>>(parameter list);//执行核函数
    
    cudaEventRecord(stop,stream);	//记录结束事件
    cudaEventSynchronize(stop);	//事件同步,等待结束事件之前的设备操作均已完成
    float elapsedTime;
    cudaEventElapsedTime(&elapsedTime,start,stop);//计算两个事件之间时长(单位为ms)

    核函数执行时间将被保存在变量elapsedTime中。通过这个值我们可以评估算法的性能。下面给一个例子,来看怎么使用计时功能。

    前面的例子规模很小,只有5个元素,处理量太小不足以计时,下面将规模扩大为1024,此外将反复运行1000次计算总时间,这样估计不容易受随机扰动影响。我们通过这个例子对比线程并行和块并行的性能如何。代码如下:

    1. #include "cuda_runtime.h"  
    2. #include "device_launch_parameters.h"  
    3. #include <stdio.h>  
    4. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);  
    5. __global__ void addKernel_blk(int *c, const int *a, const int *b)  
    6. {  
    7.     int i = blockIdx.x;  
    8.     c[i] = a[i]+ b[i];  
    9. }  
    10. __global__ void addKernel_thd(int *c, const int *a, const int *b)  
    11. {  
    12.     int i = threadIdx.x;  
    13.     c[i] = a[i]+ b[i];  
    14. }  
    15. int main()  
    16. {  
    17.     const int arraySize = 1024;  
    18.     int a[arraySize] = {0};  
    19.     int b[arraySize] = {0};  
    20.     for(int i = 0;i<arraySize;i++)  
    21.     {  
    22.         a[i] = i;  
    23.         b[i] = arraySize-i;  
    24.     }  
    25.     int c[arraySize] = {0};  
    26.     // Add vectors in parallel.  
    27.     cudaError_t cudaStatus;  
    28.     int num = 0;  
    29.     cudaDeviceProp prop;  
    30.     cudaStatus = cudaGetDeviceCount(&num);  
    31.     for(int i = 0;i<num;i++)  
    32.     {  
    33.         cudaGetDeviceProperties(&prop,i);  
    34.     }  
    35.     cudaStatus = addWithCuda(c, a, b, arraySize);  
    36.     if (cudaStatus != cudaSuccess)   
    37.     {  
    38.         fprintf(stderr, "addWithCuda failed!");  
    39.         return 1;  
    40.     }  
    41.   
    42.     // cudaThreadExit must be called before exiting in order for profiling and  
    43.     // tracing tools such as Nsight and Visual Profiler to show complete traces.  
    44.     cudaStatus = cudaThreadExit();  
    45.     if (cudaStatus != cudaSuccess)   
    46.     {  
    47.         fprintf(stderr, "cudaThreadExit failed!");  
    48.         return 1;  
    49.     }  
    50.     for(int i = 0;i<arraySize;i++)  
    51.     {  
    52.         if(c[i] != (a[i]+b[i]))  
    53.         {  
    54.             printf("Error in %d\n",i);  
    55.         }  
    56.     }  
    57.     return 0;  
    58. }  
    59. // Helper function for using CUDA to add vectors in parallel.  
    60. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)  
    61. {  
    62.     int *dev_a = 0;  
    63.     int *dev_b = 0;  
    64.     int *dev_c = 0;  
    65.     cudaError_t cudaStatus;  
    66.   
    67.     // Choose which GPU to run on, change this on a multi-GPU system.  
    68.     cudaStatus = cudaSetDevice(0);  
    69.     if (cudaStatus != cudaSuccess)   
    70.     {  
    71.         fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
    72.         goto Error;  
    73.     }  
    74.     // Allocate GPU buffers for three vectors (two input, one output)    .  
    75.     cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));  
    76.     if (cudaStatus != cudaSuccess)   
    77.     {  
    78.         fprintf(stderr, "cudaMalloc failed!");  
    79.         goto Error;  
    80.     }  
    81.     cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));  
    82.     if (cudaStatus != cudaSuccess)   
    83.     {  
    84.         fprintf(stderr, "cudaMalloc failed!");  
    85.         goto Error;  
    86.     }  
    87.     cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));  
    88.     if (cudaStatus != cudaSuccess)   
    89.     {  
    90.         fprintf(stderr, "cudaMalloc failed!");  
    91.         goto Error;  
    92.     }  
    93.     // Copy input vectors from host memory to GPU buffers.  
    94.     cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);  
    95.     if (cudaStatus != cudaSuccess)   
    96.     {  
    97.         fprintf(stderr, "cudaMemcpy failed!");  
    98.         goto Error;  
    99.     }  
    100.     cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);  
    101.     if (cudaStatus != cudaSuccess)   
    102.     {  
    103.         fprintf(stderr, "cudaMemcpy failed!");  
    104.         goto Error;  
    105.     }  
    106.     cudaEvent_t start,stop;  
    107.     cudaEventCreate(&start);  
    108.     cudaEventCreate(&stop);  
    109.     cudaEventRecord(start,0);  
    110.     for(int i = 0;i<1000;i++)  
    111.     {  
    112. //      addKernel_blk<<<size,1>>>(dev_c, dev_a, dev_b);  
    113.         addKernel_thd<<<1,size>>>(dev_c, dev_a, dev_b);  
    114.     }  
    115.     cudaEventRecord(stop,0);  
    116.     cudaEventSynchronize(stop);  
    117.     float tm;  
    118.     cudaEventElapsedTime(&tm,start,stop);  
    119.     printf("GPU Elapsed time:%.6f ms.\n",tm);  
    120.     // cudaThreadSynchronize waits for the kernel to finish, and returns  
    121.     // any errors encountered during the launch.  
    122.     cudaStatus = cudaThreadSynchronize();  
    123.     if (cudaStatus != cudaSuccess)   
    124.     {  
    125.         fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);  
    126.         goto Error;  
    127.     }  
    128.     // Copy output vector from GPU buffer to host memory.  
    129.     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);  
    130.     if (cudaStatus != cudaSuccess)   
    131.     {  
    132.         fprintf(stderr, "cudaMemcpy failed!");  
    133.         goto Error;  
    134.     }  
    135. Error:  
    136.     cudaFree(dev_c);  
    137.     cudaFree(dev_a);  
    138.     cudaFree(dev_b);      
    139.     return cudaStatus;  
    140. }  
    #include "cuda_runtime.h"
    #include "device_launch_parameters.h"
    #include <stdio.h>
    cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);
    __global__ void addKernel_blk(int *c, const int *a, const int *b)
    {
        int i = blockIdx.x;
        c[i] = a[i]+ b[i];
    }
    __global__ void addKernel_thd(int *c, const int *a, const int *b)
    {
        int i = threadIdx.x;
        c[i] = a[i]+ b[i];
    }
    int main()
    {
        const int arraySize = 1024;
        int a[arraySize] = {0};
        int b[arraySize] = {0};
    	for(int i = 0;i<arraySize;i++)
    	{
    		a[i] = i;
    		b[i] = arraySize-i;
    	}
        int c[arraySize] = {0};
        // Add vectors in parallel.
        cudaError_t cudaStatus;
    	int num = 0;
    	cudaDeviceProp prop;
    	cudaStatus = cudaGetDeviceCount(&num);
    	for(int i = 0;i<num;i++)
    	{
    		cudaGetDeviceProperties(&prop,i);
    	}
    	cudaStatus = addWithCuda(c, a, b, arraySize);
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "addWithCuda failed!");
            return 1;
        }
    
        // cudaThreadExit must be called before exiting in order for profiling and
        // tracing tools such as Nsight and Visual Profiler to show complete traces.
        cudaStatus = cudaThreadExit();
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaThreadExit failed!");
            return 1;
        }
        for(int i = 0;i<arraySize;i++)
    	{
    		if(c[i] != (a[i]+b[i]))
    		{
    			printf("Error in %d\n",i);
    		}
    	}
        return 0;
    }
    // Helper function for using CUDA to add vectors in parallel.
    cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)
    {
        int *dev_a = 0;
        int *dev_b = 0;
        int *dev_c = 0;
        cudaError_t cudaStatus;
    
        // Choose which GPU to run on, change this on a multi-GPU system.
        cudaStatus = cudaSetDevice(0);
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
            goto Error;
        }
        // Allocate GPU buffers for three vectors (two input, one output)    .
        cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMalloc failed!");
            goto Error;
        }
        cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMalloc failed!");
            goto Error;
        }
        cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMalloc failed!");
            goto Error;
        }
        // Copy input vectors from host memory to GPU buffers.
        cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMemcpy failed!");
            goto Error;
        }
        cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMemcpy failed!");
            goto Error;
        }
    	cudaEvent_t start,stop;
    	cudaEventCreate(&start);
    	cudaEventCreate(&stop);
    	cudaEventRecord(start,0);
    	for(int i = 0;i<1000;i++)
    	{
    //		addKernel_blk<<<size,1>>>(dev_c, dev_a, dev_b);
    		addKernel_thd<<<1,size>>>(dev_c, dev_a, dev_b);
    	}
    	cudaEventRecord(stop,0);
    	cudaEventSynchronize(stop);
    	float tm;
    	cudaEventElapsedTime(&tm,start,stop);
    	printf("GPU Elapsed time:%.6f ms.\n",tm);
        // cudaThreadSynchronize waits for the kernel to finish, and returns
        // any errors encountered during the launch.
        cudaStatus = cudaThreadSynchronize();
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
            goto Error;
        }
        // Copy output vector from GPU buffer to host memory.
        cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
        if (cudaStatus != cudaSuccess) 
    	{
            fprintf(stderr, "cudaMemcpy failed!");
            goto Error;
        }
    Error:
        cudaFree(dev_c);
        cudaFree(dev_a);
        cudaFree(dev_b);    
        return cudaStatus;
    }

    addKernel_blk是采用块并行实现的向量相加操作,而addKernel_thd是采用线程并行实现的向量相加操作。分别运行,得到的结果如下图所示:

    线程并行:

    块并行:

    可见性能竟然相差近16倍!因此选择并行处理方法时,如果问题规模不是很大,那么采用线程并行是比较合适的,而大问题分多个线程块处理时,每个块内线程数不要太少,像本文中的只有1个线程,这是对硬件资源的极大浪费。一个理想的方案是,分N个线程块,每个线程块包含512个线程,将问题分解处理,效率往往比单一的线程并行处理或单一块并行处理高很多。这也是CUDA编程的精髓。

    上面这种分析程序性能的方式比较粗糙,只知道大概运行时间长度,对于设备程序各部分代码执行时间没有一个深入的认识,这样我们就有个问题,如果对代码进行优化,那么优化哪一部分呢?是将线程数调节呢,还是改用共享内存?这个问题最好的解决方案就是利用Visual Profiler。下面内容摘自《CUDA_Profiler_Users_Guide》

    “Visual Profiler是一个图形化的剖析工具,可以显示你的应用程序中CPU和GPU的活动情况,利用分析引擎帮助你寻找优化的机会。”

    其实除了可视化的界面,NVIDIA提供了命令行方式的剖析命令:nvprof。对于初学者,使用图形化的方式比较容易上手,所以本节使用Visual Profiler。

    打开Visual Profiler,可以从CUDA Toolkit安装菜单处找到。主界面如下:

    我们点击File->New Session,弹出新建会话对话框,如下图所示:

    其中File一栏填入我们需要进行剖析的应用程序exe文件,后面可以都不填(如果需要命令行参数,可以在第三行填入),直接Next,见下图:

    第一行为应用程序执行超时时间设定,可不填;后面三个单选框都勾上,这样我们分别使能了剖析,使能了并发核函数剖析,然后运行分析器。

    点Finish,开始运行我们的应用程序并进行剖析、分析性能。

    上图中,CPU和GPU部分显示了硬件和执行内容信息,点某一项则将时间条对应的部分高亮,便于观察,同时右边详细信息会显示运行时间信息。从时间条上看出,cudaMalloc占用了很大一部分时间。下面分析器给出了一些性能提升的关键点,包括:低计算利用率(计算时间只占总时间的1.8%,也难怪,加法计算复杂度本来就很低呀!);低内存拷贝/计算交叠率(一点都没有交叠,完全是拷贝——计算——拷贝);低存储拷贝尺寸(输入数据量太小了,相当于你淘宝买了个日记本,运费比实物价格还高!);低存储拷贝吞吐率(只有1.55GB/s)。这些对我们进一步优化程序是非常有帮助的。

    我们点一下Details,就在Analysis窗口旁边。得到结果如下所示:

    通过这个窗口可以看到每个核函数执行时间,以及线程格、线程块尺寸,占用寄存器个数,静态共享内存、动态共享内存大小等参数,以及内存拷贝函数的执行情况。这个提供了比前面cudaEvent函数测时间更精确的方式,直接看到每一步的执行时间,精确到ns。

    在Details后面还有一个Console,点一下看看。

    这个其实就是命令行窗口,显示运行输出。看到加入了Profiler信息后,总执行时间变长了(原来线程并行版本的程序运行时间只需4ms左右)。这也是“测不准定理”决定的,如果我们希望测量更细微的时间,那么总时间肯定是不准的;如果我们希望测量总时间,那么细微的时间就被忽略掉了。

    后面Settings就是我们建立会话时的参数配置,不再详述。

    转自:http://blog.csdn.net/augusdi/article/details/12833235

    展开全文
  • CUDA从入门到精通(大神之作)膜拜

    万次阅读 多人点赞 2018-07-27 13:22:50
    CUDA从入门到精通(零):写在前面   在老板的要求下,本博主从2012年上高性能计算课程开始接触CUDA编程,随后将该技术应用到了实际项目中,使处理程序加速超过1K,可见基于图形显示器的并行计算对于追求速度的...
  • 在进行CUDA Fortran程序的测试时,发现结果不对,但是又不存在语法和逻辑错误,所以十分需要获取程序错误信息,但是Nvidia和CUDA提供了很多的可视化(Nsight、NVVP)和命令行(cuda-gdb)的形式,但是这些都支持CUDA C...
  • CUDA Fortran和CUDA C有什么区别呢,看这张表你就知道啦: C和Fortran区别.jpg (86.13 KB, 下载次数: 0) 下载附件 保存到相册 2016-11-8 13:33 上传 原文发布时间为:2016-11-8 13:43:46 原文由:emily ....
  • CUDA编程入门极简教程

    万次阅读 多人点赞 2018-04-07 13:59:02
    码字不易,欢迎给个赞! 欢迎交流与转载,文章会同步发布在公众号:机器...2006年,NVIDIA公司发布了CUDACUDA是建立在NVIDIA的CPUs上的一个通用并行计算平台和编程模型,基于CUDA编程可以利用GPUs的并行计算引擎...
  • 作者: 叶 虎 编辑:李雪冬 前 言2006年,NVIDIA公司发布了CUDA(http://docs.nvidia.com/cuda/),CUDA是建立在NVI
  • cudafortran高效编程实践,译者,小小河。资源来自网络,侵权请通知删除。
  • 如何查看windows的CUDA版本

    万次阅读 多人点赞 2019-04-12 00:27:13
    最近开始学习一些关于图像处理的计算机视觉...只需要我们自己选择相应的系统,下载使用的工具方式,python的版本,以及CUDA版本号,就可以获得到相应的命令进行下载安装。 这时遇到一个问题,CUDA是什么? CUDA(C...
  • tensorflow各个版本的CUDA以及Cudnn版本对应关系

    万次阅读 多人点赞 2019-04-08 18:05:37
    CUDA是NVIDIA推出的用于自家GPU的并行计算框架,也就是说CUDA只能在NVIDIA的GPU上运行,而且只有当要解决的计算问题是可以大量并行计算的时候才能发挥CUDA的作用。 CUDA的本质是一个工具包(ToolKit);但是二者...
  • win10下CUDA和CUDNN的安装(超详细)!亲测有效!

    万次阅读 多人点赞 2019-04-09 22:13:55
    CUDA8安装配置 CUDA8的安装包可直接从NVIDIA官网下载。根据相应的系统选项,我选择的是cuda_8.0.61_win10.exe(大小为1.3G),安装的时候建议选择 自定义 而不是“精简”(从下面的英文解释可以看出,其实这里的精简写...
  • CUDA从入门到精通

    万次阅读 多人点赞 2013-10-17 11:22:30
    CUDA从入门到精通(零):写在前面 在老板的要求下,本博主从2012年上高性能计算课程开始接触CUDA编程,随后将该技术应用到了实际项目中,使处理程序加速超过1K,可见基于图形显示器的并行计算对于追求速度的应用来...
  • windows 7 下cuda 9.0 卸载、cuda8.0 安装

    万次阅读 多人点赞 2018-03-27 10:29:34
    安装cuda9.0之后,电脑原来的NVIDIA图形驱动会被更新,NVIDIA Physx系统软件也会被更新(安装低版cuda可能不会被更新)。卸载时候要注意了,别动这2个。 2.卸载: 1.前言: 杀毒软件别用来卸载这个了,不好找...
  • 查看 CUDA cudnn 版本

    万次阅读 多人点赞 2017-08-01 15:27:09
    https://medium.com/@changrongko/nv-how-to-check-cuda-and-cudnn-version-e05aa21daf6ccuda 版本 cat /usr/local/cuda/version.txtcudnn 版本 cat /usr/local/cuda/include/cudnn.h | grep CUDNN_MAJOR -A 2
  • Cuda安装详细步骤

    万次阅读 多人点赞 2017-01-09 12:01:19
    为了学习Caffe新买了笔记本,安装cuda时遇到了很多问题,不过好在都成功的解决了。 网上有很多cuda的安装教程,每个人的电脑配置不同遇到的问题也不一样,现在就我自己的安装配置情况总结一下具体的安装步骤,因为...
  • 深度学习配置CUDA8.0/9.0及对应版本cuDNN安装

    万次阅读 多人点赞 2017-12-26 01:56:13
    本人为中科院测地所博士生,所研究专业为自然地理学(遥感数据分析方向),研究课题偏向于深度学习。 由于本人不是计算机专业,故有关...本文主要讲解CUDA8.0/cuDNN5.1(CUDA9.1/cuDNN7.0)的安装及环境配置步骤。
  • Win10安装CUDA10和cuDNN

    万次阅读 多人点赞 2018-10-17 21:52:30
    官方安装教程 CUDA:https://docs.nvidia.com/cuda/cuda-installation-guide-microsoft-windows/index.html cuDNN:...WIN10安装CUDA10 CUDA ...
1 2 3 4 5 ... 20
收藏数 86,596
精华内容 34,638
关键字:

cuda