adapt(iluvatar): BI-V150 backend adaptation for vLLM 0.24.0

[physics31415926]

PR Category

Vendor Backend Adaptation

PR Type

Bug Fix / Compatibility

Description

Adapt vllm-plugin-FL for Iluvatar BI-V150 GPUs with vLLM 0.24.0. The Iluvatar corex triton compiler (3.2.x) has several incompatibilities with vllm's triton kernels that prevent inference from running. This PR adds three targeted runtime patches applied at module import time.

Environment

Item	Value
Hardware	Iluvatar BI-V150, 32 GB VRAM × 16 cards
Driver	corex 4.5.0
Container image	dev-community-acr-registry.cn-shanghai.cr.aliyuncs.com/dev-community/dev-community:vllm-py3.12-corex.4.5.0-ubuntu24.04
Python	3.12
torch	2.10.0 (Iluvatar custom build)
triton	3.2.0 (Iluvatar corex build, iluvatar backend only)
vLLM	0.24.0
FlagGems commit	b10f1d474
vllm-plugin-FL	this branch

Changes

vllm_fl/dispatch/backends/vendor/iluvatar/iluvatar.py

Function	Condition	Description
patch_triton_language_for_iluvatar()	triton < 3.3	Adds a tl.make_tensor_descriptor stub. vllm 0.24.0's triton_unified_attention.py references this API inside @triton.jit functions; triton's DependencyFinder walks all code paths at JIT compile time and raises AttributeError even when USE_TD=False.
patch_triton_perf_model_for_iluvatar()	iluvatar only	Replaces triton.ops.matmul_perf_model.get_clock_rate_in_khz() with a fixed value (1500 MHz for BI-V150). The default implementation calls nvsmi / pynvml which are unavailable on Iluvatar hardware, causing a crash during triton autotuning.
patch_triton_chained_or_for_iluvatar()	triton < 3.3	Rewrites chained boolean or expressions (A or B or C → (A or B) or C) in vllm's triton_attention_helpers.py. Iluvatar triton 3.2.x raises UnsupportedLanguageConstruct for chained operators inside @triton.jit functions. Patches the file in-place (idempotent, marker-guarded), clears __pycache__, and evicts the module from sys.modules.
get_device_uuid()	iluvatar only	Guards NVML call behind try/except and returns a synthetic UUID (NVML not available on Iluvatar).

All patch functions are invoked at module import time (module-level calls after function definitions), so both the main process and all TP Worker subprocesses receive the patches automatically.

vllm_fl/platform.py

Location	Change
pre_register_and_update()	Calls patch_triton_chained_or_for_iluvatar() explicitly from the main process before Worker subprocesses are spawned, ensuring the vllm source file is patched on disk before any Worker imports it.

Testing

Results

Model	TP	max_model_len	gpu_memory_utilization	enforce_eager	compilation_config	Result	Output
Qwen3.6-27B	8	4096	0.75	True	default	✅ PASS	' John. I am a 25-year-old'
Qwen3.6-35B-A3B	8	4096	0.75	True	default	✅ PASS	' John. I am a 25-year-old'
Qwen3.6-27B	8	4096	0.70	False	cudagraph_capture_sizes=[1,2,4,8,16,32]	✅ PASS	' John. I am a 25-year-old male. I have been experiencing a persistent cough for'
Qwen3.6-35B-A3B	8	4096	0.70	False	cudagraph_capture_sizes=[1,2,4,8,16,32]	✅ PASS	' John. I am a 25-year-old male. I have been experiencing a persistent cough for'

Related Issues

N/A

Checklist

• I have run the existing tests and they pass
• I have added tests for my changes (if applicable)
• I have updated the documentation (if applicable)