其他 AI 加速器¶
vLLM 是一个支持以下 AI 加速器的 Python 库。选择您的 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。
AWS Neuron 是用于在 AWS Inferentia 和 AWS Trainium 驱动的 Amazon EC2 实例和 UltraServers（Inf1、Inf2、Trn1、Trn2 和 Trn2 UltraServer）上运行深度学习和生成式 AI 工作负载的软件开发工具包 (SDK)。Trainium 和 Inferentia 都由称为 NeuronCores 的完全独立的异构计算单元驱动。此选项卡描述了如何设置环境以在 Neuron 上运行 vLLM。
警告
此设备没有预构建的 wheels 或镜像，因此您必须从源代码构建 vLLM。
要求¶
- Google Cloud TPU VM
- TPU 版本：v6e, v5e, v5p, v4
- Python：3.10 或更新版本
配置 Cloud TPU¶
您可以使用 Cloud TPU API 或 排队资源（queued resources） 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
配置新环境¶
使用 queued resource 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 可用区。您使用的值 |
| ACCELERATOR_TYPE | 您想使用的 TPU 版本。指定 TPU 版本，例如 |
| RUNTIME_VERSION | 要使用的 TPU VM 运行时版本。例如，对于加载了一个或多个 v6e TPU 的 VM，请使用 v2-alpha-tpuv6e。更多信息请参阅 TPU VM 镜像。 |
使用 SSH 连接到您的 TPU
- 操作系统：Ubuntu 22.04 LTS
- Python：3.10
- Intel Gaudi 加速器
- Intel 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 软件栈验证。
运行 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 或更新版本
- Pytorch 2.5/2.6
- 加速器：NeuronCore-v2（在 trn1/inf2 芯片中）或 NeuronCore-v3（在 trn2 芯片中）
- AWS Neuron SDK 2.23
配置新环境¶
启动 Trn1/Trn2/Inf2 实例并验证 Neuron 依赖项¶
启动预装 Neuron 依赖项的 Trainium 或 Inferentia 实例最简单的方法是使用 Neuron 深度学习 AMI (Amazon machine image) 按照此 快速入门指南 进行操作。
- 启动实例后，按照 连接到您的实例 中的说明连接到实例
- 进入实例后，通过运行以下命令激活预装的用于推理的虚拟环境
有关使用 Docker 和手动安装依赖项等替代设置说明，请参阅 NxD Inference 设置指南。
注意
NxD Inference 是在 Neuron 上运行推理的默认推荐后端。如果您想使用旧版 transformers-neuronx 库，请参阅 Transformers NeuronX 设置。
配置新环境¶
使用 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 环境
克隆 vLLM 仓库并进入 vLLM 目录
卸载现有的 torch 和 torch_xla 包
安装构建依赖项
运行 setup 脚本
要从源代码构建和安装 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，请运行以下命令
从源代码安装 vLLM¶
按如下方式安装 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 -e .
AWS Neuron 在 https://github.com/aws-neuron/upstreaming-to-vllm/tree/neuron-2.23-vllm-v0.7.2 维护了一个 vLLM 的 Github 分支，其中包含 vLLM V0 中可用的额外一些特性。请利用 AWS 分支以获得以下特性
- Llama-3.2 多模态支持
- 多节点分布式推理
更多详细信息和使用示例请参阅 NxD Inference 的 vLLM 用户指南。
要安装 AWS Neuron 分支，请运行以下命令
git clone -b neuron-2.23-vllm-v0.7.2 https://github.com/aws-neuron/upstreaming-to-vllm.git
cd upstreaming-to-vllm
pip install -r requirements/neuron.txt
VLLM_TARGET_DEVICE="neuron" pip install -e .
请注意，AWS Neuron 分支仅用于支持 Neuron 硬件；与其他硬件的兼容性未经过测试。
使用 Docker 设置¶
预构建的镜像¶
有关使用官方 Docker 镜像的说明，请参阅 deployment-docker-pre-built-image，确保将镜像名称 vllm/vllm-openai 替换为 vllm/vllm-tpu。
目前没有预构建的 Intel Gaudi 镜像。
目前没有预构建的 Neuron 镜像。
从源代码构建镜像¶
您可以使用
使用以下命令运行 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 分钟。
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 镜像的说明，请参阅 deployment-docker-build-image-from-source。
请确保使用
额外信息¶
此设备没有额外信息。
支持的特性¶
- 离线推理
- 通过 OpenAI 兼容服务器 进行在线服务
- HPU 自动检测 - 无需在 vLLM 中手动选择设备
- 针对 Intel Gaudi 加速器启用算法的分页 KV 缓存
- Paged Attention、KV cache ops、prefill attention、Root Mean Square Layer Normalization、Rotary Positional Encoding 的定制 Intel Gaudi 实现
- 多卡推理的张量并行支持
- 使用 HPU 图 加速低批量延迟和吞吐量的推理
- 带线性偏置的注意力 (ALiBi)
不支持的特性¶
- Beam 搜索
- LoRA 适配器
- 量化
- 预填充分块（混合批量推理）
支持的配置¶
以下配置已验证可在 Gaudi2 设备上运行。未列出的配置可能工作，也可能不工作。
- meta-llama/Llama-2-7b 在单个 HPU 上，或在 2 倍和 8 倍 HPU 上进行张量并行，BF16 数据类型，使用随机或贪婪采样
- meta-llama/Llama-2-7b-chat-hf 在单个 HPU 上，或在 2 倍和 8 倍 HPU 上进行张量并行，BF16 数据类型，使用随机或贪婪采样
- meta-llama/Meta-Llama-3-8B 在单个 HPU 上，或在 2 倍和 8 倍 HPU 上进行张量并行，BF16 数据类型，使用随机或贪婪采样
- meta-llama/Meta-Llama-3-8B-Instruct 在单个 HPU 上，或在 2 倍和 8 倍 HPU 上进行张量并行，BF16 数据类型，使用随机或贪婪采样
- meta-llama/Meta-Llama-3.1-8B 在单个 HPU 上，或在 2 倍和 8 倍 HPU 上进行张量并行，BF16 数据类型，使用随机或贪婪采样
- meta-llama/Meta-Llama-3.1-8B-Instruct 在单个 HPU 上，或在 2 倍和 8 倍 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 for HPU 中，我们支持四种执行模式，取决于所选的 HPU PyTorch Bridge 后端（通过 PT_HPU_LAZY_MODE 环境变量）和 --enforce-eager 标志。
PT_HPU_LAZY_MODE |
enforce_eager |
执行模式 |
|---|---|---|
| 0 | 0 | torch.compile |
| 0 | 1 | PyTorch eager 模式 |
| 1 | 0 | HPU 图 |
警告
在 1.18.0 版本中，所有使用 PT_HPU_LAZY_MODE=0 的模式都处于高度实验阶段，仅应用于验证功能正确性。其性能将在后续版本中得到改进。为了在 1.18.0 中获得最佳性能，请使用 HPU 图或 PyTorch lazy 模式。
分桶机制¶
Intel Gaudi 加速器在处理具有固定张量形状的模型时效果最佳。Intel Gaudi 图编译器 负责生成优化的二进制代码，用于在 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 会连续乘以二的幂，直到达到 step。我们称之为爬升阶段（ramp-up phase），它用于以最小浪费处理较低的 batch size，同时允许在较大的 batch size 上进行较大的填充。
示例（带爬升）
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)
在日志场景中，为 prompt（预填充）运行生成了 24 个桶，为 decode 运行生成了 48 个桶。每个桶对应于给定模型具有指定张量形状的单独优化的设备二进制文件。每当处理一批请求时，它会在 batch 和 sequence length 维度上被填充到最小的可能桶大小。
警告
如果请求在任何维度上超出最大桶大小，它将不进行填充处理，并且其处理可能需要进行图编译，这可能会显著增加端到端延迟。桶的边界可以通过环境变量由用户配置，可以增加桶的上限以避免这种情况。
例如，如果一个包含 3 个序列、最大序列长度为 412 的请求进入空闲的 vLLM 服务器，它将被填充并作为 (4, 512) 的 prefill 桶执行，因为 batch_size（序列数量）将被填充到 4（大于 3 的最接近的 batch_size 维度），并且最大序列长度将被填充到 512（大于 412 的最接近的 sequence length 维度）。在 prefill 阶段之后，它将作为 (4, 512) 的 decode 桶执行，并一直保持在该桶中，直到 batch 维度发生变化（由于请求完成）——在这种情况下，它将变为 (2, 512) 的桶，或者上下文长度增加超过 512 个 token，在这种情况下，它将变为 (4, 640) 的桶。
注意
分桶对客户端是透明的——序列长度维度上的填充永远不会返回给客户端，并且 batch 维度上的填充不会创建新的请求。
预热¶
预热是 vLLM 服务器开始监听之前可选但强烈推荐的步骤。它使用 dummy 数据为每个桶执行一次正向传递。目标是预编译所有图，避免在服务器运行时在桶边界内产生任何图编译开销。每个预热步骤都会在 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 图 是目前 vLLM 在 Intel Gaudi 上性能最佳的执行方法。启用 HPU 图后，执行图将在预热后提前追踪（记录），以便稍后在推理期间重放，从而显著减少主机开销。记录可能需要大量内存，在分配 KV cache 时需要考虑这一点。启用 HPU 图会影响可用 KV cache block 的数量，但 vLLM 提供了用户可配置的变量来控制内存管理。
使用 HPU 图时，它们与 KV cache 共享共同的内存池（“可用内存”），该内存池由 gpu_memory_utilization 标志（默认为 0.9）决定。在分配 KV cache 之前，模型权重会加载到设备上，并在 dummy 数据上执行模型的正向传递以估算内存使用量。仅在此之后，才会使用 gpu_memory_utilization 标志——在其默认值下，它会将该点 90% 的空闲设备内存标记为可用。接下来，分配 KV cache，模型进行预热，并捕获 HPU 图。环境变量 VLLM_GRAPH_RESERVED_MEM 定义了为 HPU 图捕获保留的内存比例。在其默认值 (VLLM_GRAPH_RESERVED_MEM=0.1) 下，10% 的可用内存将保留用于图捕获（后文称为“可用图内存”），剩余 90% 将用于 KV cache。环境变量 VLLM_GRAPH_PROMPT_RATIO 决定了为 prefill 和 decode 图保留的可用图内存比例。默认情况下 (VLLM_GRAPH_PROMPT_RATIO=0.3)，两个阶段具有相等的内存限制。较低的值对应于为 prefill 阶段保留的可用图内存较少，例如 VLLM_GRAPH_PROMPT_RATIO=0.2 将为 prefill 图保留 20% 的可用图内存，为 decode 图保留 80% 的可用图内存。
注意
gpu_memory_utilization 不对应 HPU 的绝对内存使用量。它指定了加载模型并执行性能分析运行后的内存余量。如果设备总内存为 100 GiB，在加载模型权重和执行性能分析运行后有 50 GiB 可用内存，那么在其默认值下，gpu_memory_utilization 将标记 50 GiB 中的 90% 为可用，留下 5 GiB 的余量，而与设备总内存无关。
用户还可以分别为 prompt 和 decode 阶段配置捕获 HPU 图的策略。策略影响图的捕获顺序。目前实现了两种策略
max_bs- 图捕获队列将按其 batch size 降序排序。具有相同 batch size 的桶按 sequence length 升序排序（例如(64, 128),(64, 256),(32, 128),(32, 256),(1, 128),(1,256)），这是 decode 的默认策略min_tokens- 图捕获队列将按每个图处理的 token 数量 (batch_size*sequence_length) 升序排序，这是 prompt 的默认策略
当有大量请求待处理时，vLLM 调度程序会尽快尝试填满 decode 的最大 batch size。当请求完成后，decode batch size 会减少。发生这种情况时，vLLM 将尝试为等待队列中的请求安排 prefill 迭代，以将 decode batch size 填充回其之前的状态。这意味着在满负载场景下，decode batch size 通常处于最大值，这使得捕获大 batch size 的 HPU 图至关重要，这一点由 max_bs 策略体现。另一方面，prefill 将最频繁地以非常小的 batch size (1-4) 执行，这反映在 min_tokens 策略中。
注意
VLLM_GRAPH_PROMPT_RATIO 并未对每个阶段（prefill 和 decode）的图占用的内存设置硬性限制。vLLM 会首先尝试用尽全部可用的 prefill 图内存（可用图内存 * VLLM_GRAPH_PROMPT_RATIO）来捕获 prefill HPU 图，接下来会尝试对 decode 图和可用 decode 图内存池执行相同的操作。如果一个阶段已完全捕获，并且可用图内存池中仍有未使用的内存，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)
推荐的 vLLM 参数¶
- 我们建议在 Gaudi 2 上使用 BF16 数据类型时，将
block_size设置为 128 进行推理。使用默认值 (16, 32) 可能导致性能不佳，原因是矩阵乘法引擎利用率不足（参阅 Gaudi 架构）。 - 为了在 Llama 7B 上获得最大吞吐量，我们建议在使用 HPU Graphs 时，以 128 或 256 的 batch size 和 2048 的最大上下文长度运行。如果遇到内存不足问题，请参阅故障排除部分。
环境变量¶
诊断和性能分析旋钮
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，在每个 vLLM 引擎步骤中发生 CPU 回退时记录 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：用于 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}- 配置分桶机制范围的 12 个环境变量集合 -
{phase}是PROMPT或DECODE -
{dim}是BS、SEQ或BLOCK -
{param}是MIN、STEP或MAX -
默认值
- Prompt
- batch size 最小值 (
VLLM_PROMPT_BS_BUCKET_MIN)：1 - batch size 步长 (
VLLM_PROMPT_BS_BUCKET_STEP)：min(max_num_seqs, 32) - batch size 最大值 (
VLLM_PROMPT_BS_BUCKET_MAX)：min(max_num_seqs, 64) - sequence length 最小值 (
VLLM_PROMPT_SEQ_BUCKET_MIN)：block_size - sequence length 步长 (
VLLM_PROMPT_SEQ_BUCKET_STEP)：block_size - sequence length 最大值 (
VLLM_PROMPT_SEQ_BUCKET_MAX)：max_model_len - Decode
- batch size 最小值 (
VLLM_DECODE_BS_BUCKET_MIN)：1 - batch size 步长 (
VLLM_DECODE_BS_BUCKET_STEP)：min(max_num_seqs, 32) - batch size 最大值 (
VLLM_DECODE_BS_BUCKET_MAX)：max_num_seqs - sequence length 最小值 (
VLLM_DECODE_BLOCK_BUCKET_MIN)：block_size - sequence length 步长 (
VLLM_DECODE_BLOCK_BUCKET_STEP)：block_size - sequence length 最大值 (
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 图¶
如果您遇到设备内存不足问题或想尝试在更高的 batch size 下进行推理，请尝试通过以下方法调整 HPU 图
- 调整
gpu_memory_utilization旋钮。这将减少 KV cache 的分配，为捕获更大 batch size 的图留出一些空间。默认情况下，gpu_memory_utilization设置为 0.9。它尝试在短暂的性能分析运行后，将剩余约 90% 的 HBM 分配给 KV cache。请注意，降低此值会减少可用的 KV cache block 数量，从而降低您在给定时间可以处理的有效最大 token 数量。 - 如果此方法效率不高，您可以完全禁用
HPUGraph。禁用 HPU 图后，您牺牲了较低批量下的延迟和吞吐量，以换取较高批量下可能更高的吞吐量。您可以通过向服务器添加--enforce-eager标志（用于在线服务）或向 LLM 构造函数传递enforce_eager=True参数（用于离线推理）来实现此目的。
通过 NxD Inference 后端支持的特性¶
当前的 vLLM 和 Neuron 集成依赖于 neuronx-distributed-inference（推荐）或 transformers-neuronx 后端来执行大部分繁重工作，包括 PyTorch 模型初始化、编译和运行时执行。因此，大多数 Neuron 支持的特性 也通过 vLLM 集成提供。
要通过 vLLM 入口点配置 NxD Inference 特性，请使用 override_neuron_config 设置。将您要覆盖的配置作为字典（或从 CLI 启动 vLLM 时作为 JSON 对象）提供。例如，要禁用自动分桶，请包含
或者，用户可以直接调用 NxDI 库来追踪和编译模型，然后将预编译的 artifacts（通过 NEURON_COMPILED_ARTIFACTS 环境变量）加载到 vLLM 中运行推理工作负载。
已知限制¶
- EAGLE 推测解码：NxD Inference 要求 EAGLE 草稿 checkpoint 包含来自目标模型的 LM head 权重。请参阅此 指南，了解如何转换预训练的 EAGLE 模型 checkpoint 以使其与 NxDI 兼容。
- 量化：vLLM 中原生的量化流程在 NxD Inference 上支持不佳。建议遵循此 Neuron 量化指南，使用 NxD Inference 量化并编译模型，然后将编译后的 artifacts 加载到 vLLM 中。
- 多 LoRA 服务：NxD Inference 仅支持在服务器启动时加载 LoRA 适配器。目前不支持在运行时动态加载 LoRA 适配器。请参阅 multi-lora 示例
- 多模态支持：多模态支持仅通过 AWS Neuron 分支提供。此特性尚未上游到 vLLM 主仓库，因为 NxD Inference 目前依赖于对 vLLM 核心逻辑的某些适配来支持此特性。
- 多节点支持：跨多个 Trainium/Inferentia 实例的分布式推理仅在 AWS Neuron 分支上受支持。请参阅此 多节点示例 运行。请注意，张量并行（跨 NeuronCore 的分布式推理）在 vLLM 主仓库中可用。
- 推测解码中已知的边缘情况 bug：当序列长度接近最大模型长度时（例如，请求最大 token 数量达到最大模型长度并忽略 eos），推测解码中可能会发生边缘情况故障。在这种情况下，vLLM 可能尝试分配额外的 block 以确保有足够的内存用于 lookahead slot 数量，但由于我们对 paged attention 支持不佳，vLLM 没有另一个 Neuron block 可分配。AWS Neuron 分支中实现了一个变通修复（提前 1 次迭代终止），但由于它修改了 vLLM 核心逻辑，因此尚未上游到 vLLM 主仓库。
环境变量¶
NEURON_COMPILED_ARTIFACTS：设置此环境变量指向您预编译的模型 artifact 目录，以避免服务器初始化时的编译时间。如果未设置此变量，Neuron 模块将执行编译并将 artifacts 保存在模型路径下的neuron-compiled-artifacts/{unique_hash}/子目录中。如果设置了此环境变量，但目录不存在或内容无效，Neuron 也将回退到新的编译并将 artifacts 存储在此指定路径下。NEURON_CONTEXT_LENGTH_BUCKETS：用于上下文编码的桶大小。（仅适用于transformers-neuronx后端）。NEURON_TOKEN_GEN_BUCKETS：用于 token 生成的桶大小。（仅适用于transformers-neuronx后端）。