其他 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
後端)。