异构计算--CUDA架构

1.CUDA是什么?

  CUDA(Compute Unified Device Architecture),是显卡厂商NVIDIA推出的运算平台,是一种通用并行计算架构,该架构使GPU能够解决复杂的计算问题。说白了就是我们可以使用GPU来并行完成像神经网络、图像处理算法这些在CPU上跑起来比较吃力的程序。通过GPU和高并行,我们可以大大提高这些算法的运行速度。

2.CPU&CUDA架构

  处理器结构有2个指标要经常考虑的:延迟和吞吐。延迟指从发出指令到返回最终结果中间经历的时间间隔;吞吐指单位时间内处理的指令的条数。由于CPU以处理计算和控制为主要任务,所以设计理念是延迟导向内核;由于GPU以并行处理为主要任务,所以设计理念是吞吐导向内核。

(1)CPU

  CPU(CentralProcessing Unit)中央处理器,是一块超大规模的集成电路,主要逻辑架构包括控制单元Control,运算单元ALU和高速缓冲存储器(Cache)及实现它们之间联系的数据(Data)、控制及状态的总线(Bus)。简单说,就是计算单元、控制单元和存储单元。
架构图如下所示:
在这里插入图片描述

  CPU遵循的是冯·诺依曼架构,其核心是存储程序/数据、串行顺序执行。因此CPU的架构中需要大量的空间去放置存储单元(Cache)和控制单元(Control),相比之下计算单元(ALU)只占据了很小的一部分,所以CPU在进行大规模并行计算方面受到限制,相对而言更擅长于处理逻辑控制。CPU无法做到大量数据并行计算的能力,但GPU可以。

(2)GPU

  GPU(GraphicsProcessing Unit),即图形处理器,是一种由大量运算单元组成的大规模并行计算架构,早先由CPU中分出来专门用于处理图像并行计算数据,专为同时处理多重并行计算任务而设计。GPU中也包含基本的计算单元、控制单元和存储单元,但GPU的架构与CPU有很大不同,其架构图如下所示。
在这里插入图片描述

  与CPU相比,CPU芯片空间的不到20%是ALU,而GPU芯片空间的80%以上是ALU。即GPU拥有更多的ALU用于数据并行处理。这就是为什么GPU可以具备强大的并行计算能力的原因。
从硬件架构分析来看,CPU和GPU似乎很像,都有内存、Cache、ALU、CU,都有着很多的核心,但是CPU的核心占比比较重,相对计算单元ALU很少,可以用来处理非常复杂的控制逻辑,预测分支、乱序执行、多级流水任务等等。相对而言GPU的核心就是比较轻,用于优化具有简单控制逻辑的数据并行任务,注重并行程序的吞吐量。
  简单来说就是CPU的核心擅长完成多重复杂任务,重在逻辑,重在串行程序;GPU的核心擅长完成具有简单的控制逻辑的 任务,重在计算,重在并行。

3.NVIDIA显卡硬件架构:SM、SP、Warp

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

在这里插入图片描述
每个SM包含的SP数量依据GPU架构而不同,Fermi架构GF100是32个,GF10X是48个,Kepler架构都是192个,Maxwell都是128个。当一个kernel启动后,thread会被分配到很多SM中执行。大量的thread可能会被分配到不同的SM,但是同一个block中的thread必然在同一个SM中并行执行。
Warp调度
一个SP可以执行一个thread,但是实际上并不是所有的thread能够在同一时刻执行。Nvidia把32个threads组成一个warp,warp是调度和运行的基本单元。warp中所有threads并行的执行相同的指令。warp由SM的硬件warp scheduler负责调度,一个SM同一个时刻可以执行多个warp,这取决于warp scheduler的数量。目前每个warp包含32个threads(Nvidia保留修改数量的权利)。

4.软件架构:Kernel、Grid、Block

具体到我们如何调用GPU上的线程实现我们的算法,则是通过Kernel实现的。在GPU上调用的函数成为CUDA核函数(Kernel function),核函数会被GPU上的多个线程执行。我们可以通过如下方式来定义一个kernel:

kernel_function<<<grid, block>>>(param1, param2, param3....);

如下图所示为kernel、grid、block的组织方式以及对应的硬件部分。
在这里插入图片描述
Grid: 由一个单独的kernel启动的所有线程组成一个grid,grid中所有线程共享global memory。Grid由很多Block组成,可以是一维二维或三维。

Block: 一个grid由许多block组成,block由许多线程组成,同样可以有一维、二维或者三维。block内部的多个线程可以同步(synchronize),可访问共享内存(share memory)。

5.CUDA内存模型

CUDA中的内存模型分为以下几个层次:

  • 每个线程都用自己的registers(寄存器)
  • 每个线程都有自己的local memory(局部内存)
  • 每个线程块内都有自己的shared memory(共享内存),所有线程块内的所有线程共享这段内存资源
  • 每个grid都有自己的global memory(全局内存),不同线程块的线程都可使用
  • 每个grid都有自己的constant memory(常量内存)和texture memory(纹理内存),),不同线程块的线程都可使用
  • 线程访问这几类存储器的速度是register > local memory >shared memory > global memory

下面这幅图表示就是这些内存在计算机架构中的所在层次。

在这里插入图片描述

6.异构计算

  所谓异构计算,是指CPU+ GPU或者CPU+ 其它设备(如FPGA等)协同计算。⼀般我们的程序,是在CPU上计算。但是,当⼤量的数据需要计算时,CPU显得⼒不从⼼。那么,是否可以找寻其它的⽅法来解决计算速度呢?那就是异构计算。例如可利⽤CPU(Central Processing Unit)、GPU(Graphic Processing Unit)、甚⾄APU(Accelerated Processing Units, CPU与GPU的融合)等计算设备的计算能⼒从⽽来提⾼系统的速度。异构系统越来越普遍,对于⽀持这种环境的计算⽽⾔,也正受到越来越多的关注。
  ⽬前异构计算使⽤最多的是利⽤GPU来加速。主流GPU都采⽤了统⼀架构单元,凭借强⼤的可编程流处理器阵容,GPU在单精度浮点运算⽅⾯将CPU远远甩在⾝后。如图所示为CPU+GPU异构计算的一个示意图,其中GPU主要负责并行计算。
在这里插入图片描述

7.OpenCL与CUDA的关系

  NVIDIA的CUDA架构和KHRONOS制定的OpenCL并不冲突,他们之间的关系是API与执行架构之间的关系,举个简单的例子:我们熟悉的X86架构是一种CPU架构,而各种编程语言,如:汇编语言、C语言等低级语言或高级语言仅仅是建立在X86运算架构之上的一种编程环境。那么,CUDA架构和OpenCL之间的关系和X86与编程语言的关系是相同的。
  CUDA架构是OpenCL的运行平台之一,因此他们之间并不存在谁取代谁的关系。OpenCL仅仅是为CUDA架构提供了一个可编程的API而已。
在这里插入图片描述
参考文献
[1] http://t.zoukankan.com/liuyufei-p-13259264.html
[2] https://zhuanlan.zhihu.com/p/266633373?utm_id=0
[3] https://betheme.net/a/21450178.html?action=onClick
[4] https://blog.csdn.net/chongbin007/article/details/123667504

  • 0
    点赞
  • 16
    收藏
    觉得还不错? 一键收藏
  • 2
    评论
目目目 录录录 目录 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · i 第一章 导论 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 1 1.1 从图形处理到通用并行计算 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 1 1.2 CUDATM :一种通用并行计算架构 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 3 1.3 一种可扩展的编程模型· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 3 1.4 文档结构 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 4 第二章 编程模型 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 7 2.1 内核· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 7 2.2 线程层次 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 8 2.3 存储器层次 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 11 2.4 异构编程 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 11 2.5 计算能力 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 11 第三章 编程接口 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 15 3.1 用nvcc编译 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 15 3.1.1 编译流程 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 16 3.1.1.1 离线编译 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 16 3.1.1.2 即时编译 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 16 3.1.2 二进制兼容性· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 17 3.1.3 PTX兼容性· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 17 3.1.4 应用兼容性 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 18 3.1.5 C/C++兼容性· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 19 3.1.6 64位兼容性 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 19 3.2 CUDA C运行时· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 19 ii CUDA编程指南5.0中文版 3.2.1 初始化 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 20 3.2.2 设备存储器 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 20 3.2.3 共享存储器 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 24 3.2.4 分页锁定主机存储器 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 32 3.2.4.1 可分享存储器(portable memory) · · · · · · · · · · · · · · · · 34 3.2.4.2 写结合存储器· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 34 3.2.4.3 被映射存储器· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 34 3.2.5 异步并发执行· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 35 3.2.5.1 主机和设备间异步执行· · · · · · · · · · · · · · · · · · · · · · · · · · 35 3.2.5.2 数据传输和内核执行重叠 · · · · · · · · · · · · · · · · · · · · · · · 36 3.2.5.3 并发内核执行· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 36 3.2.5.4 并发数据传输· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 36 3.2.5.5 流 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 37 3.2.5.6 事件· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 41 3.2.5.7 同步调用 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 42 3.2.6 多设备系统 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 42 3.2.6.1 枚举设备 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 42 3.2.6.2 设备指定 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 42 3.2.6.3 流和事件行为· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 43 3.2.6.4 p2p存储器访问 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 44 3.2.6.5 p2p存储器复制 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 45 3.2.6.6 统一虚拟地址空间 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 45 3.2.6.7 错误检查 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 46 3.2.7 调用栈 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 47 3.2.8 纹理和表面存储器 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 47 3.2.8.1 纹理存储器 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 47 3.2.8.2 表面存储器(surface) · · · · · · · · · · · · · · · · · · · · · · · · · · · · 60 3.2.8.3 CUDA 数组 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 65 目录 iii 3.2.8.4 读写一致性 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 66 3.2.9 图形学互操作性 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 66 3.2.9.1 OpenGL互操作性 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 67 3.2.9.2 Direct3D互操作性 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 70 3.2.9.3 SLI(速力)互操作性· · · · · · · · · · · · · · · · · · · · · · · · · · · 82 3.3 版本和兼容性· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 82 3.4 计算模式 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 83 3.5 模式切换 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 84 3.6 Windows上的Tesla计算集群模式 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 85 第四章 硬件实现 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 87 4.1 SIMT 架构 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 87 4.2 硬件多线程 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 88 第五章 性能指南 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 91 5.1 总体性能优化策略 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 91 5.2 最大化利用率· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 91 5.2.1 应用层次 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 91 5.2.2 设备层次 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 92 5.2.3 多处理器层次· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 92 5.3 最大化存储器吞吐量 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 94 5.3.1 主机和设备的数据传输· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 95 5.3.2 设备存储器访问 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 96 5.3.2.1 全局存储器 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 96 5.3.2.2 本地存储器 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 98 5.3.2.3 共享存储器 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 99 5.3.2.4 常量存储器 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 100 5.3.2.5 纹理和表面存储器 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 100 5.4 最大化指令吞吐量 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 100 iv CUDA编程指南5.0中文版 5.4.1 算术指令 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 101 5.4.2 控制流指令 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 104 5.4.3 同步指令 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 105 附录 A 支持CUDA的GPU · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 107 附录 B C语言扩展 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 109 B.1 函数类型限定符 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 109 B.1.1 device · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 109 B.1.2 global · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 109 B.1.3 host · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 109 B.1.4 noinline 和 forceinline · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 110 B.2 变量类型限定符 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 110 B.2.1 device · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 111 B.2.2 constant · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 111 B.2.3 shared · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 112 B.2.4 restrict · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 113 B.3 内置变量类型· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 115 B.3.1 char1、uchar1、char2、uchar2、char3、uchar3、char4、 uchar4、short1、ushort1、short2、ushort2、short3、ushort3、 short4、ushort4、int1、uint1、int2、uint2、int3、uint3、 int4、uint4、long1、ulong1、long2、ulong2、long3、ulong3、 long4、ulong4、float1、float2、float3、float4、double2 · · · 115 B.3.2 dim3类型 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 115 B.4 内置变量 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 115 B.4.1 gridDim · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 115 B.4.2 blockIdx · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 115 B.4.3 blockDim · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 117 B.4.4 threadIdx · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 117 B.4.5 warpSize · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 117 目录 v B.5 存储器栅栏函数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 117 B.6 同步函数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 119 B.7 数学函数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 120 B.8 纹理函数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 120 B.8.1 纹理对象函数· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 120 B.8.1.1 tex1Dfetch() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 120 B.8.1.2 tex1D()· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 121 B.8.1.3 tex2D()· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 121 B.8.1.4 tex3D()· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 121 B.8.1.5 tex1DLayered() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 121 B.8.1.6 tex2DLayered() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 122 B.8.1.7 texCubemap() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 122 B.8.1.8 texCubemapLayered()· · · · · · · · · · · · · · · · · · · · · · · · · · · 122 B.8.1.9 tex2Dgather() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 123 B.8.2 纹理参考函数· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 123 B.8.2.1 tex1Dfetch() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 123 B.8.2.2 tex1D()· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 124 B.8.2.3 tex2D()· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 124 B.8.2.4 tex3D()· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 125 B.8.2.5 tex1DLayered() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 125 B.8.2.6 tex2DLayered() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 125 B.8.2.7 texCubemap() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 125 B.8.2.8 texCubemapLayered()· · · · · · · · · · · · · · · · · · · · · · · · · · · 126 B.8.2.9 tex2Dgather() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 126 B.9 表面函数(surface)· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 126 B.9.1 表面对象函数· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 127 B.9.1.1 surf1Dread() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 127 B.9.1.2 surf1Dwrite() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 127 vi CUDA编程指南5.0中文版 B.9.1.3 surf2Dread() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 127 B.9.1.4 surf2Dwrite() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 128 B.9.1.5 surf3Dread() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 128 B.9.1.6 surf3Dwrite() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 128 B.9.1.7 surf1DLayeredread() · · · · · · · · · · · · · · · · · · · · · · · · · · · · 129 B.9.1.8 surf1DLayeredwrite() · · · · · · · · · · · · · · · · · · · · · · · · · · · 129 B.9.1.9 surf2DLayeredread() · · · · · · · · · · · · · · · · · · · · · · · · · · · · 129 B.9.1.10 surf2DLayeredwrite() · · · · · · · · · · · · · · · · · · · · · · · · · · · 130 B.9.1.11 surfCubemapread() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 130 B.9.1.12 surfCubemapwrite()· · · · · · · · · · · · · · · · · · · · · · · · · · · · · 131 B.9.1.13 surfCubemapLayeredread() · · · · · · · · · · · · · · · · · · · · · · 131 B.9.1.14 surfCubemapLayeredwrite() · · · · · · · · · · · · · · · · · · · · · 131 B.9.2 表面引用API · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 132 B.9.2.1 surf1Dread() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 132 B.9.2.2 surf1Dwrite() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 132 B.9.2.3 surf2Dread() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 132 B.9.2.4 surf2Dwrite() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 133 B.9.2.5 surf3Dread() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 133 B.9.2.6 surf3Dwrite() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 133 B.9.2.7 surf1DLayeredread() · · · · · · · · · · · · · · · · · · · · · · · · · · · · 134 B.9.2.8 surf1DLayeredwrite() · · · · · · · · · · · · · · · · · · · · · · · · · · · 134 B.9.2.9 surf2DLayeredread() · · · · · · · · · · · · · · · · · · · · · · · · · · · · 135 B.9.2.10 surf2DLayeredwrite() · · · · · · · · · · · · · · · · · · · · · · · · · · · 135 B.9.2.11 surfCubemapread() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 135 B.9.2.12 surfCubemapwrite()· · · · · · · · · · · · · · · · · · · · · · · · · · · · · 136 B.9.2.13 surfCubemapLayeredread() · · · · · · · · · · · · · · · · · · · · · · 136 B.9.2.14 surfCubemapLayeredwrite() · · · · · · · · · · · · · · · · · · · · · 137 B.10 时间函数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 137 目录 vii B.11 原子函数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 137 B.11.1 数学函数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 138 B.11.1.1 atomicAdd() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 138 B.11.1.2 atomicSub() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 139 B.11.1.3 atomicExch() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 139 B.11.1.4 atomicMin() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 140 B.11.1.5 atomicMax()· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 140 B.11.1.6 atomicInc() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 140 B.11.1.7 atomicDec() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 141 B.11.1.8 atomicCAS() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 141 B.11.2 位逻辑函数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 141 B.11.2.1 atomicAnd() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 141 B.11.2.2 atomicOr() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 142 B.11.2.3 atomicXor() · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 142 B.12 束表决(warp vote)函数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 142 B.13 束洗牌函数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 143 B.13.1 概览 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 143 B.13.2 在束内广播一个值 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 144 B.13.3 计算8个线程的前缀和· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 145 B.13.4 束内求和 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 146 B.14 取样计数器函数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 146 B.15 断言 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 147 B.16 格式化输出 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 148 B.16.1 格式化符号 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 149 B.16.2 限制 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 149 B.16.3 相关的主机端API · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 150 B.16.4 例程 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 151 B.17 动态全局存储器分配 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 152 viii CUDA编程指南5.0中文版 B.17.1 堆存储器分配 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 153 B.17.2 与设备存储器API的互操作 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 154 B.17.3 例程 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 154 B.17.3.1 每个线程的分配 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 154 B.17.3.2 每个线程块的分配 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 155 B.17.3.3 在内核启动之间持久的分配 · · · · · · · · · · · · · · · · · · · · · 156 B.18 执行配置 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 159 B.19 启动绑定 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 160 B.20 #pragma unroll · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 162 B.21 SIMD 视频指令 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 163 附录 C 数学函数· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 165 C.1 标准函数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 165 C.1.1 单精度浮点函数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 165 C.1.2 双精度浮点函数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 168 C.2 内置函数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 171 C.2.1 单精度浮点函数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 172 C.2.2 双精度浮点函数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 172 附录 D C++语言支持· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 175 D.1 代码例子 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 175 D.1.1 数据类 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 175 D.1.2 派生类 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 176 D.1.3 类模板 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 177 D.1.4 函数模板 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 178 D.1.5 函子类 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 178 D.2 限制· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 180 D.2.1 预处理符号 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 180 D.2.2 限定符 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 180 目录 ix D.2.2.1 设备存储器限定符 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 180 D.2.2.2 Volatile限定符 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 182 D.2.3 指针· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 182 D.2.4 运算符 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 183 D.2.4.1 赋值运算符 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 183 D.2.4.2 地址运算符 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 183 D.2.5 函数· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 183 D.2.5.1 编译器生成的函数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 183 D.2.5.2 函数参数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 184 D.2.5.3 函数内静态变量 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 184 D.2.5.4 函数指针 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 184 D.2.5.5 函数递归 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 185 D.2.6 类 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 185 D.2.6.1 数据成员 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 185 D.2.6.2 函数成员 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 185 D.2.6.3 虚函数 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 185 D.2.6.4 虚基类 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 185 D.2.6.5 Windows相关 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 185 D.2.7 模板· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 186 附录 E 纹理获取· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 187 E.1 最近点取样 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 187 E.2 线性滤波 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 187 E.3 查找表 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 189 附录 F 计算能力 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 191 F.1 特性和技术规范 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 191 F.2 浮点标准 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 195 F.3 计算能力1.x · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 198 x CUDA编程指南5.0中文版 F.3.1 架构· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 198 F.3.2 全局存储器 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 199 F.3.2.1 计算能力1.0和1.1的设备 · · · · · · · · · · · · · · · · · · · · · · · · 199 F.3.2.2 计算能力1.2和1.3的设备 · · · · · · · · · · · · · · · · · · · · · · · · 199 F.3.3 共享存储器 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 201 F.3.3.1 32位步长访问· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 201 F.3.3.2 32位广播访问· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 202 F.3.3.3 8位和16位访问 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 205 F.3.3.4 大于32位访问· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 205 F.4 计算能力2.x · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 206 F.4.1 架构· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 206 F.4.2 全局存储器 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 208 F.4.3 共享存储器 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 209 F.4.3.1 32位步长访问· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 209 F.4.3.2 大于32位访问· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 210 F.4.4 常量存储器 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 211 F.5 计算能力3.x · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 211 F.5.1 架构· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 211 F.5.2 全局存储器访问 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 212 F.5.3 共享存储器 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 213 F.5.3.1 64位模式 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 213 F.5.3.2 32位模式 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 213 附录 G 驱动API · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 215 G.1 上下文 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 218 G.2 模块· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 219 G.3 内核执行 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 220 G.4 运行时API和驱动API的互操作性 · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 222 G.5 注意· · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · · 223 表表表 格格格 5.1 原生算术指令吞吐量/每时钟每流多处

“相关推荐”对你有帮助么?

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论 2
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值