Tim Northover 06b700153b SelectionDAG: accommodate atomic floating stores. ...

We were applying a pointer truncation to floating types, which crashed LLVM.
That is Not A Good Thing(TM).

git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@360421 91177308-0d34-0410-b5e6-96231b3b80d8

2019-05-10 11:23:04 +00:00

..

access-non-generic.ll

[NVPTX] Use addrspacecast instead of target-specific intrinsics in NVPTXGenericToNVVM.

2018-02-28 23:57:48 +00:00

add-128bit.ll

[DAGCombiner] add missing folds for scalar select of {-1,0,1}

2017-02-24 17:17:33 +00:00

addrspacecast-gvar.ll

[NVPTX] Handle addrspacecast constant expressions in aggregate initializers

2015-04-28 17:18:30 +00:00

addrspacecast.ll

[NVPTX] Added a feature to use short pointers for const/local/shared AS.

2018-05-09 23:46:19 +00:00

aggr-param.ll

[NVPTX] Fix emitting aggregate parameters

2014-01-28 18:35:29 +00:00

aggregate-return.ll

[NVPTX] Unify vectorization of load/stores of aggregate arguments and return values.

2017-02-21 22:56:05 +00:00

alias.ll

[CUDA] Die gracefully when trying to output an LLVM alias.

2016-01-23 21:12:20 +00:00

annotations.ll

Whitespace cleanup in test/CodeGen/NVPTX/annotations.ll.

2016-12-14 22:32:55 +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] expand mul_lohi to mul_lo and mul_hi

2016-01-22 19:47:26 +00:00

atomics-sm60.ll

[NVPTX] Implement __nvvm_atom_add_gen_d builtin.

2017-11-07 22:10:54 +00:00

atomics-with-scope.ll

[NVPTX] Added intrinsics for atom.gen.{sys|cta}.* instructions.

2016-09-28 17:25:38 +00:00

atomics.ll

Add some tests for NVPTX lowering of cmpxchg

2014-07-21 22:54:44 +00:00

barrier.ll

[NVPTX] Implemented bar.warp.sync, barrier.sync, and vote{.sync} instructions/intrinsics/builtins.

2017-09-21 18:44:49 +00:00

bfe.ll

[NVPTX] Add isel patterns for bit-field extract (bfe)

2014-06-27 18:35:27 +00:00

branch-fold.ll

Roll forward r242871

2015-07-29 18:59:09 +00:00

bug17709.ll

[NVPTX] Don't flag StoreParam/LoadParam memory chain operands as ReadMem/WriteMem (PR32146)

2017-05-15 17:17:44 +00:00

bug21465.ll

[NVPTX] Renamed NVPTXLowerKernelArgs -> NVPTXLowerArgs. NFC.

2016-07-20 21:44:07 +00:00

bug22246.ll

[NVPTX] Generate a more optimal sequence for select of i1

2015-01-26 19:52:20 +00:00

bug22322.ll

DAG: Check no-signed-zeros instead of unsafe-fp-math

2018-08-12 19:09:12 +00:00

bug26185-2.ll

[NVPTX] Fix sign/zero-extending ldg/ldu instruction selection

2016-05-02 18:12:02 +00:00

bug26185.ll

[NVPTX] Handle ldg created from sign-/zero-extended load

2016-04-05 12:38:01 +00:00

bypass-div.ll

Use 32-bit divides instead of 64-bit divides where possible.

2015-08-11 22:16:34 +00:00

call-with-alloca-buffer.ll

Fix NVPTX/call-with-alloca-buffer.ll after r276777.

2016-07-26 18:28:33 +00:00

callchain.ll

…

calling-conv.ll

IR: Make metadata typeless in assembly

2014-12-15 19:07:53 +00:00

calls-with-phi.ll

[NVPTX] Allow libcalls that are defined in the current module.

2018-12-26 19:12:31 +00:00

combine-min-max.ll

[NVPTX] Implement min/max in tablegen, rather than with custom DAGComine logic.

2017-01-18 00:09:01 +00:00

compare-int.ll

[NVPTX] Rename registers %fl -> %fd and %rl -> %rd

2014-07-16 16:26:58 +00:00

constant-vectors.ll

…

convergent-mir-call.ll

[NVPTX] Use different, convergent MIs for convergent calls.

2016-03-01 19:24:03 +00:00

convert-fp.ll

[NVPTX] Add fptosi tests to convert-fp.ll.

2017-01-15 16:55:54 +00:00

convert-int-sm20.ll

[NVPTX] Rename registers %fl -> %fd and %rl -> %rd

2014-07-16 16:26:58 +00:00

ctlz.ll

NFC - Various typo fixes in tests

2018-07-04 13:28:39 +00:00

ctpop.ll

[NVPTX] Don't flag StoreRetVal memory chain operands as ReadMem (PR32146)

2017-05-12 19:56:43 +00:00

cttz.ll

[NVPTX] Don't flag StoreRetVal memory chain operands as ReadMem (PR32146)

2017-05-12 19:56:43 +00:00

disable-opt.ll

[NVPTX] Disable performance optimizations when OptLevel==None

2016-02-04 04:15:36 +00:00

div-ri.ll

[NVPTX] Add missing patterns for div.approx with immediate denominator

2014-01-21 14:40:05 +00:00

divrem-combine.ll

[NVPTX] Compute 'rem' using the result of 'div', if possible.

2016-10-28 21:44:00 +00:00

envreg.ll

[NVPTX] Add support for envreg reads

2014-06-27 18:35:21 +00:00

extloadv.ll

[NVPTX] expand extload/truncstore for vectors of floats

2015-07-01 21:32:42 +00:00

f16-instructions.ll

[NVPTX] Fix the codegen for llvm.round.

2019-04-01 16:10:26 +00:00

f16x2-instructions.ll

[NVPTX] Fix the codegen for llvm.round.

2019-04-01 16:10:26 +00:00

fast-math.ll

[NVPTX] Enable combineRepeatedFPDivisors for NVPTX.

2017-02-03 15:13:50 +00:00

fcos-no-fast-math.ll

[NVPTX] Only lower sin/cos to approximate instructions if unsafe math is allowed.

2017-01-13 18:48:13 +00:00

fma-assoc.ll

[DAGCombine] require UnsafeFPMath for re-association of addition

2017-01-31 14:35:37 +00:00

fma-disable.ll

…

fma.ll

[NVPTX] Don't flag StoreParam/LoadParam memory chain operands as ReadMem/WriteMem (PR32146)

2017-05-15 17:17:44 +00:00

fns.ll

[NVPTX,CUDA] Added llvm.nvvm.fns intrinsic and matching __nvvm_fns builtin in clang.

2017-12-06 17:50:05 +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

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

fsin-no-fast-math.ll

[NVPTX] Only lower sin/cos to approximate instructions if unsafe math is allowed.

2017-01-13 18:48:13 +00:00

function-align.ll

[NVPTXAsmPrinter] do not print .align on function headers

2015-03-12 01:50:30 +00:00

generic-to-nvvm-ir.ll

[DebugInfo] Add DILabel metadata and intrinsic llvm.dbg.label.

2018-05-09 02:40:45 +00:00

generic-to-nvvm.ll

[NVPTX] Use addrspacecast instead of target-specific intrinsics in NVPTXGenericToNVVM.

2018-02-28 23:57:48 +00:00

global-addrspace.ll

[NVPTX] Allow undef value as global initializer

2015-08-22 05:40:26 +00:00

global-ctor-empty.ll

[CUDA] Die if we ask the NVPTX backend to emit a global ctor/dtor.

2016-01-30 01:07:38 +00:00

global-ctor.ll

[CUDA] Die if we ask the NVPTX backend to emit a global ctor/dtor.

2016-01-30 01:07:38 +00:00

global-dtor.ll

[CUDA] Die if we ask the NVPTX backend to emit a global ctor/dtor.

2016-01-30 01:07:38 +00:00

global-ordering.ll

…

global-variable-big.ll

[NVPTX] Support global variables of integer type larger than i64.

2017-01-18 00:29:53 +00:00

global-visibility.ll

[NVPTX] Do not emit .hidden or .protected directives as they are not allowed by PTX.

2016-01-15 23:57:53 +00:00

globals_init.ll

The constant initialization for globals in NVPTX is generated as an

2015-06-09 16:29:34 +00:00

globals_lowering.ll

Force relocation mode to be default, regardless of what is passed to the backend.

2015-06-30 17:18:00 +00:00

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

[NVPTX] add support for initializing fp16 arrays.

2018-04-06 22:25:08 +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

[NVPTX] Don't flag StoreParam/LoadParam memory chain operands as ReadMem/WriteMem (PR32146)

2017-05-15 17:17:44 +00:00

i128-global.ll

[NVPTX] Add lowering of i128 params.

2017-07-20 21:16:03 +00:00

i128-param.ll

[NVPTX] Add lowering of i128 params.

2017-07-20 21:16:03 +00:00

i128-retval.ll

[NVPTX] Add lowering of i128 params.

2017-07-20 21:16:03 +00:00

i128-struct.ll

[NVPTX] Add lowering of i128 numbers as struct fields

2018-12-01 00:21:52 +00:00

idioms.ll

[NVPTX] Lower integer absolute value idiom to abs instruction.

2017-01-18 00:08:44 +00:00

imad.ll

[NVPTX] Implement fma and imad contraction as target DAGCombiner patterns

2014-06-27 18:35:37 +00:00

inline-asm.ll

[NVPTX] Add 'b' asm constraint

2014-06-27 18:36:06 +00:00

inlineasm-output-template.ll

[AsmPrinter] refactor to support %c w/ GlobalAddress'

2019-04-26 18:45:04 +00:00

intrin-nocapture.ll

Reapply 239795 - [InstCombine] Propagate non-null facts to call parameters

2015-06-16 20:24:25 +00:00

intrinsic-old.ll

[FileCheck] Add -allow-deprecated-dag-overlap to failing llvm tests

2018-07-11 20:25:49 +00:00

intrinsics.ll

[NVPTX] Some nvvm.read.ptx.sreg intrinsics should have IntrInaccessibleMemOnly attribute.

2019-01-26 00:28:32 +00:00

isspacep.ll

[NVPTX] Add support for isspacep instruction

2014-06-27 18:35:24 +00:00

ld-addrspace.ll

[NVPTX] Added a feature to use short pointers for const/local/shared AS.

2018-05-09 23:46:19 +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

ld-st-addrrspace.py

Python compat - print statement

2019-01-03 14:11:33 +00:00

ldg-invariant.ll

[NVPTX] Fixed vectorized LDG for f16.

2018-04-06 21:10:24 +00:00

ldparam-v4.ll

[NVPTX] Unify vectorization of load/stores of aggregate arguments and return values.

2017-02-21 22:56:05 +00:00

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

libcall-fulfilled.ll

[NVPTX] Allow libcalls that are defined in the current module.

2018-12-26 19:12:31 +00:00

libcall-instruction.ll

[NVPTX] Allow libcalls that are defined in the current module.

2018-12-26 19:12:31 +00:00

libcall-intrinsic.ll

[NVPTX] Allow libcalls that are defined in the current module.

2018-12-26 19:12:31 +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

load-store.ll

SelectionDAG: accommodate atomic floating stores.

2019-05-10 11:23:04 +00:00

load-with-non-coherent-cache.ll

[NVPTX] Use LDG for pointer induction variables.

2015-08-05 23:11:57 +00:00

LoadStoreVectorizer.ll

[NVPTX] Added support for .f16x2 instructions.

2017-02-23 22:38:24 +00:00

local-stack-frame.ll

[NVPTX] Move NVPTXPeephole after NVPTXPrologEpilogPass

2015-07-01 20:08:06 +00:00

loop-vectorize.ll

[NVPTX] declare no vector registers

2015-07-10 04:31:56 +00:00

lower-aggr-copies.ll

Remove alignment argument from memcpy/memmove/memset in favour of alignment attributes (Step 1)

2018-01-19 17:13:12 +00:00

lower-alloca.ll

NVPTX: Move InferAddressSpaces to generic code

2017-01-31 01:10:58 +00:00

lower-kernel-ptr-arg.ll

[NVPTX] Improve lowering of byval args of device functions.

2016-07-20 18:39:47 +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

MachineSink-call.ll

[NVPTX] Annotate call machine instructions as calls.

2016-02-17 17:46:50 +00:00

MachineSink-convergent.ll

NVPTX: Replace uses of cuda.syncthreads with nvvm.barrier0

2016-07-06 20:02:45 +00:00

managed.ll

IR: Make metadata typeless in assembly

2014-12-15 19:07:53 +00:00

match.ll

[NVPTX] added match.{any,all}.sync instructions, intrinsics & builtins.

2017-09-26 17:07:23 +00:00

math-intrins.ll

[NVPTX] Fix the codegen for llvm.round.

2019-04-01 16:10:26 +00:00

minmax-negative.ll

Improve clamp recognition in ValueTracking.

2017-10-27 20:53:41 +00:00

misaligned-vector-ldst.ll

[NVPTX] Fixed lowering of unaligned loads/stores of f16 scalars and vectors.

2017-03-07 20:33:38 +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

named-barriers.ll

[NVPTX] Add intrinsics to support named barriers.

2017-01-28 16:38:15 +00:00

noduplicate-syncthreads.ll

NVPTX: Replace uses of cuda.syncthreads with nvvm.barrier0

2016-07-06 20:02:45 +00:00

nofunc.ll

[NVPTX] do not rely on cached subtarget info.

2018-12-12 18:31:04 +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-arch.ll

[NVPTX] Handle __nvvm_reflect("__CUDA_ARCH").

2018-08-03 18:05:24 +00:00

nvvm-reflect-module-flag.ll

[NVPTX] Read __CUDA_FTZ from module flags in NVVMReflect.

2016-04-01 01:09:07 +00:00

nvvm-reflect.ll

[NVPTX] Let there be One True Way to set NVVMReflect params.

2017-01-15 16:54:35 +00:00

param-align.ll

[NVPTX] Make sure we adjust alignment at all call sites

2016-07-18 21:58:48 +00:00

param-load-store.ll

[Codegen] fix typos in test case

2019-03-02 08:03:59 +00:00

pr13291-i1-store.ll

[NVPTX] roll forward r239082

2015-06-04 21:28:26 +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

proxy-reg-erasure-mir.ll

[NVPTX] Allow libcalls that are defined in the current module.

2018-12-26 19:12:31 +00:00

proxy-reg-erasure-ptx.ll

[NVPTX] Allow libcalls that are defined in the current module.

2018-12-26 19:12:31 +00:00

read-global-variable-constant.ll

[NVPTX] Lower loads from global constants using ld.global.nc (aka LDG).

2018-02-28 23:58:05 +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

reg-copy.ll

[NVPTX] allow register copy between float and int

2015-08-01 18:02:12 +00:00

reg-types.ll

[NVPTX] Use untyped (.b) integer registers in PTX.

2016-08-12 22:02:19 +00:00

rotate.ll

[NVPTX] Add support for efficient rotate instructions on SM 3.2+

2014-06-27 18:35:33 +00:00

sched1.ll

Only enable LiveRangeShrink for x86.

2017-05-17 20:18:13 +00:00

sched2.ll

Only enable LiveRangeShrink for x86.

2017-05-17 20:18:13 +00:00

sext-in-reg.ll

…

sext-params.ll

…

shfl-sync.ll

[NVPTX] Implemented shfl.sync instruction and supporting intrinsics/builtins.

2017-09-20 21:23:07 +00:00

shfl.ll

[NVPTX] Remove NVPTXFavorNonGenericAddrSpaces pass.

2016-10-31 21:51:42 +00:00

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

[NVPTX] Don't flag StoreParam/LoadParam memory chain operands as ReadMem/WriteMem (PR32146)

2017-05-15 17:17:44 +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

sm-version-60.ll

[NVPTX] Add sm_60, sm_61, sm_62 targets to LLVM.

2016-07-06 21:06:10 +00:00

sm-version-61.ll

[NVPTX] Add sm_60, sm_61, sm_62 targets to LLVM.

2016-07-06 21:06:10 +00:00

sm-version-62.ll

[NVPTX] Add sm_60, sm_61, sm_62 targets to LLVM.

2016-07-06 21:06:10 +00:00

sm-version-70.ll

[CUDA] Added rudimentary support for CUDA-9 and sm_70.

2017-09-07 18:14:32 +00:00

speculative-execution-divergent-target.ll

Move divergent-target test into CodeGen/NVPTX because it requires an NVPTX target.

2016-04-15 01:20:52 +00:00

sqrt-approx.ll

[NVPTX] Compute approx sqrt as 1/rsqrt(x) rather than x*rsqrt(x).

2017-01-31 23:08:57 +00:00

st-addrspace.ll

[NVPTX] Added a feature to use short pointers for const/local/shared AS.

2018-05-09 23:46:19 +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

[NVPTX] roll forward r239082

2015-06-04 21:28:26 +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

[NVPTX] Assign valid global names

2017-12-04 14:19:33 +00:00

TailDuplication-convergent.ll

NVPTX: Replace uses of cuda.syncthreads with nvvm.barrier0

2016-07-06 20:02:45 +00:00

tex-read-cuda.ll

[NVPTX] roll forward r239082

2015-06-04 21:28:26 +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

tid-range.ll

[SelectionDAG] Correctly transform range metadata to AssertZExt

2017-01-06 00:11:46 +00:00

tuple-literal.ll

…

vec8.ll

Revert r302938 "Add LiveRangeShrink pass to shrink live range within BB."

2017-05-18 18:50:05 +00:00

vec-param-load.ll

[NVPTX] Unify vectorization of load/stores of aggregate arguments and return values.

2017-02-21 22:56:05 +00:00

vector-args.ll

…

vector-call.ll

[NVPTX] Don't flag StoreParam/LoadParam memory chain operands as ReadMem/WriteMem (PR32146)

2017-05-15 17:17:44 +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-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

…

vectorize-misaligned.ll

[NVPTX] Implement isLegalToVectorizeLoadChain

2018-08-27 17:29:43 +00:00

vote.ll

[NVPTX] Implemented bar.warp.sync, barrier.sync, and vote{.sync} instructions/intrinsics/builtins.

2017-09-21 18:44:49 +00:00

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

wmma.py

PTX 6.3 extends wmma instruction to support s8/u8/s4/u4/b1 -> s32.

2019-04-25 22:27:57 +00:00

zeroext-32bit.ll

[NVPTX] Don't flag StoreParam/LoadParam memory chain operands as ReadMem/WriteMem (PR32146)

2017-05-15 17:17:44 +00:00