Skip to content

Build cuda.core RDC test fixtures without Bash #2335

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
rwgk wants to merge 3 commits into
base: main
Choose a base branch
from
+128 −43
Open

Build cuda.core RDC test fixtures without Bash #2335

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
447ab4b
Run nvcc fixture build without Bash
rwgk Jul 9, 2026
e449386
Document narrow nvcc skip detection
rwgk Jul 9, 2026
d7bdbd7
Merge branch 'main' into run_nvcc_without_bash
rwgk Jul 9, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 3 additions & 1 deletion .github/workflows/build-wheel.yml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -463,7 +463,9 @@ jobs:
cuda-path: "./cuda_toolkit_prev"

- name: Build cuda.core test binaries
run: bash ${{ env.CUDA_CORE_TEST_BINARIES_DIR }}/build_test_binaries.sh
run: |
nvcc --version
python "${{ env.CUDA_CORE_TEST_BINARIES_DIR }}/build_test_binaries.py"

- name: Upload cuda.core test binaries
uses: actions/upload-artifact@043fb46d1a93c77aae656e7c1c64a875d1fc6a0a # v7.0.1
Expand Down
56 changes: 56 additions & 0 deletions cuda_core/tests/test_binaries/build_test_binaries.py
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,56 @@
# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
"""Build the relocatable-device-code fixtures used by cuda.core tests."""

from __future__ import annotations

import os
import subprocess
import tempfile
from pathlib import Path


def _run(command: list[str]) -> None:
print(f"+ {subprocess.list2cmdline(command)}")
result = subprocess.run(command) # noqa: S603
if result.returncode != 0:
raise SystemExit(result.returncode)


def main() -> None:
script_dir = Path(__file__).resolve().parent
source_path = script_dir / "saxpy.cu"
final_object_path = script_dir / "saxpy.o"
final_library_path = script_dir / ("saxpy.lib" if os.name == "nt" else "saxpy.a")

nvcc_extra_flags = ["-std=c++17"]
if os.name == "nt":
nvcc_extra_flags.extend(["-Xcompiler", "/Zc:preprocessor"])

with tempfile.TemporaryDirectory(prefix="build_test_binaries-", dir=script_dir) as temp_dir:
temp_dir_path = Path(temp_dir)
object_path = temp_dir_path / final_object_path.name
library_path = temp_dir_path / final_library_path.name

_run(
[
"nvcc",

@lijinf2 lijinf2 Jul 9, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Your previous PR uses NVCC environmental variable and pathfinder instead of bare nvcc command. Do we consider that here?

@rwgk rwgk Jul 10, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, this was considered. Now that the builder is Python instead of Bash, a full nvcc.exe path would be much less problematic than it was in the earlier Bash-based version.

I still prefer bare nvcc here because it keeps the test contract simple and explicit: the test environment should put the intended CUDA toolkit's bin directory on PATH, and the host compiler environment should already be configured. That is easy to satisfy in CI, runbooks, and local development.

Using pathfinder or an NVCC override inside this unit test would add a second toolkit-selection mechanism. That can create mismatches where the test discovers one nvcc, while CUDA_PATH, PATH, runtime libraries, or the surrounding runbook point at another toolkit. For this fixture build, relying on the environment's normal executable lookup is simpler and easier to reason about.

"-dc",
*nvcc_extra_flags,
"-arch=all-major",
"-o",
str(object_path),
str(source_path),
]
)
_run(["nvcc", "-lib", "-o", str(library_path), str(object_path)])

object_path.replace(final_object_path)
library_path.replace(final_library_path)

for path in (final_object_path, final_library_path):
print(f"{path}: {path.stat().st_size} bytes")


if __name__ == "__main__":
main()
28 changes: 0 additions & 28 deletions cuda_core/tests/test_binaries/build_test_binaries.sh
View file
Open in desktop

This file was deleted.

83 changes: 69 additions & 14 deletions cuda_core/tests/test_module.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,6 +5,7 @@
import os
import pickle
import subprocess
import sys
import warnings
from pathlib import Path

Expand All @@ -16,7 +17,6 @@
from cuda.core._program import _can_load_generated_ptx
from cuda.core._utils.cuda_utils import CUDAError, driver, handle_return
from cuda.core._utils.version import binding_version, driver_version
from cuda.pathfinder import find_nvidia_binary_utility

try:
import numba
Expand Down Expand Up @@ -193,22 +193,77 @@ def _read_saxpy_rdc(kind: str) -> bytes:
raise ValueError(f"unknown saxpy RDC kind: {kind!r}")

if not rdc_path.is_file():
nvcc_path = find_nvidia_binary_utility("nvcc")
if nvcc_path is None:
pytest.skip(
f"{rdc_path.name} not found at {rdc_path} and nvcc is unavailable. "
"In CI this is downloaded from the build stage."
)
env = os.environ.copy()
env["NVCC"] = nvcc_path
subprocess.run( # noqa: S603
["bash", str(binaries_dir / "build_test_binaries.sh")], # noqa: S607
check=True,
env=env,
)
_build_saxpy_rdc(binaries_dir)
return rdc_path.read_bytes()


def _subprocess_output(result: subprocess.CompletedProcess[str]) -> str:

@lijinf2 lijinf2 Jul 9, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Could we delete this function _subprocess_output and just default to system stdout/stderr behaviors?

@rwgk rwgk Jul 10, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I would prefer to keep this helper. We intentionally capture stdout/stderr here because we need to inspect the nvcc output to distinguish known host-compiler setup failures from unexpected nvcc failures.

For unexpected failures, pytest.fail(...) includes the captured stdout/stderr in one self-contained failure message. That makes the CI/test log easier to interpret, especially when pytest output capture or parallel test output would otherwise split the context.

On success we still print the captured output back to stdout/stderr, so the normal diagnostic trail is preserved.

sections = []
if result.stdout:
sections.append(f"stdout:\n{result.stdout.rstrip()}")
if result.stderr:
sections.append(f"stderr:\n{result.stderr.rstrip()}")
return "\n".join(sections) or "<no output>"


def _host_compiler_is_unavailable(output: str) -> bool:
normalized = output.lower()
# Keep these patterns narrow. Unknown nvcc failures should fail the test and
# expose their diagnostics, not be silently reclassified as environment skips.
windows_compiler_missing = (
"cannot find compiler" in normalized and "cl.exe" in normalized and "in path" in normalized

@lijinf2 lijinf2 Jul 9, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is "cannot find compiler" keyword stable enough and will not be modified in future release of nvcc logging?

I got the following from agent and it works on linux but not sure on windows.

    for line in p.stdout.splitlines():
        if line.startswith("#$"):
            cmd = line[3:].strip().split()
            if cmd:
                return cmd[0]  # gcc / g++ / clang++ / cl.exe

@rwgk rwgk Jul 10, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Good question. I do not want to rely on this message being stable forever. The intent is almost the opposite: keep the recognized skip cases narrow.

If nvcc changes this diagnostic in a future release, this helper will stop recognizing it as a known environment setup issue. The test will then fail and include the full nvcc output, which gives us the evidence needed to decide whether to add another narrowly scoped skip pattern.

The #$ ... parsing approach can identify commands emitted by nvcc -v/dry-run style output, but it does not directly answer the question we need here: did the actual fixture build fail because the host compiler environment is missing, or did it fail for some other reason? For this test, it is safer to classify only the exact known failure modes after the real build attempt, and let everything else fail visibly.

)

@lijinf2 lijinf2 Jul 9, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

On windows is "cl.exe" the only host compiler available? What if nvcc is using another host compiler?

@rwgk rwgk Jul 10, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

For the path this PR exercises, yes, the Windows host compiler we expect is MSVC cl.exe. The fixture builder invokes bare nvcc and does not pass --compiler-bindir / -ccbin, so it is intentionally using the normal Windows nvcc setup.

If someone has a custom Windows nvcc host-compiler setup and it fails with a different diagnostic, this code will not silently skip it. It will fail the test and expose the actual output. That is intentional: the skip classification is only for the known default Windows setup failure where nvcc reports that it cannot find cl.exe in PATH.

In other words, this PR does intentionally not try to generalize all possible Windows host compiler configurations. It makes the supported test contract (see ctk-next PR 436)explicit: nvcc on PATH, with the normal host compiler environment already configured.

linux_compiler_missing = (
"no such file or directory" in normalized and "failed to preprocess host compiler properties" in normalized
)
return windows_compiler_missing or linux_compiler_missing


def _build_saxpy_rdc(binaries_dir: Path) -> None:
try:
version_result = subprocess.run(
["nvcc", "--version"], # noqa: S607 - PATH lookup is the behavior under test.
capture_output=True,
text=True,
errors="replace",
)
except FileNotFoundError:
pytest.skip("RDC test fixtures are absent and nvcc is not available on PATH")

if version_result.returncode != 0:
pytest.fail(
f"nvcc --version failed with exit code {version_result.returncode}\n{_subprocess_output(version_result)}",
pytrace=False,
)

print(version_result.stdout, end="")
if version_result.stderr:
print(version_result.stderr, end="", file=sys.stderr)

builder_path = binaries_dir / "build_test_binaries.py"
build_result = subprocess.run( # noqa: S603
[sys.executable, str(builder_path)],
capture_output=True,
text=True,
errors="replace",
)
if build_result.returncode == 0:
print(build_result.stdout, end="")
if build_result.stderr:
print(build_result.stderr, end="", file=sys.stderr)
return

output = _subprocess_output(build_result)
if _host_compiler_is_unavailable(output):
print(output, file=sys.stderr)
pytest.skip("nvcc is available, but its host compiler is not configured")

pytest.fail(
f"RDC test fixture build failed with exit code {build_result.returncode}\n{output}",
pytrace=False,
)


def test_get_kernel(init_cuda):
kernel = """extern "C" __global__ void ABC() { }"""

Expand Down
Loading