目录

CUDA学习笔记——CUDA的编程模型

CUDA 的程序结构

host 和 device

CUDA 编程模型能够让你使用少部分的、基于 C 语言的扩展,就能在异构系统上执行自己的程序。异构系统由 CPU 和 GPU 两种不同的处理器组成,两者之间用 PCIe 总线传输数据。在 CUDA 编程中,用两个名词区分他们:

  • host:CPU 和它的内存 (host memory)
  • device:GPU 和它的内存(device memory)

从 CUDA 6 开始,Nvidia 引入了 Unified Memory 的概念,它是将 CPU 内存和 GPU 内存映射到统一的地址空间,并且用统一的指针变量来访问。这样,在编程的时候,程序员可以只用一个指针,就可以控制数据对象。

需要注意的是,CPU 内存和 GPU 内存,在物理上依然是两个独立的地址空间,Unified Memory 只不过从逻辑上将它们合并在一起。而当你用一个指针来操作数据对象的时候,CUDA 在内部会自动地将数据在 CPU 内存和 GPU 内存之间传输。Unified Memory 只是简化了你的程序代码,让你用更少的代码实现原来的功能,但并没有真正实现 CPU 和 GPU 之间的共享内存。

关于 Unified Memory 细节,会在后面的系列文章里展开。

Kernel

CUDA 编程模型中一个重要的概念是 kernelkernel就是将要运行在 GPU 上的代码。作为一个程序员,你可以像写一个普通的 C 函数那样去写一个kernel,然后 CUDA 会执行调度程序,将kernel分配给多个线程,即多个线程一起执行kernel代码。

host 和 device 之间的大部分执行是异步的,也就是说,当 device 启动执行 kernel 后,控制权会立即返回给 host,host 可以继续执行接下来的其它任务。这样, CPU 和 GPU 就可以同时运行,充分用尽两者的计算能力。

典型的 CUDA 程序通常包含两种代码:串行执行的代码和并行执行的代码。如图:

串行代码,以及任务并行(task parallel)代码,通常由 host 来执行,数据并行(data parallel)代码,则由 device 来执行。host 代码就时我们平时写的 C/C++ 程序,而 device 代码就是用 CUDA C 写的kernel函数。

典型的 CUDA 程序的处理流程有三部分:

  1. 将相关数据从 CPU 复制到 GPU;
  2. 调用 kernel程序在 GPU 上执行并行计算;
  3. 将计算结果从 GPU 复制到 CPU;

 

GPU 内存

既然 GPU 有自己的内存,并且 CUDA 程序需要将数据复制到 GPU 内存后才能在 GPU 上执行计算任务,那么 GPU 就和 CPU 一样,有相关的内存管理接口。

CUDA 提供了一些 API,用来对 GPU 内存进行相关操作。这些 API 和标准 C 语言的相关 API 名字相近,只是在开头添加了 ”cuda“:

cudaMalloc用来在 GPU 上分配 全局内存(global memory)。GPU 上有多种不同类型的存储单元,它们的位置(片上/片下)、权限(读写)、访问范围等属性都不同,global memory 是最通用的存储单元,是片下(off-chip)存储,即不在 GPU 芯片内,是普通的 DRAM 内存。cudaMalloc 往往在启动 kernel 前,由 host 调用。

cudaFree 用来释放之前分配的 global memory,也是由 host 调用。

cudaMemset用于将已分配的 global memory,设置为 0。

cudaMemcpy 将数据从 host 复制到 device,或者相反。函数原型:

1
cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, cudaMemcpyKind kind)

第四个参数用来指定复制数据的方向,支持四种:

  • cudaMemcpyHostToHost:在 host 内复制数据,和 memcpy 作用一样;
  • cudaMemcpyHostToDevice:从 host 复制数据到 device;
  • cudaMemcpyDeviceToHost:从 device 复制数据到 host;
  • cudaMemcpyDeviceToDevice:在 device 内复制数据。

需要注意的是,和其他大部分 CUDA API 不同,cudaMemcpy 是同步执行函数,就是当调用时,host 会一直阻塞,直到数据传输完,函数才返回。CUDA 也提供了一个异步版本的 memcpy,通常用于 CUDA 流(stream)的开发中。关于为什么需要用同步机制,以及 stream 开发,后面的文章中会展开讲。

Nvidia 的 GPU 中,有多种不同的存储单元,有各自的特点和使用场合。其中,最通用的两种存储类型是 global memoryshared memoryglobal memory 和 CPU 系统中的主存类似,是在 GPU 芯片之外,用 DRAM 介质的存储单元。shared memory 和 CPU 的 cache 类似,在 GPU 芯片上,用 SRAM 介质存储。和 CPU 的 cache 不同的是,GPU 的 shared memory 受限于程序调度,当一个 SM 上分配的任务过多时,每个任务能获取的 shared memory 就较少。另外,不同于 CPU cache,shared memory 可以在程序中直接控制,从而带来较多的灵活性。

 

GPU 线程

kernelthread 是一体两面,kernel 是静态的 GPU 执行代码,线程是动态的执行单元。当 kernel 被启动时,CUDA 会根据配置,启动多个线程,来并行地执行 kernel 代码。CUDA 提供了一个线程层次的抽象,让程序员可以方便的管理线程,并且增加了可扩展性。

CUDA 的线程管理有两个层次:

  • 网格(grid)
  • 线程块(thread block)

gridkernel 是一对一关系,每个 kernel 启动后,对应一个 gridgrid 中的所有线程共享一块相同的 global memory。一个 grid 由 多个 thread block 组成,一个 thread block 由多个 thread 组成。blocks 和 threads 的维度由程序员在启动前设定。

一个 block 中的 threads 可以采用一些机制来合作,通常有两方面:

  • block 范围的同步:即一个 block 内的所有线程,可以同步。
  • block 范围内的数据共享:通过使用 shared memory ,一个 block 内的所有线程可以共享数据。

block 之间一般没有合作机制,属于不同 block 的线程,无法同步。

kernel 启动前,需要初始化线程的组织结构。CUDA 是用三维结构来组织 blocks 和 threads,通过 dim3 类型来设置。dim3 是一个结构体,内部有 x、y、z 三个变量分别表示三个维度的值。也可以把 blocks 和 threads 设为一维或二维结构,只需把另外的多余维度设为 1 即可。

kernel 启动后,在每个线程内部,有两组变量,分别表示线程所在的 grid/block 的维度,以及线程在该维度中的坐标。

表示维度的变量是:

  • gridDim:线程所在的 grid 的维度;
  • blockDim:线程所在的 block 的维度;

表示坐标的变量:

  • blockIdx:定位线程所属的 block 在 grid 中的坐标;
  • threadIdx:定位线程在 block 中的坐标;

这些变量的类型都是一个结构体,内部包含三个成员变量:x、y、z,用来表示维度的值或者在该维度的坐标。

在 CUDA 程序开发中,grid/block 的维度设计是很重要的,尤其是 block 的维度设计。因为像 shared memory、registers、L1 cache、cores 等资源的分配,都和 grid/block 的维度设计有关,另外,维度设计的不同,也会导致 kernel 实现的不同,从而间接地通过 unrollingbranch divergency 等特点影响性能。所以说,CUDA 的两层线程层级结构,给了程序员自主优化的权利。

需要注意的是,由于受限于 GPU 资源,block 的维度并非允许完全自由设定,具体限制可以参考 GPU 的规格说明书。


- 全文完 -

相关文章

「 您的赞赏是激励我创作和分享的最大动力! 」