《并行与分布式程序设计》课程教学参考书：CUDA C PROGRAMMING（CUDA编程指南4.0中文版）

CUDA编程指南4.0中文版 译者：风辰 由于小弟的水平所限，此文档可能存在错误，如果你觉得本文档的 某些内容可能是错误，请联系我，谢谢！ 由于这样或者那样的原因，此翻译版将可能会是“绝版”，谢谢大家 的支持。 任何人不得更改此文档内容或设置，更不能用于商业目的，本人不 承担任何责任

CUDA 编程指南 4.0 中文版 译者：风辰 由于小弟的水平所限，此文档可能存在错误，如果你觉得本文档的 某些内容可能是错误，请联系我，谢谢！ 由于这样或者那样的原因，此翻译版将可能会是“绝版”，谢谢大家 的支持。 任何人不得更改此文档内容或设置，更不能用于商业目的，本人不 承担任何责任

目录 第一章导论 1.1从图形处理到通用并行计算 1 1.2 CUDATM:一种通用并行计算架构 3 1.3一种可扩展的编程模型 1.4文档结构 第二章编程模型 7 2.1内核 2.2线程层次 .7 2.3存储器层次 10 2.4异构编程 11 2.5计算能力.… 13 第三章编程接口.… 15 3.1用nvcc编译 15 3.11编译流程.… 15 3.1.1.1离线编译 15 3.1.1.2即时编译 16 3.1.2二进制兼容性 。+ 17 3.1.3PTX兼容性 17 3.1.4应用兼容性 17 3.1.5C/C+兼容性 18 3.1.664位兼容性 18 3.2 CUDA C运行时 18 3.2.1初始化 19 3.2.2设备存储器 19 3.2.3共享存储器 22 3.2.4分页锁定主机存储器 28 3.2.4.1可分享存储器（portable memory) 28 3.2.4.2写结合存储器 28 3.2.4.3被映射存储器 28 3.2.5异步并发执行 29 3.2.5.1主机和设备间异步执行 29 3.2.5.2数据传输和内核执行重叠 30 3.2.5.3并发内核执行… 30

目录 第一章 导论......................................................................................................................... 1 1.1 从图形处理到通用并行计算................................................................................ 1 1.2 CUDATM：一种通用并行计算架构...................................................................... 3 1.3 一种可扩展的编程模型........................................................................................ 3 1.4 文档结构................................................................................................................ 5 第二章 编程模型................................................................................................................. 7 2.1 内核........................................................................................................................ 7 2.2 线程层次................................................................................................................ 7 2.3 存储器层次.......................................................................................................... 10 2.4 异构编程.............................................................................................................. 11 2.5 计算能力.............................................................................................................. 13 第三章 编程接口............................................................................................................... 15 3.1 用 nvcc 编译......................................................................................................... 15 3.1.1 编译流程................................................................................................... 15 3.1.1.1 离线编译........................................................................................ 15 3.1.1.2 即时编译........................................................................................ 16 3.1.2 二进制兼容性........................................................................................... 17 3.1.3 PTX 兼容性................................................................................................ 17 3.1.4 应用兼容性............................................................................................... 17 3.1.5 C/C++兼容性 ............................................................................................. 18 3.1.6 64 位兼容性............................................................................................... 18 3.2 CUDA C 运行时 ................................................................................................... 18 3.2.1 初始化....................................................................................................... 19 3.2.2 设备存储器............................................................................................... 19 3.2.3 共享存储器............................................................................................... 22 3.2.4 分页锁定主机存储器............................................................................... 28 3.2.4.1 可分享存储器（portable memory）.............................................. 28 3.2.4.2 写结合存储器................................................................................ 28 3.2.4.3 被映射存储器................................................................................ 28 3.2.5 异步并发执行........................................................................................... 29 3.2.5.1 主机和设备间异步执行................................................................ 29 3.2.5.2 数据传输和内核执行重叠............................................................ 30 3.2.5.3 并发内核执行................................................................................ 30

3.2.5.4并发数据传输 30 3.2.5.5流 30 3.2.5.6事件 33 3.2.5.7同步调用 34 3.2.6多设备系统 34 3.2.6.1枚举设备 34 3.2.6.2设备选择 34 3.2.6.3流和事件行为 35 3.2.6.4p2p存储器访问 35 3.2.6.5p2p存储器复制 36 3.2.7统一虚拟地址空间 37 3.2.8错误检查 37 3.2.9调用栈 38 3.2.10纹理和表面存储器 38 3.2.10.1纹理存储器 38 3.2.10.2表面存储器(surface) 44 3.2.10.3CUDA数组 46 3.2.10.4读写一致性 47 3.2.11图形学互操作性. 47 3.2.11.1 OpenGL互操作性 47 3.2.11.2 Direct3.D互操作性 50 3.2.11.3SLI(速力)互操作性 58 3.3驱动API 58 3.3.1上下文 61 3.3.2模块 62 3.3.3内核执行 63 3.34设备存储器 65 3.3.5共享存储器 69 3.3.6分页锁定主机存储器 71 3.3.7异步并发执行 71 3.3.7.1流. 72 3.3.7.2事件.… 72 3.3.7.3同步调用 73

4 3.2.5.4 并发数据传输................................................................................ 30 3.2.5.5 流.................................................................................................... 30 3.2.5.6 事件................................................................................................ 33 3.2.5.7 同步调用........................................................................................ 34 3.2.6 多设备系统............................................................................................... 34 3.2.6.1 枚举设备........................................................................................ 34 3.2.6.2 设备选择........................................................................................ 34 3.2.6.3 流和事件行为................................................................................ 35 3.2.6.4 p2p 存储器访问.............................................................................. 35 3.2.6.5 p2p 存储器复制.............................................................................. 36 3.2.7 统一虚拟地址空间................................................................................... 37 3.2.8 错误检查................................................................................................... 37 3.2.9 调用栈....................................................................................................... 38 3.2.10 纹理和表面存储器................................................................................. 38 3.2.10.1 纹理存储器.................................................................................. 38 3.2.10.2 表面存储器(surface).................................................................... 44 3.2.10.3 CUDA 数组.................................................................................. 46 3.2.10.4 读写一致性.................................................................................. 47 3.2.11 图形学互操作性..................................................................................... 47 3.2.11.1 OpenGL 互操作性........................................................................ 47 3.2.11.2 Direct3D 互操作性....................................................................... 50 3.2.11.3 SLI（速力）互操作性................................................................. 58 3.3 驱动 API .............................................................................................................. 58 3.3.1 上下文....................................................................................................... 61 3.3.2 模块........................................................................................................... 62 3.3.3 内核执行................................................................................................... 63 3.3.4 设备存储器............................................................................................... 65 3.3.5 共享存储器............................................................................................... 69 3.3.6 分页锁定主机存储器............................................................................... 71 3.3.7 异步并发执行........................................................................................... 71 3.3.7.1 流.................................................................................................... 72 3.3.7.2 事件................................................................................................ 72 3.3.7.3 同步调用........................................................................................ 73

目录 33.8多设备系统 73 3.3.8.1设备枚举 73 3.3.8.2p2p存储器访问 74 3.3.8.3p2p存储器拷贝 74 3.3.9统一虚拟地址空间 75 3.3.10错误检查 75 3.3.11调用栈 76 3.3.12纹理存储器和表面存储器 76 3.3.12.1纹理存储器 76 3.3.12.2表面存储器 78 3.3.13图形学互操作性… 80 3.3.13.1 OpenGL互操作性 80 3.3.10.2 Direct3.D互操作性 82 3.4运行时API和驱动API的互操作性 91 3.5版本和互操作性 91 3.6计算模式 92 3.7模式切换 92 3.8 Windows上的Tesla计算集群模式 93 第四章硬件实现… 95 4.1SIMT架构. 95 4.2硬件多线程 96 第五章性能指南… 99 5.1总体性能优化策略 99 5.2最大化利用率 99 5.2.1应用层次… 99 5.2.2设备层次 99 5.2.3多处理器层次 100 5.3最大化存储器吞吐量 101 5.3.1主机和设备的数据传输 102 5.3.2设备存储器访问 103 5.3.2.1全局存储器 103 5.3.2.2本地存储器 104 5.3.2.3共享存储器 .105

目录 3.3.8 多设备系统............................................................................................... 73 3.3.8.1 设备枚举........................................................................................ 73 3.3.8.2 p2p 存储器访问.............................................................................. 74 3.3.8.3 p2p 存储器拷贝.............................................................................. 74 3.3.9 统一虚拟地址空间................................................................................... 75 3.3.10 错误检查................................................................................................. 75 3.3.11 调用栈 ..................................................................................................... 76 3.3.12 纹理存储器和表面存储器..................................................................... 76 3.3.12.1 纹理存储器.................................................................................. 76 3.3.12.2 表面存储器.................................................................................. 78 3.3.13 图形学互操作性..................................................................................... 80 3.3.13.1 OpenGL 互操作性........................................................................ 80 3.3.10.2 Direct3D 互操作性 ....................................................................... 82 3.4 运行时 API 和驱动 API 的互操作性 ................................................................. 91 3.5 版本和互操作性.................................................................................................. 91 3.6 计算模式.............................................................................................................. 92 3.7 模式切换.............................................................................................................. 92 3.8 Windows 上的 Tesla 计算集群模式..................................................................... 93 第四章 硬件实现............................................................................................................... 95 4.1 SIMT 架构............................................................................................................ 95 4.2 硬件多线程.......................................................................................................... 96 第五章 性能指南............................................................................................................... 99 5.1 总体性能优化策略.............................................................................................. 99 5.2 最大化利用率...................................................................................................... 99 5.2.1 应用层次................................................................................................... 99 5.2.2 设备层次................................................................................................... 99 5.2.3 多处理器层次......................................................................................... 100 5.3 最大化存储器吞吐量........................................................................................ 101 5.3.1 主机和设备的数据传输......................................................................... 102 5.3.2 设备存储器访问..................................................................................... 103 5.3.2.1 全局存储器.................................................................................. 103 5.3.2.2 本地存储器.................................................................................. 104 5.3.2.3 共享存储器.................................................................................. 105

6 5.3.2.4常量存储器 105 5.3.2.5纹理和表面存储器 106 5.4最大化指令吞吐量 106 5.4.1算术指令 106 5.4.2控制流指令 109 5.4.3同步指令 109 附录A支持CUDA的GPU 113 附录BC语言扩展.… 117 B.1函数类型限定符 .117 B.1.1 device 117 B.1.2 global .117 B.1.3 host 117 B.1.4_noinline和 forceinline 118 B.2变量类型限定符 .118 B.2.1 device 118 B.2.2 constant .118 B.2.3 shared .119 B.2.4 restrict 120 B.3内置变量类型 .121 B.3.1 charl、uchar1、char2、uchar.2、char3、uchar.3、char4、uchar44、short1、 ushort1、short2、ushort.2、short.3、ushort3.、shorta4、ushort44、intl、uintl、int2、 uint2、int3、uint3、int4、uint4、long1、ulongl、long2、ulong2、long3、ulong3、 long4、ulong4、float1、float2、float3、float4、double2 121 B.3.2dim3类型 122 B.4内置变量 122 B.4.1 gridDim........... 122 B.4.2 blockIdx.... 122 B.4.3 blockDim 122 B.4.4 threadIdx 122 B.4.5 warpSize 122 B.5存储器栅栏函数 122 B.6同步函数 124 B.7数学函数.… 125

6 5.3.2.4 常量存储器.................................................................................. 105 5.3.2.5 纹理和表面存储器...................................................................... 106 5.4 最大化指令吞吐量............................................................................................ 106 5.4.1 算术指令................................................................................................. 106 5.4.2 控制流指令............................................................................................. 109 5.4.3 同步指令................................................................................................. 109 附录 A 支持 CUDA 的 GPU ...........................................................................................113 附录 B C 语言扩展 ...........................................................................................................117 B.1 函数类型限定符 ................................................................................................117 B.1.1 __device__................................................................................................117 B.1.2 __global__ ................................................................................................117 B.1.3 __host__....................................................................................................117 B.1.4 _noinline_ 和 __forceinline__................................................................118 B.2 变量类型限定符 ................................................................................................118 B.2.1 __device__................................................................................................118 B.2.2 __constant__.............................................................................................118 B.2.3 __shared__................................................................................................119 B.2.4 __restrict__ .............................................................................................. 120 B.3 内置变量类型 ................................................................................................... 121 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 ............................ 121 B.3.2 dim3 类型................................................................................................ 122 B.4 内置变量 ........................................................................................................... 122 B.4.1 gridDim.................................................................................................... 122 B.4.2 blockIdx ................................................................................................... 122 B.4.3 blockDim ................................................................................................. 122 B.4.4 threadIdx.................................................................................................. 122 B.4.5 warpSize .................................................................................................. 122 B.5 存储器栅栏函数 ............................................................................................... 122 B.6 同步函数 ........................................................................................................... 124 B.7 数学函数 ........................................................................................................... 125

目录 B8纹理函数… 125 B.8.1 tex1Dfetch() 125 B.8.2 tex1D() 126 B.8.3 tex2D() 126 B.8.4 tex3D() 126 B.8.5 texIDLayered() ,126 B.8.6 tex2DLayered() 127 B.9表面函数(surface) 127 B.9.1 surf1Dread() 127 B.9.2 surf1Dwrite() 127 B.9.3 surf2Dread() 127 B.9.4 surf2Dwrite() 128 B.10时间函数 128 B.11原子函数 128 B.11.1数学函数. . 129 B.11.1.1 atomicAdd() 129 B.11.1.2 atomicSub().. 129 B.11.1.3 atomicExch() 129 B.11.1.4 atomicMin(...... 130 B.11.1.5 atomicMax(). 130 B.11.1.6 atomicInc() 130 B.11.1.7 atomicDec(). 130 B.11.1.8 atomicCAS() 131 B.11.2位逻辑函数 131 B.11.2.1 atomicAnd(). 131 B.11.2.2 atomicOr(). 131 B.11.2.3 atomicXor(). 131 B.12束表决(warp vote)函数 131 B.13取样计数器函数 132 B.14格式化输出.… 132 B.14.1格式化符号 133 B.14.2限制. 133 B.14.3相关的主机端API. .134

目录 B.8 纹理函数 ........................................................................................................... 125 B.8.1 tex1Dfetch()............................................................................................. 125 B.8.2 tex1D()..................................................................................................... 126 B.8.3 tex2D()..................................................................................................... 126 B.8.4 tex3D()..................................................................................................... 126 B.8.5 tex1DLayered()........................................................................................ 126 B.8.6 tex2DLayered()........................................................................................ 127 B.9 表面函数(surface) ............................................................................................. 127 B.9.1 surf1Dread()............................................................................................. 127 B.9.2 surf1Dwrite() ........................................................................................... 127 B.9.3 surf2Dread()............................................................................................. 127 B.9.4 surf2Dwrite() ........................................................................................... 128 B.10 时间函数 ......................................................................................................... 128 B.11 原子函数.......................................................................................................... 128 B.11.1 数学函数............................................................................................... 129 B.11.1.1 atomicAdd() ................................................................................ 129 B.11.1.2 atomicSub()................................................................................. 129 B.11.1.3 atomicExch()............................................................................... 129 B.11.1.4 atomicMin() ................................................................................ 130 B.11.1.5 atomicMax()................................................................................ 130 B.11.1.6 atomicInc().................................................................................. 130 B.11.1.7 atomicDec()................................................................................. 130 B.11.1.8 atomicCAS() ............................................................................... 131 B.11.2 位逻辑函数........................................................................................... 131 B.11.2.1 atomicAnd() ................................................................................ 131 B.11.2.2 atomicOr()................................................................................... 131 B.11.2.3 atomicXor()................................................................................. 131 B.12 束表决（warp vote）函数.............................................................................. 131 B.13 取样计数器函数 ............................................................................................. 132 B.14 格式化输出 ..................................................................................................... 132 B.14.1 格式化符号 .......................................................................................... 133 B.14.2 限制 ...................................................................................................... 133 B.14.3 相关的主机端 API............................................................................... 134

8 B.14.4例程 134 B.15动态全局存储器分配 135 B.15.1堆存储器分配 136 B.15.2与设备存储器API的互操作 136 B.15.3例程 136 B.15.3.1每个线程的分配 136 B.15.3.2每个线程块的分配 137 B.15.3.3在内核启动之间持久的分配 138 B.16执行配置 140 B.17发射绑定 140 B.18 #pragma unroll. 143 附录C数学函数. .145 C.1标准函数 145 C.1.1单精度浮点函数 145 C.1.2双精度浮点函数 148 C.1.3整型函数 150 C.2内置函数 150 C.2.1单精度浮点函数 151 C.2.2双精度浮点函数 152 C.2.3整型函数 153 C.2.4类型转换函数 154 附录DC++语言支持 157 D.1代码例子. 157 D1.1数据类 157 D.1.2派生类 158 D.1.3类模板 158 D.1.4函数模板 159 D.2限制 160 D.2.1限定符 160 D.2.1.1设备存储器限定符 160 D.2.1.2 Volatile限定符 161 D.2.2指针 161 D.2.3运算符 161

8 B.14.4 例程 ...................................................................................................... 134 B.15 动态全局存储器分配 ..................................................................................... 135 B.15.1 堆存储器分配 ...................................................................................... 136 B.15.2 与设备存储器 API 的互操作.............................................................. 136 B.15.3 例程 ...................................................................................................... 136 B.15.3.1 每个线程的分配 ....................................................................... 136 B.15.3.2 每个线程块的分配 ................................................................... 137 B.15.3.3 在内核启动之间持久的分配 ................................................... 138 B.16 执行配置 ......................................................................................................... 140 B.17 发射绑定 ......................................................................................................... 140 B.18 #pragma unroll.................................................................................................. 143 附录 C 数学函数............................................................................................................. 145 C.1 标准函数 ........................................................................................................... 145 C.1.1 单精度浮点函数 .................................................................................... 145 C.1.2 双精度浮点函数 .................................................................................... 148 C.1.3 整型函数 ................................................................................................ 150 C.2 内置函数 ........................................................................................................... 150 C.2.1 单精度浮点函数 .................................................................................... 151 C.2.2 双精度浮点函数 .................................................................................... 152 C.2.3 整型函数 ................................................................................................ 153 C.2.4 类型转换函数 ........................................................................................ 154 附录 D C++语言支持 ...................................................................................................... 157 D.1 代码例子........................................................................................................... 157 D.1.1 数据类.................................................................................................... 157 D.1.2 派生类.................................................................................................... 158 D.1.3 类模板.................................................................................................... 158 D.1.4 函数模板................................................................................................ 159 D.2 限制................................................................................................................... 160 D.2.1 限定符.................................................................................................... 160 D.2.1.1 设备存储器限定符..................................................................... 160 D.2.1.2 Volatile 限定符............................................................................. 161 D.2.2 指针........................................................................................................ 161 D.2.3 运算符.................................................................................................... 161

目录 D.2.3.1赋值运算符 161 D.2.3.2地址运算符 161 D.2.4函数 162 D.2.4.1函数参数 162 D.2.4.2函数内静态变量 162 D.2.4.3函数指针 162 D.2.4.4函数递归 162 D.2.4.5函数定义 162 D2.5类 162 D.2.5.1数据成员 162 D.2.5.2函数成员 162 D.2.5.3构造器和析构器 162 D.2.5.4虚函数 162 D.2.5.5虚基类 163 D.2.6模板 163 附录E纹理获取… .165 E.1最近点取样 165 E.2线性滤波 166 E.3查找表… 167 附录F计算能力… …169 F1特性和技术规范 169 F2浮点标准 171 F.3计算能力1x 172 F.3.1架构 172 F.3.2全局存储器 173 F.3.2.1计算能力1.0和1.1的设备 173 F.3.2.2计算能力1.2和1.3的设备 173 F3.3共享存储器 174 F.3.3.132位步长访问 174 F.3.3.232位广播访问 174 F.3.3.38位和16位访问 175 F.3.3.4大于32位访问 ,175 F4计算能力2.X… .176

目录 D.2.3.1 赋值运算符 ................................................................................. 161 D.2.3.2 地址运算符 ................................................................................. 161 D.2.4 函数 ........................................................................................................ 162 D.2.4.1 函数参数 ..................................................................................... 162 D.2.4.2 函数内静态变量 ......................................................................... 162 D.2.4.3 函数指针 ..................................................................................... 162 D.2.4.4 函数递归 ..................................................................................... 162 D.2.4.5 函数定义 ..................................................................................... 162 D.2.5 类 ............................................................................................................ 162 D.2.5.1 数据成员 ..................................................................................... 162 D.2.5.2 函数成员 ..................................................................................... 162 D.2.5.3 构造器和析构器 ......................................................................... 162 D.2.5.4 虚函数 ......................................................................................... 162 D.2.5.5 虚基类 ......................................................................................... 163 D.2.6 模板 ........................................................................................................ 163 附录 E 纹理获取 ............................................................................................................. 165 E.1 最近点取样........................................................................................................ 165 E.2 线性滤波............................................................................................................ 166 E.3 查找表................................................................................................................ 167 附录 F 计算能力.............................................................................................................. 169 F.1 特性和技术规范 ................................................................................................ 169 F.2 浮点标准 ............................................................................................................ 171 F.3 计算能力 1.x ...................................................................................................... 172 F.3.1 架构 ......................................................................................................... 172 F.3.2 全局存储器 ............................................................................................. 173 F.3.2.1 计算能力 1.0 和 1.1 的设备 ........................................................ 173 F.3.2.2 计算能力 1.2 和 1.3 的设备 ........................................................ 173 F.3.3 共享存储器 ............................................................................................. 174 F.3.3.1 32 位步长访问 .............................................................................. 174 F.3.3.2 32 位广播访问 .............................................................................. 174 F.3.3.3 8 位和 16 位访问 .......................................................................... 175 F.3.3.4 大于 32 位访问 ............................................................................ 175 F.4 计算能力 2.x ...................................................................................................... 176

10 F.4.1架构 176 F.4.2全局存储器 177 F.4.3共享存储器 178 F.4.3.132位步长访问 179 F.4.3.2大于32位访问 179 F4.4常量存储器.1 80

10 F.4.1 架构......................................................................................................... 176 F.4.2 全局存储器............................................................................................. 177 F.4.3 共享存储器............................................................................................. 178 F.4.3.1 32 位步长访问.............................................................................. 179 F.4.3.2 大于 32 位访问............................................................................ 179 F.4.4 常量存储器............................................................................................. 180

第一章导论 第一章导论 1.1从图形处理到通用并行计算 市场对实时、高清晰度的三维图形具有无法满足的需求，由于这种需求的推动， 可编程图形处理器(GPU)己经演化成高并行度，多线程，拥有强大计算能力和极高 存储器带宽的多核处理器，如图11所示： Theoretical GFLOP/s 1750 GeForceGTX 580 NVIDIA GPU Single Precision 1500 NVIDIA GPU Double Precision GeForceGTX 480 Intel CPU Single Precision Intel CPU Double Precision 1250 1000 GeForceGTX 280 750 GeForce8800GTX Tesla C2050 500 GeForce7800 GTX 250 Westmere GeForce6800 Ultra Bloomfield Woodcrest Tesla C1060 GeForceFX 5800 Harpertown Sep-01Penti3 Jun-04 Oct-05 Mar-07 Jul-08 Dec-09

第一章 导论 1 第一章 导论 1.1 从图形处理到通用并行计算 市场对实时、高清晰度的三维图形具有无法满足的需求，由于这种需求的推动， 可编程图形处理器（GPU）已经演化成高并行度，多线程，拥有强大计算能力和极高 存 储 器 带 宽 的 多 核 处 理 器 ，如图 1-1 所 示 ：