diff --git a/.bazelrc.local b/.bazelrc.local new file mode 100644 index 00000000..b1770157 --- /dev/null +++ b/.bazelrc.local @@ -0,0 +1 @@ +common --repo_env=HERMETIC_PYTHON_VERSION=3.12 diff --git a/MODULE.bazel b/MODULE.bazel index e1de7e91..a89507db 100644 --- a/MODULE.bazel +++ b/MODULE.bazel @@ -146,6 +146,12 @@ use_repo(sycl_configure, "local_config_sycl", "oneapi", "level_zero", "zero_load hipcc_configure = use_extension("//extensions:hipcc_configure.bzl", "hipcc_configure_ext") use_repo(hipcc_configure, "config_rocm_hipcc") +############################################################## +# ROCm configuration + +hipcc_configure = use_extension("//extensions:hipcc_configure.bzl", "hipcc_configure_ext") +use_repo(hipcc_configure, "config_rocm_hipcc") + ############################################################## # Hermetic toolchain configuration diff --git a/WORKSPACE b/WORKSPACE index cea8f4e1..c4cb08a2 100644 --- a/WORKSPACE +++ b/WORKSPACE @@ -245,9 +245,3 @@ load("//gpu/rocm:hipcc_configure.bzl", "hipcc_configure") hipcc_configure(name = "config_rocm_hipcc") -############################################################## -# Local sysroot configuration - -load("//cc/sysroots:local_sysroot_configure.bzl", "local_sysroot_configure") - -local_sysroot_configure(name = "local_sysroot_config") diff --git a/cc/features/BUILD b/cc/features/BUILD index 93eb6a02..0564f16d 100644 --- a/cc/features/BUILD +++ b/cc/features/BUILD @@ -84,35 +84,11 @@ cc_feature( cc_toolchain_sysroot_feature( name = "sysroot_linux_x86_64", - enabled = True, + enabled = True, # Always enabled for hermetic toolchain sysroot = "@sysroot_linux_x86_64//:sysroot", target = "x86_64-linux-gnu", ) -# Local sysroot feature - use system root instead of hermetic sysroot -# Generic for all Linux architectures -cc_feature( - name = "local_sysroot_linux", - enabled = True, - compiler_flags = ["--sysroot=/"], - linker_flags = ["--sysroot=/"], - env_sets = { - "HOST_SYSROOT": "/", - }, -) - -# When using local sysroot, we need to explicitly link default system libraries -# since we don't use hermetic_libraries feature -cc_feature( - name = "local_sysroot_default_libs", - enabled = True, - linker_flags = [ - "-lstdc++", - "-lm", - "-lpthread", - "-ldl", - ], -) cc_toolchain_sysroot_feature( name = "sysroot_linux_aarch64", @@ -190,9 +166,9 @@ cc_feature( compiler_flags = [ "-O0", ], + expand_if_not_available = "opt", #implies = ["generate_pdb_file"], provides = ["compilation_mode"], - expand_if_not_available = "opt", ) cc_feature( @@ -214,15 +190,37 @@ cc_feature( provides = ["compilation_mode"], ) +# Use hermetic libstdc++ instead of libc++ +# Provides GLIBCXX_3.4.25 (GCC 8) from hermetic sysroot cc_feature( - name = "libc++", + name = "use_libstdcxx", compiler_flags = [ - # Allows the use of GNU-specific extension keywords that are not part of the standard C or C++ specifications - "-fgnu-keywords", + "-stdlib=libstdc++", ], + enabled = select({ + "//cc/sysroots:libstdcxx_enabled": True, + "//conditions:default": False, + }), linker_flags = [ - "-stdlib=libc++", + "-stdlib=libstdc++", + ], +) + +# Statically link hermetic libstdc++ to avoid C++ ABI issues +# Use this when you need to isolate from system libstdc++ (e.g., ROCm compatibility) +# When enabled, also enables use_libstdcxx +cc_feature( + name = "static_libstdcxx", + compiler_flags = [ + "-stdlib=libstdc++", + ], + enabled = select({ + "//cc/sysroots:static_libstdcxx_enabled": True, + "//conditions:default": False, + }), + linker_flags = [ + "-stdlib=libstdc++", + "-static-libstdc++", ], - enabled = True, ) diff --git a/cc/impls/linux_x86_64_linux_x86_64/BUILD b/cc/impls/linux_x86_64_linux_x86_64/BUILD index b2717c2c..cd56c190 100644 --- a/cc/impls/linux_x86_64_linux_x86_64/BUILD +++ b/cc/impls/linux_x86_64_linux_x86_64/BUILD @@ -198,6 +198,8 @@ FEATURES_ESSENTIAL = [ "//third_party/rules_cc_toolchain/features:language", "//third_party/rules_cc_toolchain/features:use_lld", "//cc/features:sysroot_linux_x86_64", + "//cc/features:use_libstdcxx", + "//cc/features:static_libstdcxx", ":startup_libs_feature", "//third_party/rules_cc_toolchain/features:coverage", #"//cc/features:clang19", # TODO: Add a selection mechanism based on the Clang version diff --git a/cc/impls/linux_x86_64_linux_x86_64_rocm/BUILD b/cc/impls/linux_x86_64_linux_x86_64_rocm/BUILD index dcb9848e..ad59716f 100644 --- a/cc/impls/linux_x86_64_linux_x86_64_rocm/BUILD +++ b/cc/impls/linux_x86_64_linux_x86_64_rocm/BUILD @@ -1,5 +1,3 @@ -load("@local_sysroot_config//:defs.bzl", "local_sysroot_info") - # Copyright 2026 Google LLC # # Licensed under the Apache License, Version 2.0 (the "License"); @@ -16,12 +14,17 @@ load("@local_sysroot_config//:defs.bzl", "local_sysroot_info") # ============================================================================== # ============================================================================== -# Tools: Linux x86_64, Sysroot: Linux x86_64, ROCm +# ROCm cc_toolchain with hipcc compiler # ============================================================================== +# +# This toolchain provides: +# 1. hipcc as the C++ compiler (for rocm_compile rule to use) +# 2. ROCm include directories (clang headers, HIP headers) +# 3. ROCm compiler flags via rocm_hipcc_feature +# 4. Standard linking capability for combining compiled .o files load("@config_rocm_hipcc//rocm:build_defs.bzl", "hipcc_config") load("@rules_cc//cc:defs.bzl", "cc_toolchain") -load("//cc/rocm/features:rocm_hipcc_feature.bzl", "rocm_hipcc_feature") load("//third_party/rules_cc_toolchain:toolchain_config.bzl", "cc_toolchain_config") load("//third_party/rules_cc_toolchain/features:cc_toolchain_import.bzl", "cc_toolchain_import") load("//third_party/rules_cc_toolchain/features:features.bzl", "cc_toolchain_import_feature") @@ -32,32 +35,13 @@ package( ], ) -ROCM_TOOLS = { - "gcc": "wrappers/hipcc_wrapper", - "cpp": "wrappers/clang++", - "ld": "wrappers/ld", - "ar": "wrappers/ar", - "gcov": "wrappers/idler", - "llvm-cov": "wrappers/idler", - "nm": "wrappers/idler", - "objdump": "wrappers/idler", - "strip": "wrappers/idler", -} - # ============================================================================= -# ROCm HIPcc Feature +# Wrappers # ============================================================================= -rocm_hipcc_feature( - name = "rocm_hipcc_feature", - amdgpu_targets = hipcc_config().gpu_architectures, - clang_version = hipcc_config().clang_version, - enabled = True, - hipcc_path = hipcc_config().hipcc_path, - host_compiler = "@llvm_linux_x86_64//:clang", - rocm_path = hipcc_config().rocm_path, - rocm_toolkit = "@config_rocm_hipcc//rocm:rocm_root", - version = str(hipcc_config().version_number), +filegroup( + name = "wrappers", + srcs = glob(["wrappers/*"]), visibility = ["//visibility:public"], ) @@ -65,32 +49,9 @@ rocm_hipcc_feature( # Filegroups # ============================================================================= -filegroup( - name = "wrappers", - srcs = ["//cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers:all"], - visibility = ["//visibility:public"], -) - -filegroup( - name = "rocm_toolkit_files", - srcs = [ - "@config_rocm_hipcc//rocm:toolchain_data", - ], -) - -cc_toolchain_import( - name = "empty_startup_libs", - deps = [], -) - alias( name = "startup_libs", - actual = select({ - # Local sysroot: startup libs via --sysroot flag - "//cc/sysroots:local_sysroot_enabled": ":empty_startup_libs", - # Hermetic sysroot: use hermetic startup libs - "//conditions:default": "@sysroot_linux_x86_64//:startup_libs", - }), + actual = "@sysroot_linux_x86_64//:startup_libs", visibility = ["//visibility:public"], ) @@ -100,28 +61,19 @@ cc_toolchain_import_feature( toolchain_import = ":startup_libs", ) -# buildifier: leave-alone cc_toolchain_import( name = "imports", + visibility = ["//visibility:public"], deps = [ - # Hermetic LLVM headers - "@llvm_linux_x86_64//:compiler_incs", - "@llvm_linux_x86_64//:libclang_rt", - # Always include hermetic sysroot headers - needed for GPU compilation - # GPU compilation requires hermetic headers (system headers incompatible with ROCm clang) - # Host compilation will use system headers when HOST_SYSROOT is set (via wrapper override) + # IMPORTANT: std_incs must come before compiler includes for correct stdatomic.h resolution + # The sysroot C++ stdatomic.h wrapper uses #include_next to find clang's built-in stdatomic.h "@sysroot_linux_x86_64//:std_incs", + "@llvm_linux_x86_64//:compiler_incs", "@sysroot_linux_x86_64//:sys_incs", - ] + select({ - # Local sysroot: don't include hermetic libs (use system libs via --sysroot=/) - "//cc/sysroots:local_sysroot_enabled": [], - # Hermetic sysroot: use hermetic libs - "//conditions:default": [ - "@sysroot_linux_x86_64//:std_libs", - "@sysroot_linux_x86_64//:sys_libs", - ], - }), - visibility = ["//visibility:public"], + "@llvm_linux_x86_64//:libclang_rt", + "@sysroot_linux_x86_64//:std_libs", + "@sysroot_linux_x86_64//:sys_libs", + ], ) cc_toolchain_import_feature( @@ -130,62 +82,58 @@ cc_toolchain_import_feature( toolchain_import = ":imports", ) -# buildifier: leave-alone filegroup( - name = "all", + name = "compiler", srcs = [ ":imports", - ":rocm_toolkit_files", - ":startup_libs", ":wrappers", - "@llvm_linux_x86_64//:all", + "@config_rocm_hipcc//rocm:toolchain_data", + "@llvm_linux_x86_64//:distro_libs", ], ) -# buildifier: leave-alone filegroup( - name = "compiler", + name = "linker", srcs = [ ":imports", - ":rocm_toolkit_files", + ":startup_libs", ":wrappers", - "@llvm_linux_x86_64//:clang", - "@llvm_linux_x86_64//:clang++", - "@llvm_linux_x86_64//:distro_libs", + "@config_rocm_hipcc//rocm:toolchain_data", ], ) -# buildifier: leave-alone filegroup( - name = "linker", + name = "ar", srcs = [ - ":compiler", - ":rocm_toolkit_files", - ":startup_libs", ":wrappers", - "@llvm_linux_x86_64//:distro_libs", - "@llvm_linux_x86_64//:ld", + "@config_rocm_hipcc//rocm:toolchain_data", ], ) -# buildifier: leave-alone filegroup( - name = "ar", + name = "strip", srcs = [ - ":rocm_toolkit_files", ":wrappers", - "@llvm_linux_x86_64//:ar", - "@llvm_linux_x86_64//:distro_libs", + "@config_rocm_hipcc//rocm:toolchain_data", ], ) -# buildifier: leave-alone filegroup( - name = "strip", + name = "objcopy", srcs = [ ":wrappers", - "@llvm_linux_x86_64//:distro_libs", - "@llvm_linux_x86_64//:strip", + "@config_rocm_hipcc//rocm:toolchain_data", + ], +) + +filegroup( + name = "all", + srcs = [ + ":ar", + ":compiler", + ":linker", + ":objcopy", + ":strip", ], ) @@ -193,50 +141,29 @@ filegroup( # Features # ============================================================================= -FEATURES_ESSENTIAL = select({ - # Local sysroot: DO NOT use hermetic_libraries (no -nostdinc/-nostdinc++) - # We want to use system headers and libraries - "//cc/sysroots:local_sysroot_enabled": [ - "//cc/features:local_sysroot_default_libs", - ], - # Hermetic sysroot: use hermetic_libraries feature - "//conditions:default": ["//third_party/rules_cc_toolchain/features:hermetic_libraries"], -}) + [ +FEATURES_ESSENTIAL = [ # Library imports ":imports_feature", - ":startup_libs_feature", - # ROCm-specific features - ":rocm_hipcc_feature", - "//cc/rocm/features:rocm_defines", - "//cc/rocm/features:rocm_hermetic", - "//cc/rocm/features:rocm_warnings", + # Sysroot and C++ library (always hermetic for ROCm) + ":startup_libs_feature", + "//cc/features:sysroot_linux_x86_64", + "//cc/features:use_libstdcxx", + "//cc/features:static_libstdcxx", # Toolchain configuration "//third_party/rules_cc_toolchain/features:warnings", "//third_party/rules_cc_toolchain/features:reproducible", "//third_party/rules_cc_toolchain/features:use_lld", -] + select({ - "//cc/sysroots:local_sysroot_enabled": ["//cc/features:local_sysroot_linux"], - "//conditions:default": ["//cc/features:sysroot_linux_x86_64"], -}) + [ - # PIC / PIE flags "//third_party/rules_cc_toolchain/features:supports_pic", "//third_party/rules_cc_toolchain/features:position_independent_code", - "//third_party/rules_cc_toolchain/features:position_independent_executable", - - # Optimization flags "//third_party/rules_cc_toolchain/features:garbage_collect_symbols", "//cc/features:constants_merge", "//cc/features:dbg", "//cc/features:fastbuild", "//cc/features:opt", "//cc/features:detect_issues", - - # C++ standard flags "//third_party/rules_cc_toolchain/features:c++17", - - # Other flags "//cc/features:allow_shlib_undefined", "//cc/features:supports_dynamic_linker", "//cc/features:supports_start_end_lib_feature", @@ -252,30 +179,23 @@ cc_toolchain_config( c_compiler = "@llvm_linux_x86_64//:clang", cc_compiler = "@llvm_linux_x86_64//:clang++", compiler_features = FEATURES_ESSENTIAL, - cxx_builtin_include_directories = select({ - # Local sysroot: use dynamically detected system include paths - "//cc/sysroots:local_sysroot_enabled": local_sysroot_info.include_directories, - # Hermetic sysroot: no system directories - "//conditions:default": [], - }), linker = "@llvm_linux_x86_64//:ld", strip_tool = "@llvm_linux_x86_64//:strip", target_cpu = "x86_64", target_system_name = "local", - tool_paths = ROCM_TOOLS, ) cc_toolchain( name = "toolchain", all_files = ":all", ar_files = ":ar", - as_files = ":compiler", compiler_files = ":compiler", dwp_files = ":all", linker_files = ":linker", - objcopy_files = ":all", + objcopy_files = ":objcopy", strip_files = ":strip", supports_param_files = 1, toolchain_config = ":config", toolchain_identifier = "toolchain_linux_x86_64_linux_x86_64_rocm_id", + visibility = ["//visibility:public"], ) diff --git a/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/ar b/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/ar index a8dc00ce..8217aaca 100755 --- a/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/ar +++ b/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/ar @@ -1,5 +1,4 @@ #!/bin/bash - # Copyright 2026 Google LLC # # Licensed under the Apache License, Version 2.0 (the "License"); @@ -13,10 +12,6 @@ # WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. # See the License for the specific language governing permissions and # limitations under the License. -# ============================================================================== - -# allow passing the output of one command as input to the next -set -o pipefail -# run specified tool with all arguments -$AR_PATH $@ +# Wrapper that calls llvm-ar from ROCm distribution +exec external/config_rocm_hipcc/rocm/rocm_dist/llvm/bin/llvm-ar "$@" diff --git a/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/clang b/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/clang index a97f8a7b..e1c87959 100755 --- a/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/clang +++ b/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/clang @@ -1,5 +1,4 @@ #!/bin/bash - # Copyright 2026 Google LLC # # Licensed under the Apache License, Version 2.0 (the "License"); @@ -13,10 +12,6 @@ # WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. # See the License for the specific language governing permissions and # limitations under the License. -# ============================================================================== - -# allow passing the output of one command as input to the next -set -o pipefail -# run specified tool with all arguments -$GCC_PATH $@ +# Wrapper that calls hipcc +exec external/config_rocm_hipcc/rocm/rocm_dist/bin/hipcc "$@" diff --git a/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/ld.lld b/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/ld.lld new file mode 100755 index 00000000..726d4930 --- /dev/null +++ b/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/ld.lld @@ -0,0 +1,17 @@ +#!/bin/bash +# Copyright 2026 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Wrapper that calls ld.lld from ROCm distribution +exec external/config_rocm_hipcc/rocm/rocm_dist/llvm/bin/ld.lld "$@" diff --git a/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/llvm-ar b/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/llvm-ar new file mode 100755 index 00000000..8217aaca --- /dev/null +++ b/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/llvm-ar @@ -0,0 +1,17 @@ +#!/bin/bash +# Copyright 2026 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Wrapper that calls llvm-ar from ROCm distribution +exec external/config_rocm_hipcc/rocm/rocm_dist/llvm/bin/llvm-ar "$@" diff --git a/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/llvm-objcopy b/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/llvm-objcopy new file mode 100755 index 00000000..8f68ef0f --- /dev/null +++ b/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/llvm-objcopy @@ -0,0 +1,17 @@ +#!/bin/bash +# Copyright 2026 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Wrapper that calls llvm-objcopy from ROCm distribution +exec external/config_rocm_hipcc/rocm/rocm_dist/llvm/bin/llvm-objcopy "$@" diff --git a/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/llvm-strip b/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/llvm-strip new file mode 100755 index 00000000..c70037f9 --- /dev/null +++ b/cc/impls/linux_x86_64_linux_x86_64_rocm/wrappers/llvm-strip @@ -0,0 +1,17 @@ +#!/bin/bash +# Copyright 2026 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Wrapper that calls llvm-strip from ROCm distribution +exec external/config_rocm_hipcc/rocm/rocm_dist/llvm/bin/llvm-strip "$@" diff --git a/cc/rocm/BUILD b/cc/rocm/BUILD new file mode 100644 index 00000000..1263e7f4 --- /dev/null +++ b/cc/rocm/BUILD @@ -0,0 +1,17 @@ +# Copyright 2025 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +package(default_visibility = ["//visibility:public"]) + +exports_files(["rocm_compile.bzl"]) diff --git a/cc/rocm/features/BUILD b/cc/rocm/features/BUILD index 3eaeb922..6d98edb8 100644 --- a/cc/rocm/features/BUILD +++ b/cc/rocm/features/BUILD @@ -15,37 +15,4 @@ """ROCm-specific toolchain features.""" -load( - "//third_party/rules_cc_toolchain/features:features.bzl", - "cc_feature", -) - package(default_visibility = ["//visibility:public"]) - -# ROCm-specific preprocessor definitions -cc_feature( - name = "rocm_defines", - compiler_flags = [ - "-DTENSORFLOW_USE_ROCM=1", - "-D__HIP_PLATFORM_AMD__", - "-DEIGEN_USE_HIP", - "-DUSE_ROCM", - ], - enabled = True, -) - -# Hermetic ROCm includes - prevents using system HIP headers -cc_feature( - name = "rocm_hermetic", - compiler_flags = ["-nohipinc"], - enabled = True, -) - -# Warning settings specific to ROCm compilation -cc_feature( - name = "rocm_warnings", - compiler_flags = [ - "-Wno-unused-result", - ], - enabled = True, -) diff --git a/cc/rocm/rocm_compile.bzl b/cc/rocm/rocm_compile.bzl new file mode 100644 index 00000000..5a77c535 --- /dev/null +++ b/cc/rocm/rocm_compile.bzl @@ -0,0 +1,199 @@ +# Copyright 2025 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +"""ROCm compilation rule that compiles GPU code into .pic.o files.""" + +load("@config_rocm_hipcc//rocm:build_defs.bzl", "hipcc_config") + +def _rocm_compile_impl(ctx): + """Compiles ROCm sources into .pic.o files with embedded GPU kernels (fat binaries).""" + + # Get the ROCm toolchain provided as an attribute + cc_toolchain = ctx.attr._cc_toolchain[cc_common.CcToolchainInfo] + + # Get GPU architectures from ROCm configuration + gpu_architectures = hipcc_config().gpu_architectures + + # Collect compilation contexts from dependencies + cc_infos = [dep[CcInfo] for dep in ctx.attr.deps if CcInfo in dep] + compilation_contexts = [cc_info.compilation_context for cc_info in cc_infos] + + # Merge compilation context from dependencies + merged_compilation_context = cc_common.merge_compilation_contexts( + compilation_contexts = compilation_contexts, + ) + + # Get feature configuration from toolchain + feature_configuration = cc_common.configure_features( + ctx = ctx, + cc_toolchain = cc_toolchain, + requested_features = ctx.features, + unsupported_features = ctx.disabled_features, + ) + + # Get compiler flags from features + compile_variables = cc_common.create_compile_variables( + feature_configuration = feature_configuration, + cc_toolchain = cc_toolchain, + user_compile_flags = ctx.attr.copts, + use_pic = True, # Enable position-independent code + ) + compiler_flags = cc_common.get_memory_inefficient_command_line( + feature_configuration = feature_configuration, + action_name = "c++-compile", + variables = compile_variables, + ) + + # Get environment variables from feature configuration + # These include HIPCC_PATH, ROCM_PATH, etc. set by rocm_hipcc_feature + env_vars = cc_common.get_environment_variables( + feature_configuration = feature_configuration, + action_name = "c++-compile", + variables = compile_variables, + ) + + # Get compiler path from toolchain + hipcc_path = cc_common.get_tool_for_action( + feature_configuration = feature_configuration, + action_name = "c++-compile", + ) + + # Compile each source file + objects = [] + for src in ctx.files.srcs: + # Use target-specific directory to avoid conflicts when same source compiled with different defines + obj = ctx.actions.declare_file("_objs/" + ctx.label.name + "/" + src.basename + ".pic.o") + + # Build compilation command for hipcc + args = ctx.actions.args() + args.add("-x", "hip") # Use standard HIP language mode + args.add("-c") + + # Add GPU architecture targets for device code generation + for arch in gpu_architectures: + args.add("--offload-arch=" + arch) + + # Add compiler flags from toolchain features + args.add_all(compiler_flags) + + # Add include paths from dependencies + args.add_all(merged_compilation_context.includes, before_each = "-I") + args.add_all(merged_compilation_context.quote_includes, before_each = "-iquote") + args.add_all(merged_compilation_context.system_includes, before_each = "-isystem") + + # Add defines from dependencies + args.add_all(merged_compilation_context.defines, format_each = "-D%s") + + # Add user-provided compiler options + args.add_all(ctx.attr.copts) + + # Add source and output + args.add(src) + args.add("-o", obj) + + ctx.actions.run( + executable = hipcc_path, + arguments = [args], + inputs = depset( + direct = [src] + ctx.files.hdrs, + transitive = [ + merged_compilation_context.headers, + cc_toolchain.all_files, + ], + ), + outputs = [obj], + env = env_vars, + mnemonic = "RocmCompile", + ) + + objects.append(obj) + + # Create static library with ar + static_library = ctx.actions.declare_file("lib" + ctx.label.name + ".a") + + # Get archiver path from toolchain + ar_path = cc_common.get_tool_for_action( + feature_configuration = feature_configuration, + action_name = "c++-link-static-library", + ) + + # Build ar command: ar rcs libname.a obj1.o obj2.o ... + ar_args = ctx.actions.args() + ar_args.add("rcs") + ar_args.add(static_library) + ar_args.add_all(objects) + + ctx.actions.run( + executable = ar_path, + arguments = [ar_args], + inputs = depset(direct = objects), + outputs = [static_library], + mnemonic = "RocmArchive", + progress_message = "Creating static library %s" % static_library.short_path, + ) + + # Create library_to_link for static library + library_to_link = cc_common.create_library_to_link( + actions = ctx.actions, + feature_configuration = feature_configuration, + cc_toolchain = cc_toolchain, + static_library = static_library, + alwayslink = True, # Always link the entire archive (needed for GPU kernels) + ) + + # Create linking context + linking_context = cc_common.create_linking_context( + linker_inputs = depset(direct = [ + cc_common.create_linker_input( + owner = ctx.label, + libraries = depset(direct = [library_to_link]), + ), + ]), + ) + + # Create compilation context for headers + compilation_context = cc_common.create_compilation_context( + headers = depset(ctx.files.hdrs), + ) + + # Return CcInfo so this can be used as a dep in cc_library + return [ + DefaultInfo(files = depset([static_library])), + CcInfo( + compilation_context = compilation_context, + linking_context = linking_context, + ), + ] + +rocm_compile = rule( + implementation = _rocm_compile_impl, + attrs = { + "srcs": attr.label_list( + allow_files = [".cu.cc", ".cc"], + mandatory = True, + ), + "hdrs": attr.label_list( + allow_files = [".h", ".cu.h"], + ), + "deps": attr.label_list( + providers = [CcInfo], + ), + "copts": attr.string_list(), + "_cc_toolchain": attr.label( + default = "//cc/impls/linux_x86_64_linux_x86_64_rocm:toolchain", + ), + }, + fragments = ["cpp"], + toolchains = ["@bazel_tools//tools/cpp:toolchain_type"], +) diff --git a/cc/rocm/rocm_library.bzl b/cc/rocm/rocm_library.bzl new file mode 100644 index 00000000..3f09408b --- /dev/null +++ b/cc/rocm/rocm_library.bzl @@ -0,0 +1,56 @@ +# Copyright 2026 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +"""ROCm library rule that compiles GPU code and wraps in cc_library.""" + +load("@rules_cc//cc:defs.bzl", "cc_library") +load("//cc/rocm:rocm_compile.bzl", "rocm_compile") + +def rocm_library(name, srcs = [], hdrs = [], copts = [], deps = [], linkopts = [], local_defines = [], **kwargs): + """Compiles ROCm sources into static library with GPU kernels embedded. + + Args: + name: Name of the library + srcs: Source files to compile (.cu.cc, .cc) + hdrs: Header files (.h, .cu.h) + copts: Compiler options + deps: Dependencies + linkopts: Linker options (currently unused) + local_defines: Preprocessor defines (will be passed as -DNAME) + **kwargs: Additional arguments passed to rocm_compile + """ + # Convert local_defines to copts with -D prefix + define_copts = ["-D" + d for d in local_defines] + all_copts = copts + define_copts + + if srcs: + # rocm_compile handles everything: compilation, archiving, and CcInfo creation + # Note: alwayslink is always True for GPU kernels (filtered out from kwargs) + # linkstatic is also filtered as it's not applicable (always produces static .a) + rocm_compile( + name = name, + srcs = srcs, + hdrs = hdrs, + deps = deps, + copts = all_copts, + **{k: v for k, v in kwargs.items() if k not in ["linkopts", "local_defines", "alwayslink", "linkstatic"]} + ) + else: + # Header-only library + cc_library( + name = name, + hdrs = hdrs, + deps = deps, + **kwargs + ) diff --git a/cc/sysroots/BUILD b/cc/sysroots/BUILD index 019c58b1..d57e7dd5 100644 --- a/cc/sysroots/BUILD +++ b/cc/sysroots/BUILD @@ -5,19 +5,31 @@ load(":collect_cc_libraries.bzl", "collect_cc_libraries") load("@bazel_skylib//rules:common_settings.bzl", "bool_flag") -# Export template file for local_sysroot_configure repository rule -exports_files(["defs.bzl.tpl"]) +# Flag to control whether to use sysroot's libstdc++ instead of hermetic libc++ +# This is needed for compatibility with prebuilt libraries like ROCm that use libstdc++ +bool_flag( + name = "use_libstdcxx", + build_setting_default = False, + visibility = ["//visibility:public"], +) + +config_setting( + name = "libstdcxx_enabled", + flag_values = {":use_libstdcxx": "True"}, + visibility = ["//visibility:public"], +) -# Flag to control whether to use local (system) sysroot instead of hermetic +# Flag to control whether to statically link libstdc++ +# This is useful to avoid C++ ABI issues with system libraries bool_flag( - name = "use_local_sysroot", + name = "use_static_libstdcxx", build_setting_default = False, visibility = ["//visibility:public"], ) config_setting( - name = "local_sysroot_enabled", - flag_values = {":use_local_sysroot": "True"}, + name = "static_libstdcxx_enabled", + flag_values = {":use_static_libstdcxx": "True"}, visibility = ["//visibility:public"], ) diff --git a/cc/sysroots/sysroot_linux.BUILD.tpl b/cc/sysroots/sysroot_linux.BUILD.tpl index 0b6e5083..4708c67e 100644 --- a/cc/sysroots/sysroot_linux.BUILD.tpl +++ b/cc/sysroots/sysroot_linux.BUILD.tpl @@ -67,3 +67,9 @@ alias( actual = "@@%{sysroot_repo_name}//:openmp_copyright", visibility = ["//visibility:public"], ) + +alias( + name = "stdc++", + actual = "@@%{sysroot_repo_name}//:stdc++", + visibility = ["//visibility:public"], +) diff --git a/cc/tests/gpu/rocm/BUILD b/cc/tests/gpu/rocm/BUILD index 50e178b0..930a971d 100644 --- a/cc/tests/gpu/rocm/BUILD +++ b/cc/tests/gpu/rocm/BUILD @@ -14,20 +14,18 @@ # ============================================================================== load("@bazel_skylib//rules:build_test.bzl", "build_test") -load( - "@config_rocm_hipcc//rocm:build_defs.bzl", - "rocm_cc_test", - "rocm_library", -) +load("@rules_cc//cc:defs.bzl", "cc_test") +load("//cc/rocm:rocm_library.bzl", "rocm_library") rocm_library( name = "vector_hip", srcs = ["vector_hip.hip.cc"], hdrs = ["vector_hip.hip.h"], tags = ["rocm-only"], + deps = ["@config_rocm_hipcc//rocm:hip_runtime"], ) -rocm_cc_test( +cc_test( name = "vector_hip_test", srcs = [ "vector_hip_test.cc", @@ -39,6 +37,7 @@ rocm_cc_test( ], deps = [ ":vector_hip", + "@config_rocm_hipcc//rocm:hip_runtime", "@gtest//:gtest_main", ], ) diff --git a/gpu/rocm/BUILD.tpl b/gpu/rocm/BUILD.tpl index eeb399bb..63c981da 100644 --- a/gpu/rocm/BUILD.tpl +++ b/gpu/rocm/BUILD.tpl @@ -73,22 +73,33 @@ filegroup( visibility = ["//visibility:public"], ) -filegroup( - name = "all_files", - srcs = glob(["%{rocm_root}/**"], allow_empty = True), +# Tool aliases for convenience - all just point to toolchain_data +alias( + name = "hipcc", + actual = ":toolchain_data", visibility = ["//visibility:public"], ) -filegroup( - name = "rocm_root", - srcs = [":all_files"], +alias( + name = "ar", + actual = ":toolchain_data", + visibility = ["//visibility:public"], +) + +alias( + name = "ld", + actual = ":toolchain_data", + visibility = ["//visibility:public"], +) + +alias( + name = "strip", + actual = ":toolchain_data", visibility = ["//visibility:public"], ) -config_setting( - name = "using_hipcc", - define_values = { - "using_rocm": "true", - }, +alias( + name = "objcopy", + actual = ":toolchain_data", visibility = ["//visibility:public"], ) diff --git a/gpu/rocm/build_defs.bzl.tpl b/gpu/rocm/build_defs.bzl.tpl index 516f846a..7fbaa74a 100644 --- a/gpu/rocm/build_defs.bzl.tpl +++ b/gpu/rocm/build_defs.bzl.tpl @@ -27,85 +27,28 @@ def hipcc_config(): Returns: A struct with the following fields: - gpu_architectures: List of AMDGPU target architectures (e.g., ["gfx906", "gfx908"]) - - version_number: ROCm version as integer (e.g., 60000 for ROCm 6.0.0) - hipcc_path: Path to the hipcc compiler binary - - rocm_path: Path to the ROCm installation root - rocm_root: Path to the ROCm root within the distribution (for file labels) - - miopen_version: MIOpen version number - - hipruntime_version: HIP runtime version number - - clang_version: Clang version string - lib_paths: List of library paths (for multiple ROCm paths setup) + - version_number: ROCm version number + - clang_version: Clang version string + - include_paths: System include paths (computed from rocm_root and clang_version) """ + rocm_root = "%{rocm_root}" + clang_version = "%{clang_version}" + + # System include paths - these depend on ROCm installation location + include_paths = [ + rocm_root + "/llvm/lib/clang/" + clang_version + "/include", + rocm_root + "/include", + ] + return struct( gpu_architectures = %{rocm_gpu_architectures}, - version_number = %{rocm_version_number}, hipcc_path = "%{hipcc_path}", - rocm_path = "%{rocm_path}", - rocm_root = "%{rocm_root}", - miopen_version = %{miopen_version_number}, - hipruntime_version = %{hipruntime_version_number}, - clang_version = "%{clang_version}", + rocm_root = rocm_root, + clang_version = clang_version, lib_paths = %{rocm_lib_paths}, - ) - -# Alias for compatibility -rocm_config = hipcc_config - -def if_rocm(if_true, if_false = []): - """Shorthand for select()'ing on whether we're building with ROCm. - - Returns a select statement which evaluates to if_true if we're building - with ROCm enabled. Otherwise, the select statement evaluates to if_false. - - Args: - if_true: Value to return when building with ROCm. - if_false: Value to return when building without ROCm (default: []). - - Returns: - A select expression. - """ - return select({ - "@config_rocm_hipcc//rocm:using_hipcc": if_true, - "//conditions:default": if_false, - }) - -def rocm_default_copts(): - """Default compiler options for ROCm/HIP compilation. - - Returns: - List of compiler options to use with HIP code. - """ - return if_rocm([ - "-D__HIP_PLATFORM_AMD__=1", - "-DTENSORFLOW_USE_ROCM=1", - ]) - -def rocm_library(copts = [], deps = [], **kwargs): - """Wrapper over cc_library which adds default ROCm/HIP options. - - Args: - copts: Additional compiler options. - deps: Library dependencies. - **kwargs: Other arguments passed to cc_library. - """ - cc_library( - copts = rocm_default_copts() + ["-x", "rocm"] + copts, - deps = deps + ["@config_rocm_hipcc//rocm:hip_runtime"], - **kwargs - ) - -def rocm_cc_test(copts = [], deps = [], **kwargs): - """Wrapper over cc_test which adds default ROCm/HIP options. - - Args: - copts: Additional compiler options. - deps: Test dependencies. - **kwargs: Other arguments passed to cc_test. - """ - cc_test( - copts = copts + if_rocm(["-D__HIP_PLATFORM_AMD__=1"]), - deps = deps + [ - "@config_rocm_hipcc//rocm:hip_runtime", - ], - **kwargs + version_number = %{rocm_version_number}, + include_paths = include_paths, ) diff --git a/gpu/rocm/hipcc_configure.bzl b/gpu/rocm/hipcc_configure.bzl index 8e4303f6..75294c0f 100644 --- a/gpu/rocm/hipcc_configure.bzl +++ b/gpu/rocm/hipcc_configure.bzl @@ -43,6 +43,7 @@ def _enable_rocm(repository_ctx): enable_rocm = repository_ctx.os.environ.get("TF_NEED_ROCM") if enable_rocm == "1": return True + # Also enable if ROCM_PATH is set and non-empty (non-hermetic ROCm) rocm_path = repository_ctx.os.environ.get("ROCM_PATH", "") if rocm_path and rocm_path.strip(): @@ -255,6 +256,7 @@ def _setup_rocm_from_multiple_paths(repository_ctx, multiple_paths, bash_bin): if llvm_path: repository_ctx.symlink(llvm_path, _DISTRIBUTION_PATH + "/llvm") repository_ctx.symlink(llvm_path, _DISTRIBUTION_PATH + "/lib/llvm") + # Only create amdgcn symlink if it exists amdgcn_path = llvm_path + "/amdgcn" if files_exist(repository_ctx, [amdgcn_path], bash_bin)[0]: @@ -283,6 +285,7 @@ def _setup_rocm_distro_dir(repository_ctx): # Use system ROCm installation by symlinking it into the repository auto_configure_warning("Using non-hermetic ROCm from ROCM_PATH: {}".format(rocm_path)) repository_ctx.symlink(rocm_path, _DISTRIBUTION_PATH) + # Use the symlinked path (_DISTRIBUTION_PATH) for all operations, not the absolute path return _get_rocm_config(repository_ctx, bash_bin, _DISTRIBUTION_PATH, rocm_path) @@ -302,7 +305,7 @@ def _setup_rocm_distro_dir(repository_ctx): if rocm_distro_version not in rocm_redist: fail("Unknown ROCM_DISTRO_VERSION: {}. Available versions: {}".format( rocm_distro_version, - ", ".join(rocm_redist.keys()) + ", ".join(rocm_redist.keys()), )) repository_ctx.report_progress("Downloading hermetic ROCm distribution: {}".format(rocm_distro_version)) @@ -310,10 +313,11 @@ def _setup_rocm_distro_dir(repository_ctx): def _create_dummy_repository(repository_ctx): """Creates a stub ROCm repository when ROCm is not enabled.""" + # Create stub repository using templates with empty values repository_ctx.file("rocm/empty/.keep", "") stub_dict = { - "%{rocm_root}": "empty", + "%{rocm_root}": "rocm_dist", "%{rocm_gpu_architectures}": "[]", "%{rocm_version_number}": "0", "%{miopen_version_number}": "0", @@ -334,31 +338,19 @@ def _setup_rocm_repository(repository_ctx): miopen_version_number = int(rocm_config.miopen_version_number) hipruntime_version_number = int(rocm_config.hipruntime_version_number) - # Handle hermetic vs non-hermetic ROCm - if rocm_config.install_path: - # Non-hermetic: symlink already created in _setup_rocm_distro_dir - # Use "rocm_dist" (relative to rocm/ directory where BUILD file is) - rocm_toolkit_path = "rocm_dist" - else: - # Hermetic: files already extracted to rocm/rocm_dist - rocm_toolkit_path = _remove_root_dir(rocm_config.rocm_toolkit_path, "rocm") - - # Always use relative paths (either symlink or hermetic dist) - rocm_path_relative = "rocm_dist" - hipcc_path_relative = rocm_path_relative + "/bin/hipcc" + rocm_toolkit_path = _remove_root_dir(rocm_config.rocm_toolkit_path, "rocm") bash_bin = get_bash_bin(repository_ctx) clang_offload_bundler_path = rocm_toolkit_path + "/llvm/bin/clang-offload-bundler" repository_dict = { - "%{rocm_root}": rocm_toolkit_path, + "%{rocm_root}": _DISTRIBUTION_PATH, "%{rocm_gpu_architectures}": str(rocm_config.amdgpu_targets), "%{rocm_version_number}": str(rocm_version_number), "%{miopen_version_number}": str(miopen_version_number), "%{hipruntime_version_number}": str(hipruntime_version_number), - "%{hipcc_path}": hipcc_path_relative, - "%{rocm_path}": rocm_path_relative, + "%{hipcc_path}": _DISTRIBUTION_PATH + "/bin/hipcc", "%{clang_version}": rocm_config.clang_version, "%{rocm_lib_paths}": str(rocm_config.rocm_lib_paths), }