CUDA编程CUDA入门笔记

Posted 2022-02-10 非晚非晚

文章目录

• 1. CUDA结构
• 2. SP、SM与warp
• 3. CUDA线程索引的方式
• 4 CUDA的内存
• 5 编程模型
• • （1）怎么确定是在CPU还是GPU上跑？
  • （2）CPU与GPU的数据传输
  • （3）怎么用代码表示线程组织模型

1. CUDA结构

显卡内部，有三级结构：网格（grid）、块（block）、线程（thread）。每个显卡只有很少的网格，一个核函数只能运行在一个网格中，而一个网格里有多个块，每个块包含了若干线程。

kernel在device上执行时实际上是启动很多线程，一个kernel所启动的所有线程称为一个网格（grid），同一个网格上的线程共享相同的全局内存空间，grid是线程结构的第一层次，而网格又可以分为很多线程块（block），一个线程块里面包含很多线程，这是第二个层次。线程两层组织结构如下图所示，这是一个gird和block均为2-dim的线程组织。

介绍说明：

1. Thread：线程，并行的基本单元
2. Thread Block：线程块，互相合作的线程组。它有以下特点：

• 允许彼此同步
• 可以通过共享内存快速交换数据
• 以1维、2维或3维组织(里面的thread)

3. Grid：一组线程块

• 以1维、2维和3维度组织(里面是block)
• 共享全局内存

4. Kernel：在GPU上执行的核心程序，这个kernel函数是运行在某个Grid上的。

每一个block和每个thread都有自己的ID，我们通过相应的索引找到相应的线程和线程块。

• threadIdx，blockIdx
• Block ID: 1D or 2D
• Thread ID: 1D, 2D or 3D

2. SP、SM与warp

SP（streaming Process），SM（streaming multiprocessor）是硬件（GPU hardware）概念。而thread，block，grid，warp是软件上的（CUDA）概念。

• SP(streaming processor)：最基本的处理单元，也称为CUDA core。最后具体的指令和任务都是在SP上处理的。GPU进行并行计算，也就是很多个SP同时做处理。

• SM(streaming multiprocessor)：多个SP加上其他的一些资源组成一个streaming multiprocessor。也叫GPU大核，其他资源如：warp scheduler，register，shared memory等。SM可以看做GPU的心脏（对比CPU核心），register和shared memory是SM的稀缺资源。CUDA将这些资源分配给所有驻留在SM中的threads。因此，这些有限的资源就使每个SM中active warps有非常严格的限制，也就限制了并行能力。

• warp(线程束)：GPU执行程序时的调度单位，目前cuda的warp的大小为32，同在一个warp的线程，以不同数据资源执行相同的指令,这就是所谓 SIMT。

需要指出，每个SM包含的SP数量依据GPU架构而不同，Fermi架构GF100是32个，GF10X是48个，Kepler架构都是192个，Maxwell都是128个。

简而言之，SP是线程执行的硬件单位，SM中包含多个SP，一个GPU可以有多个SM（比如16个），最终一个GPU可能包含有上千个SP。

下图展示了逻辑层面和硬件层面的对应关系：

• 每个线程由每个线程处理器（SP）执行
• 线程块由多核处理器（SM）执行
• 一个kernel其实由一个grid来执行，一个kernel一次只能在一个GPU上执行
• 关于warp的解释：

• 一个SM中可以同时有多个warp，这些warp被称为active warp，不同warp处于不同状态，挂起，就绪，执行。但是一个SM上正在被执行的就只有一个warp，这个正在被执行的warp叫做resident warp。 active warp是指已经分配给SM的warp，并且该warp需要的资源（寄存器）也已经分配。
• 一个SP可以执行一个thread，但是实际上并不是所有的thread能够在同一时刻执行。Nvidia把32个threads组成一个warp，warp是调度和运行的基本单元。warp中所有threads并行的执行相同的指令。一个warp需要占用一个SM运行，多个warps需要轮流进入SM。由SM的硬件warp scheduler负责调度。目前每个warp包含32个threads（Nvidia保留修改数量的权利）。所以，一个GPU上resident thread最多只有 SM*warp个。

block是软件概念，一个block只会由一个sm调度，程序员在开发时，通过设定block的属性，告诉GPU硬件，我有多少个线程，线程怎么组织。而具体怎么调度由sm的warps scheduler负责，block一旦被分配好SM，该block就会一直驻留在该SM中，直到执行结束。一个SM可以同时拥有多个blocks，但需要序列执行。下图显示了GPU内部的硬件架构：

3. CUDA线程索引的方式

CUDA使用多级索引的方式访问线程。

• 定位Block：第一级索引是(grid.xIdx, grid.yIdy)，通过它我们就能找到了这个线程块的位置。
• 定位thread：第二级索引(block.xIdx, block.yIdx, block.zIdx)来定位到指定的线程。

grid和block都是定义为dim3类型的变量，dim3可以看成是包含三个无符号整数（x，y，z）成员的结构体变量，在定义时，缺省值初始化为1。因此grid和block可以灵活地定义为1-dim，2-dim以及3-dim结构，对于图中结构（水平方向为x轴），定义的grid和block如下所示

dim3 grid(3, 2);
dim3 block(5, 3);
kernel_fun<<< grid, block >>>(prams...);

定义图解如下：

一个线程需要两个内置的坐标变量（blockIdx，threadIdx）来唯一标识，它们都是dim3类型变量，其中blockIdx指明线程所在grid中的位置，而threaIdx指明线程所在block中的位置，如图中的Thread (1,1)满足：

threadIdx.x = 1
threadIdx.y = 1
blockIdx.x = 1
blockIdx.y = 1

一个线程块上的线程是放在同一个流式多处理器（SM)上的，但是单个SM的资源有限，这导致线程块中的线程数是有限制的，现代GPUs的线程块可支持的线程数可达1024个。有时候，我们要知道一个线程在blcok中的全局ID，此时就必须还要知道block的组织结构，这是通过线程的内置变量blockDim来获得。它获取线程块各个维度的大小。对于一个2-dim的$block(D_x,D_y)$，线程$(x,y)$的ID值为$(x+y\ast D_x)$，如果是3-dim的$block(D_x,D_y,D_z)$，线程$(x,y,z)$的ID值为$(x+y\ast D_x+z\ast D_x\ast D_y)$。另外线程还有内置变量gridDim，用于获得网格块各个维度的大小。

一个Grid可以包含多个Blocks，Blocks的组织方式可以是一维的，二维或者三维的。block包含多个Threads，这些Threads的组织方式也可以是一维，二维或者三维的。

CUDA中每一个线程都有一个唯一的标识ID—ThreadIdx，这个ID随着Grid和Block的划分方式的不同而变化，这里给出Grid和Block不同划分方式下线程索引ID的计算公式。

1、 grid划分成1维，block划分为1维

    int threadId = blockIdx.x *blockDim.x + threadIdx.x;  

2、 grid划分成1维，block划分为2维

int threadId = blockIdx.x * blockDim.x * blockDim.y+ threadIdx.y * blockDim.x + threadIdx.x;  

3、 grid划分成1维，block划分为3维

    int threadId = blockIdx.x * blockDim.x * blockDim.y * blockDim.z  
                       + threadIdx.z * blockDim.y * blockDim.x  
                       + threadIdx.y * blockDim.x + threadIdx.x;  

4、 grid划分成2维，block划分为1维

    int blockId = blockIdx.y * gridDim.x + blockIdx.x;  
    int threadId = blockId * blockDim.x + threadIdx.x;  

5、 grid划分成2维，block划分为2维

    int blockId = blockIdx.x + blockIdx.y * gridDim.x;  
    int threadId = blockId * (blockDim.x * blockDim.y)  
                       + (threadIdx.y * blockDim.x) + threadIdx.x;  

6、 grid划分成2维，block划分为3维

    int blockId = blockIdx.x + blockIdx.y * gridDim.x;  
    int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)  
                       + (threadIdx.z * (blockDim.x * blockDim.y))  
                       + (threadIdx.y * blockDim.x) + threadIdx.x;  

7、 grid划分成3维，block划分为1维

    int blockId = blockIdx.x + blockIdx.y * gridDim.x  
                     + gridDim.x * gridDim.y * blockIdx.z;  
    int threadId = blockId * blockDim.x + threadIdx.x;  

8、 grid划分成3维，block划分为2维

    int blockId = blockIdx.x + blockIdx.y * gridDim.x  
                     + gridDim.x * gridDim.y * blockIdx.z;  
    int threadId = blockId * (blockDim.x * blockDim.y)  
                       + (threadIdx.y * blockDim.x) + threadIdx.x;  

9、 grid划分成3维，block划分为3维

    int blockId = blockIdx.x + blockIdx.y * gridDim.x  
                     + gridDim.x * gridDim.y * blockIdx.z;  
    int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)  
                       + (threadIdx.z * (blockDim.x * blockDim.y))  
                       + (threadIdx.y * blockDim.x) + threadIdx.x;     

4 CUDA的内存

CUDA中的内存模型分为以下几个层次：

• 每个线程都用自己的registers（寄存器）和local memory（局部内存）
• 每个线程块(block)内都有自己的shared memory（共享内存），所有线程块内的所有线程共享这段内存资源
• 每个grid都有自己的global memory（全局内存），constant memory（常量内存）和texture memory（纹理内存），不同线程块的线程都可使用。

线程访问这几类存储器的速度是：register > local memory >shared memory > global memory。

下面这幅图表示就是这些内存在计算机架构中的所在层次。

5 编程模型

在CUDA中，host和device是两个重要的概念，我们用host指代CPU及其内存，而用device指代GPU及其内存。CUDA程序中既包含host程序，又包含device程序，它们分别在CPU和GPU上运行。同时，host与device之间可以进行通信，这样它们之间可以进行数据拷贝。

典型的CUDA程序的执行流程如下：

1. 分配host内存，并进行数据初始化；
2. 分配device内存，并从host将数据拷贝到device上；
3. 调用CUDA的核函数在device上完成指定的运算；
4. 将device上的运算结果拷贝到host上；
5. 释放device和host上分配的内存。

核函数用__global__符号声明，在调用时需要用<<<grid, block>>>来指定kernel要执行的线程数量，在CUDA中，每一个线程都要执行核函数，并且每个线程会分配一个唯一的线程号thread ID，这个ID值可以通过核函数的内置变量threadIdx来获得。

（1）怎么确定是在CPU还是GPU上跑？

由于GPU实际上是异构模型，所以需要区分host和device上的代码，在CUDA中是通过函数类型限定词开区别host和device上的函数，主要的三个函数类型限定词如下：

• __global__：在device上执行，从host中调用（一些特定的GPU也可以从device上调用），返回类型必须是void，不支持可变参数参数，不能成为类成员函数。注意用__global__定义的kernel是异步的，这意味着host不会等待kernel执行完就执行下一步。
• __device__：在device上执行，仅可以从device中调用，不可以和__global__同时用。
• __host__：在host上执行，仅可以从host上调用，一般省略不写，不可以和__global__同时用，但可和__device__，此时函数会在device和host都编译。

通过关键字就可以表示某个程序在CPU上跑还是在GPU上跑。

函数	执行位置	调用位置
__device__ float DeviceFunc()	device	device
__global void KernelFunc()	device	host
__host__ float HostFunc()	host	host

（2）CPU与GPU的数据传输

首先介绍在GPU内存分配回收内存的函数接口：

• cudaMalloc(): 在设备端分配global memory
• cudaFree(): 释放存储空间

CPU的数据和GPU端数据做数据传输的函数接口是一样的，他们通过传递的函数实参（枚举类型）来表示传输方向：

cudaMemcpy(void *dst, void *src, size_t nbytes, enum cudaMemcpyKind direction);

enum cudaMemcpyKind的类型：

• cudaMemcpyHostToDevice（CPU到GPU）
• cudaMemcpyDeviceToHost（GPU到CPU）
• cudaMemcpyDeviceToDevice（GPU到GPU）

（3）怎么用代码表示线程组织模型

我们可以用dim3类来表示网格和线程块的组织方式，网格grid可以表示为一维和二维格式，线程块block可以表示为一维、二维和三维的数据格式。

dim3 DimGrid(100, 50);  //5000个线程块，维度是100*50
dim3 DimBlock(4, 8, 8);  //每个线层块内包含256个线程，线程块内的维度是4*8*8

---

参考：

https://blog.csdn.net/xiaohu2022/article/details/79599947