From 1bd227bafa4e7f01ffbce474658f289554f047cb Mon Sep 17 00:00:00 2001 From: Yang Sheng Date: Wed, 10 Apr 2024 14:03:14 -0700 Subject: [PATCH] Add sycl build target #10244 PiperOrigin-RevId: 623599115 --- .bazelrc | 5 + third_party/tsl/.bazelrc | 5 + third_party/tsl/opensource_only.files | 7 + .../third_party/gpus/crosstool/BUILD.sycl.tpl | 71 +++ .../bin/crosstool_wrapper_driver_sycl.tpl | 64 ++ .../sycl_cc_toolchain_config.bzl.tpl | 598 ++++++++++++++++++ .../tsl/third_party/gpus/find_sycl_config.py | 139 ++++ .../third_party/gpus/rocm/build_defs.bzl.tpl | 6 +- .../tsl/third_party/gpus/rocm_configure.bzl | 9 +- third_party/tsl/third_party/gpus/sycl/BUILD | 0 .../tsl/third_party/gpus/sycl/BUILD.tpl | 63 ++ .../third_party/gpus/sycl/build_defs.bzl.tpl | 39 ++ .../tsl/third_party/gpus/sycl_configure.bzl | 531 ++++++++++++++++ third_party/tsl/workspace2.bzl | 2 + xla/stream_executor/build_defs.bzl | 10 +- 15 files changed, 1544 insertions(+), 5 deletions(-) create mode 100644 third_party/tsl/third_party/gpus/crosstool/BUILD.sycl.tpl create mode 100644 third_party/tsl/third_party/gpus/crosstool/clang/bin/crosstool_wrapper_driver_sycl.tpl create mode 100644 third_party/tsl/third_party/gpus/crosstool/sycl_cc_toolchain_config.bzl.tpl create mode 100644 third_party/tsl/third_party/gpus/find_sycl_config.py create mode 100644 third_party/tsl/third_party/gpus/sycl/BUILD create mode 100644 third_party/tsl/third_party/gpus/sycl/BUILD.tpl create mode 100644 third_party/tsl/third_party/gpus/sycl/build_defs.bzl.tpl create mode 100644 third_party/tsl/third_party/gpus/sycl_configure.bzl diff --git a/.bazelrc b/.bazelrc index 4cdeb1065203e..d7ae76f096431 100644 --- a/.bazelrc +++ b/.bazelrc @@ -293,6 +293,11 @@ build:rocm --define=using_rocm_hipcc=true build:rocm --define=tensorflow_mkldnn_contraction_kernel=0 build:rocm --repo_env TF_NEED_ROCM=1 +build:sycl --crosstool_top=@local_config_sycl//crosstool:toolchain +build:sycl --define=using_sycl=true +build:sycl --define=tensorflow_mkldnn_contraction_kernel=0 +build:sycl --repo_env TF_NEED_SYCL=1 + # Options to disable default on features build:noaws --define=no_aws_support=true build:nogcp --define=no_gcp_support=true diff --git a/third_party/tsl/.bazelrc b/third_party/tsl/.bazelrc index 4cdeb1065203e..d7ae76f096431 100644 --- a/third_party/tsl/.bazelrc +++ b/third_party/tsl/.bazelrc @@ -293,6 +293,11 @@ build:rocm --define=using_rocm_hipcc=true build:rocm --define=tensorflow_mkldnn_contraction_kernel=0 build:rocm --repo_env TF_NEED_ROCM=1 +build:sycl --crosstool_top=@local_config_sycl//crosstool:toolchain +build:sycl --define=using_sycl=true +build:sycl --define=tensorflow_mkldnn_contraction_kernel=0 +build:sycl --repo_env TF_NEED_SYCL=1 + # Options to disable default on features build:noaws --define=no_aws_support=true build:nogcp --define=no_gcp_support=true diff --git a/third_party/tsl/opensource_only.files b/third_party/tsl/opensource_only.files index 2920ef7997ad9..2853de8a621c6 100644 --- a/third_party/tsl/opensource_only.files +++ b/third_party/tsl/opensource_only.files @@ -22,11 +22,14 @@ third_party/git/BUILD: third_party/git/git_configure.bzl: third_party/gpus/BUILD: third_party/gpus/crosstool/BUILD.rocm.tpl: +third_party/gpus/crosstool/BUILD.sycl.tpl: third_party/gpus/crosstool/BUILD.tpl: third_party/gpus/crosstool/BUILD: third_party/gpus/crosstool/LICENSE: third_party/gpus/crosstool/clang/bin/crosstool_wrapper_driver_is_not_gcc.tpl: third_party/gpus/crosstool/clang/bin/crosstool_wrapper_driver_rocm.tpl: +third_party/gpus/crosstool/clang/bin/crosstool_wrapper_driver_sycl.tpl: +third_party/gpus/crosstool/sycl_cc_toolchain_config.bzl.tpl: third_party/gpus/crosstool/windows/msvc_wrapper_for_nvcc.py.tpl: third_party/gpus/cuda/BUILD.tpl: third_party/gpus/cuda/BUILD.windows.tpl: @@ -42,6 +45,10 @@ third_party/gpus/rocm/BUILD: third_party/gpus/rocm/build_defs.bzl.tpl: third_party/gpus/rocm/rocm_config.h.tpl: third_party/gpus/rocm_configure.bzl: +third_party/gpus/sycl/BUILD.tpl: +third_party/gpus/sycl/BUILD: +third_party/gpus/sycl/build_defs.bzl.tpl: +third_party/gpus/sycl_configure.bzl: third_party/grpc/BUILD: third_party/implib_so/BUILD: third_party/implib_so/get_symbols.py: diff --git a/third_party/tsl/third_party/gpus/crosstool/BUILD.sycl.tpl b/third_party/tsl/third_party/gpus/crosstool/BUILD.sycl.tpl new file mode 100644 index 0000000000000..a8db760d0c695 --- /dev/null +++ b/third_party/tsl/third_party/gpus/crosstool/BUILD.sycl.tpl @@ -0,0 +1,71 @@ +# This file is expanded from a template sycl_configure.bzl +# Update sycl_configure.bzl#verify_build_defines when adding new variables. + +load(":cc_toolchain_config.bzl", "cc_toolchain_config") + +licenses(["restricted"]) + +package(default_visibility = ["//visibility:public"]) + +toolchain( + name = "toolchain-linux-x86_64", + exec_compatible_with = [ + "@bazel_tools//platforms:linux", + "@bazel_tools//platforms:x86_64", + ], + target_compatible_with = [ + "@bazel_tools//platforms:linux", + "@bazel_tools//platforms:x86_64", + ], + toolchain = ":cc-compiler-local", + toolchain_type = "@bazel_tools//tools/cpp:toolchain_type", +) + +cc_toolchain_suite( + name = "toolchain", + toolchains = { + "local|compiler": ":cc-compiler-local", + "k8": ":cc-compiler-local", + }, +) + +cc_toolchain( + name = "cc-compiler-local", + all_files = ":crosstool_wrapper_driver_is_not_gcc", + compiler_files = ":crosstool_wrapper_driver_is_not_gcc", + ar_files = ":crosstool_wrapper_driver_is_not_gcc", + as_files = ":crosstool_wrapper_driver_is_not_gcc", + dwp_files = ":empty", + linker_files = ":crosstool_wrapper_driver_is_not_gcc", + objcopy_files = ":empty", + strip_files = ":empty", + # To support linker flags that need to go to the start of command line + # we need the toolchain to support parameter files. Parameter files are + # last on the command line and contain all shared libraries to link, so all + # regular options will be left of them. + supports_param_files = 1, + toolchain_identifier = "local_linux", + toolchain_config = ":cc-compiler-local-config", +) + +cc_toolchain_config( + name = "cc-compiler-local-config", + cpu = "local", + builtin_include_directories = [%{cxx_builtin_include_directories}], + extra_no_canonical_prefixes_flags = [%{extra_no_canonical_prefixes_flags}], + host_compiler_path = "%{host_compiler_path}", + host_compiler_prefix = "%{host_compiler_prefix}", + host_unfiltered_compile_flags = [%{unfiltered_compile_flags}], + linker_bin_path = "%{linker_bin_path}", + compiler = "unknown", +) + +filegroup( + name = "empty", + srcs = [], +) + +filegroup( + name = "crosstool_wrapper_driver_is_not_gcc", + srcs = ["clang/bin/crosstool_wrapper_driver_is_not_gcc"] +) diff --git a/third_party/tsl/third_party/gpus/crosstool/clang/bin/crosstool_wrapper_driver_sycl.tpl b/third_party/tsl/third_party/gpus/crosstool/clang/bin/crosstool_wrapper_driver_sycl.tpl new file mode 100644 index 0000000000000..75474521110df --- /dev/null +++ b/third_party/tsl/third_party/gpus/crosstool/clang/bin/crosstool_wrapper_driver_sycl.tpl @@ -0,0 +1,64 @@ +#!/usr/bin/env python +"""Crosstool wrapper for compiling SYCL program +SYNOPSIS: + crosstool_wrapper_driver_sycl [options passed in by cc_library() + or cc_binary() rule] +DESCRIPTION: + This script is expected to be called by the cc_library() or cc_binary() bazel + rules. When the option "sycl_compile" is present in the list of arguments passed + to this script, it invokes the sycl compiler. When "sycl_compile" is not + present, this wrapper invokes gcc with the input arguments as is. +""" + +from __future__ import print_function +from argparse import ArgumentParser +import os +import subprocess +import sys +import shlex + +CPU_COMPILER = ('%{cpu_compiler}') + +def system(cmd): + """Invokes cmd with os.system()""" + + ret = os.system(cmd) + if os.WIFEXITED(ret): + return os.WEXITSTATUS(ret) + else: + return -os.WTERMSIG(ret) + +def call_compiler(argv): + parser = ArgumentParser() + parser.add_argument('-c', nargs=1, action='append') + parser.add_argument('-o', nargs=1, action='append') + args, leftover = parser.parse_known_args(argv) + + flags = leftover + + common_flags = [] + common_flags.append("-fno-finite-math-only") + common_flags.append("-fno-fast-math") + common_flags.append("-fexceptions") + + in_files, out_files = [], [] + if args.c: + in_files.append('-c') + in_files.extend(args.c[0]) + if args.o: + out_files.append('-o') + out_files.extend(args.o[0]) + flags += (common_flags + in_files + out_files) + print("cmd: ", " ".join([CPU_COMPILER] + flags)) + return subprocess.call([CPU_COMPILER] + flags) + +def main(): + parser = ArgumentParser() + parser = ArgumentParser(fromfile_prefix_chars='@') + parser.add_argument('-sycl_compile', action='store_true') + args, leftover = parser.parse_known_args(sys.argv[1:]) + + return call_compiler(leftover) + +if __name__ == '__main__': + sys.exit(main()) diff --git a/third_party/tsl/third_party/gpus/crosstool/sycl_cc_toolchain_config.bzl.tpl b/third_party/tsl/third_party/gpus/crosstool/sycl_cc_toolchain_config.bzl.tpl new file mode 100644 index 0000000000000..1b4e1c6def04f --- /dev/null +++ b/third_party/tsl/third_party/gpus/crosstool/sycl_cc_toolchain_config.bzl.tpl @@ -0,0 +1,598 @@ +"""cc_toolchain_config rule for configuring SYCL toolchains on Linux.""" + +load( + "@bazel_tools//tools/cpp:cc_toolchain_config_lib.bzl", + "action_config", + "artifact_name_pattern", + "env_entry", + "env_set", + "feature", + "feature_set", + "flag_group", + "flag_set", + "tool", + "tool_path", + "variable_with_value", + "with_feature_set", +) +load("@bazel_tools//tools/build_defs/cc:action_names.bzl", "ACTION_NAMES") + +def all_assembly_actions(): + return [ + ACTION_NAMES.assemble, + ACTION_NAMES.preprocess_assemble, + ] + +def all_compile_actions(): + return [ + ACTION_NAMES.assemble, + ACTION_NAMES.c_compile, + ACTION_NAMES.cpp_compile, + ACTION_NAMES.cpp_header_parsing, + ACTION_NAMES.cpp_module_codegen, + ACTION_NAMES.cpp_module_compile, + ACTION_NAMES.linkstamp_compile, + ACTION_NAMES.preprocess_assemble, + ] + +def all_c_compile_actions(): + return [ + ACTION_NAMES.c_compile, + ] + +def all_cpp_compile_actions(): + return [ + ACTION_NAMES.cpp_compile, + ACTION_NAMES.cpp_header_parsing, + ACTION_NAMES.cpp_module_codegen, + ACTION_NAMES.cpp_module_compile, + ACTION_NAMES.linkstamp_compile, + ] + +def all_preprocessed_actions(): + return [ + ACTION_NAMES.c_compile, + ACTION_NAMES.cpp_compile, + ACTION_NAMES.cpp_header_parsing, + ACTION_NAMES.cpp_module_codegen, + ACTION_NAMES.cpp_module_compile, + ACTION_NAMES.linkstamp_compile, + ACTION_NAMES.preprocess_assemble, + ] + +def all_link_actions(): + return [ + ACTION_NAMES.cpp_link_executable, + ACTION_NAMES.cpp_link_dynamic_library, + ACTION_NAMES.cpp_link_nodeps_dynamic_library, + ] + +def all_executable_link_actions(): + return [ + ACTION_NAMES.cpp_link_executable, + ] + +def all_shared_library_link_actions(): + return [ + ACTION_NAMES.cpp_link_dynamic_library, + ACTION_NAMES.cpp_link_nodeps_dynamic_library, + ] + +def all_archive_actions(): + return [ACTION_NAMES.cpp_link_static_library] + +def all_strip_actions(): + return [ACTION_NAMES.strip] + +def _library_to_link(flag_prefix, value, iterate = None): + return flag_group( + flags = [ + "{}%{{libraries_to_link.{}}}".format( + flag_prefix, + iterate if iterate else "name", + ), + ], + iterate_over = ("libraries_to_link." + iterate if iterate else None), + expand_if_equal = variable_with_value( + name = "libraries_to_link.type", + value = value, + ), + ) + +def _surround_static_library(prefix, suffix): + return [ + flag_group( + flags = [prefix, "%{libraries_to_link.name}", suffix], + expand_if_true = "libraries_to_link.is_whole_archive", + ), + flag_group( + flags = ["%{libraries_to_link.name}"], + expand_if_false = "libraries_to_link.is_whole_archive", + ), + ] + +def _prefix_static_library(prefix): + return [ + flag_group( + flags = ["%{libraries_to_link.name}"], + expand_if_false = "libraries_to_link.is_whole_archive", + ), + flag_group( + flags = [prefix + "%{libraries_to_link.name}"], + expand_if_true = "libraries_to_link.is_whole_archive", + ), + ] + +def _static_library_to_link(alwayslink_prefix, alwayslink_suffix = None): + if alwayslink_suffix: + flag_groups = _surround_static_library(alwayslink_prefix, alwayslink_suffix) + else: + flag_groups = _prefix_static_library(alwayslink_prefix) + return flag_group( + flag_groups = flag_groups, + expand_if_equal = variable_with_value( + name = "libraries_to_link.type", + value = "static_library", + ), + ) + +def _iterate_flag_group(iterate_over, flags = [], flag_groups = []): + return flag_group( + iterate_over = iterate_over, + expand_if_available = iterate_over, + flag_groups = flag_groups, + flags = flags, + ) + +def _libraries_to_link_group(flavour): + if flavour == "linux": + return _iterate_flag_group( + iterate_over = "libraries_to_link", + flag_groups = [ + flag_group( + flags = ["-Wl,--start-lib"], + expand_if_equal = variable_with_value( + name = "libraries_to_link.type", + value = "object_file_group", + ), + ), + _library_to_link("", "object_file_group", "object_files"), + flag_group( + flags = ["-Wl,--end-lib"], + expand_if_equal = variable_with_value( + name = "libraries_to_link.type", + value = "object_file_group", + ), + ), + _library_to_link("", "object_file"), + _library_to_link("", "interface_library"), + _static_library_to_link("-Wl,-whole-archive", "-Wl,-no-whole-archive"), + _library_to_link("-l", "dynamic_library"), + _library_to_link("-l:", "versioned_dynamic_library"), + ], + ) + +def _action_configs_with_tool(path, actions): + return [ + action_config( + action_name = name, + enabled = True, + tools = [tool(path = path)], + ) + for name in actions + ] + +def _action_configs(assembly_path, c_compiler_path, cc_compiler_path, archiver_path, linker_path, strip_path): + return _action_configs_with_tool( + assembly_path, + all_assembly_actions(), + ) + _action_configs_with_tool( + c_compiler_path, + all_c_compile_actions(), + ) + _action_configs_with_tool( + cc_compiler_path, + all_cpp_compile_actions(), + ) + _action_configs_with_tool( + archiver_path, + all_archive_actions(), + ) + _action_configs_with_tool( + linker_path, + all_link_actions(), + ) + _action_configs_with_tool( + strip_path, + all_strip_actions(), + ) + +def _tool_paths(cpu, ctx): + if cpu == "local": + return [ + tool_path(name = "gcc", path = ctx.attr.host_compiler_path), + tool_path(name = "ar", path = ctx.attr.host_compiler_prefix + "/ar"), + tool_path(name = "compat-ld", path = ctx.attr.host_compiler_prefix + "/ld"), + tool_path(name = "cpp", path = ctx.attr.host_compiler_prefix + "/cpp"), + tool_path(name = "dwp", path = ctx.attr.host_compiler_prefix + "/dwp"), + tool_path(name = "gcov", path = ctx.attr.host_compiler_prefix + "/gcov"), + tool_path(name = "ld", path = ctx.attr.host_compiler_prefix + "/ld"), + tool_path(name = "nm", path = ctx.attr.host_compiler_prefix + "/nm"), + tool_path(name = "objcopy", path = ctx.attr.host_compiler_prefix + "/objcopy"), + tool_path(name = "objdump", path = ctx.attr.host_compiler_prefix + "/objdump"), + tool_path(name = "strip", path = ctx.attr.host_compiler_prefix + "/strip"), + ] + else: + fail("Unreachable") + +def _sysroot_group(): + return flag_group( + flags = ["--sysroot=%{sysroot}"], + expand_if_available = "sysroot", + ) + +def _no_canonical_prefixes_group(extra_flags): + return flag_group( + flags = [ + "-no-canonical-prefixes", + ] + extra_flags, + ) + +def _nologo(): + return flag_group(flags = ["/nologo"]) + +def _features(cpu, compiler, ctx): + if cpu == "local": + return [ + feature(name = "no_legacy_features"), + feature( + name = "all_compile_flags", + enabled = True, + flag_sets = [ + flag_set( + actions = all_compile_actions(), + flag_groups = [ + flag_group( + flags = ["-MD", "-MF", "%{dependency_file}"], + expand_if_available = "dependency_file", + ), + flag_group( + flags = ["-gsplit-dwarf"], + expand_if_available = "per_object_debug_info_file", + ), + ], + ), + flag_set( + actions = all_preprocessed_actions(), + flag_groups = [ + flag_group( + flags = ["-frandom-seed=%{output_file}"], + expand_if_available = "output_file", + ), + _iterate_flag_group( + flags = ["-D%{preprocessor_defines}"], + iterate_over = "preprocessor_defines", + ), + _iterate_flag_group( + flags = ["-include", "%{includes}"], + iterate_over = "includes", + ), + _iterate_flag_group( + flags = ["-iquote", "%{quote_include_paths}"], + iterate_over = "quote_include_paths", + ), + _iterate_flag_group( + flags = ["-I%{include_paths}"], + iterate_over = "include_paths", + ), + _iterate_flag_group( + flags = ["-isystem", "%{system_include_paths}"], + iterate_over = "system_include_paths", + ), + _iterate_flag_group( + flags = ["-F", "%{framework_include_paths}"], + iterate_over = "framework_include_paths", + ), + ] + ([ + flag_group(flags = ctx.attr.host_unfiltered_compile_flags), + ] if ctx.attr.host_unfiltered_compile_flags else []), + ), + flag_set( + actions = all_cpp_compile_actions(), + flag_groups = [ + flag_group(flags = [ + "-fmerge-all-constants", + ]), + ] if compiler == "clang" else [], + ), + flag_set( + actions = all_compile_actions(), + flag_groups = [ + flag_group( + flags = [ + "-Wno-builtin-macro-redefined", + "-D__DATE__=\"redacted\"", + "-D__TIMESTAMP__=\"redacted\"", + "-D__TIME__=\"redacted\"", + ], + ), + flag_group( + flags = ["-fPIC"], + expand_if_available = "pic", + ), + flag_group( + flags = ["-fPIE"], + expand_if_not_available = "pic", + ), + flag_group( + flags = [ + "-U_FORTIFY_SOURCE", + "-D_FORTIFY_SOURCE=1", + "-fstack-protector", + "-Wall", + ] + ctx.attr.host_compiler_warnings + [ + "-fno-omit-frame-pointer", + ], + ), + _no_canonical_prefixes_group( + ctx.attr.extra_no_canonical_prefixes_flags, + ), + ], + ), + flag_set( + actions = all_compile_actions(), + flag_groups = [flag_group(flags = ["-DNDEBUG"])], + with_features = [with_feature_set(features = ["disable-assertions"])], + ), + flag_set( + actions = all_compile_actions(), + flag_groups = [ + flag_group( + flags = [ + "-g0", + "-O2", + "-ffunction-sections", + "-fdata-sections", + ], + ), + ], + with_features = [with_feature_set(features = ["opt"])], + ), + flag_set( + actions = all_compile_actions(), + flag_groups = [flag_group(flags = ["-g"])], + with_features = [with_feature_set(features = ["dbg"])], + ), + ] + [ + flag_set( + actions = all_compile_actions(), + flag_groups = [ + _iterate_flag_group( + flags = ["%{user_compile_flags}"], + iterate_over = "user_compile_flags", + ), + _sysroot_group(), + flag_group( + expand_if_available = "source_file", + flags = ["-c", "%{source_file}"], + ), + flag_group( + expand_if_available = "output_assembly_file", + flags = ["-S"], + ), + flag_group( + expand_if_available = "output_preprocess_file", + flags = ["-E"], + ), + flag_group( + expand_if_available = "output_file", + flags = ["-o", "%{output_file}"], + ), + ], + ), + ], + ), + feature( + name = "all_archive_flags", + enabled = True, + flag_sets = [ + flag_set( + actions = all_archive_actions(), + flag_groups = [ + flag_group( + expand_if_available = "linker_param_file", + flags = ["@%{linker_param_file}"], + ), + flag_group(flags = ["rcsD"]), + flag_group( + flags = ["%{output_execpath}"], + expand_if_available = "output_execpath", + ), + flag_group( + iterate_over = "libraries_to_link", + flag_groups = [ + flag_group( + flags = ["%{libraries_to_link.name}"], + expand_if_equal = variable_with_value( + name = "libraries_to_link.type", + value = "object_file", + ), + ), + flag_group( + flags = ["%{libraries_to_link.object_files}"], + iterate_over = "libraries_to_link.object_files", + expand_if_equal = variable_with_value( + name = "libraries_to_link.type", + value = "object_file_group", + ), + ), + ], + expand_if_available = "libraries_to_link", + ), + ], + ), + ], + ), + feature( + name = "all_link_flags", + enabled = True, + flag_sets = [ + flag_set( + actions = all_shared_library_link_actions(), + flag_groups = [flag_group(flags = ["-shared"])], + ), + flag_set( + actions = all_link_actions(), + flag_groups = ([ + flag_group(flags = ["-Wl,-no-as-needed"]), + ] if cpu == "local" else []) + ([ + flag_group(flags = ["-B" + ctx.attr.linker_bin_path]), + ] if ctx.attr.linker_bin_path else []) + [ + flag_group( + flags = ["@%{linker_param_file}"], + expand_if_available = "linker_param_file", + ), + _iterate_flag_group( + flags = ["%{linkstamp_paths}"], + iterate_over = "linkstamp_paths", + ), + flag_group( + flags = ["-o", "%{output_execpath}"], + expand_if_available = "output_execpath", + ), + _iterate_flag_group( + flags = ["-L%{library_search_directories}"], + iterate_over = "library_search_directories", + ), + _iterate_flag_group( + iterate_over = "runtime_library_search_directories", + flags = [ + "-Wl,-rpath,$ORIGIN/%{runtime_library_search_directories}", + ] if cpu == "local" else [ + "-Wl,-rpath,@loader_path/%{runtime_library_search_directories}", + ], + ), + _libraries_to_link_group("linux"), + _iterate_flag_group( + flags = ["%{user_link_flags}"], + iterate_over = "user_link_flags", + ), + flag_group( + flags = ["-Wl,--gdb-index"], + expand_if_available = "is_using_fission", + ), + flag_group( + flags = ["-Wl,-S"], + expand_if_available = "strip_debug_symbols", + ), + flag_group(flags = ["-lstdc++"]), + _no_canonical_prefixes_group( + ctx.attr.extra_no_canonical_prefixes_flags, + ), + ], + ), + flag_set( + actions = all_executable_link_actions(), + flag_groups = [flag_group(flags = ["-pie"])], + ), + ] + ([ + flag_set( + actions = all_link_actions(), + flag_groups = [flag_group(flags = [ + "-Wl,-z,relro,-z,now", + ])], + ), + ]) + ([ + flag_set( + actions = all_link_actions(), + flag_groups = [ + flag_group(flags = ["-Wl,--gc-sections"]), + flag_group( + flags = ["-Wl,--build-id=md5", "-Wl,--hash-style=gnu"], + ), + ], + ), + ]) + [ + flag_set( + actions = all_link_actions(), + flag_groups = [ + _sysroot_group(), + ], + ), + ], + ), + feature(name = "disable-assertions"), + feature( + name = "opt", + implies = ["disable-assertions"], + ), + feature(name = "fastbuild"), + feature(name = "dbg"), + feature(name = "supports_dynamic_linker", enabled = True), + feature(name = "pic", enabled = True), + feature(name = "supports_pic", enabled = True), + feature(name = "has_configured_linker_path", enabled = True), + ] + else: + fail("Unreachable") + +def _impl(ctx): + cpu = ctx.attr.cpu + compiler = ctx.attr.compiler + + if (cpu == "local"): + toolchain_identifier = "local_linux" + target_cpu = "local" + target_libc = "local" + action_configs = _action_configs( + assembly_path = ctx.attr.host_compiler_path, + c_compiler_path = ctx.attr.host_compiler_path, + cc_compiler_path = ctx.attr.host_compiler_path, + archiver_path = ctx.attr.host_compiler_prefix + "/ar", + linker_path = ctx.attr.host_compiler_path, + strip_path = ctx.attr.host_compiler_prefix + "/strip", + ) + artifact_name_patterns = [] + else: + fail("Unreachable") + + out = ctx.actions.declare_file(ctx.label.name) + ctx.actions.write(out, "Fake executable") + return [ + cc_common.create_cc_toolchain_config_info( + ctx = ctx, + features = _features(cpu, compiler, ctx), + action_configs = action_configs, + artifact_name_patterns = artifact_name_patterns, + cxx_builtin_include_directories = ctx.attr.builtin_include_directories, + toolchain_identifier = toolchain_identifier, + host_system_name = "local", + target_system_name = "local", + target_cpu = target_cpu, + target_libc = target_libc, + compiler = compiler, + abi_version = "local", + abi_libc_version = "local", + tool_paths = _tool_paths(cpu, ctx), + make_variables = [], + builtin_sysroot = ctx.attr.builtin_sysroot, + cc_target_os = None, + ), + DefaultInfo( + executable = out, + ), + ] + +cc_toolchain_config = rule( + implementation = _impl, + attrs = { + "cpu": attr.string(mandatory = True, values = ["local"]), + "compiler": attr.string(values = ["clang", "unknown"], default = "unknown"), + "builtin_include_directories": attr.string_list(), + "extra_no_canonical_prefixes_flags": attr.string_list(), + "host_compiler_path": attr.string(), + "host_compiler_prefix": attr.string(), + "host_compiler_warnings": attr.string_list(), + "host_unfiltered_compile_flags": attr.string_list(), + "linker_bin_path": attr.string(), + "builtin_sysroot": attr.string(), + }, + provides = [CcToolchainConfigInfo], + executable = True, +) diff --git a/third_party/tsl/third_party/gpus/find_sycl_config.py b/third_party/tsl/third_party/gpus/find_sycl_config.py new file mode 100644 index 0000000000000..5db84c32ea908 --- /dev/null +++ b/third_party/tsl/third_party/gpus/find_sycl_config.py @@ -0,0 +1,139 @@ +# Copyright 2024 The TensorFlow Authors. All Rights Reserved. +# +# 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 +# +# http://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. +# ============================================================================== +"""Prints SYCL library and header directories and versions found on the system. + +The script searches for SYCL library and header files on the system, inspects +them to determine their version and prints the configuration to stdout. The path +to inspect is specified through an environment variable (SYCL_PATH). If no valid +configuration is found, the script prints to stderr and returns an error code. +The script takes the directory specified by the SYCL_PATH environment variable. +The script looks for headers and library files in a hard-coded set of +subdirectories from base path of the specified directory. If SYCL_PATH is not +specified, then "/opt/sycl" is used as it default value +""" + +import io +import os +import re +import sys + + +class ConfigError(Exception): + pass + + +def _get_default_sycl_toolkit_path(): + return "/opt/intel/oneapi/compiler/latest" + + +def _get_toolkit_path(): + """Determines and returns the SYCL installation path.""" + sycl_toolkit_path = None + sycl_toolkit_path = _get_default_sycl_toolkit_path() + if "SYCL_TOOLKIT_PATH" in os.environ: + sycl_toolkit_path = os.environ["SYCL_TOOLKIT_PATH"] + return os.path.realpath(sycl_toolkit_path) + + +def _get_basekit_path(): + return _get_toolkit_path().split("/compiler/")[0] + + +def _get_basekit_version(): + return _get_toolkit_path().split("/compiler/")[1].split("/")[0] + + +def _get_composite_version_number(major, minor, patch): + return 10000 * major + 100 * minor + patch + + +def _get_header_version(path, name): + """Returns preprocessor defines in C header file.""" + for line in io.open(path, "r", encoding="utf-8"): + match = re.match(r"#define %s +(\d+)" % name, line) + if match: + value = match.group(1) + return int(value) + + raise ConfigError( + '#define "{}" is either\n'.format(name) + + " not present in file {} OR\n".format(path) + + " its value is not an integer literal" + ) + + +def _find_sycl_config(basekit_path): + # pylint: disable=missing-function-docstring + + def sycl_version_numbers(path): + possible_version_files = [ + "compiler/latest/linux/include/sycl/version.hpp", + "compiler/latest/include/sycl/version.hpp", + ] + version_file = None + for f in possible_version_files: + version_file_path = os.path.join(path, f) + if os.path.exists(version_file_path): + version_file = version_file_path + break + if not version_file: + raise ConfigError( + "SYCL version file not found in {}".format(possible_version_files) + ) + + major = _get_header_version(version_file, "__LIBSYCL_MAJOR_VERSION") + minor = _get_header_version(version_file, "__LIBSYCL_MINOR_VERSION") + patch = _get_header_version(version_file, "__LIBSYCL_PATCH_VERSION") + return major, minor, patch + + major, minor, patch = sycl_version_numbers(basekit_path) + + sycl_config = { + "sycl_version_number": _get_composite_version_number(major, minor, patch), + "sycl_basekit_version_number": _get_basekit_version(), + } + + return sycl_config + + +def find_sycl_config(): + """Returns a dictionary of SYCL components config info.""" + basekit_path = _get_basekit_path() + toolkit_path = _get_toolkit_path() + if not os.path.exists(basekit_path): + raise ConfigError( + 'Specified SYCL_TOOLKIT_PATH "{}" does not exist'.format(basekit_path) + ) + + result = {} + + result["sycl_basekit_path"] = basekit_path + result["sycl_toolkit_path"] = toolkit_path + result.update(_find_sycl_config(basekit_path)) + + return result + + +def main(): + try: + for key, value in sorted(find_sycl_config().items()): + print("%s: %s" % (key, value)) + except ConfigError as e: + sys.stderr.write("\nERROR: {}\n\n".format(str(e))) + sys.exit(1) + + +if __name__ == "__main__": + main() diff --git a/third_party/tsl/third_party/gpus/rocm/build_defs.bzl.tpl b/third_party/tsl/third_party/gpus/rocm/build_defs.bzl.tpl index 339733755d6f1..231a8a1913e7a 100644 --- a/third_party/tsl/third_party/gpus/rocm/build_defs.bzl.tpl +++ b/third_party/tsl/third_party/gpus/rocm/build_defs.bzl.tpl @@ -39,6 +39,10 @@ def rocm_version_number(): return %{rocm_version_number} def if_gpu_is_configured(if_true, if_false = []): + """Tests if ROCm or CUDA or SYCL was enabled during the configure process.""" + return select({"//conditions:default": %{gpu_is_configured}}) + +def if_cuda_or_rocm(if_true, if_false = []): """Tests if ROCm or CUDA was enabled during the configure process. Unlike if_rocm() or if_cuda(), this does not require that we are building @@ -46,7 +50,7 @@ def if_gpu_is_configured(if_true, if_false = []): code to depend on ROCm or CUDA libraries. """ - return select({"//conditions:default": %{gpu_is_configured}}) + return select({"//conditions:default": %{cuda_or_rocm}}) def if_rocm_is_configured(x): """Tests if the ROCm was enabled during the configure process. diff --git a/third_party/tsl/third_party/gpus/rocm_configure.bzl b/third_party/tsl/third_party/gpus/rocm_configure.bzl index 66a53276a194c..89a5c21ebf928 100644 --- a/third_party/tsl/third_party/gpus/rocm_configure.bzl +++ b/third_party/tsl/third_party/gpus/rocm_configure.bzl @@ -29,6 +29,10 @@ load( "make_copy_files_rule", "to_list_of_strings", ) +load( + ":sycl_configure.bzl", + "enable_sycl", +) _GCC_HOST_COMPILER_PATH = "GCC_HOST_COMPILER_PATH" _GCC_HOST_COMPILER_PREFIX = "GCC_HOST_COMPILER_PREFIX" @@ -450,7 +454,8 @@ def _create_dummy_repository(repository_ctx): "rocm:build_defs.bzl", { "%{rocm_is_configured}": "False", - "%{gpu_is_configured}": "if_true" if enable_cuda(repository_ctx) else "if_false", + "%{gpu_is_configured}": "if_true" if enable_cuda(repository_ctx) or enable_sycl(repository_ctx) else "if_false", + "%{cuda_or_rocm}": "if_true" if enable_cuda(repository_ctx) else "if_false", "%{rocm_extra_copts}": "[]", "%{rocm_gpu_architectures}": "[]", "%{rocm_version_number}": "0", @@ -637,6 +642,7 @@ def _create_local_rocm_repository(repository_ctx): { "%{rocm_is_configured}": "True", "%{gpu_is_configured}": "if_true", + "%{cuda_or_rocm}": "if_true", "%{rocm_extra_copts}": _compute_rocm_extra_copts( repository_ctx, rocm_config.amdgpu_targets, @@ -766,6 +772,7 @@ def _create_remote_rocm_repository(repository_ctx, remote_config_repo): { "%{rocm_is_configured}": "True", "%{gpu_is_configured}": "if_true", + "%{cuda_or_rocm}": "if_true", "%{rocm_extra_copts}": _compute_rocm_extra_copts( repository_ctx, [], #_compute_capabilities(repository_ctx) diff --git a/third_party/tsl/third_party/gpus/sycl/BUILD b/third_party/tsl/third_party/gpus/sycl/BUILD new file mode 100644 index 0000000000000..e69de29bb2d1d diff --git a/third_party/tsl/third_party/gpus/sycl/BUILD.tpl b/third_party/tsl/third_party/gpus/sycl/BUILD.tpl new file mode 100644 index 0000000000000..effbbc8879939 --- /dev/null +++ b/third_party/tsl/third_party/gpus/sycl/BUILD.tpl @@ -0,0 +1,63 @@ +package(default_visibility = ["//visibility:public"]) + +config_setting( + name = "using_sycl", + values = { + "define": "using_sycl=true", + }, +) + +cc_library( + name = "sycl_headers", + hdrs = [ + %{sycl_headers} + ], + includes = [ + ".", + "sycl/include", + "sycl/include/sycl", + ], + visibility = ["//visibility:public"], +) + +cc_library( + name = "sycl", + srcs = [ + %{core_sycl_libs} + ], + data = [ + %{core_sycl_libs} + ], + includes = [ + ".", + "sycl/include", + ], + linkopts = ["-lze_loader"], + linkstatic = 1, + visibility = ["//visibility:public"], +) + +cc_library( + name = "mkl", + srcs = [ + "sycl/lib/%{mkl_intel_ilp64_lib}", + "sycl/lib/%{mkl_sequential_lib}", + "sycl/lib/%{mkl_core_lib}", + %{mkl_sycl_libs} + ], + data = [ + "sycl/lib/%{mkl_intel_ilp64_lib}", + "sycl/lib/%{mkl_sequential_lib}", + "sycl/lib/%{mkl_core_lib}", + %{mkl_sycl_libs} + ], + includes = [ + ".", + "sycl/include", + ], + # linkopts = ["-Wl,-Bstatic,-lsvml,-lirng,-limf,-lirc,-lirc_s,-Bdynamic"], + linkstatic = 1, + visibility = ["//visibility:public"], +) + +%{copy_rules} diff --git a/third_party/tsl/third_party/gpus/sycl/build_defs.bzl.tpl b/third_party/tsl/third_party/gpus/sycl/build_defs.bzl.tpl new file mode 100644 index 0000000000000..d678f84ca88a0 --- /dev/null +++ b/third_party/tsl/third_party/gpus/sycl/build_defs.bzl.tpl @@ -0,0 +1,39 @@ +# Macros for building SYCL code. +def if_sycl(if_true, if_false = []): + """Shorthand for select()'ing on whether we're building with SYCL. + + Returns a select statement which evaluates to if_true if we're building + with SYCL enabled. Otherwise, the select statement evaluates to if_false. + + """ + return select({ + "@local_config_sycl//sycl:using_sycl": if_true, + "//conditions:default": if_false, + }) + +def sycl_default_copts(): + """Default options for all SYCL compilations.""" + return if_sycl(["-x", "sycl"]) + +def sycl_build_is_configured(): + """Returns true if SYCL compiler was enabled during the configure process.""" + return %{sycl_build_is_configured} + +def if_sycl_is_configured(x): + """Tests if the SYCL was enabled during the configure process. + + Unlike if_sycl(), this does not require that we are building with + --config=sycl. Used to allow non-SYCL code to depend on SYCL libraries. + """ + if %{sycl_is_configured}: + return select({"//conditions:default": x}) + return select({"//conditions:default": []}) + +def if_sycl_build_is_configured(x, y): + if sycl_build_is_configured(): + return x + return y + +def sycl_library(copts = [], **kwargs): + """Wrapper over cc_library which adds default SYCL options.""" + native.cc_library(copts = sycl_default_copts() + copts, **kwargs) diff --git a/third_party/tsl/third_party/gpus/sycl_configure.bzl b/third_party/tsl/third_party/gpus/sycl_configure.bzl new file mode 100644 index 0000000000000..05330b2fe5319 --- /dev/null +++ b/third_party/tsl/third_party/gpus/sycl_configure.bzl @@ -0,0 +1,531 @@ +"""Repository rule for SYCL autoconfiguration. +`sycl_configure` depends on the following environment variables: + * `TF_NEED_SYCL`: Whether to enable building with SYCL. + * `GCC_HOST_COMPILER_PATH`: The GCC host compiler path +""" + +load( + "//third_party/remote_config:common.bzl", + "err_out", + "execute", + "files_exist", + "get_bash_bin", + "get_host_environ", + "get_python_bin", + "raw_exec", + "realpath", + "which", +) +load( + ":cuda_configure.bzl", + "make_copy_dir_rule", + "make_copy_files_rule", + "to_list_of_strings", +) + +_GCC_HOST_COMPILER_PATH = "GCC_HOST_COMPILER_PATH" +_GCC_HOST_COMPILER_PREFIX = "GCC_HOST_COMPILER_PREFIX" + +def _mkl_path(sycl_config): + return sycl_config.sycl_basekit_path + "/mkl/" + sycl_config.sycl_basekit_version_number + +def _sycl_header_path(repository_ctx, sycl_config, bash_bin): + sycl_header_path = sycl_config.sycl_basekit_path + "/compiler/" + sycl_config.sycl_basekit_version_number + include_dir = sycl_header_path + "/include" + if not files_exist(repository_ctx, [include_dir], bash_bin)[0]: + sycl_header_path = sycl_header_path + "/linux" + include_dir = sycl_header_path + "/include" + if not files_exist(repository_ctx, [include_dir], bash_bin)[0]: + auto_configure_fail("Cannot find sycl headers in {}".format(include_dir)) + return sycl_header_path + +def _sycl_include_path(repository_ctx, sycl_config, bash_bin): + """Generates the cxx_builtin_include_directory entries for sycl inc dirs. + Args: + repository_ctx: The repository context. + sycl_config: The path to the gcc host compiler. + Returns: + A string containing the Starlark string for each of the gcc + host compiler include directories, which can be added to the CROSSTOOL + file. + """ + inc_dirs = [] + + inc_dirs.append(_mkl_path(sycl_config) + "/include") + inc_dirs.append(_sycl_header_path(repository_ctx, sycl_config, bash_bin) + "/include") + inc_dirs.append(_sycl_header_path(repository_ctx, sycl_config, bash_bin) + "/include/sycl") + + return inc_dirs + +def enable_sycl(repository_ctx): + """Returns whether to build with SYCL support.""" + return int(get_host_environ(repository_ctx, "TF_NEED_SYCL", False)) + +def auto_configure_fail(msg): + """Output failure message when auto configuration fails.""" + red = "\033[0;31m" + no_color = "\033[0m" + fail("\n%sAuto-Configuration Error:%s %s\n" % (red, no_color, msg)) + +def find_cc(repository_ctx): + """Find the C++ compiler.""" + + # Return a dummy value for GCC detection here to avoid error + target_cc_name = "gcc" + cc_path_envvar = _GCC_HOST_COMPILER_PATH + cc_name = target_cc_name + + cc_name_from_env = get_host_environ(repository_ctx, cc_path_envvar) + if cc_name_from_env: + cc_name = cc_name_from_env + if cc_name.startswith("/"): + # Absolute path, maybe we should make this supported by our which function. + return cc_name + cc = which(repository_ctx, cc_name) + if cc == None: + fail(("Cannot find {}, either correct your path or set the {}" + + " environment variable").format(target_cc_name, cc_path_envvar)) + return cc + +def find_sycl_root(repository_ctx, sycl_config): + sycl_name = str(repository_ctx.path(sycl_config.sycl_toolkit_path.strip()).realpath) + if sycl_name.startswith("/"): + return sycl_name + fail("Cannot find SYCL compiler, please correct your path") + +def find_sycl_include_path(repository_ctx, sycl_config): + base_path = find_sycl_root(repository_ctx, sycl_config) + bin_path = repository_ctx.path(base_path + "/" + "bin" + "/" + "icpx") + icpx_extra = "" + if not bin_path.exists: + bin_path = repository_ctx.path(base_path + "/" + "bin" + "/" + "clang") + if not bin_path.exists: + fail("Cannot find SYCL compiler, please correct your path") + else: + icpx_extra = "-fsycl" + gcc_path = repository_ctx.which("gcc") + gcc_install_dir = repository_ctx.execute([gcc_path, "-print-libgcc-file-name"]) + gcc_install_dir_opt = "--gcc-install-dir=" + str(repository_ctx.path(gcc_install_dir.stdout.strip()).dirname) + cmd_out = repository_ctx.execute([bin_path, icpx_extra, gcc_install_dir_opt, "-xc++", "-E", "-v", "/dev/null", "-o", "/dev/null"]) + outlist = cmd_out.stderr.split("\n") + real_base_path = str(repository_ctx.path(base_path).realpath).strip() + include_dirs = [] + for l in outlist: + if l.startswith(" ") and l.strip().startswith("/") and str(repository_ctx.path(l.strip()).realpath) not in include_dirs: + include_dirs.append(str(repository_ctx.path(l.strip()).realpath)) + return include_dirs + +def _lib_name(lib, version = "", static = False): + """Constructs the name of a library on Linux. + Args: + lib: The name of the library, such as "mkl" + version: The version of the library. + static: True the library is static or False if it is a shared object. + Returns: + The platform-specific name of the library. + """ + if static: + return "lib%s.a" % lib + else: + if version: + version = ".%s" % version + return "lib%s.so%s" % (lib, version) + +def _sycl_lib_paths(repository_ctx, lib, basedir): + file_name = _lib_name(lib, version = "", static = False) + return [ + repository_ctx.path("%s/lib/%s" % (basedir, file_name)), + repository_ctx.path("%s/lib/intel64/%s" % (basedir, file_name)), + ] + +def _batch_files_exist(repository_ctx, libs_paths, bash_bin): + all_paths = [] + for _, lib_paths in libs_paths: + for lib_path in lib_paths: + all_paths.append(lib_path) + return files_exist(repository_ctx, all_paths, bash_bin) + +def _select_sycl_lib_paths(repository_ctx, libs_paths, bash_bin): + test_results = _batch_files_exist(repository_ctx, libs_paths, bash_bin) + + libs = {} + i = 0 + for name, lib_paths in libs_paths: + selected_path = None + for path in lib_paths: + if test_results[i] and selected_path == None: + # For each lib select the first path that exists. + selected_path = path + i = i + 1 + if selected_path == None: + auto_configure_fail("Cannot find sycl library %s in %s" % (name, path)) + + libs[name] = struct(file_name = selected_path.basename, path = realpath(repository_ctx, selected_path, bash_bin)) + + return libs + +def _find_libs(repository_ctx, sycl_config, bash_bin): + """Returns the SYCL libraries on the system. + Args: + repository_ctx: The repository context. + sycl_config: The SYCL config as returned by _get_sycl_config + bash_bin: the path to the bash interpreter + Returns: + Map of library names to structs of filename and path + """ + mkl_path = _mkl_path(sycl_config) + sycl_path = _sycl_header_path(repository_ctx, sycl_config, bash_bin) + libs_paths = [ + (name, _sycl_lib_paths(repository_ctx, name, path)) + for name, path in [ + ("sycl", sycl_path), + ("OpenCL", sycl_path), + ("mkl_intel_ilp64", mkl_path), + ("mkl_sequential", mkl_path), + ("mkl_core", mkl_path), + ] + ] + if sycl_config.sycl_basekit_version_number < "2024": + libs_paths.append(("mkl_sycl", _sycl_lib_paths(repository_ctx, "mkl_sycl", mkl_path))) + else: + libs_paths.append(("mkl_sycl_blas", _sycl_lib_paths(repository_ctx, "mkl_sycl_blas", mkl_path))) + libs_paths.append(("mkl_sycl_lapack", _sycl_lib_paths(repository_ctx, "mkl_sycl_lapack", mkl_path))) + libs_paths.append(("mkl_sycl_sparse", _sycl_lib_paths(repository_ctx, "mkl_sycl_sparse", mkl_path))) + libs_paths.append(("mkl_sycl_dft", _sycl_lib_paths(repository_ctx, "mkl_sycl_dft", mkl_path))) + libs_paths.append(("mkl_sycl_vm", _sycl_lib_paths(repository_ctx, "mkl_sycl_vm", mkl_path))) + libs_paths.append(("mkl_sycl_rng", _sycl_lib_paths(repository_ctx, "mkl_sycl_rng", mkl_path))) + libs_paths.append(("mkl_sycl_stats", _sycl_lib_paths(repository_ctx, "mkl_sycl_stats", mkl_path))) + libs_paths.append(("mkl_sycl_data_fitting", _sycl_lib_paths(repository_ctx, "mkl_sycl_data_fitting", mkl_path))) + return _select_sycl_lib_paths(repository_ctx, libs_paths, bash_bin) + +def find_sycl_config(repository_ctx): + """Returns SYCL config dictionary from running find_sycl_config.py""" + python_bin = get_python_bin(repository_ctx) + exec_result = execute(repository_ctx, [python_bin, repository_ctx.attr._find_sycl_config]) + if exec_result.return_code: + auto_configure_fail("Failed to run find_sycl_config.py: %s" % err_out(exec_result)) + + # Parse the dict from stdout. + return dict([tuple(x.split(": ")) for x in exec_result.stdout.splitlines()]) + +def _get_sycl_config(repository_ctx, bash_bin): + """Detects and returns information about the SYCL installation on the system. + Args: + repository_ctx: The repository context. + bash_bin: the path to the path interpreter + """ + config = find_sycl_config(repository_ctx) + sycl_basekit_path = config["sycl_basekit_path"] + sycl_toolkit_path = config["sycl_toolkit_path"] + sycl_version_number = config["sycl_version_number"] + sycl_basekit_version_number = config["sycl_basekit_version_number"] + return struct( + sycl_basekit_path = sycl_basekit_path, + sycl_toolkit_path = sycl_toolkit_path, + sycl_version_number = sycl_version_number, + sycl_basekit_version_number = sycl_basekit_version_number, + ) + +def _tpl_path(repository_ctx, labelname): + return repository_ctx.path(Label("//third_party/gpus/%s.tpl" % labelname)) + +def _tpl(repository_ctx, tpl, substitutions = {}, out = None): + if not out: + out = tpl.replace(":", "/") + repository_ctx.template( + out, + _tpl_path(repository_ctx, tpl), + substitutions, + ) + +_INC_DIR_MARKER_BEGIN = "#include <...>" + +def _cxx_inc_convert(path): + """Convert path returned by cc -E xc++ in a complete path.""" + path = path.strip() + return path + +def _normalize_include_path(repository_ctx, path): + """Normalizes include paths before writing them to the crosstool. + If path points inside the 'crosstool' folder of the repository, a relative + path is returned. + If path points outside the 'crosstool' folder, an absolute path is returned. + """ + path = str(repository_ctx.path(path)) + crosstool_folder = str(repository_ctx.path(".").get_child("crosstool")) + + if path.startswith(crosstool_folder): + # We drop the path to "$REPO/crosstool" and a trailing path separator. + return "\"" + path[len(crosstool_folder) + 1:] + "\"" + return "\"" + path + "\"" + +def _get_cxx_inc_directories_impl(repository_ctx, cc, lang_is_cpp): + """Compute the list of default C or C++ include directories.""" + if lang_is_cpp: + lang = "c++" + else: + lang = "c" + + result = raw_exec(repository_ctx, [ + cc, + "-no-canonical-prefixes", + "-E", + "-x" + lang, + "-", + "-v", + ]) + stderr = err_out(result) + index1 = stderr.find(_INC_DIR_MARKER_BEGIN) + if index1 == -1: + return [] + index1 = stderr.find("\n", index1) + if index1 == -1: + return [] + index2 = stderr.rfind("\n ") + if index2 == -1 or index2 < index1: + return [] + index2 = stderr.find("\n", index2 + 1) + if index2 == -1: + inc_dirs = stderr[index1 + 1:] + else: + inc_dirs = stderr[index1 + 1:index2].strip() + + return [ + str(repository_ctx.path(_cxx_inc_convert(p))) + for p in inc_dirs.split("\n") + ] + +def get_cxx_inc_directories(repository_ctx, cc): + """Compute the list of default C and C++ include directories.""" + + # For some reason `clang -xc` sometimes returns include paths that are + # different from the ones from `clang -xc++`. (Symlink and a dir) + # So we run the compiler with both `-xc` and `-xc++` and merge resulting lists + includes_cpp = _get_cxx_inc_directories_impl(repository_ctx, cc, True) + includes_c = _get_cxx_inc_directories_impl(repository_ctx, cc, False) + + includes_cpp_set = depset(includes_cpp) + return includes_cpp + [ + inc + for inc in includes_c + if inc not in includes_cpp_set.to_list() + ] + +_DUMMY_CROSSTOOL_BZL_FILE = """ +def error_gpu_disabled(): + fail("ERROR: Building with --config=sycl but TensorFlow is not configured " + + "to build with GPU support. Please re-run ./configure and enter 'Y' " + + "at the prompt to build with GPU support.") + native.genrule( + name = "error_gen_crosstool", + outs = ["CROSSTOOL"], + cmd = "echo 'Should not be run.' && exit 1", + ) + native.filegroup( + name = "crosstool", + srcs = [":CROSSTOOL"], + output_licenses = ["unencumbered"], + ) +""" + +_DUMMY_CROSSTOOL_BUILD_FILE = """ +load("//crosstool:error_gpu_disabled.bzl", "error_gpu_disabled") +error_gpu_disabled() +""" + +def _create_dummy_repository(repository_ctx): + # Set up BUILD file for sycl/. + _tpl(repository_ctx, "sycl:build_defs.bzl") + _tpl( + repository_ctx, + "sycl:BUILD", + { + "%{mkl_intel_ilp64_lib}": _lib_name("mkl_intel_ilp64"), + "%{mkl_sequential_lib}": _lib_name("mkl_sequential"), + "%{mkl_core_lib}": _lib_name("mkl_core"), + "%{mkl_sycl_libs}": "", + "%{core_sycl_libs}": "", + "%{copy_rules}": "", + "%{sycl_headers}": "", + }, + ) + + # If sycl_configure is not configured to build with SYCL support, and the user + # attempts to build with --config=sycl, add a dummy build rule to intercept + # this and fail with an actionable error message. + repository_ctx.file( + "crosstool/error_gpu_disabled.bzl", + _DUMMY_CROSSTOOL_BZL_FILE, + ) + repository_ctx.file("crosstool/BUILD", _DUMMY_CROSSTOOL_BUILD_FILE) + + _tpl( + repository_ctx, + "sycl:build_defs.bzl", + { + "%{sycl_is_configured}": "False", + "%{sycl_build_is_configured}": "False", + }, + ) + +def _create_local_sycl_repository(repository_ctx): + tpl_paths = {labelname: _tpl_path(repository_ctx, labelname) for labelname in [ + "sycl:build_defs.bzl", + "sycl:BUILD", + "crosstool:BUILD.sycl", + "crosstool:sycl_cc_toolchain_config.bzl", + "crosstool:clang/bin/crosstool_wrapper_driver_sycl", + ]} + + bash_bin = get_bash_bin(repository_ctx) + sycl_config = _get_sycl_config(repository_ctx, bash_bin) + + # Copy header and library files to execroot. + copy_rules = [ + make_copy_dir_rule( + repository_ctx, + name = "sycl-include", + src_dir = _sycl_header_path(repository_ctx, sycl_config, bash_bin) + "/include", + out_dir = "sycl/include", + ), + ] + copy_rules.append(make_copy_dir_rule( + repository_ctx, + name = "mkl-include", + src_dir = _mkl_path(sycl_config) + "/include", + out_dir = "sycl/include", + )) + + sycl_libs = _find_libs(repository_ctx, sycl_config, bash_bin) + sycl_lib_srcs = [] + sycl_lib_outs = [] + for lib in sycl_libs.values(): + sycl_lib_srcs.append(lib.path) + sycl_lib_outs.append("sycl/lib/" + lib.file_name) + copy_rules.append(make_copy_files_rule( + repository_ctx, + name = "sycl-lib", + srcs = sycl_lib_srcs, + outs = sycl_lib_outs, + )) + + # Set up BUILD file for sycl/ + repository_ctx.template( + "sycl/build_defs.bzl", + tpl_paths["sycl:build_defs.bzl"], + { + "%{sycl_is_configured}": "True", + "%{sycl_build_is_configured}": "True", + }, + ) + + if sycl_config.sycl_basekit_version_number < "2024": + mkl_sycl_libs = '"{}"'.format( + "sycl/lib/" + sycl_libs["mkl_sycl"].file_name, + ) + else: + mkl_sycl_libs = '"{}",\n"{}",\n"{}",\n"{}",\n"{}",\n"{}",\n"{}",\n"{}"'.format( + "sycl/lib/" + sycl_libs["mkl_sycl_blas"].file_name, + "sycl/lib/" + sycl_libs["mkl_sycl_lapack"].file_name, + "sycl/lib/" + sycl_libs["mkl_sycl_sparse"].file_name, + "sycl/lib/" + sycl_libs["mkl_sycl_dft"].file_name, + "sycl/lib/" + sycl_libs["mkl_sycl_vm"].file_name, + "sycl/lib/" + sycl_libs["mkl_sycl_rng"].file_name, + "sycl/lib/" + sycl_libs["mkl_sycl_stats"].file_name, + "sycl/lib/" + sycl_libs["mkl_sycl_data_fitting"].file_name, + ) + core_sycl_libs = '"{}",\n"{}"'.format( + "sycl/lib/" + sycl_libs["sycl"].file_name, + "sycl/lib/" + sycl_libs["OpenCL"].file_name, + ) + repository_dict = { + "%{mkl_intel_ilp64_lib}": sycl_libs["mkl_intel_ilp64"].file_name, + "%{mkl_sequential_lib}": sycl_libs["mkl_sequential"].file_name, + "%{mkl_core_lib}": sycl_libs["mkl_core"].file_name, + "%{mkl_sycl_libs}": mkl_sycl_libs, + "%{core_sycl_libs}": core_sycl_libs, + "%{copy_rules}": "\n".join(copy_rules), + "%{sycl_headers}": ('":mkl-include",\n":sycl-include",\n'), + } + repository_ctx.template( + "sycl/BUILD", + tpl_paths["sycl:BUILD"], + repository_dict, + ) + + # Set up crosstool/ + + cc = find_cc(repository_ctx) + + host_compiler_includes = get_cxx_inc_directories(repository_ctx, cc) + + host_compiler_prefix = get_host_environ(repository_ctx, _GCC_HOST_COMPILER_PREFIX, "/usr/bin") + + sycl_defines = {} + + sycl_defines["%{host_compiler_prefix}"] = host_compiler_prefix + sycl_defines["%{host_compiler_path}"] = "clang/bin/crosstool_wrapper_driver_is_not_gcc" + + sycl_defines["%{cpu_compiler}"] = str(cc) + sycl_defines["%{linker_bin_path}"] = "/usr/bin" + + sycl_internal_inc_dirs = find_sycl_include_path(repository_ctx, sycl_config) + cxx_builtin_includes_list = sycl_internal_inc_dirs + _sycl_include_path(repository_ctx, sycl_config, bash_bin) + host_compiler_includes + + sycl_defines["%{cxx_builtin_include_directories}"] = to_list_of_strings(cxx_builtin_includes_list) + sycl_defines["%{extra_no_canonical_prefixes_flags}"] = "\"-fno-canonical-system-headers\"" + sycl_defines["%{unfiltered_compile_flags}"] = to_list_of_strings([ + "-DTENSORFLOW_USE_SYCL=1", + "-DMKL_ILP64", + "-fPIC", + ]) + sycl_defines["%{sycl_compiler_root}"] = str(sycl_config.sycl_toolkit_path) + sycl_defines["%{SYCL_ROOT_DIR}"] = str(sycl_config.sycl_toolkit_path) + sycl_defines["%{basekit_path}"] = str(sycl_config.sycl_basekit_path) + sycl_defines["%{basekit_version}"] = str(sycl_config.sycl_basekit_version_number) + sycl_defines["%{MKL_PATH}"] = _mkl_path(sycl_config) + + # Only expand template variables in the BUILD file + repository_ctx.template( + "crosstool/BUILD", + tpl_paths["crosstool:BUILD.sycl"], + sycl_defines, + ) + + # No templating of cc_toolchain_config - use attributes and templatize the + # BUILD file. + repository_ctx.template( + "crosstool/cc_toolchain_config.bzl", + tpl_paths["crosstool:sycl_cc_toolchain_config.bzl"], + sycl_defines, + ) + + repository_ctx.template( + "crosstool/clang/bin/crosstool_wrapper_driver_is_not_gcc", + tpl_paths["crosstool:clang/bin/crosstool_wrapper_driver_sycl"], + sycl_defines, + ) + +def _sycl_autoconf_imp(repository_ctx): + """Implementation of the sycl_autoconf rule.""" + if not enable_sycl(repository_ctx): + _create_dummy_repository(repository_ctx) + else: + _create_local_sycl_repository(repository_ctx) + +sycl_configure = repository_rule( + # Detects and configures the local SYCL toolchain. + # Add the following to your WORKSPACE FILE: + # ```python + # sycl_configure(name = "local_config_sycl") + # ``` + # Args: + # name: A unique name for this workspace rule. + implementation = _sycl_autoconf_imp, + local = True, + attrs = { + "_find_sycl_config": attr.label( + default = Label("//third_party/gpus:find_sycl_config.py"), + ), + }, +) diff --git a/third_party/tsl/workspace2.bzl b/third_party/tsl/workspace2.bzl index e23dcc3a4c7ad..d5004e732eece 100644 --- a/third_party/tsl/workspace2.bzl +++ b/third_party/tsl/workspace2.bzl @@ -20,6 +20,7 @@ load("//third_party/gemmlowp:workspace.bzl", gemmlowp = "repo") load("//third_party/git:git_configure.bzl", "git_configure") load("//third_party/gpus:cuda_configure.bzl", "cuda_configure") load("//third_party/gpus:rocm_configure.bzl", "rocm_configure") +load("//third_party/gpus:sycl_configure.bzl", "sycl_configure") load("//third_party/hwloc:workspace.bzl", hwloc = "repo") load("//third_party/implib_so:workspace.bzl", implib_so = "repo") load("//third_party/llvm:setup.bzl", "llvm_setup") @@ -76,6 +77,7 @@ def _tf_toolchains(): syslibs_configure(name = "local_config_syslibs") python_configure(name = "local_config_python") rocm_configure(name = "local_config_rocm") + sycl_configure(name = "local_config_sycl") remote_execution_configure(name = "local_config_remote_execution") # For windows bazel build diff --git a/xla/stream_executor/build_defs.bzl b/xla/stream_executor/build_defs.bzl index a937a57edc31d..88f8b4b02ef1e 100644 --- a/xla/stream_executor/build_defs.bzl +++ b/xla/stream_executor/build_defs.bzl @@ -1,7 +1,11 @@ """Configurations for StreamExecutor builds""" load("@local_config_cuda//cuda:build_defs.bzl", "if_cuda_is_configured") -load("@local_config_rocm//rocm:build_defs.bzl", _if_gpu_is_configured = "if_gpu_is_configured") +load( + "@local_config_rocm//rocm:build_defs.bzl", + _if_cuda_or_rocm = "if_cuda_or_rocm", + _if_gpu_is_configured = "if_gpu_is_configured", +) load( "@tsl//tsl/platform:rules_cc.bzl", "cc_library", @@ -23,8 +27,8 @@ def tf_additional_cudnn_plugin_copts(): def if_gpu_is_configured(if_true, if_false = []): return _if_gpu_is_configured(if_true, if_false) -def if_cuda_or_rocm(x): - return if_gpu_is_configured(x) +def if_cuda_or_rocm(if_true, if_false = []): + return _if_cuda_or_rocm(if_true, if_false) # nvlink is not available via the pip wheels, disable it since it will create # unnecessary dependency