diff --git a/tensorflow/compiler/xla/service/cpu/tests/cpu_noalias_test.cc b/tensorflow/compiler/xla/service/cpu/tests/cpu_noalias_test.cc index 8393ff4bb02..951098eb104 100644 --- a/tensorflow/compiler/xla/service/cpu/tests/cpu_noalias_test.cc +++ b/tensorflow/compiler/xla/service/cpu/tests/cpu_noalias_test.cc @@ -53,6 +53,8 @@ TEST_F(CpuNoAliasTest, Concat) { HloInstruction* concat2 = builder.AddInstruction(HloInstruction::CreateConcatenate( ShapeUtil::MakeShape(F32, {2, 6}), {concat1, param_x}, 1)); + HloInstruction* add = builder.AddInstruction(HloInstruction::CreateBinary( + ShapeUtil::MakeShape(F32, {2, 6}), HloOpcode::kAdd, concat2, concat2)); std::unique_ptr computation = builder.Build(); @@ -117,11 +119,21 @@ TEST_F(CpuNoAliasTest, Concat) { ->setName("read_concat2_array"); } + { + llvm::Value* concat2_val = ir_module.getOrInsertGlobal("add", array2d_type); + auto shape = ShapeUtil::MakeShape(F32, {2, 6}); + llvm_ir::IrArray add_array(concat2_val, shape); + aa.AddAliasingInformationToIrArray(*add, &add_array); + llvm_ir::IrArray::Index zero_2d({zero, zero}, shape, zero->getType()); + add_array.EmitReadArrayElement(zero_2d, &b)->setName("read_add_array"); + } + // Check the AA info in the loads. const char* filecheck_pattern = R"( CHECK: %read_param_x_array = load {{.*}} !noalias [[param_x_noalias:![0-9]+]] CHECK: %read_concat1_array = load {{.*}} !alias.scope [[concat1_scope:![0-9]+]], !noalias [[concat1_noalias:![0-9]+]] CHECK: %read_concat2_array = load {{.*}} !alias.scope [[concat1_noalias]], !noalias [[concat1_scope]] + CHECK: %read_add_array = load {{.*}} !alias.scope [[concat1_noalias]]{{$}} CHECK-DAG: [[buf_size32:![0-9]+]] = !{!"buffer:{{.*}} size:32 CHECK-DAG: [[buf_size48:![0-9]+]] = !{!"buffer:{{.*}} size:48 CHECK-DAG: [[param_x_noalias]] = !{[[buf_size48]], [[buf_size32]]} diff --git a/tensorflow/compiler/xla/service/llvm_ir/alias_analysis.cc b/tensorflow/compiler/xla/service/llvm_ir/alias_analysis.cc index c915a472707..0e1d392363b 100644 --- a/tensorflow/compiler/xla/service/llvm_ir/alias_analysis.cc +++ b/tensorflow/compiler/xla/service/llvm_ir/alias_analysis.cc @@ -63,7 +63,7 @@ void AliasAnalysis::AddAliasingInformationToIrArray(const HloInstruction& hlo, } if (module_.config().debug_options().xla_llvm_enable_noalias_metadata()) { - llvm::MDNode*& noalias_md = noalias_metadata_[buffer_slice]; + llvm::MDNode*& noalias_md = noalias_metadata_[{buffer_slice, &hlo}]; if (noalias_md == nullptr) { noalias_md = GetNoaliasMetadataForBuffer(buffer_slice, GetAliasDomain(), assignment_, hlo); diff --git a/tensorflow/compiler/xla/service/llvm_ir/alias_analysis.h b/tensorflow/compiler/xla/service/llvm_ir/alias_analysis.h index 12e2f449e23..7e7a6f6f820 100644 --- a/tensorflow/compiler/xla/service/llvm_ir/alias_analysis.h +++ b/tensorflow/compiler/xla/service/llvm_ir/alias_analysis.h @@ -79,9 +79,11 @@ class AliasAnalysis { absl::flat_hash_map alias_scope_metadata_; - // A map from a buffer slice to metadata corresponding to its noalias - // metadata. - absl::flat_hash_map noalias_metadata_; + // A map from a buffer slice and producer to metadata corresponding to its + // noalias metadata. + absl::flat_hash_map, + llvm::MDNode*> + noalias_metadata_; }; } // namespace llvm_ir