diff --git a/tensorflow/compiler/xla/service/cpu/ir_emitter.cc b/tensorflow/compiler/xla/service/cpu/ir_emitter.cc
index d2b2b4534a5..81bee71360d 100644
--- a/tensorflow/compiler/xla/service/cpu/ir_emitter.cc
+++ b/tensorflow/compiler/xla/service/cpu/ir_emitter.cc
@@ -241,7 +241,8 @@ llvm::Constant* IrEmitter::EmitGlobalForLiteral(const Literal& literal) {
       /*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 @@ Status IrEmitter::EmitXfeedTransfer(XfeedKind kind, const Shape& shape,
 
   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 @@ Status IrEmitter::HandleSort(HloInstruction* hlo) {
           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 @@ Status IrEmitter::HandleAllReduceSingleReplica(HloInstruction* crs) {
     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 @@ void IrEmitter::EmitTransferElements(llvm::Value* target, llvm::Value* source,
                      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 @@ Status IrEmitter::EmitMemcpy(const HloInstruction& source,
   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 6c5fe891360..b8154b0e157 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 @@ Status IrEmitterUnnested::EmitConstantGlobals() {
         /*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 93dc66f9ac1..daf98478194 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 @@ void EmitTupleSelect(const IrArray& select, const IrArray& pred,
   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,