CUDA 编程模型

概述

一、 主机与设备

CUDA 编程模型将 CPU 作为主机 (Host) ， GPU 作为协处理器 (co-processor) 或者设备 (Device). 在一个系统中可以存在一个主机和多个设备。 CPU 主要负责进行逻辑性强的事物处理和串行计算， GPU 则专注于执行高度线程化的并行处理任务。 CPU 、 GPU 各自拥有相互独立的存储器地址空间：主机端的内存和设备端的显存。 CUDA 对内存的操作与一般的 C 程序基本相同，但增加了一种新的 pinned memory ；

二、运行在 GPU 上的 CUDA 并行计算函数称为 kernel( 内核函数 ) 。一个 kernel 函数并不是一个完整的程序，而是整个 CUDA 程序中的一个可以被并行执行的步骤。

三、 CPU 串行代码完成的工作包括在 kernel 启动前进行数据准备和设备初始化的工作，以及在 kernel 之间进行一些串行计算。理想情况是， CPU 串行代码的作用应该只是清理上一个内核函数，并启动下一个内核函数。这样，就可以在设备上完成尽可能多的工作，减少主机与设备之间的数据传输。

Kernel 函数的定义与调用

一、 运行在 GPU 上的程序称为 kernel( 内核函数 ) 。

内核函数必须通过 _global_ 函数类型限定符定义，并且只能在主机端代码中调用。在调用时，必须声明内核函数的执行参数。

注意： __global__ 下划线在Visual studio中的写法。

使用一种新 <<<…>>> 执行配置语法指定执行某一指定内核调用的线程数。必须先为 Kernel 中用到的数组或变量分配好足够的空间，再调用 kernel 函数

二、 在设备端运行的线程之间是并行执行的，每个线程有自己的 blockID 和 threadID 用于与其他线程相区分。 BlockID 和 threadID 只能在 kernel 中通过内建变量访问。

三、 内建变量不需由程序员定义，是由设备中的专用寄存器提供的。所以，内建变量是只读的，并且只能在 GPU 端得 kernel 函数中使用。

线程结构

Kernel 是以 block 为单位执行的， CUDA 引入 grid 来表示一系列可以被并行执行的 block 的集合。各 block 是并行执行的， block 之间无法通信，也没有执行顺序。

 

block 内通信原理

在同一个 block 中的线程通过共享存储器 (shared memory) 交换数据，并通过栅栏同步保证线程间能够正确地共享数据。具体来说，可以在 kernel 函数中需要同步的位置调用 _syncthreads() 函数。

为了保证线程块中的各个线程能够有效协作，访问共享存储器的延迟必须很小。所以在 GPU 中，共享存储器与执行单元的物理距离必须很小，处于同一个处理核心中。而为了在硬件上用很小的代价就能实现 _syncthreads() 函数，一个 block 中所有线程的数据都必须交由同一处理核心进行处理。所以，这导致每个线程块中的线程数量、共享存储器大小、寄存器数量都要受到处理核心硬件资源的限制。目前，每个 block 里最多只能有 512 个线程。

计算单元

在 GPU 内部， SM 代表流多处理器，即计算核心。每个 SM 中又包含 8 个标量流处理器 SP 以及少量的其他计算单元。实际上， SP 只是执行单元，并不是完整的处理核心。处理核心必须包含取指、解码、分发逻辑和执行单元。隶属同一 SM 的 8 个 SP 共用同一套取指和发射单元，也共用一块共享存储器。

一个 block 必须被分配到一个 SM 中，但是一个 SM 中同一时刻可以有多个活动线程块等待执行。这可以更好地利用执行单元的资源，当一个 block 进行同步或者访问显存等高延迟操作时，另一个 block 就可以占用 GPU 执行资源。

目前，一个内核函数只有一个 grid ，但在支持 DirectX 11 的硬件中，这一限制将会解除。

真正的执行单元

在实际运行中， block 会被分割为更小的线程束 (warp) ， warp 的大小由硬件的计算能力版本决定。在采用 Tesla 架构的 GPU 中，一个线程束由连续的 32 个线程组成。 Warp 中的线程只与线程 thread  ID 有关。在每发射一条 warp 指令， SM 中的 8 个 SP 将执行这条指令 4 遍。

执行模型

CUDA 采用了 SIMT( 单指令多线程 ) 执行模型。在 SIMT 模型中，如果需要控制单个线程的行为，这会大大降低效率。