aboutsummaryrefslogtreecommitdiffhomepage
path: root/tensorflow/compiler/xla/service/gpu/ir_emitter_nested.cc
diff options
context:
space:
mode:
Diffstat (limited to 'tensorflow/compiler/xla/service/gpu/ir_emitter_nested.cc')
-rw-r--r--tensorflow/compiler/xla/service/gpu/ir_emitter_nested.cc16
1 files changed, 6 insertions, 10 deletions
diff --git a/tensorflow/compiler/xla/service/gpu/ir_emitter_nested.cc b/tensorflow/compiler/xla/service/gpu/ir_emitter_nested.cc
index c9574c87a3..5c827e5f9c 100644
--- a/tensorflow/compiler/xla/service/gpu/ir_emitter_nested.cc
+++ b/tensorflow/compiler/xla/service/gpu/ir_emitter_nested.cc
@@ -70,10 +70,10 @@ llvm::Function* IrEmitterNested::EmitBasePointersForNestedComputation(
argument_dereferenceable_bytes.push_back(root_size);
}
// The base pointer of the memory block for all pre-allocated temp buffers.
- argument_types.push_back(ir_builder_.getInt8PtrTy());
+ argument_types.push_back(b_.getInt8PtrTy());
llvm::FunctionType* function_type =
- llvm::FunctionType::get(ir_builder_.getVoidTy(), argument_types, false);
+ llvm::FunctionType::get(b_.getVoidTy(), argument_types, false);
llvm::Function* function = llvm::Function::Create(
function_type, // The function type.
llvm::GlobalValue::InternalLinkage, // The linkage type.
@@ -96,8 +96,7 @@ llvm::Function* IrEmitterNested::EmitBasePointersForNestedComputation(
llvm::BasicBlock::Create(function->getContext(), "entry", function);
// Emit a "return void" at entry_bb's end, and sets the insert point before
// that return instruction.
- ir_builder_.SetInsertPoint(
- llvm::ReturnInst::Create(function->getContext(), entry_bb));
+ b_.SetInsertPoint(llvm::ReturnInst::Create(function->getContext(), entry_bb));
std::vector<const HloInstruction*> non_io_hlos;
for (const auto* hlo : nested_computation.instructions()) {
@@ -127,20 +126,17 @@ Status IrEmitterNested::EmitTargetElementLoop(
target_arrays.push_back(GetIrArray(hlo, hlo, {i}));
}
TF_RETURN_IF_ERROR(
- llvm_ir::LoopEmitter(element_generator, target_arrays, &ir_builder_)
- .EmitLoop());
+ llvm_ir::LoopEmitter(element_generator, target_arrays, &b_).EmitLoop());
std::vector<llvm::Value*> tuple_operand_ptrs;
tuple_operand_ptrs.reserve(num_elems);
for (const llvm_ir::IrArray& array : target_arrays) {
tuple_operand_ptrs.push_back(array.GetBasePointer());
}
- llvm_ir::EmitTuple(GetIrArray(hlo, hlo), tuple_operand_ptrs, &ir_builder_,
- module_);
+ llvm_ir::EmitTuple(GetIrArray(hlo, hlo), tuple_operand_ptrs, &b_, module_);
return Status::OK();
}
- return llvm_ir::LoopEmitter(element_generator, GetIrArray(hlo, hlo),
- &ir_builder_)
+ return llvm_ir::LoopEmitter(element_generator, GetIrArray(hlo, hlo), &b_)
.EmitLoop();
}