maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/llvm/test/CodeGen/NVPTX
History
Jingyue Wu 995dde2799 [NVPTXFavorNonGenericAddrSpaces] recursively trace into GEP and BitCast 
Summary:
This patch allows NVPTXFavorNonGenericAddrSpaces to remove addrspacecast
from longer chains consisting of GEPs and BitCasts. For example, it can
now optimize

  %0 = addrspacecast [10 x float] addrspace(3)* @a to [10 x float]*
  %1 = gep [10 x float]* %0, i64 0, i64 %i
  %2 = bitcast float* %1 to i32*
  %3 = load i32* %2 ; emits ld.u32

to

  %0 = gep [10 x float] addrspace(3)* @a, i64 0, i64 %i
  %1 = bitcast float addrspace(3)* %0 to i32 addrspace(3)*
  %3 = load i32 addrspace(3)* %1 ; emits ld.shared.f32

Test Plan: @ld_int_from_global_float in access-non-generic.ll

Reviewers: broune, eliben, jholewinski, meheff

Subscribers: jholewinski, llvm-commits

Differential Revision: http://reviews.llvm.org/D10074

llvm-svn: 238574
 		2015-05-29 17:00:27 +00:00
..
access-non-generic.ll [NVPTXFavorNonGenericAddrSpaces] recursively trace into GEP and BitCast 2015-05-29 17:00:27 +00:00
add-128bit.ll Revert revisions r234755, r234759, r234760 2015-04-13 17:47:15 +00:00
addrspacecast-gvar.ll [NVPTX] Handle addrspacecast constant expressions in aggregate initializers 2015-04-28 17:18:30 +00:00
addrspacecast.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
aggr-param.ll [NVPTX] Fix emitting aggregate parameters 2014-01-28 18:35:29 +00:00
annotations.ll IR: Make metadata typeless in assembly 2014-12-15 19:07:53 +00:00
arg-lowering.ll [NVPTX] Clean up argument lowering code and properly handle alignment for structs and vectors 2014-06-27 18:35:44 +00:00
arithmetic-fp-sm20.ll [NVPTX] Improve handling of FP fusion 2014-07-17 18:10:09 +00:00
arithmetic-int.ll [NVPTX] Rename registers %fl -> %fd and %rl -> %rd 2014-07-16 16:26:58 +00:00
atomics.ll Add some tests for NVPTX lowering of cmpxchg 2014-07-21 22:54:44 +00:00
bfe.ll [NVPTX] Add isel patterns for bit-field extract (bfe) 2014-06-27 18:35:27 +00:00
bug17709.ll Remove the linker_private and linker_private_weak linkages. 2014-03-13 23:18:37 +00:00
bug21465.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
bug22246.ll [NVPTX] Generate a more optimal sequence for select of i1 2015-01-26 19:52:20 +00:00
bug22322.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
call-with-alloca-buffer.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
callchain.ll
…
calling-conv.ll IR: Make metadata typeless in assembly 2014-12-15 19:07:53 +00:00
compare-int.ll [NVPTX] Rename registers %fl -> %fd and %rl -> %rd 2014-07-16 16:26:58 +00:00
constant-vectors.ll
…
convert-fp.ll [NVPTX] Rename registers %fl -> %fd and %rl -> %rd 2014-07-16 16:26:58 +00:00
convert-int-sm20.ll [NVPTX] Rename registers %fl -> %fd and %rl -> %rd 2014-07-16 16:26:58 +00:00
ctlz.ll
…
ctpop.ll
…
cttz.ll
…
div-ri.ll
…
envreg.ll [NVPTX] Add support for envreg reads 2014-06-27 18:35:21 +00:00
fast-math.ll
…
fma-assoc.ll Check that the TLI callback enableAggressiveFMAFusion has the desired effect on FMA folding. 2015-01-14 15:36:28 +00:00
fma-disable.ll
…
fma.ll Check that the TLI callback enableAggressiveFMAFusion has the desired effect on FMA folding. 2015-01-14 15:36:28 +00:00
fp-contract.ll [NVPTX] Improve handling of FP fusion 2014-07-17 18:10:09 +00:00
fp-literals.ll [NVPTX] Improve handling of FP fusion 2014-07-17 18:10:09 +00:00
fp16.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
function-align.ll [NVPTXAsmPrinter] do not print .align on function headers 2015-03-12 01:50:30 +00:00
generic-to-nvvm.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
global-ordering.ll
…
gvar-init.ll [NVPTX] Error out if initializer is given for variable in an address space that does not support initialization 2014-06-27 18:36:01 +00:00
half.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
i1-global.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
i1-int-to-fp.ll
…
i1-param.ll IR: Make metadata typeless in assembly 2014-12-15 19:07:53 +00:00
i8-param.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
imad.ll [NVPTX] Implement fma and imad contraction as target DAGCombiner patterns 2014-06-27 18:35:37 +00:00
implicit-def.ll [NVPTX] Improve handling of FP fusion 2014-07-17 18:10:09 +00:00
inline-asm.ll [NVPTX] Add 'b' asm constraint 2014-06-27 18:36:06 +00:00
intrin-nocapture.ll
…
intrinsic-old.ll [NVPTX] Rename registers %fl -> %fd and %rl -> %rd 2014-07-16 16:26:58 +00:00
intrinsics.ll [NVPTX] Rename registers %fl -> %fd and %rl -> %rd 2014-07-16 16:26:58 +00:00
isspacep.ll [NVPTX] Add support for isspacep instruction 2014-06-27 18:35:24 +00:00
ld-addrspace.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
ld-generic.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
ldparam-v4.ll
…
ldu-i8.ll [NVPTX] Make the alignment an explicit argument to ldu/ldg 2014-08-29 15:30:20 +00:00
ldu-ldg.ll [NVPTX] Make the alignment an explicit argument to ldu/ldg 2014-08-29 15:30:20 +00:00
ldu-reg-plus-offset.ll [opaque pointer type] Add textual IR support for explicit type parameter to getelementptr instruction 2015-02-27 19:29:02 +00:00
lit.local.cfg Reduce verbiage of lit.local.cfg files 2014-06-09 22:42:55 +00:00
load-sext-i1.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
local-stack-frame.ll [NVPTX] Rename registers %fl -> %fd and %rl -> %rd 2014-07-16 16:26:58 +00:00
machine-sink.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
managed.ll IR: Make metadata typeless in assembly 2014-12-15 19:07:53 +00:00
misaligned-vector-ldst.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
module-inline-asm.ll
…
mulwide.ll [NVPTX] Add some extra tests for mul.wide to test non-power-of-two source types 2014-07-23 20:23:49 +00:00
noduplicate-syncthreads.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
nounroll.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
nvcl-param-align.ll [NVPTX] Fix bugs related to isSingleValueType 2014-12-17 17:59:04 +00:00
nvvm-reflect.ll Add support for __nvvm_reflect changes in libdevice in CUDA-7.0 2015-03-19 17:05:35 +00:00
param-align.ll
…
pr13291-i1-store.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
pr16278.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
pr17529.ll [opaque pointer type] Add textual IR support for explicit type parameter to getelementptr instruction 2015-02-27 19:29:02 +00:00
refl1.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
rotate.ll [NVPTX] Add support for efficient rotate instructions on SM 3.2+ 2014-06-27 18:35:33 +00:00
rsqrt.ll
…
sched1.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
sched2.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
sext-in-reg.ll
…
sext-params.ll
…
shift-parts.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
simple-call.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
sm-version-20.ll
…
sm-version-21.ll
…
sm-version-30.ll [NVPTX] Associate a minimum PTX version for each SM architecture 2015-03-30 19:30:55 +00:00
sm-version-32.ll [NVPTX] Associate a minimum PTX version for each SM architecture 2015-03-30 19:30:55 +00:00
sm-version-35.ll [NVPTX] Associate a minimum PTX version for each SM architecture 2015-03-30 19:30:55 +00:00
sm-version-37.ll [NVPTX] Associate a minimum PTX version for each SM architecture 2015-03-30 19:30:55 +00:00
sm-version-50.ll [NVPTX] Associate a minimum PTX version for each SM architecture 2015-03-30 19:30:55 +00:00
sm-version-52.ll [NVPTX] Associate a minimum PTX version for each SM architecture 2015-03-30 19:30:55 +00:00
sm-version-53.ll [NVPTX] Associate a minimum PTX version for each SM architecture 2015-03-30 19:30:55 +00:00
st-addrspace.ll [NVPTX] Rename registers %fl -> %fd and %rl -> %rd 2014-07-16 16:26:58 +00:00
st-generic.ll [NVPTX] Rename registers %fl -> %fd and %rl -> %rd 2014-07-16 16:26:58 +00:00
surf-read-cuda.ll IR: Make metadata typeless in assembly 2014-12-15 19:07:53 +00:00
surf-read.ll IR: Make metadata typeless in assembly 2014-12-15 19:07:53 +00:00
surf-write-cuda.ll IR: Make metadata typeless in assembly 2014-12-15 19:07:53 +00:00
surf-write.ll IR: Make metadata typeless in assembly 2014-12-15 19:07:53 +00:00
symbol-naming.ll [opaque pointer type] Add textual IR support for explicit type parameter to the call instruction 2015-04-16 23:24:18 +00:00
tex-read-cuda.ll IR: Make metadata typeless in assembly 2014-12-15 19:07:53 +00:00
tex-read.ll IR: Make metadata typeless in assembly 2014-12-15 19:07:53 +00:00
texsurf-queries.ll IR: Make metadata typeless in assembly 2014-12-15 19:07:53 +00:00
tuple-literal.ll
…
vec-param-load.ll
…
vec8.ll
…
vector-args.ll
…
vector-call.ll [NVPTX] Add missing .v4 qualifier on vector store instruction 2014-07-17 16:58:56 +00:00
vector-compare.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
vector-global.ll [NVPTX] Fix bugs related to isSingleValueType 2014-12-17 17:59:04 +00:00
vector-loads.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
vector-return.ll [NVPTX] aligned byte-buffers for vector return types 2014-10-25 03:46:16 +00:00
vector-select.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
vector-stores.ll
…
weak-global.ll [opaque pointer type] Add textual IR support for explicit type parameter to load instruction 2015-02-27 21:17:42 +00:00
weak-linkage.ll [NVPTX] Do not emit .weak symbols for NVPTX 2014-12-01 21:16:17 +00:00