本博文是根据中科大信息学院谭立湘老师的课件加上自己的理解整理出来的

************************************************************************************

NVIDIA在2007年推出CUDA这个统一计算架构

CUDA的基本思想是支持大量的线程级并行，并在硬件中动态地调度和执行这些线程

 

CUDA软件体系可以分为三层结构

• CUDA函数库（CUDA Library）

• CUDA运行时API（Runtime API）https://blog.csdn.net/qq_41598072/article/details/81030272

• CUDA驱动API（Driver API）

Difference between the driver and runtime APIs

https://docs.nvidia.com/cuda/cuda-driver-api/driver-vs-runtime-api.html#driver-vs-runtime-api

CUDA软件环境：

CUDA支持Windows、Linux、MacOS三种主流操作系统，支持CUDA C及CUDA Fortran等多种语言。无论使用何种语言或接口，指令最终都会被驱动程序转换成PTX（ParallelThread Execution，并行线程执行，CUDA架构中的指令集，类似于汇编语言）代码，交由GPU核心计算。CUDA最主要的包含两个方面：ISA指令集架构与硬件计算引擎；实际上是硬件和指令集。见下图中的绿色部分，CUDA 架构的组件组成是：

(1)NVIDIA GPU中的并行计算引擎；

(2)对硬件初始化、配置的OS内核级支持；

(3)用户模式的驱动，为开发者的PTX 指令集架构(ISA，Instructionset architecture)

Kernel

Kernel函数：

Kernel函数是指为GPU设备编译的一个函数。也就是一个编译好的、在GPU上并行运行的计算函数。Kernel在GPU上以多个线程的方式被执行。

运行在GPU上的CUDA并行计算函数称为kernel函数（内核函数）。一个完整的CUDA程序是由一系列的设备端kernel函数并行部分和主机端的串行处理部分共同组成的。这些处理步骤会按照程序中相应语句的顺序依次执行，满足顺序一致性。

CUDA编程中的术语:

• Host：宿主，指CPU，系统的CPU。负责启动应用程序，运行程序的串行部分，将程序的并行、计算密集的部分offload到GPU上运行，并最终返回程序的运行结果。

• Device：设备，指GPU，CPU的协处理器。负责程序的并行、计算密集部分的处理，并将处理结果返回给Host。

Block：线程块

——执行Kernel的一组线程组成一个线程块。(一个Kernel只做同一件事)

一个线程块最多可包含1024个并行执行的线程，线程之间通过共享内存有效地共享数据，并实现线程的通信和栅栏同步。

线程ID：线程在线程块中的线程号（唯一标识）

基于线程ID的复杂寻址，应用程序可以将线程块指定为任意大小的二维或三维数组，并使用2个或3个索引来标识每个线程。

• 对于大小是（Dx，Dy）的二维线程块，索引为（x，y）的线程的线程ID为（x+y*Dx）

• 对于大小为（Dx，Dy，Dz）的三维线程块，索引为（x，y，z）的线程的线程ID为：

（x+y*Dx+z*Dx*Dy）

Grid：线程块组成的线程网格（最多2^32 个blocks）

执行相同Kernel、具有相同维数和大小的线程块可以组合到一个网格中。这样单个Kernel调用中启动的线程数就可以很大。同一网格中的不同线程块中的线程不能互相通信和同步。

Grid 是一个线程块阵列，执行相同的内核，从全局内存读取输入数据，将计算结果写入全局内存。

Block ID：线程块ID

线程块ID是线程块在Grid中的块号。实现基于块ID的复杂寻址，应用程序可以将Grid指定为任意大小的二维数组，并用2个索引来标识每个线程块。对于大小为（Dx，Dy）的二维线程块，索引为（x，y）的线程块的ID为（x+y*Dx）。现已支持三维

Wrap：线程束

一个线程块中连续的固定数量（32）的线程组。

将线程块中的线程划分成wrap的方式是：每个wrap包含线程ID连续递增的32个线程，从线程0开始递增到线程31。

Stream：

CUDA的一个Stream表示一个按特定顺序执行的GPU操作序列。诸如kernel启动、内存拷贝、事件启动和停止等操作可以排序放置到一个Stream中。

一个Stream包含了一系列Grids，并且可以多个Stream并行执行。

在CUDA 架构下，GPU芯片执行时的最小单位是thread。

若干个thread可以组成一个线程块（block）。一个block中的thread能存取同一块共享内存，可以快速进行同步和通信操作。

每一个block 所能包含的thread 数目是有限的。执行相同程序的block，可以组成grid。不同block 中的thread 无法存取同一共享内存，因此无法直接通信或进行同步。

不同的grid可以执行不同的程序（kernel）。

Grid是由线程块组成的网格。每个线程都执行该kernel，应用程序指定了Grid和线程块的维数，Grid的布局可以是一维、二维或三维的。

每个线程块有一个唯一的线程块ID，线程块中的每个线程具有唯一的线程ID。同一个线程块中的线程可以协同访问共享内存，实现线程之间的通信和同步。

每个线程块最多可以包含的线程的个数为1024个，线程块中的线程以32个线程为一组的Wrap的方式进行分时调度。每个线程在数据的不同部分并行地执行相同的操作。

CUDA处理流程：

在CUDA 的架构下，一个程序分为两个部份：Host 端和Device 端。Host 端是指在CPU 上执行的部份，而Device 端则是在GPU上执行的部份。Device端的程序又称为kernel函数。

通常Host 端程序会将数据准备好后，复制到GPU的内存中，再由GPU执行Device 端程序，完成后再由Host 端程序将结果从GPU的内存中取回。

CPU 存取GPU 内存时只能通过PCI-E 接口，速度有限。

• 1)从系统内存中复制数据到GPU内存

• 2)CPU指令驱动GPU运行；

• 3)GPU 的每个CUDA核心并行处理

• 4)GPU 将CUDA处理的最终结果返回到系统的内存

 

CUDA编程模型：

• CPU作为主机端只能有一个

• GPU作为设备端可以有多个

• CPU主要负责逻辑处理

• GPU负责密集型的并行计算

        完整的CUDA程序包括主机端和设备端两部分代码，主机端代码在CPU上执行。

设备端代码（kernel函数）运行在GPU上。其中一个kernel函数对应一个grid，每个grid根据需要配置不同的block数量和thread数量。

CUDA包含两个并行逻辑层：block层和thread层。

在执行时block映射到SM

thread映射到SP（Core）

如何在实际应用程序中高效地开发这两个层次的并行是CUDA编程与优化的关键之一。

Stream > Grid > Block > Warp > Thread

学校        年级      班级       小组        学生

 

Kernel的启动参数

• cuda程序执行流程：

• 单显卡只需要考虑红色的，多显卡要七步曲

• 1)cudaSetDevice(0); //获取设备；只有一个GPU时或默认使用0号GPU时可以省略
  
  2)cudaMalloc((void**) &d_a,sizeof(float)*n); //分配显存
  
  3)cudaMemcpy(d_a,a,sizeof(float)*n,cudaMemcpyHostToDevice); //数据传输
  
  4)gpu_kernel<<>>(***); //kernel函数
  
  5)cudaMemcpy(a,d_a,sizeof(float)*n,cudaMemcpyDeviceToHost); //数据传输
  
  6)cudaFree(d_a); //释放显存空间
  
  7)cudaDeviceReset( ); //重置设备；可以省略

完整的向量点积CUDA程序

/*
a = [a1, a2, …an], b = [b1, b2, …bn]
a*b = a1*b1 + a2*b2 + … + an*bn
*/#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include 
#include 
#include 
#define N 10
__global__ void Dot(int *a, int *b, int *c) //声明kernel函数
{__shared__ int temp[N]; // 声明在共享存储中的变量temp[threadIdx.x] = a[threadIdx.x] * b[threadIdx.x];//__syncthreads();if (0 == threadIdx.x){//Kernel函数中利用threadIdx.x 获得线程索引号//threadIdx是内建变量，它指定block内thread索引号int sum = 0;for (int i = 0; i < N; i++)sum += temp[i];*c = sum;printf("sum Calculated on Device:%d
", *c);}
}void random_ints(int *a, int n)
{for (int i = 0; i< n; i++)*(a + i) = rand() % 10;
}int main()
{int *a, *b, *c;int *d_a, *d_b, *d_c;int size = N * sizeof(int);cudaMalloc((void **)&d_a, size);cudaMalloc((void **)&d_b, size);cudaMalloc((void **)&d_c, sizeof(int));a = (int *)malloc(size); random_ints(a, N);b = (int *)malloc(size); random_ints(b, N);c = (int *)malloc(sizeof(int));printf("Array a[N]:
");for (int i = 0; i < N; i++) printf("%d ", a[i]);printf("
");printf("Array b[N]:
");for (int i = 0; i < N; i++) printf("%d ", b[i]);printf("

");cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);Dot << <1, N >> >(d_a, d_b, d_c); //单block多threadcudaMemcpy(c, d_c, sizeof(int), cudaMemcpyDeviceToHost);int sumHost = 0;for (int i = 0; i < N; i++)sumHost += a[i] * b[i];printf("sum Calculated on Host=%d
", sumHost);printf("Device to Host: a*b=%d
", *c);free(a); free(b); free(c);cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);return 0;
}