Commit Graph

16 Commits

Author SHA1 Message Date
Eli Bendersky 3468d9d929 Move all CUDA testing inputs to Inputs/ subdirectory inside the tests.
llvm-svn: 207453
2014-04-28 22:21:28 +00:00
Eli Bendersky 8578c8f1e3 Add test case for r206302
llvm-svn: 206303
2014-04-15 16:57:53 +00:00
Eli Bendersky cb39943f6f Proper handling of static local variables with address space qualifiers.
Similar to the implementation for globals in r157167.

Patch by Jingyue Wu.

llvm-svn: 204677
2014-03-24 22:05:38 +00:00
Hans Wennborg c9bd88e681 Remove the -cxx-abi command-line flag.
This makes the C++ ABI depend entirely on the target: MS ABI for -win32 triples,
Itanium otherwise. It's no longer possible to do weird combinations.

To be able to run a test with a specific ABI without constraining it to a
specific triple, new substitutions are added to lit: %itanium_abi_triple and
%ms_abi_triple can be used to get the current target triple adjusted to the
desired ABI. For example, if the test suite is running with the i686-pc-win32
target, %itanium_abi_triple will expand to i686-pc-mingw32.

Differential Revision: http://llvm-reviews.chandlerc.com/D2545

llvm-svn: 199250
2014-01-14 19:35:09 +00:00
Hans Wennborg 9125b08b52 Update tests in preparation for using the MS ABI for Win32 targets
In preparation for making the Win32 triple imply MS ABI mode,
make all tests pass in this mode, or make them use the Itanium
mode explicitly.

Differential Revision: http://llvm-reviews.chandlerc.com/D2401

llvm-svn: 199130
2014-01-13 19:48:13 +00:00
Matt Arsenault 00e65b2afe Fix test failures after addrspacecast added.
Bitcasts between address spaces are no longer allowed.

llvm-svn: 194765
2013-11-15 02:19:52 +00:00
Stephen Lin 4362261b00 CHECK-LABEL-ify some code gen tests to improve diagnostic experience when tests fail.
llvm-svn: 188447
2013-08-15 06:47:53 +00:00
Justin Holewinski 368374308d Use kernel metadata to differentiate between kernel and device
functions for the NVPTX target.

llvm-svn: 178418
2013-03-30 14:38:24 +00:00
Peter Collingbourne c6b0857e95 CUDA: give static storage class to __shared__ and __constant__
variables without a storage class within a function, to implement
CUDA B.2.5: "__shared__ and __constant__ variables have implied static
storage [duration]."

llvm-svn: 162788
2012-08-28 20:37:50 +00:00
Peter Collingbourne ee0502d551 CUDA: give correct address space to globals declared in functions
llvm-svn: 162787
2012-08-28 20:37:10 +00:00
Justin Holewinski 83e9668133 Replace PTX back-end with NVPTX back-end in all places where Clang cares
NV_CONTRIB

llvm-svn: 157403
2012-05-24 17:43:12 +00:00
Peter Collingbourne f44bdf9c5f CUDA: add CodeGen support for global variable address spaces.
Because in CUDA types do not have associated address spaces,
globals are declared in their "native" address space, and accessed
by bitcasting the pointer to address space 0.  This relies on address
space 0 being a unified address space.

llvm-svn: 157167
2012-05-20 21:08:35 +00:00
Peter Collingbourne fa4d6033a3 CUDA: IR generation support for device stubs
llvm-svn: 141304
2011-10-06 18:51:56 +00:00
Peter Collingbourne a9455ec9f8 CUDA: add -fcuda-is-device flag
This frontend-only flag is used by the IR generator to determine
whether to filter CUDA declarations for the host or for the device.

llvm-svn: 141301
2011-10-06 18:29:46 +00:00
Peter Collingbourne fe88342240 CUDA: IR generation support for kernel call expressions
llvm-svn: 141300
2011-10-06 18:29:37 +00:00
Peter Collingbourne 5bad4afa2f CUDA: set proper calling conventions for PTX
llvm-svn: 141296
2011-10-06 16:49:54 +00:00