OSSForks/tensorflow
1
0
Fork 0
mirror of https://github.com/zebrajr/tensorflow.git synced 2025-12-06 12:20:11 +01:00
Code Issues Actions 5 Packages Projects Releases Wiki Activity

[XLA] Remove verify_unique_channel_ids verifier option.

Browse Source 
The functionality has been removed previously, but the option was never cleaned up. This does not remove the xla_ignore_channel_id debug option because it also has a non-verifier use.

PiperOrigin-RevId: 821737613
This commit is contained in:
Michael Kuperstein 2025-10-20 11:47:03 -07:00 committed by TensorFlower Gardener
parent 16e1567819
commit b824d4e187
5 changed files with 3 additions and 14 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
third_party/xla/xla/service/gpu/fusion_pipeline.cc vendored
View File

@ -53,7 +53,6 @@ HloPassPipeline FusionPipeline(
HloVerifierOpts opts = HloVerifierOpts opts =
HloVerifierOpts().MakeLayoutSensitive().WithInstructionCanChangeLayout( HloVerifierOpts().MakeLayoutSensitive().WithInstructionCanChangeLayout(
LayoutAssignment::InstructionCanChangeLayout); LayoutAssignment::InstructionCanChangeLayout);
opts.verify_unique_channel_ids = !debug_options.xla_ignore_channel_id();
fusion.AddInvariantCheckerDebug<HloVerifier>( fusion.AddInvariantCheckerDebug<HloVerifier>(
std::make_unique<CpuGpuVerifierMetadata>(std::move(opts)), std::make_unique<CpuGpuVerifierMetadata>(std::move(opts)),
"hlo verifier (debug)"); "hlo verifier (debug)");

11
third_party/xla/xla/service/gpu/gpu_compiler.cc vendored
View File

@ -571,9 +571,7 @@ GpuCompiler::GpuCompiler(se::Platform::Id platform_id,
namespace { namespace {
// Adds the HloVerifier for GPU to the given pipeline. // Adds the HloVerifier for GPU to the given pipeline.
void AddHloVerifier(HloPassPipeline* pipeline, void AddHloVerifier(HloPassPipeline* pipeline,
bool verify_unique_channel_ids = false,
HloVerifierOpts&& opts = {}, bool debug_only = false) { HloVerifierOpts&& opts = {}, bool debug_only = false) {
opts.verify_unique_channel_ids = verify_unique_channel_ids;
opts.verify_no_collective_deadlocks = true; opts.verify_no_collective_deadlocks = true;
std::unique_ptr<TargetVerifierMetadata> verifier_metadata = std::unique_ptr<TargetVerifierMetadata> verifier_metadata =
std::make_unique<CpuGpuVerifierMetadata>(std::move(opts)); std::make_unique<CpuGpuVerifierMetadata>(std::move(opts));
@ -746,7 +744,7 @@ absl::Status RunOptimizationPasses(
gpu_target_config.device_description.gpu_compute_capability(); gpu_target_config.device_description.gpu_compute_capability();
HloPassPipeline pipeline("optimization"); HloPassPipeline pipeline("optimization");
AddHloVerifier(&pipeline, !debug_options.xla_ignore_channel_id()); AddHloVerifier(&pipeline);
if (debug_options.xla_detect_unstable_reductions() != if (debug_options.xla_detect_unstable_reductions() !=
DebugOptions::UNSTABLE_REDUCTION_DETECTION_MODE_NONE) { DebugOptions::UNSTABLE_REDUCTION_DETECTION_MODE_NONE) {
pipeline.AddPass<UnstableReductionDetector>(); pipeline.AddPass<UnstableReductionDetector>();
@ -902,8 +900,7 @@ absl::Status RunOptimizationPasses(
// point. // point.
[&, &pipeline = [&, &pipeline =
pipeline.AddPass<HloPassFix<HloPassPipeline>>("simplification")] { pipeline.AddPass<HloPassFix<HloPassPipeline>>("simplification")] {
AddHloVerifier(&pipeline, !debug_options.xla_ignore_channel_id(), AddHloVerifier(&pipeline, HloVerifierOpts{}, /*debug_only=*/true);
HloVerifierOpts{}, /*debug_only=*/true);
// BatchNormExpander can create zero-sized ops, so zero-sized HLO // BatchNormExpander can create zero-sized ops, so zero-sized HLO
// elimination has to come after that pass. // elimination has to come after that pass.
@ -1856,7 +1853,7 @@ absl::Status GpuCompiler::OptimizeHloPostLayoutAssignment(
} }
HloPassPipeline pipeline("post-layout_assignment"); HloPassPipeline pipeline("post-layout_assignment");
AddHloVerifier(&pipeline, !debug_options.xla_ignore_channel_id(), AddHloVerifier(&pipeline,
HloVerifierOpts{} HloVerifierOpts{}
.MakeLayoutSensitive() .MakeLayoutSensitive()
.WithInstructionCanChangeLayout( .WithInstructionCanChangeLayout(
@ -1963,7 +1960,6 @@ absl::Status GpuCompiler::OptimizeHloPostLayoutAssignment(
LayoutAssignment::InstructionCanChangeLayout) LayoutAssignment::InstructionCanChangeLayout)
.VerifyBroadcastDimensionsOrder() .VerifyBroadcastDimensionsOrder()
.VerifyReshapeIsBitcast(); .VerifyReshapeIsBitcast();
opts.verify_unique_channel_ids = !debug_options.xla_ignore_channel_id();
pipeline.AddPass<HloVerifier>( pipeline.AddPass<HloVerifier>(
std::make_unique<DefaultVerifierMetadata>(std::move(opts)), std::make_unique<DefaultVerifierMetadata>(std::move(opts)),
"end-of-post-layout_assignment"); "end-of-post-layout_assignment");
@ -3056,7 +3052,6 @@ absl::Status GpuCompiler::RunPostSchedulingPipelines(
if (module->config().debug_options().xla_gpu_pgle_accuracy_checker() == if (module->config().debug_options().xla_gpu_pgle_accuracy_checker() ==
DebugOptions::PGLE_STRICTNESS_LEVEL_ERROR) { DebugOptions::PGLE_STRICTNESS_LEVEL_ERROR) {
AddHloVerifier(&main_pipeline, AddHloVerifier(&main_pipeline,
module->config().debug_options().xla_ignore_channel_id(),
HloVerifierOpts{}.VerifyInstructionNameUnchanged()); HloVerifierOpts{}.VerifyInstructionNameUnchanged());
} }
return main_pipeline.Run(module, {HloInstruction::kMainExecutionThread}) return main_pipeline.Run(module, {HloInstruction::kMainExecutionThread})

1
third_party/xla/xla/service/gpu/pre_scheduling_copy_insertion_pipeline.cc vendored
View File

@ -52,7 +52,6 @@ HloPassPipeline PreSchedulingCopyInsertionPipeline(
HloVerifierOpts opts = HloVerifierOpts opts =
HloVerifierOpts{}.MakeLayoutSensitive().WithInstructionCanChangeLayout( HloVerifierOpts{}.MakeLayoutSensitive().WithInstructionCanChangeLayout(
LayoutAssignment::InstructionCanChangeLayout); LayoutAssignment::InstructionCanChangeLayout);
opts.verify_unique_channel_ids = !debug_options.xla_ignore_channel_id();
std::unique_ptr<TargetVerifierMetadata> verifier_metadata = std::unique_ptr<TargetVerifierMetadata> verifier_metadata =
std::make_unique<CpuGpuVerifierMetadata>(std::move(opts)); std::make_unique<CpuGpuVerifierMetadata>(std::move(opts));
pipeline.AddInvariantCheckerDebug<HloVerifier>(std::move(verifier_metadata), pipeline.AddInvariantCheckerDebug<HloVerifier>(std::move(verifier_metadata),

3
third_party/xla/xla/service/hlo_verifier.h vendored
View File

@ -175,9 +175,6 @@ struct HloVerifierOpts {
// cloned (".clone" suffix) or rematted (".remat"); // cloned (".clone" suffix) or rematted (".remat");
bool verify_instruction_name_unchanged = false; bool verify_instruction_name_unchanged = false;
// Check if channel instructions all have unique channel ids.
bool verify_unique_channel_ids = true;
// Check if a shape has a host memory space color // Check if a shape has a host memory space color
bool verify_no_host_memory_space = false; bool verify_no_host_memory_space = false;

1
third_party/xla/xla/service/hlo_verifier_test.cc vendored
View File

@ -3658,7 +3658,6 @@ TEST_F(HloVerifierTest, NoErrorOnDuplicateChannelId) {
TF_ASSERT_OK_AND_ASSIGN(auto module, TF_ASSERT_OK_AND_ASSIGN(auto module,
ParseAndReturnUnverifiedModule(hlo_string)); ParseAndReturnUnverifiedModule(hlo_string));
HloVerifierOpts opts{}; HloVerifierOpts opts{};
opts.verify_unique_channel_ids = false;
HloVerifier verifier(std::move(opts)); HloVerifier verifier(std::move(opts));
ASSERT_IS_OK(verifier.Run(module.get()).status()); ASSERT_IS_OK(verifier.Run(module.get()).status());
} }