[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

10411

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

10411

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

10407–10409

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

10407–10409

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

10407–10409

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/Sema/Sema.h

public:
   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.
   ///
   /// 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) {
+    if (llvm::any_of(ExprEvalContexts,
+                     [](const ExpressionEvaluationContextRecord &C) {
+                       return C.ExprContext ==
+                              ExpressionEvaluationContextRecord::EK_Decltype;
+                     }))
+      return true;
     return IdentifyCUDAPreference(Caller, 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);
 

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