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

[XLA] Stop using deprecated overloads of llvm::IRBuilder::CreateMemCpy

Browse Source 
They're going away in
39cfba9e33

PiperOrigin-RevId: 305066143
Change-Id: I142d91eab14ccb322c33446f604c63db05869bb3
This commit is contained in:
Benjamin Kramer 2020-04-06 10:48:02 -07:00 committed by TensorFlower Gardener
parent b57674fe1d
commit 1e8fffdfb1
3 changed files with 22 additions and 18 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -241,7 +241,8 @@ llvm::Constant* IrEmitter::EmitGlobalForLiteral(const Literal& literal) {
/*Linkage=*/llvm::GlobalValue::PrivateLinkage, /*Linkage=*/llvm::GlobalValue::PrivateLinkage,
/*Initializer=*/initializer, /*Initializer=*/initializer,
/*Name=*/""); /*Name=*/"");
result_global->setAlignment(MinimumAlignmentForShape(literal.shape())); result_global->setAlignment(
llvm::Align(MinimumAlignmentForShape(literal.shape())));
result_global->setUnnamedAddr(llvm::GlobalVariable::UnnamedAddr::Global); result_global->setUnnamedAddr(llvm::GlobalVariable::UnnamedAddr::Global);
return llvm::ConstantExpr::getBitCast( return llvm::ConstantExpr::getBitCast(
result_global, IrShapeType(literal.shape())->getPointerTo()); result_global, IrShapeType(literal.shape())->getPointerTo());
@ -520,12 +521,14 @@ Status IrEmitter::EmitXfeedTransfer(XfeedKind kind, const Shape& shape,
if (kind == XfeedKind::kInfeed) { if (kind == XfeedKind::kInfeed) {
// Copy to the program buffer address from the acquired buffer. // Copy to the program buffer address from the acquired buffer.
MemCpy(program_buffer_address, /*DstAlign=*/1, acquired_pointer, MemCpy(program_buffer_address, /*DstAlign=*/llvm::Align(1),
/*SrcAlign=*/1, length_32); acquired_pointer,
/*SrcAlign=*/llvm::Align(1), length_32);
} else { } else {
// Outfeed -- copy from the in-program address to the acquired buffer. // Outfeed -- copy from the in-program address to the acquired buffer.
MemCpy(acquired_pointer, /*DstAlign=*/1, program_buffer_address, MemCpy(acquired_pointer, /*DstAlign=*/llvm::Align(1),
/*SrcAlign=*/1, length_32); program_buffer_address,
/*SrcAlign=*/llvm::Align(1), length_32);
} }
Call(release_func, {GetExecutableRunOptionsArgument(), b_.getInt32(length_32), Call(release_func, {GetExecutableRunOptionsArgument(), b_.getInt32(length_32),
@ -612,9 +615,9 @@ Status IrEmitter::HandleSort(HloInstruction* hlo) {
ShapeUtil::ByteSizeOfPrimitiveType(operand->shape().element_type()); ShapeUtil::ByteSizeOfPrimitiveType(operand->shape().element_type());
auto source_buffer = GetEmittedValueFor(operand); auto source_buffer = GetEmittedValueFor(operand);
int64 size = ByteSizeOf(operand->shape()); int64 size = ByteSizeOf(operand->shape());
MemCpy(destination_addresses[i], /*DstAlign=*/primitive_type_size, MemCpy(destination_addresses[i],
source_buffer, /*DstAlign=*/llvm::Align(primitive_type_size), source_buffer,
/*SrcAlign=*/primitive_type_size, size); /*SrcAlign=*/llvm::Align(primitive_type_size), size);
} }
} }
@ -1401,8 +1404,8 @@ Status IrEmitter::HandleAllReduceSingleReplica(HloInstruction* crs) {
operand_ptrs.push_back(EmitBufferPointer(out_slice, operand_shape)); operand_ptrs.push_back(EmitBufferPointer(out_slice, operand_shape));
// TODO(b/63762267): Be more aggressive about specifying alignment. // TODO(b/63762267): Be more aggressive about specifying alignment.
MemCpy(operand_ptrs.back(), /*DstAlign=*/1, in_ptr, MemCpy(operand_ptrs.back(), /*DstAlign=*/llvm::Align(1), in_ptr,
/*SrcAlign=*/1, ShapeUtil::ByteSizeOf(operand_shape)); /*SrcAlign=*/llvm::Align(1), ShapeUtil::ByteSizeOf(operand_shape));
} }
llvm_ir::EmitTuple(GetIrArrayFor(crs), operand_ptrs, &b_); llvm_ir::EmitTuple(GetIrArrayFor(crs), operand_ptrs, &b_);
return Status::OK(); return Status::OK();
@ -2746,9 +2749,10 @@ void IrEmitter::EmitTransferElements(llvm::Value* target, llvm::Value* source,
element_alignment); element_alignment);
target_array.AnnotateLoadStoreInstructionWithMetadata(store_instruction); target_array.AnnotateLoadStoreInstructionWithMetadata(store_instruction);
} else { } else {
auto* memcpy_instruction = MemCpy( auto* memcpy_instruction =
target, /*DstAlign=*/element_alignment, source, MemCpy(target, /*DstAlign=*/llvm::Align(element_alignment), source,
/*SrcAlign=*/element_alignment, element_count * primitive_type_size); /*SrcAlign=*/llvm::Align(element_alignment),
element_count * primitive_type_size);
// The memcpy does the load and the store internally. The aliasing related // The memcpy does the load and the store internally. The aliasing related
// metadata has to reflect that. // metadata has to reflect that.
@ -3316,8 +3320,8 @@ Status IrEmitter::EmitMemcpy(const HloInstruction& source,
llvm::Value* destination_value = GetEmittedValueFor(&destination); llvm::Value* destination_value = GetEmittedValueFor(&destination);
int64 source_size = ByteSizeOf(source.shape()); int64 source_size = ByteSizeOf(source.shape());
// TODO(b/63762267): Be more aggressive about specifying alignment. // TODO(b/63762267): Be more aggressive about specifying alignment.
MemCpy(destination_value, /*DstAlign=*/1, source_value, MemCpy(destination_value, /*DstAlign=*/llvm::Align(1), source_value,
/*SrcAlign=*/1, source_size); /*SrcAlign=*/llvm::Align(1), source_size);
return Status::OK(); return Status::OK();
} }

2
tensorflow/compiler/xla/service/gpu/ir_emitter_unnested.cc
View File

@ -3268,7 +3268,7 @@ Status IrEmitterUnnested::EmitConstantGlobals() {
/*TLMode=*/llvm::GlobalValue::NotThreadLocal, /*TLMode=*/llvm::GlobalValue::NotThreadLocal,
/*AddressSpace=*/global_address_space, /*AddressSpace=*/global_address_space,
/*isExternallyInitialized=*/false); /*isExternallyInitialized=*/false);
global_for_const->setAlignment(kConstantBufferAlignBytes); global_for_const->setAlignment(llvm::Align(kConstantBufferAlignBytes));
ir_emitter_context_->llvm_module()->getGlobalList().push_back( ir_emitter_context_->llvm_module()->getGlobalList().push_back(
global_for_const); global_for_const);
} }

4
tensorflow/compiler/xla/service/llvm_ir/tuple_ops.cc
View File

@ -54,8 +54,8 @@ void EmitTupleSelect(const IrArray& select, const IrArray& pred,
llvm::Value* dst = select.GetBasePointer(); llvm::Value* dst = select.GetBasePointer();
int64 table_size = ShapeUtil::ByteSizeOfTupleIndexTable( int64 table_size = ShapeUtil::ByteSizeOfTupleIndexTable(
select.GetShape(), module->getDataLayout().getPointerSize()); select.GetShape(), module->getDataLayout().getPointerSize());
b->CreateMemCpy(dst, /*DstAlign=*/1, src, /*SrcAlign=*/1, b->CreateMemCpy(dst, /*DstAlign=*/llvm::Align(1), src,
b->getInt64(table_size)); /*SrcAlign=*/llvm::Align(1), b->getInt64(table_size));
} }
void EmitTuple(const IrArray& tuple, absl::Span<llvm::Value* const> operands, void EmitTuple(const IrArray& tuple, absl::Span<llvm::Value* const> operands,