[hip] Relax CUDA call restriction within `decltype` context.

[hliao]

Summary

• Within decltype, expressions are only type-inspected. The restriction on CUDA calls should be relaxed.

[tra]

Perhaps we should allow this in all unevaluated contexts?
I.e. int s = sizeof(foo(x)); should also work.

clang/include/clang/Sema/Sema.h

10972

I think you want return llvm::any_of(ExprEvalContexts, ...) here and you can fold it directly into if() below.

[hliao]

In D61458#1488523, @tra wrote:

Perhaps we should allow this in all unevaluated contexts?
I.e. int s = sizeof(foo(x)); should also work.

good point, do we have a dedicated context for sizeof? that make the checking easier.

clang/include/clang/Sema/Sema.h

10972

yeah, that's much simpler, I will make the change.

[hliao]

simplify the logic using llvm::any_of.

[tra]

In D61458#1488550, @hliao wrote:

In D61458#1488523, @tra wrote:

Perhaps we should allow this in all unevaluated contexts?
I.e. int s = sizeof(foo(x)); should also work.

good point, do we have a dedicated context for sizeof? that make the checking easier.

Sema::isUnevaluatedContext() may be able to do the job.

[tra]

clang/include/clang/Sema/Sema.h

10969–10975

One more thing. The idea of this function is that we're checking if the Caller is allowed to call the Callee.
However here, you're checking the current context, which may not necessarily be the same as the caller's. I.e. someone could potentially call it way after the context is gone.

Currently all uses of this function obtain the caller from CurContext, but if we start relying on other properties of the current context other than the caller function, then we may neet to pass the context explicitly, or only pass the Callee and check if it's callable from the current context.

[jlebar]

Here's one for you:

__host__ float bar();
__device__ int bar();
__host__ __device__ auto foo() -> decltype(bar()) {}

What is the return type of foo? :)

I don't believe the right answer is, "float when compiling for host, int when compiling for device."

I'd be happy if we said this was an error, so long as it's well-defined what exactly we're disallowing. But I bet @rsmith can come up with substantially more evil testcases than this.

[hfinkel]

In D61458#1488970, @jlebar wrote:

Here's one for you:

__host__ float bar();
__device__ int bar();
__host__ __device__ auto foo() -> decltype(bar()) {}

What is the return type of foo? :)

I don't believe the right answer is, "float when compiling for host, int when compiling for device."

So, actually, I wonder if that's not the right answer. We generally allow different overloads to have different return types. What if, for example, the return type on the host is __float128 and on the device it's MyLongFloatTy?

I'd be happy if we said this was an error, so long as it's well-defined what exactly we're disallowing. But I bet @rsmith can come up with substantially more evil testcases than this.

[rjmccall]

In D61458#1488972, @hfinkel wrote:

In D61458#1488970, @jlebar wrote:

Here's one for you:

__host__ float bar();
__device__ int bar();
__host__ __device__ auto foo() -> decltype(bar()) {}

What is the return type of foo? :)

I don't believe the right answer is, "float when compiling for host, int when compiling for device."

So, actually, I wonder if that's not the right answer. We generally allow different overloads to have different return types.

Only if they also differ in some other way. C++ does not (generally) have return-type-based overloading. The two functions described would even mangle the same way if CUDA didn't include host/device in the mangling.

(Function templates can differ only by return type, but if both return types successfully instantiate for a given set of (possibly inferred) template arguments then the templates can only be distinguished when taking their address, not when calling.)

I think I've said before that adding this kind of overloading is not a good idea, but since it's apparently already there, you should consult the specification (or at least existing practice) to figure out what you're supposed to do.

[hfinkel]

Only if they also differ in some other way. C++ does not (generally) have return-type-based overloading. The two functions described would even mangle the same way if CUDA didn't include host/device in the mangling.

Certainly. I didn't mean to imply otherwise.

[hliao]

In D61458#1488970, @jlebar wrote:

Here's one for you:

__host__ float bar();
__device__ int bar();
__host__ __device__ auto foo() -> decltype(bar()) {}

What is the return type of foo? :)

I don't believe the right answer is, "float when compiling for host, int when compiling for device."

I'd be happy if we said this was an error, so long as it's well-defined what exactly we're disallowing. But I bet @rsmith can come up with substantially more evil testcases than this.

This patch is introduced to allow function or template function from std library to be used with device function. By allowing different-side candidates with a context only caring type inspection, we have new issue as there are extra beyond the regular rule for C++ overloadable resolution. We need an extra policy to figure out which is one the best candidate by considering CUDA attributes. Says the case you proposed, we may consider the following order to choose an overloadable candidate, e.g.

SAME-SIDE (with the same CUDA attribute)
NATIVE (without any CUDA attribute)
WRONG-SIDE (with the opposite CUDA attribute)

or just

SAME-SIDE
NATIVE

It that a reasonable change?

[hliao]

In D61458#1488981, @rjmccall wrote:

In D61458#1488972, @hfinkel wrote:

In D61458#1488970, @jlebar wrote:

Here's one for you:

__host__ float bar();
__device__ int bar();
__host__ __device__ auto foo() -> decltype(bar()) {}

What is the return type of foo? :)

I don't believe the right answer is, "float when compiling for host, int when compiling for device."

So, actually, I wonder if that's not the right answer. We generally allow different overloads to have different return types.

Only if they also differ in some other way. C++ does not (generally) have return-type-based overloading. The two functions described would even mangle the same way if CUDA didn't include host/device in the mangling.

(Function templates can differ only by return type, but if both return types successfully instantiate for a given set of (possibly inferred) template arguments then the templates can only be distinguished when taking their address, not when calling.)

I think I've said before that adding this kind of overloading is not a good idea, but since it's apparently already there, you should consult the specification (or at least existing practice) to figure out what you're supposed to do.

BTW, just check similar stuff with nvcc, with more than one candidates, it accepts the following code

float bar(); // This line could be replaced by appendig `__host` or `__device__`, all of them are accepted.
__host__ __device__ auto foo() -> decltype(bar()) {}

however, if there are more than one candidates differenct on the return type (without or with CUDA attibute difference), it could raise the error

foo.cu(4): error: cannot overload functions distinguished by return type alone

it seems to me that that's also an acceptable policy to handle the issue after we allow different-side candidates in type-only context.

[hliao]

In D61458#1488970, @jlebar wrote:

Here's one for you:

__host__ float bar();
__device__ int bar();
__host__ __device__ auto foo() -> decltype(bar()) {}

What is the return type of foo? :)

I don't believe the right answer is, "float when compiling for host, int when compiling for device."

I'd be happy if we said this was an error, so long as it's well-defined what exactly we're disallowing. But I bet @rsmith can come up with substantially more evil testcases than this.

At from CUDA 10, that's not acceptable as we are declaring two functions only differ from the return type. It seems CUDA attributes do not contribute to the function signature. clang is quite different here.

[hliao]

clang/include/clang/Sema/Sema.h

10969–10975

as the expression within decltype may be quite complicated, the idea here is to relax that rule within decltype context, not only for a particular pair of caller/callee.

[jlebar]

At [nvcc] from CUDA 10, that's not acceptable as we are declaring two functions only differ from the return type. It seems CUDA attributes do not contribute to the function signature. clang is quite different here.

Yes, this is an intentional and more relaxed semantics in clang. It's also sort of the linchpin of our mixed-mode compilation strategy, which is very different from nvcc's source-to-source splitting strategy.

Back in the day you could trick nvcc into allowing host/device overloading on same-signature functions by slapping a template on one or both of them. Checking just now it seems they fixed this, but I suspect there are still dark corners where nvcc relies on effectively the same behavior as we get in clang via true overloading.

[tra]

clang/include/clang/Sema/Sema.h

10969–10975

I understand the idea, but in this case the argument was more about the code style.

Currently the contract is that the function's decision is derived from its arguments (and could, perhaps, be a static method). With this patch you start relying on the context, but it's not obvious from the function signature. Replacing Caller with context, or removing the caller altogether would bring the function signature closer to what the function does.

[hliao]

This patch is revived with more changes addressing the previous concerns.

Back to Justin's example:

__host__ float bar();
__device__ int bar();
__host__ __device__ auto foo() -> decltype(bar()) { return bar(); }

Even without this patch, that example already passed the compilation without
either errors or warnings. Says

clang -std=c++11 -x cuda -nocudainc -nocudalib --cuda-gpu-arch=sm_60 --cuda-device-only -S -emit-llvm -O3 foo.cu

In c++14, that example could be even simplified without decltype but the same ambiguity.

__host__ float bar();
__device__ int bar();
__host__ __device__ auto foo() { return bar(); }

Without any change, clang also compiles the code as well and uses different return types between host-side and device-side compilation.[^1]

[^1]: The first example has the same return type between host-side and device-side but that seems incorrect or unreasonable to me.

The ambiguity issue is in fact not introduced by relaxing decltype. That's an inherent one as we allow overloading over target attributes. Issuing warnings instead of errors seems more reasonable to me for such cases.

In this patch, besides relaxing the CUDA call rule under decltype, it also generates warning during function overloading if there are more than candidates with different return types.

clang/include/clang/Basic/DiagnosticSemaKinds.td

def err_cuda_host_shared : Error<
     "__shared__ local variables not allowed in "
     "%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
 def err_cuda_nonglobal_constant : Error<"__constant__ variables must be global">;
 def err_cuda_ovl_target : Error<
   "%select{__device__|__global__|__host__|__host__ __device__}0 function %1 "
   "cannot overload %select{__device__|__global__|__host__|__host__ __device__}2 function %3">;
 def note_cuda_ovl_candidate_target_mismatch : Note<
     "candidate template ignored: target attributes do not match">;
+def warn_decltype_ambiguous_return_type : Warning<
+    "return type of %0 in 'decltype' is ambiguous and may not be expected">;
+def note_decltype_ambiguous_function_chosen : Note<
+    "use this definition of %0">;
+def note_decltype_ambiguous_function_other : Note<
+    "other definition of %0">;
 
 def warn_non_pod_vararg_with_format_string : Warning<
   "cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic "
   "%select{function|block|method|constructor}2; expected type from format "
   "string was %3">, InGroup<NonPODVarargs>, DefaultError;
 // The arguments to this diagnostic should match the warning above.
 def err_cannot_pass_objc_interface_to_vararg_format : Error<
   "cannot pass object with interface type %1 by value to variadic "

clang/include/clang/Sema/Sema.h

   /// Determines whether we are currently in a context that
   /// is not evaluated as per C++ [expr] p5.
   bool isUnevaluatedContext() const {
     assert(!ExprEvalContexts.empty() &&
            "Must be in an expression evaluation context");
     return ExprEvalContexts.back().isUnevaluated();
   }
 
+  bool underDecltypeContext() const {
+    return llvm::any_of(ExprEvalContexts,
+                        [](const ExpressionEvaluationContextRecord &C) {
+                          return C.ExprContext ==
+                                 ExpressionEvaluationContextRecord::EK_Decltype;
+                        });
+  }
+
   /// RAII class used to determine whether SFINAE has
   /// trapped any errors that occur during template argument
   /// deduction.
   class SFINAETrap {
     Sema &SemaRef;
     unsigned PrevSFINAEErrors;
     bool PrevInNonInstantiationSFINAEContext;
     bool PrevAccessCheckingSFINAE;
public:
   /// \param Caller function which needs address of \p Callee.
   ///               nullptr in case of global context.
   /// \param Callee target function
   ///
   /// \returns preference value for particular Caller/Callee combination.
   CUDAFunctionPreference IdentifyCUDAPreference(const FunctionDecl *Caller,
                                                 const FunctionDecl *Callee);
 
-  /// Determines whether Caller may invoke Callee, based on their CUDA
-  /// host/device attributes.  Returns false if the call is not allowed.
+  /// Determines, under the current context, whether Callee may be invokable,
+  /// based on their CUDA host/device attributes. Returns false if the call is
+  /// not allowed.
   ///
   /// Note: Will return true for CFP_WrongSide calls.  These may appear in
   /// semantically correct CUDA programs, but only if they're never codegen'ed.
-  bool IsAllowedCUDACall(const FunctionDecl *Caller,
-                         const FunctionDecl *Callee) {
-    return IdentifyCUDAPreference(Caller, Callee) != CFP_Never;
+  bool isCUDACallAllowed(const FunctionDecl *Callee) {
+    // Under `decltype`, the rule is relaxed.
+    if (underDecltypeContext())
+      return true;
+    return IdentifyCUDAPreference(dyn_cast<FunctionDecl>(CurContext), Callee) !=
+           CFP_Never;
   }
 
   /// May add implicit CUDAHostAttr and CUDADeviceAttr attributes to FD,
   /// depending on FD and the current compilation settings.
   void maybeAddCUDAHostDeviceAttrs(FunctionDecl *FD,
                                    const LookupResult &Previous);
 
 public:

clang/lib/Sema/SemaOverload.cpp

void Sema::AddOverloadCandidate(
 
   // (CUDA B.1): Check for invalid calls between targets.
   if (getLangOpts().CUDA)
     if (const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext))
       // Skip the check for callers that are implicit members, because in this
       // case we may not yet know what the member's target is; the target is
       // inferred for the member automatically, based on the bases and fields of
       // the class.
-      if (!Caller->isImplicit() && !IsAllowedCUDACall(Caller, Function)) {
+      if (!Caller->isImplicit() && !isCUDACallAllowed(Function)) {
         Candidate.Viable = false;
         Candidate.FailureKind = ovl_fail_bad_target;
         return;
       }
 
   // Determine the implicit conversion sequences for each of the
   // arguments.
   for (unsigned ArgIdx = 0; ArgIdx < Args.size(); ++ArgIdx) {
if (Candidate.Conversions[ConvIdx].isBad()) {
       Candidate.FailureKind = ovl_fail_bad_conversion;
       return;
     }
   }
 
   // (CUDA B.1): Check for invalid calls between targets.
   if (getLangOpts().CUDA)
     if (const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext))
-      if (!IsAllowedCUDACall(Caller, Method)) {
+      if (!isCUDACallAllowed(Method)) {
         Candidate.Viable = false;
         Candidate.FailureKind = ovl_fail_bad_target;
         return;
       }
 
   // Determine the implicit conversion sequences for each of the
   // arguments.
   for (unsigned ArgIdx = 0; ArgIdx < Args.size(); ++ArgIdx) {
OverloadCandidateSet::BestViableFunction(Sema &S, SourceLocation Loc,
   // If we found more than one best candidate, this is ambiguous.
   if (Best == end())
     return OR_Ambiguous;
 
   // Best is the best viable function.
   if (Best->Function && Best->Function->isDeleted())
     return OR_Deleted;
 
+  // Issue a warning of return type resolution under `decltype`.
+  if (S.getLangOpts().CUDA && Best->Function && S.underDecltypeContext()) {
+    SmallVector<const OverloadCandidate *, 16> AmbiSet;
+    QualType BestReturnType = Best->Function->getReturnType();
+    for (auto &Cand : this->Candidates) {
+      if (!Cand.Viable || !Cand.Function)
+        continue;
+      if (BestReturnType != Cand.Function->getReturnType())
+        AmbiSet.push_back(&Cand);
+    }
+    if (!AmbiSet.empty()) {
+      S.Diag(Loc, diag::warn_decltype_ambiguous_return_type) << Best->Function;
+      S.Diag(Best->Function->getLocation(),
+             diag::note_decltype_ambiguous_function_chosen)
+          << Best->Function;
+      for (auto C : AmbiSet)
+        S.Diag(C->Function->getLocation(),
+               diag::note_decltype_ambiguous_function_other)
+            << C->Function;
+    }
+  }
+
   if (!EquivalentCands.empty())
     S.diagnoseEquivalentInternalLinkageDeclarations(Loc, Best->Function,
                                                     EquivalentCands);
 
   return OR_Success;
 }
 
 namespace {
if (CXXMethodDecl *Method = dyn_cast<CXXMethodDecl>(Fn)) {
         return false;
     }
     else if (TargetTypeIsNonStaticMemberFunction)
       return false;
 
     if (FunctionDecl *FunDecl = dyn_cast<FunctionDecl>(Fn)) {
       if (S.getLangOpts().CUDA)
         if (FunctionDecl *Caller = dyn_cast<FunctionDecl>(S.CurContext))
-          if (!Caller->isImplicit() && !S.IsAllowedCUDACall(Caller, FunDecl))
+          if (!Caller->isImplicit() && !S.isCUDACallAllowed(FunDecl))
             return false;
       if (FunDecl->isMultiVersion()) {
         const auto *TA = FunDecl->getAttr<TargetAttr>();
         if (TA && !TA->isDefaultVersion())
           return false;
       }
 
       // If any candidate has a placeholder return type, trigger its deduction

clang/test/CodeGenCUDA/function-overload.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 
 // Make sure we handle target overloads correctly.  Most of this is checked in
 // sema, but special functions like constructors and destructors are here.
 //
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s \
 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST %s
 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s \
 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE %s
+// RUN: %clang_cc1 -std=c++11 -DCHECK_DECLTYPE -triple amdgcn -fcuda-is-device -emit-llvm -o - %s \
+// RUN:     | FileCheck -check-prefix=CHECK-DECLTYPE %s
 
 #include "Inputs/cuda.h"
 
 // Check constructors/destructors for D/H functions
 int x;
 struct s_cd_dh {
   __host__ s_cd_dh() { x = 11; }
   __device__ s_cd_dh() { x = 12; }

 
 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdC2Ev(
 // CHECK-BOTH:   store i32 31,
 // CHECK-BOTH: ret void
 
 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdD2Ev(
 // CHECK-BOTH: store i32 32,
 // CHECK-BOTH: ret void
+
+#if defined(CHECK_DECLTYPE)
+int foo(float);
+// CHECK-DECLTYPE-LABEL: @_Z3barf
+// CHECK-DECLTYPE: fptosi
+// CHECK-DECLTYPE: sitofp
+__device__ float bar(float x) {
+  decltype(foo(x)) y = x;
+  return y + 3.f;
+}
+#endif

clang/test/Misc/warning-flags.c

(2) It prevents us adding new warnings to Clang that have no -W flag.  All
     new warnings should have -W flags.
 
     If you add a new warning without a flag, this test will fail.  To fix
     this test, simply add a warning group to that warning.
 
 
 The list of warnings below should NEVER grow.  It should gradually shrink to 0.
 
-CHECK: Warnings without flags (74):
+CHECK: Warnings without flags (75):
 CHECK-NEXT:   ext_excess_initializers
 CHECK-NEXT:   ext_excess_initializers_in_char_array_initializer
 CHECK-NEXT:   ext_expected_semi_decl_list
 CHECK-NEXT:   ext_explicit_specialization_storage_class
 CHECK-NEXT:   ext_initializer_string_for_char_array_too_long
 CHECK-NEXT:   ext_missing_declspec
 CHECK-NEXT:   ext_missing_whitespace_after_macro_name
 CHECK-NEXT:   ext_new_paren_array_nonconst

 CHECK-NEXT:   warn_call_wrong_number_of_arguments
 CHECK-NEXT:   warn_case_empty_range
 CHECK-NEXT:   warn_char_constant_too_large
 CHECK-NEXT:   warn_collection_expr_type
 CHECK-NEXT:   warn_conflicting_variadic
 CHECK-NEXT:   warn_conv_to_base_not_used
 CHECK-NEXT:   warn_conv_to_self_not_used
 CHECK-NEXT:   warn_conv_to_void_not_used
+CHECK-NEXT:   warn_decltype_ambiguous_return_type
 CHECK-NEXT:   warn_delete_array_type
 CHECK-NEXT:   warn_double_const_requires_fp64
 CHECK-NEXT:   warn_drv_assuming_mfloat_abi_is
 CHECK-NEXT:   warn_drv_clang_unsupported
 CHECK-NEXT:   warn_drv_pch_not_first_include
 CHECK-NEXT:   warn_dup_category_def
 CHECK-NEXT:   warn_enum_value_overflow
 CHECK-NEXT:   warn_expected_qualified_after_typename

clang/test/SemaCUDA/function-overload.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify %s
+// RUN: %clang_cc1 -std=c++11 -DCHECK_DECLTYPE -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
+// RUN: %clang_cc1 -std=c++11 -DCHECK_DECLTYPE -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify %s
 
 #include "Inputs/cuda.h"
 
 // Opaque return types used to check that we pick the right overloads.
 struct HostReturnTy {};
 struct HostReturnTy2 {};
 struct DeviceReturnTy {};
 struct DeviceReturnTy2 {};

 __host__ __device__ int constexpr_overload(const T &x, const T &y) {
   return x - y;
 }
 
 // Verify that function overloading doesn't prune candidate wrongly.
 int test_constexpr_overload(C2 &x, C2 &y) {
   return constexpr_overload(x, y);
 }
+
+#if defined(CHECK_DECLTYPE)
+#if defined(__CUDA_ARCH__)
+// expected-note@+6 {{other definition of 't0'}}
+// expected-note@+6 {{use this definition of 't0'}}
+#else
+// expected-note@+3 {{use this definition of 't0'}}
+// expected-note@+3 {{other definition of 't0'}}
+#endif
+__host__ float t0();
+__device__ int t0();
+
+__host__ __device__ void dt0() {
+  // expected-warning@+1 {{return type of 't0' in 'decltype' is ambiguous and may not be expected}}
+  decltype(t0()) ret;
+}
+
+__host__ float t1();
+
+__device__ void dt1() {
+  decltype(t1()) ret; // OK. `decltype` is relaxed.
+}
+
+__host__ __device__ void dt2() {
+  decltype(t1()) ret; // OK. `decltype` is relaxed.
+}
+#endif