笔记-CUDA入门

GPU

GPU的性能指标：1. 核心数 2. GPU显存容量 3. GPU计算峰值 4. 显存带宽

GPU不能单独计算，CPU + GPU组成异构计算架构

主机（Host, 即CPU）和设备（Device，即GPU）之间的内存访问一般通过PCle总线链接。

CUDA编译开发环境

CUDA核函数

1. 核函数在 GPU 上 进行 并行执行

2. 注意：

   1. 限定词 __global__ 修饰核函数
   2. 返回值必须是 void

3. 形式：

   __global__ void kernel_function(argument arg)
   {
       printf("Hello World form the GPU!\n");
   } // 或交换 __global__ 和 void，顺序无关
   

核函数注意事项

1. 核函数只能访问GPU内存
2. 核函数不能使用变长参数
3. 核函数不能使用静态变量
4. 核函数不能使用函数指针
5. 核函数具有异步性

CUDA程序编写流程

int main(void)
{
	主机代码; // 对 GPU 配置，预处理
	核函数调用;
	主机代码; // 回传GPU处理后的数据到主机，对CPU,GPU内存释放
	return 0;
}

注意：核函数不支持 C++ 的 iostream

核函数程序示例

#include <stdio.h>
__global__ void hello_from_gpu() // 除了限定词与平时写的C++代码无任何区别
{
	printf("Hello World from the GPU!\n")
}

int main() 
{
	hello_from_gpu
<<<1,1>>>
(); // 核函数的调用与 C++ 函数有区别
	cudaDeviceSynchronize(); // 同步主机与设备，促使缓冲区刷新，打印 Hello World 到终端
													 // 调用同步函数, 让 CPU 等待 GPU 设备执行完毕再进行统一数据处理
	return 0;
}

线程模型

• grid 网格
• block 线程块

线程分块是逻辑上的划分，物理上线程不分块

配置线程 <<<grid.size, block.size>>>

最大线程块大小：1024

最大允许网格大小 2^31 - 1 (针对一维网格）

#include <stdio.h>
__global__ void hello_from_gpu() 
{
	printf("Hello World from the GPU!\n")
}

int main() 
{
	hello_from_gpu<<<2,4>>>(); 
	cudaDeviceSynchronize(); 
	return 0;
}
// 预期输出 8 条 Hello World

一维线程模型

1. 每个线程在核函数中都有一个唯一的身份标识
2. 每个线程的唯一标识由 <<<grid_size, block_size>>> 确定，grid_size，block_size 保存在内建变量（目前考虑的是一维的情况）
   • gridDim.x：该变量的数值等于执行配置中变量 grid_size 的值
   • gridDim.y: 该变量的数值等于制定配置中变量block_size 的值
3. 线程索引保存内建变量
   1. blockIdx.x: 该变量指定一个线程在一个网格中的线程块索引值，范围为 0 - gridDim.x - 1;
   2. threadIdx.x: 该变量指定一个线程在一个线程块中的线程索引值，范围为 0- blockDim.x - 1;

线程的唯一标识：Idx = threadIdx.x + blockIdx.x * blockDim.x

推广到多维线程

1. CUDA 可以组织三维的网格和线程块；
2. blockIdx 和 threadIdx 是 类型 uint3 的变量，该类型是一个结构体，具有 x, y, z 三个成员
3. girdDim 和 blockDim 是 类型为 dim3 的变量，该类型是一个结构体，具有 x, y, z 三个成员gridDim 和 blockDIm 没有指定的维度默认为 1.

定义多维网格和线程块（C++ 构造函数语法）：

dim3 grid_size(Gx, Gy, Gz);
dim3 block_size(Bx, By, Bz);

网格大小限制：

• gridDim.x 最大值 —— 2^31 - 1
• gridDim.y 最大值 —— 2^16 - 1
• girdDim.z 最大值 —— 2^16 - 1

线程块大小限制：

• blockDim.x 最大值 —— 1024
• blockDim.y 最大值 —— 1024
• blockDim.z 最大值 —— 16

注意：线程块总的大小最大为 1024！

nvcc 编译流程

1. nvcc 分离全部源代码为：(1) 主机代码 （2）设备代码
2. 主机（Hsot）代码是 C/C++ 语法，设备（device）是 C/C++ 扩展语言编写的
3. nvcc 先将设备代码编译为 PTX（Parallel Thread Execition）伪汇编代码，再将 PTX 代码编译为二进制的 cubin 目标代码
4. 在将源代码编译为 PTX 代码时，需要用选项 -arch=compute_XY 指定一个虚拟架构的计算能力，用以确认代码中能够使用的 CUDA 功能
5. 在将 PTX 代码编译为 cubin 代码时，需要用选项 -code=sm_ZW 指定一个真实架构的计算能力，用以确定可执行文件能使用的 GPU

PTX

• PTX 是 CUDA 平台为基于 GPU 的通用计算而定义的虚拟机和指令集
• nvcc 编译命令总是使用两个体系结构：一个是虚拟的中间体系结构，另一个是实际的 GPU 体系结构
• 虚拟架构更像是对应用所需的 GPU 功能的声明
• 虚拟架构应该尽可能地选择低——适配更多实际GPU
• 真是架构应该尽可能地选择高——充分发挥GPU性能

CUDA 程序基本框架

#include<头文件>
__global__ void 函数名(参数...)
{
	核函数内容
}

int main(void)
{
	设置 GPU 设备 // setGPU();
	分配主机和设备内存; // size_t stByteCount = iElemCount * sizeof(float);
	初始化主机中的数据;
	数据从主机复制到设备;
	调用和函数在设备中进行计算;
	将计算得到的数据从设备传给主机;
	释放主机和设备内存;
}

设置 GPU 设备

1. 获取GPU设备 __host__ __device__ cudaError

   int iDeviceCount = 0;
   cudaGetDeviceCount(&DeviceCount); // 返回cudaSuccess
   

2. 设备 GPU 执行时使用的设备 __host__ cudaError

   int iDev = 0;
   cudaSetDevice(iDev);
   

内存管理

• CUDA 通过内存分配，数据传递，内存初始化，内存释放进行内存管理

标准 C 语言内存管理函数	CUDA 内存管理函数
malloc	cudaMalloc
memcpy	cudaMemcpy
memset	cudaMemset
free	cudaFree

1. 内存分配

• 主机分配内存: extern void *malloc(unsigned int num_bytes);

  代码：

  float *fpHost_A;
  fpHost_A = (float *)malloc(nBytes);
  

• 设备分配内存: __host__ __device__ cudaError_t cudaMalloc(void **devPtr``, size_t size);

  代码：

  float *fpDevice_A;
  cudaMalloc((float**)&fpDevice_A, nBytes);
  

2. 数据拷贝

• 主机数据拷贝: void *memcpy(void *dest, const void *src, size_t n);

  代码：

  memcpy((void*) d, (void*)s, nBytes);
  

• 设备数据拷贝: __host__ cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, cudaMemcpyKind kind)

  代码：

  cudaMemcpy(Device_A, Host_A, nBytes, cudaMemcpyHostToHost);
  

  kind的类型：

  • cudaMemcpyHostToHost : 主机 → 主机
  • cudaMemcpyHostToDevice : 主机 → 设备
  • cudaMemcpyDeviceToHost : 设备 → 主机
  • cudaMemcpyDeviceToDevice : 设备 → 设备
  • cudaMemcpyDefault : 默认
    • 默认方式只允许在支持统一虚拟寻址的系统上使用

3. 内存初始化

• 主机内存初始化：void *memset(void *str, int c, size_t n);

  代码：

  memset(fpHost_A, 0, nBytes);
  

• 设备内存初始化：__host__ cudaError_t cudaMemset(void *str, int c, size_t n);

  代码：

  cudaMemset(fpDevice_A, 0, nBytes);
  

初始化的原因：当我们分配了一段内存地址，如果不进行初始化，这段内存地址可能是无意义的，当程序进行访问时会报错。

4. 内存释放

• 释放主机内存

  代码：

  free(pHost);
  

• 释放设备内存

  代码：

  cudaFree(pDevice_A);
  

自定义设备函数

1. 设备函数
   1. 定义只能执行在GPU设备上的函数为设备函数
   2. 设备函数只能被和函数和其他设备函数调用
   3. 设备函数用 __device__ 修饰
2. 核函数
   1. 用 __global__ 修饰的函数为核函数，一般由主机调用，在设备中执行
   2. __global__ 修饰符既不能和 __host__ 同时使用，也不可与 __device__ 同时使用
3. 主机函数
   1. 主机端的普通 C++ 函数可用 __host__ 修饰
   2. 对于主机端的函数 __host__ 修饰符可忽略
   3. 可以用 __host__ 和 __device__ 同时修饰一个函数减少冗余代码，编译器会针对主机和设备分别编译该函数

CUDA错误检查

cudaError-t

运行时API成功，返回 cudaSuccess

enum cudaError_t 是个枚举变量

错误检查函数

获取名称：cudaGetErrorName

获取描述信息：cudaGetErrorString

cudaError_t ErrorCheck(cudaError_t error_code, const char *filename, int lineNumber)
{
	if (error_code != cudaSuccess)
	{
		printf("CUDA error:\r\ncode=%d, name=%s, description=%s\r\nfile=%s, line=%d\r\n",
		error_code, cudaGetErrorName(error_code), cudaGetErrorString(error_code), filename, lineNumber);	
	}
	return error_code;
}

1. 在调用CUDA运行时API时，调用ErrorCheck函数进行包装
2. 参数 filename一般使用__FILE__; 参数 lineNumber 一般使用 __LINE__
3. 错误函数返回运行时API调用的错误代码

核函数错误检查

• 错误检测函数问题：不能捕捉调用核函数的相关错误（核函数返回空）
• 捕捉调用核函数可能的发生错误的方法

ErrorCheck(cudaGetLastError(), __FILE__, __LINE__);
ErrorCheck(cudaDeviceSynchronize(), __FILE__, __LINE__); // 调用同步函数, 让 CPU 等待 GPU 设备执行完毕再进行统一数据处理

CUDA计时

使用CUDA事件(event)计时方式：

cudaEvent_t start, stop;
ErrorCheck(cudaEventCreate(&start), __FILE__, __LINE__);
ErrorCheck(cudaEventCreate(&stop), __FILE__, __LINE__);
ErrorCheck(cudaEventRecord(&start), __FILE__, __LINE__);

cudaEventQuery(start);

 
 // your function ...
 
 ErrorCheck(cudaEventRecord(stop), __FILE__, __LINE__);
 ErrorCheck(cudaEventSynchronize(stop), __FILE__, __LINE__);
 float elapsed_time;
 ErrorCheck(cudaEventElapsedTime(&elapsed_time, start, stop), __FILE__, __LINE__);
 printf("Time = %g ms\n", elapsed_time);
 
 ErrorCheck(cudaEventDestroy(start), __FILE__, __LINE__);
 ErrorCheck(cudaEventDestroy(stop), __FILE__, __LINE__);

nvprof 性能刨析

执行命令 nvprof ./exe_name

运行时GPU查询

• 运行时API查询GPU信息

  • 调用 cudaDeviceProp prop; 之恩那个在主机调用

    ErrorCheck(cudaGetDeviceProperties(&prop, device_id), __FILE__, __LINE__);
    

  • 查询GPU计算核心数量

    根据GPU计算能力进行查询

组织线程模型

• 二维网格二维线程块
• 二维网格一维线程块
• 一维网格一维线程块

二维网格二维线程块

每个线程可负责一个矩阵的计算任务

线程与二维矩阵映射关系：

ix = threadIdx.x + blockIdx.x * blockDim.x;

iy = threadIdx.y + blcokIdx.y * blockDim.y;

二维矩阵与全局索引之间的关系：

idx = iy * nx + ix

二维网格一维线程块和二维网格二位线程块类似

一维网格一维线程块

每个线程负责矩阵一列的计算

编写核函数时，需要使用循环→限制性能，不是真正的并行计算

线程与二维矩阵的映射关系：

ix = threadIdx.x + blockIdx.x * blockDim.x;
iy 由循环指定；

GPU 硬件资源

流多处理器 - SM

GPU并行依靠流多处理器SM（streaming multiprocessor）来完成

一个GPU由多个SM构成，Fermi架构SM关键资源如下：

1. CUDA 核心（CUDA Core）
2. 共享内存 / L1 缓存（shared memory / L1 Core）
3. 寄存器文件（Register File）
4. 加载和存储单元（Load / Store Units）
5. 特殊函数单元（Special Function Unit）
6. Warps 调度（Warps Scheduler）

GPU 每个 SM 都可以支持数百个线程并发执行

以线程块 block 为单位，向 SM 分配线程块，多个线程块可被同时分配到一个可用的 SM 上。

当一个线程块被分配好SM后，就不可以再分配到其他SM上了。

软件	硬件
Thread	CUDA Core
Thread Block	SM

线程块内所有线程分配到同一个SM中执行，但每个SM上可以被分配多个线程块

线程块分配到SM中后，会以 32 个线程为一组进行分割，每个组成为一个wrap

每个线程束中只能包含同一线程块中的线程。

每个线程束中只能包含同一线程块中的线程。

每个线程束包含32个线程。

线程束在GPU 硬件上真正做到了并行

CUDA 内存模型概述

寄存器

寄存器是片上（on-chip）内存，具有GPU上最快的访问速度。

寄存器仅线程内可见，生命周期也与所属线程一致

核函数中定义的不加任何限定符的变量一般存放在寄存器中；内建变量存放在寄存器中

核函数中定义的不加任何限定符的数组有可能存放于寄存器中，但也有可能存放在本地内存中。

寄存器都是32位的，保存一个 double 需要两个寄存器，寄存器保存在SM的寄存器文件中。

计算能力 5.0 - 9.0 的 GPU，每个 SM 都有 64K 的寄存器数量。

Fermi架构只有 32K

每个线程使用的最大数量不同架构是不同的，计算能力 6.1 是 64K

每个线程的最大寄存器数量是 255 个 ，Fermi 架构是 63 个

本地内存

寄存器放不下的内存会存放在本地内存

• 索引值不能在编译时确定的数组
• 可能占用大量寄存器空间的较大本地结构体和数组
• 任何不满足核函数寄存器限定条件的变量

每个线程最多可使用 512KB 的本地内存

本地内存从硬件角度看只是全局内存的一部分，延迟也很高，本地内存的过多使用，会降低程序的性能

对于计算能力 2.0 以上的设备，本地内存可以存储在每个 SM 的一级缓存中和设备的二级缓存中。

寄存器溢出

和函数所需寄存器超出硬件支持时，数据则会保存在本地内存中：

1. 一个SM并行运行多个数据块/线程束。总的需求寄存器容量大于 64 KB
2. 单个线程运行所需寄存器数量超过 255 个

寄存器溢出会降低程序性能

全局内存

全局内存在片外（容量最大，延迟最大，使用最多），全局内存的数据所有线程可见，Host段可见，且具有与程序相同的生命周期：

全局内存初始化

• 动态全局内存：主机代码中使用 CUDA 运行时 API cudaMalloc 动态声明内存空间，由 cudaFree 释放内存
• 静态全局内存：使用 __device__ 关键字静态声明全局内存

主机中不能直接访问静态全局变量

可以通过 cudaMemcpyToSymbol 和 cudaMemcpyFromSymbol 访问

__host__ cudaError_t cudaMemcpyToSymbol(

const void *symbol
, const void *src, size_t count, size_t offset, cudaMemcpyKind kind)
__host__ cudaError_t cudaMemcpyToSymbol(
void *dst, const void *symbol, size_t count, size_t offset, cudaMemcpyKind kind)

__device__ int d_x = 1; 必须声明在所有函数（Host，Device，核函数）之外

共享内存

• 共享内存在片上，仅次于寄存器的速度，有更高的带宽和更低的延迟

• 共享内存中的数据在线程块内所有线程可见，可用线程间通信，共享内存的生命周期也与所属线程块一致。

• 使用 __share__ 修饰的变量存放于共享内存中，共享内存可定义为动态与静态两种

• 每个SM的共享内存数量是一定的。也就是说如果在单个线程块中分配过度的共享内存，将会限制活跃线程束的数量。

• 访问共享内存必须加入同步机制。

  线程块内同步 void __synchreads()

不同计算能力的架构，每个SM中拥有的共享内存大小是不同的

全局内存的作用：

• 减少全局内存访问次数，提高访问效率。
• 改变全局内存访问内存的内存事务方式，提高数据访问带宽

静态全局内存 __shared__

作用域

• 核函数中：核函数内
• 和函数外：所有核函数均有效

静态共享内存在编译时就要确定内存大小。

使用动态共享内存

extern
 __shared__ float 
s_array
[]
;

调用核函数时： kernel_1 <<<grid, block, 32>>>(d_A, nElems);

常量内存

常量内存作用

• 常量内存时有常量缓存的全局内存，数量有限，64KB，线程在读取相同的常量内存数据时，访问速度比全局内存块；
• 常量内存中的数据对同一编译单元内所有线程可见
• 使用 __constant__ 修饰变量，不能定义在核函数中，且常量内存时静态定义的
• 常量内存仅可读，不可写
• 给核函数传递数值参数时，这个变量就存放于常量内存
• 常量内存必须在主机端使用 cudaMemcpyToSymbol 进行初始化
• 线程束中所有线程要访问同一地址时，常量内存表现最好。因为线程束中所有的线程都需要读取同一地址空间的系数数据，因此只需要读取一次，广播给线程束中的所有线程