Commit Graph

137 Commits

Author SHA1 Message Date
Yaxun Liu e5d17c511f [CUDA][HIP] Fix hostness of defaulted constructor
Clang does not respect the explicit device host attributes of defaulted special members.
Also clang does not respect the hostness of special members determined by their
first declarations.
Clang also adds duplicate implicit device or host attributes in certain cases.
This patch fixes that.
Differential Revision: https://reviews.llvm.org/D67509

llvm-svn: 372394
2019-09-20 14:28:09 +00:00
Michael Liao b8fc6a9116 [CUDA][HIP] Re-apply part of r372318.
- r372318 causes violation of `use-of-uninitialized-value` detected by
  MemorySanitizer. Once `Viable` field is set to false, `FailureKind`
  needs setting as well as it will be checked during destruction if
  `Viable` is not true.
- Revert the part trying to skip `std::vector` erasing.

llvm-svn: 372356
2019-09-19 21:26:18 +00:00
Mitch Phillips 08f938bd1a Revert "[CUDA][HIP] Fix typo in `BestViableFunction`"
Broke the msan buildbots (see comments on rL372318 for more details).

This reverts commit eb231d1582.

llvm-svn: 372353
2019-09-19 21:11:28 +00:00
Michael Liao eb231d1582 [CUDA][HIP] Fix typo in `BestViableFunction`
Summary:
- Should consider viable ones only when checking SameSide candidates.
- Replace erasing with clearing viable flag to reduce data
  moving/copying.
- Add one and revise another one as the diagnostic message are more
  relevant compared to previous one.

Reviewers: tra

Subscribers: cfe-commits, yaxunl

Tags: #clang

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

llvm-svn: 372318
2019-09-19 13:14:03 +00:00
Sunil Srivastava 85d667fcb6 Renamed and changed the wording of warn_cconv_ignored
As discussed in D64780 the wording of this warning message is being
changed to say 'is not supported' instead of 'ignored', and the
diag ID itself is being changed to warn_cconv_not_supported.

llvm-svn: 366368
2019-07-17 20:41:26 +00:00
Yaxun Liu c3dfe9082b [HIP] Support attribute hip_pinned_shadow
This patch introduces support of hip_pinned_shadow variable for HIP.

A hip_pinned_shadow variable is a global variable with attribute hip_pinned_shadow.
It has external linkage on device side and has no initializer. It has internal
linkage on host side and has initializer or static constructor. It can be accessed
in both device code and host code.

This allows HIP runtime to implement support of HIP texture reference.

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

llvm-svn: 364381
2019-06-26 03:47:37 +00:00
Erich Keane 505427cb2f Permit redeclarations of a builtin to specify calling convention.
After https://reviews.llvm.org/rL355317 we noticed that quite a decent
amount of code redeclares builtins (memcpy in particular, I believe
reduced from an MSVC header) with a calling convention specified.
This gets particularly troublesome when the user specifies a new
'default' calling convention on the command line.

When looking to add a diagnostic for this case, it was noticed that we
had 3 other diagnostics that differed only slightly.  This patch ALSO
unifies those under a 'select'.  Unfortunately, the order of words in
ONE of these diagnostics was reversed ("'thiscall' calling convention"
vs "calling convention 'thiscall'"), so this patch also standardizes on
the former.

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

Change-Id: I79f99fe7c2301640755ffdd774b46eb44526bb22
llvm-svn: 356663
2019-03-21 13:30:56 +00:00
Yaxun Liu c5be267003 [CUDA][HIP][Sema] Fix template kernel with function as template parameter
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
2019-03-05 18:19:35 +00:00
Yaxun Liu fa49c3a888 [CUDA][HIP] Check calling convention based on function target
MSVC header files using vectorcall to differentiate overloaded functions, which
causes failure for AMDGPU target. This is because clang does not check function
calling convention based on function target.

This patch checks calling convention using the proper target info.

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

llvm-svn: 354929
2019-02-26 22:24:49 +00:00
Alexey Bataev 305b6b9647 [OPENMP][CUDA]Do not emit warnings for variables in late-reported asm
statements.

If the assembler instruction is not generated and the delayed diagnostic
is emitted, we may end up with extra warning message for variables used
in the asm statement. Since the asm statement is not built, the
variables may be left non-referenced and it may produce a warning about
a use of the non-initialized variables.

llvm-svn: 354928
2019-02-26 21:51:16 +00:00
Michael Liao 7557afa000 [AMDGPU] Allow using integral non-type template parameters
Summary:
- Allow using integral non-type template parameters in the following
  attributes

  __attribute__((amdgpu_flat_work_group_size(<min>, <max>)))
  __attribute__((amdgpu_waves_per_eu(<min>[, <max>])))

Reviewers: kzhuravl, yaxunl

Subscribers: jvesely, wdng, nhaehnle, dstuttard, tpr, t-tye, jdoerfert, cfe-commits

Tags: #clang

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

llvm-svn: 354909
2019-02-26 18:49:36 +00:00
Alexey Bataev e69f94e022 [OPENMP] Delayed diagnostics for VLA support.
Generalized processing of the deferred diagnostics for OpenMP/CUDA code.

llvm-svn: 354690
2019-02-22 20:36:10 +00:00
Alexey Bataev bbd5c55c66 Revert "[OPENMP] Delayed diagnostics for VLA support."
This reverts commit r354679 to fix the problem with the Windows
buildbots

llvm-svn: 354680
2019-02-22 17:16:50 +00:00
Alexey Bataev b09bcf8efd [OPENMP] Delayed diagnostics for VLA support.
Generalized processing of the deferred diagnostics for OpenMP/CUDA code.

llvm-svn: 354679
2019-02-22 16:49:13 +00:00
Alexey Bataev 3167b3035e [CUDA]Delayed diagnostics for the asm instructions.
Adapted targetDiag for the CUDA and used for the delayed diagnostics in
asm constructs. Works for both host and device compilation sides.

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

llvm-svn: 354671
2019-02-22 14:42:48 +00:00
Alexey Bataev 12a21e4b69 Revert "[CUDA]Delayed diagnostics for the asm instructions."
This reverts commit r354593 to fix the problem with the crash on
windows.

llvm-svn: 354596
2019-02-21 16:40:21 +00:00
Alexey Bataev 16d3e1a4d2 [CUDA]Delayed diagnostics for the asm instructions.
Summary:
Adapted targetDiag for the CUDA and used for the delayed diagnostics in
asm constructs. Works for both host and device compilation sides.

Reviewers: tra, jlebar

Subscribers: jdoerfert, cfe-commits

Tags: #clang

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

llvm-svn: 354593
2019-02-21 15:51:30 +00:00
Artem Belevich c62214da3d [CUDA] add support for the new kernel launch API in CUDA-9.2+.
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
2019-01-31 21:34:03 +00:00
Yaxun Liu 95f2ca541f [HIP] Fix size_t for MSVC environment
In 64 bit MSVC environment size_t is defined as unsigned long long.
In single source language like HIP, data layout should be consistent
in device and host compilation, therefore copy data layout controlling
fields from Aux target for AMDGPU target.

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

llvm-svn: 352620
2019-01-30 12:26:54 +00:00
Yaxun Liu d442500f5d [CUDA][HIP] Do not diagnose use of _Float16
r352221 caused regressions in CUDA/HIP since device function may use _Float16 whereas host does not support it.
In this case host compilation should not diagnose usage of _Float16 in device functions or variables.

For now just do not diagnose _Float16 for CUDA/HIP. In the future we should have more precise check.

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

llvm-svn: 352488
2019-01-29 13:20:23 +00:00
Yaxun Liu a461174cfd [CUDA][HIP] Fix ShouldDeleteSpecialMember for inherited constructors
ShouldDeleteSpecialMember is called upon inherited constructors.
It calls inferCUDATargetForImplicitSpecialMember.

Normally the special member enum passed to ShouldDeleteSpecialMember
matches the constructor. However this is not true when inherited
constructor is passed, where DefaultConstructor is passed to treat
the inherited constructor as DefaultConstructor. However
inferCUDATargetForImplicitSpecialMember expects the special
member enum argument to match the constructor, which results
in assertion when this expection is not satisfied.

This patch checks whether the constructor is inherited. If true it will
get the real special member enum for the constructor and pass it
to inferCUDATargetForImplicitSpecialMember.

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

llvm-svn: 344057
2018-10-09 15:53:14 +00:00
Yaxun Liu 9767089d00 [HIP] Support early finalization of device code for -fno-gpu-rdc
This patch renames -f{no-}cuda-rdc to -f{no-}gpu-rdc and keeps the original
options as aliases. When -fgpu-rdc is off,
clang will assume the device code in each translation unit does not call
external functions except those in the device library, therefore it is possible
to compile the device code in each translation unit to self-contained kernels
and embed them in the host object, so that the host object behaves like
usual host object which can be linked by lld.

The benefits of this feature is: 1. allow users to create static libraries which
can be linked by host linker; 2. amortized device code linking time.

This patch modifies HIP action builder to insert actions for linking device
code and generating HIP fatbin, and pass HIP fatbin to host backend action.
It extracts code for constructing command for generating HIP fatbin as
a function so that it can be reused by early finalization. It also modifies
codegen of HIP host constructor functions to embed the device fatbin
when it is available.

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

llvm-svn: 343611
2018-10-02 17:48:54 +00:00
Richard Smith 9b2c5e7c44 [cxx2a] P0641R2: (Some) type mismatches on defaulted functions only
render the function deleted instead of rendering the program ill-formed.

This change also adds an enabled-by-default warning for the case where
an explicitly-defaulted special member function of a non-template class
is implicitly deleted by the type checking rules. (This fires either due
to this language change or due to pre-C++20 reasons for the member being
implicitly deleted). I've tested this on a large codebase and found only
bugs (where the program means something that's clearly different from
what the programmer intended), so this is enabled by default, but we
should revisit this if there are problems with this being enabled by
default.

llvm-svn: 343285
2018-09-28 01:16:43 +00:00
Artem Belevich 78929efb4d [CUDA] Ignore uncallable functions when we check for usual deallocators.
Previously clang considered function variants from both sides of
compilation and that resulted in picking up wrong deallocation function.

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

llvm-svn: 342749
2018-09-21 17:29:33 +00:00
Artem Belevich fe5b1ac142 Revert the tests that should've been reverted in rL341115
llvm-svn: 341118
2018-08-30 20:53:15 +00:00
Jonas Hahnfeld 931939bf92 [CUDA/OpenMP] Define only some host macros during device compilation
When compiling CUDA or OpenMP device code Clang parses header files
that expect certain predefined macros from the host architecture. To
make this work the compiler passes the host triple via the -aux-triple
argument and (until now) pulls in all macros for that "auxiliary triple"
unconditionally.

However this results in defines like __SSE_MATH__ that will trigger
inline assembly making use of the "advertised" target features. See
the discussion of D47849 and PR38464 for a detailed explanation of
the encountered problems.

Instead of blacklisting "known bad" examples this patch starts adding
defines that are needed for certain headers like bits/wordsize.h and
bits/mathinline.h.
The disadvantage of this approach is that it decouples the definitions
from their target toolchain. However in my opinion it's more important
to keep definitions for one header close together. For one this will
include a clear documentation why these particular defines are needed.
Furthermore it simplifies maintenance because adding defines for a new
header or support for a new aux-triple only needs to touch one piece
of code.

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

llvm-svn: 340681
2018-08-25 13:42:40 +00:00
Yaxun Liu a4005e13f7 [CUDA][HIP] Allow function-scope static const variable
CUDA 8.0 E.3.9.4 says: Within the body of a __device__ or __global__
function, only __shared__ variables or variables without any device
memory qualifiers may be declared with static storage class.

It is unclear how a function-scope non-const static variable
without device memory qualifier is implemented, therefore only static
const variable without device memory qualifier is allowed, which
can be emitted as a global variable in constant address space.

Currently clang only allows function-scope static variable with
__shared__ qualifier.

This patch also allows function-scope static const variable without
device memory qualifier and emits it as a global variable in constant
address space.

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

llvm-svn: 338188
2018-07-28 03:05:25 +00:00
Yaxun Liu aa24601f98 [CUDA][HIP] Allow CUDA __global__ functions to have amdgpu kernel attributes
There are HIP applications e.g. Tensorflow 1.3 using amdgpu kernel attributes, however
currently they are only allowed on OpenCL kernel functions.

This patch will allow amdgpu kernel attributes to be applied to CUDA/HIP __global__
functions.

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

llvm-svn: 334561
2018-06-12 23:58:59 +00:00
Artem Belevich e9fa53a09b [CUDA] Check initializers of instantiated template variables.
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
2018-06-06 22:37:25 +00:00
Justin Lebar 5489f85fda [CUDA] Allow "extern __shared__ Foo foo[]" within anon. namespaces.
Summary:
Previously this triggered a -Wundefined-internal warning.  But it's not
an undefined variable -- any variable of this form is a pointer to the
base of GPU core's shared memory.

Reviewers: tra

Subscribers: sanjoy, rsmith

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

llvm-svn: 332621
2018-05-17 16:15:07 +00:00
Alexander Kornienko 2a8c18d991 Fix typos in clang
Found via codespell -q 3 -I ../clang-whitelist.txt
Where whitelist consists of:

  archtype
  cas
  classs
  checkk
  compres
  definit
  frome
  iff
  inteval
  ith
  lod
  methode
  nd
  optin
  ot
  pres
  statics
  te
  thru

Patch by luzpaz! (This is a subset of D44188 that applies cleanly with a few
files that have dubious fixes reverted.)

Differential revision: https://reviews.llvm.org/D44188

llvm-svn: 329399
2018-04-06 15:14:32 +00:00
Artem Belevich 67d22c8a84 Revert "[CUDA] Check initializers of instantiated template variables."
This (temporarily) reverts commit r329127 due to the problems
it exposed in TensorFlow.

llvm-svn: 329229
2018-04-04 20:48:42 +00:00
Artem Belevich d9189d1e76 [CUDA] Check initializers of instantiated template variables.
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
2018-04-03 22:41:06 +00:00
Richard Trieu b402580616 Fix some handling of AST nodes with diagnostics.
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
2018-03-28 04:16:13 +00:00
Artem Belevich e2ae8b5510 [CUDA] Fixed false error reporting in case of calling H->G->HD->D.
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
2018-03-23 19:49:03 +00:00
Jonas Hahnfeld ee47d8cb96 [CUDA] Allow external variables in separate compilation
According to the CUDA Programming Guide this is prohibited in
whole program compilation mode. This makes sense because external
references cannot be satisfied in that mode anyway. However,
such variables are allowed in separate compilation mode which
is a valid use case.

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

llvm-svn: 325136
2018-02-14 16:04:03 +00:00
Artem Belevich 50e6e54587 [CUDA] Report "unsupported VLA" errors only on device side.
This fixes erroneously reported CUDA compilation errors
in host-side code during device-side compilation.

I've also restricted OpenMP-specific checks to trigger only
if we're compiling with OpenMP enabled.

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

llvm-svn: 319201
2017-11-28 18:51:42 +00:00
Aaron Ballman adf66b6174 Determine the attribute subject for diagnostics based on declarative information in DeclNodes.td. This greatly reduces the number of enumerated values used for more complex diagnostics; these are now only required when the "attribute only applies to" diagnostic needs to be generated manually as part of semantic processing.
This also clarifies some terminology used by the diagnostic (methods -> Objective-C methods, fields -> non-static data members, etc).

Many of the tests needed to be updated in multiple places for the diagnostic wording tweaks. The first instance of the diagnostic for that attribute is fully specified and subsequent instances cut off the complete list (to make it easier if additional subjects are added in the future for the attribute).

llvm-svn: 319002
2017-11-26 20:01:12 +00:00
Justin Lebar 78137ec868 [CUDA] When compilation fails, print the compilation mode.
Summary:
That is, instead of "1 error generated", we now say "1 error generated
when compiling for sm_35".

This (partially) solves a usability foogtun wherein e.g. users call a
function that's only defined on sm_60 when compiling for sm_35, and they
get an unhelpful error message.

Reviewers: tra

Subscribers: sanjoy, cfe-commits

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

llvm-svn: 312736
2017-09-07 18:37:16 +00:00
Richard Smith 6c716116df PR34163: Don't cache an incorrect key function for a class if queried between
the class becoming complete and its inline methods being parsed.

This replaces the hack of using the "late parsed template" flag to track member
functions with bodies we've not parsed yet; instead we now use the "will have
body" flag, which carries the desired implication that the function declaration
*is* a definition, and that we've just not parsed its body yet.

llvm-svn: 310776
2017-08-12 01:46:03 +00:00
Justin Lebar 86c4e63ff9 [CUDA] Let NVPTX inherit the host's calling conventions.
Summary:
When compiling device code, we may still see host code with explicit
calling conventions.  NVPTX needs to claim that it supports these CCs,
so that (a) we don't raise noisy warnings, and (b) we don't break
existing code which relies on the existence of these CCs when
specializing templates.  (If a CC doesn't exist, clang ignores it, so
two template specializations which are different only insofar as one
specifies a CC are considered identical and therefore are an error if
that CC is not supported.)

Reviewers: tra

Subscribers: cfe-commits

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

llvm-svn: 291136
2017-01-05 16:53:38 +00:00
Justin Lebar 0203f2c26e [CUDA] Add __declspec spellings for CUDA attributes.
Summary: CUDA attributes are spelled __declspec(__foo__) on Windows.

Reviewers: tra

Subscribers: cfe-commits, rnk

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

llvm-svn: 291134
2017-01-05 16:53:04 +00:00
Artem Belevich 64135c35f7 [CUDA] Ignore implicit target attributes during function template instantiation.
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
2016-12-08 19:38:13 +00:00
Artem Belevich 13e9b4d768 [CUDA] Improve target attribute checking for function templates.
* __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
2016-12-07 19:27:16 +00:00
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 2b42ccc78b [CUDA] [AST] Allow isInlineDefinitionExternallyVisible to be called on functions without bodies.
Summary:
In CUDA compilation, we call isInlineDefinitionExternallyVisible (via
getGVALinkageForFunction) on functions while parsing their definitions.

At the point in time when we call getGVALinkageForFunction, we haven't
yet added the body to the function, so we trip this assert.  But as far
as I can tell, this is harmless.

To work around this, we add a new flag to FunctionDecl, "WillHaveBody".

There was other code that was working around the existing assert with a
really awful hack -- this change lets us get rid of that hack.

Reviewers: rsmith, tra

Subscribers: aemerson, cfe-commits

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

llvm-svn: 285410
2016-10-28 16:26:26 +00:00
Justin Lebar 4d38a5cf74 [CUDA] Simplify some repeated diagnostic expectations in CUDA tests.
Instead of repeating the diagnostic, use "expected-note N".

Test-only change.

llvm-svn: 284882
2016-10-21 20:50:47 +00:00
Artem Belevich 07db5cf6c8 Declare H and H new/delete.
llvm-svn: 284879
2016-10-21 20:34:05 +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