yaxunl updated this revision to Diff 308032.
yaxunl retitled this revision from "[CUDA][HIP] Fix implicit HD function
resolution" to "[CUDA][HIP] Fix HD function resolution".
yaxunl edited the summary of this revision.
yaxunl added a comment.
If -fgpu-defer-diags is off, keep original behavior.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D80450/new/
https://reviews.llvm.org/D80450
Files:
clang/include/clang/Sema/Overload.h
clang/lib/Sema/SemaOverload.cpp
clang/test/SemaCUDA/deferred-oeverload.cu
clang/test/SemaCUDA/function-overload.cu
Index: clang/test/SemaCUDA/function-overload.cu
===================================================================
--- clang/test/SemaCUDA/function-overload.cu
+++ clang/test/SemaCUDA/function-overload.cu
@@ -1,8 +1,10 @@
// REQUIRES: x86-registered-target
// REQUIRES: nvptx-registered-target
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify=host,expected %s
-// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify=dev,expected %s
+// RUN: %clang_cc1 -std=c++14 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify=host,hostdefer,devdefer,expected %s
+// RUN: %clang_cc1 -std=c++14 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify=dev,devnodeferonly,hostdefer,devdefer,expected %s
+// RUN: %clang_cc1 -fgpu-defer-diag -DDEFER=1 -std=c++14 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify=host,hostdefer,expected %s
+// RUN: %clang_cc1 -fgpu-defer-diag -DDEFER=1 -std=c++14 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify=dev,devdeferonly,devdefer,expected %s
#include "Inputs/cuda.h"
@@ -76,37 +78,37 @@
// Helper functions to verify calling restrictions.
__device__ DeviceReturnTy d() { return DeviceReturnTy(); }
// host-note@-1 1+ {{'d' declared here}}
-// expected-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
+// hostdefer-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
// expected-note@-3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
__host__ HostReturnTy h() { return HostReturnTy(); }
// dev-note@-1 1+ {{'h' declared here}}
-// expected-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
+// devdefer-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
// expected-note@-3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
-// expected-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}
+// devdefer-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}
__global__ void g() {}
// dev-note@-1 1+ {{'g' declared here}}
-// expected-note@-2 1+ {{candidate function not viable: call to __global__ function from __device__ function}}
+// devdefer-note@-2 1+ {{candidate function not viable: call to __global__ function from __device__ function}}
// expected-note@-3 0+ {{candidate function not viable: call to __global__ function from __host__ __device__ function}}
-// expected-note@-4 1+ {{candidate function not viable: call to __global__ function from __global__ function}}
+// devdefer-note@-4 1+ {{candidate function not viable: call to __global__ function from __global__ function}}
extern "C" __device__ DeviceReturnTy cd() { return DeviceReturnTy(); }
// host-note@-1 1+ {{'cd' declared here}}
-// expected-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
+// hostdefer-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
// expected-note@-3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
extern "C" __host__ HostReturnTy ch() { return HostReturnTy(); }
// dev-note@-1 1+ {{'ch' declared here}}
-// expected-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
+// devdefer-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
// expected-note@-3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
-// expected-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}
+// devdefer-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}
__host__ void hostf() {
DeviceFnPtr fp_d = d; // host-error {{reference to __device__ function 'd' in __host__ function}}
- DeviceReturnTy ret_d = d(); // expected-error {{no matching function for call to 'd'}}
+ DeviceReturnTy ret_d = d(); // hostdefer-error {{no matching function for call to 'd'}}
DeviceFnPtr fp_cd = cd; // host-error {{reference to __device__ function 'cd' in __host__ function}}
- DeviceReturnTy ret_cd = cd(); // expected-error {{no matching function for call to 'cd'}}
+ DeviceReturnTy ret_cd = cd(); // hostdefer-error {{no matching function for call to 'cd'}}
HostFnPtr fp_h = h;
HostReturnTy ret_h = h();
@@ -130,9 +132,9 @@
DeviceReturnTy ret_cd = cd();
HostFnPtr fp_h = h; // dev-error {{reference to __host__ function 'h' in __device__ function}}
- HostReturnTy ret_h = h(); // expected-error {{no matching function for call to 'h'}}
+ HostReturnTy ret_h = h(); // devdefer-error {{no matching function for call to 'h'}}
HostFnPtr fp_ch = ch; // dev-error {{reference to __host__ function 'ch' in __device__ function}}
- HostReturnTy ret_ch = ch(); // expected-error {{no matching function for call to 'ch'}}
+ HostReturnTy ret_ch = ch(); // devdefer-error {{no matching function for call to 'ch'}}
DeviceFnPtr fp_dh = dh;
DeviceReturnTy ret_dh = dh();
@@ -140,7 +142,7 @@
DeviceReturnTy ret_cdh = cdh();
GlobalFnPtr fp_g = g; // dev-error {{reference to __global__ function 'g' in __device__ function}}
- g(); // expected-error {{no matching function for call to 'g'}}
+ g(); // devdefer-error {{no matching function for call to 'g'}}
g<<<0,0>>>(); // dev-error {{reference to __global__ function 'g' in __device__ function}}
}
@@ -151,9 +153,9 @@
DeviceReturnTy ret_cd = cd();
HostFnPtr fp_h = h; // dev-error {{reference to __host__ function 'h' in __global__ function}}
- HostReturnTy ret_h = h(); // expected-error {{no matching function for call to 'h'}}
+ HostReturnTy ret_h = h(); // devdefer-error {{no matching function for call to 'h'}}
HostFnPtr fp_ch = ch; // dev-error {{reference to __host__ function 'ch' in __global__ function}}
- HostReturnTy ret_ch = ch(); // expected-error {{no matching function for call to 'ch'}}
+ HostReturnTy ret_ch = ch(); // devdefer-error {{no matching function for call to 'ch'}}
DeviceFnPtr fp_dh = dh;
DeviceReturnTy ret_dh = dh();
@@ -161,7 +163,7 @@
DeviceReturnTy ret_cdh = cdh();
GlobalFnPtr fp_g = g; // dev-error {{reference to __global__ function 'g' in __global__ function}}
- g(); // expected-error {{no matching function for call to 'g'}}
+ g(); // devdefer-error {{no matching function for call to 'g'}}
g<<<0,0>>>(); // dev-error {{reference to __global__ function 'g' in __global__ function}}
}
@@ -184,7 +186,7 @@
#if defined(__CUDA_ARCH__)
// expected-error@-5 {{reference to __host__ function 'h' in __host__ __device__ function}}
// expected-error@-5 {{reference to __host__ function 'h' in __host__ __device__ function}}
- // expected-error@-5 {{reference to __host__ function 'ch' in __host__ __device__ function}}
+ // devdefer-error@-5 {{reference to __host__ function 'ch' in __host__ __device__ function}}
// expected-error@-5 {{reference to __host__ function 'ch' in __host__ __device__ function}}
#endif
@@ -331,9 +333,7 @@
// If we have a mix of HD and H-only or D-only candidates in the overload set,
// normal C++ overload resolution rules apply first.
template <typename T> TemplateReturnTy template_vs_hd_function(T arg)
-#ifdef __CUDA_ARCH__
-//expected-note@-2 {{declared here}}
-#endif
+// devnodeferonly-note@-1{{'template_vs_hd_function<int>' declared here}}
{
return TemplateReturnTy();
}
@@ -342,11 +342,14 @@
}
__host__ __device__ void test_host_device_calls_hd_template() {
- HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
- TemplateReturnTy ret2 = template_vs_hd_function(1);
-#ifdef __CUDA_ARCH__
- // expected-error@-2 {{reference to __host__ function 'template_vs_hd_function<int>' in __host__ __device__ function}}
+#if __CUDA_ARCH__ && DEFER
+ typedef HostDeviceReturnTy ExpectedReturnTy;
+#else
+ typedef TemplateReturnTy ExpectedReturnTy;
#endif
+ HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
+ ExpectedReturnTy ret2 = template_vs_hd_function(1);
+ // devnodeferonly-error@-1{{reference to __host__ function 'template_vs_hd_function<int>' in __host__ __device__ function}}
}
__host__ void test_host_calls_hd_template() {
@@ -367,14 +370,14 @@
__device__ DeviceReturnTy device_only_function(int arg) { return DeviceReturnTy(); }
__device__ DeviceReturnTy2 device_only_function(float arg) { return DeviceReturnTy2(); }
#ifndef __CUDA_ARCH__
- // expected-note@-3 {{'device_only_function' declared here}}
- // expected-note@-3 {{'device_only_function' declared here}}
+ // expected-note@-3 2{{'device_only_function' declared here}}
+ // expected-note@-3 2{{'device_only_function' declared here}}
#endif
__host__ HostReturnTy host_only_function(int arg) { return HostReturnTy(); }
__host__ HostReturnTy2 host_only_function(float arg) { return HostReturnTy2(); }
#ifdef __CUDA_ARCH__
- // expected-note@-3 {{'host_only_function' declared here}}
- // expected-note@-3 {{'host_only_function' declared here}}
+ // expected-note@-3 2{{'host_only_function' declared here}}
+ // expected-note@-3 2{{'host_only_function' declared here}}
#endif
__host__ __device__ void test_host_device_single_side_overloading() {
@@ -392,6 +395,37 @@
#endif
}
+// wrong-sided overloading should not cause diagnostic unless it is emitted.
+// This inline function is not emitted.
+inline __host__ __device__ void test_host_device_wrong_side_overloading_inline_no_diag() {
+ DeviceReturnTy ret1 = device_only_function(1);
+ DeviceReturnTy2 ret2 = device_only_function(1.0f);
+ HostReturnTy ret3 = host_only_function(1);
+ HostReturnTy2 ret4 = host_only_function(1.0f);
+}
+
+// wrong-sided overloading should cause diagnostic if it is emitted.
+// This inline function is emitted since it is called by an emitted function.
+inline __host__ __device__ void test_host_device_wrong_side_overloading_inline_diag() {
+ DeviceReturnTy ret1 = device_only_function(1);
+ DeviceReturnTy2 ret2 = device_only_function(1.0f);
+#ifndef __CUDA_ARCH__
+ // expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}}
+ // expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}}
+#endif
+ HostReturnTy ret3 = host_only_function(1);
+ HostReturnTy2 ret4 = host_only_function(1.0f);
+#ifdef __CUDA_ARCH__
+ // expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}}
+ // expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}}
+#endif
+}
+
+__host__ __device__ void test_host_device_wrong_side_overloading_inline_diag_caller() {
+ test_host_device_wrong_side_overloading_inline_diag();
+ // expected-note@-1 {{called by 'test_host_device_wrong_side_overloading_inline_diag_caller'}}
+}
+
// Verify that we allow overloading function templates.
template <typename T> __host__ T template_overload(const T &a) { return a; };
template <typename T> __device__ T template_overload(const T &a) { return a; };
@@ -419,3 +453,187 @@
int test_constexpr_overload(C2 &x, C2 &y) {
return constexpr_overload(x, y);
}
+
+// Verify no ambiguity for new operator.
+void *a = new int;
+__device__ void *b = new int;
+// expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+// Verify no ambiguity for new operator.
+template<typename _Tp> _Tp&& f();
+template<typename _Tp, typename = decltype(new _Tp(f<_Tp>()))>
+void __test();
+
+void foo() {
+ __test<int>();
+}
+
+// Test resolving implicit host device candidate vs wrong-sided candidate.
+// In device compilation, implicit host device caller choose implicit host
+// device candidate and wrong-sided candidate with equal preference.
+// Resolution result should not change with/without pragma.
+namespace ImplicitHostDeviceVsWrongSided {
+HostReturnTy callee(double x);
+#pragma clang force_cuda_host_device begin
+HostDeviceReturnTy callee(int x);
+inline HostReturnTy implicit_hd_caller() {
+ return callee(1.0);
+}
+#pragma clang force_cuda_host_device end
+}
+
+// Test resolving implicit host device candidate vs same-sided candidate.
+// In host compilation, implicit host device caller choose implicit host
+// device candidate and same-sided candidate with equal preference.
+// Resolution result should not change with/without pragma.
+namespace ImplicitHostDeviceVsSameSide {
+HostReturnTy callee(int x);
+#pragma clang force_cuda_host_device begin
+HostDeviceReturnTy callee(double x);
+inline HostDeviceReturnTy implicit_hd_caller() {
+ return callee(1.0);
+}
+#pragma clang force_cuda_host_device end
+}
+
+// Test resolving explicit host device candidate vs. wrong-sided candidate.
+// When -fgpu-defer-diag is off, wrong-sided candidate is not excluded, therefore
+// the first callee is chosen.
+// When -fgpu-defer-diag is on, wrong-sided candidate is excluded, therefore
+// the second callee is chosen.
+namespace ExplicitHostDeviceVsWrongSided {
+HostReturnTy callee(double x);
+__host__ __device__ HostDeviceReturnTy callee(int x);
+#if __CUDA_ARCH__ && DEFER
+typedef HostDeviceReturnTy ExpectedRetTy;
+#else
+typedef HostReturnTy ExpectedRetTy;
+#endif
+inline __host__ __device__ ExpectedRetTy explicit_hd_caller() {
+ return callee(1.0);
+}
+}
+
+// In the implicit host device function 'caller', the second 'callee' should be
+// chosen since it has better match, even though it is an implicit host device
+// function whereas the first 'callee' is a host function. A diagnostic will be
+// emitted if the first 'callee' is chosen since deduced return type cannot be
+// used before it is defined.
+namespace ImplicitHostDeviceByConstExpr {
+template <class a> a b;
+auto callee(...);
+template <class d> constexpr auto callee(d) -> decltype(0);
+struct e {
+ template <class ad, class... f> static auto g(ad, f...) {
+ return h<e, decltype(b<f>)...>;
+ }
+ struct i {
+ template <class, class... f> static constexpr auto caller(f... k) {
+ return callee(k...);
+ }
+ };
+ template <class, class... f> static auto h() {
+ return i::caller<int, f...>;
+ }
+};
+class l {
+ l() {
+ e::g([] {}, this);
+ }
+};
+}
+
+// Implicit HD candidate competes with device candidate.
+// a and b have implicit HD copy ctor. In copy ctor of b, ctor of a is resolved.
+// copy ctor of a should win over a(short), otherwise there will be ambiguity
+// due to conversion operator.
+namespace TestImplicitHDWithD {
+ struct a {
+ __device__ a(short);
+ __device__ operator unsigned() const;
+ __device__ operator int() const;
+ };
+ struct b {
+ a d;
+ };
+ void f(b g) { b e = g; }
+}
+
+// Implicit HD candidate competes with host candidate.
+// a and b have implicit HD copy ctor. In copy ctor of b, ctor of a is resolved.
+// copy ctor of a should win over a(short), otherwise there will be ambiguity
+// due to conversion operator.
+namespace TestImplicitHDWithH {
+ struct a {
+ a(short);
+ __device__ operator unsigned() const;
+ __device__ operator int() const;
+ };
+ struct b {
+ a d;
+ };
+ void f(b g) { b e = g; }
+}
+
+// Implicit HD candidate competes with HD candidate.
+// a and b have implicit HD copy ctor. In copy ctor of b, ctor of a is resolved.
+// copy ctor of a should win over a(short), otherwise there will be ambiguity
+// due to conversion operator.
+namespace TestImplicitHDWithHD {
+ struct a {
+ __host__ __device__ a(short);
+ __device__ operator unsigned() const;
+ __device__ operator int() const;
+ };
+ struct b {
+ a d;
+ };
+ void f(b g) { b e = g; }
+}
+
+// HD candidate competes with H candidate.
+// HD has type mismatch whereas H has type match.
+// In device compilation, H wins when -fgpu-defer-diag is off and HD wins
+// when -fgpu-defer-diags is on. In both cases the diagnostic should be
+// deferred.
+namespace TestDeferNoMatchingFunc {
+ template <typename> struct a {};
+ namespace b {
+ struct c : a<int> {};
+ template <typename d> void ag(d);
+ } // namespace b
+ template <typename ae>
+ __attribute__((host)) __attribute__((device))
+ void ag(a<ae>) {
+ ae e;
+ ag(e);
+ }
+ void f() { (void)ag<b::c>; }
+}
+
+// Two HD candidates competes with H candidate.
+// HDs have type mismatch whereas H has type match.
+// In device compilation, H wins when -fgpu-defer-diag is off and two HD win
+// when -fgpu-defer-diags is on. In both cases the diagnostic should be
+// deferred.
+namespace TestDeferAmbiguity {
+ template <typename> struct a {};
+ namespace b {
+ struct c : a<int> {};
+ template <typename d> void ag(d, int);
+ } // namespace b
+ template <typename ae>
+ __attribute__((host)) __attribute__((device))
+ void ag(a<ae>, float) {
+ ae e;
+ ag(e, 1);
+ }
+ template <typename ae>
+ __attribute__((host)) __attribute__((device))
+ void ag(a<ae>, double) {
+ }
+ void f() {
+ b::c x;
+ ag(x, 1);
+ }
+}
Index: clang/test/SemaCUDA/deferred-oeverload.cu
===================================================================
--- clang/test/SemaCUDA/deferred-oeverload.cu
+++ clang/test/SemaCUDA/deferred-oeverload.cu
@@ -54,7 +54,7 @@
// This fails to substitue for A but no diagnostic
// should be emitted.
template<typename T, typename T::foo* = nullptr>
-__host__ __device__ void sfinae(T t) { // com-note {{candidate template ignored: substitution failure [with T = B]}}
+__host__ __device__ void sfinae(T t) { // host-note {{candidate template ignored: substitution failure [with T = B]}}
t.x = 1;
}
@@ -64,13 +64,13 @@
// file scope.
template<typename T, typename T::isA* = nullptr>
-__host__ __device__ void sfinae(T t) { // com-note {{candidate template ignored: substitution failure [with T = B]}}
+__host__ __device__ void sfinae(T t) { // host-note {{candidate template ignored: substitution failure [with T = B]}}
t.x = 1;
}
void test_sfinae() {
sfinae(A());
- sfinae(B()); // com-error{{no matching function for call to 'sfinae'}}
+ sfinae(B()); // host-error{{no matching function for call to 'sfinae'}}
}
// Make sure throw is diagnosed in OpenMP parallel region in host function.
Index: clang/lib/Sema/SemaOverload.cpp
===================================================================
--- clang/lib/Sema/SemaOverload.cpp
+++ clang/lib/Sema/SemaOverload.cpp
@@ -9616,6 +9616,68 @@
else if (!Cand1.Viable)
return false;
+ // [CUDA] A function with 'never' preference is marked not viable, therefore
+ // is never shown up here. The worst preference shown up here is 'wrong side',
+ // e.g. a host function called by a device host function in device
+ // compilation. This is valid AST as long as the host device function is not
+ // emitted, e.g. it is an inline function which is called only by a host
+ // function. A deferred diagnostic will be triggered if it is emitted.
+ // However a wrong-sided function is still a viable candidate here.
+ //
+ // If Cand1 can be emitted and Cand2 cannot be emitted in the current
+ // context, Cand1 is better than Cand2. If Cand1 can not be emitted and Cand2
+ // can be emitted, Cand1 is not better than Cand2. This rule should have
+ // precedence over other rules.
+ //
+ // If both Cand1 and Cand2 can be emitted, or neither can be emitted, then
+ // other rules should be used to determine which is better. This is because
+ // host/device based overloading resolution is mostly for determining
+ // viability of a function. If two functions are both viable, other factors
+ // should take precedence in preference, e.g. the standard-defined preferences
+ // like argument conversion ranks or enable_if partial-ordering. The
+ // preference for pass-object-size parameters is probably most similar to a
+ // type-based-overloading decision and so should take priority.
+ //
+ // If other rules cannot determine which is better, CUDA preference will be
+ // used again to determine which is better.
+ //
+ // TODO: Currently IdentifyCUDAPreference does not return correct values
+ // for functions called in global variable initializers due to missing
+ // correct context about device/host. Therefore we can only enforce this
+ // rule when there is a caller. We should enforce this rule for functions
+ // in global variable initializers once proper context is added.
+ if (S.getLangOpts().CUDA && Cand1.Function && Cand2.Function) {
+ if (FunctionDecl *Caller = dyn_cast<FunctionDecl>(S.CurContext)) {
+ bool IsCallerImplicitHD = Sema::isCUDAImplicitHostDeviceFunction(Caller);
+ bool IsCand1ImplicitHD =
+ Sema::isCUDAImplicitHostDeviceFunction(Cand1.Function);
+ bool IsCand2ImplicitHD =
+ Sema::isCUDAImplicitHostDeviceFunction(Cand2.Function);
+ auto P1 = S.IdentifyCUDAPreference(Caller, Cand1.Function);
+ auto P2 = S.IdentifyCUDAPreference(Caller, Cand2.Function);
+ assert(P1 != Sema::CFP_Never && P2 != Sema::CFP_Never);
+ // The implicit HD function may be a function in a system header which
+ // is forced by pragma. In device compilation, if we prefer HD candidates
+ // over wrong-sided candidates, overloading resolution may change, which
+ // may result in non-deferrable diagnostics. As a workaround, we let
+ // implicit HD candidates take equal preference as wrong-sided candidates.
+ // This will preserve the overloading resolution.
+ auto EmitThreshold =
+ (S.getLangOpts().CUDAIsDevice && IsCallerImplicitHD &&
+ (IsCand1ImplicitHD || IsCand2ImplicitHD)) ||
+ (!S.getLangOpts().GPUDeferDiag && P1 < Sema::CFP_SameSide &&
+ P2 < Sema::CFP_SameSide)
+ ? Sema::CFP_Never
+ : Sema::CFP_WrongSide;
+ auto Cand1Emittable = P1 > EmitThreshold;
+ auto Cand2Emittable = P2 > EmitThreshold;
+ if (Cand1Emittable && !Cand2Emittable)
+ return true;
+ if (!Cand1Emittable && Cand2Emittable)
+ return false;
+ }
+ }
+
// C++ [over.match.best]p1:
//
// -- if F is a static member function, ICS1(F) is defined such
@@ -9850,12 +9912,6 @@
return Cmp == Comparison::Better;
}
- if (S.getLangOpts().CUDA && Cand1.Function && Cand2.Function) {
- FunctionDecl *Caller = dyn_cast<FunctionDecl>(S.CurContext);
- return S.IdentifyCUDAPreference(Caller, Cand1.Function) >
- S.IdentifyCUDAPreference(Caller, Cand2.Function);
- }
-
bool HasPS1 = Cand1.Function != nullptr &&
functionHasPassObjectSizeParams(Cand1.Function);
bool HasPS2 = Cand2.Function != nullptr &&
@@ -9863,8 +9919,21 @@
if (HasPS1 != HasPS2 && HasPS1)
return true;
- Comparison MV = isBetterMultiversionCandidate(Cand1, Cand2);
- return MV == Comparison::Better;
+ auto MV = isBetterMultiversionCandidate(Cand1, Cand2);
+ if (MV == Comparison::Better)
+ return true;
+ if (MV == Comparison::Worse)
+ return false;
+
+ // If other rules cannot determine which is better, CUDA preference is used
+ // to determine which is better.
+ if (S.getLangOpts().CUDA && Cand1.Function && Cand2.Function) {
+ FunctionDecl *Caller = dyn_cast<FunctionDecl>(S.CurContext);
+ return S.IdentifyCUDAPreference(Caller, Cand1.Function) >
+ S.IdentifyCUDAPreference(Caller, Cand2.Function);
+ }
+
+ return false;
}
/// Determine whether two declarations are "equivalent" for the purposes of
@@ -9950,33 +10019,6 @@
std::transform(begin(), end(), std::back_inserter(Candidates),
[](OverloadCandidate &Cand) { return &Cand; });
- // [CUDA] HD->H or HD->D calls are technically not allowed by CUDA but
- // are accepted by both clang and NVCC. However, during a particular
- // compilation mode only one call variant is viable. We need to
- // exclude non-viable overload candidates from consideration based
- // only on their host/device attributes. Specifically, if one
- // candidate call is WrongSide and the other is SameSide, we ignore
- // the WrongSide candidate.
- if (S.getLangOpts().CUDA) {
- const FunctionDecl *Caller = dyn_cast<FunctionDecl>(S.CurContext);
- bool ContainsSameSideCandidate =
- llvm::any_of(Candidates, [&](OverloadCandidate *Cand) {
- // Check viable function only.
- return Cand->Viable && Cand->Function &&
- S.IdentifyCUDAPreference(Caller, Cand->Function) ==
- Sema::CFP_SameSide;
- });
- if (ContainsSameSideCandidate) {
- auto IsWrongSideCandidate = [&](OverloadCandidate *Cand) {
- // Check viable function only to avoid unnecessary data copying/moving.
- return Cand->Viable && Cand->Function &&
- S.IdentifyCUDAPreference(Caller, Cand->Function) ==
- Sema::CFP_WrongSide;
- };
- llvm::erase_if(Candidates, IsWrongSideCandidate);
- }
- }
-
// Find the best viable function.
Best = end();
for (auto *Cand : Candidates) {
@@ -11620,26 +11662,34 @@
return Cands;
}
-/// When overload resolution fails, prints diagnostic messages containing the
-/// candidates in the candidate set.
-void OverloadCandidateSet::NoteCandidates(PartialDiagnosticAt PD,
- Sema &S, OverloadCandidateDisplayKind OCD, ArrayRef<Expr *> Args,
- StringRef Opc, SourceLocation OpLoc,
- llvm::function_ref<bool(OverloadCandidate &)> Filter) {
-
+bool OverloadCandidateSet::shouldDeferDiags(Sema &S, ArrayRef<Expr *> Args,
+ SourceLocation OpLoc) {
bool DeferHint = false;
if (S.getLangOpts().CUDA && S.getLangOpts().GPUDeferDiag) {
- // Defer diagnostic for CUDA/HIP if there are wrong-sided candidates.
+ // Defer diagnostic for CUDA/HIP if there are wrong-sided candidates or
+ // host device candidates.
auto WrongSidedCands =
CompleteCandidates(S, OCD_AllCandidates, Args, OpLoc, [](auto &Cand) {
- return Cand.Viable == false &&
- Cand.FailureKind == ovl_fail_bad_target;
+ return (Cand.Viable == false &&
+ Cand.FailureKind == ovl_fail_bad_target) ||
+ (Cand.Function->template hasAttr<CUDAHostAttr>() &&
+ Cand.Function->template hasAttr<CUDADeviceAttr>());
});
DeferHint = WrongSidedCands.size();
}
+ return DeferHint;
+}
+
+/// When overload resolution fails, prints diagnostic messages containing the
+/// candidates in the candidate set.
+void OverloadCandidateSet::NoteCandidates(
+ PartialDiagnosticAt PD, Sema &S, OverloadCandidateDisplayKind OCD,
+ ArrayRef<Expr *> Args, StringRef Opc, SourceLocation OpLoc,
+ llvm::function_ref<bool(OverloadCandidate &)> Filter) {
+
auto Cands = CompleteCandidates(S, OCD, Args, OpLoc, Filter);
- S.Diag(PD.first, PD.second, DeferHint);
+ S.Diag(PD.first, PD.second, shouldDeferDiags(S, Args, OpLoc));
NoteCandidates(S, Args, Cands, Opc, OpLoc);
@@ -11691,7 +11741,9 @@
}
if (I != E)
- S.Diag(OpLoc, diag::note_ovl_too_many_candidates) << int(E - I);
+ S.Diag(OpLoc, diag::note_ovl_too_many_candidates,
+ shouldDeferDiags(S, Args, OpLoc))
+ << int(E - I);
}
static SourceLocation
Index: clang/include/clang/Sema/Overload.h
===================================================================
--- clang/include/clang/Sema/Overload.h
+++ clang/include/clang/Sema/Overload.h
@@ -1051,6 +1051,9 @@
void destroyCandidates();
+ /// Whether diagnostics should be deferred.
+ bool shouldDeferDiags(Sema &S, ArrayRef<Expr *> Args, SourceLocation OpLoc);
+
public:
OverloadCandidateSet(SourceLocation Loc, CandidateSetKind CSK,
OperatorRewriteInfo RewriteInfo = {})
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
