From 6b883259d9bc19a901515deb16459234f672fed6 Mon Sep 17 00:00:00 2001 From: Cloud Han Date: Tue, 25 Aug 2026 07:40:50 +0000 Subject: [PATCH] fix: version CUDA toolchain implementations --- MODULE.bazel | 4 +- cuda/extensions.bzl | 47 +++++++- cuda/private/repositories.bzl | 34 +++++- cuda/private/template_helper.bzl | 114 ++++++++++++++++-- cuda/private/templates/BUILD.cuda_shared | 2 +- .../templates/BUILD.nvcc_nvvm_embedded | 27 +++++ cuda/private/templates/registry.bzl | 2 +- cuda/private/toolchain.bzl | 4 +- tests/integration/test_all.sh | 73 +++++++++-- .../BUILD.bazel | 15 +++ .../MODULE.bazel | 29 +++++ 11 files changed, 317 insertions(+), 34 deletions(-) create mode 100644 cuda/private/templates/BUILD.nvcc_nvvm_embedded create mode 100644 tests/integration/toolchain_redist_json_cross_major/BUILD.bazel create mode 100644 tests/integration/toolchain_redist_json_cross_major/MODULE.bazel diff --git a/MODULE.bazel b/MODULE.bazel index e38f9a53..4913dfb7 100644 --- a/MODULE.bazel +++ b/MODULE.bazel @@ -16,9 +16,7 @@ cuda.toolkit( use_repo(cuda, "cuda") register_toolchains( - "@cuda//toolchain:nvcc-local-toolchain", - "@cuda//toolchain/clang:clang-local-toolchain", - "@cuda//toolchain/disabled:disabled-local-toolchain", + "@cuda//toolchain/...", ) bazel_dep(name = "rules_cuda_examples", dev_dependency = True) diff --git a/cuda/extensions.bzl b/cuda/extensions.bzl index 9a4d1984..a2eab085 100644 --- a/cuda/extensions.bzl +++ b/cuda/extensions.bzl @@ -35,6 +35,7 @@ cuda_component_tag = tag_class(attrs = { "If all downloads fail, the rule will fail.", ), "version": attr.string(doc = "A unique version number for component."), + "toolkit_version": attr.string(doc = "CUDA Toolkit release containing this component."), }) cuda_redist_json_tag = tag_class(attrs = { @@ -134,6 +135,7 @@ def _register_redist_components(module_ctx, attr, component_entries): component_attr = {key: value for key, value in spec.items()} component_repo_name = repo_name + "_" + platform.replace("-", "_") + "_" + redist_ver.replace(".", "_") component_attr["name"] = component_repo_name + component_attr["toolkit_version"] = redist_ver dedupe_key = (spec["component_name"], platform, redist_ver) existing_entry = component_entries.get(dedupe_key) @@ -250,7 +252,49 @@ def _impl(module_ctx): if components_mapping != None: sorted_redist_versions = sorted(redist_versions, key = _version_sort_key) - # Always use the maximum version so the toolkit includes all components. + # Generate a static implementation repository for each toolkit version. Each + # repository receives component aliases fixed to that version, so load-time + # BUILD decisions and compiler-file ownership cannot leak from another CUDA. + toolchain_repositories = {} + for redist_version in sorted_redist_versions: + exact_components_mapping = {} + version_label = redist_version.replace(".", "_").replace("-", "_") + for component_name in redist_components_mapping.keys(): + component_platforms = [ + platform + for platform in SUPPORTED_PLATFORMS + if platform in versioned_repos[component_name] and redist_version in versioned_repos[component_name][platform] + ] + if not component_platforms: + continue + + if len(component_platforms) == 1: + exact_components_mapping[component_name] = "@" + versioned_repos[component_name][component_platforms[0]][redist_version] + else: + alias_name = "{}_{}_toolchain_{}".format(toolkit.name, component_name, version_label) + platform_repo_kwargs = {} + for platform in SUPPORTED_PLATFORMS: + repos = {} + if platform in versioned_repos[component_name] and redist_version in versioned_repos[component_name][platform]: + repos[redist_version] = versioned_repos[component_name][platform][redist_version] + platform_repo_kwargs[_platform_repos_attr(platform)] = repos + platform_alias_repo( + name = alias_name, + component_name = component_name, + versions = [redist_version], + **platform_repo_kwargs + ) + exact_components_mapping[component_name] = "@" + alias_name + + toolchain_repo_name = "{}_toolchain_{}".format(toolkit.name, version_label) + cuda_toolkit( + name = toolchain_repo_name, + components_mapping = exact_components_mapping, + version = redist_version, + ) + toolchain_repositories[redist_version] = "@" + toolchain_repo_name + + # Always use the maximum version so the public toolkit facade includes all components. # Components that don't exist in older versions will fall back to dummy. toolkit_version = sorted_redist_versions[-1] @@ -263,6 +307,7 @@ def _impl(module_ctx): components_mapping = components_mapping, version = toolkit_version, toolkit_versions = sorted_redist_versions, + toolchain_repositories = toolchain_repositories, ) else: cuda_toolkit(**_module_tag_to_dict(toolkit)) diff --git a/cuda/private/repositories.bzl b/cuda/private/repositories.bzl index 3787f355..9b80297a 100644 --- a/cuda/private/repositories.bzl +++ b/cuda/private/repositories.bzl @@ -177,7 +177,10 @@ def _detect_deliverable_cuda_toolkit(repository_ctx): cicc = None libdevice = None - if "nvvm" in repository_ctx.attr.components_mapping: + if int(cuda_version_major) < 13: + cicc = "{}//:cicc".format(nvcc_repo) + libdevice = "{}//:libdevice.10.bc".format(nvcc_repo) + elif "nvvm" in repository_ctx.attr.components_mapping: nvvm_repo = repository_ctx.attr.components_mapping["nvvm"] cicc = "{}//:cicc".format(nvvm_repo) libdevice = "{}//:libdevice.10.bc".format(nvvm_repo) @@ -263,8 +266,12 @@ def config_cuda_toolkit_and_nvcc(repository_ctx, cuda): # Generate @cuda//defs.bzl template_helper.generate_defs_bzl(repository_ctx, cuda.version_major, cuda.version_minor, is_local_ctk == True) - # Generate @cuda//toolchain/BUILD - template_helper.generate_toolchain_build(repository_ctx, cuda) + # Generate toolchain implementations, or the stable @cuda facade that points at + # independently generated implementations for each configured toolkit version. + if repository_ctx.attr.toolchain_repositories: + template_helper.generate_toolchain_facade(repository_ctx) + else: + template_helper.generate_toolchain_build(repository_ctx, cuda) def detect_clang(repository_ctx): """Detect local clang installation. @@ -322,8 +329,9 @@ def config_clang(repository_ctx, cuda, clang_path_or_label): if len(repository_ctx.attr.components_mapping) != 0: is_local_ctk = False - # Generate @cuda//toolchain/clang/BUILD - template_helper.generate_toolchain_clang_build(repository_ctx, cuda, clang_path_or_label) + # The versioned facade generated above includes its own clang declarations. + if not repository_ctx.attr.toolchain_repositories: + template_helper.generate_toolchain_clang_build(repository_ctx, cuda, clang_path_or_label) def config_disabled(repository_ctx): repository_ctx.symlink(Label("//cuda/private:templates/BUILD.toolchain_disabled"), "toolchain/disabled/BUILD") @@ -355,6 +363,9 @@ cuda_toolkit = repository_rule( "generated toolchain selects among them on @rules_cuda//cuda:version instead of " + "hardcoding `version`.", ), + "toolchain_repositories": attr.string_dict( + doc = "Internal mapping from exact CUDA versions to repositories containing their toolchain implementations.", + ), }, configure = True, local = True, @@ -445,14 +456,24 @@ def _cuda_component_impl(repository_ctx): _patch_nvcc_profile_post(repository_ctx, patch_nvcc_profile) _patch_nvvm(repository_ctx, component_name) + toolkit_version = repository_ctx.attr.toolkit_version or repository_ctx.attr.version + defs_label = "//:defs.bzl" if toolkit_version else "@cuda//:defs.bzl" + template_helper.generate_build( repository_ctx, libpath = "lib", components = {component_name: repository_ctx.name}, + defs_label = defs_label, is_cuda_repo = False, is_deliverable = True, ) + if toolkit_version: + version_parts = toolkit_version.split(".") + if len(version_parts) < 2: + fail("toolkit_version must contain at least a major and minor version") + template_helper.generate_defs_bzl(repository_ctx, version_parts[0], version_parts[1], False) + desc_name = repository_ctx.attr.descriptive_name or repository_ctx.attr.component_name repository_ctx.file( "{}/version.json".format(component_name), @@ -509,6 +530,9 @@ cuda_component = repository_rule( "If all downloads fail, the rule will fail.", ), "version": attr.string(doc = "A unique version number for component. Store in version.json file"), + "toolkit_version": attr.string( + doc = "CUDA Toolkit release containing this component. Used for version-gated BUILD evaluation.", + ), }, ) diff --git a/cuda/private/template_helper.bzl b/cuda/private/template_helper.bzl index eb962aaf..ffc3c0bf 100644 --- a/cuda/private/template_helper.bzl +++ b/cuda/private/template_helper.bzl @@ -159,9 +159,11 @@ def _expand_dctk_component(repository_ctx, component): def _component_owns_cuda_repo_alias(component, target, components): if target == "culibos_a" and "culibos" in components: return component == "culibos" + if target in ["cicc", "libdevice", "libdevice.10.bc"] and "nvvm" in components: + return component == "nvvm" return True -def _generate_build_impl(repository_ctx, libpath, components, is_cuda_repo, is_deliverable): +def _generate_build_impl(repository_ctx, libpath, components, is_cuda_repo, is_deliverable, defs_label): # stitch template fragment fragments = [ Label("//cuda/private:templates/BUILD.cuda_shared"), @@ -178,6 +180,8 @@ def _generate_build_impl(repository_ctx, libpath, components, is_cuda_repo, is_d comp = components.keys()[0] fragments.append(_expand_dctk_component(repository_ctx, comp)) fragments.append(Label("//cuda/private:templates/BUILD.{}".format(comp))) + if comp == "nvcc": + fragments.append(Label("//cuda/private:templates/BUILD.nvcc_nvvm_embedded")) else: fail("unreachable") @@ -208,11 +212,12 @@ def _generate_build_impl(repository_ctx, libpath, components, is_cuda_repo, is_d substitutions = { "%{component_name}": "cuda" if is_cuda_repo else components.keys()[0], + "%{defs_label}": defs_label, "%{libpath}": libpath, } repository_ctx.template("BUILD", template_path, substitutions = substitutions, executable = False) -def _generate_build(repository_ctx, libpath, components = None, is_cuda_repo = True, is_deliverable = False): +def _generate_build(repository_ctx, libpath, components = None, is_cuda_repo = True, is_deliverable = False, defs_label = "//:defs.bzl"): """Generate `@cuda//BUILD` or `@cuda_//BUILD` Notes: @@ -237,7 +242,7 @@ def _generate_build(repository_ctx, libpath, components = None, is_cuda_repo = T if c not in REGISTRY: fail("{} is not a valid component") - _generate_build_impl(repository_ctx, libpath, components, is_cuda_repo, is_deliverable) + _generate_build_impl(repository_ctx, libpath, components, is_cuda_repo, is_deliverable, defs_label) def _generate_defs_bzl(repository_ctx, version_major, version_minor, is_local_ctk): tpl_label = Label("//cuda/private:templates/defs.bzl.tpl") @@ -268,6 +273,7 @@ def _generate_redist_bzl(repository_ctx, component_specs, redist_version): strip_prefix = {strip_prefix}, urls = {urls}, version = "{version}", + toolkit_version = "{toolkit_version}", )""" for spec in component_specs: @@ -282,6 +288,7 @@ def _generate_redist_bzl(repository_ctx, component_specs, redist_version): strip_prefix = repr(spec["strip_prefix"]), urls = repr(spec["urls"]), version = spec["version"], + toolkit_version = redist_version, ), ) mapping[spec["component_name"]] = "@" + repo_name @@ -299,7 +306,7 @@ def _generate_toolchain_build(repository_ctx, cuda): "//cuda/private:templates/BUILD.toolchain_" + ("nvcc" if _is_linux(repository_ctx) else "nvcc_msvc"), ) - compiler_files = ["@cuda//:compiler_deps"] + compiler_files = ["//:compiler_deps"] if cuda.cicc_label != None: compiler_files.append(cuda.cicc_label) if cuda.libdevice_label != None: @@ -375,20 +382,20 @@ def _generate_toolchain_clang_build(repository_ctx, cuda, clang_path_or_label): cuda_path_for_subst = "" path_data = None if cuda.path: # local installation - compiler_files.append("@cuda//:compiler_deps") + compiler_files.append("//:compiler_deps") cuda_path_for_subst = _to_forward_slash(cuda.path) else: # scattered components - cuda_path_for_subst = "$(location @cuda//:compiler_root)" - path_data = ["@cuda//:compiler_root"] + cuda_path_for_subst = "$(location //:compiler_root)" + path_data = ["//:compiler_root"] compiler_files.extend([ - "@cuda//:nvcc_all_files", - "@cuda//:cccl_all_files", - "@cuda//:cudart_all_files", - "@cuda//:curand_all_files", + "//:nvcc_all_files", + "//:cccl_all_files", + "//:cudart_all_files", + "//:curand_all_files", ]) if int(cuda.version_major) >= 13: compiler_files.extend([ - "@cuda//:nvvm_all_files", + "//:nvvm_all_files", ]) path_data_line = "path_data = " + repr(path_data) + "," compiler_files_line = "compiler_files = " + repr(compiler_files) + "," @@ -434,10 +441,93 @@ def _generate_toolchain_clang_build(repository_ctx, cuda, clang_path_or_label): executable = False, ) +def _toolchain_declaration(name, implementation, compiler_setting, version_setting, exec_constraints = []): + lines = [ + "toolchain(", + ' name = "{}",'.format(name), + ] + if exec_constraints: + lines.extend([ + " exec_compatible_with = {},".format(repr(exec_constraints)), + " target_compatible_with = {},".format(repr(exec_constraints)), + ]) + lines.extend([ + " target_settings = [", + ' "@rules_cuda//cuda:is_enabled",', + ' "{}",'.format(compiler_setting), + ' ":{}",'.format(version_setting), + " ],", + ' toolchain = "{}",'.format(implementation), + ' toolchain_type = "@rules_cuda//cuda:toolchain_type",', + ' visibility = ["//visibility:public"],', + ")", + ]) + return "\n".join(lines) + +def _generate_toolchain_facade(repository_ctx): + versions = sorted(repository_ctx.attr.toolchain_repositories.keys()) + default_version = repository_ctx.attr.version + os_constraint = "@platforms//os:windows" if _is_windows(repository_ctx) else "@platforms//os:linux" + + nvcc = [] + clang = [] + for version in versions: + label = _version_label(version) + setting = "toolkit_version_is_{}".format(label) + definition = [ + "config_setting(", + ' name = "{}",'.format(setting), + ' flag_values = {{"@rules_cuda//cuda:version": "{}"}},'.format(version), + ")", + ] + nvcc.extend(definition) + clang.extend(definition) + repo = repository_ctx.attr.toolchain_repositories[version] + nvcc.append(_toolchain_declaration( + "nvcc-{}-toolchain".format(label), + repo + "//toolchain:nvcc-local", + "@rules_cuda//cuda:compiler_is_nvcc", + setting, + [os_constraint], + )) + clang.append(_toolchain_declaration( + "clang-{}-toolchain".format(label), + repo + "//toolchain/clang:clang-local", + "@rules_cuda//cuda:compiler_is_clang", + setting, + )) + + default_setting = [ + "config_setting(", + ' name = "toolkit_version_is_default",', + ' flag_values = {"@rules_cuda//cuda:version": ""},', + ")", + ] + nvcc.extend(default_setting) + clang.extend(default_setting) + default_repo = repository_ctx.attr.toolchain_repositories[default_version] + nvcc.append(_toolchain_declaration( + "nvcc-local-toolchain", + default_repo + "//toolchain:nvcc-local", + "@rules_cuda//cuda:compiler_is_nvcc", + "toolkit_version_is_default", + [os_constraint], + )) + clang.append(_toolchain_declaration( + "clang-local-toolchain", + default_repo + "//toolchain/clang:clang-local", + "@rules_cuda//cuda:compiler_is_clang", + "toolkit_version_is_default", + )) + + repository_ctx.file("toolchain/BUILD", "\n\n".join(nvcc) + "\n") + repository_ctx.file("toolchain/clang/BUILD", "\n\n".join(clang) + "\n") + template_helper = struct( generate_build = _generate_build, generate_defs_bzl = _generate_defs_bzl, generate_redist_bzl = _generate_redist_bzl, generate_toolchain_build = _generate_toolchain_build, generate_toolchain_clang_build = _generate_toolchain_clang_build, + generate_toolchain_facade = _generate_toolchain_facade, ) diff --git a/cuda/private/templates/BUILD.cuda_shared b/cuda/private/templates/BUILD.cuda_shared index c09b3727..bb035d57 100644 --- a/cuda/private/templates/BUILD.cuda_shared +++ b/cuda/private/templates/BUILD.cuda_shared @@ -1,7 +1,7 @@ load("@bazel_skylib//rules:common_settings.bzl", "bool_setting") # @unused -load("@cuda//:defs.bzl", "additional_header_deps", "any_or_first", "if_cuda_toolkit_version_ge", "if_local_cuda_toolkit") # @unused load("@rules_cc//cc:defs.bzl", "cc_binary", "cc_import", "cc_library") load("@rules_cuda//cuda:defs.bzl", "cc_import_versioned_sos", "if_linux", "if_windows") # @unused +load("%{defs_label}", "additional_header_deps", "any_or_first", "if_cuda_toolkit_version_ge", "if_local_cuda_toolkit") # @unused package( default_visibility = ["//visibility:public"], diff --git a/cuda/private/templates/BUILD.nvcc_nvvm_embedded b/cuda/private/templates/BUILD.nvcc_nvvm_embedded new file mode 100644 index 00000000..02325817 --- /dev/null +++ b/cuda/private/templates/BUILD.nvcc_nvvm_embedded @@ -0,0 +1,27 @@ +nvcc_cicc_bin = any_or_first( + "%{component_name}/nvvm/bin/cicc", + "%{component_name}/nvvm/bin/cicc.exe", +) + +alias( + name = "cicc", + actual = if_cuda_toolkit_version_ge( + (13, 0), + "@rules_cuda//cuda/dummy:cicc", + nvcc_cicc_bin, + ), +) + +alias( + name = "libdevice.10.bc", + actual = if_cuda_toolkit_version_ge( + (13, 0), + "@rules_cuda//cuda/dummy:libdevice.10.bc", + "%{component_name}/nvvm/libdevice/libdevice.10.bc", + ), +) + +alias( + name = "libdevice", + actual = ":libdevice.10.bc", +) diff --git a/cuda/private/templates/registry.bzl b/cuda/private/templates/registry.bzl index dd9535cb..5ce91d86 100644 --- a/cuda/private/templates/registry.bzl +++ b/cuda/private/templates/registry.bzl @@ -1,7 +1,7 @@ # map short component name to consumable targets REGISTRY = { "cudart": ["cudart_all_files", "cudart_license", "cudart_headers", "cuda", "cuda_runtime", "cuda_runtime_static", "cudadevrt_a", "cudadevrt_lib", "culibos_a"], # device-runtime targets are not for end users. - "nvcc": ["nvcc_all_files", "nvcc_license", "nvcc_headers", "compiler_root", "compiler_deps", "nvptxcompiler", "nvcc", "nvlink", "ptxas", "bin2c", "fatbinary", "link.stub"], + "nvcc": ["nvcc_all_files", "nvcc_license", "nvcc_headers", "compiler_root", "compiler_deps", "nvptxcompiler", "nvcc", "nvlink", "ptxas", "bin2c", "fatbinary", "link.stub", "cicc", "libdevice", "libdevice.10.bc"], "nvvm": ["nvvm_all_files", "nvvm_license", "nvvm_headers", "cicc", "libdevice", "libdevice.10.bc"], "cccl": ["cccl_all_files", "cccl_license", "cccl_headers", "libcudacxx", "cub", "thrust"], "crt": ["crt_all_files", "crt_license", "crt_headers", "crt"], diff --git a/cuda/private/toolchain.bzl b/cuda/private/toolchain.bzl index b5fadd41..af229c11 100644 --- a/cuda/private/toolchain.bzl +++ b/cuda/private/toolchain.bzl @@ -172,7 +172,5 @@ def register_detected_cuda_toolchains(): User can setup their own toolchain if needed and ignore the detected ones by not calling this macro. """ native.register_toolchains( - "@cuda//toolchain:nvcc-local-toolchain", - "@cuda//toolchain/clang:clang-local-toolchain", - "@cuda//toolchain/disabled:disabled-local-toolchain", + "@cuda//toolchain/...", ) diff --git a/tests/integration/test_all.sh b/tests/integration/test_all.sh index 8a7a2818..335ed33a 100755 --- a/tests/integration/test_all.sh +++ b/tests/integration/test_all.sh @@ -12,6 +12,7 @@ skip_redist_json=false skip_redist_json_multi=false skip_redist_json_collision=false skip_redist_json_version_gate=false +skip_redist_json_cross_major=false while [[ $# -gt 0 ]]; do case $1 in @@ -35,6 +36,8 @@ while [[ $# -gt 0 ]]; do skip_redist_json_collision=true; shift ;; --no-redist-version-gate) skip_redist_json_version_gate=true; shift ;; + --no-redist-cross-major) + skip_redist_json_cross_major=true; shift ;; *) echo "Unknown option: $1" >&2; shift ;; esac @@ -53,6 +56,52 @@ if [[ "$RUNNER_OS" == "Windows" ]] || [[ "$(uname -s 2>/dev/null)" =~ MINGW|MSYS ) fi +# Component BUILD files are evaluated in the component repository's own CUDA +# version context. In particular, CCCL moved its public include root in CUDA 13. +if [ "$skip_redist_json_cross_major" = false ]; then +cat <<- EOF + +============================================================ +=== TEST: TOOLCHAIN WITH REDISTRIB.JSON (BZLMOD CROSS-MAJOR) +============================================================ +EOF +pushd "$this_dir/toolchain_redist_json_cross_major" + # CcInfo is native under Bazel 7 and Starlark-defined under Bazel 9 with the + # rules_cc toolchain, so match its provider key and inspect both include fields. + cccl_includes_expr='[v for k, v in providers(target).items() if k.endswith("CcInfo")][0].compilation_context.includes.to_list() + [v for k, v in providers(target).items() if k.endswith("CcInfo")][0].compilation_context.system_includes.to_list()' + cccl_12=$(env -u CUDA_REDIST_VERSION_OVERRIDE bazel cquery \ + '@cuda//:thrust' --enable_bzlmod --output=starlark \ + "--starlark:expr=${cccl_includes_expr}" \ + --@rules_cuda//cuda:version=12.8.1 "${redist_platform_args[@]}") + if ! [[ $cccl_12 == *'/cccl/include"'* ]]; then exit 1; fi + if [[ $cccl_12 == *'/cccl/include/cccl"'* ]]; then exit 1; fi + + cccl_13=$(env -u CUDA_REDIST_VERSION_OVERRIDE bazel cquery \ + '@cuda//:thrust' --enable_bzlmod --output=starlark \ + "--starlark:expr=${cccl_includes_expr}" \ + --@rules_cuda//cuda:version=13.0.1 "${redist_platform_args[@]}") + if ! [[ $cccl_13 == *'/cccl/include/cccl"'* ]]; then exit 1; fi + + default_toolchain=$(env -u CUDA_REDIST_VERSION_OVERRIDE bazel query \ + '@cuda//toolchain:nvcc-local-toolchain' --enable_bzlmod --output=build) + if ! [[ $default_toolchain == *'cuda_toolchain_13_0_1//toolchain:nvcc-local'* ]]; then exit 1; fi + cuda_12_toolchain=$(env -u CUDA_REDIST_VERSION_OVERRIDE bazel query \ + '@cuda//toolchain:nvcc-12_8_1-toolchain' --enable_bzlmod --output=build) + if ! [[ $cuda_12_toolchain == *'cuda_toolchain_12_8_1//toolchain:nvcc-local'* ]]; then exit 1; fi + + cross_major_args=(--enable_bzlmod --@rules_cuda//cuda:compiler=nvcc "${redist_platform_args[@]}") + env -u CUDA_REDIST_VERSION_OVERRIDE bazel build //:kernel_lib \ + --@rules_cuda//cuda:version=12.8.1 "${cross_major_args[@]}" + env -u CUDA_REDIST_VERSION_OVERRIDE bazel build //:kernel_lib \ + --@rules_cuda//cuda:version=13.0.1 "${cross_major_args[@]}" + env -u CUDA_REDIST_VERSION_OVERRIDE bazel build //:kernel_lib --nobuild --enable_bzlmod \ + --@rules_cuda//cuda:compiler=clang --@rules_cuda//cuda:version=12.8.1 "${redist_platform_args[@]}" + env -u CUDA_REDIST_VERSION_OVERRIDE bazel build //:kernel_lib --nobuild --enable_bzlmod \ + --@rules_cuda//cuda:compiler=clang --@rules_cuda//cuda:version=13.0.1 "${redist_platform_args[@]}" + bazel clean && bazel shutdown +popd +fi + # toolchain configured by the root module of the user if [ "$skip_root" = false ]; then cat <<- EOF @@ -203,14 +252,18 @@ cat <<- EOF ============================================================ EOF pushd "$this_dir/toolchain_redist_json_multi" - bazel build --enable_bzlmod //... --@rules_cuda//cuda:enable=False "${redist_platform_args[@]}" - bazel build --enable_bzlmod //... --@rules_cuda//cuda:enable=True "${redist_platform_args[@]}" - bazel build --enable_bzlmod //:optionally_use_rule --@rules_cuda//cuda:enable=False "${redist_platform_args[@]}" - bazel build --enable_bzlmod //:optionally_use_rule --@rules_cuda//cuda:enable=True --@rules_cuda//cuda:version=12.6.3 "${redist_platform_args[@]}" - bazel build --enable_bzlmod //:optionally_use_rule --@rules_cuda//cuda:enable=True --@rules_cuda//cuda:version=11.7.0 "${redist_platform_args[@]}" - bazel build --enable_bzlmod //:use_library "${redist_platform_args[@]}" - bazel build --enable_bzlmod //:use_rule --@rules_cuda//cuda:version=12.6.3 "${redist_platform_args[@]}" - bazel build --enable_bzlmod //:use_rule --@rules_cuda//cuda:version=11.7.0 "${redist_platform_args[@]}" + env -u CUDA_REDIST_VERSION_OVERRIDE bazel build --enable_bzlmod //... --@rules_cuda//cuda:enable=False "${redist_platform_args[@]}" + env -u CUDA_REDIST_VERSION_OVERRIDE bazel build --enable_bzlmod //... --@rules_cuda//cuda:enable=True "${redist_platform_args[@]}" + env -u CUDA_REDIST_VERSION_OVERRIDE bazel build --enable_bzlmod //:optionally_use_rule --@rules_cuda//cuda:enable=False "${redist_platform_args[@]}" + # Explicit versions are analysis-only: the CI host compiler can be outside an + # older toolkit's support window or emit PTX newer than that toolkit accepts. + # Analysis still exercises exact toolchain resolution without coupling this + # test to the host compiler/toolkit compatibility window. + env -u CUDA_REDIST_VERSION_OVERRIDE bazel build --nobuild --enable_bzlmod //:optionally_use_rule --@rules_cuda//cuda:enable=True --@rules_cuda//cuda:version=12.6.3 "${redist_platform_args[@]}" + env -u CUDA_REDIST_VERSION_OVERRIDE bazel build --nobuild --enable_bzlmod //:optionally_use_rule --@rules_cuda//cuda:enable=True --@rules_cuda//cuda:version=11.7.0 "${redist_platform_args[@]}" + env -u CUDA_REDIST_VERSION_OVERRIDE bazel build --enable_bzlmod //:use_library "${redist_platform_args[@]}" + env -u CUDA_REDIST_VERSION_OVERRIDE bazel build --nobuild --enable_bzlmod //:use_rule --@rules_cuda//cuda:version=12.6.3 "${redist_platform_args[@]}" + env -u CUDA_REDIST_VERSION_OVERRIDE bazel build --nobuild --enable_bzlmod //:use_rule --@rules_cuda//cuda:version=11.7.0 "${redist_platform_args[@]}" # Keep the override-only dedupe probe isolated so it cannot pollute later versioned builds. bazel clean && bazel shutdown @@ -242,8 +295,12 @@ cat <<- EOF EOF pushd "$this_dir/toolchain_redist_json_version_gate" version_gate_args=(--@rules_cuda//cuda:compiler=nvcc "${redist_platform_args[@]}") + env -u CUDA_REDIST_VERSION_OVERRIDE bazel build --enable_bzlmod //:kernel_lib "${version_gate_args[@]}" env -u CUDA_REDIST_VERSION_OVERRIDE bazel build --enable_bzlmod //:kernel_lib --@rules_cuda//cuda:version=12.9.1 "${version_gate_args[@]}" env -u CUDA_REDIST_VERSION_OVERRIDE bazel build --enable_bzlmod //:kernel_lib --@rules_cuda//cuda:version=12.8.1 "${version_gate_args[@]}" + ERR=$(env -u CUDA_REDIST_VERSION_OVERRIDE bazel build --enable_bzlmod //:kernel_lib \ + --@rules_cuda//cuda:version=12.7.1 "${version_gate_args[@]}" 2>&1 || true) + if ! [[ $ERR == *"No matching toolchains found"* ]]; then exit 1; fi bazel clean && bazel shutdown popd fi diff --git a/tests/integration/toolchain_redist_json_cross_major/BUILD.bazel b/tests/integration/toolchain_redist_json_cross_major/BUILD.bazel new file mode 100644 index 00000000..a6c430b3 --- /dev/null +++ b/tests/integration/toolchain_redist_json_cross_major/BUILD.bazel @@ -0,0 +1,15 @@ +load("@bazel_skylib//rules:write_file.bzl", "write_file") +load("@rules_cuda//cuda:defs.bzl", "cuda_library") + +write_file( + name = "kernel", + out = "kernel.cu", + content = ["__global__ void noop() {}"], +) + +cuda_library( + name = "kernel_lib", + srcs = [":kernel"], + tags = ["manual"], + deps = ["@cuda//:cuda_runtime"], +) diff --git a/tests/integration/toolchain_redist_json_cross_major/MODULE.bazel b/tests/integration/toolchain_redist_json_cross_major/MODULE.bazel new file mode 100644 index 00000000..caae1c92 --- /dev/null +++ b/tests/integration/toolchain_redist_json_cross_major/MODULE.bazel @@ -0,0 +1,29 @@ +module(name = "bzlmod_redist_json_cross_major") + +bazel_dep(name = "bazel_skylib", version = "1.7.1") +bazel_dep(name = "rules_cc", version = "0.2.14") +bazel_dep(name = "rules_cuda", version = "0.0.0") +local_path_override( + module_name = "rules_cuda", + path = "../../..", +) + +cuda = use_extension("@rules_cuda//cuda:extensions.bzl", "toolchain") +cuda.redist_json( + name = "cuda_13_0_1", + platforms = [ + "linux-x86_64", + "windows-x86_64", + ], + version = "13.0.1", +) +cuda.redist_json( + name = "cuda_12_8_1", + platforms = [ + "linux-x86_64", + "windows-x86_64", + ], + version = "12.8.1", +) +cuda.toolkit(name = "cuda") +use_repo(cuda, "cuda")