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
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
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
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
ee0502d551
CUDA: give correct address space to globals declared in functions
...
llvm-svn: 162787
2012-08-28 20:37:10 +00:00
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
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
fa4d6033a3
CUDA: IR generation support for device stubs
...
llvm-svn: 141304
2011-10-06 18:51:56 +00:00
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
fe88342240
CUDA: IR generation support for kernel call expressions
...
llvm-svn: 141300
2011-10-06 18:29:37 +00:00
5bad4afa2f
CUDA: set proper calling conventions for PTX
...
llvm-svn: 141296
2011-10-06 16:49:54 +00:00
Matt Arsenault