| Commit message (Collapse) | Author | Age | Files | Lines |
|---|
| |
|
|
|
|
|
|
| |
The reversion apparently deleted the test/Transforms directory.
Will be re-reverting again.
llvm-svn: 358552
|
| |
|
|
|
|
|
|
| |
As it's causing some bot failures (and per request from kbarton).
This reverts commit r358543/ab70da07286e618016e78247e4a24fcb84077fda.
llvm-svn: 358546
|
| |
|
|
|
|
| |
Differential Revision: https://reviews.llvm.org/D54848
llvm-svn: 348570
|
| |
|
|
| |
llvm-svn: 342131
|
| |
|
|
|
|
|
|
| |
This reapplies commit r337489 reverted by r337541
Additionally, this commit contains a speculative fix to the issue reported in r337541
(the report does not contain an actionable reproducer, just a stack trace)
llvm-svn: 337606
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
This reverts commit r337489.
It causes asserts to fire in some TensorFlow tests, e.g.
tensorflow/compiler/tests/gather_test.py on GPU.
Example stack trace:
Start test case: GatherTest.testHigherRank
assertion failed at third_party/llvm/llvm/lib/Support/APInt.cpp:819 in llvm::APInt llvm::APInt::trunc(unsigned int) const: width && "Can't truncate to 0 bits"
@ 0x5559446ebe10 __assert_fail
@ 0x55593ef32f5e llvm::APInt::trunc()
@ 0x55593d78f86e (anonymous namespace)::Vectorizer::lookThroughComplexAddresses()
@ 0x55593d78f2bc (anonymous namespace)::Vectorizer::areConsecutivePointers()
@ 0x55593d78d128 (anonymous namespace)::Vectorizer::isConsecutiveAccess()
@ 0x55593d78c926 (anonymous namespace)::Vectorizer::vectorizeInstructions()
@ 0x55593d78c221 (anonymous namespace)::Vectorizer::vectorizeChains()
@ 0x55593d78b948 (anonymous namespace)::Vectorizer::run()
@ 0x55593d78b725 (anonymous namespace)::LoadStoreVectorizer::runOnFunction()
@ 0x55593edf4b17 llvm::FPPassManager::runOnFunction()
@ 0x55593edf4e55 llvm::FPPassManager::runOnModule()
@ 0x55593edf563c (anonymous namespace)::MPPassManager::runOnModule()
@ 0x55593edf5137 llvm::legacy::PassManagerImpl::run()
@ 0x55593edf5b71 llvm::legacy::PassManager::run()
@ 0x55593ced250d xla::gpu::IrDumpingPassManager::run()
@ 0x55593ced5033 xla::gpu::(anonymous namespace)::EmitModuleToPTX()
@ 0x55593ced40ba xla::gpu::(anonymous namespace)::CompileModuleToPtx()
@ 0x55593ced33d0 xla::gpu::CompileToPtx()
@ 0x55593b26b2a2 xla::gpu::NVPTXCompiler::RunBackend()
@ 0x55593b21f973 xla::Service::BuildExecutable()
@ 0x555938f44e64 xla::LocalService::CompileExecutable()
@ 0x555938f30a85 xla::LocalClient::Compile()
@ 0x555938de3c29 tensorflow::XlaCompilationCache::BuildExecutable()
@ 0x555938de4e9e tensorflow::XlaCompilationCache::CompileImpl()
@ 0x555938de3da5 tensorflow::XlaCompilationCache::Compile()
@ 0x555938c5d962 tensorflow::XlaLocalLaunchBase::Compute()
@ 0x555938c68151 tensorflow::XlaDevice::Compute()
@ 0x55593f389e1f tensorflow::(anonymous namespace)::ExecutorState::Process()
@ 0x55593f38a625 tensorflow::(anonymous namespace)::ExecutorState::ScheduleReady()::$_1::operator()()
*** SIGABRT received by PID 7798 (TID 7837) from PID 7798; ***
llvm-svn: 337541
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
This is mostly a preparation work for adding a limited support for
select instructions. It proved to be difficult to do due to size and
irregularity of Vectorizer::isConsecutiveAccess, this is fixed here I
believe.
It also turned out that these changes make it simpler to finish one of
the TODOs and fix a number of other small issues, namely:
1. Looking through bitcasts to a type of a different size (requires
careful tracking of the original load/store size and some math
converting sizes in bytes to expected differences in indices of GEPs).
2. Reusing partial analysis of pointers done by first attempt in proving
them consecutive instead of starting from scratch. This added limited
support for nested GEPs co-existing with difficult sext/zext
instructions. This also required a careful handling of negative
differences between constant parts of offsets.
3. Handing a case where the first pointer index is not an add, but
something else (a function parameter for instance).
I observe an increased number of successful vectorizations on a large
set of shader programs. Only few shaders are affected, but those that
are affected sport >5% less loads and stores than before the patch.
Reviewed By: rampitec
Differential-Revision: https://reviews.llvm.org/D49342
llvm-svn: 337489
|
|
|
Differential Revisison: https://reviews.llvm.org/D32101
llvm-svn: 301343
|
