1. Nvidia GPU 参考文献:
HeKun-NVIDIA. CUDA-Programming-Guide-in-Chinese
\[EB/OL\]. https://github.com/HeKun-NVIDIA/CUDA-Programming-Guide-in-Chinese.
Nvidia. Nvidia CUDA编程指南
\[M\]. 1.1. Nvidia, 2008.
1.1 GPU架构 软件层面架构:Nvidia GPU的架构在进行CUDA软件操作时可以抽象为:Thread->Block->Grid三层,通过指定这三个参数可以精确的控制目标运算放在目标的Thread上进行计算。
硬件层面架构:在硬件层可以讲GPU抽象为:SP(Streaming Processor)->SM(Streaming Multiprocessor)->Device三层,基本对应CUDA软件编程层面的Thread->Block->Grid,Nvidia的GPU在指令集上属于是SIMD(Single Instruction Multiple Data)架构,可以实现单条指令在多个数据的并行处理。
集群层面架构:在集群层面,一般按照Host-Device结构划分,其中一个Host可以视作一个由CPU、GPU、内存、硬盘等组成的主机,Device一般特指某一个指定的GPU设备,也就是软件层面的Grid。
图. Nvidia的GPU架构
A. SP / Thread GPU上每个SP主要由ALU(逻辑控制单元)与FPU(浮点运算单元)组成,可以进行简单的逻辑运算与浮点运算,每个SP在对应软件层面可以操作的一个Thread,也是开发者可以操作的最小单元,用户可以借助CUDA在Thread上实现SPMD(Single Program Multiple Data)操作。
B. MP / Block 多个Thread组成一个Block,同一个Block上的Thread共享同一块内存(Shared Memory),这种方法可以实现同一个Block上并行的Thread可以并行的对同一块内存中的数据进行操作。
在进行CUDA编程时,开发者可以借助一个SM控制多个Block中的Thread进行运算,SM中还包含一些必要的寄存器、共享内存、L1Cache、Scheduler、SPU、LD/ST单元等。
C. Device / Grid 多个Block一起共同组成一个Grid,也就是硬件层面与集群层面的Device,借助该概念可以实现集群中多个GPU之间的数据传输和并行计算。
1.2 GPU编程 Nvidia的GPU上有多个Thread可以并行的完成矩阵运算,例如在进行通用矩阵乘(GEMM,General matrix multiply)时,不同Thread在共享内存中取出对应位置的数据进行计算,并讲结果重新放置在共享内存中以实现GEMM的并行计算。
图. 使用共享内存并行计算矩阵乘
在代码层面的实现过程大致如下:
// Host层面的矩阵乘,假定矩阵的维度是BLOCK_SIZE的整数倍 void MatMul(const Matrix A, const Matrix B, Matrix C) { // 从显存中加载A和B两个矩阵的数据 cudaMalloc(..., size); cudaMemcpy(..., ..., size, cudaMemcpyHostToDevice); // 从设备显存中申请储存乘法运算结果的矩阵C cudaMalloc(&d_C.elements, size); // 调用内核 MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C); // <<<dimGrid, dimBlock>>> // 从显存中获取计算结果C cudaMemcpy(..., ..., size, cudaMemcpyDeviceToHost); // 释放显存 cudaFree(...); } // 调用MatMul()以实现Kernel层的矩阵乘 __global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) { // __global__ // 每个Thread计算矩阵C中的一个元素,并将结果累加到CValue中 float Cvalue = 0; int row = blockIdx.y * blockDim.y + threadIdx.y; int col = blockIdx.x * blockDim.x + threadIdx.x; for (int e = 0; e < A.width; ++e) Cvalue += A.elements[row * A.width + e] * B.elements[e * B.width + col]; C.elements[row * C.width + col] = Cvalue; } 2. Huawei NPU Huawei的NPU设备Ascend 910B相较于Nvidia的GPU做了一些有针对性的优化,同时由于NPU主要针对神经网络训练与推理,因此其在芯片设计架构上与Nvidia的GPU拥有完全不同的设计思路。
...