experiments/STT-tensorflow
1
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

[XLA:CPU] Always set the unnamed_addr bit on constants

Browse Source 
This allows LLVM (and the linker) to merge or pool the constants. In many cases
LLVM can derive unnamed_addr by itself, but not if it's a constant being passed
to a runtime function.

PiperOrigin-RevId: 236298140
This commit is contained in:
Benjamin Kramer 2019-03-01 05:00:23 -08:00 committed by TensorFlower Gardener
parent 29d87c6c77
commit a9b3864070
5 changed files with 13 additions and 11 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
tensorflow/compiler/xla/service/cpu/ir_emitter.cc
View File

@ -189,6 +189,7 @@ llvm::Constant* IrEmitter::EmitGlobalForLiteral(const Literal& literal) {
/*Initializer=*/initializer,
/*Name=*/"");
result_global->setAlignment(MinimumAlignmentForShape(literal.shape()));
result_global->setUnnamedAddr(llvm::GlobalVariable::UnnamedAddr::Global);
return llvm::ConstantExpr::getBitCast(
result_global, IrShapeType(literal.shape())->getPointerTo());
}

8
tensorflow/compiler/xla/service/cpu/tests/cpu_external_constants_test.cc
View File

@ -56,8 +56,8 @@ class CpuExternalConstantsTest : public CpuCodegenTest {
TEST_F(CpuExternalConstantsTest, Basic) {
TestWithArray(/*rows=*/1024, /*cols=*/1024, R"(
CHECK-NOT: @constant_global_0 = external constant [1024 x [1024 x float]], align 16
CHECK: @0 = private constant [4194304 x i8] {{.*}}, align 16
CHECK-NOT: @constant_global_0 = external unnamed_addr constant [1024 x [1024 x float]], align 16
CHECK: @0 = private unnamed_addr constant [4194304 x i8] {{.*}}, align 16
)");
}
@ -65,8 +65,8 @@ TEST_F(CpuExternalConstantsTest, BasicNegative) {
// The constant array in this test case is small enough that there is no need
// to externalize it.
TestWithArray(/*rows=*/4, /*cols=*/4, R"(
CHECK-NOT: @constant_global_0 = external constant [16 x float], align 8
CHECK: @0 = private constant [64 x i8] {{.*}}, align 8
CHECK-NOT: @constant_global_0 = external unnamed_addr constant [16 x float], align 8
CHECK: @0 = private unnamed_addr constant [64 x i8] {{.*}}, align 8
)");
}
} // namespace

12
tensorflow/compiler/xla/service/cpu/tests/cpu_literal_caching_test.cc
View File

@ -56,8 +56,8 @@ ENTRY main {
)";
string filecheck_pattern = R"(
CHECK: private constant [48 x i8]
CHECK-NOT: private constant [48 x i8]
CHECK: private unnamed_addr constant [48 x i8]
CHECK-NOT: private unnamed_addr constant [48 x i8]
)";
TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr<HloModule> module,
@ -102,10 +102,10 @@ ENTRY main {
)";
string filecheck_pattern = R"(
CHECK-DAG: private constant [4 x i8]
CHECK-DAG: private constant [8 x i8]
CHECK-NOT: private constant [4 x i8]
CHECK-NOT: private constant [8 x i8]
CHECK-DAG: private unnamed_addr constant [4 x i8]
CHECK-DAG: private unnamed_addr constant [8 x i8]
CHECK-NOT: private unnamed_addr constant [4 x i8]
CHECK-NOT: private unnamed_addr constant [8 x i8]
)";
TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr<HloModule> module,

2
tensorflow/compiler/xla/service/cpu/tests/cpu_outfeed_test.cc
View File

@ -38,7 +38,7 @@ ENTRY main {
)";
string filecheck_pattern = R"(
CHECK: private constant [48 x i8]
CHECK: private unnamed_addr constant [48 x i8]
)";
TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr<HloModule> module,

1
tensorflow/compiler/xla/service/llvm_ir/fused_ir_emitter.cc
View File

@ -82,6 +82,7 @@ Status FusedIrEmitter::HandleConstant(HloInstruction* constant) {
/*Linkage=*/llvm::GlobalValue::PrivateLinkage,
/*Initializer=*/initializer,
/*Name=*/"");
global->setUnnamedAddr(llvm::GlobalVariable::UnnamedAddr::Global);
llvm::Constant* shape_constant = llvm::ConstantExpr::getBitCast(
global,
llvm_ir::ShapeToIrType(literal.shape(), module_)->getPointerTo());