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

Integrate LLVM at llvm/llvm-project@bf6380c096

Browse Source 
Updates LLVM usage to match
[bf6380c0966b](https://github.com/llvm/llvm-project/commit/bf6380c0966b)

PiperOrigin-RevId: 359728311
Change-Id: I154af13cedad41b7ce411964dcedce519d102632
This commit is contained in:
A. Unique TensorFlower 2021-02-26 04:03:59 -08:00 committed by TensorFlower Gardener
parent 362f9fc585
commit 7e1c942100
2 changed files with 10 additions and 8 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
tensorflow/compiler/xla/service/gpu/ir_emitter.cc
View File

@ -299,6 +299,7 @@ bool IrEmitter::MaybeEmitDirectAtomicOperation(
(f64_atomic_add_supported && element_type == F64); (f64_atomic_add_supported && element_type == F64);
if (atomic_add_supported) { if (atomic_add_supported) {
AtomicRMW(llvm::AtomicRMWInst::FAdd, output_address, source, AtomicRMW(llvm::AtomicRMWInst::FAdd, output_address, source,
llvm::MaybeAlign(),
llvm::AtomicOrdering::SequentiallyConsistent); llvm::AtomicOrdering::SequentiallyConsistent);
return true; return true;
} }
@ -307,6 +308,7 @@ bool IrEmitter::MaybeEmitDirectAtomicOperation(
if (is_atomic_integral) { if (is_atomic_integral) {
// integral + integral // integral + integral
AtomicRMW(llvm::AtomicRMWInst::Add, output_address, source, AtomicRMW(llvm::AtomicRMWInst::Add, output_address, source,
llvm::MaybeAlign(),
llvm::AtomicOrdering::SequentiallyConsistent); llvm::AtomicOrdering::SequentiallyConsistent);
return true; return true;
} }
@ -318,7 +320,7 @@ bool IrEmitter::MaybeEmitDirectAtomicOperation(
auto opcode = primitive_util::IsSignedIntegralType(element_type) auto opcode = primitive_util::IsSignedIntegralType(element_type)
? llvm::AtomicRMWInst::Max ? llvm::AtomicRMWInst::Max
: llvm::AtomicRMWInst::UMax; : llvm::AtomicRMWInst::UMax;
AtomicRMW(opcode, output_address, source, AtomicRMW(opcode, output_address, source, llvm::MaybeAlign(),
llvm::AtomicOrdering::SequentiallyConsistent); llvm::AtomicOrdering::SequentiallyConsistent);
return true; return true;
} }
@ -328,7 +330,7 @@ bool IrEmitter::MaybeEmitDirectAtomicOperation(
auto opcode = primitive_util::IsSignedIntegralType(element_type) auto opcode = primitive_util::IsSignedIntegralType(element_type)
? llvm::AtomicRMWInst::Min ? llvm::AtomicRMWInst::Min
: llvm::AtomicRMWInst::UMin; : llvm::AtomicRMWInst::UMin;
AtomicRMW(opcode, output_address, source, AtomicRMW(opcode, output_address, source, llvm::MaybeAlign(),
llvm::AtomicOrdering::SequentiallyConsistent); llvm::AtomicOrdering::SequentiallyConsistent);
return true; return true;
} }
@ -476,8 +478,8 @@ Status IrEmitter::EmitAtomicOperationUsingCAS(const HloComputation& computation,
// Emit code to perform the atomicCAS operation // Emit code to perform the atomicCAS operation
// (cas_old_output, success) = atomicCAS(memory_address, cas_old_output, // (cas_old_output, success) = atomicCAS(memory_address, cas_old_output,
// cas_new_output); // cas_new_output);
llvm::Value* ret_value = llvm::Value* ret_value = AtomicCmpXchg(
AtomicCmpXchg(atomic_memory_address, cas_old_output, cas_new_output, atomic_memory_address, cas_old_output, cas_new_output, llvm::MaybeAlign(),
llvm::AtomicOrdering::SequentiallyConsistent, llvm::AtomicOrdering::SequentiallyConsistent,
llvm::AtomicOrdering::SequentiallyConsistent); llvm::AtomicOrdering::SequentiallyConsistent);

4
third_party/llvm/workspace.bzl vendored
View File

@ -4,8 +4,8 @@ load("//third_party:repo.bzl", "tf_http_archive")
def repo(name): def repo(name):
"""Imports LLVM.""" """Imports LLVM."""
LLVM_COMMIT = "5d64dd8e3c22e12e4f7b3d08ffe88fc41e727117" LLVM_COMMIT = "bf6380c0966b26a2aec7f2072efd0a1a9b6328f2"
LLVM_SHA256 = "19131ecf82681acec18780a6b9dbe2412e35132ccab104074951b6d50407d08d" LLVM_SHA256 = "b225983215a6ba47cbc344df25c158c60baf8cf2cfe5e990e13e028f7ff6dcdc"
tf_http_archive( tf_http_archive(
name = name, name = name,