[CUDA] Pass ExecConfig through BuildCallToMemberFunction

[tra]

Summary

Otherwise, we fail to compile calls to CUDA kernels that are static members.

[jlebar]

Thanks, Art.

[yaxunl]

need a test for non-template static member function as kernel. also need codegen tests.

[tra]

Added more tests

[tra]

In D108787#2968503, @yaxunl wrote:

need a test for non-template static member function as kernel. also need codegen tests.

I've added more tests for different code paths leading to the kernel call. Interestingly enough, only the a0 actually calls BuildCallToMemberFunction. Other variants go through different code paths that handle the call without it.

As for the codegen, all kernel calls that make it to clang::Sema::BuildResolvedCallExpr with launch config are handled the same way. I don't think codegen tests will buy us much.

[yaxunl]

I am concerned that there may be more places which need handling, and passing exec config expr by function arguments may not scale. Is it possible to represent the kernel call expr by a derived class of call expr and add the exec config expr as member to it?

[tra]

In D108787#2999943, @yaxunl wrote:

I am concerned that there may be more places which need handling, and passing exec config expr by function arguments may not scale.
Is it possible to represent the kernel call expr by a derived class of call expr and add the exec config expr as member to it?

I don't think it's worth it.

This config pass-through code has been around from the very early days of attempting to implement CUDA and we're already passing it around during call resolution.
AFAICT, this particular place was a relatively new addition which didn't implement the pass-through of the config.

While there may be other places where a similar issue may happen in the future (or exists as a corner case we didn't find yet), it/when we run into it, it will be diagnosed, as it was in this case.
It took us few years until we ran into this one. I'm pretty sure that this particular code path is pretty rare and the patch is not going to have a measurable impact on compiler performance.

[yaxunl]

LGTM. Thanks.

clang/include/clang/Sema/Sema.h

 //===--- Sema.h - Semantic Analysis & AST Building --------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file defines the Sema class, which performs semantic analysis and
ExprResult BuildSynthesizedThreeWayComparison(SourceLocation OpLoc,
                                                 const UnresolvedSetImpl &Fns,
                                                 Expr *LHS, Expr *RHS,
                                                 FunctionDecl *DefaultedFn);
 
   ExprResult CreateOverloadedArraySubscriptExpr(SourceLocation LLoc,
                                                 SourceLocation RLoc,
                                                 Expr *Base,Expr *Idx);
 
   ExprResult BuildCallToMemberFunction(Scope *S, Expr *MemExpr,
                                        SourceLocation LParenLoc,
                                        MultiExprArg Args,
                                        SourceLocation RParenLoc,
+                                       Expr *ExecConfig = nullptr,
+                                       bool IsExecConfig = false,
                                        bool AllowRecovery = false);
   ExprResult
   BuildCallToObjectOfClassType(Scope *S, Expr *Object, SourceLocation LParenLoc,
                                MultiExprArg Args,
                                SourceLocation RParenLoc);
 
   ExprResult BuildOverloadedArrowExpr(Scope *S, Expr *Base,
                                       SourceLocation OpLoc,

clang/lib/Sema/SemaExpr.cpp

if (getLangOpts().CPlusPlus) {
     if (Fn->getType() == Context.UnknownAnyTy) {
       ExprResult result = rebuildUnknownAnyFunction(*this, Fn);
       if (result.isInvalid()) return ExprError();
       Fn = result.get();
     }
 
     if (Fn->getType() == Context.BoundMemberTy) {
       return BuildCallToMemberFunction(Scope, Fn, LParenLoc, ArgExprs,
-                                       RParenLoc, AllowRecovery);
+                                       RParenLoc, ExecConfig, IsExecConfig,
+                                       AllowRecovery);
     }
   }
 
   // Check for overloaded calls.  This can happen even in C due to extensions.
   if (Fn->getType() == Context.OverloadTy) {
     OverloadExpr::FindResult find = OverloadExpr::find(Fn);
 
     // We aren't supposed to apply this logic if there's an '&' involved.
     if (!find.HasFormOfMemberPointer) {
       if (Expr::hasAnyTypeDependentArguments(ArgExprs))
         return CallExpr::Create(Context, Fn, ArgExprs, Context.DependentTy,
                                 VK_PRValue, RParenLoc, CurFPFeatureOverrides());
       OverloadExpr *ovl = find.Expression;
       if (UnresolvedLookupExpr *ULE = dyn_cast<UnresolvedLookupExpr>(ovl))
         return BuildOverloadedCallExpr(
             Scope, Fn, ULE, LParenLoc, ArgExprs, RParenLoc, ExecConfig,
             /*AllowTypoCorrection=*/true, find.IsAddressOfOperand);
       return BuildCallToMemberFunction(Scope, Fn, LParenLoc, ArgExprs,
-                                       RParenLoc, AllowRecovery);
+                                       RParenLoc, ExecConfig, IsExecConfig,
+                                       AllowRecovery);
     }
   }
 
   // If we're directly calling a function, get the appropriate declaration.
   if (Fn->getType() == Context.UnknownAnyTy) {
     ExprResult result = rebuildUnknownAnyFunction(*this, Fn);
     if (result.isInvalid()) return ExprError();
     Fn = result.get();

clang/lib/Sema/SemaOverload.cpp

 /// arguments to the function call (not including the object
 /// parameter). The caller needs to validate that the member
 /// expression refers to a non-static member function or an overloaded
 /// member function.
 ExprResult Sema::BuildCallToMemberFunction(Scope *S, Expr *MemExprE,
                                            SourceLocation LParenLoc,
                                            MultiExprArg Args,
                                            SourceLocation RParenLoc,
+                                           Expr *ExecConfig, bool IsExecConfig,
                                            bool AllowRecovery) {
   assert(MemExprE->getType() == Context.BoundMemberTy ||
          MemExprE->getType() == Context.OverloadTy);
 
   // Dig out the member expression. This holds both the object
   // argument and the member function we're referring to.
   Expr *NakedMemExpr = MemExprE->IgnoreParens();
 
if (isa<MemberExpr>(NakedMemExpr)) {
     if (!Succeeded)
       return BuildRecoveryExpr(chooseRecoveryType(CandidateSet, &Best));
 
     MemExprE = FixOverloadedFunctionReference(MemExprE, FoundDecl, Method);
 
     // If overload resolution picked a static member, build a
     // non-member call based on that function.
     if (Method->isStatic()) {
-      return BuildResolvedCallExpr(MemExprE, Method, LParenLoc, Args,
-                                   RParenLoc);
+      return BuildResolvedCallExpr(MemExprE, Method, LParenLoc, Args, RParenLoc,
+                                   ExecConfig, IsExecConfig);
     }
 
     MemExpr = cast<MemberExpr>(MemExprE->IgnoreParens());
   }
 
   QualType ResultType = Method->getReturnType();
   ExprValueKind VK = Expr::getValueKindForType(ResultType);
   ResultType = ResultType.getNonLValueExprType(Context);

clang/test/SemaCUDA/kernel-call.cu

int main(void) {
 
   h1<<<1, 1>>>(42); // expected-error {{kernel call to non-global function 'h1'}}
 
   int (*fp)(int) = h2;
   fp<<<1, 1>>>(42); // expected-error {{must have void return type}}
 
   g1<<<undeclared, 1>>>(42); // expected-error {{use of undeclared identifier 'undeclared'}}
 }
+
+// Make sure we can call static member kernels.
+template <typename > struct a0 {
+  template <typename T> static __global__ void Call(T);
+};
+struct a1 {
+  template <typename T> static __global__ void Call(T);
+};
+template <typename T> struct a2 {
+  static __global__ void Call(T);
+};
+struct a3 {
+  static __global__ void Call(int);
+  static __global__ void Call(void*);
+};
+
+struct b {
+  template <typename c> void d0(c arg) {
+    a0<c>::Call<<<0, 0>>>(arg);
+    a1::Call<<<0,0>>>(arg);
+    a2<c>::Call<<<0,0>>>(arg);
+    a3::Call<<<0, 0>>>(arg);
+  }
+  void d1(void* arg) {
+    a0<void*>::Call<<<0, 0>>>(arg);
+    a1::Call<<<0,0>>>(arg);
+    a2<void*>::Call<<<0,0>>>(arg);
+    a3::Call<<<0, 0>>>(arg);
+  }
+  void e() { d0(1); }
+};