目录
一、cuda编程的基本概念入门
1、GPU架构和存储结构
2、cuda编程模型
3、cuda编程流程
二、cuda向量加法实践
1、代码实现
2、代码运行和结果
有一段时间对模型加速比较感兴趣,其中的一块儿内容就是使用C++和cuda算子优化之类一起给模型推理提速。之前一直没有了解过cuda编程,也没有学习过C++相关的东西。强迫自己来学习一下cuda编程,同时也学习一下C++,扩宽一下AI相关的领域知识。主要是能够理解怎么使用cuda来提升模型的推理速度,学习的目标就是要会使用cuda编程实现基本的向量加法乘法、能使用C++和cuda混合编程实现神经网络的一些基本模块、最终能够完成C++语言和cuda混合编程(自己实现算子或者调用英伟达成熟的库)完成一个LLM模型的前向推理过程。这里是cuda编程的第一篇入门篇,了解基本概念、gpu的架构、cuda编程模型和实现cuda向量加法,对cuda编程有一个基础的了解和实践。
一、cuda编程的基本概念入门
1、GPU架构和存储结构
GPU全称图形处理器(graphics processing unit),主要是做图像和图像等涉及到并行计算的微处理器。 GPU和CPU同样有自己的架构,GPU更重计算、CPU更重逻辑控制。从硬件层面来说,GPU的内部构成如下图——详解GPU:
GPU通常包括图形显存控制器、压缩单元、BIOS、图形和计算阵列、总线接口、电源管理单元、视频管理单元、显示接口等,我们用来做深度学习就主要用到了它的图形和计算阵列模块。
GPU微架构
从微架构角度来说,GPU是有一个个SM(Streaming Multiprocessors)构成的。如下图:
这是A100 安培架构显卡的SM内部结构图,SM由L1缓存、指令缓存、寄存器(Register)和Wrap scheduler等构成。图中的绿色部分是tensor core 也可以称作SP(Streaming processor),用于浮点数的计算,它可以支持一个时钟周期完成两个16×16矩阵的乘法操作,其他版本如Volta完成两个两个4×4半精度浮点矩阵的计算、Turing完成64个半精度浮点的乘加操作,总之计算速度更慢。
内存模型
GPU的内存也是多层级结构的,具体结构如下:
通用内存DRAM目前最好的显卡采用了HBM(High Bandwidth Memory 高带宽内存);更近一级的是L2缓存,所有的SM共享;L2之上就是L1缓存,SM独有,所有的显存共享;L1之上的就是寄存器,线程独有的。如图,英伟达的cuda编程指南中给出示意图:
2、cuda编程模型
GPU其实可以看做一个超多线程处理器,一个运行多次使用不同数据执行的程序,可以使用很多不同的线程来执行,在GPU上就是把这个函数编译为设备的指令集——kernel核函数。cuda就是实现这样功能的一个代码库,可以让开发者使用高级语言来实现上述GPU的多线程并行执行,加速计算速度。
图中显示一个kernel会被grid中的线程块一起执行。这里就有几个概念,gird、block和thread。一个grid有多个block构成;一个block有多个thread构成。其中grid中的block有x/y/z三个维度,总数有最大值,每个维度上有各自的最大值,需要查阅当前的cuda规范。同时block中的线程也分x/y/z三个维度,总数有最大值,每个维度上有各自的最大值。一般来说,block中的线程数最大为1024个。线程的序号由block数目和线程在block中的位置,对于上述kernel1,thread(4,2)来说,线程Id
threadId
=(threadIdx.x+threadIdx.y*blockDim.x)+(blockIdx.x+blockIdx.y*gridDim.x)*(blockDim.x*blockDim.y)
=(4+2*5)+(1+1*3)*(5*3)=14+60=74
内存模型
线程在内存的使用是什么样的?GPU的内存模型如下:
block中的线程共用shared Mem,线程独立拥有寄存器和本地内存,其它的内存都是所有的block共享的。
3、cuda编程流程
cuda编程流程其实有点像我们使用GPU进行模型训练,模型训练中首先是模型加载和数据的处理;然后是把模型参数和数据都从CPU内存移动到GPU内存(显存)上;最后进行模型训练。那cuda编程宏观上也是这么个逻辑,这里摘抄一段知乎博主小小将博文《CUDA编程入门极简教程》总结的流程如下:
- 分配host内存,并进行数据初始化;
- 分配device内存,并从host将数据拷贝到device上;
- 调用CUDA的核函数在device上完成指定的运算;
- 将device上的运算结果拷贝到host上;
- 释放device和host上分配的内存。
cuda编程的重点和难点也在于第三步cuda核函数的设计实现(设计一个跑通的可能不难但是设计一个高效率的可能就很难了),核函数的定义如下:
__global__ void kernelFunction(float *result, float *a, float *b){
doSomething
}
使用__global__对核函数进行限定,表示该函数是一个GPU核函数,在GPU的线程中被执行。
使用一个核函数整体的代码流程如下:
__global__ void kernelFunction(float *result, float *a, float *b){
doSomething
}
int main(){
......
// 分配内存和显存
cudaMallocManaged();
//数据初始化
initWith();
// 每一个gird有多少个block 最大2^31-1 x方向最大2^31-1 y,z 方向65535
dim3 gridDim(x,y,z);
// 每一个block有多少个线程 最大1024 x,y方向最大1024 z最大64
dim3 blockDim(x,y,z);
//执行核函数
kernelFunction <<< gridDim, blockDim >>>();
cudaDeviceSynchronize(); // 同步
......
}
使用kernelFunction<<<gridDim,blockDim>>>()来指定对应的gridDim和blockDim并且启动和函数。根据wiki的数据显示:每一个gird有多少个block 最大2^31-1 x方向最大2^31-1 y,z 方向65535; 每一个block有多少个线程 最大1024 x,y方向最大1024 z最大64。
二、cuda向量加法实践
1、代码实现
接下来基于cuda来实现两个一维矩阵(一维向量)的加法。按照上述cuda编程流程,首先需要进行数据初始化,然后把数据传输到GPU上,然后进行cuda核函数的计算,最后得到结果释放资源。首先看一下数据怎么在CPU和GPU上灵活的传输,新版本的cuda有如下API:
cudaError_t cudaMallocManaged(void** ptr, size_t size)
该函数运行我们在内存和显存开辟size_t大小的空间,并智能的进行数据在CPU和GPU上的移动。
现在需要设计核函数,简单起见我们设置grid和block都为一维的,核函数逻辑就可以按照如下设计:
__global__ void addVectorskernel(float *result, float *a, float *b, int N){
int index = threadIdx.x + blockIdx.x * blockDim.x;
int stride = blockDim.x * gridDim.x;
for (int i=index; i<N; i+=stride){
result[i] = a[i] + b[i]; // 元素a[i] + 元素 b[i]
}
}
其中blockDim就表示一个block中有多少个线程,gridDim表示一个grid(一个gpu)中有多个block,那么总线程数就是blockDim.x * gridDim.x,每个线程处理的向量元素就是N/(blockDim.x * gridDim.x),因此就会有内部的循环,循环的步长也是总线程数blockDim.x * gridDim.x。当N=102400000,blockDim.x= 256,gridDim.x = 10,a矩阵的值全为3.0,b矩阵的值全为4.0,那么就可以得到如下代码nvcc_vector_add.cu:
#include<stdio.h>
#include<assert.h>
#include<cstdio>
#include<sys/time.h>
#include<iostream>
// 编译加链接
// nvcc -o nvcc_vector_add.cu nvcc_vector_add.o
// 直接运行即可
// 向量加法核函数
__global__ void addVectorskernel(float *result, float *a, float *b, int N){
int index = threadIdx.x + blockIdx.x * blockDim.x;
int stride = blockDim.x * gridDim.x;
for (int i=index; i<N; i+=stride){
result[i] = a[i] + b[i]; // 元素a[i] + 元素 b[i]
}
}
// 初始化数组 a
void initWith(float num, float *a, int N) {
for(int i = 0; i < N; ++i) {
a[i] = num;
}
};
int main(){
const int N = 102400000;
const int M = 10;
size_t Mem = N * sizeof(float);
float *a;
float *b;
float *c;
cudaMallocManaged(&a, Mem);
cudaMallocManaged(&b, Mem);
cudaMallocManaged(&c, Mem);
initWith(3.0, a, N); // 将数组a中所有的元素初始化为3
initWith(4.0, b, N); // 将数组b中所有的元素初始化为4
initWith(0.0, c, N); // 将数组c中所有的元素初始化为0,数组c是结果向量
for(int i=0;i<M;i++){
printf("%f ",a[i]);
}
printf("\n");
printf("******************\n");
for(int i=0;i<M;i++){
printf("%f ",b[i]);
}
printf("\n");
printf("******************\n");
for(int i=0;i<M;i++){
printf("%f ",c[i]);
}
printf("\n");
printf("******************\n");
// 配置参数
size_t threadsPerBlock = 256;
// size_t numberOfBlocks = (N + threadsPerBlock - 1) / threadsPerBlock;
size_t numberOfBlocks = 10;
struct timeval start;
struct timeval end;
gettimeofday(&start,NULL);
addVectorskernel <<< numberOfBlocks, threadsPerBlock >>> (c, a, b, N); // 执行核函数
cudaDeviceSynchronize(); // 同步,且检查执行期间发生的错误
gettimeofday(&end,NULL);
float time_use;
time_use=(end.tv_sec-start.tv_sec)*1000000+(end.tv_usec-start.tv_usec);//微秒
std::cout <<"vector_add gpu time cost is "<<time_use/1000/100<< " ms"<< std::endl;
for(int i=0;i<M;i++){
printf("%f ",a[i]);
}
printf("\n");
printf("******************\n");
for(int i=0;i<M;i++){
printf("%f ",b[i]);
}
printf("\n");
printf("******************\n");
for(int i=0;i<M;i++){
printf("%f ",c[i]);
}
printf("\n");
printf("******************\n");
return 0;
}
2、代码运行和结果
以上是一个.cu单文件,怎么运行它呢?需要使用英伟达的cuda环境进行编译和链接。在安装cuda toolkit的环境下,使用nvcc编译器进行编译:
//编译和链接
nvcc nvcc_vector_add.cu -o nvcc_vector_add
//运行
./nvcc_vector_add
参数-o表示生成可执行文件,编译结果如下:
编译后生成一个可执行文件,直接运行该可执行文件,得到如下结果:
可以看到结果正确,耗时为2ms,显存使用384M。
调整一下blockDim、gridDim的大小看看耗时的变化情况。
blockDim=256,gridDim=5
vector_add gpu time cost is 2.23614 ms
blockDim=256,gridDim=10
vector_add gpu time cost is 2.01713 ms
blockDim=256,gridDim=20
vector_add gpu time cost is 2.05501 ms
blockDim=128,gridDim=10
vector_add gpu time cost is 2.1049 ms
blockDim=512,gridDim=10
vector_add gpu time cost is 2.027 ms
以上结果具有误差因为都只跑了一次,没有多次求平均值,但是可以说明gridDim和blockDim对性能是有影响的。一般来说blockDim选择为32的倍数,因为一个wrap的线程束是32,blockDim这样设置可以减少bank conflict。
CUDA编程入门极简教程
Cuda Core VS Tensor Core