[PATCH] D25809: [CUDA] Improved target attribute-based overloading.

[Artem Belevich via cfe-commits]

tra created this revision.
tra added reviewers: jlebar, rsmith.
tra added a subscriber: cfe-commits.

Current behavior:

- __host__ __device__ (HD) functions are considered to be redeclarations of `__host__` (H) of `__device__` (D) functions with same signature.
- Target attributes are not taken into account during selection of function template during explicit instantiation and specialization.

Issues:
a) It's possible for a H or D function to inherit HD attributes from a HD declaration which results in those functions being silently treated as HD in the rest of the code which leads to clang accepting the code instead of diagnosing it as an error.
b) If we have definitions of HD and a H or D function, compiler complains about redefinition of the same function, which is misleading.
c) It is impossible to disambiguate across target-overloaded function templates during explicit instantiation/specialization.

Changes in this patch:
a) treat HD functions as overloads and add Sema checks to explicitly diagnose attempts to overload HD functions with H or D ones.
b) Require matching target attributes for explicit function template instantiation/specialization and narrow list of candidates to templates with the same target. Diagnose rejected candidates.

The changes (a) and (b) can be split into separate patches, but both must be committed simultaneously as
(a) alone further breaks function template instantiation/specialization when target attributes are involved and (b) is half-broken until (a) is in place and prevents HD attributes merging into functions with H or D attributes.

Open issues:

- It's not clear how to handle explicit specialization of constexpr function templates:
  - Implicit target attributes of constexpr functions and templates change depending on whether -fno-cuda-host-device-constexpr is in effect.
  - C++11 [dcl.constexpr]p1: An explicit specialization of a constexpr function can differ from the template declaration with respect to the constexpr specifier.

One idea is to only match explicitly written target attributes when we choose candidate templates. This makes it easier to tell which template we instantiate based only on what's in the source we compile.


https://reviews.llvm.org/D25809

Files:
  include/clang/Basic/DiagnosticSemaKinds.td
  include/clang/Sema/Sema.h
  lib/Sema/SemaCUDA.cpp
  lib/Sema/SemaDecl.cpp
  lib/Sema/SemaOverload.cpp
  lib/Sema/SemaTemplate.cpp
  test/CodeGenCUDA/launch-bounds.cu
  test/SemaCUDA/function-overload.cu
  test/SemaCUDA/function-template-overload.cu
  test/SemaCUDA/target_attr_inheritance.cu

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D25809.75238.patch
Type: text/x-patch
Size: 25035 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20161020/be6da813/attachment-0001.bin>