前言
随着深度学习的发展,AI算法对计算的需求量越来越大,传统的CPU串行编程已经不能满足企业对AI低延迟高性能要求, GPU并行编程越来越受到关注, 因此掌握一门GPU并行编程技术对于AI软件栈开发的人员非常必要.
关于GPU编程
目前Server端主流的GPU大部分采用NVIDIA GPU, 例如V100, A100等系列, 也有部分采用AMD 系列GPU, NVIDIA以及AMD都为GPU编程提供了相应的软件开发工具以及框架.
- NVIDIA: CUDA
- AMD: HIP
笔者对AMD GPU的软件栈略知一二, AMD的HIP编程基本上和NVIDIA CUDA非常相似, 目前已经有相应的工具可以将AMD HIP与NV CUDA的代码进行相互转换, 例如PyTorch提供的Hipfiy工具.
测试环境
- OS: Ubuntu 20.04
- CUDA: v11.4
- GCC: 10.3
- Docker: v20
- VSCode
Ubuntu CUDA开发环境快速搭建
Ubuntu上搭建CUDA开发环境有2种方式:
- NVIDIA官网下载CUDA安装包, 执行安装脚本
- 采用NVIDIA提供的CUDA docker环境, 开箱即用
在公司和企业中, 由于不同人员往往会交叉使用服务器资源,因此docker应用的比较广泛, docker可以提供一个标准化, 可复现的统一环境. 因此笔者决定采用NVIDIA提供的docker进行CUDA开发环境的创建.
Docker 环境检查
- 首先需要确保Ubuntu系统是否安装了docker:
docker --version, 为了方便使用GPU, 选择docker的版本>19 - 安装: nvidia-docker2
输出结果:
Docker version 20.10.11, build dea9396
获取NVIDIA CUDA docker
DockerHub提供了 nvidia/cuda 的docker 镜像:
- dockerHub: https://registry.hub.docker.com/
image.png
- nvidia/cuda镜像: 在dockerHub中搜索nvidia/cuda: https://registry.hub.docker.com/r/nvidia/cuda/tags
image.png
nvidia/cuda针对x86, ARM64等提供了各个版本的docker镜像, nvidia/cuda中的docker镜像主要包含3中不同的镜像:
hree flavors of images are provided:
base: Includes the CUDA runtime (cudart)runtime: Builds on thebaseand includes the [CUDA math libraries](https://developer.nvidia.com/gpu-> accelerated-libraries), and NCCL. Aruntimeimage that also includes cuDNN is available.devel: Builds on theruntimeand includes headers, development tools for building CUDA images. These images are particularly useful for multi-stage builds.
由于nvidia/docker提供了多种docker镜像, 因此我们根据自己的需求选择一个合适版本/处理器架构的docker镜像, 以 Ubuntu 20.04为例:
- 11.4.2-devel-ubuntu20.04
- 11.4.2-runtime-ubuntu20.04 包含cuda-base + CUDA 数学加速库(比如cublas, cufft) + cudnn
- 11.4.2-base-ubuntu20.04
- 11.4.2-cudnn8-runtime-ubuntu20.04
- 11.4.2-cudnn8-devel-ubuntu20.04 devel版本的docker镜像是基于runtime镜像创建的
笔者选择了一个比较全的docker镜像: 11.4.2-cudnn8-devel-ubuntu20.04
docker镜像下载: 在ubuntu终端中输入: docker pull nvidia/cuda:11.4.2-cudnn8-devel-ubuntu20.04
下载完成之后, 可以查看docker镜像: docker image ls
REPOSITORY TAG IMAGE ID CREATED SIZE
nvidia/cuda 11.4.2-cudnn8-devel-ubuntu20.04 b1539d83387e 3 months ago 9.14GB
创建CUDA Docker容器
Docker容器: docker容器是docker镜像的实例化, docker镜像运行之后的产物; 类似于进程和程序的概念, 程序是静态的代码, 进程是程序载入内存之后运行态的程序.
下载好 nvidia/docker 镜像之后, 开始启动一个docker容器, 并且进入docker:
简单的命令: docker run -it --name=test-cuda --gpus=all nvidia/cuda:11.4.2-cudnn8-devel-ubuntu20.04
不出意外, docker容器创建成功并且自动进入了docker, 检查一下环境:
nvidia-smi # 查看GPU
nvcc --version # 查看CUDA编译器版本
image.png
工程测试
Ubuntu中CUDA 的安装位置说明
一般情况下CUDA默认安装的目录: /usr/local/cuda, 存在如下目录:
- bin: 二进制目录,包含nvcc, nvprof. cuda-gdb等相关工具
- extras
- nsight-compute-2020.2.0
- nvvm
- src
- compute-sanitizer
-
include: CUDA提供的C/C++ 头文件, 例如:
cuda_runtime.h - nsightee_plugins
- README
- targets
- DOCS
- lib64: CUDA提供的so动态库
- nsight-systems-2020.3.4
- samples: CUDA演示的例子
- tools
- EULA.txt
- libnvvp
- nvml
- share
基于cmake 的简单CUDA测试程序
测试程序的功能: 两个数组简单相加, element-wise add
- main.cu
#include<iostream>
#include<algorithm>
#include<cmath>
#include<cuda_runtime.h>
void elmtwise_sum_cpu(int* arr1, int* arr2, int* out, int N) {
for(int i=0;i<N;i++) out[i] = arr1[i] + arr2[i];
}
__global__ void kernel_sum(int* arr1, int* arr2, int* out, int N) {
int thread_id = blockIdx.x * blockDim.x + threadIdx.x;
if(thread_id < N) {
out[thread_id] = arr1[thread_id] + arr2[thread_id];
}
}
void elmtwise_sum_gpu(int* arr1, int* arr2, int* out, int N) {
// 1. GPU端申请显存
int* d_arr1 = nullptr;
int* d_arr2 = nullptr;
int* d_out = nullptr;
cudaMalloc((void**)&d_arr1, sizeof(int)*N);
cudaMalloc((void**)&d_arr2, sizeof(int)*N);
cudaMalloc((void**)&d_out, sizeof(int)*N);
// 2. CPU Memory数据复制到GPU显存
cudaMemcpy(d_arr1, arr1, sizeof(int)*N, cudaMemcpyHostToDevice);
cudaMemcpy(d_arr2, arr2, sizeof(int)*N, cudaMemcpyHostToDevice);
// 3. 设置GPU端线程配置, launch the GPU kernel
int blk_size = 128;
int grid_size = (N + blk_size -1) / blk_size;
kernel_sum<<<grid_size, blk_size>>>(d_arr1, d_arr2, d_out, N);
// 4. Cpoy GPU result to CPU
cudaMemcpy(out, d_out, sizeof(int)*N, cudaMemcpyDeviceToHost);
// 5. Free GPU Memory
cudaFree(d_arr1);
cudaFree(d_arr2);
cudaFree(d_out);
}
int main() {
const int N = 512* 512;
int* arr1 = new int[N];
int* arr2 = new int[N];
int* out_cpu = new int[N];
int* out_gpu = new int[N];
srand(123456);
for(int i=0;i<N;i++) {
arr1[i] = rand() * 5 % 255;
arr2[i] = rand() % 128 + 5;
}
elmtwise_sum_cpu(arr1, arr2, out_cpu, N);
elmtwise_sum_gpu(arr1, arr2, out_gpu, N);
auto print_array = [](int* arr, int N, int k, const std::string& msg) -> void {
std::cout << msg << std::endl;
int n = std::min(N, k);
for(int i=0;i<n;i++) std::cout << arr[i] <<" ";
std::cout << std::endl;
};
print_array(out_cpu, N, 10, "CPU");
print_array(out_gpu, N, 10, "GPU");
// validate
int i=0;
for(i=0;i<N;i++){
if(out_cpu[i] != out_gpu[i]){
std::cout << "Error, not equal!" << std::endl;
break;
}
}
if(i==N) std::cout << "Test OK, all correct !" << std::endl;
delete[] arr1;
delete[] arr2;
delete[] out_cpu;
delete[] out_gpu;
return 0;
}
- CMakeLists.txt
project(TEST_CUDA LANGUAGES CXX CUDA)
cmake_minimum_required(VERSION 3.10)
# https://zhuanlan.zhihu.com/p/105721133
if(CUDA_ENABLE)
enable_language(CUDA)
endif()
add_executable(main "main.cu")
编译 & run:
mkdir -p build
cd build
cmake ../
make -j8
# run
./main
运行结果:5
image.png
程序分析:
-
典型的GPU程序执行流程: GPU端申请内存 ---> Copy data from CPU to GPU ---> Launch GPU kenrel ---> Copy result from GPU to CPU ---> Free GPU Memory
-
CUDA编程头文件:
<cuda_runtime.h>, 包含常用的CUDA函数, 例如cudaMalloc(), cudaMemcpy() 用于在显存分配空间以及CPU-GPU端数据拷贝传输 -
__global__ void kernel_sum: GPU上执行的核函数, kernel function,__global__修饰符表示此函数是一个GPU kernel function, 次函数在CPU端调用,在GPU端执行 -
GPU端线程配置
// 3. 设置GPU端线程配置, launch the GPU kernel
int blk_size = 128; --- block_size, 代表1个block中CUDA线程的数量,一般为2的幂数
int grid_size = (N + blk_size -1) / blk_size; --- gride_size: 代表全部计算需要的block个数, 注意这里需要向上取整
kernel_sum<<<grid_size, blk_size>>>(d_arr1, d_arr2, d_out, N); --- <<< grid_size, blk_size>>> CUDA特有的kernel启动方式
网友评论