其他AI加速器¶

张量处理单元(TPUs)是谷歌自主研发的专用集成电路(ASIC)，用于加速机器学习工作负载。TPU有多个版本，每个版本具有不同的硬件规格。有关TPU的更多信息，请参阅TPU系统架构。如需了解vLLM支持的TPU版本信息，请参阅：

• TPU v6e
• TPU v5e
• TPU v5p
• TPU v4

这些TPU版本允许您配置TPU芯片的物理排列方式。这可以提高吞吐量和网络性能。如需了解更多信息，请参阅：

• TPU v6e 拓扑结构
• TPU v5e 拓扑结构
• TPU v5p 拓扑结构
• TPU v4 拓扑结构

要使用Cloud TPU，您需要获得Google Cloud Platform项目的TPU配额。TPU配额规定了您可以在GPC项目中使用的TPU数量，具体取决于TPU版本、您想使用的TPU数量以及配额类型。更多信息请参阅TPU配额。

有关TPU的定价信息，请参阅Cloud TPU定价。

您的TPU虚拟机可能需要额外的持久存储。更多信息，请参阅Cloud TPU数据存储选项。

本标签页提供在英特尔Gaudi设备上运行vLLM的指导说明。

从vLLM 0.3.3版本开始，支持通过Neuron SDK在AWS Trainium/Inferentia上进行模型推理和服务，并具备持续批处理功能。 分页注意力机制(Paged Attention)和分块预填充(Chunked Prefill)功能目前正在开发中，即将推出。 当前Neuron SDK支持的数据类型为FP16和BF16。

• Google Cloud TPU 虚拟机
• TPU版本：v6e、v5e、v5p、v4
• Python: 3.10 或更新版本

配置云TPU¶

您可以通过Cloud TPU API或队列资源API（推荐方式）来配置Cloud TPU。本节展示如何使用队列资源API创建TPU。如需了解使用Cloud TPU API的更多信息，请参阅使用Create Node API创建Cloud TPU。队列资源允许您以排队方式请求Cloud TPU资源。当您请求队列资源时，该请求会被添加到由Cloud TPU服务维护的队列中。当请求的资源可用时，它将被分配给您的Google Cloud项目供您立即独占使用。

使用GKE配置Cloud TPU¶

有关在GKE中使用TPU的更多信息，请参阅： - https://cloud.google.com/kubernetes-engine/docs/how-to/tpus - https://cloud.google.com/kubernetes-engine/docs/concepts/tpus - https://cloud.google.com/kubernetes-engine/docs/concepts/plan-tpus

配置新环境¶

使用队列资源API配置Cloud TPU¶

创建一个包含4个TPU芯片的TPU v5e：

gcloud alpha compute tpus queued-resources create QUEUED_RESOURCE_ID \
--node-id TPU_NAME \
--project PROJECT_ID \
--zone ZONE \
--accelerator-type ACCELERATOR_TYPE \
--runtime-version RUNTIME_VERSION \
--service-account SERVICE_ACCOUNT

通过SSH连接到您的TPU：

gcloud compute tpus tpu-vm ssh TPU_NAME --zone ZONE

• 操作系统：Ubuntu 22.04 LTS
• Python: 3.10
• 英特尔Gaudi加速器
• 英特尔Gaudi软件版本1.18.0

请按照Gaudi安装指南中的说明设置执行环境。为获得最佳性能，请遵循优化训练平台指南中概述的方法。

配置新环境¶

环境验证¶

要验证Intel Gaudi软件是否正确安装，请运行：

hl-smi # verify that hl-smi is in your PATH and each Gaudi accelerator is visible
apt list --installed | grep habana # verify that habanalabs-firmware-tools, habanalabs-graph, habanalabs-rdma-core, habanalabs-thunk and habanalabs-container-runtime are installed
pip list | grep habana # verify that habana-torch-plugin, habana-torch-dataloader, habana-pyhlml and habana-media-loader are installed
pip list | grep neural # verify that neural_compressor is installed

更多详情请参阅Intel Gaudi Software Stack Verification

运行Docker镜像¶

强烈建议使用Intel Gaudi仓库中的最新Docker镜像。更多详情请参阅Intel Gaudi文档。

使用以下命令运行Docker镜像：

docker pull vault.habana.ai/gaudi-docker/1.18.0/ubuntu22.04/habanalabs/pytorch-installer-2.4.0:latest
docker run \
  -it \
  --runtime=habana \
  -e HABANA_VISIBLE_DEVICES=all \
  -e OMPI_MCA_btl_vader_single_copy_mechanism=none \
  --cap-add=sys_nice \
  --net=host \
  --ipc=host \
  vault.habana.ai/gaudi-docker/1.18.0/ubuntu22.04/habanalabs/pytorch-installer-2.4.0:latest

• 操作系统: Linux
• Python: 3.9 -- 3.11
• 加速器：NeuronCore_v2（在 trn1/inf2 实例中）
• PyTorch 2.0.1/2.1.1
• AWS Neuron SDK 2.16/2.17 (已在python 3.8上验证)

配置新环境¶

启动Trn1/Inf2实例¶

以下是启动trn1/inf2实例的步骤，以便安装PyTorch Neuron ("torch-neuronx")在Ubuntu 22.04 LTS上的配置。

• 请按照启动Amazon EC2实例的说明来启动实例。在EC2控制台选择实例类型时，请确保选择正确的实例类型。
• 要获取有关实例规格和定价的更多信息，请参阅：Trn1 web page, Inf2 web page
• 选择Ubuntu Server 22.04 TLS AMI
• 启动Trn1/Inf2实例时，请将主EBS卷大小调整为至少512GB。
• 启动实例后，按照连接到您的实例中的说明连接到实例

安装驱动程序和工具¶

如果已安装Deep Learning AMI Neuron，则无需安装驱动程序和工具。若操作系统中未安装驱动程序和工具，请按照以下步骤操作：

# Configure Linux for Neuron repository updates
. /etc/os-release
sudo tee /etc/apt/sources.list.d/neuron.list > /dev/null <<EOF
deb https://apt.repos.neuron.amazonaws.com ${VERSION_CODENAME} main
EOF
wget -qO - https://apt.repos.neuron.amazonaws.com/GPG-PUB-KEY-AMAZON-AWS-NEURON.PUB \
    | sudo apt-key add -

# Update OS packages
sudo apt-get update -y

# Install OS headers
sudo apt-get install linux-headers-$(uname -r) -y

# Install git
sudo apt-get install git -y

# install Neuron Driver
sudo apt-get install aws-neuronx-dkms=2.* -y

# Install Neuron Runtime
sudo apt-get install aws-neuronx-collectives=2.* -y
sudo apt-get install aws-neuronx-runtime-lib=2.* -y

# Install Neuron Tools
sudo apt-get install aws-neuronx-tools=2.* -y

# Add PATH
export PATH=/opt/aws/neuron/bin:$PATH

目前还没有预构建的TPU轮子。

目前，没有预构建的Intel Gaudi轮包。

目前还没有预构建的Neuron轮子。

安装Miniconda：

wget https://repo.anaconda.com/miniconda/Miniconda3-latest-Linux-x86_64.sh
bash Miniconda3-latest-Linux-x86_64.sh
source ~/.bashrc

为vLLM创建并激活Conda环境：

conda create -n vllm python=3.10 -y
conda activate vllm

克隆vLLM仓库并进入vLLM目录：

git clone https://github.com/vllm-project/vllm.git && cd vllm

卸载现有的 torch 和 torch_xla 包：

pip uninstall torch torch-xla -y

安装构建依赖项：

pip install -r requirements/tpu.txt
sudo apt-get install libopenblas-base libopenmpi-dev libomp-dev

运行安装脚本：

VLLM_TARGET_DEVICE="tpu" python -m pip install -e .

要从源代码构建并安装vLLM，请运行：

git clone https://github.com/vllm-project/vllm.git
cd vllm
pip install -r requirements/hpu.txt
python setup.py develop

目前，最新的功能和性能优化正在Gaudi的vLLM-fork中开发，我们会定期将其上游同步至vLLM主仓库。要安装最新的HabanaAI/vLLM-fork，请运行以下命令：

git clone https://github.com/HabanaAI/vllm-fork.git
cd vllm-fork
git checkout habana_main
pip install -r requirements/hpu.txt
python setup.py develop

以下说明适用于Neuron SDK 2.16及更高版本。

安装transformers-neuronx及其依赖项¶

transformers-neuronx 将作为后端支持在trn1/inf2实例上进行推理。按照以下步骤安装transformer-neuronx软件包及其依赖项。

# Install Python venv
sudo apt-get install -y python3.10-venv g++

# Create Python venv
python3.10 -m venv aws_neuron_venv_pytorch

# Activate Python venv
source aws_neuron_venv_pytorch/bin/activate

# Install Jupyter notebook kernel
pip install ipykernel
python3.10 -m ipykernel install \
    --user \
    --name aws_neuron_venv_pytorch \
    --display-name "Python (torch-neuronx)"
pip install jupyter notebook
pip install environment_kernels

# Set pip repository pointing to the Neuron repository
python -m pip config set \
    global.extra-index-url \
    https://pip.repos.neuron.amazonaws.com

# Install wget, awscli
python -m pip install wget
python -m pip install awscli

# Update Neuron Compiler and Framework
python -m pip install --upgrade neuronx-cc==2.* --pre torch-neuronx==2.1.* torchvision transformers-neuronx

从源码安装vLLM¶

安装完neuronx-cc和transformers-neuronx软件包后，我们将能够按以下方式安装vllm：

git clone https://github.com/vllm-project/vllm.git
cd vllm
pip install -U -r requirements/neuron.txt
VLLM_TARGET_DEVICE="neuron" pip install .

如果在安装过程中正确检测到neuron软件包，将安装vllm-0.3.0+neuron212。

有关使用官方Docker镜像的说明，请参阅deployment-docker-pre-built-image，确保将镜像名称vllm/vllm-openai替换为vllm/vllm-tpu。

目前还没有预构建的Intel Gaudi镜像。

目前没有预构建的Neuron镜像。

你可以使用来构建一个支持TPU的Docker镜像。

docker build -f docker/Dockerfile.tpu -t vllm-tpu .

使用以下命令运行Docker镜像：

# Make sure to add `--privileged --net host --shm-size=16G`.
docker run --privileged --net host --shm-size=16G -it vllm-tpu

from torch._C import *  # noqa: F403
ImportError: libopenblas.so.0: cannot open shared object file: No such
file or directory

sudo apt-get install libopenblas-base libopenmpi-dev libomp-dev

有关构建Docker镜像的说明，请参阅deployment-docker-build-image-from-source。

请确保使用替代默认的Dockerfile。

该设备没有额外信息。

支持的功能¶

• 离线推理
• 通过OpenAI兼容服务器实现在线服务
• HPU自动检测 - 无需在vLLM中手动选择设备
• 支持为英特尔Gaudi加速器启用算法的分页KV缓存
• 针对Intel Gaudi平台定制的分页注意力机制(Paged Attention)实现、KV缓存操作、预填充注意力机制、均方根层归一化(RMS Layer Normalization)以及旋转位置编码(Rotary Positional Encoding)
• 支持多卡推理的张量并行
• 使用HPU Graphs进行推理，以加速低批次延迟和吞吐量
• 带线性偏置的注意力机制 (ALiBi)

不支持的功能¶

• 束搜索
• LoRA适配器
• 量化
• 预填充分块（混合批次推理）

支持的配置¶

以下配置已经验证可在Gaudi2设备上正常运行。未列出的配置可能兼容也可能不兼容。

• meta-llama/Llama-2-7b 在单颗HPU上运行，或通过张量并行在2x和8x HPU上运行，使用BF16数据类型，采用随机或贪婪采样
• meta-llama/Llama-2-7b-chat-hf 在单颗HPU上运行，或通过张量并行在2颗和8颗HPU上运行，使用BF16数据类型，采用随机或贪婪采样
• meta-llama/Meta-Llama-3-8B 在单颗HPU上运行，或通过张量并行在2x和8x HPU上运行，使用BF16数据类型，采用随机或贪婪采样
• meta-llama/Meta-Llama-3-8B-Instruct 在单颗HPU上运行，或通过张量并行在2颗和8颗HPU上运行，使用BF16数据类型，采用随机或贪婪采样
• meta-llama/Meta-Llama-3.1-8B 在单颗HPU上运行，或通过2x和8x HPU张量并行，使用BF16数据类型，采用随机或贪婪采样
• meta-llama/Meta-Llama-3.1-8B-Instruct 在单颗HPU上运行，或通过2x和8x HPU张量并行，使用BF16数据类型，采用随机或贪婪采样
• meta-llama/Llama-2-70b 在8个HPU上使用张量并行，BF16数据类型，采用随机或贪婪采样
• meta-llama/Llama-2-70b-chat-hf 在8个HPU上使用张量并行，BF16数据类型，采用随机或贪婪采样
• meta-llama/Meta-Llama-3-70B 在8个HPU上使用张量并行，采用BF16数据类型，支持随机或贪婪采样
• meta-llama/Meta-Llama-3-70B-Instruct 在8个HPU上使用张量并行，采用BF16数据类型，支持随机或贪婪采样
• meta-llama/Meta-Llama-3.1-70B 在8个HPU上使用张量并行，BF16数据类型，采用随机或贪婪采样
• meta-llama/Meta-Llama-3.1-70B-Instruct 在8个HPU上使用张量并行，采用BF16数据类型，支持随机或贪婪采样

性能调优¶

执行模式¶

目前在vLLM的HPU版本中，我们支持四种执行模式，具体取决于所选的HPU PyTorch Bridge后端（通过PT_HPU_LAZY_MODE环境变量设置）以及--enforce-eager标志。

分桶机制¶

英特尔Gaudi加速器在处理具有固定张量形状的模型时表现最佳。Intel Gaudi Graph Compiler负责生成优化后的二进制代码，在Gaudi上实现给定的模型拓扑结构。在默认配置下，生成的二进制代码可能高度依赖于输入和输出张量的形状，当在同一拓扑结构中遇到不同形状的张量时，可能需要重新编译图结构。虽然生成的二进制代码能高效利用Gaudi，但编译过程本身可能会给端到端执行带来显著开销。 在动态推理服务场景中，需要尽量减少图编译的次数，并降低图编译在服务器运行时发生的风险。目前这是通过将模型的前向传播在batch_size和sequence_length两个维度上进行"分桶"来实现的。

分桶范围由3个参数决定 - min, step 和 max。这些参数可以分别为提示阶段和解码阶段设置，也可以为批处理大小和序列长度维度单独设置。这些参数可以在vLLM启动时的日志中观察到：

INFO 08-01 21:37:59 hpu_model_runner.py:493] Prompt bucket config (min, step, max_warmup) bs:[1, 32, 4], seq:[128, 128, 1024]
INFO 08-01 21:37:59 hpu_model_runner.py:499] Generated 24 prompt buckets: [(1, 128), (1, 256), (1, 384), (1, 512), (1, 640), (1, 768), (1, 896), (1, 1024), (2, 128), (2, 256), (2, 384), (2, 512), (2, 640), (2, 768), (2, 896), (2, 1024), (4, 128), (4, 256), (4, 384), (4, 512), (4, 640), (4, 768), (4, 896), (4, 1024)]
INFO 08-01 21:37:59 hpu_model_runner.py:504] Decode bucket config (min, step, max_warmup) bs:[1, 128, 4], seq:[128, 128, 2048]
INFO 08-01 21:37:59 hpu_model_runner.py:509] Generated 48 decode buckets: [(1, 128), (1, 256), (1, 384), (1, 512), (1, 640), (1, 768), (1, 896), (1, 1024), (1, 1152), (1, 1280), (1, 1408), (1, 1536), (1, 1664), (1, 1792), (1, 1920), (1, 2048), (2, 128), (2, 256), (2, 384), (2, 512), (2, 640), (2, 768), (2, 896), (2, 1024), (2, 1152), (2, 1280), (2, 1408), (2, 1536), (2, 1664), (2, 1792), (2, 1920), (2, 2048), (4, 128), (4, 256), (4, 384), (4, 512), (4, 640), (4, 768), (4, 896), (4, 1024), (4, 1152), (4, 1280), (4, 1408), (4, 1536), (4, 1664), (4, 1792), (4, 1920), (4, 2048)]

min 决定桶的最小值。step 决定桶之间的间隔，max 决定桶的上限。此外，min 和 step 之间的区间有特殊处理——min 会连续乘以2的幂次方，直到达到 step。我们称此为上升阶段，用于以最小浪费处理较小的批量大小，同时允许在较大批量大小上进行更大的填充。

示例（带预热）

min = 2, step = 32, max = 64
=> ramp_up = (2, 4, 8, 16)
=> stable = (32, 64)
=> buckets = ramp_up + stable => (2, 4, 8, 16, 32, 64)

示例（无预热阶段）

min = 128, step = 128, max = 512
=> ramp_up = ()
=> stable = (128, 256, 384, 512)
=> buckets = ramp_up + stable => (128, 256, 384, 512)

在记录的场景中，为提示（预填充）运行生成了24个桶，为解码运行生成了48个桶。每个桶对应一个针对特定模型及指定张量形状的独立优化设备二进制文件。每当处理一批请求时，系统会在批次和序列长度维度上将其填充至最小的可能桶。

例如，如果vLLM服务器空闲时收到一个包含3个序列、最大序列长度为412的请求，它将被填充为(4, 512)预填充桶执行，因为batch_size（序列数量）会被填充到4（最接近且大于3的批次大小维度），而最大序列长度会被填充到512（最接近且大于412的序列长度维度）。预填充阶段完成后，它将作为(4, 512)解码桶执行，并持续使用该桶直到发生以下情况之一：批次维度因请求完成而改变（此时会变为(2, 512)桶），或者上下文长度超过512个token（此时会变为(4, 640)桶）。

预热¶

预热是一个可选但强烈推荐的步骤，发生在vLLM服务器开始监听之前。它会使用虚拟数据为每个桶执行一次前向传播。其目的是预编译所有计算图，避免在服务器运行时因桶边界内的计算图编译而产生额外开销。vLLM启动时会记录每个预热步骤：

INFO 08-01 22:26:47 hpu_model_runner.py:1066] [Warmup][Prompt][1/24] batch_size:4 seq_len:1024 free_mem:79.16 GiB
INFO 08-01 22:26:47 hpu_model_runner.py:1066] [Warmup][Prompt][2/24] batch_size:4 seq_len:896 free_mem:55.43 GiB
INFO 08-01 22:26:48 hpu_model_runner.py:1066] [Warmup][Prompt][3/24] batch_size:4 seq_len:768 free_mem:55.43 GiB
...
INFO 08-01 22:26:59 hpu_model_runner.py:1066] [Warmup][Prompt][24/24] batch_size:1 seq_len:128 free_mem:55.43 GiB
INFO 08-01 22:27:00 hpu_model_runner.py:1066] [Warmup][Decode][1/48] batch_size:4 seq_len:2048 free_mem:55.43 GiB
INFO 08-01 22:27:00 hpu_model_runner.py:1066] [Warmup][Decode][2/48] batch_size:4 seq_len:1920 free_mem:55.43 GiB
INFO 08-01 22:27:01 hpu_model_runner.py:1066] [Warmup][Decode][3/48] batch_size:4 seq_len:1792 free_mem:55.43 GiB
...
INFO 08-01 22:27:16 hpu_model_runner.py:1066] [Warmup][Decode][47/48] batch_size:2 seq_len:128 free_mem:55.43 GiB
INFO 08-01 22:27:16 hpu_model_runner.py:1066] [Warmup][Decode][48/48] batch_size:1 seq_len:128 free_mem:55.43 GiB

本示例使用了与分桶机制部分相同的桶结构。每行输出对应单个桶的执行过程。当某个桶首次执行时，其计算图会被编译，后续可重复使用，从而跳过额外的图编译步骤。

HPU图捕获¶

HPU Graphs 目前是vLLM在英特尔Gaudi平台上性能最优的执行方式。启用HPU Graphs后，执行图将预先进行追踪记录（在完成预热后），随后在推理阶段重放这些记录，从而显著降低主机开销。记录过程可能占用大量内存，在分配KV缓存时需要将此因素纳入考量。启用HPU Graphs会影响可用KV缓存块的数量，但vLLM提供了用户可配置的变量来控制内存管理。

当使用HPU Graphs时，它们与KV缓存共享公共内存池（"可用内存"），该内存池大小由gpu_memory_utilization参数决定（默认值为0.9）。 在KV缓存分配之前，模型权重会被加载到设备上，并在虚拟数据上执行模型的前向传播以估算内存使用量。 之后才会应用gpu_memory_utilization参数——默认情况下，会将此时设备空闲内存的90%标记为可用内存。 接着分配KV缓存，预热模型，并捕获HPU Graphs。 环境变量VLLM_GRAPH_RESERVED_MEM定义了为HPU Graphs捕获保留的内存比例。 默认值（VLLM_GRAPH_RESERVED_MEM=0.1）下，10%的可用内存将保留用于graph捕获（后称"可用graph内存"），剩余90%用于KV缓存。 环境变量VLLM_GRAPH_PROMPT_RATIO决定了预填充和解码graph之间可用graph内存的分配比例。默认情况下（VLLM_GRAPH_PROMPT_RATIO=0.3），两个阶段具有相同的内存限制。 较低的值意味着预填充阶段分配的可用graph内存较少，例如VLLM_GRAPH_PROMPT_RATIO=0.2将为预填充graph保留20%的可用graph内存，为解码graph保留80%的可用graph内存。

用户还可以分别配置用于提示和解码阶段的HPU Graphs捕获策略。该策略会影响捕获图的顺序。目前已实现两种策略：

• max_bs - 图捕获队列将按批次大小降序排序。批次大小相同的桶按序列长度升序排序（例如 (64, 128), (64, 256), (32, 128), (32, 256), (1, 128), (1,256)），解码的默认策略
• min_tokens - 图捕获队列将按照每个图处理的token数量(batch_size*sequence_length)进行升序排序，这是提示的默认策略

当有大量请求处于待处理状态时，vLLM调度器会尽快尝试将解码批次大小填充至最大值。当某个请求完成时，解码批次大小会减小。此时，vLLM将尝试为等待队列中的请求调度预填充迭代，以使解码批次大小恢复到之前的状态。这意味着在满载场景下，解码批次大小通常保持最大值，因此捕获大批次HPU图至关重要，这体现在max_bs策略中。另一方面，预填充操作最常以极小的批次大小（1-4）执行，这反映在min_tokens策略中。

vLLM服务器会记录每个描述步骤的内存变化如下（负值表示内存被释放）：

INFO 08-02 17:37:44 hpu_model_runner.py:493] Prompt bucket config (min, step, max_warmup) bs:[1, 32, 4], seq:[128, 128, 1024]
INFO 08-02 17:37:44 hpu_model_runner.py:499] Generated 24 prompt buckets: [(1, 128), (1, 256), (1, 384), (1, 512), (1, 640), (1, 768), (1, 896), (1, 1024), (2, 128), (2, 256), (2, 384), (2, 512), (2, 640), (2, 768), (2, 896), (2, 1024), (4, 128), (4, 256), (4, 384), (4, 512), (4, 640), (4, 768), (4, 896), (4, 1024)]
INFO 08-02 17:37:44 hpu_model_runner.py:504] Decode bucket config (min, step, max_warmup) bs:[1, 128, 4], seq:[128, 128, 2048]
INFO 08-02 17:37:44 hpu_model_runner.py:509] Generated 48 decode buckets: [(1, 128), (1, 256), (1, 384), (1, 512), (1, 640), (1, 768), (1, 896), (1, 1024), (1, 1152), (1, 1280), (1, 1408), (1, 1536), (1, 1664), (1, 1792), (1, 1920), (1, 2048), (2, 128), (2, 256), (2, 384), (2, 512), (2, 640), (2, 768), (2, 896), (2, 1024), (2, 1152), (2, 1280), (2, 1408), (2, 1536), (2, 1664), (2, 1792), (2, 1920), (2, 2048), (4, 128), (4, 256), (4, 384), (4, 512), (4, 640), (4, 768), (4, 896), (4, 1024), (4, 1152), (4, 1280), (4, 1408), (4, 1536), (4, 1664), (4, 1792), (4, 1920), (4, 2048)]
INFO 08-02 17:37:52 hpu_model_runner.py:430] Pre-loading model weights on hpu:0 took 14.97 GiB of device memory (14.97 GiB/94.62 GiB used) and 2.95 GiB of host memory (475.2 GiB/1007 GiB used)
INFO 08-02 17:37:52 hpu_model_runner.py:438] Wrapping in HPU Graph took 0 B of device memory (14.97 GiB/94.62 GiB used) and -252 KiB of host memory (475.2 GiB/1007 GiB used)
INFO 08-02 17:37:52 hpu_model_runner.py:442] Loading model weights took in total 14.97 GiB of device memory (14.97 GiB/94.62 GiB used) and 2.95 GiB of host memory (475.2 GiB/1007 GiB used)
INFO 08-02 17:37:54 hpu_worker.py:134] Model profiling run took 504 MiB of device memory (15.46 GiB/94.62 GiB used) and 180.9 MiB of host memory (475.4 GiB/1007 GiB used)
INFO 08-02 17:37:54 hpu_worker.py:158] Free device memory: 79.16 GiB, 39.58 GiB usable (gpu_memory_utilization=0.5), 15.83 GiB reserved for HPUGraphs (VLLM_GRAPH_RESERVED_MEM=0.4), 23.75 GiB reserved for KV cache
INFO 08-02 17:37:54 hpu_executor.py:85] # HPU blocks: 1519, # CPU blocks: 0
INFO 08-02 17:37:54 hpu_worker.py:190] Initializing cache engine took 23.73 GiB of device memory (39.2 GiB/94.62 GiB used) and -1.238 MiB of host memory (475.4 GiB/1007 GiB used)
INFO 08-02 17:37:54 hpu_model_runner.py:1066] [Warmup][Prompt][1/24] batch_size:4 seq_len:1024 free_mem:55.43 GiB
...
INFO 08-02 17:38:22 hpu_model_runner.py:1066] [Warmup][Decode][48/48] batch_size:1 seq_len:128 free_mem:55.43 GiB
INFO 08-02 17:38:22 hpu_model_runner.py:1159] Using 15.85 GiB/55.43 GiB of free device memory for HPUGraphs, 7.923 GiB for prompt and 7.923 GiB for decode (VLLM_GRAPH_PROMPT_RATIO=0.3)
INFO 08-02 17:38:22 hpu_model_runner.py:1066] [Warmup][Graph/Prompt][1/24] batch_size:1 seq_len:128 free_mem:55.43 GiB
...
INFO 08-02 17:38:26 hpu_model_runner.py:1066] [Warmup][Graph/Prompt][11/24] batch_size:1 seq_len:896 free_mem:48.77 GiB
INFO 08-02 17:38:27 hpu_model_runner.py:1066] [Warmup][Graph/Decode][1/48] batch_size:4 seq_len:128 free_mem:47.51 GiB
...
INFO 08-02 17:38:41 hpu_model_runner.py:1066] [Warmup][Graph/Decode][48/48] batch_size:1 seq_len:2048 free_mem:47.35 GiB
INFO 08-02 17:38:41 hpu_model_runner.py:1066] [Warmup][Graph/Prompt][12/24] batch_size:4 seq_len:256 free_mem:47.35 GiB
INFO 08-02 17:38:42 hpu_model_runner.py:1066] [Warmup][Graph/Prompt][13/24] batch_size:2 seq_len:512 free_mem:45.91 GiB
INFO 08-02 17:38:42 hpu_model_runner.py:1066] [Warmup][Graph/Prompt][14/24] batch_size:1 seq_len:1024 free_mem:44.48 GiB
INFO 08-02 17:38:43 hpu_model_runner.py:1066] [Warmup][Graph/Prompt][15/24] batch_size:2 seq_len:640 free_mem:43.03 GiB
INFO 08-02 17:38:43 hpu_model_runner.py:1128] Graph/Prompt captured:15 (62.5%) used_mem:14.03 GiB buckets:[(1, 128), (1, 256), (1, 384), (1, 512), (1, 640), (1, 768), (1, 896), (1, 1024), (2, 128), (2, 256), (2, 384), (2, 512), (2, 640), (4, 128), (4, 256)]
INFO 08-02 17:38:43 hpu_model_runner.py:1128] Graph/Decode captured:48 (100.0%) used_mem:161.9 MiB buckets:[(1, 128), (1, 256), (1, 384), (1, 512), (1, 640), (1, 768), (1, 896), (1, 1024), (1, 1152), (1, 1280), (1, 1408), (1, 1536), (1, 1664), (1, 1792), (1, 1920), (1, 2048), (2, 128), (2, 256), (2, 384), (2, 512), (2, 640), (2, 768), (2, 896), (2, 1024), (2, 1152), (2, 1280), (2, 1408), (2, 1536), (2, 1664), (2, 1792), (2, 1920), (2, 2048), (4, 128), (4, 256), (4, 384), (4, 512), (4, 640), (4, 768), (4, 896), (4, 1024), (4, 1152), (4, 1280), (4, 1408), (4, 1536), (4, 1664), (4, 1792), (4, 1920), (4, 2048)]
INFO 08-02 17:38:43 hpu_model_runner.py:1206] Warmup finished in 49 secs, allocated 14.19 GiB of device memory
INFO 08-02 17:38:43 hpu_executor.py:91] init_cache_engine took 37.92 GiB of device memory (53.39 GiB/94.62 GiB used) and 57.86 MiB of host memory (475.4 GiB/1007 GiB used)

推荐的vLLM参数¶

• We recommend running inference on Gaudi 2 with block_size of 128 for BF16 data type. Using default values (16, 32) might lead to sub-optimal performance due to Matrix Multiplication Engine under-utilization (see Gaudi Architecture).
• For max throughput on Llama 7B, we recommend running with batch size of 128 or 256 and max context length of 2048 with HPU Graphs enabled. If you encounter out-of-memory issues, see troubleshooting section.

环境变量¶

诊断与分析调节选项：

• VLLM_PROFILER_ENABLED: 如果设为true，将启用高级性能分析器。生成的JSON跟踪文件可在perfetto.habana.ai查看。默认为false。
• VLLM_HPU_LOG_STEP_GRAPH_COMPILATION: 如果设为true，当vLLM引擎每一步发生图编译时会记录日志。强烈建议与PT_HPU_METRICS_GC_DETAILS=1配合使用。默认为false。
• VLLM_HPU_LOG_STEP_GRAPH_COMPILATION_ALL: 如果设为true，即使没有发生任何变化，也会记录每个vLLM引擎步骤的图编译过程。默认为false。
• VLLM_HPU_LOG_STEP_CPU_FALLBACKS: 如果设为true，当发生CPU回退时，会记录每个vLLM引擎步骤的CPU回退情况。默认为false。
• VLLM_HPU_LOG_STEP_CPU_FALLBACKS_ALL: 如果设为true，即使没有发生CPU回退，也会记录每个vLLM引擎步骤的CPU回退情况。默认为false。

性能调优参数：

• VLLM_SKIP_WARMUP: 如果设为true，将跳过预热阶段，默认为false

• VLLM_GRAPH_RESERVED_MEM: 为HPUGraph捕获保留的内存百分比，默认为0.1

• VLLM_GRAPH_PROMPT_RATIO: 为提示图保留的显存百分比，默认为0.3

• VLLM_GRAPH_PROMPT_STRATEGY: 决定提示图捕获顺序的策略，可选min_tokens或max_bs，默认为min_tokens

• VLLM_GRAPH_DECODE_STRATEGY: 决定解码图捕获顺序的策略，可选min_tokens或max_bs，默认为max_bs

• VLLM_{phase}_{dim}_BUCKET_{param} - 12个环境变量的集合，用于配置分桶机制的范围

• {phase} 可以是 PROMPT 或 DECODE

• {dim} 可以是 BS, SEQ 或 BLOCK

• {param} 可以是 MIN、STEP 或 MAX

• 默认值：

  • 提示词：
  • 最小批次大小 (VLLM_PROMPT_BS_BUCKET_MIN): 1
  • 批量大小步长 (VLLM_PROMPT_BS_BUCKET_STEP): min(max_num_seqs, 32)
  • 批次大小最大值 (VLLM_PROMPT_BS_BUCKET_MAX): min(max_num_seqs, 64)
  • 序列长度最小值 (VLLM_PROMPT_SEQ_BUCKET_MIN): block_size
  • 序列长度步长 (VLLM_PROMPT_SEQ_BUCKET_STEP): block_size
  • 序列长度最大值 (VLLM_PROMPT_SEQ_BUCKET_MAX): max_model_len
  • 解码：
  • 最小批处理大小 (VLLM_DECODE_BS_BUCKET_MIN): 1
  • 批处理大小步长 (VLLM_DECODE_BS_BUCKET_STEP): min(max_num_seqs, 32)
  • 最大批处理大小 (VLLM_DECODE_BS_BUCKET_MAX): max_num_seqs
  • 序列长度最小值 (VLLM_DECODE_BLOCK_BUCKET_MIN): block_size
  • 序列长度步长 (VLLM_DECODE_BLOCK_BUCKET_STEP): block_size
  • 序列长度最大值 (VLLM_DECODE_BLOCK_BUCKET_MAX): max(128, (max_num_seqs*max_model_len)/block_size)

此外，还有一些影响vLLM执行的HPU PyTorch Bridge环境变量：

• PT_HPU_LAZY_MODE: 如果设为0，将使用Gaudi的PyTorch Eager后端；如果设为1，将使用Gaudi的PyTorch Lazy后端。默认值为1。
• PT_HPU_ENABLE_LAZY_COLLECTIVES: 在使用HPU Graphs进行张量并行推理时必须设置为true

故障排除：调整HPU图¶

If you experience device out-of-memory issues or want to attempt inference at higher batch sizes, try tweaking HPU Graphs by following the below:

• 调整gpu_memory_utilization参数。这会减少KV缓存的分配，为捕获更大批次的图保留一些余量。默认情况下gpu_memory_utilization设置为0.9，它尝试在短暂性能分析后分配约90%的剩余HBM用于KV缓存。注意降低该值会减少可用的KV缓存块数量，从而降低你在给定时间内能处理的有效最大token数。
• 如果该方法效率不高，您可以完全禁用HPUGraph。禁用HPU Graphs后，您将以较低批次的延迟和吞吐量为代价，换取较高批次下可能更高的吞吐量。您可以通过向服务器添加--enforce-eager标志（用于在线服务），或向LLM构造函数传递enforce_eager=True参数（用于离线推理）来实现这一点。

该设备没有额外信息。