其他 AI 加速器

其他 AI 加速器#

vLLM 是一个 Python 库，支持以下 AI 加速器。选择您的 AI 加速器类型以查看特定于供应商的说明

张量处理单元 (TPU) 是 Google 定制开发的专用集成电路 (ASIC)，用于加速机器学习工作负载。TPU 有不同的版本，每个版本都有不同的硬件规格。有关 TPU 的更多信息，请参阅TPU 系统架构。有关 vLLM 支持的 TPU 版本的更多信息，请参阅

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

为了使用 Cloud TPU，您需要为您的 Google Cloud Platform 项目授予 TPU 配额。TPU 配额指定您可以在 GPC 项目中使用多少 TPU，并以 TPU 版本、您想要使用的 TPU 数量和配额类型来指定。有关更多信息，请参阅TPU 配额。

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

您的 TPU VM 可能需要额外的持久性存储。有关更多信息，请参阅Cloud TPU 数据的存储选项。

注意

此设备没有预构建的 wheels，因此您必须使用预构建的 Docker 镜像或从源代码构建 vLLM。

此选项卡提供有关使用 Intel Gaudi 设备运行 vLLM 的说明。

注意

此设备没有预构建的 wheels 或镜像，因此您必须从源代码构建 vLLM。

vLLM 0.3.3 及更高版本支持使用 Neuron SDK 在 AWS Trainium/Inferentia 上进行模型推理和服务，并具有连续批处理功能。分页注意力和分块预填充目前正在开发中，并将很快推出。Neuron SDK 当前支持的数据类型为 FP16 和 BF16。

注意

此设备没有预构建的 wheels 或镜像，因此您必须从源代码构建 vLLM。

要求#

  • Google Cloud TPU VM

  • TPU 版本：v6e、v5e、v5p、v4

  • Python：3.10 或更高版本

配置 Cloud 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

有关将 TPU 与 GKE 结合使用的更多信息，请参阅

  • 操作系统：Ubuntu 22.04 LTS

  • Python：3.10

  • Intel Gaudi 加速器

  • Intel Gaudi 软件版本 1.18.0

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

  • 操作系统：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 上验证）

配置新环境#

使用排队资源 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
参数描述#

参数名称

描述

QUEUED_RESOURCE_ID

排队资源请求的用户分配 ID。

TPU_NAME

TPU 的用户分配名称，在分配排队资源请求时创建。

PROJECT_ID

您的 Google Cloud 项目

ZONE

您想要在其中创建 Cloud TPU 的 GCP 区域。您使用的值取决于您使用的 TPU 版本。有关更多信息，请参阅 TPU 区域和可用区 <https://cloud.google.com/tpu/docs/regions-zones>_

ACCELERATOR_TYPE

您想要使用的 TPU 版本。指定 TPU 版本，例如 v5litepod-4 指定具有 4 个内核的 v5e TPU。有关更多信息，请参阅 TPU 版本 <https://cloud.devsite.corp.google.com/tpu/docs/system-architecture-tpu-vm#versions>_。

RUNTIME_VERSION

要使用的 TPU VM 运行时版本。有关更多信息，请参阅 TPU VM 镜像 <https://cloud.google.com/tpu/docs/runtimes>_。

SERVICE_ACCOUNT

您的服务帐户的电子邮件地址。您可以在 IAM Cloud Console 的服务帐户下找到它。例如：tpu-service-account@<your_project_ID>.iam.gserviceaccount.com

使用 SSH 连接到您的 TPU

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

环境验证

要验证 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 软件堆栈验证。

运行 Docker 镜像

强烈建议使用 Intel Gaudi vault 中的最新 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

启动 Trn1/Inf2 实例

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

  • 请按照启动 Amazon EC2 实例中的说明启动实例。在 EC2 控制台中选择实例类型时，请确保选择正确的实例类型。

  • 要获取有关实例大小和定价的更多信息，请参阅：Trn1 网页，Inf2 网页

  • 选择 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

使用 Python 设置#

预构建 wheels#

目前，没有预构建的 TPU wheels。

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

目前，没有预构建的 Neuron wheels。

从源代码构建 wheel#

安装 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 setup.py develop

要从源代码构建和安装 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 当前支持的 Pytorch 版本安装 triton 版本 2.1.0。这与 vllm >= 0.5.3 不兼容。您可能会看到错误 cannot import name 'default_dump_dir...。要解决此问题，请在安装 vLLM wheel 后运行 pip install --upgrade triton==3.0.0。

以下说明适用于 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 设置#

预构建镜像#

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

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

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

从源代码构建镜像#

您可以使用 docker/Dockerfile.tpu 构建具有 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

注意

由于 TPU 依赖于需要静态形状的 XLA，因此 vLLM 对可能的输入形状进行分桶，并为每个形状编译一个 XLA 图。首次运行时，编译时间可能需要 20~30 分钟。但是，由于 XLA 图缓存在磁盘中（默认情况下在 VLLM_XLA_CACHE_PATH 或 ~/.cache/vllm/xla_cache 中），因此编译时间随后会缩短至约 5 分钟。

提示

如果您遇到以下错误

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

使用以下命令安装 OpenBLAS

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

docker build -f docker/Dockerfile.hpu -t vllm-hpu-env  .
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 --rm vllm-hpu-env

提示

如果您观察到以下错误：docker: Error response from daemon: Unknown runtime specified habana.，请参阅Intel Gaudi 软件堆栈和驱动程序安装的“使用容器安装”部分。确保您已安装 habana-container-runtime 包，并且已注册 habana 容器运行时。

有关构建 Docker 镜像的说明，请参阅从源代码构建 vLLM 的 Docker 镜像。

确保使用 docker/Dockerfile.neuron 代替默认的 Dockerfile。

额外信息#

此设备没有额外信息。

支持的功能

  • 离线推理

  • 通过OpenAI 兼容服务器进行在线服务

  • HPU 自动检测 - 无需在 vLLM 中手动选择设备

  • 分页 KV 缓存，为 Intel Gaudi 加速器启用算法

  • Paged Attention、KV 缓存操作、预填充注意力、均方根层归一化、旋转位置编码的自定义 Intel Gaudi 实现

  • 多卡推理的张量并行支持

  • 使用HPU Graphs进行推理，以加速低批量延迟和吞吐量

  • 具有线性偏差的注意力 (ALiBi)

不支持的功能

  • 束搜索

  • LoRA 适配器

  • 量化

  • 预填充分块（混合批次推理）

支持的配置

以下配置已验证可在 Gaudi2 设备上运行。未列出的配置可能可以工作，也可能无法工作。

性能调优

执行模式

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

vLLM 执行模式#

PT_HPU_LAZY_MODE

enforce_eager

执行模式

0

0

torch.compile

0

1

PyTorch eager 模式

1

0

HPU Graphs

1

1

PyTorch lazy 模式

警告

在 1.18.0 中，所有使用 PT_HPU_LAZY_MODE=0 的模式都高度实验性，仅应用于验证功能正确性。它们的性能将在后续版本中得到改进。为了在 1.18.0 中获得最佳性能，请使用 HPU Graphs 或 PyTorch lazy 模式。

分桶机制

Intel Gaudi 加速器在处理具有固定张量形状的模型时效果最佳。Intel Gaudi Graph Compiler 负责生成优化的二进制代码，以在 Gaudi 上实现给定的模型拓扑。在其默认配置中，生成的二进制代码可能严重依赖于输入和输出张量形状，并且在同一拓扑中遇到不同形状的张量时可能需要重新编译图。虽然生成的二进制文件有效地利用了 Gaudi，但编译本身可能会在端到端执行中引入明显的开销。在动态推理服务场景中，需要最大限度地减少图编译的数量，并降低服务器运行时发生图编译的风险。目前，这通过在两个维度上“分桶”模型的正向传递来实现 - batch_size 和 sequence_length。

注意

分桶允许我们显着减少所需图的数量，但它不处理任何图编译和设备代码生成 - 这在预热和 HPUGraph 捕获阶段完成。

分桶范围由 3 个参数确定 - min、step 和 max。它们可以分别为 prompt 和 decode 阶段以及 batch size 和 sequence length 维度单独设置。这些参数可以在 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。我们称之为 ramp-up 阶段，它用于处理较低的批次大小，最大限度地减少浪费，同时允许在较大的批次大小上进行更大的填充。

示例（带 ramp-up）

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)

示例（不带 ramp-up）

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

在记录的场景中，为 prompt（预填充）运行生成了 24 个桶，为 decode 运行生成了 48 个桶。每个桶对应于具有指定张量形状的给定模型的单独优化设备二进制文件。每当处理一批请求时，都会在批次和序列长度维度上填充到尽可能小的桶。

警告

如果请求在任何维度上超过最大桶大小，则将在不填充的情况下处理该请求，并且其处理可能需要图编译，从而可能显着增加端到端延迟。桶的边界是用户可配置的，通过环境变量，并且可以增加桶的上限以避免这种情况。

例如，如果一个包含 3 个序列的请求，最大序列长度为 412，进入空闲的 vLLM 服务器，它将被填充并作为 (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

此示例使用与分桶机制部分相同的桶。每个输出行对应于单个桶的执行。当首次执行桶时，其图将被编译，并且可以在以后重用，从而跳过进一步的图编译。

提示

编译所有桶可能需要一些时间，可以使用 VLLM_SKIP_WARMUP=true 环境变量关闭。请记住，如果您这样做，则在首次执行给定桶时可能会遇到图编译。对于开发来说，禁用预热是可以的，但在部署中强烈建议启用它。

HPU 图捕获

HPU Graphs 目前是在 Intel Gaudi 上执行 vLLM 的最高效方法。启用 HPU Graphs 后，执行图将提前（在执行预热后）进行跟踪（记录），以便稍后在推理期间重放，从而显着减少主机开销。记录会占用大量内存，这需要在分配 KV 缓存时考虑在内。启用 HPU Graphs 将影响可用的 KV 缓存块的数量，但 vLLM 提供了用户可配置的变量来控制内存管理。

当使用 HPU 图时，它们会与 KV 缓存共享通用内存池（“可用内存”），这由 gpu_memory_utilization 标志（默认值为 0.9）确定。在分配 KV 缓存之前，模型权重会被加载到设备上，并对虚拟数据执行模型的前向传递，以估计内存使用量。 之后，才会使用 gpu_memory_utilization 标志 - 在其默认值下，会将此时空闲设备内存的 90% 标记为可用。接下来，分配 KV 缓存，预热模型，并捕获 HPU 图。环境变量 VLLM_GRAPH_RESERVED_MEM 定义了为 HPU 图捕获保留的内存比例。使用其默认值（VLLM_GRAPH_RESERVED_MEM=0.1），10% 的可用内存将保留用于图捕获（以下称为“可用图内存”），剩余的 90% 将用于 KV 缓存。环境变量 VLLM_GRAPH_PROMPT_RATIO 确定了为预填充和解码图保留的可用图内存比例。默认情况下（VLLM_GRAPH_PROMPT_RATIO=0.3），两个阶段都具有相等的内存约束。较低的值对应于为预填充阶段保留的较少可用图内存，例如 VLLM_GRAPH_PROMPT_RATIO=0.2 将为预填充图保留 20% 的可用图内存，为解码图保留 80% 的可用图内存。

注意

gpu_memory_utilization 并不对应于整个 HPU 的绝对内存使用量。它指定了在加载模型并执行性能分析运行后剩余的内存余量。如果设备总内存为 100 GiB，并且在加载模型权重并执行性能分析运行后剩余 50 GiB 的空闲内存，则默认值下的 gpu_memory_utilization 将把 50 GiB 中的 90% 标记为可用，留下 5 GiB 的余量，而与设备总内存无关。

用户还可以分别配置用于 prompt 和 decode 阶段的 HPU 图捕获策略。策略会影响捕获图的顺序。目前实现了两种策略：- max_bs - 图捕获队列将按批量大小降序排序。批量大小相同的 bucket 按序列长度升序排序（例如，(64, 128)、(64, 256)、(32, 128)、(32, 256)、(1, 128)、(1,256)），解码的默认策略 - min_tokens - 图捕获队列将按每个图处理的 token 数量（batch_size*sequence_length）升序排序，prompt 的默认策略

当有大量请求等待处理时，vLLM 调度器将尝试尽快填充解码的最大批量大小。当一个请求完成时，解码批量大小会减小。当这种情况发生时，vLLM 将尝试为等待队列中的请求调度一个预填充迭代，以将解码批量大小填充到之前的状态。这意味着在满负载情况下，解码批量大小通常处于最大值，这使得捕获大批量大小的 HPU 图至关重要，正如 max_bs 策略所反映的那样。另一方面，预填充将最频繁地以非常小的批量大小（1-4）执行，这反映在 min_tokens 策略中。

注意

VLLM_GRAPH_PROMPT_RATIO 并未对每个阶段（预填充和解码）的图所占用的内存设置硬性限制。vLLM 将首先尝试用尽所有可用的预填充图内存（可用图内存 * VLLM_GRAPH_PROMPT_RATIO）来捕获预填充 HPU 图，接下来它将尝试对解码图和可用的解码图内存池执行相同的操作。如果一个阶段已完全捕获，并且可用图内存池中还有未使用的内存，则 vLLM 将尝试为另一个阶段进一步捕获图，直到无法在不超出保留内存池的情况下捕获更多 HPU 图。该机制的行为可以在下面的示例中观察到。

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)

  • 我们建议在 Gaudi 2 上使用 128 的 block_size 和 BF16 数据类型运行推理。由于矩阵乘法引擎利用率不足，使用默认值 (16, 32) 可能会导致次优性能（请参阅 Gaudi 架构）。

  • 为了在 Llama 7B 上获得最大吞吐量，我们建议在启用 HPU 图的情况下，以 128 或 256 的批量大小和 2048 的最大上下文长度运行。如果您遇到内存不足问题，请参阅故障排除部分。

环境变量

诊断和性能分析旋钮

  • VLLM_PROFILER_ENABLED: 如果为 true，则将启用高级性能分析器。结果 JSON 跟踪可以在 perfetto.habana.ai 中查看。默认禁用。

  • VLLM_HPU_LOG_STEP_GRAPH_COMPILATION: 如果为 true，将记录每个 vLLM 引擎步骤的图编译，仅当存在任何编译时才记录 - 强烈建议与 PT_HPU_METRICS_GC_DETAILS=1 一起使用。默认禁用。

  • VLLM_HPU_LOG_STEP_GRAPH_COMPILATION_ALL: 如果为 true，将记录每个 vLLM 引擎步骤的图编译，始终记录，即使没有编译。默认禁用。

  • VLLM_HPU_LOG_STEP_CPU_FALLBACKS: 如果为 true，将记录每个 vLLM 引擎步骤的 CPU 回退，仅当存在任何回退时才记录。默认禁用。

  • VLLM_HPU_LOG_STEP_CPU_FALLBACKS_ALL: 如果为 true，将记录每个 vLLM 引擎步骤的 CPU 回退，始终记录，即使没有回退。默认禁用。

性能调优旋钮

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

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

  • VLLM_GRAPH_PROMPT_RATIO: 为 prompt 图预留的图内存百分比，默认为 0.3

  • VLLM_GRAPH_PROMPT_STRATEGY: 确定 prompt 图捕获顺序的策略，min_tokens 或 max_bs，默认为 min_tokens

  • VLLM_GRAPH_DECODE_STRATEGY: 确定 decode 图捕获顺序的策略，min_tokens 或 max_bs，默认为 max_bs

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

    • {phase} 是 PROMPT 或 DECODE

    • {dim} 是 BS、SEQ 或 BLOCK

    • {param} 是 MIN、STEP 或 MAX

    • 默认值

      • Prompt

        • 批量大小最小值 (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

      • Decode

        • 批量大小最小值 (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 图进行张量并行推理时，必须为 true

故障排除：调整 HPU 图

如果您遇到设备内存不足的问题，或者想尝试以更高的批量大小进行推理，请尝试通过以下方法调整 HPU 图

  • 调整 gpu_memory_utilization knob。这将减少 KV 缓存的分配，为捕获更大批量大小的图留出一些空间。默认情况下，gpu_memory_utilization 设置为 0.9。它尝试为 KV 缓存分配在短暂的性能分析运行后剩余的约 90% 的 HBM。请注意，减少此值会减少您可用的 KV 缓存块数量，从而减少您在给定时间内可以处理的有效最大 token 数量。

  • 如果此方法效率不高，您可以完全禁用 HPUGraph。禁用 HPU 图后，您将以较低批量时的延迟和吞吐量为代价，换取较高批量时可能更高的吞吐量。您可以通过向服务器添加 --enforce-eager 标志（用于在线服务），或者将 enforce_eager=True 参数传递给 LLM 构造函数（用于离线推理）来做到这一点。

此设备没有额外信息。