Commit Graph

40 Commits

Author SHA1 Message Date
Justin Lebar 2d56c26504 [CUDA] Use only the GVALinkage on function definitions.
Summary:
Previously we'd look at the GVALinkage of whatever FunctionDecl you
happened to be calling.

This is not right.  In the absence of the gnu_inline attribute, to be
handled separately, the function definition determines the function's
linkage.  So we need to wait until we get a def before we can know
whether something is known-emitted.

Reviewers: tra

Subscribers: cfe-commits, rsmith

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

llvm-svn: 286313
2016-11-08 23:45:51 +00:00
Justin Lebar 6f72737c3b [CUDA] Use FunctionDeclAndLoc for the Sema::LocsWithCUDACallDiags hashtable.
Summary: NFC

Reviewers: rnk

Subscribers: cfe-commits

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

llvm-svn: 284869
2016-10-21 20:08:52 +00:00
Artem Belevich 9fb40e3bbd Removed unused function argument. NFC.
Differential Revision: https://reviews.llvm.org/D25839

llvm-svn: 284843
2016-10-21 17:15:46 +00:00
Justin Lebar 6c86e9160d [CUDA] When we emit an error that might have been deferred, also print a callstack.
Summary:
Previously, when you did something not allowed in a host+device function
and then caused it to be codegen'ed, we would print out an error telling
you that you did something bad, but we wouldn't tell you how we decided
that the function needed to be codegen'ed.

This change causes us to print out a callstack when emitting deferred
errors.  This is immensely helpful when debugging highly-templated code,
where it's often unclear how a function became known-emitted.

We only print the callstack once per function, after we print the all
deferred errors.

This patch also switches all of our hashtables to using canonical
FunctionDecls instead of regular FunctionDecls.  This prevents a number
of bugs, some of which are caught by tests added here, in which we
assume that two FDs for the same function have the same pointer value.

Reviewers: rnk

Subscribers: cfe-commits, tra

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

llvm-svn: 284647
2016-10-19 21:15:01 +00:00
Justin Lebar 9730ae943f [CUDA] Emit errors for wrong-side calls made on the same line as non-wrong-side calls.
Summary:
This fixes two related bugs:

1) Previously, if you had a non-wrong side call at some source code
location L, we wouldn't emit errors for wrong-side calls that appeared
at L.

2) We'd only emit one wrong-side error per source code location, when we
actually want to emit it twice if we hit this line more than once due to
e.g. template instantiation.

Reviewers: tra

Subscribers: rnk, cfe-commits

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

llvm-svn: 284643
2016-10-19 21:03:38 +00:00
Justin Lebar d692dfb65e [CUDA] Fix false-positive in known-emitted handling.
Previously: When compiling for host, our constructed call graph went
*through* kernel calls.  This meant that if we had

  host calls kernel calls HD

we would incorrectly mark the HD function as known-emitted on the host
side, and thus perform host-side checks on it.

Fixing this exposed another issue, wherein when marking a function as
known-emitted, we also need to traverse the callgraph of its template,
because non-dependent calls are attached to a function's template, not
its instantiation.

llvm-svn: 284355
2016-10-17 02:25:55 +00:00
Justin Lebar 606f01f309 Add and use isDiscardableGVALinkage function.
Reviewers: rnk

Subscribers: cfe-commits

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

llvm-svn: 284159
2016-10-13 20:52:17 +00:00
Justin Lebar 23d954241b [CUDA] Emit deferred diagnostics during Sema rather than during codegen.
Summary:
Emitting deferred diagnostics during codegen was a hack.  It did work,
but usability was poor, both for us as compiler devs and for users.  We
don't codegen if there are any sema errors, so for users this meant that
they wouldn't see deferred errors if there were any non-deferred errors.
For devs, this meant that we had to carefully split up our tests so that
when we tested deferred errors, we didn't emit any non-deferred errors.

This change moves checking for deferred errors into Sema.  See the big
comment in SemaCUDA.cpp for an overview of the idea.

This checking adds overhead to compilation, because we have to maintain
a partial call graph.  As a result, this change makes deferred errors a
CUDA-only concept (whereas before they were a general concept).  If
anyone else wants to use this framework for something other than CUDA,
we can generalize at that time.

This patch makes the minimal set of test changes -- after this lands,
I'll go back through and do a cleanup of the tests that we no longer
have to split up.

Reviewers: rnk

Subscribers: cfe-commits, rsmith, tra

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

llvm-svn: 284158
2016-10-13 20:52:12 +00:00
Justin Lebar 179bdce72a [CUDA] Add Sema::CUDADiagBuilder and Sema::CUDADiagIf{Device,Host}Code().
Summary:
Together these let you easily create diagnostics that

 - are never emitted for host code
 - are always emitted for __device__ and __global__ functions, and
 - are emitted for __host__ __device__ functions iff these functions are
   codegen'ed.

At the moment there are only three diagnostics that need this treatment,
but I have more to add, and it's not sustainable to write code for emitting
every such diagnostic twice, and from a special wrapper in SemaCUDA.cpp.

While we're at it, don't emit the function name in
err_cuda_device_exceptions: It's not necessary to print it, and making
this work in the new framework in the face of a null value for
dyn_cast<FunctionDecl>(CurContext) isn't worth the effort.

Reviewers: rnk

Subscribers: cfe-commits, tra

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

llvm-svn: 284143
2016-10-13 18:45:08 +00:00
Justin Lebar 0254c46300 [CUDA] Make touching a kernel from a __host__ __device__ function a deferred error.
Previously, this was an immediate, don't pass go, don't collect $200
error.  But this precludes us from writing code like

  __host__ __device__ void launch_kernel() {
    kernel<<<...>>>();
  }

Such code isn't wrong, following our notions of right and wrong in CUDA,
unless it's codegen'ed.

llvm-svn: 283963
2016-10-12 01:30:08 +00:00
Richard Smith f75dcbef20 Aligned allocation versus CUDA: make deallocation function preference order
match other CUDA preference orders, per discussion with jlebar. We now model
this in an attempt to match overload resolution as closely as possible:

- First, we throw out all non-callable (due to CUDA host/device mismatch)
  operator delete functions.
- Then we apply sizedness / alignedness preferences based on whether the type
  is overaligned and whether the deallocation function is a member.
- Finally, we use the CUDA callability preference as a tiebreaker.

llvm-svn: 283830
2016-10-11 00:21:10 +00:00
Richard Smith b2f0f05742 Re-commit r283722, reverted in r283750, with a fix for a CUDA-specific use of
past-the-end iterator.

Original commit message:

P0035R4: Semantic analysis and code generation for C++17 overaligned
allocation.

llvm-svn: 283789
2016-10-10 18:54:32 +00:00
Daniel Jasper e9abe64816 Revert "P0035R4: Semantic analysis and code generation for C++17 overaligned allocation."
This reverts commit r283722. Breaks:
  Clang.SemaCUDA.device-var-init.cu
  Clang.CodeGenCUDA.device-var-init.cu

http://lab.llvm.org:8080/green/job/clang-stage1-cmake-RA-expensive/884/

llvm-svn: 283750
2016-10-10 14:13:55 +00:00
Richard Smith 189e52fcdf P0035R4: Semantic analysis and code generation for C++17 overaligned
allocation.

llvm-svn: 283722
2016-10-10 06:42:31 +00:00
Justin Lebar 67a78a6cc0 [CUDA] Add #pragma clang force_cuda_host_device_{begin,end} pragmas.
Summary:
These cause us to consider all functions in-between to be __host__
__device__.

You can nest these pragmas; you just can't have more 'end's than
'begin's.

Reviewers: rsmith

Subscribers: tra, jhen, cfe-commits

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

llvm-svn: 283677
2016-10-08 22:15:58 +00:00
Justin Lebar 9fdb46e71c [CUDA] Do a better job at detecting wrong-side calls.
Summary:
Move CheckCUDACall from ActOnCallExpr and BuildDeclRefExpr to
DiagnoseUseOfDecl.  This lets us catch some edge cases we were missing,
specifically around class operators.

This necessitates a few other changes:

 - Avoid emitting duplicate deferred diags in CheckCUDACall.

   Previously we'd carefully placed our call to CheckCUDACall such that
   it would only ever run once for a particular callsite.  But now this
   isn't the case.

 - Emit deferred diagnostics from a template
   specialization/instantiation's primary template, in addition to from
   the specialization/instantiation itself.  DiagnoseUseOfDecl ends up
   putting the deferred diagnostics on the template, rather than the
   specialization, so we need to check both.

Reviewers: rsmith

Subscribers: cfe-commits, tra

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

llvm-svn: 283637
2016-10-08 01:07:11 +00:00
Justin Lebar 9d4ed262ef [CUDA] Harmonize asserts in SemaCUDA, NFC.
llvm-svn: 282987
2016-09-30 23:57:38 +00:00
Justin Lebar 7a406e78b3 [CUDA] Remove incorrect comment in CUDASetLambdaAttrs.
I'd said that nvcc doesn't allow you to add __host__ or __device__
attributes on lambdas in all circumstances, but I believe this was user
error on my part.  I can't reproduce these warnings/errors if I pass
--expt-extended-lambda to nvcc.

llvm-svn: 282912
2016-09-30 19:55:59 +00:00
Justin Lebar 7ca116cacf [CUDA] Make lambdas inherit __host__ and __device__ attributes from the scope in which they're created.
Summary: NVCC compat.  Fixes bug 30567.

Reviewers: tra

Subscribers: cfe-commits, rnk

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

llvm-svn: 282880
2016-09-30 17:14:53 +00:00
Justin Lebar b17840de33 [CUDA] Disallow variable-length arrays in CUDA device code.
Reviewers: tra

Subscribers: cfe-commits, jhen

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

llvm-svn: 282647
2016-09-28 22:45:58 +00:00
Justin Lebar 2a8db34044 [CUDA] Disallow exceptions in device code.
Reviewers: tra

Subscribers: cfe-commits, jhen

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

llvm-svn: 282646
2016-09-28 22:45:54 +00:00
Justin Lebar 26bb31123a [CUDA] Fix "declared here" note on deferred wrong-side errors.
Previously we weren't deferring these "declared here" notes, which is
obviously wrong.

llvm-svn: 278767
2016-08-16 00:48:21 +00:00
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
Justin Lebar e71c08f84c [CUDA] Use the multi-element remove function in EraseUnwantedCUDAMatches.
Summary:
Bug pointed out by Benjamin Kramer in r264008.  I think the bug is
benign because by the time this is called, we should only have at most
two overloads to consider (either a host and a device overload, or a
host+device overload, but not all three).

Reviewers: tra

Subscribers: cfe-commits, bkramer

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

llvm-svn: 275233
2016-07-12 23:23:13 +00:00
Artem Belevich 3650bbeebc [CUDA] Do not allow non-empty destructors for global device-side variables.
According to Cuda Programming guide (v7.5, E2.3.1):
> __device__, __constant__ and __shared__ variables defined in namespace
> scope, that are of class type, cannot have a non-empty constructor or a
> non-empty destructor.

Clang already deals with device-side constructors (see D15305).
This patch enforces similar rules for destructors.

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

llvm-svn: 270108
2016-05-19 20:13:53 +00:00
Justin Lebar ba122ab42f [CUDA] Make unattributed constexpr functions implicitly host+device.
With this patch, by a constexpr function is implicitly host+device
unless:

 a) it's a variadic function (variadic functions are not allowed on the
    device side), or
 b) it's preceeded by a __device__ overload in a system header.

The restriction on overloading __host__ __device__ functions on the
basis of their CUDA attributes remains in place, but we use (b) to allow
us to define __device__ overloads for constexpr functions in cmath,
which would otherwise be __host__ __device__ and thus not overloadable.

You can disable this behavior with -fno-cuda-host-device-constexpr.

Reviewers: tra, rnk, rsmith

Subscribers: cfe-commits

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

llvm-svn: 264964
2016-03-30 23:30:21 +00:00
Justin Lebar 3918647a11 [CUDA] Fix order of overloading preferences in comment.
llvm-svn: 264741
2016-03-29 16:24:22 +00:00
Justin Lebar 25c4a81e79 [CUDA] Remove three obsolete CUDA cc1 flags.
Summary:
* -fcuda-target-overloads

  Previously unconditionally set to true by the driver.  Necessary for
  correct functioning of the compiler -- our CUDA headers wrapper won't
  compile without this.

* -fcuda-disable-target-call-checks

  Previously unconditionally set to true by the driver.  Necessary to
  compile almost any external CUDA code -- almost all libraries assume
  that host+device code can call host or device functions.

* -fcuda-allow-host-calls-from-host-device

  No effect when target overloading is enabled.

Reviewers: tra

Subscribers: rsmith, cfe-commits

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

llvm-svn: 264739
2016-03-29 16:24:16 +00:00
Justin Lebar e6a2cc12a3 [sema] [CUDA] Use std algorithms in EraseUnwantedCUDAMatchesImpl.
Summary: NFC

Reviewers: tra

Subscribers: cfe-commits

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

llvm-svn: 264008
2016-03-22 00:09:25 +00:00
Artem Belevich 186091094a [CUDA] Tweak attribute-based overload resolution to match nvcc behavior.
This is an artefact of split-mode CUDA compilation that we need to
mimic. HD functions are sometimes allowed to call H or D functions. Due
to split compilation mode device-side compilation will not see host-only
function and thus they will not be considered at all. For clang both H
and D variants will become function overloads visible to
compiler. Normally target attribute is considered only if C++ rules can
not determine which function is better. However in this case we need to
ignore functions that would not be present during current compilation
phase before we apply normal overload resolution rules.

Changes:
* introduced another level of call preference to better describe
  possible call combinations.
* removed WrongSide functions from consideration if the set contains
  SameSide function.
* disabled H->D, D->H and G->H calls. These combinations are
  not allowed by CUDA and we were reluctantly allowing them to work
  around device-side calls to math functions in std namespace.
  We no longer need it after r258880.

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

llvm-svn: 260697
2016-02-12 18:29:18 +00:00
Artem Belevich 97c01c35f8 [CUDA] Do not allow dynamic initialization of global device side variables.
In general CUDA does not allow dynamic initialization of
global device-side variables. One exception is that CUDA allows
records with empty constructors as described in section E2.2.1 of
CUDA 7.5 Programming guide.

This patch applies initializer checks for all device-side variables.
Empty constructors are accepted, but no code is generated for them.

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

llvm-svn: 259592
2016-02-02 22:29:48 +00:00
Justin Lebar c66a10652a [CUDA] Only allow __global__ on free functions and static member functions.
Summary:
Warn for NVCC compatibility if you declare a static member function or
inline function as __global__.

Reviewers: tra

Subscribers: jhen, echristo, cfe-commits

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

llvm-svn: 258263
2016-01-20 00:26:57 +00:00
Artem Belevich 94a55e8169 [CUDA] Allow function overloads in CUDA based on host/device attributes.
The patch makes it possible to parse CUDA files that contain host/device
functions with identical signatures, but different attributes without
having to physically split source into host-only and device-only parts.

This change is needed in order to parse CUDA header files that have
a lot of name clashes with standard include files.

Gory details are in design doc here: https://goo.gl/EXnymm
Feel free to leave comments there or in this review thread.

This feature is controlled with CC1 option -fcuda-target-overloads
and is disabled by default.

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

llvm-svn: 248295
2015-09-22 17:22:59 +00:00
Eli Bendersky 4bdc50eccb Create a frontend flag to disable CUDA cross-target call checks
For CUDA source, Sema checks that the targets of call expressions make sense
(e.g. a host function can't call a device function).

Adding a flag that lets us skip this check. Motivation: for source-to-source
translation tools that have to accept code that's not strictly kosher CUDA but
is still accepted by nvcc. The source-to-source translation tool can then fix
the code and leave calls that are semantically valid for the actual compilation
stage.

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

llvm-svn: 235049
2015-04-15 22:27:06 +00:00
Jacques Pienaar a50178c23e CUDA: Add option to allow host device functions to call host functions
Commiting code from review http://reviews.llvm.org/D7841

llvm-svn: 230385
2015-02-24 21:45:33 +00:00
Jacques Pienaar 5bdd67778f Consider calls from implict host device functions as valid in SemaCUDA.
In SemaCUDA all implicit functions were considered host device, this led to
errors such as the following code snippet failing to compile:

struct Copyable {
  const Copyable& operator=(const Copyable& x) { return *this; }
};

struct Simple {
  Copyable b;
};

void foo() {
  Simple a, b;

  a = b;
}

Above the implicit copy assignment operator was inferred as host device but
there was only a host assignment copy defined which is an error in device
compilation mode.

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

llvm-svn: 224358
2014-12-16 20:12:38 +00:00
Reid Kleckner bbc0178518 CUDA host device code with two code paths
Summary:
Allow CUDA host device functions with two code paths using __CUDA_ARCH__
to differentiate between code path being compiled.

For example:
  __host__ __device__ void host_device_function(void) {
  #ifdef __CUDA_ARCH__
    device_only_function();
  #else
    host_only_function();
  #endif
  }

Patch by Jacques Pienaar.

Reviewed By: rnk

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

llvm-svn: 223271
2014-12-03 21:53:36 +00:00
Eli Bendersky f2787a0dc0 CUDA: mark the target of implicit intrinsics properly
r218624 implemented target inference for implicit special members. However,
other entities can be implicit - for example intrinsics. These can not have
inference running on them, so they should be marked host device as before. This
is the safest and most flexible setting, since by construction these functions
don't invoke anything, and we'd like them to be invokable from both host and
device code. LLVM's intrinsics definitions (where these intrinsics come from in
the case of CUDA/NVPTX) have no notion of target, so both host and device
intrinsics can be supported this way.

llvm-svn: 218688
2014-09-30 17:38:34 +00:00
Eli Bendersky 9a220fca4a CUDA: Fix incorrect target inference for implicit members.
As PR20495 demonstrates, Clang currenlty infers the CUDA target (host/device,
etc) for implicit members (constructors, etc.) incorrectly. This causes errors
and even assertions in Clang when compiling code (assertions in C++11 mode where
implicit move constructors are added into the mix).

Fix the problem by inferring the target from the methods the implicit member
should call (depending on its base classes and fields).

llvm-svn: 218624
2014-09-29 20:38:29 +00:00
Eli Bendersky 7325e56c10 Split off CUDA-specific Sema parts to a new file
In line with SemaOpenMP.cpp, etc. CUDA-specific semantic analysis code goes into
a separate file. This is in anticipation of adding extra functionality here in
the near future.

No change in functionality.

Review: http://reviews.llvm.org/D5160
llvm-svn: 217043
2014-09-03 15:27:03 +00:00