HCCL集合通信测试环境设置

1. 快速使用 HCCL(Huawei Collective Communication Library)是基于昇腾AI处理器的高性能集合通信库,其主要功能与作用与Nvidia的NCCL库相似,主要用于集合通信,CANN库种自带一套测试工具用以分析集合通信性能。 1.1 编译环境配置 前置环境配置阶段请参考Llama2部署记录,后续的所有流程均需要在具备HCCL硬件的平台上实现,需要注意的是,建议使用的Ubuntu版本大于等于20.04LTS,否则可能会遇到VSCode不支持的情况。 A. CANN CANN的配置过程参考Llama2部署记录文档的4. 安装CANN章节,并记录安装位置,以备后续使用 完成CANN后,建议在~/.bashrc文件中添加对应的环境变量,例如: # 使用nano打开~/.bashrc文件 nano ~/.bashrc 在文件的最后添加export ASCEND_HOME=/usr/local/Ascend/ascend-toolkit/latest B. OpenMPI 在OpenMPI官网找到对应版本的安装包下载,然后解压缩安装包,以mpich-4.2.0版本为例: # mpich下载链接参考:https://www.mpich.org/static/downloads wget https://www.mpich.org/static/downloads/4.2.0/mpich-4.2.0.tar.gz # 解压缩mpich压缩包 tar -zxvf mpich-4.2.0.tar.gz 配置mpi的安装路径,以/root/mpich路径为例: # 新建文件夹 mkdir ~/mpich # 进入解压缩后的mpich文件夹 cd mpich-4.2.0 # 配置安装路径 ./configure -prefix=/root/mpich --disable-fortran 使用make命令构建mpich的可执行文件,该过程可能耗时比较长 make make install 完后后将mpich添加到环境变量种,并添加对应的链接库和帮助文件 # 修改~/.bashrc文件 nano ~/.bashrc 在最后一行添加以下内容: export MPI_HOME=/root/mpich export PATH=$MPI_HOME/bin:$PATH ...

2024年4月18日 · 2 分钟

Llama2部署记录

1. 创建docker 1.1 docker命令 显示全部的镜像:docker images 创建一个容器:docker run ...<查看1.2节内容> 显示全部的容器:docker ps -a 启动一个容器:docker start <id> 停止一个容器:docker stop <id> 进入一个正在运行的程序:docker exec -it <id> /bin/bash 删除一个容器:docker rm <id> 1.2 创建容器 创建一个docker容器的时候需要同时挂载本地工作目录和相关的程序驱动,因此创建docker容器的命令参数较多,以下为一个创建docker容器命令的基本格式: docker run -it --ipc=host --name <name> -v <local_path>:<docker_path> \ --workdir=<docker_path> \ --pids-limit 409600 \ --privileged --network=host \ --shm-size=128G \ --device=/dev/davinci0 \ --device=/dev/davinci1 \ --device=/dev/davinci2 \ --device=/dev/davinci3 \ --device=/dev/davinci4 \ --device=/dev/davinci5 \ --device=/dev/davinci6 \ --device=/dev/davinci7 \ --device=/dev/davinci_manager \ --device=/dev/devmm_svm \ --device=/dev/hisi_hdc \ -v /usr/local/Ascend/driver:/usr/local/Ascend/driver \ -v /usr/local/Ascend/firmware:/usr/local/Ascend/firmware \ -v /usr/local/Ascend/add-ons/:/usr/local/Ascend/add-ons/ \ -v /usr/local/bin/npu-smi:/usr/local/bin/npu-smi \ -v /usr/local/dcmi:/usr/local/dcmi \ -v /etc/ascend_install.info:/etc/ascend_install.info \ -v /etc/bashrc:/etc/bashrc \ -p 223:223 \ -p 224:6006 \ -p 225:8080 \ -p 226:8888 \ -u root \ <image_id> /bin/bash -it:已交互方式运行容器,并分配一个终端 ...

2024年4月9日 · 5 分钟

3D并行(3D Parallelism)

参考文献: Narayanan D, Shoeybi M, Casper J, et al. Efficient Large-Scale Language Model Training on GPU Clusters Using Megatron-LM: SC’ 21, November 14-19, 2021, St. Louis, MO, USA \[C\], 2021. ACM, 2021-01-01. 单一的MP或DP无法硬度继续增长的大模型算力需求,面对成千上万的GPU集群规模增长,结合TP、PP、DP等多种方式组成3D Parallelism的并行方式逐渐成为了主流的并行计算框架。 1. 基于Megatron-LM 1.1 研究背景 单一的DP、TP、PP方法无法满足日益增长的模型参数量,因此提出一种基于多种并行方法的PTD-P并行技术。 1.2 实现方法 图. 默认的Pipeline调度(上)和插值的Pipeline调度(下) 对每个设备上运行的网络层进行进一步拆分,每个子集称为一个Chunk,在图(下)中用深色和浅色两种颜色表示 假设一共4个Device,每个Device可以分为4层(例如:在Device 1上运行1~4层,Device 2上运行5-8),则可以将每个设备的计算过程分为两个Chunk(每个Chunk需要计算两层网络)例如:Device 1对应1,2,9,10四层,Device 2对应3,4,11,12四层。 以Device 3为例,为什么运行完Chunk 2的Batch2后需要Wait,而不是直接运行Batch 3?只是为了展示插值的概念? 1.3 实验结论 提出一种插值Pipeline调度算法,实现在3072个GPU上训练1T参数的训练,总体算力达到502PFLOPS,单个GPU的算力达到了理论峰值的52%。 相比较ZeRO-3(不带MP),因为减少了通信量,在175B和530B规模的大模型上效果要好70% 插值Pipeline调度算法可以增加计算密集度,但是对通信的开销也会增加 PP在较大的模型上效率更高,效果更好,TP的并行方式会增加不同设备之间的通信量 重算技术只是一个为了平衡内存与算力之间可选的技术,会降低内存消耗,但同时也会增加1/3的算力消耗 高度的MP可能会导致较小的矩阵乘运算,降低GPU的利用率 1.4 实验数据 本文对PTD-P模型中,GPU数量、全局Batch Size、Microbatch Size、Bubble Time、Throughput等多个参量进行多角度多变量讨论 ...

2024年4月1日 · 1 分钟

Nvidia GPU与Huawei NPU

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拥有完全不同的设计思路。 ...

2024年4月1日 · 2 分钟

流水线并行(Pipeline Parallelism)

参考文献: Huang Y, Cheng Y, Chen D, et al. GPipe: Efficient Training of Giant Neural Networks using Pipeline Parallelism: 33rd Conference on Neural Information Processing Systems (NeurIPS 2019), \[C\], Vancouver, Canada, 2018. Harlap A, Narayanan D, Phanishayee A, et al. PipeDream: Fast and Efficient Pipeline Parallel DNN Training \[J\]. ArXiv, 2018,abs/1806.03377. Kim T, Kim H, Yu G, et al. BPipe: Memory-Balanced Pipeline Parallelism for Training Large Language Models: Proceedings of the 40th International Conference on Machine Learning ...

2024年4月1日 · 2 分钟