前言

随着深度学习的发展,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 the base and includes the [CUDA math libraries](https://developer.nvidia.com/gpu-> accelerated-libraries), and NCCL. A runtime image that also includes cuDNN is available.

devel: Builds on the runtime and 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