From 82c4823d5575c467dee97980b8faed4966d011cb Mon Sep 17 00:00:00 2001 From: Li Baoming <1508269885@qq.com> Date: Thu, 22 Oct 2026 16:33:51 +0800 Subject: [PATCH 1/9] fix: resolve CI build dependencies, Iluvatar includes and Ascend SOC detection --- .github/ci_config.yml | 14 +- src/CMakeLists.txt | 2 + .../ascend/custom/cmake/config_ascend.cmake | 2 +- .../ascend/custom/cmake/detect_soc.cmake | 61 +++-- tests/build_config/test_ci_build.py | 219 ++++++++++++++++++ 5 files changed, 276 insertions(+), 22 deletions(-) create mode 100644 tests/build_config/test_ci_build.py diff --git a/.github/ci_config.yml b/.github/ci_config.yml index df124563e..ec55f4ea9 100644 --- a/.github/ci_config.yml +++ b/.github/ci_config.yml @@ -15,7 +15,7 @@ platforms: BASE_IMAGE: nvcr.io/nvidia/pytorch:24.10-py3 SKIP_APT: "1" PIP_INDEX_URL: https://pypi.tuna.tsinghua.edu.cn/simple - setup: pip install scikit-build-core pybind11 libclang pyyaml && rm -rf /tmp/infinirt-src /tmp/infinirt-build /tmp/infinirt-prefix && git clone --depth 1 ${INFINI_RT_GIT_URL:-https://gh-proxy.com/https://github.com/InfiniTensor/InfiniRT.git} /tmp/infinirt-src && cmake -S /tmp/infinirt-src -B /tmp/infinirt-build -DWITH_CPU=ON -DWITH_NVIDIA=ON -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=/tmp/infinirt-prefix && cmake --build /tmp/infinirt-build -j$(nproc) && cmake --install /tmp/infinirt-build && pip install .[dev] --no-build-isolation --config-settings=cmake.define.INFINI_OPS_SMOKE_BUILD=ON --config-settings=cmake.define.INFINI_RT_ROOT=/tmp/infinirt-prefix + setup: python -m pip install 'scikit-build-core>=1.1.1' pybind11 libclang packaging pyyaml jinja2 && rm -rf /tmp/infinirt-src /tmp/infinirt-build /tmp/infinirt-prefix && git clone --depth 1 ${INFINI_RT_GIT_URL:-https://gh-proxy.com/https://github.com/InfiniTensor/InfiniRT.git} /tmp/infinirt-src && cmake -S /tmp/infinirt-src -B /tmp/infinirt-build -DWITH_CPU=ON -DWITH_NVIDIA=ON -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=/tmp/infinirt-prefix && cmake --build /tmp/infinirt-build -j$(nproc) && cmake --install /tmp/infinirt-build && python -m pip install .[dev] --no-build-isolation --config-settings=cmake.define.INFINI_OPS_SMOKE_BUILD=ON --config-settings=cmake.define.INFINI_RT_ROOT=/tmp/infinirt-prefix jobs: gpu: type: unittest @@ -51,7 +51,7 @@ platforms: - /lib/firmware:/lib/firmware - /usr/src:/usr/src - /lib/modules:/lib/modules - setup: python -m pip install scikit-build-core packaging exceptiongroup typing-extensions pygments pybind11 libclang pyyaml && rm -rf /tmp/infinirt-src /tmp/infinirt-build /tmp/infinirt-prefix && git clone --depth 1 ${INFINI_RT_GIT_URL:-https://gh-proxy.com/https://github.com/InfiniTensor/InfiniRT.git} /tmp/infinirt-src && cmake -S /tmp/infinirt-src -B /tmp/infinirt-build -DWITH_CPU=ON -DWITH_ILUVATAR=ON -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=/tmp/infinirt-prefix -DCMAKE_CUDA_ARCHITECTURES=OFF -DCMAKE_CUDA_COMPILER=/usr/local/corex-4.3.0.20250624/bin/clang++ -DCUDAToolkit_ROOT=/usr/local/corex-4.3.0.20250624 && cmake --build /tmp/infinirt-build -j$(nproc) && cmake --install /tmp/infinirt-build && python -m pip install . --no-build-isolation --no-deps --config-settings=cmake.define.INFINI_OPS_SMOKE_BUILD=ON --config-settings=cmake.define.INFINI_RT_ROOT=/tmp/infinirt-prefix + setup: python -m pip install 'scikit-build-core>=1.1.1' exceptiongroup typing-extensions pygments pybind11 libclang packaging pyyaml jinja2 && rm -rf /tmp/infinirt-src /tmp/infinirt-build /tmp/infinirt-prefix && git clone --depth 1 ${INFINI_RT_GIT_URL:-https://gh-proxy.com/https://github.com/InfiniTensor/InfiniRT.git} /tmp/infinirt-src && cmake -S /tmp/infinirt-src -B /tmp/infinirt-build -DWITH_CPU=ON -DWITH_ILUVATAR=ON -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=/tmp/infinirt-prefix -DCMAKE_CUDA_ARCHITECTURES=OFF -DCMAKE_CUDA_COMPILER=/usr/local/corex-4.3.0.20250624/bin/clang++ -DCUDAToolkit_ROOT=/usr/local/corex-4.3.0.20250624 && cmake --build /tmp/infinirt-build -j$(nproc) && cmake --install /tmp/infinirt-build && python -m pip install . --no-build-isolation --no-deps --config-settings=cmake.define.INFINI_OPS_SMOKE_BUILD=ON --config-settings=cmake.define.INFINI_RT_ROOT=/tmp/infinirt-prefix jobs: gpu: type: unittest @@ -81,7 +81,7 @@ platforms: - "--privileged" - "--ulimit=memlock=-1" - "--ulimit=stack=67108864" - setup: pip install scikit-build-core pybind11 libclang pyyaml && rm -rf /tmp/infinirt-src /tmp/infinirt-build /tmp/infinirt-prefix && git clone --depth 1 ${INFINI_RT_GIT_URL:-https://gh-proxy.com/https://github.com/InfiniTensor/InfiniRT.git} /tmp/infinirt-src && cmake -S /tmp/infinirt-src -B /tmp/infinirt-build -DWITH_CPU=ON -DWITH_METAX=ON -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=/tmp/infinirt-prefix && cmake --build /tmp/infinirt-build -j$(nproc) && cmake --install /tmp/infinirt-build && pip install .[dev] --no-build-isolation --config-settings=cmake.define.INFINI_OPS_SMOKE_BUILD=ON --config-settings=cmake.define.INFINI_RT_ROOT=/tmp/infinirt-prefix + setup: python -m pip install 'scikit-build-core>=1.1.1' pybind11 libclang packaging pyyaml jinja2 && rm -rf /tmp/infinirt-src /tmp/infinirt-build /tmp/infinirt-prefix && git clone --depth 1 ${INFINI_RT_GIT_URL:-https://gh-proxy.com/https://github.com/InfiniTensor/InfiniRT.git} /tmp/infinirt-src && cmake -S /tmp/infinirt-src -B /tmp/infinirt-build -DWITH_CPU=ON -DWITH_METAX=ON -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=/tmp/infinirt-prefix && cmake --build /tmp/infinirt-build -j$(nproc) && cmake --install /tmp/infinirt-build && python -m pip install .[dev] --no-build-isolation --config-settings=cmake.define.INFINI_OPS_SMOKE_BUILD=ON --config-settings=cmake.define.INFINI_RT_ROOT=/tmp/infinirt-prefix jobs: gpu: type: unittest @@ -109,7 +109,7 @@ platforms: PIP_INDEX_URL: https://pypi.org/simple docker_args: - "--privileged" - setup: pip install scikit-build-core pybind11 libclang pyyaml && rm -rf /tmp/infinirt-src /tmp/infinirt-build /tmp/infinirt-prefix && git clone --depth 1 ${INFINI_RT_GIT_URL:-https://gh-proxy.com/https://github.com/InfiniTensor/InfiniRT.git} /tmp/infinirt-src && cmake -S /tmp/infinirt-src -B /tmp/infinirt-build -DWITH_CPU=ON -DWITH_MOORE=ON -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=/tmp/infinirt-prefix && cmake --build /tmp/infinirt-build -j$(nproc) && cmake --install /tmp/infinirt-build && pip install .[dev] --no-build-isolation --config-settings=cmake.define.INFINI_OPS_SMOKE_BUILD=ON --config-settings=cmake.define.INFINI_RT_ROOT=/tmp/infinirt-prefix + setup: python -m pip install 'scikit-build-core>=1.1.1' pybind11 libclang packaging pyyaml jinja2 && rm -rf /tmp/infinirt-src /tmp/infinirt-build /tmp/infinirt-prefix && git clone --depth 1 ${INFINI_RT_GIT_URL:-https://gh-proxy.com/https://github.com/InfiniTensor/InfiniRT.git} /tmp/infinirt-src && cmake -S /tmp/infinirt-src -B /tmp/infinirt-build -DWITH_CPU=ON -DWITH_MOORE=ON -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=/tmp/infinirt-prefix && cmake --build /tmp/infinirt-build -j$(nproc) && cmake --install /tmp/infinirt-build && python -m pip install .[dev] --no-build-isolation --config-settings=cmake.define.INFINI_OPS_SMOKE_BUILD=ON --config-settings=cmake.define.INFINI_RT_ROOT=/tmp/infinirt-prefix jobs: gpu: type: unittest @@ -136,7 +136,7 @@ platforms: PIP_INDEX_URL: https://pypi.org/simple docker_args: - "--privileged" - setup: pip install scikit-build-core pybind11 libclang pyyaml && rm -rf /tmp/infinirt-src /tmp/infinirt-build /tmp/infinirt-prefix && git clone --depth 1 ${INFINI_RT_GIT_URL:-https://gh-proxy.com/https://github.com/InfiniTensor/InfiniRT.git} /tmp/infinirt-src && cmake -S /tmp/infinirt-src -B /tmp/infinirt-build -DWITH_CPU=ON -DWITH_CAMBRICON=ON -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=/tmp/infinirt-prefix && cmake --build /tmp/infinirt-build -j$(nproc) && cmake --install /tmp/infinirt-build && pip install .[dev] --no-build-isolation --config-settings=cmake.define.INFINI_OPS_SMOKE_BUILD=ON --config-settings=cmake.define.INFINI_RT_ROOT=/tmp/infinirt-prefix + setup: python -m pip install 'scikit-build-core>=1.1.1' pybind11 libclang packaging pyyaml jinja2 && rm -rf /tmp/infinirt-src /tmp/infinirt-build /tmp/infinirt-prefix && git clone --depth 1 ${INFINI_RT_GIT_URL:-https://gh-proxy.com/https://github.com/InfiniTensor/InfiniRT.git} /tmp/infinirt-src && cmake -S /tmp/infinirt-src -B /tmp/infinirt-build -DWITH_CPU=ON -DWITH_CAMBRICON=ON -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=/tmp/infinirt-prefix && cmake --build /tmp/infinirt-build -j$(nproc) && cmake --install /tmp/infinirt-build && python -m pip install .[dev] --no-build-isolation --config-settings=cmake.define.INFINI_OPS_SMOKE_BUILD=ON --config-settings=cmake.define.INFINI_RT_ROOT=/tmp/infinirt-prefix jobs: gpu: type: unittest @@ -172,7 +172,7 @@ platforms: - "--group-add=video" volumes: - /opt/hyhal:/opt/hyhal:ro - setup: pip install scikit-build-core pybind11 libclang pyyaml && rm -rf /tmp/infinirt-src /tmp/infinirt-build /tmp/infinirt-prefix && git clone --depth 1 ${INFINI_RT_GIT_URL:-https://gh-proxy.com/https://github.com/InfiniTensor/InfiniRT.git} /tmp/infinirt-src && cmake -S /tmp/infinirt-src -B /tmp/infinirt-build -DWITH_CPU=ON -DWITH_HYGON=ON -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=/tmp/infinirt-prefix && cmake --build /tmp/infinirt-build -j$(nproc) && cmake --install /tmp/infinirt-build && pip install .[dev] --no-build-isolation --config-settings=cmake.define.INFINI_OPS_SMOKE_BUILD=ON --config-settings=cmake.define.INFINI_RT_ROOT=/tmp/infinirt-prefix + setup: python -m pip install 'scikit-build-core>=1.1.1' pybind11 libclang packaging pyyaml jinja2 && rm -rf /tmp/infinirt-src /tmp/infinirt-build /tmp/infinirt-prefix && git clone --depth 1 ${INFINI_RT_GIT_URL:-https://gh-proxy.com/https://github.com/InfiniTensor/InfiniRT.git} /tmp/infinirt-src && cmake -S /tmp/infinirt-src -B /tmp/infinirt-build -DWITH_CPU=ON -DWITH_HYGON=ON -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=/tmp/infinirt-prefix && cmake --build /tmp/infinirt-build -j$(nproc) && cmake --install /tmp/infinirt-build && python -m pip install .[dev] --no-build-isolation --config-settings=cmake.define.INFINI_OPS_SMOKE_BUILD=ON --config-settings=cmake.define.INFINI_RT_ROOT=/tmp/infinirt-prefix jobs: gpu: type: unittest @@ -207,7 +207,7 @@ platforms: - /usr/local/bin/npu-smi:/usr/local/bin/npu-smi:ro env: ASCEND_HOME_PATH: /usr/local/Ascend/ascend-toolkit/latest - setup: pip install scikit-build-core pybind11 libclang pyyaml && rm -rf /tmp/infinirt-src /tmp/infinirt-build /tmp/infinirt-prefix && git clone --depth 1 ${INFINI_RT_GIT_URL:-https://gh-proxy.com/https://github.com/InfiniTensor/InfiniRT.git} /tmp/infinirt-src && cmake -S /tmp/infinirt-src -B /tmp/infinirt-build -DWITH_CPU=ON -DWITH_ASCEND=ON -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=/tmp/infinirt-prefix && cmake --build /tmp/infinirt-build -j$(nproc) && cmake --install /tmp/infinirt-build && pip install .[dev] --no-build-isolation --config-settings=cmake.define.INFINI_OPS_SMOKE_BUILD=ON --config-settings=cmake.define.INFINI_RT_ROOT=/tmp/infinirt-prefix + setup: python -m pip install 'scikit-build-core>=1.1.1' pybind11 libclang packaging pyyaml jinja2 && rm -rf /tmp/infinirt-src /tmp/infinirt-build /tmp/infinirt-prefix && git clone --depth 1 ${INFINI_RT_GIT_URL:-https://gh-proxy.com/https://github.com/InfiniTensor/InfiniRT.git} /tmp/infinirt-src && cmake -S /tmp/infinirt-src -B /tmp/infinirt-build -DWITH_CPU=ON -DWITH_ASCEND=ON -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=/tmp/infinirt-prefix && cmake --build /tmp/infinirt-build -j$(nproc) && cmake --install /tmp/infinirt-build && python -m pip install .[dev] --no-build-isolation --config-settings=cmake.define.INFINI_OPS_SMOKE_BUILD=ON --config-settings=cmake.define.INFINI_RT_ROOT=/tmp/infinirt-prefix jobs: npu: type: unittest diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 1d4305c09..0b6e72703 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -1549,7 +1549,9 @@ if(GENERATE_PYTHON_BINDINGS) set(_iluvatar_dispatch_include_flags "-I${CMAKE_CURRENT_SOURCE_DIR}" "-I${PROJECT_SOURCE_DIR}" + "-I${PROJECT_SOURCE_DIR}/include" "-I${PROJECT_SOURCE_DIR}/generated" + "-I${PROJECT_SOURCE_DIR}/generated/include" ${INFINI_RT_INCLUDE_FLAGS}) foreach(_dir IN LISTS TORCH_INCLUDE_DIRS CUDAToolkit_INCLUDE_DIRS) list(APPEND _iluvatar_dispatch_include_flags "-I${_dir}") diff --git a/src/native/ascend/custom/cmake/config_ascend.cmake b/src/native/ascend/custom/cmake/config_ascend.cmake index 29bbee85f..bfbb07a6a 100644 --- a/src/native/ascend/custom/cmake/config_ascend.cmake +++ b/src/native/ascend/custom/cmake/config_ascend.cmake @@ -6,7 +6,7 @@ endif() set(ASCEND_CANN_PACKAGE_PATH ${ASCEND_HOME_PATH}) -# Auto-detect `SOC_VERSION` from `npu-smi info` if not set externally. +# Auto-detect `SOC_VERSION` from the runtime or `npu-smi` if not set externally. # Required by `CANN`'s `ascendc.cmake` for `AscendC` kernel compilation. if(NOT DEFINED SOC_VERSION OR "${SOC_VERSION}" STREQUAL "") include(${CMAKE_CURRENT_LIST_DIR}/detect_soc.cmake) diff --git a/src/native/ascend/custom/cmake/detect_soc.cmake b/src/native/ascend/custom/cmake/detect_soc.cmake index 548ead0aa..d0feb3a7d 100644 --- a/src/native/ascend/custom/cmake/detect_soc.cmake +++ b/src/native/ascend/custom/cmake/detect_soc.cmake @@ -1,9 +1,4 @@ -# Auto-detect the Ascend SOC version from `npu-smi info`. -# -# `infiniops_detect_soc()` parses the first `910*` / `310*` entry -# in `npu-smi info` and writes `Ascend` into the named variable in -# the caller's scope. Falls back to `Ascend910B4` when detection fails -# (no NPU on the host, `npu-smi` missing, output format mismatch). +# Query the runtime for the exact SOC name before consulting `npu-smi`. # # Called from both `src/CMakeLists.txt` (outer `pip install` build, to # forward `SOC_VERSION` to the standalone `build.sh` invocation) and @@ -11,14 +6,52 @@ # by that `build.sh`). function(infiniops_detect_soc out_var) - execute_process( - COMMAND bash -c "npu-smi info 2>/dev/null | awk '/910B|910A|310/ {for (i=1;i<=NF;i++) if ($i ~ /^(910|310)/) {print \"Ascend\" $i; exit}}'" - OUTPUT_VARIABLE _detected - OUTPUT_STRIP_TRAILING_WHITESPACE) - - if(_detected) - set(${out_var} "${_detected}" PARENT_SCOPE) + set(_soc_pattern "910(B[1-4](-1)?|A|ProA|ProB|PremiumA|_[0-9][0-9A-Za-z]*)|310[BP][0-9]+(vir[0-9]+)?") + if(_TORCH_PYTHON) + set(_soc_python "${_TORCH_PYTHON}") + elseif(Python_EXECUTABLE) + set(_soc_python "${Python_EXECUTABLE}") else() - set(${out_var} "Ascend910B4" PARENT_SCOPE) + find_program(_soc_python NAMES python3 python) + endif() + + if(_soc_python) + execute_process( + COMMAND "${_soc_python}" -c + "import torch, torch_npu; print(torch.npu.get_device_name(0))" + OUTPUT_VARIABLE _runtime_soc + OUTPUT_STRIP_TRAILING_WHITESPACE + RESULT_VARIABLE _runtime_result + ERROR_QUIET + TIMEOUT 30) + if(_runtime_result STREQUAL "0") + string(REGEX MATCH "^(Ascend)?(${_soc_pattern})$" + _runtime_match "${_runtime_soc}") + if(_runtime_match) + set(${out_var} "Ascend${CMAKE_MATCH_2}" PARENT_SCOPE) + return() + endif() + endif() endif() + + execute_process( + COMMAND npu-smi info + OUTPUT_VARIABLE _smi_output + RESULT_VARIABLE _smi_result + ERROR_QUIET + TIMEOUT 10) + if(_smi_result STREQUAL "0") + # Require model-qualified tokens so HBM usage such as `3103` is ignored. + string(REGEX MATCH "(^|[ \t\r\n|])(Ascend)?(${_soc_pattern})([ \t\r\n|]|$)" + _smi_match "${_smi_output}") + if(_smi_match) + set(${out_var} "Ascend${CMAKE_MATCH_3}" PARENT_SCOPE) + return() + endif() + endif() + + message(FATAL_ERROR + "Cannot determine an exact Ascend SOC version. " + "Set SOC_VERSION explicitly (pip: --config-settings=cmake.define.SOC_VERSION=). " + "Generic model names and numeric npu-smi fields are not valid SOC versions.") endfunction() diff --git a/tests/build_config/test_ci_build.py b/tests/build_config/test_ci_build.py new file mode 100644 index 000000000..483575dcc --- /dev/null +++ b/tests/build_config/test_ci_build.py @@ -0,0 +1,219 @@ +"""Hardware-independent regressions for CI build setup and SOC discovery.""" + +import os +import pathlib +import shlex +import shutil +import subprocess +import sys + +import pytest +import yaml +from packaging.requirements import Requirement + +try: + import tomllib +except ModuleNotFoundError: + import tomli as tomllib + + +ROOT = pathlib.Path(__file__).resolve().parents[2] +SOC_DIR = ROOT / "src/native/ascend/custom/cmake" +CMAKE = shutil.which("cmake") + + +@pytest.mark.parametrize( + "platform", ("nvidia", "iluvatar", "metax", "moore", "cambricon", "hygon", "ascend") +) +def test_ci_installs_declared_build_requirements(platform): + project = tomllib.loads((ROOT / "pyproject.toml").read_text()) + config = yaml.safe_load((ROOT / ".github/ci_config.yml").read_text()) + commands = config["platforms"][platform]["setup"].split("&&") + install = shlex.split(commands[0]) + assert install[:4] == ["python", "-m", "pip", "install"] + requirements = {Requirement(arg).name: Requirement(arg) for arg in install[4:]} + + for declared in project["build-system"]["requires"]: + required = Requirement(declared) + assert required.name in requirements + assert requirements[required.name].specifier == required.specifier + + for command in commands: + tokens = shlex.split(command) + + if "install" in tokens and "pip" in tokens: + assert tokens[:4] == ["python", "-m", "pip", "install"] + + +def test_wheel_metadata_accepts_sdist_configuration(tmp_path, monkeypatch): + from scikit_build_core.build import prepare_metadata_for_build_wheel + + monkeypatch.chdir(ROOT) + dist_info = prepare_metadata_for_build_wheel(str(tmp_path)) + metadata = (tmp_path / dist_info / "METADATA").read_text() + assert "Name: InfiniOps" in metadata + assert "Version: 0.1.0" in metadata + + +@pytest.fixture +def soc_probe(tmp_path, monkeypatch, fake_smi): + if not CMAKE: + pytest.skip("cmake is unavailable") + + fake_modules = tmp_path / "python" + fake_modules.mkdir() + (fake_modules / "torch_npu.py").write_text("") + (fake_modules / "torch.py").write_text( + "import os\n" + "class npu:\n" + " @staticmethod\n" + " def get_device_name(index):\n" + " name = os.environ.get('TEST_RUNTIME_SOC')\n" + " if not name:\n" + " raise RuntimeError('No NPU runtime is available.')\n" + " return name\n" + ) + smi_output = tmp_path / "smi.txt" + monkeypatch.setenv("TEST_SMI_OUTPUT", str(smi_output)) + monkeypatch.setenv("PATH", str(fake_smi) + os.pathsep + os.environ["PATH"]) + monkeypatch.setenv("PYTHONPATH", str(fake_modules)) + output = tmp_path / "soc.txt" + script = tmp_path / "probe.cmake" + script.write_text( + f'include("{SOC_DIR.as_posix()}/detect_soc.cmake")\n' + "infiniops_detect_soc(DETECTED_SOC)\n" + f'file(WRITE "{output.as_posix()}" "${{DETECTED_SOC}}")\n' + ) + + def run(runtime_soc, smi_text, *, torch_python=True): + monkeypatch.setenv("TEST_RUNTIME_SOC", runtime_soc) + smi_output.write_text(smi_text) + interpreter = "_TORCH_PYTHON" if torch_python else "Python_EXECUTABLE" + command = [CMAKE, f"-D{interpreter}={sys.executable}"] + + if torch_python: + command.append("-DPython_EXECUTABLE=/missing/python") + + result = subprocess.run( + [*command, "-P", str(script)], capture_output=True, text=True, timeout=45 + ) + name = output.read_text() if output.exists() else None + + return result, name + + return run + + +@pytest.fixture(scope="session") +def fake_smi(tmp_path_factory): + directory = tmp_path_factory.mktemp("npu-smi") + smi = directory / ("npu-smi.exe" if os.name == "nt" else "npu-smi") + + if os.name == "nt": + compiler = shutil.which("g++") or shutil.which("clang++") + if not compiler: + pytest.skip("A C++ compiler is required for the Windows npu-smi fixture") + source = directory / "smi.cc" + source.write_text( + "#include \n#include \n#include \n" + 'int main() { std::ifstream input(std::getenv("TEST_SMI_OUTPUT")); ' + "std::cout << input.rdbuf(); return 0; }\n" + ) + subprocess.run([compiler, str(source), "-o", str(smi)], check=True, timeout=30) + else: + smi.write_text('#!/bin/sh\ncat "$TEST_SMI_OUTPUT"\n') + smi.chmod(0o755) + + return directory + + +@pytest.mark.parametrize("torch_python", (True, False)) +def test_runtime_soc_takes_priority_over_smi(soc_probe, torch_python): + result, name = soc_probe( + "Ascend910_9362", "| 0 | 910B4 | 3103 |", torch_python=torch_python + ) + assert result.returncode == 0, result.stderr + assert name == "Ascend910_9362" + + +@pytest.mark.parametrize( + "model", ("910B4", "910B4-1", "Ascend910_9382", "310P3", "310P3vir02") +) +def test_smi_accepts_qualified_models(soc_probe, model): + result, name = soc_probe("", f"| 0 | {model} | OK |\n| 3103 | 3109 |") + assert result.returncode == 0, result.stderr + assert name == (model if model.startswith("Ascend") else f"Ascend{model}") + + +@pytest.mark.parametrize( + "smi_text", ("", "| Ascend910 | 3103 | 3109 |", "| 3103 |", "| 910B40garbage |") +) +def test_ambiguous_soc_requires_explicit_configuration(soc_probe, smi_text): + result, name = soc_probe("Ascend910", smi_text) + assert result.returncode != 0 + assert name is None + assert "Set SOC_VERSION explicitly" in result.stderr + + +def test_explicit_soc_bypasses_detection(tmp_path): + if not CMAKE: + pytest.skip("cmake is unavailable") + + toolkit = tmp_path / "toolkit" + ascendc = toolkit / "tools/tikcpp/ascendc_kernel_cmake" + ascendc.mkdir(parents=True) + (ascendc / "ascendc.cmake").write_text("") + script = tmp_path / "explicit.cmake" + script.write_text( + 'set(SOC_VERSION "Ascend910_9382")\n' + f'set(ASCEND_HOME_PATH "{toolkit.as_posix()}")\n' + f'include("{SOC_DIR.as_posix()}/config_ascend.cmake")\n' + 'if(NOT SOC_VERSION STREQUAL "Ascend910_9382")\n' + ' message(FATAL_ERROR "Explicit SOC version was changed.")\n' + "endif()\n" + ) + result = subprocess.run( + [CMAKE, "-P", str(script)], capture_output=True, text=True, timeout=15 + ) + assert result.returncode == 0, result.stderr + + +def test_iluvatar_dispatch_can_include_public_and_generated_headers(tmp_path): + compiler = shutil.which("g++") or shutil.which("clang++") + + if not CMAKE or not compiler: + pytest.skip("cmake and a C++ compiler are required") + + source = (ROOT / "src/CMakeLists.txt").read_text(encoding="utf-8") + start = source.index(" set(_iluvatar_dispatch_include_flags") + end = source.index(" foreach(", start) + include_flags = source[start:end] + + for directory in ("include", "generated/include"): + headers = tmp_path / directory / "infini" + headers.mkdir(parents=True) + header = ( + "public.h" if directory == "include" else "operator_call_instantiations.h" + ) + (headers / header).write_text("#pragma once\n") + + dispatch = tmp_path / "dispatch.cc" + dispatch.write_text( + '#include "infini/public.h"\n#include "infini/operator_call_instantiations.h"\n' + ) + script = tmp_path / "compile.cmake" + script.write_text( + f'set(PROJECT_SOURCE_DIR "{tmp_path.as_posix()}")\n' + f'set(CMAKE_CURRENT_SOURCE_DIR "{tmp_path.as_posix()}/src")\n' + f"{include_flags}\n" + f'execute_process(COMMAND "{pathlib.Path(compiler).as_posix()}" ' + f'${{_iluvatar_dispatch_include_flags}} -fsyntax-only "{dispatch.as_posix()}" ' + "RESULT_VARIABLE result)\n" + 'if(NOT result STREQUAL "0")\n' + ' message(FATAL_ERROR "Dispatch header lookup failed.")\n' + "endif()\n" + ) + result = subprocess.run( + [CMAKE, "-P", str(script)], capture_output=True, text=True, timeout=30 + ) + assert result.returncode == 0, result.stderr From 9a0b89e38582d642f20fe7fa67b0da8f1859717f Mon Sep 17 00:00:00 2001 From: Li Baoming <1508269885@qq.com> Date: Tue, 27 Oct 2026 17:45:53 +0800 Subject: [PATCH 2/9] fix: resolve NVIDIA dispatch and MetaX compiler configuration failures --- CMakeLists.txt | 11 +- .../cuda/nvidia/ops/quick_gelu/kernel.cu | 34 ++-- src/operator.h | 20 ++- tests/build_config/nvidia_dispatch.cu | 15 ++ tests/build_config/test_compilers.py | 147 ++++++++++++++++++ 5 files changed, 206 insertions(+), 21 deletions(-) create mode 100644 tests/build_config/nvidia_dispatch.cu create mode 100644 tests/build_config/test_compilers.py diff --git a/CMakeLists.txt b/CMakeLists.txt index 02c416958..7cbe93d18 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -18,7 +18,7 @@ if(SKBUILD_STATE STREQUAL "sdist") DESTINATION "${CMAKE_BINARY_DIR}/flash-attention-cutlass-source") return() endif() -project(InfiniOps LANGUAGES CXX) +project(InfiniOps LANGUAGES NONE) if(POLICY CMP0116) cmake_policy(SET CMP0116 NEW) @@ -210,6 +210,13 @@ if(AUTO_DETECT_DEVICES) endif() endif() +# Select the MetaX compiler before CMake initializes CXX and its cache. +if(WITH_METAX) + set(CMAKE_C_COMPILER "${CMAKE_CURRENT_SOURCE_DIR}/scripts/mxcc_wrapper.sh") + set(CMAKE_CXX_COMPILER "${CMAKE_CURRENT_SOURCE_DIR}/scripts/mxcc_wrapper.sh") +endif() +enable_language(CXX) + if(BUILD_NVIDIA_FLASH_INFERENCE AND NOT WITH_NVIDIA) message(FATAL_ERROR "Native FlashAttention requires WITH_NVIDIA") endif() @@ -567,8 +574,6 @@ if(WITH_METAX) # Normally can be found at: `/opt/maca/`. set(MACA_PATH $ENV{MACA_PATH}) - set(CMAKE_C_COMPILER ${CMAKE_CURRENT_SOURCE_DIR}/scripts/mxcc_wrapper.sh) - set(CMAKE_CXX_COMPILER ${CMAKE_CURRENT_SOURCE_DIR}/scripts/mxcc_wrapper.sh) include_directories("${MACA_PATH}/include") link_directories("${MACA_PATH}/lib") diff --git a/src/native/cuda/nvidia/ops/quick_gelu/kernel.cu b/src/native/cuda/nvidia/ops/quick_gelu/kernel.cu index b97cc21ae..b690a08e6 100644 --- a/src/native/cuda/nvidia/ops/quick_gelu/kernel.cu +++ b/src/native/cuda/nvidia/ops/quick_gelu/kernel.cu @@ -1,5 +1,3 @@ -#include "native/cuda/nvidia/ops/quick_gelu/kernel.h" - #include #include #include @@ -9,6 +7,8 @@ #include #include +#include "native/cuda/nvidia/ops/quick_gelu/kernel.h" + namespace infini::ops::nvidia_quick_gelu_detail { namespace { // The three branches preserve Core 0.2.9 quickgelu/cuda/kernel.cuh semantics. @@ -18,7 +18,12 @@ __device__ __forceinline__ T Apply(T x) { constexpr float alpha = 1.702f; if constexpr (std::is_same_v) { const __half ax = __hmul(__float2half(alpha), x); - const __half denominator = __hadd(__float2half(1.0f), hexp(__hneg(ax))); +#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 530 + const __half exp_ax = __float2half(__expf(-__half2float(ax))); +#else + const __half exp_ax = hexp(__hneg(ax)); +#endif + const __half denominator = __hadd(__float2half(1.0f), exp_ax); const __half sigmoid = hrcp(denominator); return __hmul(x, sigmoid); } else if constexpr (std::is_same_v) { @@ -44,8 +49,9 @@ template __global__ void QuickGeluKernel(const T* input, T* out, Layout layout, std::size_t count) { const std::size_t step = static_cast(blockDim.x) * gridDim.x; - for (std::size_t index = static_cast(blockIdx.x) * blockDim.x + - threadIdx.x; index < count; index += step) { + for (std::size_t index = + static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + index < count; index += step) { std::size_t remainder = index, input_offset = 0, out_offset = 0; for (unsigned axis = layout.rank; axis > 0; --axis) { const auto coordinate = remainder % layout.shape[axis - 1]; @@ -71,12 +77,20 @@ void Submit(const void* input, void* out, Layout layout, std::size_t count, void Launch(const void* input, void* out, DataType dtype, Layout layout, std::size_t count, cudaStream_t stream) { switch (dtype) { - case DataType::kFloat16: Submit<__half>(input, out, layout, count, stream); break; - case DataType::kBFloat16: Submit<__nv_bfloat16>(input, out, layout, count, stream); break; - case DataType::kFloat32: Submit(input, out, layout, count, stream); break; - default: throw std::invalid_argument("NVIDIA QuickGelu unsupported dtype"); + case DataType::kFloat16: + Submit<__half>(input, out, layout, count, stream); + break; + case DataType::kBFloat16: + Submit<__nv_bfloat16>(input, out, layout, count, stream); + break; + case DataType::kFloat32: + Submit(input, out, layout, count, stream); + break; + default: + throw std::invalid_argument("NVIDIA QuickGelu unsupported dtype"); } const auto status = cudaPeekAtLastError(); - if (status != cudaSuccess) throw std::runtime_error(cudaGetErrorString(status)); + if (status != cudaSuccess) + throw std::runtime_error(cudaGetErrorString(status)); } } // namespace infini::ops::nvidia_quick_gelu_detail diff --git a/src/operator.h b/src/operator.h index 3ec0c6733..2b4659c01 100644 --- a/src/operator.h +++ b/src/operator.h @@ -653,13 +653,14 @@ class Operator : public OperatorBase { DispatchFunc>( dispatch_device_type, - [&](auto device_tag) { + [&](auto device_tag, auto&... device_args) { constexpr Device::Type kDev = decltype(device_tag)::value; using Implementations = typename ActiveImplementations::type; const bool automatic = config.implementation_index() == Config::kAutomaticImplementation; - const auto construct = [&](auto implementation_tag) { + const auto construct = [&](auto implementation_tag, + auto&... provider_args) { constexpr std::size_t kImplementationIndex = decltype(implementation_tag)::value; using Provider = Operator; @@ -671,10 +672,11 @@ class Operator : public OperatorBase { std::is_constructible_v) { bool supported = false; if constexpr (kConfigAware) { - supported = Provider::supports(config, std::as_const(args)...); + supported = + Provider::supports(config, std::as_const(provider_args)...); } else if (!config.requires_config_aware_provider()) { supported = detail::ProviderSupports( - 0, std::as_const(args)...); + 0, std::as_const(provider_args)...); } if (!supported) { if (automatic) return; @@ -687,7 +689,8 @@ class Operator : public OperatorBase { } matched = true; selected(device_tag, implementation_tag, - std::bool_constant{}, args...); + std::bool_constant{}, + provider_args...); } else { if (!automatic && strict) { throw std::invalid_argument( @@ -701,7 +704,7 @@ class Operator : public OperatorBase { for (const auto index : detail::ListToVector(Implementations{})) { detail::DispatchImplementation(index, construct, "Operator::Make(automatic)", - Implementations{}); + Implementations{}, device_args...); if (matched) break; } } else { @@ -712,10 +715,11 @@ class Operator : public OperatorBase { } detail::DispatchImplementation( config.implementation_index(), construct, - "Operator::Make(implementation_index)", Implementations{}); + "Operator::Make(implementation_index)", Implementations{}, + device_args...); } }, - "Operator::Make"); + "Operator::Make", args...); return matched; } diff --git a/tests/build_config/nvidia_dispatch.cu b/tests/build_config/nvidia_dispatch.cu new file mode 100644 index 000000000..d6aaec982 --- /dev/null +++ b/tests/build_config/nvidia_dispatch.cu @@ -0,0 +1,15 @@ +#include "native/cuda/nvidia/ops/quick_gelu/kernel.h" + +namespace infini::ops { + +bool ProbeQuickGeluSelection(const Config& config, Tensor input, + Tensor output) { + return Operator::CanMake(config, input, output); +} + +std::unique_ptr> ProbeQuickGeluConstruction( + const Config& config, Tensor input, Tensor output) { + return Operator::Make(config, input, output); +} + +} // namespace infini::ops diff --git a/tests/build_config/test_compilers.py b/tests/build_config/test_compilers.py new file mode 100644 index 000000000..a393e280e --- /dev/null +++ b/tests/build_config/test_compilers.py @@ -0,0 +1,147 @@ +"""Compile regressions for backend compiler selection and NVIDIA dispatch.""" + +import os +import pathlib +import shlex +import shutil +import subprocess + +import pytest + + +ROOT = pathlib.Path(__file__).resolve().parents[2] + + +@pytest.mark.skipif(os.name == "nt", reason="MetaX uses a POSIX compiler wrapper") +@pytest.mark.parametrize("auto_detect", (False, True)) +def test_metax_compiler_preserves_infinirt_configuration( + tmp_path, monkeypatch, auto_detect +): + cmake = shutil.which("cmake") + compiler = shutil.which("g++") or shutil.which("clang++") + + if not cmake or not compiler: + pytest.skip("cmake and a C++ compiler are required") + + toolkit = tmp_path / "maca" + mxcc = toolkit / "mxgpu_llvm/bin/mxcc" + mxcc.parent.mkdir(parents=True) + mxcc.write_text(f'#!/bin/sh\nexec {shlex.quote(compiler)} "$@"\n') + mxcc.chmod(0o755) + libraries = toolkit / "lib" + libraries.mkdir() + + for name in ("mcruntime", "mcdnn", "mcblas"): + (libraries / f"lib{name}.so").touch() + + monkeypatch.setenv("MACA_PATH", str(toolkit)) + + for name in ( + "INFINI_RT_ROOT", + "HPCC_PATH", + "NEUWARE_HOME", + "MUSA_ROOT", + "MUSA_HOME", + "MUSA_PATH", + "ASCEND_HOME_PATH", + ): + monkeypatch.delenv(name, raising=False) + + wrapper = tmp_path / "scripts/mxcc_wrapper.sh" + wrapper.parent.mkdir() + shutil.copyfile(ROOT / "scripts/mxcc_wrapper.sh", wrapper) + wrapper.chmod(0o755) + prefix = tmp_path / "infinirt" + (prefix / "include/infini").mkdir(parents=True) + (prefix / "include/infini/rt.h").touch() + (prefix / "lib").mkdir() + (prefix / "lib/libinfinirt.so").touch() + + source = (ROOT / "CMakeLists.txt").read_text(encoding="utf-8") + initialization = source.split("if(AUTO_DETECT_BACKENDS)", 1)[0] + metax_start = source.index("if(WITH_METAX)\n add_compile_definitions") + metax = source[metax_start : source.index("if(WITH_MARS)", metax_start)] + runtime_start = source.index('set(INFINI_RT_ROOT "$ENV{INFINI_RT_ROOT}"') + runtime = source[ + runtime_start : source.index("add_subdirectory(src)", runtime_start) + ] + (tmp_path / "CMakeLists.txt").write_text( + initialization + metax + runtime + "\nadd_subdirectory(dependency)\n" + 'file(WRITE "${CMAKE_BINARY_DIR}/selected-compiler.txt" ' + '"${CMAKE_CXX_COMPILER}")\nadd_library(probe OBJECT probe.cc)\n' + ) + dependency = tmp_path / "dependency" + dependency.mkdir() + (dependency / "CMakeLists.txt").write_text( + "cmake_minimum_required(VERSION 3.18)\nproject(Dependency LANGUAGES CXX)\n" + ) + (tmp_path / "probe.cc").write_text("int probe() { return 0; }\n") + build = tmp_path / "build" + configure = [ + cmake, + "-S", + str(tmp_path), + "-B", + str(build), + f"-DINFINI_RT_ROOT={prefix}", + f"-DAUTO_DETECT_DEVICES={'ON' if auto_detect else 'OFF'}", + ] + + if not auto_detect: + configure.append("-DWITH_METAX=ON") + + for _ in range(2): + result = subprocess.run(configure, capture_output=True, text=True, timeout=45) + assert result.returncode == 0, result.stdout + result.stderr + assert "require your cache to be deleted" not in result.stdout + result.stderr + assert "Using installed InfiniRT:" in result.stdout + cache = (build / "CMakeCache.txt").read_text() + assert (build / "selected-compiler.txt").read_text() == str(wrapper) + assert f"INFINI_RT_ROOT:PATH={prefix}" in cache + + subprocess.run([cmake, "--build", str(build)], check=True, timeout=30) + + +@pytest.mark.parametrize("architecture", ("sm_52", "sm_80")) +def test_nvidia_dispatch_and_quick_gelu_compile(tmp_path, architecture): + nvcc = shutil.which("nvcc") + runtime = os.environ.get("INFINI_RT_ROOT") + + if not nvcc or not runtime: + pytest.skip("nvcc and an installed INFINI_RT_ROOT are required") + + supported = subprocess.run( + [nvcc, "--list-gpu-code"], + check=True, + capture_output=True, + text=True, + timeout=15, + ).stdout.split() + + if architecture not in supported: + pytest.skip(f"nvcc does not support {architecture}") + + for source in ( + ROOT / "tests/build_config/nvidia_dispatch.cu", + ROOT / "src/native/cuda/nvidia/ops/quick_gelu/kernel.cu", + ): + result = subprocess.run( + [ + nvcc, + "-std=c++17", + f"-arch={architecture}", + "-DWITH_CPU=1", + "-DWITH_NVIDIA=1", + f"-I{ROOT / 'src'}", + f"-I{ROOT / 'include'}", + f"-I{pathlib.Path(runtime) / 'include'}", + "-c", + str(source), + "-o", + str(tmp_path / (source.stem + ".o")), + ], + capture_output=True, + text=True, + timeout=90, + ) + assert result.returncode == 0, result.stdout + result.stderr From e56e89cac3ba2ceb94ec1684012f7681ea894b9b Mon Sep 17 00:00:00 2001 From: Li Baoming <1508269885@qq.com> Date: Tue, 27 Oct 2026 17:52:12 +0800 Subject: [PATCH 3/9] fix: use host compiler for MetaX CI test probes --- .github/ci_config.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.github/ci_config.yml b/.github/ci_config.yml index ec55f4ea9..58efcd877 100644 --- a/.github/ci_config.yml +++ b/.github/ci_config.yml @@ -96,7 +96,7 @@ platforms: junit_path: test-results.xml stages: - name: test - run: pytest tests/ -m smoke -n 1 -v --tb=short --junitxml=/workspace/results/test-results.xml + run: CXX=/usr/bin/g++ pytest tests/ -m smoke -n 1 -v --tb=short --junitxml=/workspace/results/test-results.xml moore: runner_label: Moore From b3bb5ab1053172d1a743e18218f882f2115c0873 Mon Sep 17 00:00:00 2001 From: Li Baoming <1508269885@qq.com> Date: Thu, 29 Oct 2026 15:41:51 +0800 Subject: [PATCH 4/9] fix: resolve generated dispatch and accelerator CI test failures --- .github/ci_config.yml | 2 +- scripts/generate_wrappers.py | 3 ++- tests/build_config/test_compilers.py | 39 ++++++++++++++++++++++++++++ tests/test_cast.py | 35 ++++++++++++++++++++++--- tests/test_cpp_api.py | 2 +- 5 files changed, 75 insertions(+), 6 deletions(-) diff --git a/.github/ci_config.yml b/.github/ci_config.yml index 58efcd877..3149e6163 100644 --- a/.github/ci_config.yml +++ b/.github/ci_config.yml @@ -96,7 +96,7 @@ platforms: junit_path: test-results.xml stages: - name: test - run: CXX=/usr/bin/g++ pytest tests/ -m smoke -n 1 -v --tb=short --junitxml=/workspace/results/test-results.xml + run: HOST_CXX=/usr/bin/g++ pytest tests/ -m smoke -n 1 -v --tb=short --junitxml=/workspace/results/test-results.xml moore: runner_label: Moore diff --git a/scripts/generate_wrappers.py b/scripts/generate_wrappers.py index 38cb3e8f9..a4cb10c62 100644 --- a/scripts/generate_wrappers.py +++ b/scripts/generate_wrappers.py @@ -1439,7 +1439,8 @@ def _generate_generated_dispatch_header(op_names, devices, declarations): def _generate_generated_dispatch_source(impl_paths, definitions): impl_includes = "\n".join(f'#include "{impl_path}"' for impl_path in impl_paths) - return f"""#include "generated_dispatch.h" + return f"""#include "src/operator.h" +#include "generated_dispatch.h" #include "infini/operator_call_instantiations.h" #include "host_range_profiler.h" diff --git a/tests/build_config/test_compilers.py b/tests/build_config/test_compilers.py index a393e280e..8ab588ce7 100644 --- a/tests/build_config/test_compilers.py +++ b/tests/build_config/test_compilers.py @@ -2,6 +2,7 @@ import os import pathlib +import runpy import shlex import shutil import subprocess @@ -12,6 +13,24 @@ ROOT = pathlib.Path(__file__).resolve().parents[2] +@pytest.mark.skipif(os.name == "nt", reason="the GPU compiler wrapper uses POSIX") +def test_workspace_probe_uses_host_compiler(tmp_path, monkeypatch): + compiler = shutil.which("g++") or shutil.which("clang++") + + if not compiler: + pytest.skip("a host C++ compiler is required") + + gpu_compiler = tmp_path / "gpu-cxx" + gpu_compiler.write_text( + '#!/bin/sh\necho "GPU compiler used for host probe" >&2\nexit 1\n' + ) + gpu_compiler.chmod(0o755) + monkeypatch.setenv("CXX", str(gpu_compiler)) + monkeypatch.setenv("HOST_CXX", compiler) + module = runpy.run_path(str(ROOT / "tests/test_cpp_api.py")) + module["test_ascend_workspace_pool_retains_allocations_until_destruction"](tmp_path) + + @pytest.mark.skipif(os.name == "nt", reason="MetaX uses a POSIX compiler wrapper") @pytest.mark.parametrize("auto_detect", (False, True)) def test_metax_compiler_preserves_infinirt_configuration( @@ -121,9 +140,27 @@ def test_nvidia_dispatch_and_quick_gelu_compile(tmp_path, architecture): if architecture not in supported: pytest.skip(f"nvcc does not support {architecture}") + public_header = runpy.run_path( + str(ROOT / "scripts/generate_public_operator_header.py") + ) + public_header["generate_public_operator_header"]( + ROOT / "src/operator.h", tmp_path / "operator.h" + ) + (tmp_path / "generated_dispatch.h").write_text('#include "operator.h"\n') + (tmp_path / "infini").mkdir() + (tmp_path / "infini/operator_call_instantiations.h").write_text("\n") + generator = runpy.run_path(str(ROOT / "scripts/generate_wrappers.py")) + dispatch = tmp_path / "generated_dispatch_regression.cu" + dispatch.write_text( + generator["_generate_generated_dispatch_source"]( + ["native/cuda/nvidia/ops/cat/kernel.h"], [] + ) + ) + for source in ( ROOT / "tests/build_config/nvidia_dispatch.cu", ROOT / "src/native/cuda/nvidia/ops/quick_gelu/kernel.cu", + dispatch, ): result = subprocess.run( [ @@ -132,6 +169,8 @@ def test_nvidia_dispatch_and_quick_gelu_compile(tmp_path, architecture): f"-arch={architecture}", "-DWITH_CPU=1", "-DWITH_NVIDIA=1", + *([f"-I{tmp_path}"] if source == dispatch else []), + f"-I{ROOT}", f"-I{ROOT / 'src'}", f"-I{ROOT / 'include'}", f"-I{pathlib.Path(runtime) / 'include'}", diff --git a/tests/test_cast.py b/tests/test_cast.py index 6412458ec..7c7fa39b6 100644 --- a/tests/test_cast.py +++ b/tests/test_cast.py @@ -41,9 +41,20 @@ def test_cast( ): input = randn_strided(shape, input_strides, dtype=input_dtype, device=device) out = empty_strided(shape, out_strides, dtype=out_dtype, device=device) + cast = _cast + + if ( + device == "cuda" + and infini.ops.Cast.active_implementation_indices("nvidia") + and (not input.is_contiguous() or not out.is_contiguous()) + ): + if 8 not in infini.ops.Cast.active_implementation_indices("nvidia"): + pytest.skip("strided NVIDIA Cast requires ATen slot 8") + + cast = _cast_aten return Payload( - _cast, + cast, _torch_cast, (input, out), {}, @@ -52,18 +63,36 @@ def test_cast( ) -def _cast(input, out): - infini.ops.cast(input, out, stream=get_stream(input.device)) +def _cast(input, out, **kwargs): + infini.ops.cast(input, out, stream=get_stream(input.device), **kwargs) return out +def _cast_aten(input, out): + return _cast(input, out, implementation_index=8) + + def _torch_cast(input, out): out.copy_(input.to(out.dtype)) return out +@pytest.mark.smoke +def test_native_cuda_cast_rejects_strided_layout(device): + if device != "cuda" or 0 not in infini.ops.Cast.active_implementation_indices( + "nvidia" + ): + pytest.skip("native NVIDIA Cast is not active") + + input = randn_strided((13, 4), (10, 1), device=device) + out = empty_strided((13, 4), (10, 1), dtype=torch.float16, device=device) + + with pytest.raises(ValueError, match="does not support these parameters"): + infini.ops.cast(input, out, implementation_index=0, stream=get_stream(device)) + + @pytest.fixture def require_aten_cast(device): if 8 not in infini.ops.Cast.active_implementation_indices(device): diff --git a/tests/test_cpp_api.py b/tests/test_cpp_api.py index 72951a4ea..1c2125346 100644 --- a/tests/test_cpp_api.py +++ b/tests/test_cpp_api.py @@ -182,7 +182,7 @@ def test_ascend_workspace_pool_retains_allocations_until_destruction(tmp_path): _run( [ - _compiler("CXX", "c++"), + _compiler("HOST_CXX", os.environ.get("CXX", "c++")), "-std=c++17", "-Werror", f"-I{tmp_path / 'include'}", From 99314502d6f4a0060226b663c5aea221e316f1a1 Mon Sep 17 00:00:00 2001 From: Li Baoming <1508269885@qq.com> Date: Thu, 29 Oct 2026 17:29:11 +0800 Subject: [PATCH 5/9] fix: initialize Moore compiler before CXX --- CMakeLists.txt | 9 +++--- tests/build_config/test_compilers.py | 46 ++++++++++++++++++---------- 2 files changed, 34 insertions(+), 21 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 7cbe93d18..bf0cb2f4e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -210,11 +210,15 @@ if(AUTO_DETECT_DEVICES) endif() endif() -# Select the MetaX compiler before CMake initializes CXX and its cache. +# Select backend compilers before CMake initializes CXX and its cache. if(WITH_METAX) set(CMAKE_C_COMPILER "${CMAKE_CURRENT_SOURCE_DIR}/scripts/mxcc_wrapper.sh") set(CMAKE_CXX_COMPILER "${CMAKE_CURRENT_SOURCE_DIR}/scripts/mxcc_wrapper.sh") endif() +if(WITH_MOORE) + set(CMAKE_C_COMPILER "${CMAKE_CURRENT_SOURCE_DIR}/scripts/mcc_wrapper.sh") + set(CMAKE_CXX_COMPILER "${CMAKE_CURRENT_SOURCE_DIR}/scripts/mcc_wrapper.sh") +endif() enable_language(CXX) if(BUILD_NVIDIA_FLASH_INFERENCE AND NOT WITH_NVIDIA) @@ -632,9 +636,6 @@ if(WITH_MOORE) message(STATUS "Using Moore from `${MUSA_ROOT}`.") - set(CMAKE_C_COMPILER ${CMAKE_CURRENT_SOURCE_DIR}/scripts/mcc_wrapper.sh) - set(CMAKE_CXX_COMPILER ${CMAKE_CURRENT_SOURCE_DIR}/scripts/mcc_wrapper.sh) - include_directories("${MUSA_ROOT}/include") link_directories("${MUSA_ROOT}/lib") diff --git a/tests/build_config/test_compilers.py b/tests/build_config/test_compilers.py index 8ab588ce7..a98c81ad3 100644 --- a/tests/build_config/test_compilers.py +++ b/tests/build_config/test_compilers.py @@ -31,10 +31,11 @@ def test_workspace_probe_uses_host_compiler(tmp_path, monkeypatch): module["test_ascend_workspace_pool_retains_allocations_until_destruction"](tmp_path) -@pytest.mark.skipif(os.name == "nt", reason="MetaX uses a POSIX compiler wrapper") +@pytest.mark.skipif(os.name == "nt", reason="GPU compiler wrappers use POSIX") +@pytest.mark.parametrize("backend", ("metax", "moore")) @pytest.mark.parametrize("auto_detect", (False, True)) -def test_metax_compiler_preserves_infinirt_configuration( - tmp_path, monkeypatch, auto_detect +def test_gpu_compiler_preserves_infinirt_configuration( + tmp_path, monkeypatch, auto_detect, backend ): cmake = shutil.which("cmake") compiler = shutil.which("g++") or shutil.which("clang++") @@ -42,21 +43,26 @@ def test_metax_compiler_preserves_infinirt_configuration( if not cmake or not compiler: pytest.skip("cmake and a C++ compiler are required") - toolkit = tmp_path / "maca" - mxcc = toolkit / "mxgpu_llvm/bin/mxcc" - mxcc.parent.mkdir(parents=True) - mxcc.write_text(f'#!/bin/sh\nexec {shlex.quote(compiler)} "$@"\n') - mxcc.chmod(0o755) + toolkit = tmp_path / "toolkit" + compiler_path = "mxgpu_llvm/bin/mxcc" if backend == "metax" else "bin/mcc" + gpu_compiler = toolkit / compiler_path + gpu_compiler.parent.mkdir(parents=True) + gpu_compiler.write_text(f'#!/bin/sh\nexec {shlex.quote(compiler)} "$@"\n') + gpu_compiler.chmod(0o755) libraries = toolkit / "lib" libraries.mkdir() - for name in ("mcruntime", "mcdnn", "mcblas"): + library_names = ( + ("mcruntime", "mcdnn", "mcblas") + if backend == "metax" + else ("musa", "musart", "mublas", "iomp5") + ) + for name in library_names: (libraries / f"lib{name}.so").touch() - monkeypatch.setenv("MACA_PATH", str(toolkit)) - for name in ( "INFINI_RT_ROOT", + "MACA_PATH", "HPCC_PATH", "NEUWARE_HOME", "MUSA_ROOT", @@ -66,9 +72,11 @@ def test_metax_compiler_preserves_infinirt_configuration( ): monkeypatch.delenv(name, raising=False) - wrapper = tmp_path / "scripts/mxcc_wrapper.sh" + monkeypatch.setenv("MACA_PATH" if backend == "metax" else "MUSA_HOME", str(toolkit)) + wrapper_name = "mxcc_wrapper.sh" if backend == "metax" else "mcc_wrapper.sh" + wrapper = tmp_path / "scripts" / wrapper_name wrapper.parent.mkdir() - shutil.copyfile(ROOT / "scripts/mxcc_wrapper.sh", wrapper) + shutil.copyfile(ROOT / "scripts" / wrapper_name, wrapper) wrapper.chmod(0o755) prefix = tmp_path / "infinirt" (prefix / "include/infini").mkdir(parents=True) @@ -78,14 +86,18 @@ def test_metax_compiler_preserves_infinirt_configuration( source = (ROOT / "CMakeLists.txt").read_text(encoding="utf-8") initialization = source.split("if(AUTO_DETECT_BACKENDS)", 1)[0] - metax_start = source.index("if(WITH_METAX)\n add_compile_definitions") - metax = source[metax_start : source.index("if(WITH_MARS)", metax_start)] + backend_option = f"WITH_{backend.upper()}" + backend_start = source.index(f"if({backend_option})\n add_compile_definitions") + next_backend = "WITH_MARS" if backend == "metax" else "WITH_CAMBRICON" + backend_config = source[ + backend_start : source.index(f"if({next_backend})", backend_start) + ] runtime_start = source.index('set(INFINI_RT_ROOT "$ENV{INFINI_RT_ROOT}"') runtime = source[ runtime_start : source.index("add_subdirectory(src)", runtime_start) ] (tmp_path / "CMakeLists.txt").write_text( - initialization + metax + runtime + "\nadd_subdirectory(dependency)\n" + initialization + backend_config + runtime + "\nadd_subdirectory(dependency)\n" 'file(WRITE "${CMAKE_BINARY_DIR}/selected-compiler.txt" ' '"${CMAKE_CXX_COMPILER}")\nadd_library(probe OBJECT probe.cc)\n' ) @@ -107,7 +119,7 @@ def test_metax_compiler_preserves_infinirt_configuration( ] if not auto_detect: - configure.append("-DWITH_METAX=ON") + configure.append(f"-D{backend_option}=ON") for _ in range(2): result = subprocess.run(configure, capture_output=True, text=True, timeout=45) From 8304829b44c54f6fbd0ed38392560fe38b921757 Mon Sep 17 00:00:00 2001 From: Li Baoming <1508269885@qq.com> Date: Thu, 29 Oct 2026 18:16:37 +0800 Subject: [PATCH 6/9] fix: align Moore CI runtime and reference precision --- .github/ci_config.yml | 6 +++--- tests/build_config/test_ci_build.py | 29 +++++++++++++++++++++++++++++ tests/test_gemm.py | 9 +++++++++ 3 files changed, 41 insertions(+), 3 deletions(-) diff --git a/.github/ci_config.yml b/.github/ci_config.yml index 3149e6163..e9ea187fc 100644 --- a/.github/ci_config.yml +++ b/.github/ci_config.yml @@ -104,7 +104,7 @@ platforms: image: dockerfile: images/moore/ build_args: - BASE_IMAGE: sh-harbor.mthreads.com/mcctest/vllm_musa:20251112_hygon + BASE_IMAGE: registry.mthreads.com/mcconline/inference/vllm:4.3.6_torch2.7.1_20260323_ubuntu APT_MIRROR: http://archive.ubuntu.com/ubuntu PIP_INDEX_URL: https://pypi.org/simple docker_args: @@ -116,7 +116,7 @@ platforms: resources: gpu_ids: "1" ngpus: 1 - gpu_style: none # Moore uses passthrough via `--privileged`; `MTHREADS_VISIBLE_DEVICES` controls device selection. + gpu_style: none # `MTHREADS_VISIBLE_DEVICES` loads drivers; PyTorch also needs `MUSA_VISIBLE_DEVICES`. memory: 32GB shm_size: 16g timeout: 3600 @@ -124,7 +124,7 @@ platforms: junit_path: test-results.xml stages: - name: test - run: pytest tests/ -m smoke -n 4 -v --tb=short --junitxml=/workspace/results/test-results.xml + run: MUSA_VISIBLE_DEVICES="${MTHREADS_VISIBLE_DEVICES:?CI must select a Moore GPU}" pytest tests/ -m smoke -n 4 -v --tb=short --junitxml=/workspace/results/test-results.xml cambricon: runner_label: Cambricon diff --git a/tests/build_config/test_ci_build.py b/tests/build_config/test_ci_build.py index 483575dcc..3a1cfbce4 100644 --- a/tests/build_config/test_ci_build.py +++ b/tests/build_config/test_ci_build.py @@ -45,6 +45,35 @@ def test_ci_installs_declared_build_requirements(platform): assert tokens[:4] == ["python", "-m", "pip", "install"] +@pytest.mark.skipif(os.name == "nt", reason="CI stages use a POSIX shell") +@pytest.mark.parametrize("allocated_devices", ("0", "2,5", "", None)) +def test_moore_smoke_uses_allocated_devices(tmp_path, monkeypatch, allocated_devices): + config = yaml.safe_load((ROOT / ".github/ci_config.yml").read_text()) + stage = config["platforms"]["moore"]["jobs"]["gpu"]["stages"][0]["run"] + probe = tmp_path / "pytest" + probe.write_text('#!/bin/sh\nprintf "%s" "$MUSA_VISIBLE_DEVICES"\n') + probe.chmod(0o755) + monkeypatch.setenv("PATH", str(tmp_path) + os.pathsep + os.environ["PATH"]) + monkeypatch.setenv("MUSA_VISIBLE_DEVICES", "7") + + if allocated_devices is None: + monkeypatch.delenv("MTHREADS_VISIBLE_DEVICES", raising=False) + else: + monkeypatch.setenv("MTHREADS_VISIBLE_DEVICES", allocated_devices) + + result = subprocess.run( + ["/bin/sh", "-c", stage], capture_output=True, text=True, timeout=10 + ) + + if allocated_devices: + assert result.returncode == 0, result.stderr + assert result.stdout == allocated_devices + else: + assert result.returncode != 0 + assert not result.stdout + assert "CI must select a Moore GPU" in result.stderr + + def test_wheel_metadata_accepts_sdist_configuration(tmp_path, monkeypatch): from scikit_build_core.build import prepare_metadata_for_build_wheel diff --git a/tests/test_gemm.py b/tests/test_gemm.py index a79fc5678..289f1d92d 100644 --- a/tests/test_gemm.py +++ b/tests/test_gemm.py @@ -128,6 +128,15 @@ def _gemm(a, b, c, alpha, beta, trans_a, trans_b, y, implementation_index=0): def _torch_gemm(a, b, c, alpha, beta, trans_a, trans_b, y): + if a.device.type == "musa" and a.dtype == torch.float32: + # Match native muBLAS FP32 arithmetic when computing the reference. + with torch.backends.mudnn.flags(allow_tf32=False): + return _torch_gemm_reference(a, b, c, alpha, beta, trans_a, trans_b, y) + + return _torch_gemm_reference(a, b, c, alpha, beta, trans_a, trans_b, y) + + +def _torch_gemm_reference(a, b, c, alpha, beta, trans_a, trans_b, y): if trans_a: a = a.transpose(-2, -1) From 83bfacb44e5a0149eae461ddca77cf8b088e2196 Mon Sep 17 00:00:00 2001 From: Li Baoming <1508269885@qq.com> Date: Fri, 30 Oct 2026 14:55:38 +0800 Subject: [PATCH 7/9] fix: limit NVCC argument workaround to provider dispatch --- src/operator.h | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/src/operator.h b/src/operator.h index 26b47427e..9e461f337 100644 --- a/src/operator.h +++ b/src/operator.h @@ -710,12 +710,13 @@ class Operator : public OperatorBase { DispatchFunc>( dispatch_device_type, - [&](auto device_tag, auto&... device_args) { + [&](auto device_tag) { constexpr Device::Type kDev = decltype(device_tag)::value; using Implementations = typename ActiveImplementations::type; const bool automatic = config.implementation_index() == Config::kAutomaticImplementation; + // Avoid NVCC's nested-lambda capture error for the argument pack. const auto construct = [&](auto implementation_tag, auto&... provider_args) { constexpr std::size_t kImplementationIndex = @@ -761,7 +762,7 @@ class Operator : public OperatorBase { for (const auto index : detail::ListToVector(Implementations{})) { detail::DispatchImplementation(index, construct, "Operator::Make(automatic)", - Implementations{}, device_args...); + Implementations{}, args...); if (matched) break; } } else { @@ -773,10 +774,10 @@ class Operator : public OperatorBase { detail::DispatchImplementation( config.implementation_index(), construct, "Operator::Make(implementation_index)", Implementations{}, - device_args...); + args...); } }, - "Operator::Make", args...); + "Operator::Make"); return matched; } From 345d092c24c8d4240c46d7bde6928d10ecb03529 Mon Sep 17 00:00:00 2001 From: Li Baoming <1508269885@qq.com> Date: Fri, 30 Oct 2026 15:25:11 +0800 Subject: [PATCH 8/9] fix: reuse master operator dispatch workaround --- src/operator.h | 23 ++++++++++------------- 1 file changed, 10 insertions(+), 13 deletions(-) diff --git a/src/operator.h b/src/operator.h index 9e461f337..4b19415ed 100644 --- a/src/operator.h +++ b/src/operator.h @@ -698,7 +698,7 @@ class Operator : public OperatorBase { template static bool SelectWithDevice(const Config& config, Device::Type dispatch_device_type, bool strict, - Selected&& selected, Args&&... args) { + Selected&& selected, Args&&... call_args) { bool matched = false; if (!detail::ListContains(dispatch_device_type, ActiveDevices{})) { @@ -710,15 +710,13 @@ class Operator : public OperatorBase { DispatchFunc>( dispatch_device_type, - [&](auto device_tag) { + [&](auto device_tag, auto&&... dispatched_args) { constexpr Device::Type kDev = decltype(device_tag)::value; using Implementations = typename ActiveImplementations::type; const bool automatic = config.implementation_index() == Config::kAutomaticImplementation; - // Avoid NVCC's nested-lambda capture error for the argument pack. - const auto construct = [&](auto implementation_tag, - auto&... provider_args) { + const auto construct = [&](auto implementation_tag) { constexpr std::size_t kImplementationIndex = decltype(implementation_tag)::value; using Provider = Operator; @@ -730,11 +728,11 @@ class Operator : public OperatorBase { std::is_constructible_v) { bool supported = false; if constexpr (kConfigAware) { - supported = - Provider::supports(config, std::as_const(provider_args)...); + supported = Provider::supports( + config, std::as_const(dispatched_args)...); } else if (!config.requires_config_aware_provider()) { supported = detail::ProviderSupports( - 0, std::as_const(provider_args)...); + 0, std::as_const(dispatched_args)...); } if (!supported) { if (automatic) return; @@ -748,7 +746,7 @@ class Operator : public OperatorBase { matched = true; selected(device_tag, implementation_tag, std::bool_constant{}, - provider_args...); + dispatched_args...); } else { if (!automatic && strict) { throw std::invalid_argument( @@ -762,7 +760,7 @@ class Operator : public OperatorBase { for (const auto index : detail::ListToVector(Implementations{})) { detail::DispatchImplementation(index, construct, "Operator::Make(automatic)", - Implementations{}, args...); + Implementations{}); if (matched) break; } } else { @@ -773,11 +771,10 @@ class Operator : public OperatorBase { } detail::DispatchImplementation( config.implementation_index(), construct, - "Operator::Make(implementation_index)", Implementations{}, - args...); + "Operator::Make(implementation_index)", Implementations{}); } }, - "Operator::Make"); + "Operator::Make", call_args...); return matched; } From 3d81a4134223efd123d7df0c3202fd6ea4287eac Mon Sep 17 00:00:00 2001 From: Li Baoming <1508269885@qq.com> Date: Fri, 30 Oct 2026 18:17:13 +0800 Subject: [PATCH 9/9] fix: allocate idle Ascend devices within CI queue budget --- .github/ci_config.yml | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/.github/ci_config.yml b/.github/ci_config.yml index 8a758690c..710053fdf 100644 --- a/.github/ci_config.yml +++ b/.github/ci_config.yml @@ -212,13 +212,13 @@ platforms: npu: type: unittest resources: - gpu_ids: "1" + gpu_ids: auto ngpus: 1 gpu_style: none memory: 32GB shm_size: 16g timeout: 3600 - queue_timeout: 3600 + queue_timeout: 600 junit_path: test-results.xml stages: - name: test