NVIDIA CUDA User Manual

NVIDIA CUDA 编程指南

NVIDIA

技术文档（全译文）

Jan. 2008 Version 1.1

GPU

系列技术文档

.....................................................................................................................1

NVIDIA CUDA

Chapter1

1.1

作为一个并行数据计算设备的图形处理器单元

1.2 CUDA

Chapter2

2.1

一个超多线程协处理器

2.2

线程批处理

2.2.1

2.2.2

2.3

线程块线程块栅格

内存模型

Chapter3

3.1

一组带有

3.2

执行模式

3.3

计算兼容性

3.4

多设备

3.5

模式切换

编程指南

介绍

CUDA…….....................................................................................................................11

：一个在

编程模型

.........................................................................................................................1

GPU

上计算的新架构

............................................................................................................................... 15

.....................................................................................................................15

.......................................................................................................................................15

..........................................................................................................................................16

...................................................................................................................................16

...........................................................................................................................................17

硬件实现

on-chip

................................................................................................................................18

共享内存的

SIMD

多处理器

...........................................................................................................................................19

........................................................................................................................................20

...............................................................................................................................................20

...........................................................................................................................................20

………………………….............................................11

..............................................................................................12

....................................................................................18

Chapter4

4.1

一个C语言的扩展

4.2

语言扩展

4.2.1

应用程序编程接口（

API） ...................................................................................................21

..............................................................................................................................21

...........................................................................................................................................21

函数类型限定词

......................................................................................................................... 22

4.2.1.1 __device__..................................................................................................................22

4.2.1.2 __global__..........................................................................................................22

4.2.1.3 __host__............................................................................................................22

4.2.1.4

4.2.2

限定

....................................................................................................................22

变量类型限定词

........................................................................................................................23

4.2.2.1 __device__..................................................................................................23

4.2.2.2 __constant__ .......................................................................................................23

4.2.2.3 __shared__............................................................................................................23

4.2.2.4

4.2.3

4.2.4

限定

执行配置

内置变量

.............................................................................................................................24

...........................................................................................................................25

..............................................................................................................................26

4.2.4.1 gridDim.............................................................................................................................. 26

4.2.4.2 blockIdx...........................................................................................................................26

4.2.4.3 blockDim............................................................................................................................26

- 2 -

4.2.4.4 threadIdx......................................................................................................................... 26

4.2.4.5

4.2.5 NVCC

限定

..................................................................................................................................... 26

编译

.................................................................................................................................. 26

4.2.5.1 __noinline__ ...................................................................................................................27

4.2.5.2 #pragmaunroll .................................................................................................................27

4.3

公共

4.3.1

Runtime

内置矢量类型

组件

........................................................................................................................... 28

............................................................................................................................. 28

4.3.1.1

char1,uchar1,char2,uchar2,char3,uchar3,char4,uchar4,short1,ushort1,short2,us

hort2,short3,ushort3,short4,ushort4,int1,uint1,int2,uint2,int3,uint3,int4,ui

nt4,long1,ulong1,long2,ulong2,long3,ulong3,long4,ulong4,float1,float2,float3

,float4.......................................................................................................................................... 28

4.3.1.2 dim3

4.3.2

数学函数

4.3.3

时间函数

4.3.4

纹理类型

4.3.4.1 Texture Reference

4.3.4.2 RuntimeTexture Reference

类型

........................................................................................................................... 28

..................................................................................................................................... 28

..................................................................................................................................... 29

声明

.......................................................................................................29

属性

.........................................................................................30

4.3.4.3

4.4

设备

4.4.1

4.4.2

4.4.3

4.4.4 TypeCasting

4.4.5

4.4.5.1

4.4.5.2 CUDA

4.4.6

4.5

主机

4.5.1

4.5.1.1

4.5.1.2

线性内存纹理操作对比

Runtime

数学函数同步函数类型转换函数

组件

........................................................................................................................... 31

..................................................................................................................................... 31

............................................................................................................................. 32

函数

纹理函数

..................................................................................................................................... 33

设备内存纹理操作

数组纹理操作

原子函数

Runtime

公共概念

...................................................................................................................................... 34

组件

........................................................................................................................... 34

..................................................................................................................................... 35

设备

..................................................................................................................................... 35

内存

..................................................................................................................................... 35

CUDA

数组

.......................................................................................31

........................................................................................................................32

................................................................................................................33

............................................................................................................33

4.5.1.3 OpenGL Interoperability ...................................................................................................... 36

4.5.1.4 Direct3D Interoperability ..................................................................................................... 36

4.5.1.5

异步的并发执行

...................................................................................................................37

- 3 -

4.5.2 RuntimeAPI.................................................................................................................................... 38

4.5.2.1

4.5.2.2

4.5.2.3

4.5.2.4

4.5.2.5

4.5.2.6 Texture Reference

初始化

..................................................................................................................................... 38

设备管理内存管理流管理

..................................................................................................................................... 40

事件管理

.................................................................................................................................. 38

.................................................................................................................................. 39

.................................................................................................................................. 41

管理

...........................................................................................................42

4.5.2.7 OpenGL Interoperability .......................................................................................................... 44

4.5.2.8 Direct3D Interoperability ......................................................................................................... 44

4.5.2.9

4.5.3

4.5.3.1

4.5.3.2

4.5.3.3 Context

4.5.3.4

4.5.3.5

4.5.3.6

4.5.3.7

使用设备仿真方式调试

驱动

API ........................................................................................................................................ 47

初始化设备管理

模块管理执行控制内存管理流管理

..................................................................................................................................... 47

................................................................................................................................. 47

管理

............................................................................................................................. 47

................................................................................................................................. 48

................................................................................................................................. 49

..................................................................................................................................... 51

.............................................................................................................45

4.5.3.8

4.5.3.9 Texture Reference

事件管理

................................................................................................................................. 51

管理

..........................................................................................................52

4.5.3.10 OpenGL Interoperability ...................................................................................................... 53

4.5.3.11 Direct3D Interoperability ...................................................................................................... 53

Chapter5

5.1

指令性能

5.1.1

5.1.1.1

5.1.1.2

5.1.1.3

5.1.1.4

5.1.2

5.1.2.1

5.1.2.2

5.1.2.3

5.1.2.4

5.1.2.5

性能指导

................................................................................................................................ 54

........................................................................................................................................... 54

指令吞吐量

算术指令控制流指令内存指令同步指令

内存带宽

全局内存常量内存纹理内存共享内存寄存器

................................................................................................................................. 54

............................................................................................................................. 54

.......................................................................................................................... 55

............................................................................................................................. 56

.................................................................................................................................... 56

............................................................................................................................. 57

............................................................................................................................. 62

............................................................................................................................. 63

................................................................................................................................. 70

- 4 -

5.2

每个块的线程数量

5.3

主机与设备的数据传输

.......................................................................................................................... 70

....................................................................................................................71

5.4 Texture Fetch

5.5

性能优化策略总结

Chapter6

6.1

概要

6.2

源代码

6.3

源代码解释

矩阵乘法的例子

............................................................................................................................................... 74

........................................................................................................................................... 76

对比全局或常驻内存读取

.........................................................................................71

.......................................................................................................................... 72

....................................................................................................................74

.................................................................................................................................... 78

6.3.1 Mul().................................................................................................................................. 78

6.3.2 Muld()................................................................................................................................ 79

附录A 技术规格

A.1

通用规格

A.2

浮点数标准

附录B 数学函数

B.1

公共

runtime

B.2

设备

runtime

附录C 原子函数

C.1

算法函数

................................................................................................................................... 80

....................................................................................................................................... 81

.................................................................................................................................... 82

................................................................................................................................... 83

组件

............................................................................................................................. 83

组件

........................................................................................................................... 86

................................................................................................................................... 88

....................................................................................................................................... 88

C.1.1 atomicAdd() ................................................................................................................... 88

C.1.2 atomicSub() ................................................................................................................... 88

C.1.3 atomicExch() ................................................................................................................. 88

C.1.4 atomicMin() ................................................................................................................... 88

C.1.5 atomicMax() ................................................................................................................... 89

C.1.6 atomicInc() ....................................................................................................................89

C.1.7 atomicDec() ................................................................................................................... 89

C.1.8 atomicCAS() ................................................................................................................... 89

C.2

位操作函数

.................................................................................................................................... 90

C.2.1 atomicAnd() ................................................................................................................... 90

C.2.2 atomicOr()..................................................................................................................….. 90

C.2.3 atomicXor() ................................................................................................................... 90

附录

D Runtime API Reference ............................................................................................................ 91

D.1

设备管理

........................................................................................................................................ 91

D.1.1 cudaGetDeviceCount().......................................................................................................91

D.1.2 cudaSetDevice()..................................................................................................................91

D.1.3 cudaGetDevice()..................................................................................................................91

D.1.4 cudaGetDeviceProperties() .............................................................................91

D.1.5 cudaChooseDevice() .......................................................................................................93

- 5 -

D.2

线程管理

...................................................................................................................................... 93

D.2.1 cudaThreadSynchronize()..................................................................................................93

D.2.2 cudaThreadExit()..................................................................................................................93

D.3

流管理

........................................................................................................................................... 93

D.3.1 cudaStreamCreate() ...........................................................................................................93

D.3.2 cudaStreamQuery()...............................................................................................................93

D.3.3 cudaStreamSyncronize().....................................................................................................93

D.3.4 cudaStreamDestroy() .........................................................................................................94

D.4

事件管理

....................................................................................................................................... 94

D.4.1 cudaEventCreate()...............................................................................................................94

D.4.2 cudaEventRecord()................................................................................................................94

D.4.3 cudaEventQuery()...................................................................................................................94

D.4.4 cudaEventSyncronize()........................................................................................................94

D.4.5 cudaEventDestroy() .............................................................................................................95

D.4.6 cudaEventElapsedTime()......................................................................................................95

D.5

内存管理

........................................................................................................................................ 95

D.5.1 cudaMalloc() ..........................................................................................................................95

D.5.2 cudaMallocPitch().................................................................................................................95

D.5.3 cudaFree()................................................................................................................................96

D.5.4 cudaMallocArray()................................................................................................................96

D.5.5 cudaFreeArray().....................................................................................................................96

D.5.6 cudaMallocHost()...................................................................................................................96

D.5.7 cudaFreeHost().......................................................................................................................96

D.5.8 cudaMemSet() ..........................................................................................................................97

D.5.9 cudaMemSet2D().......................................................................................................................97

D.5.10 cudaMemcpy() .........................................................................................................................97

D.5.11 cudaMemcpy2D() ...................................................................................................................98

D.5.12 cudaMemcpyToArray() .........................................................................................................98

D.5.13 cudaMemcpy2DToArray().....................................................................................................99

D.5.14 cudaMemcpyFromArray().....................................................................................................99

D.5.15 cudaMemcpy2DFromArray()................................................................................................100

D.5.16 cudaMemcpyArrayToArray()..............................................................................................100

D.5.17 cudaMemcpy2DArrayToArray() ........................................................................................101

D.5.18 cudaMemcpyToSymbol().....................................................................................................101

D.5.19 cudaMemcpyFromSymbol()..................................................................................................101

D.5.20 cudaGetSymbolAddress()..................................................................................................102

D.5.21 cudaGetSymbolSize() .......................................................................................................102

- 6 -

D.6 Texture Reference

D.6.1

低级

API ................................................................................................................................... 102

管理

................................................................................................................102

D.6.1.1 cudaCreateChannelDesc() ............................................................................................102

D.6.1.2 cudaGetChannelDesc()...................................................................................................102

D.6.1.3 cudaGetTextureReference().........................................................................................103

D.6.1.4 cudaBindTexture().........................................................................................................103

D.6.1.5 cudaBindTextureToArray()..........................................................................................103

D.6.1.6 cudaUnBindTexture()......................................................................................................103

D.6.1.7 cudaGetTextureAlignmentOffset()............................................................................104

D.6.2

高级

API .....................................................................................................................................104

D.6.2.1 cudaCreateChannelDesc()...............................................................................................104

D.6.2.2 cudaBindTexture()............................................................................................................104

D.6.2.3 cudaBindTextureToArray().............................................................................................105

D.6.2.4 cudaUnBindTexture() ......................................................................................................105

D.7

执行控制

...................................................................................................................................... 105

D.7.1 cudaConfigureCall() .........................................................................................................105

D.7.2 cudaLaunch() ........................................................................................................................105

D.7.3 cudaSetupArgument() .........................................................................................................106

D.8 OpenGL Interoperability ............................................................................................................... 106

D.8.1 cudaGLRegisterBufferObject().......................................................................................106

D.8.2 cudaGLMapBufferObject()..................................................................................................106

D.8.3 cudaGLUnMapBufferObject() ............................................................................................106

D.8.4 cudaGLUnRegisterBufferObject()...................................................................................106

D.9 Direct3D Interoperability .............................................................................................................. 107

D.9.1 cudaD3D9Begin()...................................................................................................................107

D.9.2 cudaD3D9End().......................................................................................................................107

D.9.3 cudaD3D9RegisterVertexBuffer()...................................................................................107

D.9.4 cudaD3D9MapVertexBuffer() ............................................................................................107

D.9.5 cudaD3D9UnMapVertexBuffer().........................................................................................107

D.9.6 cudaD3D9UnRegisterVertexBuffer() ............................................................................107

D.9.7 cudaD3D9GetDevice() ........................................................................................................108

D.10

错误处理

......................................................................................................................................108

D.10.1 cudaGetLastError() .........................................................................................................108

D.10.2 cudaGetErrorString()......................................................................................................108

附录

E DriverAPIReference ................................................................................................................ 109

E.1

初始化

........................................................................................................................................... 109

E.1.1 cuInit()................................................................................................................................. 109

- 7 -

E.2

设备管理

....................................................................................................................................... 109

E.2.1 cuGetDeviceCount().............................................................................................................109

E.2.2 cuDeviceGet() ......................................................................................................................109

E.2.3 cuDeviceGetName()...............................................................................................................109

E.2.4 cuDeviceTotalMem().............................................................................................................109

E.2.5 cuDeviceComputeCapability() .......................................................................................110

E.2.6 cuDeviceGetAttribute()....................................................................................................110

E.2.7 cuDeviceGetProperties()..................................................................................................111

E.3 Context

管理

...................................................................................................................................111

E.3.1 cuCtxCreate() .....................................................................................................................111

E.3.2 cuCtxAttach() .....................................................................................................................112

E.3.3 cuCtxDetach() .....................................................................................................................112

E.3.4 cuCtxGetDevice().................................................................................................................112

E.3.5 cuCtxSynchronize().............................................................................................................112

E.4

模块管理

.......................................................................................................................................112

E.4.1 cuModuleLoad().....................................................................................................................112

E.4.2 cuModuleLoadData()............................................................................................................113

E.4.3 cuModuleLoadFatBinary().................................................................................................113

E.4.4 cuModuleUnload()................................................................................................................113

E.4.5 cuModuleGetFunction().....................................................................................................113

E.4.6 cuModuleGetGlobal() .......................................................................................................113

E.4.7 cuModuleGetTexRef() .......................................................................................................114

E.5

流管理

.........................................................................................................................................114

E.5.1 cuStreamCreate()...............................................................................................................114

E.5.2 cuStreamQuery().................................................................................................................114

E.5.3 cuStreamSynchronize()....................................................................................................114

E.5.4 cuStreamDestroy().............................................................................................................114

E.6

事件管理

.....................................................................................................................................114

E.6.1 cuEventCreate().................................................................................................................114

E.6.2 cuEventRecord().................................................................................................................115

E.6.3 cuEventQuery()...................................................................................................................115

E.6.4 cuEventSynchronize() ....................................................................................................115

E.6.5 cuEventDestroy()...............................................................................................................115

E.6.6 cuEventElapsedTime() ....................................................................................................115

E.7

执行控制

.....................................................................................................................................116

E.7.1 cuFuncSetBlockShape()....................................................................................................116

E.7.2 cuFuncSetSharedSize().....................................................................................................116

- 8 -

E.7.3 cuParamSetSize()................................................................................................................116

E.7.4 cuParamSeti() .....................................................................................................................116

E.7.5 cuParamSetf() .....................................................................................................................116

E.7.6 cuParamSetv() .....................................................................................................................116

E.7.7 cuParamSetTexRef()............................................................................................................117

E.7.8 cuLaunch().............................................................................................................................117

E.7.9 cuLaunchGrid().....................................................................................................................117

E.8 Memory

管理

..................................................................................................................................117

E.8.1 cuMemGetInfo().....................................................................................................................117

E.8.2 cuMemAlloc() .......................................................................................................................118

E.8.3 cuMemAllocPitch()...............................................................................................................118

E.8.4 cuMemFree()............................................................................................................................118

E.8.5 cuMemAllocHost().................................................................................................................118

E.8.6 cuMemFreeHost()...................................................................................................................119

E.8.7 cuMemGetAddressRange()....................................................................................................119

E.8.8 cuArrayCreate()...................................................................................................................119

E.8.9 cuArrayGetDescriptor()....................................................................................................120

E.8.10 cuArrayDestroy()...............................................................................................................121

E.8.11 cuMemset()........................................................................................................................... 121

E.8.12 cuMemset2D() .......................................................................................................................121

E.8.13 cuMemcpyHtoD()...................................................................................................................121

E.8.14 cuMemcpyDtoH()..................................................................................................................122

E.8.15 cuMemcpyDtoD()...................................................................................................................122

E.8.16 cuMemcpyDtoA()...................................................................................................................122

E.8.17 cuMemcpyAtoD()...................................................................................................................123

E.8.18 cuMemcpyAtoH()...................................................................................................................123

E.8.19 cuMemcpyHtoA()...................................................................................................................123

E.8.20 cuMemcpyAtoA()...................................................................................................................124

E.8.21 cuMemcpy2D() ......................................................................................................................124

E.9 Texture Reference

管理

.................................................................................................................126

E.9.1 cuTexRefCreate().................................................................................................................126

E.9.2 cuTexRefDestroy()...............................................................................................................127

E.9.3 cuTexRefSetArray().............................................................................................................127

E.9.4 cuTexRefSetAddress() .......................................................................................................127

E.9.5 cuTexRefSetFormat() ..........................................................................................................128

E.9.6 cuTexRefSetAddressMode()................................................................................................128

E.9.7 cuTexRefSetFilterMode()..................................................................................................128

- 9 -

E.9.8 cuTexRefSetFlags().............................................................................................................129

E.9.9 cuTexRefGetAddress() ......................................................................................................129

E.9.10 cuTexRefGetArray()...........................................................................................................129

E.9.11 cuTexRefGetAddressMode()..............................................................................................129

E.9.12 cuTexRefGetFilterMode()................................................................................................129

E.9.13 cuTexRefGetFormat() .......................................................................................................130

E.9.14 cuTexRefGetFlags()...........................................................................................................130

E.10 OpenGLInteroperability .............................................................................................................. 130

E.10.1 cuGLInit()........................................................................................................................... 130

E.10.2 cuGLRegisterBufferObject() ........................................................................................130

E.10.3 cuGLMapBufferObject()....................................................................................................130

E.10.4 cuGLUnmapBufferObject()................................................................................................131

E.10.5 cuGLUnregisterBufferObject().....................................................................................131

E.11 Direct3DInteroperability ............................................................................................................. 131

E.11.1 cuD3D9Begin().....................................................................................................................131

E.11.2 cuD3D9End()..........................................................................................................................131

E.11.3 cuD3D9RegisterVertexBuffer() ...................................................................................131

E.11.4 cuD3D9MapVertexBuffer()................................................................................................131

E.11.5 cuD3D9UnmapVertexBuffer()...........................................................................................132

E.11.6 cuD3D9UnregisterVertexBuffer()................................................................................132

E.11.7 cuD3D9GetDevice()............................................................................................................132

附录

F TextureFetching .................................................................................................................... 133

F.1 Nearest-Point

F.2

线性过滤

F.3

查表法

........................................................................................................................................... 136

采样

..................................................................................................................... 134

..................................................................................................................................... 135

- 10 -

Chapter 1 介绍CUDA

1.1

作为一个并行数据计算设备的图形处理器单元仅仅几年的时间，可编程的图形处理器单元演变成为了

一匹绝对的计算悍马，正如图

1-1

带宽驱动多核处理器时，当今的

所示。当极高的内存

GPU

为图型和非图型处理提供了难以置信的资源。

图

这个演变背后的主要原因是由于

1-1。CPU 和GPU

GPU

被设计用于高密度和并行计算，更确切地说是用于图形渲染。因此

的每秒浮点运算

更多的晶体管被投入到数据处理而不是数据缓存和流量控制，图

图

1-2。 GPU

投入更多晶体管进行数据处理

1-2

所示。

- 11 -

更加具体地看，GPU 是特别适合于并行数据运算的问题－同一个程序在许多并行数据元素，

并带有高运算密度（算术运算与内存操作的比例）。由于同一个程序要执行每个数据元素，

降低了对复杂的流量控制要求; 并且，因为它执行许多数据元素并且据有高运算密度，内存

访问的延迟可以被忽略。

并行数据处理，意味着数据元素以并行线程处理。许多处理大量数据集，例如数组的应用程

序可以使用一个并行数据的编程模型来加速计算。在3D 渲染上，大的像素集和顶点被映射

到并行线程。同样，图像和媒体处理的应用程序例如着色的图像后处理，录像编码和解码，

图像缩放比例，立体视觉，以及图像识别也可以映射图像块和像素到并行处理线程。实际上，

在图像着色和处理领域外的许多算法同样可以通过并行数据处理得到加速，从一般信号处理

或物理模拟到金融计算或者生物计算。

然而直到今天，尽管强大的计算能力包装进了GPU，而它对非图形应用的有效支持依然有限：

 GPU 只能通过图型API 来编程，导致新手很难学习和非图形API 上很不充分的应用。

 GPU DRAM 可以用一般方式下读取，GPU 程序可以从任何DRAM 部分收集数据元素。

但不可写，在一般方式下的GPU 程序不能写入信息到DRAM 的任何部分，相比CPU 丧失

了很多编程的灵活性。



有些应用是由于DRAM 内存带宽而形成的瓶颈，未能充分利用GPU 的计算能力。

本文描述的是一个崭新的硬件和编程模型，它直接答复了这些问题并且展示了GPU 如何成

为一个真正的通用并行数据计算设备。

1.2 CUDA

备架构，在

构，它不需要映射到一个图型

都可以支持。操作系统的多任务机制通过几个

：一个在

GPU

上发布的一个新的硬件和软件架

GPU

上计算的新架构

API

便可在

CUDA（Compute Unified Device Architecture

GPU

上管理和进行并行数据计算。从

CUDA

和图型应用程序协调运行来管理

G80

）统一计算设

系列和以后的型号

访问

GPU

。

CUDA 软件堆栈由几层组成，如图1-3 所示：一个硬件驱动程序，一个应用程序编程接口(API)

和它的Runtime，还有二个高级的通用数学库，CUFFT 和CUBLAS。硬件被设计成支持轻

量级的驱动和

Runtime 层面，因而提高性能。

- 12 -

CUDA API

更像是C 语言的扩展，以便

散” 和“聚集”内存操作，如图

的任

何区域

进行读写数据的操作，

图

1-3

1-4

所示。从而提供

就像

。统一计算设备架构的软件

最小

化学习的时间。

在

CPU

上一样。

最大

的编程

CUDA

灵活

堆栈

提供一

般

DRAM

性。从编程的观点

内存

来看

，它可以在

寻址

方式：“发

DRAM

图

1-4。 “

聚集”和“发

散”内存操作

- 13 -

CUDA

如图

1-5

带宽的

允许

并行数据缓冲或者在

所示，应用程序可以

依赖

。

On-chip

最小

化数据到

内存共享，可以进行

快速

的常规读写存取，在线程之间共享数据。

DRAM 的overfetch 和round-trips

，从而

减少

对

DRAM

内存

图1-5。共享内存使数据更接近ALU

- 14 -

Chapter 2 编程模型

2.1

一个超多线程协处理器

当通过

协处理器。换

CUDA

编译时，

句话

GPU

可以被视为能执行非常高数量并行线程的计算设备。它作为主

说，运行在主机上的并行数据和高密度计算应用程序

更准确地讲，一个被执行许多次不同数据的应用程序

执行的函数。达到这个

主机和设备使用它

的

引擎

（

API

）从一个

效果

们自己

DRAM

，这个函数被编译成设备的指令集（

的

DRAM

，主机内存和设备内存。并可以通过利用设备高性能直接内存存取

复

制数据到

其他

DRAM

2.2

线程批处理

线程批处理就是执行一个被组织成许多线程块的

kernel

部分

。

，如图

，可以被

kernel

2-1

CPU

部分

，被

卸载

到这个设备上。

分离

成为一个有很多不同线程在设备上

程序），被

下载

到设备上。

所示。

的一个

(DMA)

主机发送一个

连续

的

kernel

调用到设备。每个

图

2-1

kernel

作为一个由线程块组成的批处理线程来执行。

。线程批处理

- 15 -

2.2.1

线程块

一个线程块是一个线程的批处理，它通过一

步它们的执行。更准确地说，它可以在

Kernel

同步点。

每条线程是由它的线程

一个应用程序可以指定一个块作为一个

指定每条线程。对于一个

一个

三维的大小

为

(

，

大小

所确定，

为

(

，

)

的块，这个线程的

是在块之内的线程编号。根据线程的

二维或三维

，

)

2.2.2

线程块栅格

些快速

二维

的共享内存有效地分享数据并且在制定的内存

中

指定同步点，一个块里的线程被

数组的任

块，线程的

索引

意大小，并且

索引

是

(

x，y，z

(x, y)

通过一个

，这个线程

)

，线程的

挂起直

可以

ID 是(x +

到它们所有都到

帮助

进行

2 -或3-

组件

y D

访问中

复杂寻址

索引代替来

)

。而对于

z DxD

，

)

同

达

。

一个块可以

此通过单一

自

同一个栅格的不同线程块中的线程彼此之不间能通讯和同步。这个模式

效

地运行在

性，或者并行地运行，如果它有很多的并行的特性，或者通常是

包含

的线程

kernel

各种

设备上而不用再编译：一个设备可以序列地运行栅格的所有块，如果它有非常少的并行

最大

发送的

数量是有限的。然而，执行同一个

请求

的线程总数可以是非常

巨大

kernel

的块可以合成一批线程块的栅格，因

的。线程协作的

二者

的组合。

减少会造

允许

kernel

成性能的

损失

用不同的并行能力有

，因为

来

特

每个块是由它的块ID 确定的，块的ID 是在栅格之内的块编号。根据块ID 可以帮助进行

复杂寻址，一个应用程序可也以指定一个栅格作为任意大小的一个二维数组，并且通过一个

2-组件索引替换来制定每个块。对于一个大小为 (

块的

ID 是(x +

y D

)

。

，

)

二维

块，这个块的索引是(x，y)，

- 16 -

2.3

内存模型

一条执行在设备上的线程，

只允许

通过如下的内存空间使用设备的

DRAM 和On-Chip

示：



读写每条线程的寄存器，



读写每条线程的本地内存，



读写每个块的共享内存，



读写每个栅格的全局内存，



只

读每个栅格的常量内存，



只

读每个栅格的纹理内存。

全局，常量，和纹理内存空间可以通过主机或者同一应用程序持续的通过

kernel

调用

全局，常量，和纹理内存空间对不同内存的用法加以优化。纹理内存同样提供不同的

特殊

的数据格式进行数据过滤。

内存，如图

来完

成读取或写入。

寻址

模式，也为一

2-2

所

些

线程通过一组不同

范围

的内存空间来使用设备的

图

2-2

。内存模型

DRAM 和On-chip

内存。

- 17 -

Chapter 3 硬件实现

3.1

一组带有

on-chip

共享内存的

SIMD

多处理器

设备可以被看作一组多处理器，如图

何给

定的时

钟周期

内，多处理器的每个处理器执行同一指令，但操作不同的数据。

每个多处理器使用四个以下类型的



每个处理器一组本地



并行数据缓存或共享内存，被所有处理器共享实现内存空间共享，



通过设备内存的一个只读



通过设备内存的一个只读

本

地和全局内存空间作为设备内存的读

每个多处理器通过纹理单元

位寄存器，

区域区域

访问

3-1

所示。每个多处理器使用单一指令，多数据架构

on-chip

内存：

，一个只读常量缓冲器被所有处理器共享，，一个只读纹理缓冲器被所有处理器共享，

写区域

纹理缓冲器，它执行

，而不被缓冲。

各种各

样的

寻址

模式和数据过滤。

(SIMD)

：在任

图

3-1

。硬件模型

- 18 -

3.2

执行模式

一个线程块栅格是通过多处理器规划执行的。每个多处理器一个接一个的处理块批处理。一个块只被一个

多处理器处理，因此可以对驻留在

一个批处理中每一个多处理器可以处理多少个块，取决于每个线程

on-chip

共享内存中的共享内存空间形成非常

中分

配了多少个寄存器和

个时钟需要多少的共享内存，因为多处理器的寄存器和内存在所有的线程中是

中

，每个多处理器没有

足够

的寄存器或共享内存可用，

那么

内核

将无法启

动。

线程块在一个批处理中被一个多处理器执行，被称作

为

warps;

执行; 线程调度程序

每一条这样的

周期

warp

包含

性地从一

数量相同的线程，

条

warp

切换到另一

active

叫做

条

warp

。每个

warp

active

大小

块被

，并且在

，以达到多处理器计算资源使用的

块被

划分

成为

warp

程0 开始递增

的方式总是相同的; 每

。

条

warp

包含连续

的线程，线程

索引从第

快速的访问

分开

的。如果在

划分

成为

SIMD

一个

。

已知

内核中每

至少

一个块

SIMD

线程组，

称

方式下通过多处理器

最大

化。

warp

包含着

的线

一个多处理器可以处理并发地几个块，通过

划分

在它

们之中

的寄存器和共享内存。更准确地说，每条线程

可使用的寄存器数量，等于每个多处理器寄存器总数除以并发的线程数量，并发线程的数量等于并发块的

数量乘以每块线程的数量

。

在一个块内的warp 次序是未定义的，但通过协调全局或者共享内存的存取，它们可以同步

的执行。如果一个通过warp 线程执行的指令写入全局或共享内存的同一位置，写的次序是

未定义的。

在一个线程块栅格内的块次序是未定义的，并且在块之间不存在同步机制，因此

同块的线程不能通过全局内存彼此安全地通讯。

来自

同一个栅格的二个不

- 19 -

3.3

计算兼容性

设备的计算兼容性由两个参数定义，主要

心

架构。在附录A 中列出

次要版本

3.4

号代表一

多设备

些改

的设备全部是

进的核心架构，比如新的特性。不同计算兼容性的技术规格见附录A。

版本

号和次要

1.x

（它们的主要

版本

号。设备拥有相同的主要

版本

号是1）。

版本

号代表相同的核

为一个应用程序使用多GPU 作为CUDA 设备，必须保证这些GPU 是一样的类型。如果系

统工作在SLI 模式下，那么只有一个GPU 可以作为CUDA 设备，由于所有的GPU 在驱动

堆栈中被底层的融合了。SLI 模式需要在控制面板中关闭，这样才能事多个GPU 作为CUDA

设备。

3.5

模式切换

GPU 指定一些DRAM 来存储被称作primary surface 的内容，这些内容被用于显示输出。如果

用户改变显示的分辨率或者色深，那么primary surface 的存储需求量将改变。比如，如果

用户将显示分辨率从1280x1024x32bit 到1600x1200x32bit ，系统必须指定7.68MB 的

primary surface 而不在是5.24MB。（使全屏抗锯齿的应用程序需要更多的primary surface

空间）。另外，比如在Windows 中使用Alt+Tab 的切换，或者Ctrl+Alt+Del 的操作同样需要

额外的primary surface 空间。

如果模式切换增加了primary surface 的内存空间，系统将占用CUDA 指定的内存空间，导

致程序崩溃。

- 20 -

Chapter 4 应用程序编程接口（API）

4.1

一个C 语言的扩展

CUDA

它



编程接口的目标是为

包括

：

一个小的C 语言扩展集，在

熟悉C 语言的用户提供一个相对简单的

第

4.2

部分描述，允许

程序

员专注

途径来编写

于在设备执行的原代码的

设备执行程序。

部分

;

 一个runtime 库分成：

 一个主机组件，在第4.5 部分描述，它在主机上运行并且提供函数来控制和访问一个

或多个计算设备;

 一个设备组件，在第4.4 部分描述，它在设备运行并且提供特定设备的函数;

 一个公共的组件，在第4.3 部分描述，它提供内置矢量类型和主机与设备编码都支持

的C 标准库的一个子集。

应该强调的是，只有来自C 标准库的函数支持在设备上运行，是由公共Runtime 的组件提供

的函数。

4.2 语言扩展

C 语言的扩展是四重的：

 函数类型限定句指定一个函数是否执行在主机或者执行在设备，和是否从主机或者从设

备上调用(第4.2.1 部分);

 变量类型限定句指定设备上一个变量的内存位置(第4.2.2 部分);

 一个新的指令指定一个来自主机的kernel 如何在设备上执行 (第4.2.3 部分);

 四个内置的变量指定栅格和块的维数，还有块和线程的指标 (第4.2.4 部分)。

每个包含CUDA 语言扩展的源文件必须通过CUDA 编译器nvcc 编译，在第4.2.5 部分简要

地加以描述。 nvcc 的一个详细的描述可以在一单独的文件中找到。

每一个扩展

一个

警告，但有些违反无

伴随

的一些限定在每一

节下面给予描述

法被查出。

。

nvcc

将

提供在一

些违反这些

限制时的一个错误或

者

- 21 -

4.2.1

函数类型限定词

4.2.1.1 __device__

__device__限定词声明一个函数是：

 在设备上执行的，

 仅可从设备调用。

4.2.1.2 __global__

__global__



在设备上执行的，



仅可从主机调用。

4.2.1.3 __host__

__host__

__host__限定词声明的函数是：

__host____host__

限定词声明一个函数作为一个存在的

__host__

__host____host__

kernel

。这样的一个函数是：

 在主机上执行的，

 仅可从主机调用。

它等于声明一个函数仅带有__host__

__global__

__global__限定词; 在其他情况下这个函数仅为主机编译。

__global____global__

然而，__host__

__host__ 限定词也可以用于与__device__

__host__ __host__

__host__限定词或者声明它没有任何__ho

__host____host__

__device__限定词的组合，这种的情况下，这个函数是

__device____device__

为主机和设备双方编译。

__host__

st__， __device__

__ho__ho

st__st__

__device__，或

__device____device__

4.2.1.4

__device__和__global__

__device__

的。

不能一起使用__global__和__host__限定词。

限定

函数不可能取得它们的地

函数不支持

函数不能声明

函数不能有自变量的一个变量数字。

递归

静态

址; 另

。

变量在它们体内。

一方面，函数指

- 22 -

向

__global__

函数是支持

__global__

在

第

4.2.3

__global__

4.2.2

变量类型限定词

4.2.2.1 __device__

__device__

最

多的一个其它类型限定词被定义在

属在

哪些



驻留在全局内存空间，



具

有应用的生存期，



从栅格内所有线程和从主机通过

函数

必须

有

void

的

部分所描述

函数参数

限定词声明驻留在设备上的一个变量。

内存空间。如果它们都不存在，这个变量：

。对一个

目前

是通过共享内存到设备的，并且被限制在

类型。任何调用到一个

__global__

下面的三项里

runtime

函数的调用是同步的，

库

，可以与

是可

访问

__global__

__device__

的。

函数

必须

指定它的执行配置，如

意味着

在设备执行完成

256

个

字节

。

一起共同用于进一步指定变量

前返回

。

归

4.2.2.2 __constant__

__constant__



驻留在常量内存空间，



具

有应用的生存期，



从栅格内所有线程和从主机通过

4.2.2.3 __shared__

__shared__

__shared__限定词，与__device__

__shared____shared__

__shared__

__shared__ __shared__

限定词，与

__device__一起选择使用，声明一个变量：

__device____device__

__device__

runtime

一

起随

机使用，声明一变量：

库

的是可

访问

的。

 驻留在线程块的共享内存空间中，

 具有块的生存期，

 只有块之内的所有线程是可访问的。

在线程中共享的变量有完全的顺序一致性。只有执行过一个__syncthreads

__syncthreads（）函数，从其他

__syncthreads__syncthreads

线程的写才保证可见。除非变量被定义为可挥发的，否则只要前一个状态到达，编译器将自

由的优化共享内存中的读写。

- 23 -

当声明一个在共享内存的变量作为一个外部数组时，例如

extern_shared_float shared[]；

数组的地址，因此在数组的变量布局

大小

是由发送时间(参见第

必须

4.2.3

通过

部分)决

offset(

定的。所有变量用这种方式声明的，开始于内存的同一个

位移量)明确地加以控制。例如，如

short array0[128];

float array1[64];

int array2[256];

在动

态分

配的共享内存，你可以用以下方式定义数组

extern __shared__ short array[];

__device__ void func() //__device__or__global__function

{

Short* array0 = (short*)array;

float* array1 = (float*)&array0[128];

int* array2 = (int*)&array1[64];

}

果你想要等

于

4.2.2.4

这些限定词不允许的一个函数之内的

这些限定词不允许的一个函数之内的这些限定词不允许的一个函数之内的

__shared__和__constant__

__device__,__shared__和__constant__

__device__和__constant__

__constant__

限定

限定限定

变量不能从设备上

struct 和union

变量

隐含了静态存储

变量

只允许

赋值

，仅可以通过主机

。

变量不能被用

在文件

范围

成员，形式参数和局部变量在主机上执行。

。

extern

runtime

关键字定义为外部

函数从主机上

4.5.3.6）。

__shared__

变量不能作为它们声明的一

部分得

到初始化。

一个不用任何限定词在设备码中声明的自变量通常驻留在一个寄存器中。尽管在

择

在局部内存确定数组用常量数量如果变量在第一个编译

ld.local 和st.local

器空间，随后编译

报告

（

lmem）。

里安

置它。通常在这个

建立了索引

阶段期

阶段仍然

情况下，大

。

ptx

汇

型结构或者数组

编代码的检查(通过

将消耗许

- ptx 或-keep

间被安置在局部内存，因此它将声明是使用

助记符访问

可以

做出另外的决

的。假如它没有这样做，即使它们发现它为目标架构

定。可以使用

--ptxas-option=-v

多寄存器空间，并且编译器不能

赋值（参见

某些情况下

选项

编译

.local

生

成局部内存使用量

使用。

4.5.2.3

编译器可以

获得

的)将指出

助记符或者

消耗太

多寄存

和

选

，

使用

- 24 -

执行在设备上的指针代码支持，当编译器可以解析它们是否指向全局内存空间或者局部内存空间。

们将

被限制只指

解除一个指针在主机上执行的全局或共享内存中的代码，或

一个未定义的行为，从而

向寻址

的内存或在全局内存中声明的空间。

产生片断

错误或者应用

终止

。

者载

设备上执行的主机内存代码，其结

否则

果产生

只能通过设备代码中的__device__,__shared__或__constant__变量来获取地址。

__device__或__constant__变量的地址只能通过主机代码获得，cudaGetSymbolAddress()

参见4.5.2.3 部分。

4.2.3

执行配置

所有

__global__

执行配置定义了通常在设备执行的函数的栅格和块的维数，同样

函数的调用

必须

指定执行配置。