Nvidia CUDA初级教程6 CUDA编程一

视频：https://www.bilibili.com/video/BV1kx411m7Fk?p=7
讲师：周斌

GPU架构概览

GPU特别使用于：
- 密集计算，高度可并行计算
- 图形学
晶体管主要被用于：
- 执行计算
- 而不是
  - 缓存数据
  - 控制指令流

在这里插入图片描述

图中分别是CPU、GPU各个部件所占的芯片面积。可以看到，CPU芯片中大量部分是缓存和控制逻辑，而GPU中则绝大部分都是计算单元。

CUDA编程相关简介

CUDA的一些信息

层次化线程集合
共享存储
同步

CUDA术语

主机端和设备端

HOST - 主机端，通常指CPU
- 采用ANSI标准C语言编程
Device - 设备端，通常指GPU（数据可并行）
- 采用ANSI标准C的扩展语言编程（CUDA C）
HOST 和 Device 拥有各自的存储器
CUDA编程
- 包括主机端和设备端两部分代码

核

Kernel 数据并行处理函数
- 通过调用 Kernel 函数在设备端创建轻量级的线程，线程由硬件负责创建并调度

类似于 OpenCL 的 shader？

核函数会在 N 个不同的 CUDA 线程上并行执行

// 定义核函数
__global__ void VecAdd(float* a, float* B, float* C) {int i = threadIdx.x;C[i] = A[i] + B[i];
}int main() {// ...// 在N个线程上调用核函数VecAdd<<<1, N>>>(A, B, C);
}

CUDA程序的执行

CUDA程序执行的流程大体上是这样的：当我们在CPU端的代码是串行执行的（这里简单地认为指令在CPU上串行执行），当遇到需要并行大量处理数据时，会调用核函数在GPU上进行计算，计算完成后将结果返回给CPU。

在这里插入图片描述

线程层次

Grid - 一维或多维线程块（block）
- 一维或二维
Block - 一维线程
- 一维，二维或三维
- 一个 Grid 中的每个 Block 的线程数是一样的
- Block 内部的每个线程可以：
  - 同步 Synchronize
  - 访问共享存储器 shared memory

一个线程可以类比为一个员工，一个 block 是一个科室，grid 是整个公司。

在这里插入图片描述

线程ID

每一个线程都有一个索引：threadIdx

一维 Block Thread ID == Thread Index
二维 Block (Dx, Dy)
- 索引为 (x, y) 的 Thread ID == x + yDy
三维 Block (Dx, Dy, Dz)
- 索引为 (x, y) 的 Thread ID == x + yDy + zDxDy

代码实例

__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {int i = threadIdx.x;int j = threadIdx.y;C[i][j] = A[i][j] + B[i][j];
}int main() {int numBlocks = 1;dim3 threadsPerBlock(N, N);MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
}

每个 Block 中的线程的索引是二维的，这在我们处理二维数据（矩阵）时可以很方便地进行对应。

线程数

Thread Block 线程块

线程的的集合
- G80 和 GT200：多达512个线程
- Fermi：多达1024个线程
位于相同的处理器核（相同的SM）
共享所在核的存储器

在这里插入图片描述

块索引

块索引：blockIdx
维度：blockDim
- 一维，二维或三维

__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {int i = blockIdx.x * blockDim.x + threadIdx.x;int j = blockIdx.y * blockDim.y + threadIdx.y;if (i < N && j < N)C[i][j] = A[i][j] + B[i][j];
}int main() {// ...dim3 threadsPerBlock(16, 16);dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
}

例如 N = 32

每个块有16x16个线程（跟N无关）
- threadIdx([0, 15], [0, 15])
Grid 里面有 2x2 个线程块 block
- blockIdx([0, 1], [0, 1])
- blockDim = 16

int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;

i = [0, 1] * 16 + [0, 15]

线程块之间

线程块彼此之间独立执行

任意顺序：并行或串行
被任意数量的处理器以任意顺序调度
- 处理器的数量具有可扩展性

一个块内部的线程

一个块内部的线程有一些重要的特性：

共享容量有限的低延迟存储器 (shared memory)
同步执行
- 合并访存
- __syncThreads()
  - barrier - 块内线程一起等待所有的线程都
  - 轻量级线程

CUDA内存传输

主机端与设备端

CUDA内存传输

在这里插入图片描述

device 端代码可以：
- 读写该线程的 registers
- 读写该线程的local memory
- 读写该线程所属的块的 shared memory
- 读写grid的 global memory
- 只读grid的 constant memory
host 端代码可以：
- 读写grid的 global memory 和 constant memory
host 可以从 device 往返传输数据
- global memory 全局存储器
- constant memory 常量存储器

CUDA内存传输函数

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

float* Md;
int size = Width * Width * sizof(float);
cudaMalloc((void**)&Md, size);
//...
cudaFree(Md);

注意这里的指针 Md 是指向 device（GPU）上的存储空间。

内存传输：cudaMemcpy()
- host to host
- host to device
- device to host
- device to device

cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);
cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);

示例：矩阵相乘 Matrix Multiply

矩阵相乘简介

向量
点乘
行优先或列优先
每次点乘结果输出一个元素

在这里插入图片描述

矩阵相乘CPU实现

void MatrixMulOnHost(float* M, float* N, float* P, int width) {for (int i=0; i<width; ++i) {for (int j=0; j<width; ++j) {float sum = 0;for (int k=0; k<width; ++k) {float a = M[i * width + k];float b = N[k * width + j]:sum += a  * b;}P[i * width + j] = sum;}}
}

CUDA算法框架

三步走：

int main(void) {// 1 分配device空间// 2 在GPU上，并行计算MatrixMulOnDevice(M, N, P, width);// 3 将结果拷贝回CPU，并释放device空间return 0;
}

伪代码如下：

void MatrixMulOnDevice(float* M, float* N, float* P, int Width) {int size = Width * Width * sizeof(float);// 1cudaMalloc(Md, size);cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);cudaMalloc(Nd, size);cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice);cudaMalloc(Pd, size);// 2 调用cuda核函数，并行计算// 3cudaMemcpy(P. Pd, size, cudaMemcpyDeivceToHost)cudaFree(Md); cudaFree(Nd); cudaFree(Pd);
}