[XLA] Stop using deprecated overloads of llvm::IRBuilder::CreateMemCpy
They're going away in
https://github.com/llvm/llvm-project/commit/39cfba9e333ce41b67c605173de9bf1575328b8a
PiperOrigin-RevId: 305066143
Change-Id: I142d91eab14ccb322c33446f604c63db05869bb3
diff --git a/tensorflow/compiler/xla/service/cpu/ir_emitter.cc b/tensorflow/compiler/xla/service/cpu/ir_emitter.cc
index d2b2b45..81bee71 100644
--- a/tensorflow/compiler/xla/service/cpu/ir_emitter.cc
+++ b/tensorflow/compiler/xla/service/cpu/ir_emitter.cc
@@ -241,7 +241,8 @@
/*Linkage=*/llvm::GlobalValue::PrivateLinkage,
/*Initializer=*/initializer,
/*Name=*/"");
- result_global->setAlignment(MinimumAlignmentForShape(literal.shape()));
+ result_global->setAlignment(
+ llvm::Align(MinimumAlignmentForShape(literal.shape())));
result_global->setUnnamedAddr(llvm::GlobalVariable::UnnamedAddr::Global);
return llvm::ConstantExpr::getBitCast(
result_global, IrShapeType(literal.shape())->getPointerTo());
@@ -520,12 +521,14 @@
if (kind == XfeedKind::kInfeed) {
// Copy to the program buffer address from the acquired buffer.
- MemCpy(program_buffer_address, /*DstAlign=*/1, acquired_pointer,
- /*SrcAlign=*/1, length_32);
+ MemCpy(program_buffer_address, /*DstAlign=*/llvm::Align(1),
+ acquired_pointer,
+ /*SrcAlign=*/llvm::Align(1), length_32);
} else {
// Outfeed -- copy from the in-program address to the acquired buffer.
- MemCpy(acquired_pointer, /*DstAlign=*/1, program_buffer_address,
- /*SrcAlign=*/1, length_32);
+ MemCpy(acquired_pointer, /*DstAlign=*/llvm::Align(1),
+ program_buffer_address,
+ /*SrcAlign=*/llvm::Align(1), length_32);
}
Call(release_func, {GetExecutableRunOptionsArgument(), b_.getInt32(length_32),
@@ -612,9 +615,9 @@
ShapeUtil::ByteSizeOfPrimitiveType(operand->shape().element_type());
auto source_buffer = GetEmittedValueFor(operand);
int64 size = ByteSizeOf(operand->shape());
- MemCpy(destination_addresses[i], /*DstAlign=*/primitive_type_size,
- source_buffer,
- /*SrcAlign=*/primitive_type_size, size);
+ MemCpy(destination_addresses[i],
+ /*DstAlign=*/llvm::Align(primitive_type_size), source_buffer,
+ /*SrcAlign=*/llvm::Align(primitive_type_size), size);
}
}
@@ -1401,8 +1404,8 @@
operand_ptrs.push_back(EmitBufferPointer(out_slice, operand_shape));
// TODO(b/63762267): Be more aggressive about specifying alignment.
- MemCpy(operand_ptrs.back(), /*DstAlign=*/1, in_ptr,
- /*SrcAlign=*/1, ShapeUtil::ByteSizeOf(operand_shape));
+ MemCpy(operand_ptrs.back(), /*DstAlign=*/llvm::Align(1), in_ptr,
+ /*SrcAlign=*/llvm::Align(1), ShapeUtil::ByteSizeOf(operand_shape));
}
llvm_ir::EmitTuple(GetIrArrayFor(crs), operand_ptrs, &b_);
return Status::OK();
@@ -2746,9 +2749,10 @@
element_alignment);
target_array.AnnotateLoadStoreInstructionWithMetadata(store_instruction);
} else {
- auto* memcpy_instruction = MemCpy(
- target, /*DstAlign=*/element_alignment, source,
- /*SrcAlign=*/element_alignment, element_count * primitive_type_size);
+ auto* memcpy_instruction =
+ MemCpy(target, /*DstAlign=*/llvm::Align(element_alignment), source,
+ /*SrcAlign=*/llvm::Align(element_alignment),
+ element_count * primitive_type_size);
// The memcpy does the load and the store internally. The aliasing related
// metadata has to reflect that.
@@ -3316,8 +3320,8 @@
llvm::Value* destination_value = GetEmittedValueFor(&destination);
int64 source_size = ByteSizeOf(source.shape());
// TODO(b/63762267): Be more aggressive about specifying alignment.
- MemCpy(destination_value, /*DstAlign=*/1, source_value,
- /*SrcAlign=*/1, source_size);
+ MemCpy(destination_value, /*DstAlign=*/llvm::Align(1), source_value,
+ /*SrcAlign=*/llvm::Align(1), source_size);
return Status::OK();
}
diff --git a/tensorflow/compiler/xla/service/gpu/ir_emitter_unnested.cc b/tensorflow/compiler/xla/service/gpu/ir_emitter_unnested.cc
index 6c5fe89..b8154b0 100644
--- a/tensorflow/compiler/xla/service/gpu/ir_emitter_unnested.cc
+++ b/tensorflow/compiler/xla/service/gpu/ir_emitter_unnested.cc
@@ -3268,7 +3268,7 @@
/*TLMode=*/llvm::GlobalValue::NotThreadLocal,
/*AddressSpace=*/global_address_space,
/*isExternallyInitialized=*/false);
- global_for_const->setAlignment(kConstantBufferAlignBytes);
+ global_for_const->setAlignment(llvm::Align(kConstantBufferAlignBytes));
ir_emitter_context_->llvm_module()->getGlobalList().push_back(
global_for_const);
}
diff --git a/tensorflow/compiler/xla/service/llvm_ir/tuple_ops.cc b/tensorflow/compiler/xla/service/llvm_ir/tuple_ops.cc
index 93dc66f..daf9847 100644
--- a/tensorflow/compiler/xla/service/llvm_ir/tuple_ops.cc
+++ b/tensorflow/compiler/xla/service/llvm_ir/tuple_ops.cc
@@ -54,8 +54,8 @@
llvm::Value* dst = select.GetBasePointer();
int64 table_size = ShapeUtil::ByteSizeOfTupleIndexTable(
select.GetShape(), module->getDataLayout().getPointerSize());
- b->CreateMemCpy(dst, /*DstAlign=*/1, src, /*SrcAlign=*/1,
- b->getInt64(table_size));
+ b->CreateMemCpy(dst, /*DstAlign=*/llvm::Align(1), src,
+ /*SrcAlign=*/llvm::Align(1), b->getInt64(table_size));
}
void EmitTuple(const IrArray& tuple, absl::Span<llvm::Value* const> operands,
