llvm-project/clang/test/CodeGenCUDA
Justin Lebar 18e2d82297 [CUDA] Raise an error if a wrong-side call is codegen'ed.
Summary:
Some function calls in CUDA are allowed to appear in
semantically-correct programs but are an error if they're ever
codegen'ed.  Specifically, a host+device function may call a host
function, but it's an error if such a function is ever codegen'ed in
device mode (and vice versa).

Previously, clang made no attempt to catch these errors.  For the most
part, they would be caught by ptxas, and reported as "call to unknown
function 'foo'".

Now we catch these errors and report them the same as we report other
illegal calls (e.g. a call from a host function to a device function).

This has a small change in error-message behavior for calls that were
previously disallowed (e.g. calls from a host to a device function).
Previously, we'd catch disallowed calls fairly early, before doing
additional semantic checking e.g. of the call's arguments.  Now we catch
these illegal calls at the very end of our semantic checks, so we'll
only emit a "illegal CUDA call" error if the call is otherwise
well-formed.

Reviewers: tra, rnk

Subscribers: cfe-commits

Differential Revision: https://reviews.llvm.org/D23242

llvm-svn: 278759
2016-08-15 23:00:49 +00:00
..
Inputs Avoid depending on test inputes that aren't in Inputs 2016-05-20 00:38:25 +00:00
address-spaces.cu [CUDA] Restrict init of local __shared__ variables to empty constructors only. 2016-05-09 22:09:56 +00:00
alias.cu [CUDA] Don't generate aliases for static extern "C" functions. 2016-01-25 22:36:37 +00:00
convergent.cu [CUDA] Conservatively mark inline asm as convergent. 2016-05-31 21:27:13 +00:00
cuda-builtin-vars.cu NVPTX: Use the nvvm builtins to read SRegs rather than the legacy ptx ones 2016-07-07 16:41:08 +00:00
device-stub.cu [CUDA] Place GPU binary into .nv_fatbin section and align it by 8. 2016-08-12 18:44:01 +00:00
device-var-init.cu Avoid depending on test inputes that aren't in Inputs 2016-05-20 00:38:25 +00:00
device-vtable.cu [CUDA] Make vtable construction aware of host/device side of CUDA compilation. 2015-12-17 18:12:36 +00:00
filter-decl.cu [CUDA] Emit host-side 'shadows' for device-side global variables 2016-03-02 18:28:50 +00:00
flush-denormals.cu [CUDA] Fix flush-denormals.cu test so that it checks what it intends to CHECK. 2016-05-10 00:34:50 +00:00
fp-contract.cu [CUDA] Enable fusing FP ops (-ffp-contract=fast) for CUDA by default. 2016-05-19 18:44:45 +00:00
function-overload.cu [CUDA] Remove three obsolete CUDA cc1 flags. 2016-03-29 16:24:16 +00:00
kernel-args-alignment.cu [CUDA] Align kernel launch args correctly when the LLVM type's alignment is different from the clang type's alignment. 2016-07-27 22:36:21 +00:00
kernel-call.cu
launch-bounds.cu [CUDA] Add implicit conversion of __launch_bounds__ arguments to rvalue. 2016-06-06 22:54:57 +00:00
link-device-bitcode.cu [CUDA] Add -disable-llvm-passes to CodeGenCUDA/link-device-bitcode.cu. NFC 2016-03-30 23:45:38 +00:00
llvm-used.cu When generating llvm.used, we may need an addrspacecast instead of a bitcast. 2015-02-02 21:05:49 +00:00
printf-aggregate.cu [CUDA] Don't crash when trying to printf a non-scalar object. 2016-02-11 02:00:52 +00:00
printf.cu [CUDA] Generate CUDA's printf alloca in its function's entry block. 2016-01-28 23:58:28 +00:00
ptx-kernels.cu [CUDA] Give templated device functions internal linkage, templated kernels external linkage. 2016-06-30 18:41:33 +00:00