[CUDA][HIP] Fix overloading resolution in global variable initializer

[yaxunl]

Summary

Currently, clang does not resolve certain overloaded functions correctly in the initializer
of global variables, e.g.

template<typename T1, typename U>
T1 mypow(T1, U);

__attribute__((device)) double mypow(double, int);

double t_extent = mypow(1.0, 2);

In the above example, mypow is supposed to resolve to the host version
but clang resolves it to the device version instead, and emits an error
(https://godbolt.org/z/17xxzaa67).

However, if the variable is assigned in a host function, there is no error.
The discrepancy in overloading resolution inside and outside of
a function is due to clang not accounting for the host/device target
when resolving functions called in the initializer of a global variable.

This patch introduces a global host/device target context for CUDA/HIP
for functions called outside of functions. For global variable initialization,
it is determined by the host/device attribute of the variable. For other
situations, a default value of host_device is sufficient.

[tra]

Same reproducer but for CUDA: https://godbolt.org/z/WhjTMffnx

clang/include/clang/Sema/Sema.h

4753

It appears that Declarator D here is only used as an attribute carrier used to identify CUDA calling target.
Should we pass CudaTarget ContextTarget instead and let the caller figure out how to find it?

I'm just thinking that we're hardcoding just one specific way to find the target, while there may potentially be more.
The current way is OK, as we have just one use case at the moment.

clang/lib/Sema/SemaCUDA.cpp

137

Style nit: no braces around single-statement body.

clang/test/CodeGenCUDA/global-initializers.cu

11–12

We don't really need templates to reproduce the issue. We just need a host function with lower overloading priority. A function requiring type conversion or with an additional default argument should do. E.g. float pow(float, int); or double X = pow(double, int, bool lower_priority_host_overload=1);

Removing template should unclutter the tests a bit.

[yaxunl]

clang/include/clang/Sema/Sema.h

4753

will do

clang/lib/Sema/SemaCUDA.cpp

137

will fix

clang/test/CodeGenCUDA/global-initializers.cu

11–12

will do

[yaxunl]

revised by comments

[yaxunl]

ping

[yaxunl]

The patch was reverted since it caused regressions on Windows for HIP. A reduced test case is:

typedef void (__stdcall* funcTy)();
void invoke(funcTy f);

static void __stdcall callee() noexcept {
}

void foo() {
   invoke(callee);
}

It is due to clang missed handling host/device attributes for calling convention at a few places.

Reopen to fix it.

[yaxunl]

Phabricator no longer allows me to update the patch. Created PR in github https://github.com/llvm/llvm-project/pull/65606

clang/include/clang/Sema/Sema.h

public:
 
   /// Determine if type T is a valid subject for a nonnull and similar
   /// attributes. By default, we look through references (the behavior used by
   /// nonnull), but if the second parameter is true, then we treat a reference
   /// type as valid.
   bool isValidPointerAttrType(QualType T, bool RefOkay = false);
 
   bool CheckRegparmAttr(const ParsedAttr &attr, unsigned &value);
+
+  /// Check validaty of calling convention attribute \p attr. If \p FD
+  /// is not null pointer, use \p FD to determine the host/device target.
+  /// Otherwise, use \p D to determiine the host/device target.
   bool CheckCallingConvAttr(const ParsedAttr &attr, CallingConv &CC,
-                            const FunctionDecl *FD = nullptr);
+                            const FunctionDecl *FD = nullptr,
+                            const Declarator *D = nullptr);
   bool CheckAttrTarget(const ParsedAttr &CurrAttr);
   bool CheckAttrNoArgs(const ParsedAttr &CurrAttr);
   bool checkStringLiteralArgumentAttr(const AttributeCommonInfo &CI,
                                       const Expr *E, StringRef &Str,
                                       SourceLocation *ArgLocation = nullptr);
   bool checkStringLiteralArgumentAttr(const ParsedAttr &Attr, unsigned ArgNum,
                                       StringRef &Str,
                                       SourceLocation *ArgLocation = nullptr);
enum CUDAVariableTarget {
     CVT_Device,  /// Emitted on device side with a shadow variable on host side
     CVT_Host,    /// Emitted on host side only
     CVT_Both,    /// Emitted on both sides with different addresses
     CVT_Unified, /// Emitted as a unified address, e.g. managed variables
   };
   /// Determines whether the given variable is emitted on host or device side.
   CUDAVariableTarget IdentifyCUDATarget(const VarDecl *D);
 
+  /// Defines different CUDA host/device contexts where a function may be
+  /// called.
+  enum CUDATargetContextKind {
+    CTCK_Unknown,       /// Unknown context
+    CTCK_InitGlobalVar, /// Function called during global variable
+                        /// initialization
+  };
+
+  struct CUDATargetContext {
+    CUDAFunctionTarget Target = CFT_HostDevice;
+    CUDATargetContextKind Kind = CTCK_Unknown;
+    Decl *D = nullptr;
+  } CurCUDATargetCtx;
+
+  struct CUDATargetContextRAII {
+    Sema &S;
+    CUDATargetContext SavedCtx;
+    CUDATargetContextRAII(Sema &S_, CUDATargetContextKind K, Decl *D);
+    ~CUDATargetContextRAII() { S.CurCUDATargetCtx = SavedCtx; }
+  };
+
   /// Gets the CUDA target for the current context.
   CUDAFunctionTarget CurrentCUDATarget() {
     return IdentifyCUDATarget(dyn_cast<FunctionDecl>(CurContext));
   }
 
   static bool isCUDAImplicitHostDeviceFunction(const FunctionDecl *D);
 
   // CUDA function call preference. Must be ordered numerically from

clang/lib/Parse/ParseDecl.cpp

if (Tok.is(tok::semi)) {
         ThisDecl =
             Actions.ActOnTemplateDeclarator(getCurScope(), FakedParamLists, D);
       }
     }
     break;
     }
   }
 
+  Sema::CUDATargetContextRAII X(Actions, Sema::CTCK_InitGlobalVar, ThisDecl);
   switch (TheInitKind) {
   // Parse declarator '=' initializer.
   case InitKind::Equal: {
     SourceLocation EqualLoc = ConsumeToken();
 
     if (Tok.is(tok::kw_delete)) {
       if (D.isFunctionDeclarator())
         Diag(ConsumeToken(), diag::err_default_delete_in_multiple_declaration)

clang/lib/Sema/SemaCUDA.cpp

Sema::IdentifyCUDATarget(const ParsedAttributesView &Attrs) {
 
   if (HasDeviceAttr)
     return CFT_Device;
 
   return CFT_Host;
 }
 
 template <typename A>
-static bool hasAttr(const FunctionDecl *D, bool IgnoreImplicitAttr) {
+static bool hasAttr(const Decl *D, bool IgnoreImplicitAttr) {
   return D->hasAttrs() && llvm::any_of(D->getAttrs(), [&](Attr *Attribute) {
            return isa<A>(Attribute) &&
                   !(IgnoreImplicitAttr && Attribute->isImplicit());
          });
 }
 
+Sema::CUDATargetContextRAII::CUDATargetContextRAII(Sema &S_,
+                                                   CUDATargetContextKind K,
+                                                   Decl *D)
+    : S(S_) {
+  SavedCtx = S.CurCUDATargetCtx;
+  assert(K == CTCK_InitGlobalVar);
+  auto *VD = dyn_cast_or_null<VarDecl>(D);
+  if (VD && VD->hasGlobalStorage() && !VD->isStaticLocal()) {
+    auto Target = CFT_Host;
+    if ((hasAttr<CUDADeviceAttr>(VD, /*IgnoreImplicit=*/true) &&
+         !hasAttr<CUDAHostAttr>(VD, /*IgnoreImplicit=*/true)) ||
+        hasAttr<CUDASharedAttr>(VD, /*IgnoreImplicit=*/true) ||
+        hasAttr<CUDAConstantAttr>(VD, /*IgnoreImplicit=*/true))
+      Target = CFT_Device;
+    S.CurCUDATargetCtx = {Target, K, VD};
+  }
+}
+
 /// IdentifyCUDATarget - Determine the CUDA compilation target for this function
 Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D,
                                                   bool IgnoreImplicitHDAttr) {
-  // Code that lives outside a function is run on the host.
-  if (D == nullptr)
-    return CFT_Host;
+  // Code that lives outside a function gets the target from CurCUDATargetCtx.
+  if (D == nullptr) {
+    return CurCUDATargetCtx.Target;
+  }
 
   if (D->hasAttr<CUDAInvalidTargetAttr>())
     return CFT_InvalidTarget;
 
   if (D->hasAttr<CUDAGlobalAttr>())
     return CFT_Global;
 
   if (hasAttr<CUDADeviceAttr>(D, IgnoreImplicitHDAttr)) {

clang/lib/Sema/SemaDeclAttr.cpp

static void handleNoRandomizeLayoutAttr(Sema &S, Decl *D,
                                         const ParsedAttr &AL) {
   if (checkAttrMutualExclusion<RandomizeLayoutAttr>(S, D, AL))
     return;
   if (!D->hasAttr<NoRandomizeLayoutAttr>())
     D->addAttr(::new (S.Context) NoRandomizeLayoutAttr(S.Context, AL));
 }
 
 bool Sema::CheckCallingConvAttr(const ParsedAttr &Attrs, CallingConv &CC,
-                                const FunctionDecl *FD) {
+                                const FunctionDecl *FD, const Declarator *D) {
   if (Attrs.isInvalid())
     return true;
 
   if (Attrs.hasProcessingCache()) {
     CC = (CallingConv) Attrs.getProcessingCache();
     return false;
   }
 
bool Sema::CheckCallingConvAttr(const ParsedAttr &Attrs, CallingConv &CC,
   TargetInfo::CallingConvCheckResult A = TargetInfo::CCCR_OK;
   const TargetInfo &TI = Context.getTargetInfo();
   // CUDA functions may have host and/or device attributes which indicate
   // their targeted execution environment, therefore the calling convention
   // of functions in CUDA should be checked against the target deduced based
   // on their host/device attributes.
   if (LangOpts.CUDA) {
     auto *Aux = Context.getAuxTargetInfo();
-    auto CudaTarget = IdentifyCUDATarget(FD);
+    assert(FD || D);
+    auto CudaTarget =
+        FD ? IdentifyCUDATarget(FD) : IdentifyCUDATarget(D->getAttributes());
     bool CheckHost = false, CheckDevice = false;
     switch (CudaTarget) {
     case CFT_HostDevice:
       CheckHost = true;
       CheckDevice = true;
       break;
     case CFT_Host:
       CheckHost = true;

clang/lib/Sema/SemaOverload.cpp

if (!AggregateCandidateDeduction && Args.size() < MinRequiredArgs &&
       !PartialOverloading) {
     // Not enough arguments.
     Candidate.Viable = false;
     Candidate.FailureKind = ovl_fail_too_few_arguments;
     return;
   }
 
   // (CUDA B.1): Check for invalid calls between targets.
-  if (getLangOpts().CUDA)
-    if (const FunctionDecl *Caller = getCurFunctionDecl(/*AllowLambda=*/true))
-      // 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)) {
-        Candidate.Viable = false;
-        Candidate.FailureKind = ovl_fail_bad_target;
-        return;
-      }
+  if (getLangOpts().CUDA) {
+    const FunctionDecl *Caller = getCurFunctionDecl(/*AllowLambda=*/true);
+    if (!(Caller && Caller->isImplicit()) &&
+        !IsAllowedCUDACall(Caller, Function)) {
+  }
 
   if (Function->getTrailingRequiresClause()) {
     ConstraintSatisfaction Satisfaction;
     if (CheckFunctionConstraints(Function, Satisfaction, /*Loc*/ {},
                                  /*ForOverloadResolution*/ true) ||
         !Satisfaction.IsSatisfied) {
       Candidate.Viable = false;
       Candidate.FailureKind = ovl_fail_constraints_not_satisfied;
if (Candidate.Conversions[FirstConvIdx].isBad()) {
       Candidate.Viable = false;
       Candidate.FailureKind = ovl_fail_bad_conversion;
       return;
     }
   }
 
   // (CUDA B.1): Check for invalid calls between targets.
   if (getLangOpts().CUDA)
-    if (const FunctionDecl *Caller = getCurFunctionDecl(/*AllowLambda=*/true))
-      if (!IsAllowedCUDACall(Caller, Method)) {
-        Candidate.Viable = false;
-        Candidate.FailureKind = ovl_fail_bad_target;
-        return;
-      }
+    if (!IsAllowedCUDACall(getCurFunctionDecl(/*AllowLambda=*/true), Method)) {
 
   if (Method->getTrailingRequiresClause()) {
     ConstraintSatisfaction Satisfaction;
     if (CheckFunctionConstraints(Method, Satisfaction, /*Loc*/ {},
                                  /*ForOverloadResolution*/ true) ||
         !Satisfaction.IsSatisfied) {
       Candidate.Viable = false;
       Candidate.FailureKind = ovl_fail_constraints_not_satisfied;
if (CXXMethodDecl *Method = dyn_cast<CXXMethodDecl>(Fn)) {
       // when converting to member pointer.
       if (Method->isStatic() == TargetTypeIsNonStaticMemberFunction)
         return false;
     }
     else if (TargetTypeIsNonStaticMemberFunction)
       return false;
 
     if (FunctionDecl *FunDecl = dyn_cast<FunctionDecl>(Fn)) {
-      if (S.getLangOpts().CUDA)
-        if (FunctionDecl *Caller = S.getCurFunctionDecl(/*AllowLambda=*/true))
-          if (!Caller->isImplicit() && !S.IsAllowedCUDACall(Caller, FunDecl))
-            return false;
+      if (S.getLangOpts().CUDA) {
+        FunctionDecl *Caller = S.getCurFunctionDecl(/*AllowLambda=*/true);
+        if (!(Caller && Caller->isImplicit()) &&
+            !S.IsAllowedCUDACall(Caller, FunDecl))
+      }
       if (FunDecl->isMultiVersion()) {
         const auto *TA = FunDecl->getAttr<TargetAttr>();
         if (TA && !TA->isDefaultVersion())
           return false;
         const auto *TVA = FunDecl->getAttr<TargetVersionAttr>();
         if (TVA && !TVA->isDefaultVersion())
           return false;
       }

clang/lib/Sema/SemaType.cpp

static CallingConv getCCForDeclaratorChunk(
   // Check for an explicit CC attribute.
   for (const ParsedAttr &AL : AttrList) {
     switch (AL.getKind()) {
     CALLING_CONV_ATTRS_CASELIST : {
       // Ignore attributes that don't validate or can't apply to the
       // function type.  We'll diagnose the failure to apply them in
       // handleFunctionTypeAttr.
       CallingConv CC;
-      if (!S.CheckCallingConvAttr(AL, CC) &&
+      if (!S.CheckCallingConvAttr(AL, CC, /*FunctionDecl=*/nullptr, &D) &&
           (!FTI.isVariadic || supportsVariadicCall(CC))) {
         return CC;
       }
       break;
     }
 
     default:
       break;

clang/test/CodeGenCUDA/global-initializers.cu

+// RUN: %clang_cc1 %s -triple x86_64-linux-unknown -emit-llvm -o - \
+// RUN:   | FileCheck -check-prefix=HOST %s
+// RUN: %clang_cc1 %s -fcuda-is-device \
+// RUN:   -emit-llvm -o - -triple nvptx64 \
+// RUN:   -aux-triple x86_64-unknown-linux-gnu | FileCheck \
+// RUN:   -check-prefix=DEV %s
+
+#include "Inputs/cuda.h"
+
+// Check host/device-based overloding resolution in global variable initializer.
+template<typename T, typename U>
+T pow(T, U) { return 1.0; }
+
+__device__ double pow(double, int) { return 2.0; }
+
+// HOST-DAG: call {{.*}}double @_Z3powIdiET_S0_T0_(double noundef 1.000000e+00, i32 noundef 1)
+double X = pow(1.0, 1);
+
+template<typename T, typename U>
+constexpr T cpow(T, U) { return 11.0; }
+
+constexpr __device__ double cpow(double, int) { return 12.0; }
+
+// HOST-DAG: @CX = global double 1.100000e+01
+double CX = cpow(11.0, 1);
+
+// DEV-DAG: @CY = addrspace(1) externally_initialized global double 1.200000e+01
+__device__ double CY = cpow(12.0, 1);
+
+struct A {
+  template<typename T, typename U>
+  T pow(T, U) { return 3.0; }
+
+  __device__ double pow(double, int) { return 4.0; }
+};
+
+A a;
+
+// HOST-DAG: call {{.*}}double @_ZN1A3powIdiEET_S1_T0_(ptr {{.*}}@a, double noundef 3.000000e+00, i32 noundef 1)
+double AX = a.pow(3.0, 1);
+
+struct CA {
+  template<typename T, typename U>
+  constexpr T cpow(T, U) const { return 13.0; }
+
+  constexpr __device__ double cpow(double, int) const { return 14.0; }
+};
+
+const CA ca;
+
+// HOST-DAG: @CAX = global double 1.300000e+01
+double CAX = ca.cpow(13.0, 1);
+
+// DEV-DAG: @CAY = addrspace(1) externally_initialized global double 1.400000e+01
+__device__ double CAY = ca.cpow(14.0, 1);

clang/test/SemaCUDA/amdgpu-windows-vectorcall.cu

 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -aux-triple x86_64-pc-windows-msvc -fms-compatibility -fcuda-is-device -fsyntax-only -verify %s
+// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fms-compatibility -fsyntax-only -verify %s
 
 __cdecl void hostf1();
 __vectorcall void (*hostf2)() = hostf1; // expected-error {{cannot initialize a variable of type 'void ((*))() __attribute__((vectorcall))' with an lvalue of type 'void () __attribute__((cdecl))'}}

clang/test/SemaCUDA/function-overload.cu

 #if defined(__CUDA_ARCH__)
   // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
 #endif
 }
 
 // Test for address of overloaded function resolution in the global context.
 HostFnPtr fp_h = h;
 HostFnPtr fp_ch = ch;
+#if defined (__CUDA_ARCH__)
+__device__
+#endif
 CurrentFnPtr fp_dh = dh;
+#if defined (__CUDA_ARCH__)
+__device__
+#endif
 CurrentFnPtr fp_cdh = cdh;
 GlobalFnPtr fp_g = g;
 
 
 // Test overloading of destructors
 // Can't mix H and unattributed destructors
 struct d_h {
   ~d_h() {} // expected-note {{previous definition is here}}

clang/test/SemaCUDA/global-initializers-host.cu

-// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-linux-unknown -fsyntax-only -o - -verify
-
-#include "Inputs/cuda.h"
-
-// Check that we get an error if we try to call a __device__ function from a
-// module initializer.
-
-struct S {
-  __device__ S() {}
-  // expected-note@-1 {{'S' declared here}}
-};
-
-S s;
-// expected-error@-1 {{reference to __device__ function 'S' in global initializer}}
-
-struct T {
-  __host__ __device__ T() {}
-};
-T t;  // No error, this is OK.
-
-struct U {
-  __host__ U() {}
-  __device__ U(int) {}
-  // expected-note@-1 {{'U' declared here}}
-};
-U u(42);
-// expected-error@-1 {{reference to __device__ function 'U' in global initializer}}
-
-__device__ int device_fn() { return 42; }
-// expected-note@-1 {{'device_fn' declared here}}
-int n = device_fn();
-// expected-error@-1 {{reference to __device__ function 'device_fn' in global initializer}}

clang/test/SemaCUDA/global-initializers.cu

+// RUN: %clang_cc1 %s -triple x86_64-linux-unknown -fsyntax-only -o - -verify
+// RUN: %clang_cc1 %s -fcuda-is-device -triple nvptx -fsyntax-only -o - -verify
+
+#include "Inputs/cuda.h"
+
+// Check that we get an error if we try to call a __device__ function from a
+// module initializer.
+
+struct S {
+  // expected-note@-1 {{candidate constructor (the implicit copy constructor) not viable: requires 1 argument, but 0 were provided}}
+  // expected-note@-2 {{candidate constructor (the implicit move constructor) not viable: requires 1 argument, but 0 were provided}}
+  __device__ S() {}
+  // expected-note@-1 {{candidate constructor not viable: call to __device__ function from __host__ function}}
+};
+
+S s;
+// expected-error@-1 {{no matching constructor for initialization of 'S'}}
+
+struct T {
+  __host__ __device__ T() {}
+};
+T t;  // No error, this is OK.
+
+struct U {
+  // expected-note@-1 {{candidate constructor (the implicit copy constructor) not viable: no known conversion from 'int' to 'const U' for 1st argument}}
+  // expected-note@-2 {{candidate constructor (the implicit move constructor) not viable: no known conversion from 'int' to 'U' for 1st argument}}
+  __host__ U() {}
+  // expected-note@-1 {{candidate constructor not viable: requires 0 arguments, but 1 was provided}}
+  __device__ U(int) {}
+  // expected-note@-1 {{candidate constructor not viable: call to __device__ function from __host__ function}}
+};
+U u(42);
+// expected-error@-1 {{no matching constructor for initialization of 'U'}}
+
+__device__ int device_fn() { return 42; }
+// expected-note@-1 {{candidate function not viable: call to __device__ function from __host__ function}}
+int n = device_fn();
+// expected-error@-1 {{no matching function for call to 'device_fn'}}
+
+// Check host/device-based overloding resolution in global variable initializer.
+template<typename T, typename U>
+T pow(T, U);
+
+__device__ double pow(double, int);
+
+double X = pow(1.0, 1);
+__device__ double Y = pow(2.0, 2); // expected-error{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+
+template<typename T, typename U>
+constexpr T cpow(T, U) { return 1.0; }
+
+constexpr __device__ double cpow(double, int) { return 2.0; }
+
+const double CX = cpow(1.0, 1);
+const __device__ double CY = cpow(2.0, 2);
+
+struct A {
+  template<typename T, typename U>
+  T pow(T, U);
+
+  __device__ double pow(double, int);
+
+  template<typename T, typename U>
+  constexpr T cpow(T, U) const { return 1.0; }
+
+  constexpr __device__ double cpow(double, int) const { return 1.0; }
+
+};
+
+A a;
+double AX = a.pow(1.0, 1);
+__device__ double AY = a.pow(2.0, 2); // expected-error{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+
+const A ca;
+const double CAX = ca.cpow(1.0, 1);
+const __device__ double CAY = ca.cpow(2.0, 2);