CUDA GPU编程

GPU史前文明

HPC高性能计算：是一个很早就有的领域. GPU和CPU的结合在HPC领域成为趋势。

冯诺依曼架构处理器

当代几乎所有的处理器都是冯诺依架构： 处理器持续从存储器取指令、指令译码、执行指令的指令周期。 CPU由控制器、存储器、运算器组成。

内存速度和时钟速度的比率是限制处理器吞吐量的重要因素。大多数程序都是内存受限型而不是时钟受限型，CPU引入三级缓存来加速。 缓存越多，芯片成本越高，成品难度也越大。

克雷向量处理器

70年代，克雷计算机，其设计结构对后续GPU有很深的影响

连接机并行指令集

80年代，连接机，其SIMD处理方式与后续处理器SIMD基本一致

SIMD: 一条指令处理多条数据，用于小碎数据的并行操作。比如i7是64位处理器，一次最多处理64位（8个字节）数据。早先处理这些数据只能用于他们的低８位，如果把64位寄存器拆成８个８位寄存器就能同时完成８个操作，计算效率提升了８倍．

CELL处理器：

20世纪初，IBM、索尼合作研发，它和之后Nvidia的GPU很类似。

多点计算：

Bonic、Folding@home、Bitcoin、Hadoop等点计算项目出现。

单节点算力受主频太高，电力和散热成本上升，收益递减。使用多个单核实现集群并行处理(网络互联、PCIE互联)，成本更低。

GPGPU编程

GPU的应用早已不限于图像领域。 使用GPU做通用目的的并行计算（而不是仅仅用于图形渲染）。 这已经跟CUDA的设计思想一致了 http://graphics.stanford.edu/courses/cs148-10-fall/lectures/programmable.pdf

典型操作

GPU产品线

GPU 首字母 Graphics， 是为了图形处理而设计的处理器。

GPU原先作为显卡的芯片处理器，随着并行计算的发展，发展到计算卡的领域。 GPU和显卡(图形卡) -> GPU和计算卡

显卡都有DVI接口、VGA接口、HDMI接口，而计算卡/加速卡只有PCIE接口。

Nvida GPU

AMD GPU

Intel GPU

ARM GPU

Apple GPU

并行设计模式

多核并发OpenMP标准

单个节点内部实现并行处理，多核处理器共享存储并行(shared-memory parallel programming)。

拆成了几个线程，然后使用事件对象等待所有线程结束。

类似技术 GCD、TBB

多机并发MPI标准

多个节点间的并行处理，计算机集群共享通信并行。

其主要瓶颈在网络通信上(Ethernet、infiniband)。

并行问题

并发时资源共享，需要引入semaphore、mutex等机制。

易并行性

有一些问题属于 “易并行” 问题：如矩阵乘法。在这类型问题中，各个运算单元输出的结果是相互独立的，这类问题能够得到很轻松的解决 (通常甚至调用几个类库就能搞定问题)。 然而，若各个运算单元之间有依赖关系，那问题就复杂了。在 CUDA 中，块内的通信通过共享内存来实现，而块间的通信，则只能通过全局内存。

局部性

操作系统原理里讲过局部性原理，简单来说就是将之前访问过的数据 (时间局部性) 和之前访问过的数据的附近数据 (空间局部性) 保存在缓存中。 计算机局部性原理指导处理器逐渐设计出了多级缓存、超线程，提供ALU的利用率。

基于任务的并行

这种并行模式将计算任务拆分成若干个小的但不同的任务，如有的运算单元负责取数，有的运算单元负责计算，有的负责…… 这样一个大的任务可以组成一道流水线。其瓶颈在于效率最低的运算单元。

基于数据的并行

这种并行模式将数据分解为多个部分，让多个运算单元分别去计算这些小块的数据，最后再将其汇总起来。

可并行化的模式

(X)PU计算单元

数据并行的简单描述：

对一个数据一个进行操作 -> 对一组数据进行一个操作。

SIMT

Nvidia GPU版的SIMD又成为SIMT，单指令多线程，其指令操作码跟CPU的不同，需要程序通过一个内核程序指定每个线程的工作内容。

SIMT 每一个core有自己的寄存器、自己的ALU、自己的data cache， 没有独立的instruction cache、没有独立的解码器、没有独立的程序计数器。

SIMD是一个单独的线程，只是这个线程能同时进行多个计算而已. 比如SIMD单元128-bit宽，包含16个128-bit的寄存器，能够被用来当做32个64-bit寄存器。这些寄存器能被当做是同等数据类型的的vector。 SIMT多个线程各有各的处理单元，和SIMD公用一个ALU不同。因而可以进行更为复杂的并行计算。

面向GPU编程思想上的差异

程序中能够并行运行的代码占多大比例？

单线程CPU程序员 vs GPU上的并行程序员

CUDA

2007年，nvidia发现了一个能使得GPU进入主流的契机，那就是为GPU制定一个标准的编程接口，这就是CUDA.

CUDA编译模型使用了和java语言一样的编译原则：基于虚拟指令集的运行时编译。

它使 NVIDIA GPU 能够执行使用 C、C++、Fortran、OpenCL、DirectCompute 和其他语言编写的程序.

CUDA的替代选择(通用并行计算平台和编程模型)

标准

概念

import numpy as np
from numba import cuda

# 创建GPU上下文 ，它用于管理GPU资源和执行GPU操作。
cuda_context = cuda.create_context()

# 创建GPU流 它用于将操作分配到不同的流中以实现并行执行。
stream = cuda.stream()

# 定义一个核函数，用于在GPU上执行并行计算
@cuda.jit
def vector_add(a, b, c):
    idx = cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x
    if idx < len(c):
        c[idx] = a[idx] + b[idx]

# 创建输入数据和输出数据
a = np.array([1, 2, 3, 4, 5])
b = np.array([6, 7, 8, 9, 10])
c = np.zeros_like(a)

# 在GPU上分配内存
d_a = cuda.to_device(a, stream)
d_b = cuda.to_device(b, stream)
d_c = cuda.to_device(c, stream)

# 在GPU上执行核函数
threads_per_block = 32
blocks_per_grid = (len(c) + (threads_per_block - 1)) // threads_per_block
vector_add[blocks_per_grid, threads_per_block, stream](d_a, d_b, d_c)

# 将结果从GPU复制回CPU
d_c.copy_to_host(c, stream)

# 同步GPU流
stream.synchronize()

# 打印结果
print(c)

# 销毁GPU上下文
cuda_context.pop()

CUDA 线程层次

一个CUDA Kernel大概可以分为这么几层（从底层到顶层）：thread < warp < block < grid PS：core和kernel两个词，core指的是硬件核心，处理器；kernel指软件核心，可重复使用的种子。

在CUDA编程模型中，一个”lane”通常指的是一个在GPU上执行的线程在其所属的warp（线程束）中的唯一索引。CUDA中的一个warp是一组同时执行的线程，NVIDIA的GPU中这个数目通常是32个线程。

shared memory

共享内存被分成了若干个等宽的存储段，这些存储段被称为“banks”，硬件上的限制是，每个bank在一个时钟周期内只能服务一个访问请求。

一个block要访问shared memory，只要能够保证以其中相邻的16个线程一组访问thread，每个线程与bank是一一对应就不会产生bank conflict。否则会产生bank conflict，访存时间成倍增加，增加的倍数由一个bank最多被多少个thread同时访问决定。有一种极端情况，就是所有的16个thread同时访问同一bank时反而只需要一个访问周期，此时产生了一次广播。

CUDA GPU 物理层次

一张GPU卡由若干个流处理簇(SM)组成，一个SM配置若干个流处理器（SP），

streaming processor(sp): 最基本的处理单元。GPU进行并行计算，也就是很多个sp同时做处理。现在SP的术语已经有点弱化了，而是直接使用thread来代替。一个SP对应一个thread.

Warp：warp是SM调度和执行的基础概念，通常一个SM中的SP(thread)会分成几个warp(也就是SP在SM中是进行分组的，物理上进行的分组)，一般每一个WARP中有32个thread.这个WARP中的32个thread(sp)是一起工作的，执行相同的指令，如果没有这么多thread需要工作，那么这个WARP中的一些thread(sp)是不工作的. （每一个线程都有自己的寄存器内存和local memory，一个warp中的线程是同时执行的，也就是当进行并行计算时，线程数尽量为32的倍数，如果线程数不上32的倍数的话；假如是1，则warp会生成一个掩码，当一个指令控制器对一个warp单位的线程发送指令时，32个线程中只有一个线程在真正执行，其他31个 进程会进入静默状态。）

tensor core 和 cuda core 的概念

用户代码  -> AI框架（PyTorch/Tensorflow/Caffe等）-> CUDA lib -> Driver -> 显卡

tensor core和cuda core 都是运算单元，是硬件名词，其主要的差异是算力和运算场景。 NVIDIA的发展历史上来讲，先是推出的cuda core，然后再推出的tensor core 场景：cuda core是全能通吃型的浮点运算单元，tensor core专门为深度学习矩阵运算设计。 算力：在高精度矩阵运算上 tensor cores吊打cuda cores。

CUDA编程

在CUDA编程中，有几种常见的数据类型，包括但不限于：

CUDA的核函数（也称为设备函数）中自动定义的变量：

__global__ void kernelFunction(int *data) {
    // 计算当前线程的全局索引
    int index = blockIdx.x * blockDim.x + threadIdx.x;

    // 使用全局索引来访问数组中的元素
    data[index] = ...; // 执行某些操作
}

// 根据实际问题来设置blockdim和griddim
dim3 blocksPerGrid(4, 1, 1); // 网格维度：4个线程块
dim3 threadsPerBlock(256, 1, 1); // 线程块维度：每个块256个线程

kernelFunction<<<blocksPerGrid, threadsPerBlock>>>(data); // <<<blocksPerGrid, threadsPerBlock>>>语法是用于在CUDA中启动核函数的特殊语法。这里，blocksPerGrid和threadsPerBlock分别指定了网格和线程块的大小。核函数内部，每个线程会使用它的全局索引来决定操作数组的哪个部分。这样可以确保所有的数组元素都被并行处理，而且每个元素只被处理一次。

#pragma unroll

CUDA Runtime API

cudaSetDevice

cudaMalloc/cudaFree

cudaMemcpy

使用__global__关键字定义CUDA内核

cudaDeviceSynchronize

cuda程序编译

main.cu一般以cu结尾后缀，使用nvcc编译

nvcc cuda_test.cu -o cuda_test

大型项目中，cmake已经完美支持cuda

enable_language(CUDA)

CUDA Library

NVIDIA HPC SDK 生态

NVIDIA Nsight 产品家族

GPU 硬件架构

CPU 被设计为以尽可能快的速度执行称为线程(thread)的一系列操作，并且可以并行执行几十个这样的线程； GPU 被设计为并行执行数千个线程（摊销较慢的单线程性能以实现更大的吞吐量）, 将更多的晶体管用于数据处理而不是数据缓存和流量控制。

CPU 一般是基于时间片轮转调度原则，每个线程固定地执行单个时间片；而 GPU 的策略则是在线程阻塞的时候迅速换入换出.

CPU

连接总线

Intel

AMD

GPU

截止2021年，按时间线顺序的8代NVIDIA GPU微架构：

GPU 算力演进

Nvidia GPU 算力（compute-capabilities）由版本号表示，有时也称为“SM 版本”。 这个版本号被用来标识 GPU 硬件支持的功能，并在运行时由应用程序使用以确定当前 GPU 上可用的硬件功能和/或指令。 算力包括主要修订号 X 和次要修订号 Y，并用 X.Y 表示。 具有相同主要修订号的设备具有相同的核心架构：

怎么查询GPU的算力版本？compute capability

$ nvidia-smi --query-gpu=compute_cap --format=csv

参考

ARM NEON 指令集

ARM从v7版本开始引入高级SIMD，称之为NEON。GCC里面有个编译选项是-fpu=neon，当这个设置时，编译器就支持了ARM SIMD，也就是neon。

GPU 硬编码

GPU 硬编码 实现 AES、H264

*****
Written by Lu.dev on 06 December 2021