CUDA C 编程权威指南 学习笔记:第二章 CUDA编程模型
详细细节参考书籍或者这个博客《CUDA C编程权威指南》——第2章 CUDA编程模型 2.1 CUDA编程模型概述
2.1.1 CUDA编程结构
CUDA编程模型还利用GPU架构的计算能力提供了以下几个特有功能。
1. 一种通过层次结构在GPU中组织线程的方法
2. 一种通过层次结构在GPU中访问内存的方法
CUDA编程模型主要是异步的,因此在GPU上进行的运算可以与主机-设备通信重叠。一个典型的CUDA程序包括由并行代码互补的串行代码。如图2-2所示,串行代码(及任务并行代码)在主机CPU上执行,而并行代码在GPU上执行。主机代码按照ANSI C标准进行编写,而设备代码使用CUDA C进行编写。一个典型的CUDA程序实现流程遵循以下模式。
- 把数据从CPU内存拷贝到GPU内存。
- 调用核函数对存储在GPU内存中的数据进行操作。
- 将数据从GPU内存传送回到CPU内存。
2.1.2 内存管理
udaMalloc与标准C语言中的malloc函数几乎一样,只是此函数在GPU的内存里分配内存。通过充分保持与标准C语言运行库中的接口一致性,可以实现CUDA应用程序的轻松接入。
cudaMemcpy函数负责主机和设备之间的数据传输,这个函数以同步方式执行,因为在cudaMemcpy函数返回以及传输操作完成之前主机应用程序是阻塞的。
2.1.3 线程管理
CUDA明确了线程层次抽象的概念以便于你组织线程。这是一个两层的线程层次结构,由线程块和线程块网格构成,如图2-5所示。
- 同步
- 共享内存
不同块内的线程不能协作。 线程依靠以下两个坐标变量来区分彼此。
- blockIdx(线程块在线程格内的索引)
- threadIdx(块内的线程索引)
网格和块的维度由下列两个内置变量指定。
- blockDim(线程块的维度,用每个线程块中的线程数来表示)
- gridDim(线程格的维度,用每个线程格中的线程数来表示)
2.1.4 启动核函数
同一个块中的线程之间可以相互协作,不同块内的线程不能协作。核函数调用结束后,控制权立刻返回给主机端。你可以调用以下函数来强制主机端程序等待所有的核函数执行结束:cudaDeviceSynchronize(). 一些CUDA运行时API在主机和设备之间是隐式同步的。当使用cudaMemcpy函数在主机和设备之间拷贝数据时,主机端隐式同步,即主机端程序必须等待数据拷贝完成后才能继续执行程序。
2.2 给核函数计时
2.2.1 使用CPU计时
由于GPU的执行是异步的,所以使用CPU计时就必须要显示做同步,cudaDeviceSynchronize()
2.2.2 使用nvprof计时
CUDA还提供了一个nvprof命令行分析工具
2.3 组织并行线程
从矩阵加法的例子中可以看出:
- 改变执行配置对内核性能有影响
- 传统的核函数实现一般不能获得最佳性能
- 对于一个给定的核函数,尝试使用不同的网格和线程块大小可以获得更好的性能