[ROCm][CI] Adding test groups for parity with upstream (#47479)

Signed-off-by: Andreas Karatzas <akaratza@amd.com>

Andreas Karatzas committed Jul 3, 2026 at 10:15 UTC 3775d5fcabf7bc5d4d92768485d860d132c6e1b6
2 files changed +189 -2
.buildkite/test-amd.yaml
+186 -1
@@ -1025,6 +1025,23 @@ steps:
1025 commands:
1026 - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-small.txt
1027
1028 +- label: MRCR Eval Small Models # TBD
1029 + timeout_in_minutes: 180
1030 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
1031 + agent_pool: mi300_1
1032 + optional: true
1033 + working_dir: "/vllm-workspace/tests"
1034 + source_file_dependencies:
1035 + - vllm/model_executor/models/
1036 + - vllm/model_executor/model_loader/
1037 + - vllm/v1/attention/backends/
1038 + - vllm/v1/attention/selector.py
1039 + - vllm/_aiter_ops.py
1040 + - vllm/platforms/rocm.py
1041 + - tests/evals/mrcr/
1042 + commands:
1043 + - pytest -s -v evals/mrcr/test_mrcr_correctness.py --config-list-file=evals/mrcr/configs/models-small.txt
1044 +
1045 - label: LM Eval Small Models (MI300) # TBD
1046 timeout_in_minutes: 180
1047 mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
@@ -1283,6 +1300,21 @@ steps:
1300
1301 #---------------------------------------------------------- mi300 · kernels ----------------------------------------------------------#
1302
1303 +- label: vLLM IR Tests # TBD
1304 + timeout_in_minutes: 30
1305 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
1306 + agent_pool: mi300_1
1307 + optional: true
1308 + working_dir: "/vllm-workspace/"
1309 + source_file_dependencies:
1310 + - vllm/ir
1311 + - vllm/kernels
1312 + - vllm/_aiter_ops.py
1313 + - vllm/platforms/rocm.py
1314 + commands:
1315 + - pytest -v -s tests/ir
1316 + - pytest -v -s tests/kernels/ir
1317 +
1318 - label: Kernels Attention Test %N # TBD
1319 timeout_in_minutes: 100
1320 mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
@@ -1317,6 +1349,21 @@ steps:
1349 commands:
1350 - pytest -v -s kernels/core --ignore=kernels/core/test_minimax_reduce_rms.py kernels/test_concat_mla_q.py kernels/test_top_k_per_row.py
1351
1352 +- label: Kernels KDA Test # TBD
1353 + timeout_in_minutes: 30
1354 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
1355 + agent_pool: mi300_1
1356 + optional: true
1357 + working_dir: "/vllm-workspace/tests"
1358 + source_file_dependencies:
1359 + - vllm/model_executor/layers/fla/ops/kda.py
1360 + - vllm/model_executor/layers/fla/ops/chunk_delta_h.py
1361 + - vllm/model_executor/layers/fla/ops/l2norm.py
1362 + - tests/kernels/test_kda.py
1363 + - vllm/platforms/rocm.py
1364 + commands:
1365 + - pytest -v -s kernels/test_kda.py
1366 +
1367 - label: Kernels MoE Test %N # TBD
1368 timeout_in_minutes: 95
1369 mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
@@ -1429,6 +1476,128 @@ steps:
1476 - pytest -v -s model_executor -m '(not slow_test)'
1477 - pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py
1478
1479 +#---------------------------------------------------- mi300 · model_runner_v2 -------------------------------------------------------#
1480 +
1481 +- label: Model Runner V2 Core Tests # TBD
1482 + timeout_in_minutes: 60
1483 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
1484 + agent_pool: mi300_1
1485 + optional: true
1486 + working_dir: "/vllm-workspace/tests"
1487 + source_file_dependencies:
1488 + - vllm/v1/worker/gpu/
1489 + - vllm/v1/worker/gpu_worker.py
1490 + - vllm/v1/core/sched/
1491 + - vllm/v1/attention/
1492 + - tests/v1/engine/test_llm_engine.py
1493 + - tests/v1/e2e/
1494 + - tests/entrypoints/llm/test_struct_output_generate.py
1495 + - vllm/platforms/rocm.py
1496 + commands:
1497 + - set -x
1498 + - export VLLM_USE_V2_MODEL_RUNNER=1
1499 + - pytest -v -s v1/engine/test_llm_engine.py -k "not test_engine_metrics"
1500 + - ENFORCE_EAGER=1 pytest -v -s v1/e2e/general/test_async_scheduling.py -k "not ngram"
1501 + - pytest -v -s v1/e2e/general/test_context_length.py
1502 + - pytest -v -s v1/e2e/general/test_min_tokens.py
1503 + - pytest -v -s entrypoints/llm/test_struct_output_generate.py -k "xgrammar and not speculative_config6 and not speculative_config7 and not speculative_config8 and not speculative_config0"
1504 +
1505 +- label: Model Runner V2 Examples # TBD
1506 + timeout_in_minutes: 60
1507 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
1508 + agent_pool: mi300_1
1509 + optional: true
1510 + working_dir: "/vllm-workspace/examples"
1511 + source_file_dependencies:
1512 + - vllm/v1/worker/gpu/
1513 + - vllm/v1/core/sched/
1514 + - vllm/v1/worker/gpu_worker.py
1515 + - examples/basic/offline_inference/
1516 + - examples/generate/multimodal/
1517 + - examples/features/
1518 + - examples/pooling/embed/vision_embedding_offline.py
1519 + - examples/features/tensorize_vllm_model.py
1520 + - examples/deployment/llm_engine_example.py
1521 + - vllm/platforms/rocm.py
1522 + commands:
1523 + - set -x
1524 + - export VLLM_USE_V2_MODEL_RUNNER=1
1525 + - pip install tensorizer
1526 + - python3 basic/offline_inference/chat.py
1527 + - python3 basic/offline_inference/generate.py --model facebook/opt-125m
1528 + - python3 generate/multimodal/audio_language_offline.py --seed 0
1529 + - python3 generate/multimodal/vision_language_offline.py --seed 0
1530 + - python3 generate/multimodal/vision_language_multi_image_offline.py --seed 0
1531 + - python3 generate/multimodal/encoder_decoder_multimodal_offline.py --model-type whisper --seed 0
1532 + - python3 pooling/embed/vision_embedding_offline.py --seed 0
1533 + - python3 features/automatic_prefix_caching/prefix_caching_offline.py
1534 + - python3 deployment/llm_engine_example.py
1535 + - python3 features/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 features/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
1536 + - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
1537 + - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
1538 +
1539 +- label: Model Runner V2 Distributed (2 GPUs) # TBD
1540 + timeout_in_minutes: 60
1541 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
1542 + agent_pool: mi300_2
1543 + num_gpus: 2
1544 + optional: true
1545 + working_dir: "/vllm-workspace/tests"
1546 + source_file_dependencies:
1547 + - vllm/v1/worker/gpu/
1548 + - vllm/v1/worker/gpu_worker.py
1549 + - tests/basic_correctness/test_basic_correctness.py
1550 + - tests/v1/distributed/test_async_llm_dp.py
1551 + - tests/v1/distributed/test_eagle_dp.py
1552 + - vllm/platforms/rocm.py
1553 + commands:
1554 + - set -x
1555 + - export VLLM_USE_V2_MODEL_RUNNER=1
1556 + - TARGET_TEST_SUITE=MI300 pytest -v -s basic_correctness/test_basic_correctness.py -m 'distributed(num_gpus=2)' -k "not ray and not True"
1557 + - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py -k "not ray"
1558 + - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
1559 +
1560 +- label: Model Runner V2 Pipeline Parallelism (4 GPUs) # TBD
1561 + timeout_in_minutes: 60
1562 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
1563 + agent_pool: mi300_4
1564 + num_gpus: 4
1565 + optional: true
1566 + working_dir: "/vllm-workspace/tests"
1567 + source_file_dependencies:
1568 + - vllm/v1/worker/gpu/
1569 + - vllm/v1/worker/gpu_worker.py
1570 + - tests/distributed/test_pipeline_parallel.py
1571 + - tests/distributed/test_pp_cudagraph.py
1572 + - vllm/platforms/rocm.py
1573 + commands:
1574 + - set -x
1575 + - export VLLM_USE_V2_MODEL_RUNNER=1
1576 + - pytest -v -s distributed/test_pipeline_parallel.py -k "not ray and not Jamba"
1577 + - pytest -v -s distributed/test_pp_cudagraph.py -k "not ray"
1578 +
1579 +- label: Model Runner V2 Spec Decode # TBD
1580 + timeout_in_minutes: 45
1581 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
1582 + agent_pool: mi300_1
1583 + optional: true
1584 + working_dir: "/vllm-workspace/tests"
1585 + source_file_dependencies:
1586 + - vllm/v1/worker/gpu/
1587 + - vllm/v1/worker/gpu_worker.py
1588 + - tests/v1/spec_decode/test_max_len.py
1589 + - tests/v1/spec_decode/test_rejection_sampler_utils.py
1590 + - tests/v1/spec_decode/test_synthetic_rejection_sampler_utils.py
1591 + - tests/v1/e2e/spec_decode/test_spec_decode.py
1592 + - vllm/platforms/rocm.py
1593 + commands:
1594 + - set -x
1595 + - export VLLM_USE_V2_MODEL_RUNNER=1
1596 + - pytest -v -s v1/spec_decode/test_max_len.py -k "eagle or mtp"
1597 + - pytest -v -s v1/spec_decode/test_rejection_sampler_utils.py
1598 + - pytest -v -s v1/spec_decode/test_synthetic_rejection_sampler_utils.py
1599 + - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "eagle or mtp"
1600 +
1601 #------------------------------------------------------ mi300 · models / basic -------------------------------------------------------#
1602
1603 - label: Basic Models Tests (Extra Initialization) %N # TBD
@@ -1745,6 +1914,22 @@ steps:
1914 - pytest -v -s plugins_tests/test_oot_registration_offline.py # it needs a clean process
1915 - pytest -v -s plugins_tests/lora_resolvers # unit tests for in-tree lora resolver plugins
1916
1917 +- label: GGUF Plugin # TBD
1918 + timeout_in_minutes: 30
1919 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
1920 + agent_pool: mi300_1
1921 + optional: true
1922 + soft_fail: true
1923 + working_dir: "/vllm-workspace/tests"
1924 + source_file_dependencies:
1925 + - vllm/model_executor/layers/quantization
1926 + - tests/plugins_tests/test_gguf_plugin.py
1927 + - tests/plugins_tests/gguf
1928 + - vllm/platforms/rocm.py
1929 + commands:
1930 + - pip install "vllm-gguf-plugin >= 0.0.2"
1931 + - pytest -v -s plugins_tests/gguf
1932 +
1933 #------------------------------------------------------- mi300 · quantization --------------------------------------------------------#
1934
1935 - label: Quantization # TBD
@@ -2770,7 +2955,7 @@ steps:
2955 #--------------------------------------------------------- mi355 · examples ----------------------------------------------------------#
2956
2957 - label: Examples # TBD
2773 - timeout_in_minutes: 45
2958 + timeout_in_minutes: 100
2959 mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
2960 agent_pool: mi355_1
2961 working_dir: "/vllm-workspace/examples"
vllm/model_executor/layers/fla/ops/chunk_delta_h.py
+3 -1
@@ -17,6 +17,8 @@ from .op import exp, exp2
17 from .utils import FLA_CHUNK_SIZE, use_cuda_graph
18
19 NUM_WARPS = [2, 4, 8, 16]
20 +# Triton's AMD backend fails to lower this kernel with num_stages=4.
21 +_CHUNK_DELTA_H_NUM_STAGES = [2, 3] if torch.version.hip else [2, 3, 4]
22
23
24 @triton.heuristics(
@@ -33,7 +35,7 @@ NUM_WARPS = [2, 4, 8, 16]
35 configs=[
36 triton.Config({"BV": BV}, num_warps=num_warps, num_stages=num_stages)
37 for num_warps in [2, 4]
36 - for num_stages in [2, 3, 4]
38 + for num_stages in _CHUNK_DELTA_H_NUM_STAGES
39 for BV in [32, 64]
40 ],
41 key=["H", "K", "V", "BT"],