| # Description: |
| # GPU-specific components in XLA service implementation. |
| |
| load( |
| "//tensorflow/core/platform:build_config.bzl", |
| "tf_proto_library_cc", |
| ) |
| load( |
| "//tensorflow/core/platform:build_config_root.bzl", |
| "tf_cuda_tests_tags", |
| ) |
| load( |
| "//tensorflow:tensorflow.bzl", |
| "tf_cc_test", |
| "tf_copts", |
| "tf_cuda_library", |
| ) |
| load("@local_config_cuda//cuda:build_defs.bzl", "if_cuda") |
| load( |
| "//tensorflow/core/platform/default:cuda_build_defs.bzl", |
| "if_cuda_is_configured", |
| ) |
| load( |
| "@local_config_rocm//rocm:build_defs.bzl", |
| "if_rocm_is_configured", |
| ) |
| load("//tensorflow:tensorflow.bzl", "if_nccl") |
| |
| package( |
| default_visibility = [":friends"], |
| licenses = ["notice"], # Apache 2.0 |
| ) |
| |
| package_group( |
| name = "friends", |
| includes = [ |
| "//tensorflow/compiler/xla:friends", |
| ], |
| ) |
| |
| # Filegroup used to collect source files for dependency checking. |
| filegroup( |
| name = "c_srcs", |
| data = glob([ |
| "**/*.cc", |
| "**/*.h", |
| ]), |
| ) |
| |
| tf_proto_library_cc( |
| name = "backend_configs", |
| srcs = ["backend_configs.proto"], |
| cc_api_version = 2, |
| protodeps = ["//tensorflow/compiler/xla:xla_data_proto"], |
| ) |
| |
| cc_library( |
| name = "gpu_constants", |
| srcs = ["gpu_constants.cc"], |
| hdrs = ["gpu_constants.h"], |
| deps = [ |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/core:framework", |
| ], |
| ) |
| |
| cc_library( |
| name = "gpu_types", |
| hdrs = ["gpu_types.h"], |
| deps = [ |
| "@com_google_absl//absl/types:variant", |
| ], |
| ) |
| |
| cc_library( |
| name = "partition_assignment", |
| srcs = [ |
| "partition_assignment.cc", |
| ], |
| hdrs = [ |
| "partition_assignment.h", |
| ], |
| deps = [ |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:stream_executor_no_cuda", |
| "@com_google_absl//absl/memory", |
| "@com_google_absl//absl/strings:str_format", |
| ], |
| ) |
| |
| # TODO(b/29140563) This target is flaky, disabled until flakiness is |
| # root-caused. Failed on 2016-06-08. |
| #tf_cc_test( |
| # name = "partition_assignment_test", |
| # srcs = [ |
| # "partition_assignment_test.cc", |
| # ], |
| # tags = tf_cuda_tests_tags(), |
| # deps = [ |
| # ":partition_assignment", |
| # "//tensorflow/core:stream_executor_no_cuda", |
| # "//tensorflow/compiler/xla:shape_util", |
| # "//tensorflow/compiler/xla:xla_data_proto", |
| # "//tensorflow/compiler/xla/service:gpu_plugin", |
| # "//tensorflow/compiler/xla/service:hlo", |
| # "//tensorflow/compiler/xla/tests:hlo_test_base", |
| # "//tensorflow/core:test_main", |
| # ], |
| #) |
| |
| tf_cc_test( |
| name = "custom_call_test", |
| srcs = if_cuda_is_configured(["custom_call_test.cc"]), |
| tags = tf_cuda_tests_tags(), |
| deps = [ |
| "//tensorflow/compiler/xla/tests:xla_internal_test_main", # fixdeps: keep |
| ] + if_cuda_is_configured([ |
| "@local_config_cuda//cuda:cuda_headers", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:test_helpers", |
| "//tensorflow/compiler/xla/client:xla_builder", |
| "//tensorflow/compiler/xla/client/lib:constants", |
| "//tensorflow/compiler/xla/service:custom_call_target_registry", |
| "//tensorflow/compiler/xla/service:gpu_plugin", |
| "//tensorflow/compiler/xla/tests:client_library_test_base", |
| "//tensorflow/core:test", |
| ]), |
| ) |
| |
| cc_library( |
| name = "stream_assignment", |
| srcs = ["stream_assignment.cc"], |
| hdrs = ["stream_assignment.h"], |
| deps = [ |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_reachability", |
| "@com_google_absl//absl/container:flat_hash_map", |
| "@com_google_absl//absl/container:flat_hash_set", |
| "@com_google_absl//absl/memory", |
| ], |
| ) |
| |
| tf_cc_test( |
| name = "stream_assignment_test", |
| srcs = [ |
| "stream_assignment_test.cc", |
| ], |
| tags = ["no_pip"], |
| deps = [ |
| ":stream_assignment", |
| "//tensorflow/compiler/xla:test_helpers", |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/tests:hlo_test_base", |
| "//tensorflow/compiler/xla/tests:test_utils", |
| "//tensorflow/compiler/xla/tests:xla_internal_test_main", |
| "//tensorflow/core:lib", |
| "@com_google_absl//absl/memory", |
| "@com_google_absl//absl/strings:str_format", |
| ], |
| ) |
| |
| cc_library( |
| name = "hlo_to_ir_bindings", |
| srcs = ["hlo_to_ir_bindings.cc"], |
| hdrs = ["hlo_to_ir_bindings.h"], |
| deps = [ |
| ":buffer_allocations", |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla/service:buffer_assignment", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service/llvm_ir:alias_analysis", |
| "//tensorflow/compiler/xla/service/llvm_ir:buffer_assignment_util", |
| "//tensorflow/compiler/xla/service/llvm_ir:ir_array", |
| "//tensorflow/compiler/xla/service/llvm_ir:llvm_util", |
| "//tensorflow/compiler/xla/service/llvm_ir:tuple_ops", |
| "//tensorflow/core:lib", |
| "@com_google_absl//absl/container:flat_hash_map", |
| "@com_google_absl//absl/container:flat_hash_set", |
| "@com_google_absl//absl/strings", |
| "@com_google_absl//absl/types:span", |
| "@llvm-project//llvm:core", |
| ], |
| ) |
| |
| cc_library( |
| name = "target_util", |
| srcs = ["target_util.cc"], |
| hdrs = ["target_util.h"], |
| deps = [ |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/service/llvm_ir:llvm_util", |
| "//tensorflow/core:lib", |
| "@com_google_absl//absl/strings", |
| "@com_google_absl//absl/types:span", |
| "@llvm-project//llvm:core", |
| "@llvm-project//llvm:support", |
| ], |
| ) |
| |
| cc_library( |
| name = "thunk_emitter", |
| srcs = ["thunk_emitter.cc"], |
| hdrs = ["thunk_emitter.h"], |
| copts = if_cuda_is_configured(["-DGOOGLE_CUDA=1"]), |
| deps = [ |
| ":backend_configs_cc", |
| ":buffer_allocations", |
| ":gpu_constants", |
| ":gpu_executable", |
| ":ir_emission_utils", |
| ":nccl_all_reduce_thunk", |
| ":thunk", |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/compiler/xla/service:buffer_assignment", |
| "//tensorflow/compiler/xla/service:custom_call_target_registry", |
| "//tensorflow/compiler/xla/service:hlo_casting_utils", |
| ], |
| ) |
| |
| cc_library( |
| name = "ir_emitter", |
| srcs = [ |
| "ir_emitter.cc", |
| "ir_emitter_nested.cc", |
| "ir_emitter_unnested.cc", |
| ], |
| hdrs = [ |
| "ir_emitter.h", |
| "ir_emitter_context.h", |
| "ir_emitter_nested.h", |
| "ir_emitter_unnested.h", |
| "kernel_mapping_scheme.h", |
| ], |
| deps = [ |
| ":backend_configs_cc", |
| ":buffer_allocations", |
| ":cudnn_batchnorm_runner", |
| ":elemental_ir_emitter", |
| ":gpu_constants", |
| ":gpu_conv_runner", |
| ":gpu_executable", |
| ":hlo_to_ir_bindings", |
| ":ir_emission_utils", |
| ":nccl_all_reduce_thunk", |
| ":parallel_loop_emitter", |
| ":partition_assignment", |
| ":target_util", |
| ":thunk", |
| ":thunk_emitter", |
| "//tensorflow/compiler/xla:literal", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla:window_util", |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/service:buffer_assignment", |
| "//tensorflow/compiler/xla/service:elemental_ir_emitter", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_casting_utils", |
| "//tensorflow/compiler/xla/service:name_uniquer", |
| "//tensorflow/compiler/xla/service:pattern_matcher", |
| "//tensorflow/compiler/xla/service:while_loop_analysis", |
| "//tensorflow/compiler/xla/service/llvm_ir:buffer_assignment_util", |
| "//tensorflow/compiler/xla/service/llvm_ir:dynamic_update_slice_util", |
| "//tensorflow/compiler/xla/service/llvm_ir:fused_ir_emitter", |
| "//tensorflow/compiler/xla/service/llvm_ir:ir_array", |
| "//tensorflow/compiler/xla/service/llvm_ir:ir_builder_mixin", |
| "//tensorflow/compiler/xla/service/llvm_ir:kernel_support_library", |
| "//tensorflow/compiler/xla/service/llvm_ir:llvm_loop", |
| "//tensorflow/compiler/xla/service/llvm_ir:llvm_util", |
| "//tensorflow/compiler/xla/service/llvm_ir:loop_emitter", |
| "//tensorflow/compiler/xla/service/llvm_ir:sort_util", |
| "//tensorflow/compiler/xla/service/llvm_ir:tuple_ops", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:stream_executor_no_cuda", |
| "@com_google_absl//absl/algorithm:container", |
| "@com_google_absl//absl/container:inlined_vector", |
| "@com_google_absl//absl/memory", |
| "@com_google_absl//absl/strings", |
| "@com_google_absl//absl/types:optional", |
| "@com_google_absl//absl/types:span", |
| "@llvm-project//llvm:core", |
| "@llvm-project//llvm:support", |
| ], |
| ) |
| |
| cc_library( |
| name = "parallel_loop_emitter", |
| srcs = ["parallel_loop_emitter.cc"], |
| hdrs = ["parallel_loop_emitter.h"], |
| deps = [ |
| ":partition_assignment", |
| ":target_util", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/service/llvm_ir:ir_array", |
| "//tensorflow/compiler/xla/service/llvm_ir:llvm_loop", |
| "//tensorflow/compiler/xla/service/llvm_ir:llvm_util", |
| "//tensorflow/compiler/xla/service/llvm_ir:loop_emitter", |
| "//tensorflow/core:lib", |
| "@llvm-project//llvm:core", |
| ], |
| ) |
| |
| cc_library( |
| name = "elemental_ir_emitter", |
| srcs = ["elemental_ir_emitter.cc"], |
| hdrs = ["elemental_ir_emitter.h"], |
| deps = [ |
| ":target_util", |
| "//tensorflow/compiler/xla:literal", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla:window_util", |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/service:elemental_ir_emitter", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_module_config", |
| "//tensorflow/compiler/xla/service/llvm_ir:ir_array", |
| "//tensorflow/compiler/xla/service/llvm_ir:llvm_loop", |
| "//tensorflow/compiler/xla/service/llvm_ir:llvm_util", |
| "//tensorflow/compiler/xla/service/llvm_ir:loop_emitter", |
| "//tensorflow/compiler/xla/service/llvm_ir:math_ops", |
| "//tensorflow/core:lib", |
| "@com_google_absl//absl/strings", |
| "@com_google_absl//absl/types:span", |
| "@llvm-project//llvm:core", |
| "@llvm-project//llvm:support", |
| ], |
| ) |
| |
| cc_library( |
| name = "buffer_allocations", |
| srcs = ["buffer_allocations.cc"], |
| hdrs = ["buffer_allocations.h"], |
| deps = [ |
| ":gpu_constants", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla/service:buffer_assignment", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:lib_internal", |
| "//tensorflow/core:stream_executor_no_cuda", |
| "//tensorflow/stream_executor:device_memory_allocator", |
| "@com_google_absl//absl/container:flat_hash_map", |
| "@com_google_absl//absl/memory", |
| "@com_google_absl//absl/types:span", |
| ], |
| ) |
| |
| cc_library( |
| name = "hlo_execution_profiler", |
| srcs = ["hlo_execution_profiler.cc"], |
| hdrs = ["hlo_execution_profiler.h"], |
| deps = [ |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_execution_profile", |
| "//tensorflow/compiler/xla/service:stream_pool", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:ptr_util", |
| "//tensorflow/core:stream_executor_no_cuda", |
| "@com_google_absl//absl/memory", |
| ], |
| ) |
| |
| cc_library( |
| name = "thunk", |
| srcs = ["thunk.cc"], |
| hdrs = ["thunk.h"], |
| deps = [ |
| ":buffer_allocations", |
| ":hlo_execution_profiler", |
| "//tensorflow/compiler/xla:executable_run_options", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:stream_executor_no_cuda", |
| ], |
| ) |
| |
| # First level of nested select. NCCL requires both if_cuda and if_nccl. |
| filegroup( |
| name = "nccl_all_reduce_thunk_src", |
| srcs = if_nccl( |
| ["nccl_all_reduce_thunk.cc"], |
| ["dummy_all_reduce_thunk.cc"], |
| ), |
| ) |
| |
| tf_cuda_library( |
| name = "nccl_all_reduce_thunk", |
| srcs = if_cuda( |
| [":nccl_all_reduce_thunk_src"], |
| ["dummy_all_reduce_thunk.cc"], |
| ), |
| hdrs = ["nccl_all_reduce_thunk.h"], |
| deps = [ |
| ":buffer_allocations", |
| ":hlo_execution_profiler", |
| ":thunk", |
| "@com_google_absl//absl/base:core_headers", |
| "//tensorflow/compiler/xla/service:pattern_matcher", |
| "//tensorflow/compiler/xla:refcounting_hash_map", |
| "//tensorflow/compiler/xla/service:collective_ops_utils", |
| "@com_google_absl//absl/container:flat_hash_set", |
| "@com_google_absl//absl/synchronization", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla/service:buffer_assignment", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:stream_executor_no_cuda", |
| "//tensorflow/stream_executor/cuda:cuda_activation", |
| "//tensorflow/stream_executor/cuda:cuda_gpu_executor", |
| ] + if_nccl([ |
| "@local_config_nccl//:nccl", |
| ]), |
| ) |
| |
| cc_library( |
| name = "gpu_debug_info_manager", |
| srcs = [ |
| "gpu_debug_info_manager.cc", |
| ], |
| hdrs = [ |
| "gpu_debug_info_manager.h", |
| ], |
| deps = [ |
| "//tensorflow/compiler/xla/service:buffer_assignment", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_proto_util", |
| "//tensorflow/core:lib", |
| "@com_google_absl//absl/container:flat_hash_map", |
| ], |
| ) |
| |
| tf_cc_test( |
| name = "gpu_debug_info_manager_test", |
| srcs = ["gpu_debug_info_manager_test.cc"], |
| tags = tf_cuda_tests_tags(), |
| deps = [ |
| ":gpu_constants", |
| ":gpu_debug_info_manager", |
| ":gpu_hlo_schedule", |
| ":stream_assignment", |
| "//tensorflow/compiler/xla/service:buffer_assignment", |
| "//tensorflow/compiler/xla/tests:hlo_test_base", |
| "//tensorflow/compiler/xla/tests:test_utils", |
| "//tensorflow/compiler/xla/tests:xla_internal_test_main", |
| "//tensorflow/core:test", |
| ], |
| ) |
| |
| cc_library( |
| name = "gpu_executable", |
| srcs = [ |
| "collective_permute_thunk.cc", |
| "conditional_thunk.cc", |
| "convolution_thunk.cc", |
| "copy_thunk.cc", |
| "cudnn_batchnorm_thunk.cc", |
| "fft_thunk.cc", |
| "for_thunk.cc", |
| "gemm_thunk.cc", |
| "gpu_executable.cc", |
| "infeed_thunk.cc", |
| "kernel_thunk.cc", |
| "memset_thunk.cc", |
| "outfeed_thunk.cc", |
| "replica_id_thunk.cc", |
| "sequential_thunk.cc", |
| "thunk_schedule.cc", |
| "triangular_solve_thunk.cc", |
| "tuple_thunk.cc", |
| "while_thunk.cc", |
| ] + if_cuda_is_configured([ |
| "cholesky_thunk.cc", |
| "custom_call_thunk.cc", |
| ]), |
| hdrs = [ |
| "collective_permute_thunk.h", |
| "conditional_thunk.h", |
| "convolution_thunk.h", |
| "copy_thunk.h", |
| "cudnn_batchnorm_thunk.h", |
| "fft_thunk.h", |
| "for_thunk.h", |
| "gemm_thunk.h", |
| "gpu_executable.h", |
| "infeed_thunk.h", |
| "kernel_thunk.h", |
| "memset_thunk.h", |
| "outfeed_thunk.h", |
| "replica_id_thunk.h", |
| "sequential_thunk.h", |
| "thunk_schedule.h", |
| "triangular_solve_thunk.h", |
| "tuple_thunk.h", |
| "while_thunk.h", |
| ] + if_cuda_is_configured([ |
| "cholesky_thunk.h", |
| "custom_call_thunk.h", |
| ]), |
| deps = [ |
| ":backend_configs_cc", |
| ":buffer_allocations", |
| ":cusolver_context", |
| ":cudnn_batchnorm_runner", |
| ":gpu_conv_runner", |
| ":gpu_debug_info_manager", |
| ":gpu_types", |
| ":hlo_execution_profiler", |
| ":infeed_manager", |
| ":ir_emission_utils", |
| ":nccl_all_reduce_thunk", # fixdeps: keep |
| ":outfeed_manager", |
| ":partition_assignment", |
| ":stream_assignment", |
| ":stream_executor_util", |
| ":thunk", |
| "//tensorflow/compiler/xla:array2d", |
| "//tensorflow/compiler/xla:literal", |
| "//tensorflow/compiler/xla:refcounting_hash_map", |
| "//tensorflow/compiler/xla:shape_tree", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:status", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/service:buffer_assignment", |
| "//tensorflow/compiler/xla/service:executable", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_casting_utils", |
| "//tensorflow/compiler/xla/service:hlo_dataflow_analysis", |
| "//tensorflow/compiler/xla/service:hlo_execution_profile", |
| "//tensorflow/compiler/xla/service:logical_buffer", |
| "//tensorflow/compiler/xla/service:shaped_buffer", |
| "//tensorflow/compiler/xla/service:transfer_manager", |
| "//tensorflow/compiler/xla/service/llvm_ir:buffer_assignment_util", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:lib_internal", |
| "//tensorflow/core:stream_executor_no_cuda", |
| "//tensorflow/core/profiler/lib:traceme", |
| "//tensorflow/core/profiler/lib:scoped_annotation", |
| "//tensorflow/stream_executor", |
| "//tensorflow/stream_executor:blas", |
| "//tensorflow/stream_executor:device_memory", |
| "//tensorflow/stream_executor:device_memory_allocator", |
| "//tensorflow/stream_executor:kernel", |
| "//tensorflow/stream_executor/gpu:gpu_stream", |
| "@com_google_absl//absl/algorithm:container", |
| "@com_google_absl//absl/base:core_headers", |
| "@com_google_absl//absl/container:flat_hash_map", |
| "@com_google_absl//absl/container:flat_hash_set", |
| "@com_google_absl//absl/memory", |
| "@com_google_absl//absl/strings", |
| "@com_google_absl//absl/strings:str_format", |
| "@com_google_absl//absl/types:optional", |
| "@com_google_absl//absl/types:span", |
| ] + if_cuda_is_configured([ |
| "//tensorflow/stream_executor/cuda:cuda_stream", |
| "//tensorflow/core/platform/default/build_config:cublas_plugin", |
| "//tensorflow/core/platform/default/build_config:cudnn_plugin", |
| "//tensorflow/core/platform/default/build_config:cufft_plugin", |
| "//tensorflow/core/platform/default/build_config:stream_executor_cuda", # build_cleaner: keep |
| "@local_config_cuda//cuda:cuda_headers", |
| ]) + if_rocm_is_configured([ |
| "//tensorflow/core/platform/default/build_config:stream_executor_rocm", |
| "@local_config_rocm//rocm:rocm_headers", |
| ]), |
| ) |
| |
| cc_library( |
| name = "ir_emission_utils", |
| srcs = ["ir_emission_utils.cc"], |
| hdrs = ["ir_emission_utils.h"], |
| deps = [ |
| ":backend_configs_cc", |
| ":target_util", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla:window_util", |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service/llvm_ir:llvm_util", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:stream_executor_no_cuda", |
| "@com_google_absl//absl/algorithm:container", |
| "@llvm-project//llvm:core", |
| ], |
| ) |
| |
| cc_library( |
| name = "gemm_rewriter", |
| srcs = ["gemm_rewriter.cc"], |
| hdrs = ["gemm_rewriter.h"], |
| deps = [ |
| ":backend_configs_cc", |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_casting_utils", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/compiler/xla/service:pattern_matcher", |
| "//tensorflow/core:lib", |
| "//tensorflow/stream_executor/lib", |
| "@com_google_absl//absl/types:optional", |
| ], |
| ) |
| |
| cc_library( |
| name = "gemm_algorithm_picker", |
| srcs = if_cuda_is_configured(["gemm_algorithm_picker.cc"]), |
| hdrs = if_cuda_is_configured(["gemm_algorithm_picker.h"]), |
| deps = if_cuda_is_configured([ |
| ":backend_configs_cc", |
| ":buffer_comparator", |
| ":gpu_conv_runner", |
| ":gpu_executable", |
| ":ir_emission_utils", |
| ":stream_executor_util", |
| "@com_google_absl//absl/types:optional", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/core:autotuning_proto_cc", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:stream_executor_no_cuda", |
| "//tensorflow/core/util/proto:proto_utils", |
| "//tensorflow/stream_executor:blas", |
| "//tensorflow/stream_executor:device_memory", |
| "//tensorflow/stream_executor:device_memory_allocator", |
| "//tensorflow/stream_executor/gpu:redzone_allocator", |
| ]), |
| ) |
| |
| cc_library( |
| name = "gpu_conv_algorithm_picker", |
| srcs = ["gpu_conv_algorithm_picker.cc"], |
| hdrs = ["gpu_conv_algorithm_picker.h"], |
| copts = if_cuda_is_configured(["-DGOOGLE_CUDA=1"]), |
| deps = [ |
| ":backend_configs_cc", |
| ":gpu_autotuning_proto_cc", |
| ":gpu_conv_runner", |
| ":gpu_executable", |
| ":hlo_algorithm_blacklist", |
| ":ir_emission_utils", |
| ":stream_executor_util", |
| "@com_google_absl//absl/algorithm:container", |
| "@com_google_absl//absl/strings", |
| "@com_google_absl//absl/strings:str_format", |
| "@com_google_absl//absl/time", |
| "@com_google_absl//absl/types:optional", |
| "//tensorflow/compiler/xla:literal_util", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla/service:compiler", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_casting_utils", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/core:autotuning_proto_cc", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:lib_internal", |
| "//tensorflow/core:stream_executor_no_cuda", |
| "//tensorflow/core/util/proto:proto_utils", |
| "//tensorflow/stream_executor:device_memory_allocator", |
| ] + if_cuda_is_configured([ |
| ":buffer_comparator", |
| "//tensorflow/stream_executor/gpu:redzone_allocator", |
| ]), |
| ) |
| |
| cc_library( |
| name = "gpu_conv_runner", |
| srcs = ["gpu_conv_runner.cc"], |
| hdrs = ["gpu_conv_runner.h"], |
| deps = [ |
| ":backend_configs_cc", |
| ":ir_emission_utils", |
| ":stream_executor_util", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:status", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/core:stream_executor_no_cuda", |
| "@com_google_absl//absl/strings", |
| "@com_google_absl//absl/types:optional", |
| ], |
| ) |
| |
| cc_library( |
| name = "cudnn_batchnorm_runner", |
| srcs = ["cudnn_batchnorm_runner.cc"], |
| hdrs = ["cudnn_batchnorm_runner.h"], |
| deps = [ |
| ":backend_configs_cc", |
| ":ir_emission_utils", |
| ":stream_executor_util", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:status", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/core:stream_executor_no_cuda", |
| "@com_google_absl//absl/strings", |
| "@com_google_absl//absl/types:optional", |
| ], |
| ) |
| |
| cc_library( |
| name = "gpu_conv_rewriter", |
| srcs = ["gpu_conv_rewriter.cc"], |
| hdrs = ["gpu_conv_rewriter.h"], |
| deps = [ |
| ":backend_configs_cc", |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:literal", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla:window_util", |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/core:lib", |
| ], |
| ) |
| |
| tf_cc_test( |
| name = "gpu_conv_rewriter_test", |
| srcs = ["gpu_conv_rewriter_test.cc"], |
| tags = tf_cuda_tests_tags(), |
| deps = [ |
| ":gpu_conv_rewriter", |
| ":ir_emission_utils", |
| "//tensorflow/compiler/jit:xla_gpu_jit", |
| "//tensorflow/compiler/xla:test", |
| "//tensorflow/compiler/xla:test_helpers", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_matchers", |
| "//tensorflow/compiler/xla/service:shape_inference", |
| "//tensorflow/compiler/xla/tests:hlo_test_base", |
| "//tensorflow/compiler/xla/tests:xla_internal_test_main", # fixdeps: keep |
| "//tensorflow/core:test", |
| ], |
| ) |
| |
| cc_library( |
| name = "cusolver_context", |
| srcs = if_cuda_is_configured(["cusolver_context.cc"]), |
| hdrs = ["cusolver_context.h"], |
| deps = [ |
| # LINT.IfChange |
| "@local_config_cuda//cuda:cublas_headers", |
| # LINT.ThenChange(//tensorflow/copy.bara.sky:cublas_headers) |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:stream_executor_no_cuda", |
| "//tensorflow/stream_executor:blas", |
| ] + if_cuda_is_configured([ |
| "@local_config_cuda//cuda:cuda_headers", |
| "//tensorflow/stream_executor/cuda:cusolver_lib", |
| ]), |
| ) |
| |
| cc_library( |
| name = "cusolver_rewriter", |
| srcs = if_cuda_is_configured(["cusolver_rewriter.cc"]), |
| hdrs = if_cuda_is_configured(["cusolver_rewriter.h"]), |
| deps = if_cuda_is_configured([ |
| ":cusolver_context", |
| ":ir_emission_utils", |
| "@com_google_absl//absl/types:optional", |
| "//tensorflow/compiler/xla:literal", |
| "//tensorflow/compiler/xla:literal_util", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:stream_executor_no_cuda", |
| "//tensorflow/stream_executor:blas", |
| "//tensorflow/stream_executor:device_memory_allocator", |
| ]), |
| ) |
| |
| cc_library( |
| name = "instruction_fusion", |
| srcs = ["instruction_fusion.cc"], |
| hdrs = ["instruction_fusion.h"], |
| deps = [ |
| ":gpu_fusible", |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/service:fusion_node_indexing_evaluation", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_query", |
| "//tensorflow/compiler/xla/service:instruction_fusion", |
| "//tensorflow/compiler/xla/service:pattern_matcher", |
| "//tensorflow/compiler/xla/service/llvm_ir:fused_ir_emitter", |
| "@com_google_absl//absl/container:flat_hash_map", |
| "@com_google_absl//absl/container:flat_hash_set", |
| ], |
| ) |
| |
| tf_cc_test( |
| name = "instruction_fusion_test", |
| srcs = ["instruction_fusion_test.cc"], |
| tags = ["no_pip"], |
| deps = [ |
| ":gpu_fusible", |
| ":instruction_fusion", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_matchers", |
| "//tensorflow/compiler/xla/service:hlo_parser", |
| "//tensorflow/compiler/xla/tests:hlo_test_base", |
| "//tensorflow/compiler/xla/tests:test_utils", |
| "//tensorflow/compiler/xla/tests:xla_internal_test_main", |
| ], |
| ) |
| |
| cc_library( |
| name = "multi_output_fusion", |
| srcs = ["multi_output_fusion.cc"], |
| hdrs = ["multi_output_fusion.h"], |
| deps = [ |
| ":gpu_fusible", |
| ":instruction_fusion", |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:debug_options_flags", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_reachability", |
| "//tensorflow/compiler/xla/service:multi_output_fusion", |
| "//tensorflow/core:lib", |
| "@com_google_absl//absl/algorithm:container", |
| "@com_google_absl//absl/container:flat_hash_set", |
| ], |
| ) |
| |
| tf_cc_test( |
| name = "multi_output_fusion_test", |
| srcs = ["multi_output_fusion_test.cc"], |
| tags = ["no_pip"], |
| deps = [ |
| ":gpu_fusible", |
| ":instruction_fusion", |
| ":multi_output_fusion", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_matchers", |
| "//tensorflow/compiler/xla/service:hlo_parser", |
| "//tensorflow/compiler/xla/tests:hlo_test_base", |
| "//tensorflow/compiler/xla/tests:xla_internal_test_main", |
| "//tensorflow/core:lib", |
| "@com_google_absl//absl/strings", |
| ], |
| ) |
| |
| cc_library( |
| name = "gpu_copy_insertion", |
| srcs = ["gpu_copy_insertion.cc"], |
| hdrs = ["gpu_copy_insertion.h"], |
| deps = [ |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla/service:call_graph", |
| "//tensorflow/compiler/xla/service:copy_insertion", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_dataflow_analysis", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/core:lib", |
| "@com_google_absl//absl/container:flat_hash_set", |
| ], |
| ) |
| |
| cc_library( |
| name = "gpu_sanitize_constant_names", |
| srcs = ["gpu_sanitize_constant_names.cc"], |
| hdrs = ["gpu_sanitize_constant_names.h"], |
| deps = [ |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/compiler/xla/service/llvm_ir:buffer_assignment_util", |
| "//tensorflow/core:lib", |
| ], |
| ) |
| |
| tf_cc_test( |
| name = "gpu_sanitize_constant_names_test", |
| srcs = ["gpu_sanitize_constant_names_test.cc"], |
| tags = tf_cuda_tests_tags(), |
| deps = [ |
| ":gpu_sanitize_constant_names", |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:shape_layout", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:test_helpers", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/service:computation_layout", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_matchers", |
| "//tensorflow/compiler/xla/service:hlo_module_config", |
| "//tensorflow/compiler/xla/service:hlo_parser", |
| "//tensorflow/compiler/xla/tests:hlo_test_base", |
| "//tensorflow/compiler/xla/tests:test_utils", |
| "//tensorflow/compiler/xla/tests:xla_internal_test_main", |
| "//tensorflow/core:test", |
| "@com_google_absl//absl/strings", |
| ], |
| ) |
| |
| cc_library( |
| name = "fusion_merger", |
| srcs = ["fusion_merger.cc"], |
| hdrs = ["fusion_merger.h"], |
| deps = [ |
| ":gpu_fusible", |
| ":instruction_fusion", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_cost_analysis", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/compiler/xla/service/llvm_ir:fused_ir_emitter", |
| "//tensorflow/core:lib", |
| "@com_google_absl//absl/algorithm:container", |
| "@com_google_absl//absl/strings", |
| ], |
| ) |
| |
| tf_cc_test( |
| name = "fusion_merger_test", |
| srcs = ["fusion_merger_test.cc"], |
| tags = ["no_pip"], |
| deps = [ |
| ":fusion_merger", |
| ":gpu_fusible", |
| ":instruction_fusion", |
| "//tensorflow/compiler/xla:test_helpers", |
| "//tensorflow/compiler/xla/service:hlo_matchers", |
| "//tensorflow/compiler/xla/service:hlo_parser", |
| "//tensorflow/compiler/xla/tests:hlo_test_base", |
| "//tensorflow/compiler/xla/tests:xla_internal_test_main", |
| "@com_google_absl//absl/types:span", |
| ], |
| ) |
| |
| cc_library( |
| name = "gpu_conv_padding_legalization", |
| srcs = ["gpu_conv_padding_legalization.cc"], |
| hdrs = ["gpu_conv_padding_legalization.h"], |
| deps = [ |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:literal", |
| "//tensorflow/compiler/xla:literal_util", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla:window_util", |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/service:hlo_casting_utils", |
| "//tensorflow/compiler/xla/service:hlo_creation_utils", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/compiler/xla/service:shape_inference", |
| "@com_google_absl//absl/memory", |
| ], |
| ) |
| |
| cc_library( |
| name = "cudnn_pad_for_convolutions", |
| srcs = ["cudnn_pad_for_convolutions.cc"], |
| hdrs = ["cudnn_pad_for_convolutions.h"], |
| deps = [ |
| ":ir_emission_utils", |
| ":stream_executor_util", |
| "//tensorflow/compiler/xla:literal_util", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla:window_util", |
| "//tensorflow/compiler/xla/service:hlo_casting_utils", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| ], |
| ) |
| |
| tf_cc_test( |
| name = "cudnn_pad_for_convolutions_test", |
| srcs = ["cudnn_pad_for_convolutions_test.cc"], |
| tags = tf_cuda_tests_tags(), |
| deps = [ |
| ":cudnn_pad_for_convolutions", |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla/service:hlo_matchers", |
| "//tensorflow/compiler/xla/service:hlo_parser", |
| "//tensorflow/compiler/xla/tests:hlo_test_base", |
| "//tensorflow/compiler/xla/tests:xla_internal_test_main", # build_cleaner: keep |
| ], |
| ) |
| |
| cc_library( |
| name = "cublas_gemm_pad_for_tensor_cores", |
| srcs = ["cublas_gemm_pad_for_tensor_cores.cc"], |
| hdrs = ["cublas_gemm_pad_for_tensor_cores.h"], |
| deps = [ |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:literal_util", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla:window_util", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_casting_utils", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/core/platform:types", |
| ], |
| ) |
| |
| tf_cc_test( |
| name = "cublas_gemm_pad_for_tensor_cores_test", |
| srcs = ["cublas_gemm_pad_for_tensor_cores_test.cc"], |
| tags = ["no_pip"], |
| deps = [ |
| ":cublas_gemm_pad_for_tensor_cores", |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla/service:hlo_matchers", |
| "//tensorflow/compiler/xla/service:hlo_parser", |
| "//tensorflow/compiler/xla/tests:hlo_test_base", |
| "//tensorflow/compiler/xla/tests:test_utils", |
| "//tensorflow/compiler/xla/tests:xla_internal_test_main", # build_cleaner: keep |
| ], |
| ) |
| |
| cc_library( |
| name = "target_constants", |
| hdrs = ["target_constants.h"], |
| ) |
| |
| cc_library( |
| name = "gpu_transfer_manager", |
| srcs = ["gpu_transfer_manager.cc"], |
| hdrs = ["gpu_transfer_manager.h"], |
| deps = [ |
| ":infeed_manager", |
| ":outfeed_manager", |
| ":target_constants", |
| "//tensorflow/compiler/xla:literal", |
| "//tensorflow/compiler/xla:literal_util", |
| "//tensorflow/compiler/xla:shape_tree", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/service:generic_transfer_manager", |
| "//tensorflow/compiler/xla/service:transfer_manager", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:stream_executor_no_cuda", |
| "@com_google_absl//absl/memory", |
| "@llvm-project//llvm:core", |
| ], |
| alwayslink = True, # Contains per-platform transfer manager registration |
| ) |
| |
| cc_library( |
| name = "gpu_compiler", |
| srcs = [ |
| "gpu_compiler.cc", |
| ], |
| hdrs = [ |
| "gpu_compiler.h", |
| ], |
| deps = [ |
| ":alias_passthrough_params", |
| ":cudnn_batchnorm_rewriter", |
| ":cudnn_pad_for_convolutions", |
| ":fusion_merger", |
| ":gpu_constants", |
| ":gpu_conv_algorithm_picker", |
| ":gpu_conv_padding_legalization", |
| ":gpu_conv_rewriter", |
| ":gpu_copy_insertion", |
| ":gpu_executable", |
| ":gpu_hlo_schedule", |
| ":gpu_layout_assignment", |
| ":gpu_sanitize_constant_names", |
| ":gpu_scatter_expander", |
| ":horizontal_fusion", |
| ":instruction_fusion", |
| ":ir_emission_utils", |
| ":ir_emitter", |
| ":multi_output_fusion", |
| ":partition_assignment", |
| ":stream_assignment", |
| ":stream_executor_util", |
| ":target_constants", |
| ":variadic_op_splitter", |
| "//tensorflow/compiler/xla:protobuf_util", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla/service:algebraic_simplifier", |
| "//tensorflow/compiler/xla/service:batchnorm_expander", |
| "//tensorflow/compiler/xla/service:buffer_assignment", |
| "//tensorflow/compiler/xla/service:call_inliner", |
| "//tensorflow/compiler/xla/service:conditional_simplifier", |
| "//tensorflow/compiler/xla/service:convolution_4d_expander", |
| "//tensorflow/compiler/xla/service:convolution_group_converter", |
| "//tensorflow/compiler/xla/service:depthwise_convolution_converter", |
| "//tensorflow/compiler/xla/service:dot_decomposer", |
| "//tensorflow/compiler/xla/service:dump", |
| "//tensorflow/compiler/xla/service:dynamic_index_splitter", |
| "//tensorflow/compiler/xla/service:executable", |
| "//tensorflow/compiler/xla/service:flatten_call_graph", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_constant_folding", |
| "//tensorflow/compiler/xla/service:hlo_cse", |
| "//tensorflow/compiler/xla/service:hlo_dataflow_analysis", |
| "//tensorflow/compiler/xla/service:hlo_dce", |
| "//tensorflow/compiler/xla/service:hlo_element_type_converter", |
| "//tensorflow/compiler/xla/service:hlo_get_dimension_size_rewriter", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/compiler/xla/service:hlo_pass_pipeline", |
| "//tensorflow/compiler/xla/service:hlo_proto_cc", |
| "//tensorflow/compiler/xla/service:hlo_proto_util", |
| "//tensorflow/compiler/xla/service:hlo_subcomputation_unification", |
| "//tensorflow/compiler/xla/service:hlo_verifier", |
| "//tensorflow/compiler/xla/service:llvm_compiler", |
| "//tensorflow/compiler/xla/service:reshape_mover", |
| "//tensorflow/compiler/xla/service:rng_bit_generator_expander", |
| "//tensorflow/compiler/xla/service:rng_expander", |
| "//tensorflow/compiler/xla/service:slice_sinker", |
| "//tensorflow/compiler/xla/service:slow_operation_alarm", |
| "//tensorflow/compiler/xla/service:sort_simplifier", |
| "//tensorflow/compiler/xla/service:stable_sort_expander", |
| "//tensorflow/compiler/xla/service:transpose_folding", |
| "//tensorflow/compiler/xla/service:tuple_simplifier", |
| "//tensorflow/compiler/xla/service:while_loop_constant_sinking", |
| "//tensorflow/compiler/xla/service:while_loop_simplifier", |
| "//tensorflow/compiler/xla/service:while_loop_trip_count_annotator", |
| "//tensorflow/compiler/xla/service:zero_sized_hlo_elimination", |
| "//tensorflow/compiler/xla/service/gpu/llvm_gpu_backend", |
| "//tensorflow/compiler/xla/service/llvm_ir:llvm_util", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:lib_internal", |
| "//tensorflow/core:regexp_internal", |
| "//tensorflow/core:stream_executor_no_cuda", |
| "//tensorflow/core/profiler/lib:traceme", |
| "//tensorflow/stream_executor:stream_executor_headers", |
| "@com_google_absl//absl/container:node_hash_map", |
| "@com_google_absl//absl/memory", |
| "@com_google_absl//absl/strings", |
| "@com_google_absl//absl/types:optional", |
| "@com_google_absl//absl/types:span", |
| "@llvm-project//llvm:core", |
| ], |
| ) |
| |
| cc_library( |
| name = "nvptx_compiler", |
| srcs = if_cuda_is_configured([ |
| "nvptx_compiler_registration.cc", |
| ]), |
| deps = if_cuda_is_configured([ |
| ":nvptx_compiler_impl", |
| ]), |
| alwayslink = True, # Contains compiler registration |
| ) |
| |
| cc_library( |
| name = "nvptx_compiler_impl", |
| srcs = if_cuda_is_configured([ |
| "nvptx_compiler.cc", |
| ]), |
| hdrs = if_cuda_is_configured([ |
| "nvptx_compiler.h", |
| ]), |
| deps = if_cuda_is_configured([ |
| ":cublas_gemm_pad_for_tensor_cores", |
| ":cudnn_fused_conv_rewriter", |
| ":cudnn_pad_for_convolutions", |
| ":cusolver_rewriter", |
| ":gemm_algorithm_picker", |
| ":gemm_rewriter", |
| ":gpu_compiler", |
| ":gpu_conv_algorithm_picker", |
| ":gpu_conv_padding_legalization", |
| ":gpu_conv_rewriter", |
| ":gpu_layout_assignment", |
| ":ir_emission_utils", |
| ":reduction_degenerate_dim_remover", |
| ":reduction_dimension_grouper", |
| ":reduction_layout_normalizer", |
| ":stream_executor_util", |
| ":target_constants", |
| ":tree_reduction_rewriter", |
| "@com_google_absl//absl/base", |
| "@com_google_absl//absl/container:node_hash_map", |
| "@com_google_absl//absl/types:optional", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla/service:algebraic_simplifier", |
| "//tensorflow/compiler/xla/service:dump", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_constant_folding", |
| "//tensorflow/compiler/xla/service:hlo_cse", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/compiler/xla/service:hlo_pass_pipeline", |
| "//tensorflow/compiler/xla/service:hlo_proto_cc", |
| "//tensorflow/compiler/xla/service:hlo_verifier", |
| "//tensorflow/compiler/xla/service:llvm_compiler", |
| "//tensorflow/compiler/xla/service:tuple_simplifier", |
| "//tensorflow/compiler/xla/service/gpu/llvm_gpu_backend", |
| "//tensorflow/compiler/xla/service/llvm_ir:llvm_util", |
| "//tensorflow/core:cuda_libdevice_path", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:lib_internal", |
| "//tensorflow/core/profiler/lib:traceme", |
| "//tensorflow/stream_executor:stream_executor_headers", |
| "//tensorflow/stream_executor/cuda:cuda_diagnostics", |
| "//tensorflow/stream_executor/gpu:asm_compiler", |
| ]), |
| ) |
| |
| cc_library( |
| name = "amdgpu_compiler", |
| srcs = [ |
| "amdgpu_compiler_registration.cc", |
| ], |
| deps = [ |
| ":amdgpu_compiler_impl", |
| ], |
| alwayslink = True, # Contains compiler registration |
| ) |
| |
| cc_library( |
| name = "amdgpu_compiler_impl", |
| srcs = [ |
| "amdgpu_compiler.cc", |
| ], |
| hdrs = [ |
| "amdgpu_compiler.h", |
| ], |
| deps = [ |
| ":gemm_rewriter", |
| ":gpu_compiler", |
| ":gpu_conv_algorithm_picker", |
| ":gpu_conv_padding_legalization", |
| ":gpu_conv_rewriter", |
| ":gpu_layout_assignment", |
| ":reduction_degenerate_dim_remover", |
| ":reduction_dimension_grouper", |
| ":reduction_layout_normalizer", |
| ":target_constants", |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla/service:algebraic_simplifier", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_constant_folding", |
| "//tensorflow/compiler/xla/service:hlo_cse", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/compiler/xla/service:hlo_pass_pipeline", |
| "//tensorflow/compiler/xla/service:hlo_verifier", |
| "//tensorflow/compiler/xla/service:tuple_simplifier", |
| "//tensorflow/compiler/xla/service/gpu/llvm_gpu_backend", |
| "//tensorflow/compiler/xla/service/llvm_ir:llvm_util", |
| "//tensorflow/core/platform:rocm_rocdl_path", |
| ], |
| ) |
| |
| cc_library( |
| name = "cudnn_batchnorm_rewriter", |
| srcs = ["cudnn_batchnorm_rewriter.cc"], |
| hdrs = ["cudnn_batchnorm_rewriter.h"], |
| deps = [ |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:literal", |
| "//tensorflow/compiler/xla:literal_util", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| ], |
| ) |
| |
| cc_library( |
| name = "xfeed_queue", |
| hdrs = ["xfeed_queue.h"], |
| deps = [ |
| "//tensorflow/core:lib", |
| "@com_google_absl//absl/base:core_headers", |
| ], |
| ) |
| |
| cc_library( |
| name = "infeed_manager", |
| srcs = ["infeed_manager.cc"], |
| hdrs = ["infeed_manager.h"], |
| deps = [ |
| ":xfeed_queue", |
| "//tensorflow/compiler/xla:shape_tree", |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/core:stream_executor_no_cuda", |
| "@com_google_absl//absl/base:core_headers", |
| "@com_google_absl//absl/memory", |
| ], |
| ) |
| |
| cc_library( |
| name = "outfeed_manager", |
| srcs = ["outfeed_manager.cc"], |
| hdrs = ["outfeed_manager.h"], |
| deps = [ |
| ":xfeed_queue", |
| "//tensorflow/compiler/xla:literal", |
| "//tensorflow/compiler/xla:shape_tree", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/core:lib", |
| "@com_google_absl//absl/memory", |
| ], |
| ) |
| |
| cc_library( |
| name = "gpu_layout_assignment", |
| srcs = ["gpu_layout_assignment.cc"], |
| hdrs = ["gpu_layout_assignment.h"], |
| deps = [ |
| ":backend_configs_cc", |
| ":ir_emission_utils", |
| ":stream_executor_util", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:window_util", |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/service:computation_layout", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_casting_utils", |
| "//tensorflow/compiler/xla/service:layout_assignment", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:stream_executor_no_cuda", |
| ], |
| ) |
| |
| tf_cc_test( |
| name = "gpu_layout_assignment_test", |
| srcs = ["gpu_layout_assignment_test.cc"], |
| tags = tf_cuda_tests_tags(), |
| deps = [ |
| ":gemm_rewriter", |
| ":gpu_layout_assignment", |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:shape_layout", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/service:computation_layout", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_matchers", |
| "//tensorflow/compiler/xla/service:hlo_parser", |
| "//tensorflow/compiler/xla/tests:hlo_test_base", |
| "//tensorflow/compiler/xla/tests:xla_internal_test_main", # build_cleaner: keep |
| "//tensorflow/stream_executor/lib", |
| "@com_google_absl//absl/strings", |
| ], |
| ) |
| |
| cc_library( |
| name = "gpu_hlo_schedule", |
| srcs = ["gpu_hlo_schedule.cc"], |
| hdrs = ["gpu_hlo_schedule.h"], |
| deps = [ |
| ":stream_assignment", |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/compiler/xla/service:buffer_value", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_memory_scheduler", |
| "//tensorflow/compiler/xla/service:hlo_ordering", |
| "//tensorflow/compiler/xla/service:hlo_reachability", |
| "@com_google_absl//absl/memory", |
| ], |
| ) |
| |
| tf_cc_test( |
| name = "gpu_hlo_schedule_test", |
| srcs = [ |
| "gpu_hlo_schedule_test.cc", |
| ], |
| tags = ["no_pip"], |
| deps = [ |
| ":gpu_hlo_schedule", |
| ":stream_assignment", |
| "//tensorflow/compiler/xla:test_helpers", |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/tests:hlo_test_base", |
| "//tensorflow/compiler/xla/tests:test_utils", |
| "//tensorflow/compiler/xla/tests:xla_internal_test_main", |
| "@com_google_absl//absl/memory", |
| "@com_google_absl//absl/strings:str_format", |
| ], |
| ) |
| |
| tf_cc_test( |
| name = "while_transformer_test", |
| srcs = ["while_transformer_test.cc"], |
| tags = ["no_pip"], |
| deps = [ |
| ":instruction_fusion", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:test", |
| "//tensorflow/compiler/xla:test_helpers", |
| "//tensorflow/compiler/xla/service:copy_insertion", |
| "//tensorflow/compiler/xla/service:hlo_verifier", |
| "//tensorflow/compiler/xla/service:while_loop_analysis", |
| "//tensorflow/compiler/xla/tests:hlo_test_base", |
| "//tensorflow/compiler/xla/tests:xla_internal_test_main", |
| "//tensorflow/core:test", |
| ], |
| ) |
| |
| cc_library( |
| name = "stream_executor_util", |
| srcs = ["stream_executor_util.cc"], |
| hdrs = ["stream_executor_util.h"], |
| copts = tf_copts(), |
| deps = [ |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_module_config", |
| "//tensorflow/core:cuda_libdevice_path", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:lib_internal", |
| "//tensorflow/core:regexp_internal", |
| "//tensorflow/core:stream_executor_no_cuda", |
| "//tensorflow/core/profiler/lib:traceme", |
| "//tensorflow/stream_executor:kernel_spec", |
| "//tensorflow/stream_executor/gpu:gpu_asm_opts", |
| "@com_google_absl//absl/memory", |
| "@com_google_absl//absl/strings", |
| "@com_google_absl//absl/types:span", |
| ], |
| ) |
| |
| cc_library( |
| name = "buffer_comparator", |
| srcs = if_cuda_is_configured(["buffer_comparator.cc"]), |
| hdrs = if_cuda_is_configured(["buffer_comparator.h"]), |
| deps = if_cuda_is_configured([ |
| ":partition_assignment", |
| ":stream_executor_util", |
| "@com_google_absl//absl/base", |
| "@com_google_absl//absl/strings", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla/service:hlo_module_config", |
| "//tensorflow/core:stream_executor_no_cuda", |
| "//tensorflow/stream_executor:stream_executor_headers", |
| "//tensorflow/stream_executor/gpu:asm_compiler", |
| ]), |
| ) |
| |
| tf_cc_test( |
| name = "buffer_comparator_test", |
| srcs = if_cuda_is_configured(["buffer_comparator_test.cc"]), |
| tags = ["no_rocm"] + tf_cuda_tests_tags(), |
| deps = [ |
| "//tensorflow/core:test_main", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/core:test", |
| ] + if_cuda_is_configured([ |
| ":buffer_comparator", |
| "//tensorflow/core/platform/default/build_config:stream_executor_cuda", # build_cleaner: keep |
| "//tensorflow/stream_executor:device_memory", |
| ]), |
| ) |
| |
| cc_library( |
| name = "gpu_fusible", |
| srcs = ["gpu_fusible.cc"], |
| hdrs = ["gpu_fusible.h"], |
| deps = [ |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla/service:hlo", |
| ], |
| ) |
| |
| tf_cc_test( |
| name = "gpu_fusible_test", |
| srcs = ["gpu_fusible_test.cc"], |
| tags = ["no_pip"], |
| deps = [ |
| ":gpu_fusible", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_parser", |
| "//tensorflow/compiler/xla/tests:hlo_test_base", |
| "//tensorflow/compiler/xla/tests:xla_internal_test_main", |
| "@com_google_absl//absl/strings", |
| ], |
| ) |
| |
| cc_library( |
| name = "cudnn_fused_conv_rewriter", |
| srcs = ["cudnn_fused_conv_rewriter.cc"], |
| hdrs = ["cudnn_fused_conv_rewriter.h"], |
| deps = [ |
| ":backend_configs_cc", |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:literal_util", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_casting_utils", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/compiler/xla/service:pattern_matcher", |
| "//tensorflow/core:stream_executor_no_cuda", |
| ], |
| ) |
| |
| tf_cc_test( |
| name = "cudnn_fused_conv_rewriter_test", |
| srcs = ["cudnn_fused_conv_rewriter_test.cc"], |
| tags = [ |
| "gpu", |
| "no_oss", |
| "noasan", |
| "nomsan", |
| "requires-gpu-sm70", |
| ], |
| deps = [ |
| ":cudnn_fused_conv_rewriter", |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:test_helpers", |
| "//tensorflow/compiler/xla/service:hlo_parser", |
| "//tensorflow/compiler/xla/service/gpu/tests:gpu_codegen_test", |
| "//tensorflow/compiler/xla/tests:hlo_test_base", |
| "//tensorflow/core:test", |
| "//tensorflow/core:test_main", |
| "@com_google_absl//absl/strings", |
| ], |
| ) |
| |
| cc_library( |
| name = "variadic_op_splitter", |
| srcs = ["variadic_op_splitter.cc"], |
| hdrs = ["variadic_op_splitter.h"], |
| deps = [ |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/core:lib", |
| "@com_google_absl//absl/strings", |
| "@com_google_absl//absl/types:span", |
| ], |
| ) |
| |
| cc_library( |
| name = "gpu_scatter_expander", |
| srcs = ["gpu_scatter_expander.cc"], |
| hdrs = ["gpu_scatter_expander.h"], |
| deps = [ |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:scatter_expander", |
| "@com_google_absl//absl/algorithm:container", |
| ], |
| ) |
| |
| tf_cc_test( |
| name = "variadic_op_splitter_test", |
| srcs = ["variadic_op_splitter_test.cc"], |
| tags = ["no_pip"], |
| deps = [ |
| ":ir_emission_utils", |
| ":variadic_op_splitter", |
| "//tensorflow/compiler/xla:literal_util", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_matchers", |
| "//tensorflow/compiler/xla/service:hlo_parser", |
| "//tensorflow/compiler/xla/service:pattern_matcher", |
| "//tensorflow/compiler/xla/tests:hlo_test_base", |
| "//tensorflow/compiler/xla/tests:xla_internal_test_main", |
| ], |
| ) |
| |
| tf_proto_library_cc( |
| name = "gpu_autotuning_proto", |
| srcs = ["gpu_autotuning.proto"], |
| cc_api_version = 2, |
| protodeps = [ |
| "//tensorflow/compiler/xla:xla_data_proto", |
| "//tensorflow/compiler/xla/service:hlo_proto", |
| "//tensorflow/core:autotuning_proto", |
| ], |
| ) |
| |
| cc_library( |
| name = "hlo_algorithm_blacklist", |
| srcs = ["hlo_algorithm_blacklist.cc"], |
| hdrs = ["hlo_algorithm_blacklist.h"], |
| deps = [ |
| ":gpu_autotuning_proto_cc", |
| "//tensorflow/compiler/xla:debug_options_flags", |
| "//tensorflow/core:autotuning_proto_cc", |
| "//tensorflow/core:stream_executor_no_cuda", |
| "@com_google_absl//absl/container:flat_hash_map", |
| ], |
| ) |
| |
| tf_cc_test( |
| name = "hlo_algorithm_blacklist_test", |
| srcs = ["hlo_algorithm_blacklist_test.cc"], |
| data = ["data/hlo_algorithm_blacklist.pbtxt"], |
| tags = ["no_pip"], |
| deps = [ |
| ":hlo_algorithm_blacklist", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:test", |
| "//tensorflow/core:test_main", |
| "//tensorflow/core/platform:resource_loader", |
| "//tensorflow/stream_executor:dnn", |
| ], |
| ) |
| |
| cc_library( |
| name = "alias_passthrough_params", |
| srcs = ["alias_passthrough_params.cc"], |
| hdrs = ["alias_passthrough_params.h"], |
| deps = [ |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| ], |
| ) |
| |
| tf_cc_test( |
| name = "alias_passthrough_params_test", |
| srcs = ["alias_passthrough_params_test.cc"], |
| tags = ["no_pip"], |
| deps = [ |
| ":alias_passthrough_params", |
| "//tensorflow/compiler/xla/tests:hlo_test_base", |
| "//tensorflow/compiler/xla/tests:test_utils", |
| "//tensorflow/compiler/xla/tests:xla_internal_test_main", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:test", |
| ], |
| ) |
| |
| cc_library( |
| name = "horizontal_fusion", |
| srcs = ["horizontal_fusion.cc"], |
| hdrs = ["horizontal_fusion.h"], |
| deps = [ |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_creation_utils", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/core:lib", |
| "//tensorflow/core:lib_internal", |
| "@com_google_absl//absl/container:flat_hash_set", |
| "@com_google_absl//absl/types:span", |
| ], |
| ) |
| |
| tf_cc_test( |
| name = "horizontal_fusion_test", |
| srcs = ["horizontal_fusion_test.cc"], |
| deps = [ |
| ":fusion_merger", |
| ":horizontal_fusion", |
| ":instruction_fusion", |
| ":multi_output_fusion", |
| "//tensorflow/compiler/jit:xla_gpu_jit", |
| "//tensorflow/compiler/xla:literal", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:test", |
| "//tensorflow/compiler/xla:test_helpers", |
| "//tensorflow/compiler/xla/service:hlo_dce", |
| "//tensorflow/compiler/xla/service:hlo_matchers", |
| "//tensorflow/compiler/xla/service:hlo_parser", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/compiler/xla/service:hlo_pass_pipeline", |
| "//tensorflow/compiler/xla/service:tuple_simplifier", |
| "//tensorflow/compiler/xla/tests:filecheck", |
| "//tensorflow/compiler/xla/tests:hlo_test_base", |
| "//tensorflow/compiler/xla/tests:xla_internal_test_main", |
| ], |
| ) |
| |
| cc_library( |
| name = "reduction_degenerate_dim_remover", |
| srcs = ["reduction_degenerate_dim_remover.cc"], |
| hdrs = ["reduction_degenerate_dim_remover.h"], |
| deps = [ |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_casting_utils", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/compiler/xla/service:pattern_matcher", |
| "//tensorflow/core:lib", |
| "//tensorflow/stream_executor/lib", |
| "@com_google_absl//absl/algorithm:container", |
| "@com_google_absl//absl/strings", |
| "@com_google_absl//absl/types:optional", |
| ], |
| ) |
| |
| cc_library( |
| name = "reduction_dimension_grouper", |
| srcs = ["reduction_dimension_grouper.cc"], |
| hdrs = ["reduction_dimension_grouper.h"], |
| deps = [ |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_casting_utils", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/compiler/xla/service:pattern_matcher", |
| "//tensorflow/core:lib", |
| "//tensorflow/stream_executor/lib", |
| "@com_google_absl//absl/algorithm:container", |
| "@com_google_absl//absl/strings", |
| "@com_google_absl//absl/types:optional", |
| ], |
| ) |
| |
| cc_library( |
| name = "reduction_layout_normalizer", |
| srcs = ["reduction_layout_normalizer.cc"], |
| hdrs = ["reduction_layout_normalizer.h"], |
| deps = [ |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_casting_utils", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/compiler/xla/service:pattern_matcher", |
| "//tensorflow/core:lib", |
| "//tensorflow/stream_executor/lib", |
| "@com_google_absl//absl/algorithm:container", |
| "@com_google_absl//absl/strings", |
| "@com_google_absl//absl/types:optional", |
| ], |
| ) |
| |
| cc_library( |
| name = "tree_reduction_rewriter", |
| srcs = ["tree_reduction_rewriter.cc"], |
| hdrs = ["tree_reduction_rewriter.h"], |
| deps = [ |
| ":ir_emission_utils", |
| "//tensorflow/compiler/xla:shape_util", |
| "//tensorflow/compiler/xla:status_macros", |
| "//tensorflow/compiler/xla:statusor", |
| "//tensorflow/compiler/xla:types", |
| "//tensorflow/compiler/xla:util", |
| "//tensorflow/compiler/xla:window_util", |
| "//tensorflow/compiler/xla:xla_data_proto_cc", |
| "//tensorflow/compiler/xla/client:padding", |
| "//tensorflow/compiler/xla/service:hlo", |
| "//tensorflow/compiler/xla/service:hlo_casting_utils", |
| "//tensorflow/compiler/xla/service:hlo_creation_utils", |
| "//tensorflow/compiler/xla/service:hlo_evaluator", |
| "//tensorflow/compiler/xla/service:hlo_pass", |
| "//tensorflow/compiler/xla/service:shape_inference", |
| "//tensorflow/core:lib", |
| "//tensorflow/stream_executor/lib", |
| "@com_google_absl//absl/algorithm:container", |
| "@com_google_absl//absl/container:flat_hash_map", |
| "@com_google_absl//absl/container:flat_hash_set", |
| "@com_google_absl//absl/container:inlined_vector", |
| "@com_google_absl//absl/memory", |
| "@com_google_absl//absl/strings", |
| "@com_google_absl//absl/types:optional", |
| "@com_google_absl//absl/types:span", |
| ], |
| ) |
