If a kernel template has a function as its template parameter, a device function should be
allowed as template argument since a kernel can call a device function. However,
currently if the kernel template is instantiated in a host function, clang will emit an error
message saying the device function is an invalid candidate for the template parameter.
This happens because clang checks the reference to the device function during parsing
the template arguments. At this point, the template is not instantiated yet. Clang incorrectly
assumes the device function is called by the host function and emits the error message.
This patch fixes the issue by disabling checking of device function during parsing template
arguments and deferring the check to the instantion of the template. At that point, the
template decl is already available, therefore the check can be done against the instantiated
function template decl.
Differential Revision: https://reviews.llvm.org/D56411
llvm-svn: 355421
Summary:
Deferred diagnostic interface is going to be used for OpenMP device
compilation. Generalized previously existed deferred diagnostic
interface for CUDA to be used with OpenMP and, possibly, other models.
Reviewers: rjmccall, tra
Subscribers: caomhin, cfe-commits, kkwli0
Tags: #clang
Differential Revision: https://reviews.llvm.org/D57908
llvm-svn: 353456
Instead of calling CUDA runtime to arrange function arguments,
the new API constructs arguments in a local array and the kernels
are launched with __cudaLaunchKernel().
The old API has been deprecated and is expected to go away
in the next CUDA release.
Differential Revision: https://reviews.llvm.org/D57488
llvm-svn: 352799
to reflect the new license.
We understand that people may be surprised that we're moving the header
entirely to discuss the new license. We checked this carefully with the
Foundation's lawyer and we believe this is the correct approach.
Essentially, all code in the project is now made available by the LLVM
project under our new license, so you will see that the license headers
include that license only. Some of our contributors have contributed
code under our old license, and accordingly, we have retained a copy of
our old license notice in the top-level files in each project and
repository.
llvm-svn: 351636
All of the other constructors already take a reference to the AST context.
This avoids calling Decl::getASTContext in most cases. Additionally move
the definition of the constructor from Expr.h to Expr.cpp since it is calling
DeclRefExpr::computeDependence. NFC.
llvm-svn: 349901
Basically, "AttributeList" loses all list-like mechanisms, ParsedAttributes is
switched to use a TinyPtrVector (and a ParsedAttributesView is created to
have a non-allocating attributes list). DeclaratorChunk gets the later kind,
Declarator/DeclSpec keep ParsedAttributes.
Iterators are added to the ParsedAttribute types so that for-loops work.
llvm-svn: 336945
We were already performing checks on non-template variables,
but the checks on templated ones were missing.
Differential Revision: https://reviews.llvm.org/D45231
llvm-svn: 334143
This is similar to the LLVM change https://reviews.llvm.org/D46290.
We've been running doxygen with the autobrief option for a couple of
years now. This makes the \brief markers into our comments
redundant. Since they are a visual distraction and we don't want to
encourage more \brief markers in new code either, this patch removes
them all.
Patch produced by
for i in $(git grep -l '\@brief'); do perl -pi -e 's/\@brief //g' $i & done
for i in $(git grep -l '\\brief'); do perl -pi -e 's/\\brief //g' $i & done
Differential Revision: https://reviews.llvm.org/D46320
llvm-svn: 331834
HIP is a language similar to CUDA (https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_kernel_language.md ).
The language syntax is very similar, which allows a hip program to be compiled as a CUDA program by Clang. The main difference
is the host API. HIP has a set of vendor neutral host API which can be implemented on different platforms. Currently there is open source
implementation of HIP runtime on amdgpu target (https://github.com/ROCm-Developer-Tools/HIP).
This patch adds support of input kind and language standard hip.
When hip file is compiled, both LangOpts.CUDA and LangOpts.HIP is turned on. This allows compilation of hip program as CUDA
in most cases and only special handling of hip program is needed LangOpts.HIP is checked.
This patch also adds support of kernel launching of HIP program using HIP host API.
When -x hip is not specified, there is no behaviour change for CUDA.
Patch by Greg Rodgers.
Revised and lit test added by Yaxun Liu.
Differential Revision: https://reviews.llvm.org/D44984
llvm-svn: 330790
We were already performing checks on non-template variables,
but the checks on templated ones were missing.
Differential Revision: https://reviews.llvm.org/D45231
llvm-svn: 329127
The diagnostic system for Clang can already handle many AST nodes. Instead
of converting them to strings first, just hand the AST node directly to
the diagnostic system and let it handle the output. Minor changes in some
diagnostic output.
llvm-svn: 328688
Launching a kernel from the host code does not generate code for the
kernel itself. This fixes an issue with clang erroneously reporting
an error for a HD->D call from within the kernel.
Differential Revision: https://reviews.llvm.org/D44837
llvm-svn: 328362
While a function body is being parsed, the function declaration is not considered
as a definition because it does not have a body yet. In some cases it leads to
incorrect interpretation, the case is presented in
https://bugs.llvm.org/show_bug.cgi?id=14785:
```
template<typename T> struct Somewhat {
void internal() const {}
friend void operator+(int const &, Somewhat<T> const &) {}
};
void operator+(int const &, Somewhat<char> const &x) { x.internal(); }
```
When statement `x.internal()` in the body of global `operator+` is parsed, the type
of `x` must be completed, so the instantiation of `Somewhat<char>` is started. It
instantiates the declaration of `operator+` defined inline, and makes a check for
redefinition. The check does not detect another definition because the declaration
of `operator+` is still not defining as does not have a body yet.
To solves this problem the function `isThisDeclarationADefinition` considers
a function declaration as a definition if it has flag `WillHaveBody` set.
This change fixes PR14785.
Differential Revision: https://reviews.llvm.org/D30375
This is a recommit of 305379, reverted in 305381, with small changes.
llvm-svn: 305903
While a function body is being parsed, the function declaration is not considered
as a definition because it does not have a body yet. In some cases it leads to
incorrect interpretation, the case is presented in
https://bugs.llvm.org/show_bug.cgi?id=14785:
```
template<typename T> struct Somewhat {
void internal() const {}
friend void operator+(int const &, Somewhat<T> const &) {}
};
void operator+(int const &, Somewhat<char> const &x) { x.internal(); }
```
When statement `x.internal()` in the body of global `operator+` is parsed, the type
of `x` must be completed, so the instantiation of `Somewhat<char>` is started. It
instantiates the declaration of `operator+` defined inline, and makes a check for
redefinition. The check does not detect another definition because the declaration
of `operator+` is still not defining as does not have a body yet.
To solves this problem the function `isThisDeclarationADefinition` considers
a function declaration as a definition if it has flag `WillHaveBody` set.
This change fixes PR14785.
Differential Revision: https://reviews.llvm.org/D30375
llvm-svn: 305379
Some functions and templates are treated as __host__ __device__ even
when they don't have explicitly specified target attributes.
What's worse, this treatment may change depending on command line
options (-fno-cuda-host-device-constexpr) or
#pragma clang force_cuda_host_device.
Combined with strict checking for matching function target that comes
with D25809(r288962), it makes it hard to write code which would
explicitly instantiate or specialize some functions regardless of
pragmas or command line options in effect.
This patch changes the way we match target attributes of base template
vs attributes used in explicit instantiation or specialization so that
only explicitly specified attributes are considered. This makes base
template selection behave consistently regardless of pragma of command
line options that may affect CUDA target.
Differential Revision: https://reviews.llvm.org/D25845
llvm-svn: 289091
* __host__ __device__ functions are no longer considered to be
redeclarations of __host__ or __device__ functions. This prevents
unintentional merging of target attributes across them.
* Function target attributes are not considered (and must match) during
explicit instantiation and specialization of function templates.
Differential Revision: https://reviews.llvm.org/D25809
llvm-svn: 288962
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
Yaxun Liu